CUDA编程之线程模型

Posted 2021-04-13 AI异构

CUDA编程之线程模型

CUDA线程模型概述

线程模型

CUDA线程层次

线程层次——二维Block

线程层次——三维Block

通过上图线程层次可划分为：

网格（Grid）

• 一Kernel映射一网格

• 网格在设备上执行

• 划分为线程块

线程块（Block）

• 发射到SM上执行

• 利用共享存储器通信

• 划分为线程

线程（Thread）

• 映射到SP上执行

五个内建变量

运行时获得网格和块的尺寸及线程索引等信息。

• gridDim：包含三个元素x, y, z的结构体，表示网格在x,y,z方向上的尺寸，对应于执行配置中的第一个参数。

• blockDim：包含三个元素x, y, z的结构体，表示块在x,y,z方向上的尺寸，对应于执行配置的第二个参数

• blockIdx：包含三个元素x, y, z的结构体，分别表示当前线程所在块在网格中x, y, z方向上的索引

• threadIdx：包含三个元素x, y, z的结构体，分别表示当前线程在其所在块中x, y, z方向上的索引

• warpSize：表明warp的尺寸，在计算能力1.0的设备中，这个值是24，在1.0以上的设备中，这个值是32。

Kernel分配线程

一个kernel结构如下：Kernel<< >>(param1, param2, …)

• Dg：grid的尺寸，说明一个grid含有多少个block，为dim3类型，一个grid最多含有65535 * 65535 * 65535个block，Dg.x，Dg.y，Dg.z最大值为65535；

• Db：block的尺寸，说明一个block含有多少个thread，为dim3类型，一个block最多含有1024(cuda2.x版本)个threads，Db.x和Db.y最大值为1024，Db.z最大值64；（举个例子，一个block的尺寸可以是：1024 * 1 * 1 | 256 * 2 * 2 | 1 * 1024 * 1 | 2 * 8 * 64 | 4 * 4 * 64等）

• Ns：可选参数，如果kernel中由动态分配内存的shared memory，需要在此指定大小，以字节为单位；

• S：可选参数，表示该kernel处在哪个流当中。

CUDA向量加法深入理解grid、block、thread的关系及thread索引的计算

CUDA编程流程

• CPU在GPU上分配内存：cudaMalloc；

• CPU把数据发送到GPU：cudaMemcpy；

• CPU在GPU上启动内核（kernel），它是自己写的一段程序，在每个线程上运行；

• CPU把数据从GPU取回：cudaMemcpy；

• CPU释放GPU上的内存。

CUDA向量加法源代码

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <stdlib.h>

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

int main()
{
      // 定义向量长度
    const int arraySize = 16;
      // 向量a与b初始化填值
    int a[arraySize], b[arraySize], c[arraySize];
    for (int i = 0; i < arraySize; i++) {
        a[i] = b[i] = i;
    }
      // 向量c初始化为零
    int c[arraySize] = { 0 };

    // 向量加法并行化计算
    cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addWithCuda failed!");
        return 1;
    }

        //打印输出结果
    for (int i = 0; i < arraySize; i++) {  // 打印出来方便观察
        cout << c[i] << " ";
    }
    cout << endl;

    //退出之前调用cudaDeviceReset，以便分析和跟踪工具（如Nsight和Visual Profiler）显示完整的跟踪。
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }
    system("pause");
    return 0;
}

//CUDA实现向量加法操作.
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
      // 初始化设备端地址
    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;
    cudaError_t cudaStatus;

    // 选择0号GPU进行计算
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    // 为两个输入向量和一个输出向量申请显存
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    // 将主机端的内存拷贝到设备端的显存
    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    // 在GPU上启动一个内核，为每个元素启动一个线程
        dim3 grid(1, 1, 1), block(size, 1, 1);
    addKernel<<<grid, block>>>(dev_c, dev_a, dev_b);

    // 检查启动内核的任何错误
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s
", cudaGetErrorString(cudaStatus));
        goto Error;
    }

    // cudaDeviceSynchronize等待内核完成并返回（同步）
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!
", cudaStatus);
        goto Error;
    }

    // 将GPU显存上的输出向量复制到主机内存。
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

Error:
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);

    return cudaStatus;
}

基于源码的线程分配分析

方式一：grid(1, 1, 1), block(size, 1, 1)

首先定义一个线程如下图所示：

线程示意图

然后直观解释程序中的线程设置

dim3 grid(1, 1, 1), block(size, 1, 1); // 设置参数

在这段代码中，我们设置参数为线程格（grid）中只有一个一维的块（block），该block的x维度上有16个线程。图示如下：

因此，按照此线程分配情况执行代码如下：

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

方式二：grid(1, 1, 1), block(8, 2, 1)

dim3 grid(1, 1, 1), block(8, 2, 1); // 设置参数

直观图示：

索引执行代码

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.y * blockDim.x +  threadIdx.x;  // 使用了threadIdx.x, threadIdx.x, blockDim.x
    c[i] = a[i] + b[i];
}

方式三： grid(16, 1, 1), block(1, 1, 1)

dim3 grid(16, 1, 1), block(1, 1, 1);  // 设置参数

直观图示：

索引执行代码

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = blockIdx.x;
    c[i] = a[i] + b[i];
}

方式四： grid(4, 1, 1), block(4, 1, 1)

dim3 grid(4, 1, 1), block(4, 1, 1);    // 设置参数

直观图示：

索引执行代码

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = blockIdx.x * gridDim.x + threadIdx.x;
    c[i] = a[i] + b[i];
}

方式五： grid(2, 2, 1), block(2, 2, 1)

dim3 grid(2, 2, 1), block(2, 2, 1);    // 设置参数

直观图示：

索引执行代码

__global__ void addKernel(int *c, const int *a, const int *b)
{
    // 在第几个块中 * 块的大小 + 块中的x, y维度（几行几列）
    int i =  (blockIdx.y * gridDim.x + blockIdx.x) * (blockDim.x * blockDim.y) + threadIdx.y * blockDim.y + threadIdx.x;
    c[i] = a[i] + b[i];
}

参考

[CUDA基础（1）：操作流程与kernel概念]https://www.cnblogs.com/hankeyyh/p/6580427.html

[【CUDA】grid、block、thread的关系及thread索引的计算]https://blog.csdn.net/hujingshuang/article/details/53097222

－长按关注－