title: CUDA racecheck, shared memory array and cudaDeviceSynchronize()
body: <p>I recently discovered the <strong>racecheck</strong> tool of <strong>cuda-memcheck</strong>, available in CUDA 5.0 (<code>cuda-memcheck --tool racecheck</code>, see the <a href="http://docs.nvidia.com/cuda/cuda-memcheck/index.html#using-racecheck" rel="nofollow">NVIDIA doc</a>). This tool can detect race conditions with shared memory in a CUDA kernel.</p>
  
  <p>In debug mode, this tool does not detect anything, which is apparently normal. However, in release mode (-O3), I get errors depending on the parameters of the problem.</p>
  
  <p>Here is an error example (initialization of shared memory on line 22, assignment on line 119):</p>
  
  <blockquote>
    <blockquote>
      <p>========= ERROR: Potential WAW hazard detected at <strong>shared</strong> 0x0 in block (35, 0, 0) :
      =========     Write Thread (32, 0, 0) at 0x00000890 in ....h:119:void kernel_test3(Data*)
      =========     Write Thread (0, 0, 0) at 0x00000048 in ....h:22:void kernel_test3(Data*)<br>
      =========     Current Value : 13, Incoming Value : 0</p>
    </blockquote>
  </blockquote>
  
  <ol>
  <li>The first thing that surprised me is the thread ids. When I first encountered the error, each block contained 32 threads (ids 0 to 31). So why is there a problem with the thread id 32? I even added an extra check on <code>threadIdx.x</code>, but this changed nothing.</li>
  <li>I use shared memory as a temporary buffer, and each thread deals with its own parameters of a multidimensional array, e.g. <code>__shared__ float arr[SIZE_1][SIZE_2][NB_THREADS_PER_BLOCK]</code>. I do not really understand how there could be any race conditions, since each thread deals with its own part of shared memory.</li>
  <li>Reducing the grid size from 64 blocks to 32 blocks seemed to solve the issue (with 32 threads per block). I do not understand why.</li>
  </ol>
  
  <p>In order to understand what was happening, I tested with some simpler kernels.
  Let me show you an example of a kernel that creates that kind of error. Basically, this kernel uses <code>SIZE_X*SIZE_Y*NTHREADS*sizeof(float)</code> B of shared memory, and I can use 48KB of shared memory per SM.</p>
  
  <p><strong>test.cu</strong></p>
  
  <pre><code>template &lt;unsigned int NTHREADS&gt;
  __global__ void kernel_test()
  {
      const int SIZE_X = 4;
      const int SIZE_Y = 4;
  
      __shared__ float tmp[SIZE_X][SIZE_Y][NTHREADS];
  
      for (unsigned int i = 0; i &lt; SIZE_X; i++)
          for (unsigned int j = 0; j &lt; SIZE_Y; j++)
              tmp[i][j][threadIdx.x] = threadIdx.x;
  }
  
  int main()
  {
    const unsigned int NTHREADS = 32;
  
    //kernel_test&lt;NTHREADS&gt;&lt;&lt;&lt;32, NTHREADS&gt;&gt;&gt;(); // ---&gt; works fine
    kernel_test&lt;NTHREADS&gt;&lt;&lt;&lt;64, NTHREADS&gt;&gt;&gt;();
  
    cudaDeviceSynchronize(); // ---&gt; gives racecheck errors if NBLOCKS &gt; 32
  }
  </code></pre>
  
  <p><strong>Compilation</strong>:</p>
  
  <p><code>nvcc test.cu --ptxas-options=-v -o test</code></p>
  
  <p><strong>If we run the kernel</strong>:</p>
  
  <p><code>cuda-memcheck --tool racecheck test</code></p>
  
  <ul>
  <li><code>kernel_test&lt;32&gt;&lt;&lt;&lt;32, 32&gt;&gt;&gt;();</code> : 32 blocks, 32 threads => does not lead to any apparent racecheck error.</li>
  <li><code>kernel_test&lt;32&gt;&lt;&lt;&lt;64, 32&gt;&gt;&gt;();</code> : 64 blocks, 32 threads => leads to WAW hazards (threadId.x = 32?!) and errors.</li>
  </ul>
  
  <blockquote>
    <blockquote>
      <p>========= ERROR: Potential WAW hazard detected at <strong>shared</strong> 0x6 in block (57, 0, 0) :<br>
      =========     Write Thread (0, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void)<br>
      =========     Write Thread (1, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void)<br>
      =========     Current Value : 0, Incoming Value : 128  </p>
      
      <p>========= INFO:(Identical data being written) Potential WAW hazard detected at <strong>shared</strong> 0x0 in block (47, 0, 0) :<br>
      =========     Write Thread (32, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void)<br>
      =========     Write Thread (0, 0, 0) at 0x00000048 in ....h:403:void kernel_test(void)<br>
      =========     Current Value : 0, Incoming Value : 0  </p>
    </blockquote>
  </blockquote>
  
  <p>So what am I missing here? Am I doing something wrong with shared memory? (I am still a beginner with this)</p>
  
  <h1>** UPDATE **</h1>
  
  <p>The problem seems to be coming from <code>cudaDeviceSynchronize()</code> when <code>NBLOCKS &gt; 32</code>. Why is this happening?</p>