CUDA学习笔记——入门例子
虽然配置CUDA的目的主要是支持PyTorch的机器学习,但是如果需要直接手写CUDA代码,并在Windows上使用nvcc顺利编译和执行纯CUDA程序,还需要Visual Studio相关工具的支持。(MINGW似乎不太行,或者说鼓捣起来太麻烦) 鉴于我在电脑上安装了VS2022,并且CUDA自动给VS2022加了插件适配,因此尝试学一点CUDA,并写几个简单的HelloWorld程序。
基本概念
CUDA编程模型是一个异构模型,需要CPU和GPU协同工作,通常CPU负责处理逻辑复杂的串行程序,而GPU重点处理数据密集型的并行计算程序,从而发挥最大功效。 在CUDA中,用host指代CPU及其内存,用device指代GPU及其内存。CUDA程序中既包含host程序,又包含device程序,它们分别在CPU和GPU上运行。 host与device之间可以进行通信,在它们之间需要进行数据拷贝。 一个典型的CUDA程序的执行流程如下:
- 分配host内存,并进行数据初始化;
- 分配device内存,并从host将数据拷贝到device上;
- 调用CUDA的核函数在device上完成指定的运算;
- 将device上的运算结果拷贝到host上;
- 释放device和host上分配的内存。
在CUDA编程中需要区分在host和device上执行的代码,对于函数需要通过特殊限定词来区分,主要的三个函数类型限定词如下:
__global__
:函数在device执行,仅在host上调用,这种标记的函数又称为核函数(kernel)__device__
:函数在device执行,仅在device上调用__host__
:函数在host执行,仅在host上调用
注意:
__host__
是不含任何限定词时的默认情形,此时调用和执行都在host进行,相当于常规的C语言函数,因此很少单独使用这个标记__host__
可以和__device__
组合使用,此时代码将在host和device各自编译一份
对于__global__
和__device__
,由于需要在device执行或调用,因此相比于普通的C语言函数由更多的要求:
- 不支持递归
- 不支持在函数体内声明静态变量
- 不支持可变长参数
最重要的概念是__global__
标记的核函数,对核函数还有更多的要求:
__global__
限定词不能与其它限定词组合使用- 返回类型必须是void,因此获取返回值必须通过函数参数指针实现
- 不能作为类的成员函数
- ...
核函数的定义和调用示例如下:
核函数由host调用,但是由device上线程负责并行执行,一个最简单的核函数如下,这里不仅返回值是void,甚至没有函数参数,只是打印一个字符串。
1
2
3
4__global__ void hello_world_kernel()
{
printf("Hello GPU!\n");
}在调用核函数时,需要加上
<<<grid, block>>>
来指定执行核函数的线程数量,<<<grid,block>>>
表示启动一个grid个线程块,每个线程块中包含block个线程的CUDA网格,其中所有的线程一起执行核函数,示例如下,这里cudaDeviceReset()
用于重置当前设备上的所有资源状态,建议在程序的最后调用这个函数,以确保释放所有由CUDA运行时分配的资源1
2
3
4
5int main()
{
hello_world_kernel <<<1,8>>>();
cudaDeviceReset();
}每一个线程在各自执行核函数时,会分配一个唯一的线程号thread ID,这个ID值可以通过核函数的内置变量threadIdx来获得,从而可以让不同的线程根据拓扑结构或线程号相互合作,例如
1
2
3
4
5__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = a[i] + b[i];
}
注意:核函数都是异步执行,即当核函数在GPU中唤醒后,不论函数是否执行完毕,控制权都会立刻返回给host。
例子
例一
创建一个名为hello_world_gpu.cu的文件,写入如下内容,可以看出它在语法上基本上是C语言加上了一些定制扩展,因此nvcc需要MSVC的支持
1 |
|
编译执行CUDA程序需要特殊的指令,基于VS的GUI的操作略,对于命令行操作,在Developer
PowerShell for Visual Studio 2022中执行如下语句 1
$ nvcc hello_world_gpu.cu -o hello_world_gpu.exe
可以编译得到 hello_world_gpu.exe,运行结果如下
1 | Hello GPU! |
这里编译必须在Visual Studio配置好的命令行环境中,否则会报错找不到cl.exe等,但是编译之后的程序运行并不需要特殊的命令行环境。
例二
第一个程序过于简单,没有返回值,下面的例子是VS2022中自带的项目,只有一个kernel.cu文件,完整代码如下
1 |
|
下面来消化理解一下这个例子。
首先关注核函数的部分,这个程序的核心功能就是让GPU来计算一个向量加法,把每一个元素的加法分配给一个线程来执行,这里的 c 是返回值向量,a 和 b 是输入向量。
1 | __global__ void addKernel(int *c, const int *a, const int *b) |
然后关注main函数部分,首先定义了向量a和向量b,以及存储结果的向量c,然后调用addWithCuda
函数执行,输出计算结果c,最后释放相应资源。
1 | int main() |
剩下的部分是addWithCuda
函数,它就是在给GPU分配这个向量加法任务并处理调用GPU计算前后的细节:
cudaMalloc
申请GPU内存,内存大小对应a,b,c的尺寸cudaMemcpy
将a,b,c
从CPU内存拷贝到GPU内存中,得到副本dev_a,dev_b,dev_c
- 执行核函数
addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
- 将计算结果的c从GPU内存拷贝到CPU内存中,并返回
这是CUDA官方提供的例子,代码的健壮性很高,因此包含了很多的错误处理,但是主要内容还是简单清晰的。