title: cudamemcpy error:"the launch timed out and was terminated"
body: <p>My code is a parallel implmentation that calculates the nth digit of pi. When I finish the kernel and try to copy the memory back to the host I get a "the launch timed out and was terminated" error. 
  I used this code for error checking for each cudamalloc, cudamemcpy, and kernal launch.</p>
  
  <pre><code>std::string error = cudaGetErrorString(cudaGetLastError());
  printf("%s\n", error);
  </code></pre>
  
  <p>These calls were saying everything was fine until the first cudamemcpy call after returning from the kernel. the error happens in the line "cudaMemcpy(avhost, avdev, size, cudaMemcpyDeviceToHost);" in main. Any help is appreciated.</p>
  
  <pre><code>#include &lt;stdlib.h&gt;
  #include &lt;stdio.h&gt;
  #include &lt;math.h&gt;
  
  #define mul_mod(a,b,m) fmod( (double) a * (double) b, m)
  ///////////////////////////////////////////////////////////////////////////////////////
  ///////////////////////////////////////////////////////////////////////////////////////
  /* return the inverse of x mod y */
  __device__ int inv_mod(int x,int y) {
    int q,u,v,a,c,t;
  
    u=x;
    v=y;
    c=1;
    a=0;
    do {
      q=v/u;
  
      t=c;
      c=a-q*c;
      a=t;
  
      t=u;
      u=v-q*u;
      v=t;
    } while (u!=0);
    a=a%y;
    if (a&lt;0) a=y+a;
    return a;
  }
  ///////////////////////////////////////////////////////////////////////////////////////
  ///////////////////////////////////////////////////////////////////////////////////////
  /* return the inverse of u mod v, if v is odd */
  __device__ int inv_mod2(int u,int v) {
    int u1,u3,v1,v3,t1,t3;
  
    u1=1;
    u3=u;
  
    v1=v;
    v3=v;
  
    if ((u&amp;1)!=0) {
      t1=0;
      t3=-v;
      goto Y4;
    } else {
      t1=1;
      t3=u;
    }
  
    do {
  
      do {
        if ((t1&amp;1)==0) {
      t1=t1&gt;&gt;1;
      t3=t3&gt;&gt;1;
        } else {
      t1=(t1+v)&gt;&gt;1;
      t3=t3&gt;&gt;1;
        }
        Y4:;
      } while ((t3&amp;1)==0);
  
      if (t3&gt;=0) {
        u1=t1;
        u3=t3;
      } else {
        v1=v-t1;
        v3=-t3;
      }
      t1=u1-v1;
      t3=u3-v3;
      if (t1&lt;0) {
        t1=t1+v;
      }
    } while (t3 != 0);
    return u1;
  }
  
  
  /* return (a^b) mod m */
  __device__ int pow_mod(int a,int b,int m)
  {
    int r,aa;
  
    r=1;
    aa=a;
    while (1) {
      if (b&amp;1) r=mul_mod(r,aa,m);
      b=b&gt;&gt;1;
      if (b == 0) break;
      aa=mul_mod(aa,aa,m);
    }
    return r;
  }
  ///////////////////////////////////////////////////////////////////////////////////////
  ///////////////////////////////////////////////////////////////////////////////////////
  /* return true if n is prime */
  int is_prime(int n)
  {
     int r,i;
     if ((n % 2) == 0) return 0;
  
     r=(int)(sqrtf(n));
     for(i=3;i&lt;=r;i+=2) if ((n % i) == 0) return 0;
     return 1;
  }
  ///////////////////////////////////////////////////////////////////////////////////////
  ///////////////////////////////////////////////////////////////////////////////////////
  /* return the prime number immediatly after n */
  int next_prime(int n)
  {
     do {
        n++;
     } while (!is_prime(n));
     return n;
  }
  ///////////////////////////////////////////////////////////////////////////////////////
  ///////////////////////////////////////////////////////////////////////////////////////
  #define DIVN(t,a,v,vinc,kq,kqinc)       \
  {                       \
    kq+=kqinc;                    \
    if (kq &gt;= a) {                \
      do { kq-=a; } while (kq&gt;=a);        \
      if (kq == 0) {              \
        do {                  \
      t=t/a;                  \
      v+=vinc;                \
        } while ((t % a) == 0);           \
      }                       \
    }                     \
  }
  
  ///////////////////////////////////////////////////////////////////////////////////////
  ///////////////////////////////////////////////////////////////////////////////////////
  
  __global__ void digi_calc(int *s, int *av, int *primes, int N, int n, int nthreads){
      int a,vmax,num,den,k,kq1,kq2,kq3,kq4,t,v,i,t1, h;
      unsigned int tid = blockIdx.x*blockDim.x + threadIdx.x;
  // GIANT LOOP
      for (h = 0; h&lt;1; h++){
      if(tid &gt; nthreads) continue;
      a = primes[tid];
      vmax=(int)(logf(3*N)/logf(a));
      if (a==2) {
        vmax=vmax+(N-n);
        if (vmax&lt;=0) continue;
      }
      av[tid]=1;
      for(i=0;i&lt;vmax;i++) av[tid]*= a;
  
      s[tid]=0;
      den=1;
      kq1=0;
      kq2=-1;
      kq3=-3;
      kq4=-2;
      if (a==2) {
        num=1;
        v=-n; 
      } else {
        num=pow_mod(2,n,av[tid]);
        v=0;
      }
  
      for(k=1;k&lt;=N;k++) {
  
        t=2*k;
        DIVN(t,a,v,-1,kq1,2);
        num=mul_mod(num,t,av[tid]);
  
        t=2*k-1;
        DIVN(t,a,v,-1,kq2,2);
        num=mul_mod(num,t,av[tid]);
  
        t=3*(3*k-1);
        DIVN(t,a,v,1,kq3,9);
        den=mul_mod(den,t,av[tid]);
  
        t=(3*k-2);
        DIVN(t,a,v,1,kq4,3);
        if (a!=2) t=t*2; else v++;
        den=mul_mod(den,t,av[tid]);
  
        if (v &gt; 0) {
      if (a!=2) t=inv_mod2(den,av[tid]);
      else t=inv_mod(den,av[tid]);
      t=mul_mod(t,num,av[tid]);
      for(i=v;i&lt;vmax;i++) t=mul_mod(t,a,av[tid]);
      t1=(25*k-3);                                                                                                                                                                                                                                                                                                                                                                       
      t=mul_mod(t,t1,av[tid]);
      s[tid]+=t;
      if (s[tid]&gt;=av[tid]) s-=av[tid];
        }
      }
  
      t=pow_mod(5,n-1,av[tid]);
      s[tid]=mul_mod(s[tid],t,av[tid]);
      }
      __syncthreads();
  }
  ///////////////////////////////////////////////////////////////////////////////////////
  ///////////////////////////////////////////////////////////////////////////////////////
  int main(int argc,char *argv[])
  {
    int N,n,i,totalp, h;
    double sum;
    const char *error;
    int *sdev, *avdev, *shost, *avhost, *adev, *ahost;
      argc = 2;
      argv[1] = "2";
    if (argc&lt;2 || (n=atoi(argv[1])) &lt;= 0) {
      printf("This program computes the n'th decimal digit of pi\n"
         "usage: pi n , where n is the digit you want\n"
         );
      exit(1);
    }
      sum = 0;
      N=(int)((n+20)*logf(10)/logf(13.5));
      totalp=(N/logf(N))+10;
      ahost = (int *)calloc(totalp, sizeof(int));
      i = 0;
      ahost[0]=2;
      for(i=1; ahost[i-1]&lt;=(3*N); ahost[i+1]=next_prime(ahost[i])){
          i++;
      }
      // allocate host memory
      size_t size = i*sizeof(int);
      shost = (int *)malloc(size);
      avhost = (int *)malloc(size);
  
    //allocate memory on device
      cudaMalloc((void **) &amp;sdev, size);
      cudaMalloc((void **) &amp;avdev, size);
      cudaMalloc((void **) &amp;adev, size);
      cudaMemcpy(adev, ahost, size, cudaMemcpyHostToDevice);
  
      if (i &gt;= 512){
          h = 512;
      }
      else h = i;
      dim3 dimGrid(((i+512)/512),1,1);                   
      dim3 dimBlock(h,1,1);
  
      // launch kernel
      digi_calc &lt;&lt;&lt;dimGrid, dimBlock &gt;&gt;&gt; (sdev, avdev, adev, N, n, i);
  
      //copy memory back to host
      cudaMemcpy(avhost, avdev, size, cudaMemcpyDeviceToHost);
      cudaMemcpy(shost, sdev, size, cudaMemcpyDeviceToHost);
  
    // end malloc's, memcpy's, kernel calls
      for(h = 0; h &lt;=i; h++){
      sum=fmod(sum+(double) shost[h]/ (double) avhost[h],1.0);
      }
    printf("Decimal digits of pi at position %d: %09d\n",n,(int)(sum*1e9));
      //free memory
      cudaFree(sdev);
      cudaFree(avdev);
      cudaFree(adev);
      free(shost);
      free(avhost);
      free(ahost);
    return 0;
  }
  </code></pre>