<html>
<head>
<base href="https://llvm.org/bugs/" />
</head>
<body><table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Bug ID</th>
<td><a class="bz_bug_link
bz_status_NEW "
title="NEW --- - Optimizer bug causes infinite loop in cuda code"
href="https://llvm.org/bugs/show_bug.cgi?id=30564">30564</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>Optimizer bug causes infinite loop in cuda code
</td>
</tr>
<tr>
<th>Product</th>
<td>clang
</td>
</tr>
<tr>
<th>Version</th>
<td>trunk
</td>
</tr>
<tr>
<th>Hardware</th>
<td>PC
</td>
</tr>
<tr>
<th>OS</th>
<td>Linux
</td>
</tr>
<tr>
<th>Status</th>
<td>NEW
</td>
</tr>
<tr>
<th>Severity</th>
<td>normal
</td>
</tr>
<tr>
<th>Priority</th>
<td>P
</td>
</tr>
<tr>
<th>Component</th>
<td>CUDA
</td>
</tr>
<tr>
<th>Assignee</th>
<td>unassignedclangbugs@nondot.org
</td>
</tr>
<tr>
<th>Reporter</th>
<td>crtrott@sandia.gov
</td>
</tr>
<tr>
<th>CC</th>
<td>llvm-bugs@lists.llvm.org
</td>
</tr>
<tr>
<th>Classification</th>
<td>Unclassified
</td>
</tr></table>
<p>
<div>
<pre>Created <span class=""><a href="attachment.cgi?id=17377" name="attach_17377" title="Reproducer">attachment 17377</a> <a href="attachment.cgi?id=17377&action=edit" title="Reproducer">[details]</a></span>
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;
}</pre>
</div>
</p>
<hr>
<span>You are receiving this mail because:</span>
<ul>
<li>You are on the CC list for the bug.</li>
</ul>
</body>
</html>