Wie zu erklären Inline-PTX internen Compiler-Fehler von CUDA

Question

Ich möchte das Cache-Verhalten von GPU globalen Speicher und unten ist die Mikro-Benchmark, die ich Design. Was ich tun möchte, ist von der globalen Speicheradresse r_add0 zu laden und in den gemeinsamen Speicher s_tvalue [0] zu speichern. Aus irgendeinem Grund muss ich den Ladebefehl vom globalen Speicher durch Inline-PTX-Code ersetzen.

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

jedoch, wenn ich es mit nvcc kompilieren, es Beschwerde bei der Kompilierung Fehler

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

Hat jemand den Grund meinen Codes kennt.

komplette Codes siehe unten:

__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

Nach meiner Erfahrung Token sind Null basiert. Da Sie nur 2 Parameter haben, wären das% 0 und% 1. Sie verwenden% 2, was "größer ist als die Anzahl der bereitgestellten asm-Operanden".