CUDA学习第一天: 基础概念扫盲
CUDA基础学习开始,以下内容是来自于各个博客和官方文档,如果我写不够全面,大家可以去往相应的参考链接进行学习。
参考链接:https://zhuanlan.zhihu.com/p/34587739
参考链接:https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#application-compatibility
参考链接:https://www.cnblogs.com/skyfsm/p/9673960.html
CUDA编程模型是一个异构模型,需要CPU和GPU协同工作。
在CUDA中,host和device是两个重要的概念,我们用host
指代CPU及其内存,而用device
指代GPU及其内存。
CUDA程序中既包含host程序, 又包含device程序,他们分别在CPU和GPU上运行,同时, host与device之间可以进行通信,这样他们之间可以进行数据拷贝。
-
典型的CUDA程序的执行流程
1.分配host内存,并进行数据初始化;---- CPU初始化
2.分配device内存,并从host将数据拷贝到device上; — GPU初始化
3.调用CUDA的核函数在device上完成指定的运算;----GPU并行运算(核心)
4.将device上的运算结果拷贝到host上;— 将GPU结果传回CPU
5.释放device和host上分配的内存。---- 初始化清空
-
CUDA中的核函数 — Kernel
1.kernel是在device上线程中并行执行的函数,核函数用
__global__符号声明
,在调用时需要用<<<grid, block>>>
来指定kernel要执行的线程数量;2.在CUDA中,每一个线程都要执行核函数,并且
每个线程会分配一个唯一的线程号thread ID
,这个ID值可以通过核函数的内置变量threadIdx
来获得。 -
CPU 与 GPU的结构不一, 所以调用时需要区分host 与device上的代码
在CUDA中是通过函数类型限定词开区别host和device上的函数,主要的三个函数类型限定词如下:
-
global:在device上执行,从host中调用(一些特定的GPU也可以从device上调用),返回类型必须是void,不支持可变参数参数,不能成为类成员函数。注意用__global__定义的kernel是异步的,这意味着host不会等待kernel执行完就执行下一步。
-
device:在device上执行,单仅可以从device中调用,不可以和__global__同时用;
-
host:在host上执行,仅可以从host上调用, 一般省略不写, 不可以和__global__同时用,但可和__device__,此时函数会在device和host都编译;
-
-
深刻理解Kernel
- 网格 – grid
GPU上有很多并行化的轻量级线程。kernel在device上执行时实际上是启动很多线程,一个kernel所启动的所有线程称为一个网格(grid),同一个网格上的线程共享相同的全局内存空间,grid是线程结构的第一层次。
- 线程块 – block
网格可以分为很多线程块(block),一个线程块里面包含很多线程,这是第二个层次。
- 数据类型定义 — dim3
dim3可以看成是包含三个无符号整数(x,y,z)成员的结构体变量,在定义时,缺省值初始化为1; 对于图中结构(主要水平方向为x轴)。
定义的grid和block示例如下图所示,kernel在调用时也必须通过执行配置<<<grid, block>>>来指定kernel所使用的线程数及结构。
dim3 grid(3, 2);
dim3 block(5, 3);
kernel_fun<<< grid, block >>>(prams...);
所以,一个线程需要两个内置的坐标变量**(blockIdx,threadIdx)**来唯一标识,它们都是dim3类型变量
,其中blockIdx
指明线程所在grid中的位置
,而threaIdx
指明线程所在block中的位置
。
-
线程块上的线程变量
一个线程块上的线程是放在同一个流式多处理器(SM)上的,但是单个SM的资源有限,这导致线程块中的线程数是有限制的,现代GPUs的线程块可支持的线程数可达1024个。
有时候,我们要知道一个线程在blcok中的全局ID,此时就必须还要知道block的组织结构。
- 内置变量 – blockDim : 获取线程块各个维度的大小
对于一个2-dim的block(Dx, Dy), 线程(x, y)的ID值为:(x+y*Dx);
对于一个3-dim的block(Dx, Dy, Dz), 线程(x, y, z)的ID值为:(x+yDx+ zDx*Dy);
- 内置变量 – gridDim: 获取网格块各个维度的大小
-
CUDA编程示例1
kernel的这种线程组织结构天然适合vector,matrix等运算,如我们将利用上图2-dim结构实现两个矩阵的加法,每个线程负责处理每个位置的两个元素相加,代码如下所示。线程块大小为(16, 16),然后将N*N大小的矩阵均分为不同的线程块来执行加法运算。
// Kernel定义
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
int j = blockIdx.y * blockDim.y + threadIdx.y;
if (i < N && j < N)
C[i][j] = A[i][j] + B[i][j];
}
int main()
{
...
// Kernel 线程配置
dim3 threadsPerBlock(16, 16);
dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
// kernel调用
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
...
}