我最近开始使用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之间?
在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)可能有效,但是由于非理想的内存访问模式,速度会慢得多。