cudamemcpy त्रुटि: "लॉन्च का समय समाप्त हो गया और इसे समाप्त कर दिया गया"

Question

मेरा कोड एक समानांतर प्रत्यारोपण है जो पीआई के nth अंक की गणना करता है। जब मैं कर्नेल को समाप्त करता हूं और मेमोरी को मेजबान पर वापस कॉपी करने का प्रयास करता हूं तो मुझे "लॉन्च टाइम आउट और समाप्त कर दिया गया" त्रुटि मिलती है। मैंने प्रत्येक कोडडॉलोक, कुडामेम्पी और कर्नल लॉन्च के लिए त्रुटि जांच के लिए इस कोड का उपयोग किया।

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

ये कॉल कह रहे थे कि कर्नेल से लौटने के बाद पहले cudamemcpy कॉल तक सब कुछ ठीक था। त्रुटि "cudaMemcpy (avhost, avdev, आकार, 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 में पूछा था। कर्नेल को ड्राइवर द्वारा जल्दी समाप्त कर दिया जा रहा है क्योंकि इसे खत्म करने में बहुत लंबा समय लग रहा है। आप इन क्रम एपीआई कार्यों से किसी के लिए दस्तावेज़ पढ़ें, तो आप निम्नलिखित टिप्पणी देखेंगे:

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() कॉल दिखाएगा कि कर्नेल निष्पादित करते समय कोई त्रुटि उत्पन्न हुई थी या नहीं।

समाधान पिछले प्रश्न में उल्लिखित है: शायद सबसे आसान तरीका कोड को फिर से डिजाइन करना है, इसलिए यह "पुनः प्रवेश" है ताकि आप कई कर्नेल लॉन्च पर काम को विभाजित कर सकें, प्रत्येक कर्नेल डिस्प्ले के तहत सुरक्षित रूप से लॉन्च हो चालक वॉचडॉग टाइमर सीमा।

Answer 2

कूडा किसी भी तरह वैश्विक स्मृति पर सभी पढ़ने/लिखने के संचालन को बफर करता है। तो आप कुछ कर्नेल के साथ कुछ लूप में ऑपरेशन बैच कर सकते हैं, और वास्तव में इसमें कोई समय नहीं लगेगा। फिर, जब आप memcpy पर कॉल करते हैं, तो सभी बफर किए गए ऑपरेशन किए जाते हैं, और यह समय-समय पर हो सकता है। साथ जाने का तरीका, पुनरावृत्तियों के बीच cudaThreadSynchronize प्रक्रिया को कॉल करना है।

तो याद रखें: यदि कर्नेल रन केवल गणना करने के लिए नैनोसेकंड लेता है - इसका मतलब यह नहीं है कि यह बहुत तेज़ है - वैश्विक स्मृति में कुछ लिखते हैं, memcpy या threadsynchronize कहा जाता है।