dace icon indicating copy to clipboard operation
dace copied to clipboard

Codegen: Invalid control flow in generated code for BFS sample

Open phschaad opened this issue 4 years ago • 0 comments

Description The generated code of the included BFS sample contains invalid control flow, where the else clause of a conditional is always taken due to a missing else statement.

Steps to reproduce:

  1. Run the BFS sample (python samples/graph/bfs_bsp.py)
  2. Computes the wrong output -> high error/diff in most cases
  3. Generated code in .daceche/bfs_bsp/src/cpu/bfs_bsp.cpp contains wrong control flow.

Generated code:

/* DaCe AUTO-GENERATED FILE. DO NOT MODIFY */
#include <dace/dace.h>

struct bfs_bsp_t {

};

void __program_bfs_bsp_internal(bfs_bsp_t *__state, unsigned int * __restrict__ G_col, unsigned int * __restrict__ G_row, unsigned int * __restrict__ depth, int E, int V, unsigned int srcnode)
{
    unsigned int fsz0;
    unsigned int *frontier;
    frontier = new unsigned int DACE_ALIGN(64)[V];
    unsigned int *frontier2;
    frontier2 = new unsigned int DACE_ALIGN(64)[V];
    unsigned int fsz1;
    long long d;

    {


        {
            unsigned int d;
            unsigned int f;
            unsigned int ofsz;

            ///////////////////
            // Tasklet code (init_depth)
            d = 0;
            f = srcnode;
            ofsz = 1;
            ///////////////////

            depth[srcnode] = d;
            frontier[0] = f;
            fsz0 = ofsz;
        }

    }
    d = 1;
    __state_0_reset1:;
    {


        {
            unsigned int ofsz;

            ///////////////////
            // Tasklet code (reset_1)
            ofsz = 0;
            ///////////////////

            fsz1 = ofsz;
        }

    }
    {
        dace::ArrayStreamView<unsigned int> out_stream0 (frontier2);


        {
            unsigned int fsz0 = fsz0;
            #pragma omp parallel for
            for (auto f = 0; f < fsz0; f += 1) {
                unsigned int rowb0;
                unsigned int rowe0;
                {
                    unsigned int in_f = frontier[f];
                    unsigned int* in_row = &G_row[0];
                    unsigned int ob;
                    unsigned int oe;

                    ///////////////////
                    // Tasklet code (indirection)
                    ob = in_row[in_f];
                    oe = in_row[(in_f + 1)];
                    ///////////////////

                    rowb0 = ob;
                    rowe0 = oe;
                }
                {
                    unsigned int rowb0 = rowb0;
                    unsigned int rowe0 = rowe0;
                    for (auto nid = rowb0; nid < rowe0; nid += 1) {
                        {
                            unsigned int* const &in_col = &G_col[0];
                            unsigned int* const &in_depth = &depth[0];
                            unsigned int* out_depth = depth;
                            unsigned int out_fsz;

                            ///////////////////
                            // Tasklet code (update_and_push)
                            auto node = in_col[nid];
                            auto dep = in_depth[node];
                            if ((d < dep)) {out_depth[node] = d;out_stream0.push(node);
                                out_fsz = dace::wcr_fixed<dace::ReductionType::Sum, unsigned int>::reduce_atomic(&fsz1, 1);
                            }
                            ///////////////////

                        }
                    }
                }
            }
        }

    }
    if ((fsz1 > 0)) {
    }
    goto __state_exit_0;
    d = (d + 1);
    {


        {
            unsigned int ofsz;

            ///////////////////
            // Tasklet code (reset_0)
            ofsz = 0;
            ///////////////////

            fsz0 = ofsz;
        }

    }
    {
        dace::ArrayStreamView<unsigned int> out_stream1 (frontier);


        {
            unsigned int fsz1 = fsz1;
            #pragma omp parallel for
            for (auto f = 0; f < fsz1; f += 1) {
                unsigned int rowb1;
                unsigned int rowe1;
                {
                    unsigned int in_f = frontier2[f];
                    unsigned int* in_row = &G_row[0];
                    unsigned int ob;
                    unsigned int oe;

                    ///////////////////
                    // Tasklet code (indirection)
                    ob = in_row[in_f];
                    oe = in_row[(in_f + 1)];
                    ///////////////////

                    rowb1 = ob;
                    rowe1 = oe;
                }
                {
                    unsigned int rowb1 = rowb1;
                    unsigned int rowe1 = rowe1;
                    for (auto nid = rowb1; nid < rowe1; nid += 1) {
                        {
                            unsigned int* const &in_col = &G_col[0];
                            unsigned int* const &in_depth = &depth[0];
                            unsigned int* out_depth = depth;
                            unsigned int out_fsz;

                            ///////////////////
                            // Tasklet code (update_and_push)
                            auto node = in_col[nid];
                            auto dep = in_depth[node];
                            if ((d < dep)) {out_depth[node] = d;out_stream1.push(node);
                                out_fsz = dace::wcr_fixed<dace::ReductionType::Sum, unsigned int>::reduce_atomic(&fsz0, 1);
                            }
                            ///////////////////

                        }
                    }
                }
            }
        }

    }
    if ((fsz0 > 0)) {
    }
    goto __state_exit_0;
    d = (d + 1);
    goto __state_0_reset1;
    __state_exit_0:;
    delete[] frontier;
    delete[] frontier2;
}

DACE_EXPORTED void __program_bfs_bsp(bfs_bsp_t *__state, unsigned int * __restrict__ G_col, unsigned int * __restrict__ G_row, unsigned int * __restrict__ depth, int E, int V, unsigned int srcnode)
{
    __program_bfs_bsp_internal(__state, G_col, G_row, depth, E, V, srcnode);
}

DACE_EXPORTED bfs_bsp_t *__dace_init_bfs_bsp(int E, int V, unsigned int srcnode)
{
    int __result = 0;
    bfs_bsp_t *__state = new bfs_bsp_t;



    if (__result) {
        delete __state;
        return nullptr;
    }
    return __state;
}

DACE_EXPORTED void __dace_exit_bfs_bsp(bfs_bsp_t *__state)
{
    delete __state;
}

Erroneous control flow: Screenshot 2021-01-21 145517

Expected control flow: Screenshot 2021-01-21 145517

Expected behavior: Sample runs with an error/diff of 0.0.

Environment:

  • OS: Windows 10 with WSL2 (Ubuntu20.04)

phschaad avatar Jan 21 '21 14:01 phschaad