Rodinia's streamcluster causes a kernel compiler crash

[Pekka Jääskeläinen]

build:

opencl/streamcluster$ make OPENCL_INC=$HOME/repo/pocl/include CC_FLAGS="-DCL_USE_DEPRECATED_OPENCL_1_1_APIS" OPENCL_LIB="$HOME/repo/pocl/lib/CL/.libs -lpthread -Wno-deprecated -Xlinker -rpath $HOME/repo/pocl/lib/CL/.libs -lOpenCL -lpthread" CXX_FLAGS="-Wno-deprecated"

modify 'gpu' in the 'run' script to 'cpu'.

./run crashes with a module verification issue with lots of "PHI node entries do not match predecessors!". Thus, the Phi-node fixing is (still) broken.

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

It has one of these "suspicious" kernels where there's a barrier guarded by an if with a thread_id check and the limit num is given as a kernel argument. That check is illegal unless it's guaranteed thread_id is always or never less than num. In this case the 'else' for that if is empty. However, the kernel compiler shouldn't of course crash in this case but assume the check is false or true for all thread_ids.

 /* block ID and global thread ID */
 const int thread_id = get_global_id(0);
 const int local_id = get_local_id(0);

 if(thread_id<num){
   // coordinate mapping of point[x] to shared mem
   if(local_id == 0)
     for(int i=0; i<dim; i++){
      coord_s[i] = coord_d[i*num + x];
     }
   barrier(CLK_LOCAL_MEM_FENCE);
         ...

[Carlos Sánchez de La Lama (csanchezdll)]

Seems to work, probably fixed as side-effect of some other fix. Closing the bug.