CUDA核函数
在GPU上执行的函数称为CUDA核函数(Kernel Function),核函数会被GPU上多个线程执行,我们可以在核函数中获取当前线程的ID。
// CUDA核函数的定义
__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = a[i] + b[i];
}
// CUDA核函数调用
addKernel<<<Dg,Db, Ns, S>>>(c, a, b);
可以总结出CUDA核函数的使用方式:
- 使用关键字
__global__
来标识,在CPU上调用,在GPU上执行,返回值为void - 使用
<<< >>>
来指定线程索引方式 - 核函数相对于CPU是异步的,在核函数执行完之前就会返回,这样CPU可以不用等待核函数的完成,继续执行后续代码
- 核函数不支持可变参数,不支持静态变量,不支持函数指针
线程配置
这是刚刚接触GPU编程最为困惑的地方,到底应该如何去索引GPU的线程?首先要对GPU中的基本概念有所了解,可参考另一篇文章《GPU中的基本概念》
在Host端核函数的调用方式为:
kernel<<<Dg, Db, Ns, S>>>(param list);
其中,
-
Dg
:int型或者dim3类型(x,y,z),用于定义一个Grid中Block是如何组织的,如果是int型,则表示一维组织结构 -
Db
:int型或者dim3类型(x,y,z),用于定义一个Block中Thread是如何组织的,如果是int型,则表示一维组织结构 -
Ns
:size_t类型,可缺省,默认为0; 用于设置每个block除了静态分配的共享内存外,最多能动态分配的共享内存大小,单位为byte。 0表示不需要动态分配。 -
S
:cudaStream_t类型,可缺省,默认为0。 表示该核函数位于哪个流。
Grid/Block/Thread都是软件的组织结构,并不是硬件的,因此理论上我们可以以任意的维度(一维、二维、三维)去排列Thread;在硬件上就是一个个的SP,并没有维度这一说,只是软件上抽象成了具有维度的概念。
当使用dim3
类型时,比如:
dim3 grid(3,2,1), block(4,3,1);
kernel_name<<<grid, block>>>(...);
表示一个Grid中有3x2x1=6个Block,在(x,y,z)三个方向上的排布方式分别是3、2、1;一个Block中有4x3x1=12个Thread,在(x,y,z)三个方向上的排布方式分别是4、3、1。
当使用int
类型时,表示一维排布,比如:
kernel_name<<<5,8>>>(...);
表示一个Grid中有5个Block,在(x,y,z)三个方向上的排布方式分别是5、1、1;一个Block中有8个Thread,在(x,y,z)三个方向上的排布方式分别是8、1、1。
在CUDA上可以使用内置变量来获取Thread ID和Block ID:
-
threadIdx.[x, y, z]
表示Block内Thread的编号 -
blockIdx.[x, y, z]
表示Gird内Block的编号 -
blockDim.[x, y, z]
表示Block的维度,也就是Block中每个方向上的Thread的数目 -
gridDim.[x, y, z]
表示Gird的维度,也就是Grid中每个方向上Block的数目
下面我们举几个例子来说明
一维Grid 一维Block
kernel_name<<<4, 8>>>(...)
具体的线程索引方式如下图所示,blockIdx从0到3,threadIdx从0到7.
当我们要计算下图中红色的Thread的索引时,可以看出,它的blockIdx.x是2,threadIdx.x是1,因此它的threadId索引计算方式为:
int threadId = blockIdx.x * blockDim.x + threadIdx.x = 2 * 4 + 1 = 9
二维Grid 二维Block
无论是几维的,索引的原则是一样的,先求出这个Thread前面的所有Block中Thread的数量,再求出该Thread在本Block中的序号,两个相加即可。
下面为了画图方便,我们以将Block的维度设为(4,1,1),其实是一维Block了,但计算公式是一样的:
dim grid(4,1,1), block(2,2,1);
kernel_name<<<grid, block>>>(...)
需要注意的是,二维排序中,Thread(0,1)表示第1行第0列的Thread,这跟我们传统中理解的横坐标和纵坐标不太一样;我们定义grid(4,2)表示第一维度有4个索引值,第二个维度有2个索引值,即2行4列
具体排列方式如下图所示,blockidx从0到3,Threadidx从(0,0)到(1,2)
Idx的具体索引方式如下公式,idx
表示当前Thread在全局索引中的序号
int blockId = blockIdx.x + blockId.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y *blockDim.x) + threadIdx.x;
三维Grid 三维Block
三维的图画起来有点复杂,我们就不画图,直接给出计算公式:
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;
int threadIc = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x) + threadIdx.x;
三维Grid 三维Thread是参数最多的组合,对于其他维度的组合,只需要把相应的参数置1即可,下面我们将各种组合总结一下,以后大家可以直接使用:(一维的话我们默认是使用x维度,二维的话默认使用(x,y)维度)
- 一维Grid 一维Block
blockId = blockIdx.x
threadId = blockIdx.x *blockDim.x + threadIdx.x
- 一维Grid 二维Block
blockId = blockIdx.x
threadId = blockIdx.x * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x
- 一维Grid 三维Block
blockId = blockIdx.x
threadId = blockIdx.x * blockDim.x * blockDim.y * blockDim.z
+ threadIdx.z * blockDim.y * blockDim.x
+ threadIdx.y * blockDim.x + threadIdx.x
- 二维Grid 一维Block
int blockId = blockIdx.y * gridDim.x + blockIdx.x;
int threadId = blockId * blockDim.x + threadIdx.x;
- 二维Grid 二维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;
- 二维Grid 三维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x) + threadIdx.x;
- 三维Grid 一维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * blockDim.x + threadIdx.x;
- 三维Grid 二维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * (blockDim.x * blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;
- 三维Grid 三维Block
int blockId = blockIdx.x + blockIdx.y * gridDim.x + gridDim.x * gridDim.y * blockIdx.z;
int threadId = blockId * (blockDim.x * blockDim.y * blockDim.z)
+ (threadIdx.z * (blockDim.x * blockDim.y))
+ (threadIdx.y * blockDim.x) + threadIdx.x;