02《CUDA编程基础与实践》-CUDA中的线程组织

第二章 CUDA中的线程组织

2.1 CUDA中的Hello World程序

2.1.1 只有主机函数的CUDA程序
一般来说一个标准的 CUDA 程序中既有纯粹的 C++ 代码，又有不属于 C++的真正的CUDA 代码。CUDA 程序的编译器驱动 nvcc 在编译一个 CUDA程序时，会将纯粹的 C++ 代码交给 C++ 的编译器(如前面提到的 g++ 或 cl)去处理它自己则负责编译剩下的部分。CUDA 程序源文件的扩展名是 cu，所以我们可以先将上面写好的源文件更名为 hello1.cu，然后用 vcc 编译:$ nvcc hello1.cu编译好之后即可运行。运行结果与 C++ 序的运行结果相同。
2.1.2 使用核函数的CUDA程序
GPU 只是一个设备，要它工作的话还需要有一个主机给它下达命令。这个主机就是 CPU。所以，一个真正利用了 GPU 的 CUDA 程序有主机代码(在程序 hellolcu 中的所有代码都是主机代码)，又有设备代码(可以理解为需要设备执行的代码)。主机对设备的调用是通过核函数(kernel function)来实现的。所以，一个典型的、简单的 CUDA 程序的结构具有下面的形式:

int main(void)
{
	主机代码
	核函数的调用
	主机代码
	return 0
}

CUDA 中的核函数与 C++ 中的函数是类似的，但一个显著的差别是它必须被限定词(qualifer)__global__修饰。其中，global 前后是双下划线。另外，核函数的返回类型必须是空类型，即 void。我们先写一个打印字符串的核函数

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

限定__global__和void的次序可随意。也就是说上述核函数也可以写为如下形式:

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

我们来写一个主函数来调用这个核函数：

#include<stdio.h>
#include "cuda_runtime.h"
__global__ void hello_from_gpu()
{
    printf("Hello World from the GPU!\n");
}
int main(void)
{
    hello_from_gpu << <1, 1 >> > ();
    cudaDeviceSynchronize();
    return 0;
}

上述程序有 3个地方需要做进一步解释：

