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)
- kPrintf就是一个核函数,可以看到它的定义前有一个修饰词__global__,在任意函数前加上__global__就是定义该函数为核函数,就这么简单。
- 核函数的返回类型必须是void。
- 核函数的调用时比普通函数多了一个三括号<<<2, 3 >>>,这就是GPU多线程的意思,2*3=6 而例子中也打印了六个“Hello World!”, 接下来用<<<g, b >>>作为例子详细讲讲。
- GPU的目的就是多线程加速计算,所以在使用核函数的时候需要告知GPU需要调用多少核心多少线程来处理这个函数。
- 从小到大看,程序由b个线程和g个线程块来完成,所以使用的总线程就是b*g个。
- 每个线程块由b个线程组成,所以线程块的大小就是b。
- 所有的线程块组成一个网格(grid),而线程块有g个,所以g也被称为网格的网格大小。
- blockIdx和threadIdx为核函数的内置变量,就是为了区分哪个线程块的哪个线程来处理的当前数据,这里埋个伏笔为什么用.x。
- 还有两个常用的内置变量gridDim和blockDim,分别就是代表的网格大小和线程块大小,本例中分别就是2,3。
- 本例子中需要注意一点的是,核函数支持printf,也就是支持stdio.h但不支持iostream。
- 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来说:
- 网格大小在x,y,z三个维度上的最大值为2^32-1,65535,65535;
- 线程块大小在x,y,z三个维度上的最大值为1024,1024,64;
- 线程块的总大小不能大于1024,也就是三个维度的乘积要小于1024。