title: CUDA Unable to see shared memory values in Nsight debugging
body: <p>I've been struggling for some time a problem I can't seem to find a solution to.
  The problem is that when I try to debug my CUDA code using Nvidia Nsight under Visual Studio 2008 I get strange results when using shared memory.</p>
  
  <p>My code is:</p>
  
  <pre><code>template&lt;typename T&gt;
  __device__
  T integrate()
  {
     extern __shared__ T s_test[]; // Dynamically allocated shared memory
     /**** Breakpoint (1) here ****/
     int index = threadIdx.x + threadIdx.y * blockDim.x; // Local index in block. Column major ordering
     if(index &lt; 64 &amp;&amp; blockIdx.x==0) { // Only work on a few values. Just testing
        s_test[index] = (T)index;
        /* Some other irelevant code here */
     }
     return v;
  }
  </code></pre>
  
  <p>When I reach breakpoint 1 and inspect the shared memory inside Visual Studio Watch window only the first 8 values of the array change and the others remain null. I would expect all first 64 to do so.
  <img src="https://i.stack.imgur.com/VNIVK.png" alt="Watch window from Visual Studio"></p>
  
  <p>I thought it might have something to do with all warps not executing simultaneously. So I tried synchronizing them. I added this code inside <code>integrate()</code></p>
  
  <pre><code>template&lt;typename T&gt;
  __device__
  T integrate()
  {
     /* Old code is still here */
  
     __syncthreads();
     /**** Breakpoint (2) here ****/
     if(index &lt; 64 &amp;&amp; blockIdx.x==0) {
        T tmp = s_test[index]; // Write to tmp variable so I can inspect it inside Nsight Watch window
        v = tmp + index; // Use `tmp` and `index` somehow so that the compiler doesn't optimize it out of existence
     }
  return v;
  }
  </code></pre>
  
  <p>But the problem is still there. Furthermore the rest of the values inside tmp are not <code>0</code> as the Watch window form VS is indicating. 
  <img src="https://i.stack.imgur.com/ZGzal.png" alt="Watch window from Nsight"></p>
  
  <p>I must mention that it takes a lot of steps to step over <code>__syncthreads()</code>, so when I reach it I just jump to breakpoint 2. What the heck is going on!?</p>
  
  <hr>
  
  <p><strong>EDIT</strong> Information about the system/launch configuration</p>
  
  <p><strong>System</strong></p>
  
  <ul>
  <li>Name Intel(R) Core(TM)2 Duo CPU     E7300  @ 2.66GHz</li>
  <li>Architecture x86</li>
  <li>Frequency    2.666 MHz</li>
  <li>Number of Cores  2</li>
  <li>Page Size    4.096</li>
  <li>Total Physical Memory    3.582,00 MB</li>
  <li>Available Physical Memory    1.983,00 MB</li>
  <li>Version Name Windows 7 Ultimate</li>
  <li>Version Number   6.1.7600</li>
  </ul>
  
  <p><strong>Device</strong> GeForce 9500 GT</p>
  
  <ul>
  <li>Driver Version   301.42</li>
  <li>Driver Model WDDM</li>
  <li>CUDA Device Index    0</li>
  <li>GPU Family   G96</li>
  <li>Compute Capability   1.1</li>
  <li>Number of SMs    4</li>
  <li>Frame Buffer Physical Size (MB)  512</li>
  <li>Frame Buffer Bandwidth (GB/s)    16</li>
  <li>Frame Buffer Bus Width (bits)    128</li>
  <li>Frame Buffer Location    Dedicated</li>
  <li>Graphics Clock (Mhz) 812</li>
  <li>Memory Clock (Mhz)   500</li>
  <li>Processor Clock (Mhz)    1625</li>
  <li>RAM Type DDR2</li>
  </ul>
  
  <p><strong>IDE</strong></p>
  
  <ul>
  <li>Microsoft Visual Studio Team System 2008</li>
  <li>NVIDIA Nsight Visual Studio Edition, Version 2.2 Build No. 2.2.0.12255</li>
  </ul>
  
  <p><strong>Compiler</strong> comands</p>
  
  <p>1> "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\bin\nvcc.exe"  -G  -gencode=arch=compute_10,code=\"sm_10,compute_10\"   --machine 32 -ccbin "C:\Program Files\Microsoft Visual Studio 9.0\VC\bin" -D_NEXUS_DEBUG -g  -D_DEBUG -Xcompiler "/EHsc /W3 /nologo /Od /Zi /RTC1 /MDd  " -I"inc" -I"C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v4.2\include" -maxrregcount=0  --compile -o "Debug/process_f2f.cu.obj" process_f2f.cu </p>
  
  <p><strong>Launch configuration</strong>. The shared memory size and doesn't seem to matter. I've tried several versions. The one I've worked with the most is:</p>
  
  <ul>
  <li>Shared memory 2048 Bytes </li>
  <li>Grid/block sizes : {101, 101, 1} , {16, 16, 1}</li>
  </ul>