CUDA代码中‘ld’和‘add’指令参数不匹配问题

在CUDA编程中，有时会遇到内联汇编指令‘ld’和‘add’参数不匹配的问题。具体来说，以下代码段展示了这一问题：

__global__ void access(double *posArray) {
uint32_t tid = threadIdx.x;
double sink = 0;
for(uint32_t i = tid; i double* ptr = posArray + i;
asm volatile("{ \n"
".reg .f32 data;\n\t"
"ld.global.ca.f64 data, [%1];\n\t"
"add.f64 %0, data, %0;\n\t"
"}" : "+d"(sink) : "l"(ptr) : "memory");
}
// 同步所有线程
asm volatile("bar.sync 0;");
for(uint32_t i = 0; i double* ptr = posArray + i;
// 每个warp加载L1缓存中的所有数据
for(uint32_t j = 0; j uint32_t offset = (tid + j) % THREADS_NUM;
asm volatile("{ \n"
".reg .f64 data;\n\t"
"ld.global.ca.f64 data, [%1];\n\t"
"add.f64 %0, %0, data;\n\t"
"}" : "+d"(sink) : "l"(ptr + offset) : "memory");
}
}
}

编译时会出现以下错误：

ptxas /tmp/tmpxft_00003451_00000000-5_l1.ptx, line 63; error: Arguments mismatch for instruction 'ld'
ptxas /tmp/tmpxft_00003451_00000000-5_l1.ptx, line 64; error: Arguments mismatch for instruction 'add'

这些错误提示‘ld’和‘add’指令的参数不匹配。具体原因在于以下行：

.reg .f32 data;\n\t

应更改为：

.reg .f64 data;\n\t

因为‘ld.global.ca.f64’和‘add.f64’指令都操作64位浮点数（.f64），而不是32位浮点数（.f32）。正确的代码如下：

__global__ void access(double *posArray) {
uint32_t tid = threadIdx.x;
double sink = 0;
for(uint32_t i = tid; i double* ptr = posArray + i;
asm volatile("{ \n"
".reg .f64 data;\n\t"
"ld.global.ca.f64 data, [%1];\n\t"
"add.f64 %0, data, %0;\n\t"
"}" : "+d"(sink) : "l"(ptr) : "memory");
}
// 同步所有线程
asm volatile("bar.sync 0;");
for(uint32_t i = 0; i double* ptr = posArray + i;
// 每个warp加载L1缓存中的所有数据
for(uint32_t j = 0; j uint32_t offset = (tid + j) % THREADS_NUM;
asm volatile("{ \n"
".reg .f64 data;\n\t"
"ld.global.ca.f64 data, [%1];\n\t"
"add.f64 %0, %0, data;\n\t"
"}" : "+d"(sink) : "l"(ptr + offset) : "memory");
}
}
}

通过将寄存器类型从.f32更改为.f64，可以解决参数不匹配的问题，使代码能够正确编译和运行。