GPU高性能计算CUDA编程：线程块的维度

GPU高性能计算CUDA编程：线程块的维度

声明：本文不做商用

你的直觉一定会告诉你，线程块的维度不会只有1D的情况，它们可以是2D或3D的。例如，如果要在每个块中启动256个线程，可以用大小为16x16的2D线程数组或大小为8x8x4的3D线程数组来启动它们。在这种情况下，线程ID的范围如下:

• 当线程组是大小为256的1D数组时，则blockDim.x=256，blockDim.y=1，block-Dim.z=1，线程ID的范围是:threadIdx.x=0…255,threadIdx.y=0，threadIdx.z=0。
• 当线程组是大小为16x16的2D数组时，则blockDim.x=16，blockDim.y=16，线程ID的范围是:threadIdx.x=0…15,threadIdx.y=0 … 15,threadIdx.z=0。
• 当线程组是大小为8x8x4的3D数组时，则blockDim.x=8，blockDim.y-8，blockDim.z=4，线程ID的范围是:threadIdx.x=0…7,threadIdx.y=0…7,threadIdx.z=0…3.

在这三种情况中，每一个的线程块都启动了256个线程。这三种情况的不同之处在于，2D或3D线程数组分别相当于在两层或三层嵌套for循环内执行线程，而不是在一层for循环内，并将for循环变量(threadIdx.x、threadIdx.y和threadIdx.z)传递给核函数。这意味着程序员不必担心与线程有关的for循环。这是GPU硬件实现的功能。换句话说，你可以免费实现循环。但是，这并不意味着循环立即免费。每个核函数仍然需要检查这个庞大的变量列表来查看自己是谁。参见【0voice C++】
继续用网格x和y尺寸为768x217的2D网格举例，内层循环会执行完整的线程块块ID为 blockIdx.x和 blockldx.y。执行一个线程块意味着执行该线程块内的每个线程。每个块内的线程数量在另一个参数blockDim.x中，此处为256。每一个线程中执行的都是相同的核函数 Vflp()(如代码6.7所示)，所以就好像正在一个for 循环中运行 Vflip()函数循环 256次，如下所示:

for(blockIdx.x=0;blockIdx.x<=767，blockIdx.x++){
    for(blockIdx.y=0;blockIdx.y<=216，blockIdx.y++){
        //在此处执行线程块(blockldx.x,blockldx.y)
        //该线程块将能访问参数gridDim,x=768和gridDim.y=217.
        //并将它们传送给属于该线程块的线程
        //执行该线程块意味着执行256个线程(8个线程束)
        //线程块的大小可以通过blockDim.x、blockDim.y和blockDim.z获得
        for(threadIdx.x=0;threadIdx.x<=7,threadIdx.x++){
            for(threadIdx.y=0;threadIdx.y<=7,threadIdx.y++){
                for(threadIdx.z=0:threadIdx.z<=3,threadIdx.z++){
                    //执行该线程块的1个线程，该线程块有8*8*4=256个线程
                    //该线程将继承qridDim.x、gridDim.y
                    //该线程将继承blockIdx.x、blockIdx.y
                    //该线程将继承blockDim.x、blockDim.y、blockDim.z
                    //所有的GPU核函数参数将被传递给下面的函数
                    Vflip(...);
                }
            }
        }
    }

}

该伪代码显示了以二维线程块启动核函数时的情况，每个线程块都有一个2D索引。此外，每个块由总计256个线程的3D线程数组组成，组织为8x8x4的数组，每个线程均具有3D线程ID。注意，在这种情况下，我们免费得到五层for()循环，因为Nvidia GPU的硬件在启动线程块和线程时会生成所有的五层循环变量。