HelloWorld!¶

参考书目： CUDA编程基础与实践

章节：第2章 CUDA中的线程组织

学习cuda第一件事当然是helloworld啦！

hello world in c++¶

首先复习一下c++的开发过程

写下一段代码 ⬇️

#include <stdio.h>

int main()
{
    printf("Hello World!\n");
    return 0;
}

终端使用g++可以对其进行预处理，编译，汇编，链接成一个可执行文件

g++ ./helloworld.cpp

此时可以在当前路径得到a.out的文件，可以执行./a.out

./a.out
## 👉 Hello World!

option

g++ -S helloworld.cpp # 只编译，不进行汇编链接 得到.s文件
g++ -c helloworld.cpp # 进行汇编、编译，不进行链接 得到.o文件
g++ -o main.exe hello.cpp # 编译链接生成可执行文件(a.out by default)

hello world in cuda¶

我们上面写的代码也可以被cuda代码编译器nvcc编译链接，由于一个cuda程序会有c++部分和cuda部分，nvcc在编译.cu文件时，会将c++代码交给c++编译器，cuda代码则会被nvcc编译

但此时代码并没有调用GPU,也就是全部在主机上进行计算的(host)，对GPU(也叫设备device)的调用需要使用核函数，核函数需要特定关键字__global__,返回类型必须为void

接下来写一段helloworld代码

#include <stdio.h>

__global__ void test(){
    int j = threadIdx.x;
    int i = blockIdx.x;
    printf("Hello World from Cuda in block%d, thread%d\n", i, j);
}

int main( int argc, char** argv )
{
    test<<<2,5>>>();//定义blocks和threads, 可以自行更改体会
    cudaDeviceSynchronize();//同步主机与设备,否则无法打印
    return 0;
}

$ nvcc helloworld.cu
$ ./a.out

Hello World from Cuda in block0, thread0
Hello World from Cuda in block0, thread1
Hello World from Cuda in block0, thread2
Hello World from Cuda in block0, thread3
Hello World from Cuda in block0, thread4
Hello World from Cuda in block1, thread0
Hello World from Cuda in block1, thread1
Hello World from Cuda in block1, thread2
Hello World from Cuda in block1, thread3
Hello World from Cuda in block1, thread4

调用核函数格式为<<<grid_size, block_size>>>,其中的两个数字代表线程块(block)数量和每个线程块的线程(thread)数,一个核函数的全部线程块构成一个网格(grid),线程块个数就是网格大小(最大为 $2^{31}-1$ ),一个线程块中的线程数量就是线程块大小(最大为1024),示例代码的例子中,网格大小为2,线程块大小为5,总的线程数为2*5=10,所以打印10条信息

每个线程在核函数中都有唯一的身份标识,也就是blockIdx和threadIdx,类似于索引,它的范围为[0,gridDim.x-1]和[0,blockDim.x-1],gridDim和blockDim时结构体,除了x也有y,z成员变量表示另外两个维度,与x类似. 符合正常逻辑, x维度在最内层,变化最快,z维度在最外层,变化最慢

所以我们可以用dim3构建多维网格和线程块,缺省参数就默认为1

dim3 grid_size(Gx, Gy, Gz);
dim3 block_size(Bx, By, Bz);

一个线程块中的线程还可以分为不同线程束(thread warp),目前GPU架构的线程束都为32

大小的限制

网格大小的x,y,z最大不超过 $2^{31}-1, 65535, 65535$ .线程块的x,y,z最大不超过 $1024,1024,64$ ,但线程块总大小不超过1024,也就是blockDim.x,blockDim.y,blockDim.z乘积不超过1024

cuda基础¶

一般在开发cuda程序时,nvcc会将源代码分离为主机代码和设备代码, 主机代码完整支持C++, 设备代码部分支持c++, nvcc会先将设备代码编译为PTX(parallel thread execution)伪汇编代码,再将PTX编译为二进制的cubin目标代码,编译到PTX的过程需要optionarch=compute_XY指定一个虚拟架构的计算能力,将PTX代码编译到cubin时需要option-code=sm_ZW指定一个真实架构的计算能力,以确定可执行文件能够使用GPU,真实架构计算能力必须大于等于虚拟架构计算能力,可执行文件只能再主版本号为Z,次版本号大于等于W的GPU中运行. 举个例子-arch=compute_35 -code=sm_35生成的可执行文件只能在计算能力为3.5和3.7的GPU中执行,如果希望在更多的GPU中执行,可以用-gencode同时指定多组计算能力,例如

-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

此option生成的可执行文件包含4个二进制版本,分别对应开普勒,麦克斯韦,帕斯卡和伏特架构,被称为fatbinary,GPU会自动选择对应的文件

nvcc有即时编译的机制(just-in-time compilation),可以在运行可执行文件时从其中保留的PTX代码临时编译一个cubin目标代码,用以下方式指定所保留的PTX代码的虚拟架构

-gencode arch=compute_XY,code=compute_XY

两个计算能力都是虚拟架构的,必须完全一致

在学习CUDA编程时,简化编译选项为-arch=sm_XY,等价于

-gencode arch=compute_XY,code=sm_XY
-gencode arch=compute_XY,code=compute_XY

不指定计算能力则编译器会选择默认的计算能力