CUDA内存复制错误:“启动超时并被终止”

9

我的代码是一个并行实现,用于计算π的第n位。当我完成内核并尝试将内存复制回主机时,出现“启动超时并被终止”的错误。 我使用了这段代码来检查每个cudamalloc、cudamemcpy和kernal launch的错误。

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;
}
2个回答

8
这正是你在这个问题中询问的相同问题。由于内核运行时间过长,驱动程序会提前终止内核,导致出现问题。如果你阅读任何这些运行时API函数的文档,你会看到以下说明:

注意: 请注意,此函数还可能返回之前异步启动的错误代码。

所有发生的事情都是内核启动后第一个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()调用返回的错误代码将显示内核执行时是否生成了任何错误。
解决方案与上一个问题中概述的完全相同:可能最简单的方法是重新设计代码,使其“可重入”,这样您就可以将工作分成几个内核启动,每个内核启动都在显示驱动程序看门狗定时器限制下安全地运行。

1
啊,我以为它至少有点不同,因为在内核执行完之后,我做了一个cudagetlasterror,它说没有错误。在另一个问题中,内核实际上运行了5秒钟,然后被看门狗关闭,但是这个内核在不到一秒钟的时间内就完成了。 - zetatr
我添加了你建议的代码,cudaPeekAtLastError没有出现错误,但是cudaThreadSynchronize超时并被终止,因为它持续了5秒以上。 - zetatr
这是预料中的。例如,如果您使用无效的内核参数,cudaPeekAtLastError会返回错误。 cudaThreadSynchronize阻止主机,直到内核完成或终止,并提供在cudaPeekAtLastError调用和内核结束之间发生的任何错误。 - talonmies
1
我意识到我经常从全局内存中读取相同的变量,决定只读一次并使用本地变量来存储它是更明智的选择。现在的问题是将最终结果写回全局内存会导致与之前相同的错误。我尝试注释掉两个全局写入,就没有任何错误了。我很难相信每个线程对全局内存进行两次写入会影响我的内核执行时间。 - zetatr
这是编译器优化。如果您不进行写操作,编译器足够聪明,可以确定产生写操作的所有代码都是多余的,并将删除该代码作为“死代码”。因此,如果没有写操作,您的内核可能为空。如果您编译为PTX并查看汇编程序,则可以确认此内容。当注释掉写操作时,指令应该会少得多。 - talonmies

0

Cuda会在全局内存中缓冲所有读写操作。因此,您可以使用某个循环和某个核函数批量操作,实际上不需要时间。然后,当您调用memcpy时,所有缓冲的操作都将完成,它可能会超时。推荐方法是在迭代之间调用cudaThreadSynchronize过程。

因此,请记住:如果一个内核运行只需要几纳秒的计算时间 - 这并不意味着它如此快速 - 一些对全局内存的写入是在调用memcpythreadsynchronize时完成的。


网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接