可能是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的教科书"编程大规模并行处理器"第9章第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]吗？在一次迭代之后,"stride"然后变为2,我们将看到threadIdx.x = 0的temp [-2]？

这对我来说没有多大意义,虽然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;
}

你会发现控制台里什么都没有.