CUDA入门学习笔记：二 第一个核函数

CUDA入门学习笔记：二 第一次写核函数

---

一、什么是核函数

一个程序的流程是由CPU控制的，GPU只是起到计算的作用，CPU指定GPU来计算的函数就是核函数。

CUDA编程中核函数需要写在.cu文件中，一般的C++程序可以写在.cpp或者.cu中。也就是说.cu文件是支持普通C++程序的，编译时CUDA会针对普通C++程序调用C++的编译器，CUDA程序由CUDA自己编译完成。

二、初次遇到核函数

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>

__global__ void kPrintf()
{
	int the_block = blockIdx.x;
	int the_thread = threadIdx.x;
	printf("Hello World! (this block: %d   this thread: %d)\n", the_block, the_thread);
}

int main()
{
	kPrintf<<<2, 3 >>> ();
	cudaDeviceSynchronize();

	return 0;
}

上图就是一个使用CUDA实现打印Hello World的例子，运行结果如下：

Hello World! (this block: 0 this thread: 0)
Hello World! (this block: 0 this thread: 1)
Hello World! (this block: 0 this thread: 2)
Hello World! (this block: 1 this thread: 0)
Hello World! (this block: 1 this thread: 1)
Hello World! (this block: 1 this thread: 2)

1. kPrintf就是一个核函数，可以看到它的定义前有一个修饰词__global__，在任意函数前加上__global__就是定义该函数为核函数，就这么简单。
2. 核函数的返回类型必须是void。
3. 核函数的调用时比普通函数多了一个三括号<<<2, 3 >>>，这就是GPU多线程的意思，2*3=6 而例子中也打印了六个“Hello World！”， 接下来用<<<g, b >>>作为例子详细讲讲。
4. GPU的目的就是多线程加速计算，所以在使用核函数的时候需要告知GPU需要调用多少核心多少线程来处理这个函数。
5. 从小到大看，程序由b个线程和g个线程块来完成，所以使用的总线程就是b*g个。
6. 每个线程块由b个线程组成，所以线程块的大小就是b。
7. 所有的线程块组成一个网格（grid），而线程块有g个，所以g也被称为网格的网格大小。
8. blockIdx和threadIdx为核函数的内置变量，就是为了区分哪个线程块的哪个线程来处理的当前数据，这里埋个伏笔为什么用.x。
9. 还有两个常用的内置变量gridDim和blockDim，分别就是代表的网格大小和线程块大小，本例中分别就是2，3。
10. 本例子中需要注意一点的是，核函数支持printf，也就是支持stdio.h但不支持iostream。
11. cudaDeviceSynchronize()是用来刷新缓存区用的，如果不进行刷新，打印的结果会卡在缓冲区，所以就无法输出信息。

三、令人头大的多维网格

上个例子中可以看到blockIdx和threadIdx都使用了成员变量.x，这里解释一下blockIdx和threadIdx的数据类型，他们都是uint3类型的变量，uint3的定义如下：

struct __device_builtin__ uint3
{
    unsigned int x, y, z;
};

可以看到blockIdx和threadIdx都有三个成员变量，gridDim和blockDim是dim3类型的变量也有xyz三个成员变量，只不过比uint3多了几个内置的函数。那么yz是做什么的？

还记得核函数的调用时需要加上的<<<g, b >>>吗，其实g和b不仅可以是一个数，还可以是一个最多表示三维的dim3变量。

例如

dim3 block_size(2, 4, 1);
kPrintf<<<2, block_size>>> ();

这样就是表示用一个网格大小为2，线程块大小为2* 4* 1的一个三维线程块来运行核函数。

可以看到block_size的第三维为1，此时就和一个2* 4的二维线程块没有区别，所以在这种情况下也可以写成如下形式：

dim3 block_size(2, 4);
kPrintf<<<2, block_size>>> ();

其实不管是二维还是三维最终都可以转为一维，线程ID如下：

threadIdx.z * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x

线程块ID如下：

blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x

搞清楚了降维索引后其实多维网络非常简单明了。

四、线程的使用规范

对于目前主流的GPU来说：

1. 网格大小在x,y,z三个维度上的最大值为2^32-1，65535，65535；
2. 线程块大小在x,y,z三个维度上的最大值为1024，1024，64；
3. 线程块的总大小不能大于1024，也就是三个维度的乘积要小于1024。