title: CUDA program causes nvidia driver to crash
body: <p>My monte carlo pi calculation CUDA program is causing my nvidia driver to crash when I exceed around 500 trials and 256 full blocks. It seems to be happening in the monteCarlo kernel function.Any help is appreciated.</p>
  
  <pre><code>#include &lt;stdio.h&gt;
  #include &lt;stdlib.h&gt;
  #include &lt;cuda.h&gt;
  #include &lt;curand.h&gt;
  #include &lt;curand_kernel.h&gt;
  
  
  #define NUM_THREAD 256
  #define NUM_BLOCK 256
  
  
  
  ///////////////////////////////////////////////////////////////////////////////////////////
  ///////////////////////////////////////////////////////////////////////////////////////////
  
  // Function to sum an array
  __global__ void reduce0(float *g_odata) {
  extern __shared__ int sdata[];
  
  // each thread loads one element from global to shared mem
  unsigned int tid = threadIdx.x;
  unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
  sdata[tid] = g_odata[i];
  __syncthreads();
  
  // do reduction in shared mem
  for (unsigned int s=1; s &lt; blockDim.x; s *= 2) { // step = s x 2
      if (tid % (2*s) == 0) { // only threadIDs divisible by the step participate
          sdata[tid] += sdata[tid + s];
      }
      __syncthreads();
  }
  
  // write result for this block to global mem
  if (tid == 0) g_odata[blockIdx.x] = sdata[0];
  }
  
  ///////////////////////////////////////////////////////////////////////////////////////////
  ///////////////////////////////////////////////////////////////////////////////////////////
  __global__ void monteCarlo(float *g_odata, int  trials, curandState *states){
  //  unsigned int tid = threadIdx.x;
      unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
      unsigned int incircle, k;
      float x, y, z;
      incircle = 0;
  
      curand_init(1234, i, 0, &amp;states[i]);
  
      for(k = 0; k &lt; trials; k++){
          x = curand_uniform(&amp;states[i]);
          y = curand_uniform(&amp;states[i]);
          z =(x*x + y*y);
          if (z &lt;= 1.0f) incircle++;
      }
      __syncthreads();
      g_odata[i] = incircle;
  }
  ///////////////////////////////////////////////////////////////////////////////////////////
  ///////////////////////////////////////////////////////////////////////////////////////////
  int main() {
  
      float* solution = (float*)calloc(100, sizeof(float));
      float *sumDev, *sumHost, total;
      const char *error;
      int trials; 
      curandState *devStates;
  
      trials = 500;
      total = trials*NUM_THREAD*NUM_BLOCK;
  
      dim3 dimGrid(NUM_BLOCK,1,1); // Grid dimensions
      dim3 dimBlock(NUM_THREAD,1,1); // Block dimensions
      size_t size = NUM_BLOCK*NUM_THREAD*sizeof(float); //Array memory size
      sumHost = (float*)calloc(NUM_BLOCK*NUM_THREAD, sizeof(float));
  
      cudaMalloc((void **) &amp;sumDev, size); // Allocate array on device
      error = cudaGetErrorString(cudaGetLastError());
      printf("%s\n", error);
  
  
      cudaMalloc((void **) &amp;devStates, (NUM_THREAD*NUM_BLOCK)*sizeof(curandState));
      error = cudaGetErrorString(cudaGetLastError());
      printf("%s\n", error);
  
  
      // Do calculation on device by calling CUDA kernel
      monteCarlo &lt;&lt;&lt;dimGrid, dimBlock&gt;&gt;&gt; (sumDev, trials, devStates);
      error = cudaGetErrorString(cudaGetLastError());
      printf("%s\n", error);
  
          // call reduction function to sum
      reduce0 &lt;&lt;&lt;dimGrid, dimBlock, (NUM_THREAD*sizeof(float))&gt;&gt;&gt; (sumDev);
      error = cudaGetErrorString(cudaGetLastError());
      printf("%s\n", error);
  
      dim3 dimGrid1(1,1,1);
      dim3 dimBlock1(256,1,1);
      reduce0 &lt;&lt;&lt;dimGrid1, dimBlock1, (NUM_THREAD*sizeof(float))&gt;&gt;&gt; (sumDev);
      error = cudaGetErrorString(cudaGetLastError());
      printf("%s\n", error);
  
      // Retrieve result from device and store it in host array
      cudaMemcpy(sumHost, sumDev, sizeof(float), cudaMemcpyDeviceToHost);
      error = cudaGetErrorString(cudaGetLastError());
      printf("%s\n", error);
  
  
      *solution = 4*(sumHost[0]/total);
      printf("%.*f\n", 1000, *solution);
      free (solution);
      free(sumHost);
      cudaFree(sumDev);
      cudaFree(devStates);
      //*solution = NULL;
      return 0;
  }
  </code></pre>