HIP
HIP copied to clipboard
ROCm
updates to shared memory variable not seen by all threads
Around line 190 in tmpRadixSort5.h, there is an update to shared variable ibs, which is the control variable in a while loop. It appears the update to ibs is not seen by all threads.
Using ibs - either by printing it out, or by assigning it to another shared variable - will fix the problem, and the program runs correctly.
__shared__ int ibs;
__shared__ int ibs2;
...
ibs = size - 1;
__syncthreads();
while (__syncthreads_and(ibs>0)) {
...
if (threadIdx.x == 0) {
ibs -= sb;
// uncomment the next line to fix the hang or error
//ibs2 = ibs;
}
__syncthreads();
}
More info in Readme. Build command in "compile.sh" reproduce_sync_failure.tar.gz
One thread may be needed to update ibs.
// broadcast
if (threadIdx.x == 0) ibs = size - 1;
