memory - cudamemcpy 错误 :"the launch timed out and was terminated"

Question

我的代码是计算 pi 的第 n 位的并行实现。当我完成内核并尝试将内存复制回主机时，我收到“启动超时并被终止”错误。我使用此代码对每个 cudamalloc、cudamemcpy 和内核启动进行错误检查。

std::string error = cudaGetErrorString(cudaGetLastError());
printf("%s\n", error);

这些调用表明一切正常，直到从内核返回后的第一个 cudamemcpy 调用。错误发生在“cudaMemcpy(avhost, avdev, size, cudaMemcpyDeviceToHost);”行中在主要。感谢您的帮助。

#include <stdlib.h>
#include <stdio.h>
#include <math.h>

#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<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&1)!=0) {
    t1=0;
    t3=-v;
    goto Y4;
  } else {
    t1=1;
    t3=u;
  }

  do {

    do {
      if ((t1&1)==0) {
    t1=t1>>1;
    t3=t3>>1;
      } else {
    t1=(t1+v)>>1;
    t3=t3>>1;
      }
      Y4:;
    } while ((t3&1)==0);

    if (t3>=0) {
      u1=t1;
      u3=t3;
    } else {
      v1=v-t1;
      v3=-t3;
    }
    t1=u1-v1;
    t3=u3-v3;
    if (t1<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&1) r=mul_mod(r,aa,m);
    b=b>>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<=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 >= a) {                \
    do { kq-=a; } while (kq>=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<1; h++){
    if(tid > nthreads) continue;
    a = primes[tid];
    vmax=(int)(logf(3*N)/logf(a));
    if (a==2) {
      vmax=vmax+(N-n);
      if (vmax<=0) continue;
    }
    av[tid]=1;
    for(i=0;i<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<=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 > 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<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]>=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<2 || (n=atoi(argv[1])) <= 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]<=(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 **) &sdev, size);
    cudaMalloc((void **) &avdev, size);
    cudaMalloc((void **) &adev, size);
    cudaMemcpy(adev, ahost, size, cudaMemcpyHostToDevice);

    if (i >= 512){
        h = 512;
    }
    else h = i;
    dim3 dimGrid(((i+512)/512),1,1);                   
    dim3 dimBlock(h,1,1);

    // launch kernel
    digi_calc <<<dimGrid, dimBlock >>> (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 <=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;
}

Answer 1

这与您在 this question 中询问的问题完全相同.内核被驱动程序提前终止，因为它完成时间太长。如果您阅读任何这些运行时 API 函数的文档，您将看到以下注释:

Note: Note that this function may also return error codes from previous, asynchronous launches.

所有发生的事情是内核启动后的第一个 API 调用返回内核运行时发生的错误 - 在本例中为 cudaMemcpy 调用。您可以自己确认这一点的方法是在内核启动后直接执行类似的操作:

// launch kernel
digi_calc <<<dimGrid, dimBlock >>> (sdev, avdev, adev, N, n, i);
std::string error = cudaGetErrorString(cudaPeekAtLastError());
printf("%s\n", error);
error = cudaGetErrorString(cudaThreadSynchronize());
printf("%s\n", error);

cudaPeekAtLastError() 调用将显示内核启动是否有任何错误，cudaThreadSynchronize() 调用返回的错误代码将显示是否有任何错误内核执行时产生错误。

解决方案与上一个问题中概述的完全相同:可能最简单的方法是重新设计代码，使其“可重入”，这样您就可以将工作分配给多个内核启动，每个内核都在显示驱动程序下安全启动看门狗定时器限制。