Miscompile with a kernel with id-dependent computation before exit and also early exit

Bug #940549 reported by Pekka Jääskeläinen
6
This bug affects 1 person
Affects Status Importance Assigned to Milestone
pocl
Fix Released
Critical
Pekka Jääskeläinen

Bug Description

These curious cases keep popping up ;)

"kernel \n"
"void test_kernel(__global float *input, \n"
" __global int *result) {\n"
" size_t gid = get_global_id(0);\n"
" if (input[0] > 2) return;\n"
" result[gid] = 0;\n"
" barrier(CLK_GLOBAL_MEM_FENCE);\n"
" result[gid] += input[gid];\n"
" if (gid == 0) {\n"
" result[gid] = 42;\n"
" }\n"
"}\n";

Produces:
F(1: 42 != 1)

While only gid == 0 should have '42' so the if is always executed.

Hard to say what is the actual cause of this case but at least it has 3 entries to the exit node and there is thread id dependent computation in the last block so it can be related to the tail replication somehow.

Related branches

Changed in pocl:
status: New → Fix Released
To post a comment you must log in.
This report contains Public information  
Everyone can see this information.

Other bug subscribers

Remote bug watches

Bug watches keep track of this bug in other bug trackers.