https://llvm.org/bugs/show_bug.cgi?id=30564
Bug ID: 30564
Summary: Optimizer bug causes infinite loop in cuda code
Product: clang
Version: trunk
Hardware: PC
OS: Linux
Status: NEW
Severity: normal
Priority: P
Component: CUDA
Assignee: unassignedclangb...@nondot.org
Reporter: crtr...@sandia.gov
CC: llvm-bugs@lists.llvm.org
Classification: Unclassified
Created attachment 17377
--> https://llvm.org/bugs/attachment.cgi?id=17377&action=edit
Reproducer
This bug is present in 3.9 and current head.
I have some code which essentially does a lock loop with a while construct to
implement arbitrary atomic operations. We use that in production using NVCC,
but the code infinitely loops at optimization level O2 or higher with clang. I
can workaround the bug by adding another check on the loop variable. Putting
print statements in shows the check is never triggered and behaviour of the
lock loop is as expected.
A reproducer is attached. The build.clang and build.cuda files have build
lines, please change paths so they work for you. Removing the -DBUG puts in the
workaround (i.e. the additional check and printf). Also change on the clang
build -O2 to -O1 as another way to make it work.
This is the offending function (the lock functions just do a simple hash and
try to grep a lock from a global array):
template < typename T >
__inline__ __device__
T atomic_fetch_add( volatile T * const dest , const T & val )
{
T return_val;
// This is a way to (hopefully) avoid dead lock in a warp
int done = 1;
while ( done>0 ) {
done++;
bool locked = lock_address_cuda_space( (void*) dest );
if( locked ) {
return_val = *dest;
*dest = return_val + val;
unlock_address_cuda_space( (void*) dest );
done = 0;
}
#ifndef BUG
printf("Done: %i %i %i\n",threadIdx.y,done,locked?1:0);
if(done==100) done = 0;
#endif
}
return return_val;
}
--
You are receiving this mail because:
You are on the CC list for the bug.
_______________________________________________
llvm-bugs mailing list
llvm-bugs@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-bugs
