0

我想测量 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]);  
}
4

1 回答 1

3

根据我的经验，代币是从零开始的。由于您只有 2 个参数，那就是 %0 和 %1。您正在使用 %2，它“大于提供的 asm 操作数的数量”。

于 2016-06-07T04:06:06.997 回答