CUDA为什么要分线程块和线程网格？

1、线程

该处的线程和 CPU 上是不同的概念，CPU 上一般是执行不同任务，GPU 执行的这是相同的核心函数，是同时执行。当然 GPU 上每个线程的 thresholdidx 是唯一的。

1 CUDA 框架

Device：指的是 GPU 芯片。
Grid：对应 Device 级别的调度单位，一组block，一个grid中的block可以在多个SM中执行。
Block：对应 SM（Streaming Multiprocessor）级别的调度单位，一组thread，同block中的thread可以协作。
Thread：对应 CUDA Core 级别的调度单位，最小执行单元。

上一篇文章介绍的：CUDA函数用的<<<x,y>>>: x 即是grid_size， grid 的 block 数量； y 表示block_size， block 的 threshold 数量。因此<<<2, 3>>> 即表示分配2x3=6个线程运行某一核函数。

2 线程变量

Grid 和 Block 可以用多维矩阵表示，按照上图演示，定义线程的方式如下，z维度默认为1：

dim3 grid(3, 2, 1), block(5, 4, 3)

CUDA内置了一些变量方便我们调取线程的坐标：

threadIdx.[x y z]- 线程坐标。框架图中， Thread(2,1) 的 threadIdx.x = 2, threadIdx.y = 1。
blockIdx.[x y z]- block坐标。框架图中， Block 的 blockIdx.x = 1, blockIdx.y = 1。
blockDim.[x y z]- block的维度，表示一个block中包含多少个thread。框架图中，blockDim.x = 5, blockDim.y = 3。
gridDim.[x y z] - grid的维度，表示一个grid中包含多少个block。框架图中，gridDim.x = 3, gridDim.y = 2。

理解线程的索引方式非常重要！特别是后面用线程对存储进行操作。其中xy表示的和数学的矩阵坐标是相反的：x表示横着数第几列，有表示竖着数第几行。

网格大小（gridDim）：

x： 2^31−1
y：65535
z：65535

线程块大小（blockDim）：

x：1024
y：1024
z：64

另外，还要求线程块总的大小，即 blockDim.x、blockDim.y 和 blockDim.z 的乘积不能大于1024。也就是说，不管如何定义，一个线程块最多只能有1024个线程。这些限制是必须牢记的。

2 软硬件对照

图的阅读可以看出硬件的组成关系，也可以看出软件和硬件的对那个关系。这里解释下后者：

一个 thread 一定对应一个 CUDA Core，但是CUDA Core可能对应多个 thread。
一个Block内的线程一定会在同一个SM（Streaming Multiprocessor，注意不是后面经常提到的Shared Memory）内，一个SM可以运行多个Block。
每一个block内的thread会以warp为单位进行运算，一个warp对应一条指令流，一个warp内的thread是真正同步的，同一个warp内的thread可以读取其他warp的值。

思考：为何我们需要设计 Block 呢，增加一个抽象，增加了复杂度。

简单回复：

硬件设计架构决定的；
没有block的，会导致全局的同步开销过大，这么设计使得可扩展性更好；

4 CUDA函数执行流程

加载核函数
将Grid分配到一个Device
根据<<<..>>>内的执行设置的第一个参数，Giga threads engine将block分配到SM中。一个Block内的线程一定会在同一个SM内，一个SM可以有很多个Block。
根据<<<..>>>内的执行设置的第二个参数，Warp调度器会调用线程。
Warp调度器为了提高运行效率，会将每32个线程分为一组，称作一个warp。
每个Warp会被分配到32个core上运行。

2、线程索引

线程索引和数据地址的转换。CUDA核函数工作时，所有线程同时执行，以下面并行加法为例，需要为每个线程匹配对应的数据地址，因此需要将线程索引转换为数据地址，方法如上图。

3、线程分配

1 如何设置 Grid 和 Block 的大小：

block_size = 128;
grid_size = (N + block_size - 1) / block_size;

N为数据维度。最大可配置的Block大小为1024，详细可按照如下命令查询硬件：

cd /usr/local/cuda/samples/1_Utilities/deviceQuery
make
./deviceQuery

因为block内的线程是按照warp来调度的，所以blocksize尽量设置为32的倍数。

提问：block_size（尖括号第二位）不是32倍数的后果是什么，也可以通过的

解答：正确使用了内置的索引和形状变量(threadIdx, blockDim; blockIdx, gridDim)不会有正确性的影响。但可能会导致性能下降。

追问：改成了33就分成2个warp？

解答：是的。超出哪怕1个线程，也会分配一个warp（浪费31/32的潜在执行能力）。

2 数据大于线程数的处理方法

当数据过大，超过线程数的时候，可以采用如下方式处理：

__global__ add(const double *x, const double *y, double *z, int n)
{
	int index = blockDim.x * blockIdx.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;
	for(; index <n; index +=stride)
		z[index] = x[index] + y[index];
}

其中index为线程匹配的数据地址（下面会详细介绍），stride为总线程数。

4、案例代码

CUDA函数流程：

利用数据索引编写CUDA函数；
利用cudaMalloc为变量分配CUDA地址；
cudaMemcpy 将数据从内存拷贝到GPU；
<<<grid_size, block_size>>>调用CUDA函数；
cudaMemcpy 将数据从GPU拷贝回内存；
cudaFree 释放资源。

#include <math.h>
#include <stdio.h>

void __global__ add(const double *x, const double *y, double *z, int count)
{
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
	if( n < count)
	{
	    z[n] = x[n] + y[n];
	}

}
void check(const double *z, const int N)
{
    bool error = false;
    for (int n = 0; n < N; ++n)
    {
        if (fabs(z[n] - 3) > (1.0e-10))
        {
            error = true;
        }
    }
    printf("%s\n", error ? "Errors" : "Pass");
}


int main(void)
{
    const int N = 1000;
    const int M = sizeof(double) * N;
    double *h_x = (double*) malloc(M);
    double *h_y = (double*) malloc(M);
    double *h_z = (double*) malloc(M);

    for (int n = 0; n < N; ++n)
    {
        h_x[n] = 1;
        h_y[n] = 2;
    }

    double *d_x, *d_y, *d_z;
    cudaMalloc((void **)&d_x, M);
    cudaMalloc((void **)&d_y, M);
    cudaMalloc((void **)&d_z, M);
    cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);

    const int block_size = 128;
    const int grid_size = (N + block_size - 1) / block_size;
    add<<<grid_size, block_size>>>(d_x, d_y, d_z, N);

    cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
    check(h_z, N);

    free(h_x);
    free(h_y);
    free(h_z);
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    return 0;
}

代码的执行情况：可见相关链接的github代码库的 02_2.2。

5、问题回顾

1）每个block应该申请多少个线程呢？

底层是以warp为单位申请。如果blockDim为160，则正好申请5个warp。如果blockDim为161，则不得不申请6个warp。

2）如何设置Gridsize和blocksize呢？

对于一维的情况：

block_size=128;

grid_size = (N+ block_size-1)/block_size;

（没有设成什么值是最好的）

注意：代码可见相关链接的github代码库的 02_2.2。

CUDA为什么要分线程块和线程网格？

CUDA为什么要分线程块和线程网格？

1、线程

1 CUDA 框架

2 线程变量

2 软硬件对照

4 CUDA函数执行流程

2、线程索引

3、线程分配

1 如何设置 Grid 和 Block 的大小：

2 数据大于线程数的处理方法

4、案例代码

5、问题回顾

相关链接

Published

Category

Tags