目录
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可以由下面的公式计算得到:
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代码时,都有一个默认计算能力。