理解和优化 pyCUDA 中的线程、块和网格

Posted

技术标签:

【中文标题】理解和优化 pyCUDA 中的线程、块和网格【英文标题】:Understanding and optimising threads, blocks, and grids in pyCUDA 【发布时间】:2017-06-04 05:35:09 【问题描述】:

我对 GPU 编程和 pyCUDA 非常陌生,并且在我的知识上存在相当大的差距。我花了很多时间搜索 SO,查看示例代码并阅读 CUDA/pyCUDA 的支持文档,但在解释中没有发现太多的多样性,并且无法理解一些事情。

我无法正确定义块和网格尺寸。我目前正在运行的代码如下,旨在将数组a 乘以浮点数b

from __future__ import division
import pycuda.gpuarray as gpuarray
import pycuda.driver as cuda
import pycuda.autoinit
from pycuda.compiler import SourceModule
import numpy as np

rows = 256
cols = 10
a = np.ones((rows, cols), dtype=np.float32)
a_gpu = cuda.mem_alloc(a.nbytes)
cuda.memcpy_htod(a_gpu, a)

b = np.float32(2)

mod = SourceModule("""
  __global__ void MatMult(float *a, float b)
  
    const int i = threadIdx.x + blockDim.x * blockIdx.x;
    const int j = threadIdx.y + blockDim.y * blockIdx.y;
    int Idx = i + j*gridDim.x;
    a[Idx] *= b;
  
  """)

func = mod.get_function("MatMult")

xBlock = np.int32(np.floor(1024/rows))
yBlock = np.int32(cols)

bdim = (xBlock, yBlock, 1)
dx, mx = divmod(rows, bdim[0])
dy, my = divmod(cols, bdim[1])
gdim = ( (dx + (mx>0)) * bdim[0], (dy + (my>0)) * bdim[1])
print "bdim=",bdim, ", gdim=", gdim

func(a_gpu, b, block=bdim, grid=gdim)
a_doubled = np.empty_like(a)
cuda.memcpy_dtoh(a_doubled, a_gpu)
print a_doubled - 2*a

代码应打印块尺寸bdim 和网格尺寸gdim,以及一个零数组。

这适用于较小的数组大小,例如,如果 rows=256cols=10(如上例所示)输出如下:

bdim= (4, 10, 1) , gdim= (256, 10)
[[ 0.  0.  0. ...,  0.  0.  0.]
 [ 0.  0.  0. ...,  0.  0.  0.]
 [ 0.  0.  0. ...,  0.  0.  0.]
 ..., 
 [ 0.  0.  0. ...,  0.  0.  0.]
 [ 0.  0.  0. ...,  0.  0.  0.]
 [ 0.  0.  0. ...,  0.  0.  0.]]

但是,如果我增加rows=512,我会得到以下输出:

bdim= (2, 10, 1) , gdim= (512, 10)
[[ 0.  0.  0. ...,  0.  0.  0.]
 [ 0.  0.  0. ...,  0.  0.  0.]
 [ 0.  0.  0. ...,  0.  0.  0.]
 ..., 
 [ 2.  2.  2. ...,  2.  2.  2.]
 [ 2.  2.  2. ...,  2.  2.  2.]
 [ 2.  2.  2. ...,  2.  2.  2.]]

表示对于数组的某些元素,乘法发生了两次。

但是,如果我将块尺寸强制为 bdim = (1,1,1),则问题不再出现,并且对于较大的数组大小,我得到以下(正确)输出:

bdim= (1, 1, 1) , gdim= (512, 10)
[[ 0.  0.  0. ...,  0.  0.  0.]
 [ 0.  0.  0. ...,  0.  0.  0.]
 [ 0.  0.  0. ...,  0.  0.  0.]
 ..., 
 [ 0.  0.  0. ...,  0.  0.  0.]
 [ 0.  0.  0. ...,  0.  0.  0.]
 [ 0.  0.  0. ...,  0.  0.  0.]]

我不明白这一点。这里发生了什么,这意味着随着数组大小的增加,这种定义块和网格尺寸的方法不再合适?另外,如果块的尺寸为(1,1,1),这是否意味着计算正在串行执行?

提前感谢您的任何指点和帮助!

【问题讨论】:

【参考方案1】:

您在 2D 块的 2D 网格上进行操作。在您的内核中,您似乎假设gridDim.x 将返回网格的x 维度中的线程数。

__global__ void MatMult(float *a, float b)

    const int i = threadIdx.x + blockDim.x * blockIdx.x;
    const int j = threadIdx.y + blockDim.y * blockIdx.y;
    int Idx = i + j*gridDim.x;
    a[Idx] *= b;

gridDim.x 返回块数 r x 网格方向,而不是线程数。为了获得给定方向的线程数,您应该将块中的线程数乘以同一方向的网格中的块数:

int Idx = i + j * blockDim.x * gridDim.x

【讨论】:

非常感谢您的回答。我想我的下一个问题是:有没有办法对所有数组形状进行概括?您在上面给出的解决方案非常有用,并且允许我使用更大的数组,但是如果我不断增加数组大小,我仍然会遇到同样的问题。另外,每个块中的线程是并行运行的,而每个块是串行运行的,还是网格中的每个块都是并行运行的?例如,(1,1,1) 的块形状会导致计算按顺序进行吗? 很高兴我能帮上忙,@FeeJF。我建议您阅读 this answer 和 this answer 并将其应用于您的解决方案。 太好了,非常感谢。除了链接之外,我应该能够取得相当大的进展并开始了解我在这里做什么!

以上是关于理解和优化 pyCUDA 中的线程、块和网格的主要内容,如果未能解决你的问题,请参考以下文章

来自 CUDA 代码的 100% GPU 使用率导致屏幕滞后

(py)CUDA中的网格和块尺寸[重复]

如何在 CUDA 中自动计算 2D 图像的块和网格大小?

CUDA 学习(十九)优化策略4:线程使用计算和分支

CUDA 学习(十九)优化策略4:线程使用计算和分支

交叉验证与网格搜索