在本章节中以最简单的CUDA程序:从GPU中输出“Hello World!”字符串开始CUDA编程的学习。将在WIN10系统上通过Command Prompt调用nvcc来进行编译和学习。
2.1 C++语言中的Hello World程序
首先回归一下C++程序的开发过程:
(1)用文本编辑器写一个源代码。
(2)使用编译器对源代码进行预编译、编译、汇编并链接目标文件可得到执行文件。
(3)运行可执行文件即可。
首先用编辑器写下这段代码并命名为hello.cpp,并用g++进行编译,g++ hello.cpp,编译完成后会生成一个a.exe的可执行文件，并调用生成的a.exe,可以看到屏幕上出现:
Hello World!
#include<stdio.h>
int main()
{
printf("Hello World\n");
return 0;
}2.2 CUDA中的Hello World程序
熟悉了C++语言中的Hello World程序之后，下面介绍CUDA中的Hello World程序。
2.2.1 只有主机函数的CUDA 程序
其实本质上已经写好了一个CUDA程序。原因在于CUDA程序的编译器驱动(compiler driver)nvcc支持编译不包含CUDA代码的纯C++函数。且一个CUDA源代码包含C++代码和CUDA代码两部分。在使用nvcc编译一个CUDA程序时，会将存粹的C++代码交给C++编译器(例如g++和cl)进行处理，nvcc则负责编译CUDA部分代码，而CUDA程序的源文件后缀名为.cu,所以直接将上述hello.cpp更名为hello1.cu并用nvcc进行编译:
nvcc hello1.cu
编译之后运行，与之前效果一致。
2.2.2 使用核函数的CUDA程序
虽然2.2.1节中使用nvcc进行编译的源程序，但是源程序中仅包含纯粹的C++代码，没有包含CUDA代码部分，所以也就没有用到GPU，本章节中将在源程序中添加CUDA代码，这样就是一个完整的CUDA程序，并调用GPU输出Hello World!。
首先GPU是一个设备，如果需要工作需要主机下达命令。这个主机就是CPU。从而一个完整的CUDA程序包含主机代码，又有设备代码(可以理解为需要设备执行的代码)。主机对设备的调用是通过核函数来实现的，一个典型的CUDA程序框架具有以下形式:
int main(void)
{
//主机代码
//核函数的调用
//主机代码
return 0;
}从而核函数以及设备调用是在主机代码之间。
CUDA中的核函数与C++中的函数是类似的，但一个显著的区别在于它必须有以下两个限制:
(1)在函数声明和定义时必须加__global__限定符(qualifier),其中global前后都是双下划线，且不能返回任何值，返回值必须为空void，根据这两个要求，书写一个打印字符串的核函数:
//写法一
__global__ void hello_from_gpu()
{
printf("Hello World from the GPU");
}
//写法二
__global__ void hello_from_gpu()
{
printf("Hello World from the GPU");
}在声明和定义了核函数之后也需要调用，写一个主函数调用核函数的完整CUDA代码，用以下命令进行编译：
nvcc hello2.cu
然后运行得到的可执行文件就可以从屏幕上看到以下输出:
Hello Wolrd from the GPU!
#include<stdio.h>
__global__ void hello_from_gpu()
{
printf("Hello World from the GPU!\n");
}
int main(void)
{
hello_from_gpu<<<1,1>>>();
cudaDeviceSynchronize();
return 0;
} 上述程序有三个地方需要进一步解释：
(1)调用核函数的格式:
hello_from_gpu<<<1,1>>>();
这个调用与普通C++函数的调用格式是有区别的。首先在函数名hello_from_gpu和括号()之间有一对三括号<<<>>>，其中还有用逗号分隔开的数字(1,1)。调用核函数时，必须包含这对三括号中的信息，现在来讲解一下为什么需要三括号中的信息。原因在于一块GPU中有很多计算核心(Tesla V100中有5120个)，可以支持很多线程(thread)。主机代码在调用核函数时必须指明在设备中指派多少个线程，否则设备不知道如何工作。三括号<<<>>>中数字就是用来指明核函数中的线程数目和排列情况。
核函数中线程常组织为若干线程块(thread block)：三括号中第一个数字可以看成是线程块的数目。第二个数字可以看成每个线程块中的线程数。全部线程块(thread block)构成一个网格(grid),线程块(thread block)的个数记为网格大小(grid size)。每个线程块(thread block)中都有相同数目的线程，该数目称为(block size)线程块大小。所以核函数中总的线程数目等于网格大小(线程块数目)乘以线程块大小(一个线程块中线程数目)，从而三括号中两个数字分别为<<<网格大小,线程块大小>>>从而在上述程序中，只指派了一个线程块，网格数为1，一个线程块中只有一个线程，线程块大小为1，从而一共只指派了一个线程。
(2)核函数中使用printf()函数的使用方式和C++库基本一样，#include<stdio.h>也可以写成#include <cstdio.h> ，但并不支持#include <iostream>。
(3)在调用完核函数之后，有如下一行语句:
cudaDeviceSynchronize();
这行语句调用了一个CUDA的运行时API函数，去掉这个函数将不能正确输出字符串。原因在于调用输出函数时，输出流先存放在缓冲区，而缓冲区不会自动刷新。只有程序遇到某种同步操作时缓冲区才会刷新。函数cudaDeviceSynchronize()作用是同步主机和设备，促使缓冲区刷新。
2.3 CUDA中的线程组织
2.3.1 使用多个线程的核函数
核函数中允许指派多个线程，这是一个必然的特征。一个GPU一共有几千个核心，从何总的线程数必须大于等于核心数目才能充分利用GPU的全部计算资源。实际上线程数大于计算核心时才会充分利用GPU中的计算资源，可以让计算、内存访问、不同计算之间合理的重叠，从而减少核心空闲的时间。
所以，所以在调用核函数时需要指定多个线程，以下代码hello3.cu指定了包含两个线程块的网格，而且每个线程块的大小是4。
#include<stdio.h>
__global__ void hello_from_gpu()
{
printf("Hello World from the GPU!\n");
}
int main()
{
hello_from_gpu<<<2,4>>>();
cudaDeviceSynchronize();
}因为网格大小为2(线程块数目为2)，线程块大小为4(网格中一个线程块中线程的数目),故总的线程数为2×4=8。也就是说主机指派该核函数调用时一共用了八个线程。核函数中的代码执行方式是“单指令-多线程模式”，即每一个线程都执行同一段代码。从而执行该程序会打印八行一样的文字：
Hello World from the GPU!
但是需要知道每一行是由哪一个线程所输出的。
2.3.2 使用线程索引
根据前面的介绍可以知道，可为一个核函数指派多个线程，这些线程是由执行配置(execution configuration)
<<<grid_size,block_size>>>
决定的。这里的grid_size(网格大小)核block_size(线程块大小)一般是用一个结构体类型的变量，但也可以是普通的int型变量，这两个数的乘积就是总的线程数目。
从开普勒架构开始，最大允许的线程块数目是-1(最大的网格大小)，最大的线程块大小为1024(一个线程块中允许指定的线程数目)，两万亿远大于编程中需要的线程数目，一般来说比GPU的核心数多几倍就可以充分利用GPU中的全部计算资源。核函数可以指定的线程数是非常巨大的，但在执行时同时活跃(其他的进行等待)的线程数由硬件(CUDA核心数目)和软件(核函数)共同决定。
每个线程都有唯一的身份标识符。由于通过线程块数量和线程块大小共同指定，所以灭个线程需要表示其在网格(grid)中的第几个，和这个线程块(block)中的第几个线程,此外调用核函数指派的执行配置参数grid_size和block_size保存在内建变量(built-in variable)中，可以直接引用。
(1)gridDim.x:该变量的数值等于执行配置中变量grid_size的数值。
(2)blockDim.x:该变量的数值等于执行配置中变量block_size的值。
类似的在核函数中预定义了标识线程的内建变量:
(1)blockIdx.x:该变量标识一个线程在一个网格中的线程块指标，其取值在0-gridDim.x-1之间。
(2)threadIdx.x:该变量标识一个线程在一个线程块中的线程指标，其范围是0-blockDim.x-1之间。
在下面的代码中，指派八个线程，并且输出每个线程的身份标识符：
#include <stdio.h>
__global__ void hello_from_gpu()
{
const int bid=blockIdx.x;
const int tid=threadIdx.x;
printf("Hello World from block %d and thread %d!\n");
}
int main()
{
hello_from_gpu<<<2,4>>>();
cudaDeviceSynchronize();
return 0;
}用nvcc编译连接后生成exe文件并执行
也就是说，有时候线程块的执行计算是相互独立的，每个线程块中的每个线程都会进行一次计算。
2.3.3 推广到多维网络
前面介绍了四个内建(built-in variables)变量,gridDim(网格大小)、blockDim(线程块大小)、blockIdx(当前线程所在线程块)、threadIdx(当前线程所在线程块中位置)均使用C++中的结构体或者变量，其中:
blockIdx和threadIdx是类型为uint3类型的变量，该类型是一个结构体，具有x，y，z三个成员。所以blockIdx.x只是这三个成员中的一个，另外两个成员分别是blockIdx.y和blockIdx.z,类似的threadIdx也有这三个成员，其中结构体uint3在vector_types.h中的定义为：
struct __device_builtin__ uint3
{
unsigned int x,y,z;
};
typedef __device_builtin__ struct uint3 uint3; 也就是说该结构体由三个无符号整型组成。
(2)gridDim和blockDim是类型为dim3的变量，该类型是一个结构体，一共有x,y,z三个成员，其中成员数与uint3完全一致，不过还定义了一些成员函数。
前面介绍过，网格大小(grid_size、gridDim)和线程块大小(block_size、blockDim)的大小是在调用核函数时通过执行配置指定的。在之前的例子中，配置仅仅制定了两个变量：
<<<grid_size,block_size>>>
这两个整数是分别赋予gridDim.x和blockDim.x，在这种情况下线程块都是一维的。
可以用结构体dim3定义多维的网格和线程块
dim3 grid_size(x,y,z)
dim3 block_size(x,y,z)
若要定义一个2×2×1的网格以及3×2×1的线程块，可以定义如下:
dim3 drid_size(2,2,1)和dim3 block_size(3,2,1)也可以等价为：
dim3 grid_size(2,2)和dim block_size(3,2),省略部分可认定为1
多维网格和线程块本质上是一维的，就像本质上多维数组也是一维数组一样。与一个多维线程块指标threadIdx.x,threadIdx.y,threadIdx.z对应的一维指标是：
int tid=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y
对于某些问题，如第七章引入的矩阵转置问题，有时使用如下复合线程索引更合适：
int nx=threadIdx.x+blockDim.x*blockIdx.x:
int ny=threadIdx.y+blockDim.y*blockIdx.y:
int nz=threadIdx.z+blockDim.z*blockIdx.z;
一个线程块中的线程还可以细分为线程束(thread warp)。一个线程束(即一束线程)是同一个线程块中相邻的warpSize个线程。warpSize也是一个内建变量，表示线程束的大小，其值对于目前的GPU架构来说都是32。所以一个线程束就是连续的32个线程。具体来说，一个线程块中第0-31个线程属于第0个线程束，第32-63个线程属于第一个线程束，可以通过修改Hello World程序来展示使用多维线程块的核函数组织情况。程序的输出如下:
#include <stdio.h>
__global__ void hello_from_gpu()
{
const int b=blockIdx.x;
const int tx=threadIdx.x;
const int ty=threadIdx.y;
printf("Hello World from block %d and thread (%d,%d)!\n",b,tx,ty);
}
int main()
{
const dim3 block_size(2,4);
hello_from_gpu<<<1,block_size>>>();
cudaDeviceSynchronize();
return 0;
}因为线程块的大小为2×4，所以知道在核函数中，blockDim.x值为2，blockDim.y的值为4.可以看到threadIdx.x的取值为0到1,threadIdx.y的取值为0到3.此外由于gridDim.x的取值为1，所以blockIdx.x的值只能取为0。x维度的线程指标threadIdx.x是最内层的。
2.3.4网格与线程块大小的限制
CUDA中能够定义的网格大小和线程块大小做了限制。网格大小在x,y,z三个方向的最大允许值分别为-1、65535、65535；线程块大小在x、y、z这三个方向的最大允许值分别为1024、1024、64，此外还要求总的线程块大小，即blockDim.x、blockDim.y、blockDim.z的乘积不能大于1024，也就是说不论怎么定义，一个线程块最多只能由1024个线程。
2.4 CUDA中的头文件
在编写C++程序的时候，会包含C++的头文件<stdio.h>,但并没有包含任何CUDA的头文件。CUDA中也有一些头文件，但是在使用nvcc编译驱动编译.cu文件时，将自动包含必要的CUDA头文件，如<cuda.h>、<cuda_runtime.h>。因为<cuda.h>包含<stdlib.h>从而使用nvcc编译cuda程序时不需要在.cu文件中包含<stdlib.h>。但使用一些利用CUDA进行加速的应用程序库的时候，需要包含一些必要的头文件，有些时候还需要指定链接选项。
2.5 用nvcc编译CUDA程序
CUDA的编译器驱动(compiler driver)nvcc在处理源码的过程中先将源代码分离成主机代码和设备代码。主机代码完整支持C++语法，但设备代码只部分的支持C++。nvcc先将设备代码编译为PTX(parallel thread execution)伪汇编代码，在将PTX码转换为cubin目标代码。再将源代码编译为PTX代码时，需要用选项-arch=compute_XY指定一个虚拟架构的计算能力，用以确定代码中能够使用的CUDA功能。再将PTX代码编译为cubin代码时，需要用选项-code=sm_ZW指定一个真实架构的计算能力，用于确定可执行文件能够使用的GPU，真实架构的计算能力必须大于或等于虚拟架构的计算能力。例如，可以用选项
-arch=compute_35 -code=sm_60
进行编译，但不能用选项:
-arch=compute_60 -code=sm_35
编译时编译器会报错。如果仅仅针对一个GPU编译程序，一般情况下建议将两个计算能力都指定为一个。
用以上几个方式可执行文件只能在少数几个GPU中才能运行。选项-code=sm_ZW指定了GPU的真实架构为Z,W。对应可执行文件只能在主版本号为Z，次版本大于或等于W的GPU中运行，例如:
-arch=compute_35 -code=sm_35
编译出来的可执行文件只能在计算能力为3.5和3.7的GPU中执行。如果希望编译出来的可执行文件能够在更多的GPU中运行，可以同时指定多组计算能力，每一组用如下形式的编译选项：
-gencode arch=compute_35,code=sm_35
例如用选项：
-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
编译出来四个二进制版本分别对应开普勒架构、帕斯卡架构、伏特架构。这种可执行文件称为胖二进制文件(fatbinary)。在不同架构的GPU中运行时会自动选择对应的二进制版本。
nvcc还有一种即时编译(just-in-time compilation)的机制，可以在运行可执行文件时从其中保留的PTX代码临时编译出一个cubin目标代码。要在可执行文件时从中保留PTX代码临时编译出一个cubin代码，就必须用以下方式指定所保留的PTX代码的虚拟架构:
-gencode arch=compute_XY,code=compute_XY
这里的两个计算能力都是虚拟架构的计算能力，必须完全一致。在学习CUDA时有一个简化的编译选项可以调用：
-arch=sm_XY
其等价于:
-gencode arch=compute_XY,code=sm_XY
以下是一些cuda的计算能力
(1)CUDA6.0及更早：默认的计算能力是1.0
(2)CUDA6.5-CUDA8.0:2.0
(3)CUDA9.0-CUDA10.2:3.0
本章所用CUDA版本是10.1，从而本章使用了3.0的计算能力
关于更多nvcc编译器的更多介绍参照：
https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc
145
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