(1)用核函数的格
hello_from_gpu<<<1，1>>>();这个调用格式与普通的 C++ 函数的调用格式是有区别的。我们看到，在函数名hello_from_gpu 和括号(之间有一对三括号 <<<>>>，其中还有用逗号隔开的两个数字 (1,1)。调用核函数时为什么需要这对三括号里面的信息呢?这是因为一块GPU 中有很多(例如，Tesla V100 中有 5120个)计算核心，可以支持很多线程(thread)。主机在调用一个核函数时，必须指明需要在设备中指派多少个线程否则设备不知道如何工作。三括号中的数就是用来指明核函数中的线程数目及排列情况的。核函数中的线程常组织为若干线程块 (thread block):三括号中的第一个数字可以看作线程块的个数，第二个数字可以看作每个线程块中的线程数。一个核函数的全部线程块构成一个网格(grid)，而线程块的个数就记为网格大小(gridsize)。每个线程块中含有同样数目的线程该数目称为线程块大小(block size)。所以，核函数中总的线程数就等于网格大小乘以线程块大小，而三括号中的两个数字分别为网格大小和线程块大小，即<<<网格大小,线程块大小>>>。所以，在上述程序中，主机只指派了设备的一个线程，网格大小和线程块大小都是 1,即 1x1=1。

(2)核函数中的 printf函的使用方 C++库(者说 C++从C中继承的库)中的 printf(函数的使用方式基本上是一样的。而且，在核函数中使用printf()函数时也需要包含头文件 <stdio.h>(也可以写成 <cstdio>)。需要注意的是，核函数中不支持 C++的iostream(读者可亲自测试)。

(3)在调用核函数之后，有如下一行语句:
cudaDeviceSynchronize();这行语句调用了一个 CUDA 的运行时 API 函数。去这个函数将不能输出字符串(请读者亲自尝试)。这是因为调用输出函数时，输出流是先存放在缓冲区的，而缓冲区不会自动刷新。只有程序遇到某种同步操作时缓冲区才会刷新。函数 cudaDeviceSynchronize()的作用是同步主机与设备，所以能够促使缓冲区刷新。读者现在不需要弄明白这个函数到底是什么，因为我们这里的主要目的是介绍CUDA中的线程组织。

2.2 CUDA中的线程组织

2.2.1 使用多个线程的核函数
一个 GPU 往往有几千个计算核心，而总的线程数必须至少等于计算核心数时才有可能充分利用GPU 中的全部计算资源。实际上，总的线程数大于计算核心数时才能更充分地利用 GPU 中的计算资源，因为这会让计算和内存访问之间及不同的计算之间合理地重叠，从而减小计算核心空闲的时间。所以，根据需要，在调用核函数时可以指定使用多个线程：

#include<stdio.h>
#include "cuda_runtime.h"
__global__ void hello_from_gpu()
{
    printf("Hello World from the GPU!\n");
}
int main(void)
{
    hello_from_gpu << <2, 4 >> > ();
    cudaDeviceSynchronize();
    return 0;
}

2.2.2 使用线程索引
一个核函数中虽然可以指派如此巨大数目的线程数，但在执行时能够同时活跃（不活跃的线程处于等待状态）的线程数是由硬件（主要是CUDA核心数）和软件（即核函数中的代码）决定的。
每个线程在核函数中都有—个唯—的身份标识。由于我们用两个参数指定了线程数目，那么自然地，每个线程的身份可由两个参数确定。在核函数内部,程序是知道执行配置参数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。
下面,举一个具体的例子来进行说明。假如某个核函数的执行配置是<<<10000,256>>>，那么网格大小 gridDim.x 的值为10000，线程块大小 blockDim.x的值为256。线程块指标 blockIdx.x可以取 0-9999 范围内的值，而每一个线程块中的线程指标 threadIdx.x 可以取 0~255 范围内的值。当 blockIdx.x 等于0 时，所有256 个 threadIdxx的值对应第 0个线程块当 blockIdx.x 等于1时，所有 256个 threadIdx.x的值对应于第个线程块;以此类推,举个例子：

#include<stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
__global__ void hello_from_gpu()
{
    const int bid = blockIdx.x;
    const int tid = threadIdx.x;
    printf("Hello World from the block %d and thread%d\n",bid,tid);
}
int main(void)
{
    hello_from_gpu << <2, 4 >> > ();
    cudaDeviceSynchronize();
    return 0;
}

2.2.3 推广至多维网格
例如，如果要定义—个2×2×1的网格及3×2×1的线程块，可将执行配置中的grid_size和block_size分别定义为如下结构体变量:

dim3 grid_size(2,2);//等价于dim3 grid_size(2,2,1)
dim3 block_size(3,2);//等价于dim3 block_size(3,2,1)

#include<stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.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 the block-%d and thread-(%d,%d)\n",b,tx,ty);
}
int main(void)
{
    const dim3 block_size(2, 4);
    hello_from_gpu << <1, block_size >> > ();
    cudaDeviceSynchronize();
    return 0;
}

2.3.4 网络与线程块大小的限制
CUDA中对能够定义的网格大小和线程块大小做了限制。对任何从开普勒到图灵架构的 GPU 来说，网格大小在x,y,z 这 3个方向的最大取值分别为2^31-1，65535和65535;线程块大小在x，y和z这3个方向的最大取值分别为1024、1024 和 64。另外，还要求线程块总的大小，即 blockDim.x、blockDim.y和blockDim.z 的乘积不能大于 1024。也就是说，不管如何定义，一个线程块最多只能有 1024 个线程。这些限制是必须牢记的。

2.3 CUDA中的头文件

我们知道，在编写 C++ 程序时，往往需要在源文件中包含一些标准的头文件。读者也许注意到了，本章程序包含 C++ 的头文件 <stdio.h>，但并没有包含任何 CUDA 相关的头文件。CUDA 中也有一些头文件，但是在使用 nvcc 编译器驱动编译 .cu 文件时，将自动包含必要的 CUDA 头文件，如 <cuda.h>和<cuda_runtime.h>。因为 <cuda,h> 包含<stdlib.h>，故用nvcc 编译 CUDA 程序时甚至不需要在 .cu 文件中包含 <stdlib.h>。当然，用户依然可以在 cu 文件中包含 <stdlib.h>，因为(正确编写的)头文件不会在一个编译单元内被包含多次本书会从第 4 章开始使用一个用户自定义头文件。

2.4 用nvcc编译CUDA程序