title: CUDA:cudaDeviceSynchronize return error code 30
body: <p>The Code below is about Sha-1 algorithm.</p>
  
  <p>There is a big Loop in kernel(HashRounds=0x40000,262144 times).
  Now the situation is ：</p>
  
  <p>1.when running under debug mode,it will report error 30.but if i reduce the cycle time,like 50000 times,no error will be reported.</p>
  
  <p>2.when running under release mode,it is normal.but when i increase the thread scale,like block = 48,thread = 192, it will report the same problem.</p>
  
  <p>Environment:GTX560Ti+Win8+Visual Studio 2012+Cuda5.5</p>
  
  <p>Pleading for your help！！</p>
  
  <pre><code>#include &lt;stdio.h&gt;
  #include &lt;string.h&gt;
  #include "cuda_runtime.h"
  #include "device_launch_parameters.h"
  typedef unsigned int    uint32;
  typedef unsigned char   byte;
  typedef struct {
      uint32 state[5];
      uint32 count[2];
      byte buffer[64];
      byte workspace[64]; // Temporary buffer.
  } hash_context;
  
  #define rol(value, bits) (((value) &lt;&lt; (bits)) | ((value) &gt;&gt; (32 - (bits))))
  /* blk0() and blk() perform the initial expand. */
  /* I got the idea of expanding during the round function from SSLeay */
  
  #define blk0(i) (block-&gt;l[i] = (rol(block-&gt;l[i],24)&amp;0xFF00FF00) \
      |(rol(block-&gt;l[i],8)&amp;0x00FF00FF))
  
  #define blk(i) (block-&gt;l[i&amp;15] = rol(block-&gt;l[(i+13)&amp;15]^block-&gt;l[(i+8)&amp;15] \
      ^block-&gt;l[(i+2)&amp;15]^block-&gt;l[i&amp;15],1))
  
  /* (R0+R1), R2, R3, R4 are the different operations used in SHA1 */
  #define R0(v,w,x,y,z,i) {z+=((w&amp;(x^y))^y)+blk0(i)+0x5A827999+rol(v,5);w=rol(w,30);}
  #define R1(v,w,x,y,z,i) {z+=((w&amp;(x^y))^y)+blk(i)+0x5A827999+rol(v,5);w=rol(w,30);}
  #define R2(v,w,x,y,z,i) {z+=(w^x^y)+blk(i)+0x6ED9EBA1+rol(v,5);w=rol(w,30);}
  #define R3(v,w,x,y,z,i) {z+=(((w|x)&amp;y)|(w&amp;x))+blk(i)+0x8F1BBCDC+rol(v,5);w=rol(w,30);}
  #define R4(v,w,x,y,z,i) {z+=(w^x^y)+blk(i)+0xCA62C1D6+rol(v,5);w=rol(w,30);}
  
  cudaError_t addWithCuda();
  __global__ void cryptKernel();
  __device__ void hash_initial(hash_context* context);
  __device__ void hash_process( hash_context * context, byte * data, size_t len);
  __device__ void SHA1Transform(uint32 state[5], byte workspace[64], byte buffer[64]);
  
  int main()
  {
      cudaError_t cudaStatus = addWithCuda();
      if (cudaStatus != cudaSuccess) {
          fprintf(stderr, "addWithCuda failed!");
          return 1;
      }
      // cudaDeviceReset must be called before exiting in order for profiling and
      // tracing tools such as Nsight and Visual Profiler to show complete traces.
      cudaStatus = cudaDeviceReset();
      if (cudaStatus != cudaSuccess) {
          fprintf(stderr, "cudaDeviceReset failed!");
          return 1;
      }
      printf("over");
      getchar();
      return 0;
  }
  // Helper function for using CUDA to add vectors in parallel.
  cudaError_t addWithCuda()
  {
      cudaError_t cudaStatus;
      // Choose which GPU to run on, change this on a multi-GPU system.
      cudaStatus = cudaSetDevice(0);
      if (cudaStatus != cudaSuccess) {
          fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
          goto Error;
      }
      cryptKernel&lt;&lt;&lt;1,1&gt;&gt;&gt;();
      // Check for any errors launching the kernel
      cudaStatus = cudaGetLastError();
      if (cudaStatus != cudaSuccess) {
          fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
          goto Error;
      }
      // cudaDeviceSynchronize waits for the kernel to finish, and returns
      // any errors encountered during the launch.
      cudaStatus = cudaDeviceSynchronize();
      if (cudaStatus != cudaSuccess) {
          fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
          goto Error;
      }
  Error:
      return cudaStatus;
  }
  
  __global__ void cryptKernel()
  {
      int i = blockIdx.x*blockDim.x+threadIdx.x;
      byte RawPsw[24] = {'\0'};
      hash_context c;
      hash_initial(&amp;c);
      const int HashRounds=0x40000;
      for (int I=0;I&lt;HashRounds;I++)
      {
          hash_process( &amp;c, RawPsw, 24);
      }
  }
  
  __device__ void SHA1Transform(uint32 state[5], byte workspace[64], byte buffer[64])
  {
      uint32 a, b, c, d, e;
      typedef union {
          byte c[64];
          uint32 l[16];
      } CHAR64LONG16;
      CHAR64LONG16* block;
      block = (CHAR64LONG16*)buffer;
      /* Copy context-&gt;state[] to working vars */
      a = state[0];
      b = state[1];
      c = state[2];
      d = state[3];
      e = state[4];
      /* 4 rounds of 20 operations each. Loop unrolled. */
      R0(a,b,c,d,e, 0); R0(e,a,b,c,d, 1); R0(d,e,a,b,c, 2); R0(c,d,e,a,b, 3);
      R0(b,c,d,e,a, 4); R0(a,b,c,d,e, 5); R0(e,a,b,c,d, 6); R0(d,e,a,b,c, 7);
      R0(c,d,e,a,b, 8); R0(b,c,d,e,a, 9); R0(a,b,c,d,e,10); R0(e,a,b,c,d,11);
      R0(d,e,a,b,c,12); R0(c,d,e,a,b,13); R0(b,c,d,e,a,14); R0(a,b,c,d,e,15);
      R1(e,a,b,c,d,16); R1(d,e,a,b,c,17); R1(c,d,e,a,b,18); R1(b,c,d,e,a,19);
      R2(a,b,c,d,e,20); R2(e,a,b,c,d,21); R2(d,e,a,b,c,22); R2(c,d,e,a,b,23);
      R2(b,c,d,e,a,24); R2(a,b,c,d,e,25); R2(e,a,b,c,d,26); R2(d,e,a,b,c,27);
      R2(c,d,e,a,b,28); R2(b,c,d,e,a,29); R2(a,b,c,d,e,30); R2(e,a,b,c,d,31);
      R2(d,e,a,b,c,32); R2(c,d,e,a,b,33); R2(b,c,d,e,a,34); R2(a,b,c,d,e,35);
      R2(e,a,b,c,d,36); R2(d,e,a,b,c,37); R2(c,d,e,a,b,38); R2(b,c,d,e,a,39);
      R3(a,b,c,d,e,40); R3(e,a,b,c,d,41); R3(d,e,a,b,c,42); R3(c,d,e,a,b,43);
      R3(b,c,d,e,a,44); R3(a,b,c,d,e,45); R3(e,a,b,c,d,46); R3(d,e,a,b,c,47);
      R3(c,d,e,a,b,48); R3(b,c,d,e,a,49); R3(a,b,c,d,e,50); R3(e,a,b,c,d,51);
      R3(d,e,a,b,c,52); R3(c,d,e,a,b,53); R3(b,c,d,e,a,54); R3(a,b,c,d,e,55);
      R3(e,a,b,c,d,56); R3(d,e,a,b,c,57); R3(c,d,e,a,b,58); R3(b,c,d,e,a,59);
      R4(a,b,c,d,e,60); R4(e,a,b,c,d,61); R4(d,e,a,b,c,62); R4(c,d,e,a,b,63);
      R4(b,c,d,e,a,64); R4(a,b,c,d,e,65); R4(e,a,b,c,d,66); R4(d,e,a,b,c,67);
      R4(c,d,e,a,b,68); R4(b,c,d,e,a,69); R4(a,b,c,d,e,70); R4(e,a,b,c,d,71);
      R4(d,e,a,b,c,72); R4(c,d,e,a,b,73); R4(b,c,d,e,a,74); R4(a,b,c,d,e,75);
      R4(e,a,b,c,d,76); R4(d,e,a,b,c,77); R4(c,d,e,a,b,78); R4(b,c,d,e,a,79);
      /* Add the working vars back into context.state[] */
      state[0] += a;
      state[1] += b;
      state[2] += c;
      state[3] += d;
      state[4] += e;
  }
  __device__ void hash_initial(hash_context* context)
  {
      /* SHA1 initialization constants */
      context-&gt;state[0] = 0x67452301;
      context-&gt;state[1] = 0xEFCDAB89;
      context-&gt;state[2] = 0x98BADCFE;
      context-&gt;state[3] = 0x10325476;
      context-&gt;state[4] = 0xC3D2E1F0;
      context-&gt;count[0] = context-&gt;count[1] = 0;
  }
  /* Run your data through this. */
  __device__ void hash_process( hash_context * context, byte * data, size_t len)
  {
      unsigned int i, j;
      uint32 blen = ((uint32)len)&lt;&lt;3;
  
      j = (context-&gt;count[0] &gt;&gt; 3) &amp; 63;
      if ((context-&gt;count[0] += blen) &lt; blen ) context-&gt;count[1]++;
      context-&gt;count[1] += (uint32)(len &gt;&gt; 29);
      if ((j + len) &gt; 63) {
          memcpy(&amp;context-&gt;buffer[j], data, (i = 64-j));
          SHA1Transform(context-&gt;state, context-&gt;workspace, context-&gt;buffer);
          for ( ; i + 63 &lt; len; i += 64) {
  
              SHA1Transform(context-&gt;state, context-&gt;workspace, &amp;data[i]);
  
          }
          j = 0;
      }
      else i = 0;
      if (len &gt; i)
        memcpy(&amp;context-&gt;buffer[j], &amp;data[i], len - i);
  }
  </code></pre>