CUDA의 인라인 PTX 내부 컴파일러 오류를 설명하는 방법

Question

gpu 전역 메모리의 캐시 동작을 측정하고 싶습니다. 아래에서 제가 설계 한 마이크로 벤치 마크입니다. 내가하고 싶은 일은 전역 메모리 주소 r_add0에서로드하여 공유 메모리 s_tvalue [0]에 저장하는 것입니다. 어떤 이유인지, 전역 ​​메모리에서 인라인 명령을 인라인 PTX 코드로 대체해야합니다. 그러나

i = *r_addr0; 
//asm("ldu.global.f64.cs %1, [%2];":"=l"(i):"l"(r_addr0)); 
s_tvalue[0] = i; 

, 내가 NVCC로 컴파일, 컴파일 오류

error: Internal Compiler Error (codegen): "asm operand index requested is larger than the number of asm operands provided!" 

이 사람이 내 코드에 대한 이유를 알고 있는가와는 불만.

전체 코드는 아래 참조 :

내 경험에

__global__ void global_latency (long long * my_array, 
           long long array_length, int position, 
           long long *d_time) 
{ 

    unsigned int start_time, end_time; 

    __shared__ long long s_tvalue[2];//2: number of threads per block 

    int k; 
    long long i, j; 
    for(k=0; k<2; k++) 
     s_tvalue[k] = 0L; 
    long long addr0,addr1; 

    addr0=(long long)my_array; 

    addr1 = (addr0^(1 << position)); 

    long long *r_addr0, *r_addr1; 
    r_addr0 = (long long *)addr0; 
    r_addr1 = (long long *)addr1; 

    start_time = clock(); 
    //i = *r_addr0; 
    asm("ldu.global.f64.cs %1, [%2];":"=l"(i):"l"(r_addr0)); 

    s_tvalue[0] = i; 
    //j = *r_addr1; 
    asm("ld.global.f64.cs %3, [%4];" : "=l"(j):"l"(r_addr1)); 
    s_tvalue[1] = j; 


    end_time = clock(); 

    d_time[0] = end_time-start_time; 
    d_time[1] = s_tvalue[0]; 
    printf("[%p]=%lld\n",addr0,d_time[1]); 
    d_time[2] = s_tvalue[1]; 
    printf("[%p]=%lld\n",addr1,d_time[2]); 
} 

Answer 1

, 토큰 기반의 제로입니다. 매개 변수가 2 개인 경우 % 0 및 % 1이됩니다. % 2을 (를) 사용하고 있습니다. "제공된 asm 피연산자 수보다 큽니다."