CUDA编程之线程模型

AI异构

发布于 2020-07-29 14:55:48

2.2K0

发布于 2020-07-29 14:55:48

文章被收录于专栏：AI异构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\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }

    // cudaDeviceSynchronize等待内核完成并返回（同步）
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", 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];
}