cudamemcpy error: "se agotó el tiempo de espera del inicio y finalizó"

Question

Mi código es una implementación paralela que calcula el enésimo dígito de pi. Cuando termino el kernel y trato de copiar la memoria de vuelta al host, aparece el error "El tiempo de espera se agotó y finalizó". Utilicé este código para la comprobación de errores para cada lanzamiento de cundamalloc, cudamemcpy y kernal.

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

Estas llamadas se decía que todo estaba bien hasta que la primera llamada cudamemcpy después de regresar del núcleo. el error ocurre en la línea "cudaMemcpy (avhost, avdev, size, cudaMemcpyDeviceToHost);" en principal. Cualquier ayuda es apreciada.

#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

Este es exactamente el mismo problema que preguntaste en this question. El kernel está siendo cancelado antes por el conductor porque tarda demasiado en terminar. Si usted lee la documentación de cualquiera de estas funciones de API en tiempo de ejecución, verá la siguiente nota:

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

todo lo que sucede es que la primera llamada a la API después del lanzamiento del kernel está devolviendo el error incurrido mientras que el kernel se ejecuta - en este caso, la llamada cudaMemcpy. La forma en que puede confirmar esto por sí mismo es hacer algo como esto inmediatamente después de la puesta en marcha del núcleo:

// 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); 

La llamada cudaPeekAtLastError() le mostrará si hay algún error en el lanzamiento del kernel, y el código de error devuelto por el cudaThreadSynchronize() llamada mostrará si se generaron errores mientras el núcleo se estaba ejecutando.

La solución es exactamente la descrita en la pregunta anterior: probablemente la forma más simple es rediseñar el código para que sea "reentrante" para poder dividir el trabajo en varios lanzamientos de kernel, con cada kernel abierto de manera segura debajo de la pantalla Límite del temporizador de vigilancia del conductor.

Answer 2

Cuda de alguna manera almacena en búfer todas las operaciones de lectura/escritura en la memoria global. Entonces puedes agrupar las operaciones en un bucle con algo de kernel, y realmente NO TENDRÁ TIEMPO. Luego, cuando llame al memcpy, todas las operaciones de búfer se realizan y puede agotar el tiempo de espera. El método para ir es llamar al procedimiento cudaThreadSynchronize entre iteraciones.

Recuerde: si una ejecución del kernel tarda solo un nanosegundo en calcular, no significa que sea tan rápido, algunas de las escrituras en la memoria global se realizan cuando se llama al memcpy o threadsynchronize.