CUDA 中的线程组织

正在努力的小破孩

已于 2023-10-09 13:30:37 修改

阅读量225

点赞数

分类专栏： CUDA编程文章标签： c++

于 2023-10-09 13:29:57 首次发布

本文链接：https://blog.csdn.net/weixin_54137865/article/details/133691757

版权

CUDA编程专栏收录该内容

2 篇文章 0 订阅

订阅专栏

1 CUDA 中的 Hello World 程序

CUDA 程序的编译器驱动 nvcc 支持编译纯粹的 C++ 代码。一般来说，一个标准的CUDA 程序中既有纯粹的 C++ 代码，又有不属于 C++ 的真正的 CUDA 代码。CUDA 程序的编译器驱动 nvcc 在编译一个CUDA 程序时，会将纯粹的 C++ 代码交给 C++的编译器去处理，它自己则负责编译剩下的部分。CUDA 程序源文件的扩展名是 .cu。

GPU 只是一个设备，要他工作的话，需要由主机给他下达命令。这个主机就是 CPU。所以，一个正真正利用了 GPU 的 CUDA 程序既有主机代码，又有设备代码。主机对设备的调用是通过核函数来实现的。核函数在 GPU 上进行并行运行。

CUDA 中的核函数与 C++ 中的函数是类似的，但一个显著的差别是它必须被限定词 __global__ 修饰。另外，核函数的返回值必须是空类型 void。【void 和限定词的次序可随意】

核函数注意事项：

核函数只能访问 GPU 内存
核函数不能使用变长参数
核函数不能使用静态变量
核函数不能使用函数指针
核函数具有异步性
核函数不支持 C++ 的 iostream

#include<stdio.h>

__global__ void hello_from_gpu(){
    printf("Hello World from the GPU!\n");
}

int main(){
    hello_from_gpu<<<1, 1>>>();
    // 同步函数，等待GPU执行完毕
    cudaDeviceSynchronize(); 
    return 0;
}

2 CUDA 中的线程组织

2.1 线程模型结构

核函数的调用格式为

hello_from_gpu<<<1, 1>>>

主机在调用一个核函数时，必须指明需要在设备中指派多少个线程，否则设备不知道该如何工作。三括号就是用来指明核函数中的线程数目及排列情况的。
核函数中的线程常组织为若干线程块：三括号中的第一个数字可以看做是线程块的个数，第二个数字可以看做每个线程块中的线程数。一个核函数的全部线程块构成一个网格，而线程块的个数就记为网格大小。每个线程块中含有同样数目的线程，该数目称为线程块大小。
所以，核函数中总的线程数就等于网格大小乘以线程块大小，而三括号中的两个数字分别为网格大小和线程块大小，即<<<网格大小，线程块大小>>>。核函数中代码的执行方式为“单指令-多线程”，即每一个线程都执行同一串指令。

2.2 一维线程模型

每个线程在核函数中都有一个唯一的身份标识。每个线程的唯一标识由这两个<<<网格大小，线程块大小>>>确定。这两个值（网格大小：grid_size，线程块大小：block_size）保存在内建变量中，即：

gridDim.x：该变量的数值等于执行配置中变量 grid_size 的值

blockDim.x：该变量的数值等于执行配置中变量 block_size 的值

类似的，在核函数中预定义了如下标识线程的内建变量：

blockIdx.x：该变量指定一个线程在一个网格中的线程块索引值，范围为0~ gridDim.x-1

threadIdx.x：该变量指定一个线程在一个线程块中的线程索引值，范围为0~ blockDim.x-1

#include<stdio.h>

__global__ void hello_from_gpu(){
    const int bid = blockIdx.x;
    const int tid = threadIdx.x;

    const int id = threadIdx.x + blockDim.x * blockIdx.x;
    printf("Hello World from block %d and thread %d, global id %d!\n", bid, tid, id);
}

int main(){
    hello_from_gpu<<<2, 4>>>;
    cudaDeviceSynchronize();
    return 0;
}

一维线程模型的唯一ID可以由下面的公式计算得到：

$id = threadIdx.x+blockIdx.x\times blockDim.x$

2.3 推广至多维网格

前述所说的四个内建变量，都使用的结构体的成员变量的语法，其中：

blockIdx 和 threadIdx 是类型为 uint3 的变量。该类型是一个结构体，具有x、y、z 这3个成员。所以，blockIdx.x 只是3 个成员中的一个，另外两个分别是 blockId.y 和 blockIdx.z。类似的，threadIdx.x 只是3 个成员中的一个，另外两个分别是 threadId.y 和 threadIdx.z。
gridDim 和 blockDim是类型为 dim3 的变量。该类型是一个结构体，具有x、y、z 这3个成员。所以，gridDim .x 只是3 个成员中的一个，另外两个分别是 gridDim .y 和 gridDim .z。类似的，blockDim.x 只是3 个成员中的一个，另外两个分别是 blockDim.y 和 blockDim.z。除了和结构体 uint3 有同样的3个成员外，还在使用C++程序的情况下定义了一些成员函数。

这些内建变量都只在核函数中有效，且满足如下的关系：

2.4 线程全局索引计算方式

2.4.1 一维网格一维线程块

定义网格（grid）和线程块（block）的尺寸：

// 其他维度，不写则默认为1
dim3 grid_size(4);
dim3 block_size(8);

全局唯一标识为

int id = blockDim.x * blockIdx.x + threadIdx.x;

2.4.2 二维网格二维线程块

定义网格（grid）和线程块（block）的尺寸：

// 其他维度，不写则默认为1
dim3 grid_size(2,2);
dim3 block_size(4,4);

