Multi-level (2) for-loop with a barrier inside crashes

[Pekka Jääskeläinen]

This can be reproduced in tests/regression/multi-level-loops-with-barriers.

Its kernel:

"kernel \n"
"void test_kernel(__global float *input, \n"
" __global int *result,\n"
" int a) {\n"
" int gid = get_global_id(0);\n"
" int i, j;\n"
" for (i = 0; i < 32; ++i) {\n"
" result[gid] = input[gid];\n"
" for (j = 0; j < 32; ++j) {\n"
" result[gid] = input[gid] * input[gid + j];\n"
" barrier(CLK_GLOBAL_MEM_FENCE);\n"
" }\n"
" }\n"
"}\n";

Crashes the kernel compiler with:

opt: WorkitemReplication.cc:134: bool pocl::WorkitemReplication::ProcessFunction(llvm::Function&): Assertion `(non_backedge_predecessors == 1) && "Could not determine parallel region entry!"' failed.

[Pekka Jääskeläinen (pekka-jaaskelainen)]

This case actually crashed because of computation after the loop exit block's "loop back" branch, not due to the multi-level for-loop itself.

[Pekka Jääskeläinen (pekka-jaaskelainen)]

OK, seems my verification code was disabled ;) It doesn't crash but miscomputed. The inner loop is not rolling correctly.