Warm tip: This article is reproduced from serverfault.com, please click

indexing-是否为CUDA中的每个内核调用保证唯一线程ID?

(indexing - Is Unique Thread Id guaranteed for each Kernel Call in CUDA?)

发布于 2020-12-06 12:31:34

我最近开始使用Cuda,在C ++,Java和Python上有多线程,多进程编码的经验。

有了PyCuda,我看到了这样的示例代码,

ker = SourceModule("""
__global__ void scalar_multiply_kernel(float *outvec, float scalar, float *vec)
{
     int i = threadIdx.x;
     outvec[i] = scalar*vec[i];
}
""")

线程ID本身似乎参与了代码的逻辑。然后的问题是,是否有足够的线程ID覆盖我的整个数组(显然我需要为其建立索引的索引必须到达那里的所有元素),以及如果更改数组的大小会发生什么情况。

索引是否总是在0到N之间?

Questioner
BBSysDyn
Viewed
11
Paul G. 2020-12-15 19:24:06

在CUDA中,线程ID在每个所谓的线程块中都是唯一的,这意味着你的示例内核仅在一个块工作的情况下才做正确的事情。可以在早期示例中完成此操作,以使你容易理解,但就性能而言,通常这是一件非常糟糕的事情:

对于一个块,你只能利用GPU中的多个流式多处理器(SM)之一,即使SM在等待时有足够的并行工作要做时也只能隐藏内存访问等待时间。

如果你的内核不包含循环,则单个线程块还限制了线程数,因此也限制了问题的大小,因此每个线程可以计算多个元素。

内核执行在层次结构上看得很清楚:为简单起见,我们将自己限制为一维索引,内核在所谓的gridDim.x线程块网格上执行,每个线程块包含的blockDim.x线程数由threadIdx.x每个块编号,而每个块的编号均由编号blockIdx.x

要获取线程的唯一ID(最好使用硬件从数组中加载元素的方式),你必须采用blockIdx.x * blockDim.x + threadIdx.x如果每个线程都要计算一个以上的元素,则使用以下形式的循环

for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < InputSize; i += gridDim.x * blockDim.x) { 
/* ... */
}

这称为网格跨度循环,因为它gridDim.x * blockDim.x是内核上所有工作线程的数量。不同的步幅(特别是让一个线程在连续元素上工作:步幅= 1)可能有效,但是由于非理想的内存访问模式,速度会慢得多。