Steven Huang Steven Huang - 1 month ago 17
C++ Question

How to explain inline PTX Internal Compiler Error of CUDA

I want to measure the cache behavior of gpu global memory and below is the micro-benchmark that i design. What i want to do is to load from global memory address r_add0 and store it into shared memory s_tvalue[0]. For some reason, i need to replace the loading instruction from global memory with inline PTX code.

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


However, when i compile it with nvcc, it complaint with compile error

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


Does anybody knows the reason to my codes.

Complete codes see below:

__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

In my experience, tokens are zero based. Since you only have 2 parameters, that would be %0 and %1. You are using %2, which "is larger than the number of asm operands provided."