title: CUDA kernel seemingly ignoring "if" statement
body: <p>What follows is the part of my kernel that does not behave properly, then an explanation of what I've found while debugging.</p>
  
  <pre><code>__global__ void Mangler(float *matrix, int *map)
  {
      __shared__ signed int localMap[N];
  
      if(0 == threadIdx.x) 
      {
          for(int i=0; i&lt;N; i++) 
              localMap[i] = -1;
      }
  
      __syncthreads();
  
      int fn = ...; // a lot of code goes into this number, skipped for clarity
      int rnumber = threadIdx.x;
  
      int X = atomicCAS(&amp;localMap[fn], -1, rnumber); // Spot of bother 1
  
      if(X == -1) // Spot of bother 2
      {
          // some code
      }
      else 
      {
          // other code
      }
  }
  </code></pre>
  
  <p>I've found in the documentation that <code>atomicCAS(*address, compare, value)</code> basically returns (and saves to the given address) the result of <code>(old == compare ? value : old)</code>, where old is the value at the address before executing the function.</p>
  
  <p>Going with this, I believe that executing <code>int X = atomicCAS(&amp;localMap[fn], -1, rnumber);</code> should have two possible outcomes (according to NVidia Cuda C Programming Guide):</p>
  
  <ul>
  <li>if <code>localMap[fn] == -1</code> then <code>X</code> should have a value of <code>rnumber</code> and <code>localMap[fn]</code> should have a value of <code>rnumber</code>. <em>This does not happen.</em></li>
  <li>if <code>localMap[fn] != -1</code> then <code>X</code> should be set to the value of <code>localMap[fn]</code> and said value should be left intact.</li>
  </ul>
  
  <p>What happens instead, as debugging with NSight has shown me, is that <code>X</code> is being assigned -1, while <code>localMap[fn]</code> is being assigned the value of <code>rnumber</code>. I do not understand that, but as you can see in my code, I've changed the <code>if</code> to catch this situation.</p>
  
  <p>Which brings me to spot of bother number 2: though NSight shows the value of <code>X</code> as -1, the <code>if {}</code> is being completely skipped (no breakpoints within hit whatsoever) and execution jumps straight to <code>else</code>.</p>
  
  <p>My questions: </p>
  
  <ul>
  <li><s>Do I misunderstand <code>atomicCAS</code> completely?</s> <em>yes, I did</em></li>
  <li>What could cause and <code>if</code> which should evaluate as true to jump straight into <code>else</code> in device code?</li>
  </ul>
  
  <p>I'm using NVidia CUDA 5.5, Visual Studio 2012 x64 on Windows 8, NVidia Nsight Monitor Visual Studio Edition 3.1. The GPU for the machine is NVidia GeForce GTX 550 Ti.</p>
  
  <p>I've tried changing the syntax to <code>if(X!=-1)</code>; the true branch of the if is still not being executed.</p>