Rodinia's streamcluster causes a kernel compiler crash
Bug #966214 reported by
Pekka Jääskeläinen
This bug affects 1 person
Affects | Status | Importance | Assigned to | Milestone | |
---|---|---|---|---|---|
pocl |
Invalid
|
Critical
|
Carlos Sánchez de La Lama |
Bug Description
build:
opencl/
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.
Changed in pocl: | |
milestone: | none → 0.6 |
Changed in pocl: | |
assignee: | nobody → Carlos Sánchez de La Lama (csanchezdll) |
To post a comment you must log in.
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){ CLK_LOCAL_ MEM_FENCE) ;
// 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(
...