CUDA 线程块中可能有负索引？

Question

我有一个非常简单的 1D CUDA 内核做一个包含和，也就是说，如果我们有一个输入 1D 数组

[x_0, x_1, x_2,..., x_n-1]

输出将是

[ x_0, x_0+x_1, x_0+x_1+x_2, ..., x_0+x_1+...x_n-1 ]。

下面显示的内核实际上并没有完全完成这项工作，另一方面它在每个块内完成了它的工作。无论如何，我的问题不是关于如何完全实现包容性总和，而是我认为在线程计算期间可能存在负索引错误。

__global__ void parallel_scan_inefficient(float* input, float* output){
// num_threads and max_i are globalled defined  
__shared__ float temp[num_threads];

int i = blockIdx.x*blockDim.x+threadIdx.x;//global index

if (i<max_i)
{
    temp[threadIdx.x]=input[i];
}

for (unsigned int stride=1;stride<=threadIdx.x; stride*=2)
{
    __syncthreads();
    temp[threadIdx.x]+=temp[threadIdx.x-stride];
}

output[i]=temp[threadIdx.x];

}

这段程序实际上来自 Hwu&Kirk 的教科书“Programming Massively Parallel Processors”第 9 章 pp.203。

正如你在for循环中看到的那样

for (unsigned int stride=1;stride<=threadIdx.x; stride*=2)
{
    __syncthreads();
    temp[threadIdx.x]+=temp[threadIdx.x-stride];
}

因为每个块的“threadIdx.x”从 0 开始，但“stride”从 1 开始。我们不会看到例如 temp[-1] 用于块中的第一个元素吗？同样经过一次迭代，“步幅”变为 2，我们将看到 temp[-2] for threadIdx.x=0 ？

这对我来说不太有意义，尽管 CUDA 编译器没有报告任何错误 - 我为这个内核运行了 cuda-memcheck，它仍然没问题。结果也是正确的（当然每个块都是正确的，正如我所说，这个内核只是部分完成了包容和）

我认为我可能会犯一个非常愚蠢的错误，但我就是无法发现它。任何光都将不胜感激。非常感谢。

Answer 1

如果你有这样的代码：

for (unsigned int stride=1;stride<=threadIdx.x; stride*=2)
{
    __syncthreads();
    temp[threadIdx.x]+=temp[threadIdx.x-stride];
}

然后对于 threadIdx.x == 0 的线程，将完全跳过 for 循环。尝试在 main 中运行以下代码：

for (unsigned int stride=1;stride<=0; stride*=2)
{
    cout << "I am running" << endl;
}

您会看到控制台中没有任何内容。