[llvm-bugs] [Bug 30564] New: Optimizer bug causes infinite loop in cuda code

[via llvm-bugs]

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: unassignedclangbugs at nondot.org
          Reporter: crtrott at sandia.gov
                CC: llvm-bugs at 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.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-bugs/attachments/20160929/96b69b98/attachment.html>