全局唯一标识为

int blockId = blockIdx.x + blockId.y * gridDim.x;

int threadId = threadIdx.y * blockDim.x + threadIdx.x;

int id = blockId * (blockDim.x * blockDim.y) + threadId;

2.4.3 三维网格三维线程块

定义网格（grid）和线程块（block）的尺寸：

// 其他维度，不写则默认为1
dim3 grid_size(2,2,2);
dim3 block_size(4,4,2);

全局唯一标识为

int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;

int threadId = (threadIdx.z * (blockDim.x * blockDim.y)) + (threadIdx.y * blockDim.x) + threadIdx.x;

int id = blockId * (blockDim.x *blockDim.y * blockDim.z) + threadId;

2.4.4 其余组合方式

2.5 网格和线程块大小的限制

网格大小限制：

线程块大小限制：

不管如何定义，一个线程块最多只能有1024个线程。

3 CUDA 中的头文件

在使用 nvcc 编译器驱动编译 .cu 文件时，将自动包含必要的 CUDA 头文件，如<cuda.h>和<cuda_runtime.h>。因为 <cuda.h> 包含 <stdlib.h>，顾用 nvcc 编译 CUDA 程序时甚至不用在 .cu文件中包含<stdlib.h>。

4 用 nvcc 编译 CUDA 程序

4.1 nvcc 编译流程

nvcc 先将全部源码分离为主机代码和设备代码；
nvcc 先将设备代码编译为PTX （Parallel Thread Execution）伪汇编代码，再将PTX代码编译为二进制的 cubin 目标代码；
在将源代码编译为 PTX 代码时，需要用选项 -arch=compute_XY 指定一个虚拟架构的计算能力，用以确定代码中能够使用的 CUDA 功能；
在将 PTX 代码编译为 cubin 代码时，需要用选项 -code=sm_ZW 指定一个真实架构的计算能力，用以确定可执行文件能够使用的 GPU。
真实架构的计算能力一定要大于或等于虚拟架构的计算能力

4.2 PTX

PTX （Parallel Thread Execution）是 CUDA 平台为基于 GPU 的通用计算而定义的虚拟机和指令集。

nvcc 编译命令总是使用两个体系结构：一个是虚拟的中间体系结构，另一个是实际的 GPU 体系结构。

虚拟架构更像是对应用所需的GPU功能的声明。虚拟架构应该尽可能选择低——适配更多实际 GPU，真实架构应该尽可能选择高——充分发挥 GPU 性能。

4.3 GPU 架构与计算能力

每款 GPU 都有用于标识“计算能力”的版本号。形式为 X.Y，X标识主版本号，代表 GPU 的架构， Y标识次版本号。【GPU 的型号与 GPU的架构是两种不同的概念，型号为Tesla，其架构可以为伏特架构】。

并非 GPU 的计算能力越高，性能就越强。

4.4 CUDA 程序兼容性问题

C/C++源码编译为PTX时，可以指定虚拟架构的计算能力，用来确定代码中能够使用的CUDA功能。C/C++源码转化为PTX这一步骤与GPU硬件无关。

4.4.1 指定虚拟架构计算能力

编译指令：

-arch=compute_XY

XY：第一个数字X代表计算能力的主版本号，第二个数字Y代表计算能力的次版本号。主版本用于确定 GPU 的架构，次版本号用于确定对应版本的计算能力。

4.4.2 指定真实架构计算能力

-code=sm_XY

注意事项：

二进制cubin代码，大版本之间不兼容!!!
指定真实架构计算能力的时候必须指定虚拟架构计算能力!!!
指定的真实架构计算能力必须大于或等于虚拟架构计算能力!!!

4.4.3 指定多个版本编译

可以同时指定多种计算能力，是的编译出来的可执行文件可以在多 GPU 中执行。

同时指定多组计算能力：

-gencode=arch=compute_35, code=sm_35

-gencode=arch=compute_50, code=sm_50

-gencode=arch=compute_60, code=sm_60

-gencode=arch=compute_70, code=sm_70

编译出来的可执行文件将包括 4 个二进制版本，这样的二进制文件称为胖二进制文件。

4.4.4 nvcc 即时编译

在运行可执行文件时，从保留的PTX代码临时编译出cubin文件

在可执行文件中保留PTX代码，nvcc编译指令指定所保留的PTX代码虚拟架构，指令为

-gencode arch=compute_60, code=compute_60

注意：

（1）两个计算能力都是虚拟架构计算能力

（2）两个虚拟架构计算能力必须一致

简化编译选项

-arch=sm_XY

---->

-gencode arch=compute_XY, code=sm_XY

-gencode arch=compute_XY, code=compute_XY

4.4.5 nvcc 编译默认计算能力

不同版本CUDA编译器在编译CUDA代码时，都有一个默认计算能力。

正在努力的小破孩

关注

0
点赞
踩
3

收藏

觉得还不错? 一键收藏
2
评论
CUDA 中的线程组织

CUDA 程序的编译器驱动 nvcc 支持编译纯粹的 C++ 代码。一般来说，一个标准的CUDA 程序中既有纯粹的 C++ 代码，又有不属于 C++ 的真正的 CUDA 代码。CUDA 程序的编译器驱动 nvcc 在编译一个CUDA 程序时，会将纯粹的 C++ 代码交给 C++的编译器去处理，它自己则负责编译剩下的部分。CUDA 程序源文件的扩展名是 .cu。GPU 只是一个设备，要他工作的话，需要由主机给他下达命令。这个主机就是 CPU。
复制链接

扫一扫