Crash with a for-loop with a variable iteration count and a barrier

Bug #938883 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

this is the regression test for-with-var-iteration-count

"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;\n"
" float sum = 0.0f;\n"
" for (i = 0; i < a; ++i) {\n"
" sum += input[i]; \n"
" barrier(CLK_GLOBAL_MEM_FENCE);\n"
" }\n"
" result[gid] = sum;\n"
"}\n";

This causes a branch from the first block to the exit block of the kernel function because a can be zero which causes the whole for-loop to be jumped over. That branch is not reproduced by pocl thus causing a crash because of invalid phi-nodes in the block preceeding the exit:

PHINode should have one entry for each predecessor of its parent basic block!
  %sum.0.lcssa = phi i32 [ %phitmp, %for.cond.for.end_crit_edge.wi_0_0_0 ], [ 0, %entry.wi_0_0_0 ]
PHINode should have one entry for each predecessor of its parent basic block!
  %sum.0.lcssa.wi_1_0_0 = phi i32 [ %phitmp, %for.cond.for.end_crit_edge.wi_0_0_0 ], [ 0, %for.end.wi_0_0_0 ]
PHINode should have one entry for each predecessor of its parent basic block!
  %14 = phi i32 [ %phitmp, %for.cond.for.end_crit_edge.wi_0_0_0 ], [ 0, %entry.barrier ]
PHINode should have one entry for each predecessor of its parent basic block!
  %18 = phi i32 [ %phitmp.wi_1_0_0, %for.end.btr.wi_0_0_0 ], [ 0, %entry.barrier ]
Broken module found, compilation aborted!

Related branches

Revision history for this message
Pekka Jääskeläinen (pekka-jaaskelainen) wrote :

This still crashes but in another position. Checking...

Revision history for this message
Pekka Jääskeläinen (pekka-jaaskelainen) wrote :

The test case passes now.

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.