I have been battling with atomics for past 2 days. I am trying to implement a simple thread-level spinlock on CUDA and am hitting roadblocks.
So, Here is the crux of my issue. Testing this requires compute 1.2 hardware (as it uses smem atomics).
The following code hangs plain on TESLA C1060. The kernel is called with 1 block, 32 threads.
ANd, if I replace the shared memory lock with “g_lock” – it all works fine.
My guess : The “totalParticipation” shared memory variable is NOT seen correctly by threads although it is volatile and proper “threadfence” and “__syncthreads” are done to avoid compiler optimization.
However, if I use g_lock, everything works.
Can some1 confirm this on your hardware and tell me what the solution is???
Can some1 tell me what is the mistake I am making in this code?
#include <stdio.h>
__device__ int result = -123;
__device__ int g_lock=-1;
__global__ void checkKernel(void)
{
__shared__ int sharedLock;
__shared__ volatile int totalParticipation;
if (threadIdx.x == 0)
{
atomicExch(&sharedLock,-1);
totalParticipation = 0;
__threadfence();
}
__syncthreads();
for(int i=0; totalParticipation<5; i++)
{
if(atomicCAS(&sharedLock, -1, (int)threadIdx.x) == -1)
{
totalParticipation++;
__threadfence();
atomicExch(&sharedLock, -1);
}
__threadfence();
__syncthreads();
}
if (threadIdx.x == 0)
{
result = totalParticipation;
}
__syncthreads();
return;
}
int main()
{
cudaError_t err;
int host=-123;
checkKernel<<< 1, 32>>>();
err = cudaThreadSynchronize();
if (err != 0)
{
printf("Error launching kernel, err = %d\n",err);
return -1;
}
cudaMemcpyFromSymbol(&host, "result", sizeof(int), 0, cudaMemcpyDeviceToHost);
printf("host = %d\n", host);
}