在CUDA核函数中声明变量

Question

假设您在CUDA内核中声明了一个新变量，然后在多个线程中使用它，如下所示：

__global__ void kernel(float* delt, float* deltb) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
float a;
a = delt[i] + deltb[i];
a += 1;
}

内核调用看起来像下面这样，有多个线程和块：

int threads = 200;
uint3 blocks = make_uint3(200,1,1);
kernel<<<blocks,threads>>>(d_delt, d_deltb);

1. "a" 存储在堆栈上吗？
2. 每个线程初始化时，是否为其创建了一个新的 "a"？
3. 还是每个线程将独立访问 "a"，在未知时间可能会干扰算法的情况下？

- John W.

Answer 1

任何未带有外部定义符的核函数内声明的变量（标量或数组）都是对每个线程局部的，也就是每个线程都有自己的“副本”，不会发生线程之间的数据竞争！

编译器根据所执行的转换和优化选择本地变量是否驻留在寄存器或本地存储器（实际上是全局存储器）中。

有关哪些变量放在本地存储器中的详细信息，请参阅 NVIDIA CUDA 用户指南的第5.3.2.2章节。

- Luca Ferraro

Answer 2

以上均不正确。CUDA编译器足够聪明且具有优化能力，它可以检测到变量a没有使用，因此完整的代码可以被优化掉。您可以通过将选项-Xptxas=-v用于编译内核，并查看资源计数来确认这一点，该计数应基本上没有寄存器、本地存储器或堆栈。

在一个较为复杂的示例中，变量a可能会存储在每个线程的寄存器或本地存储器（off-die DRAM）中。

- talonmies