虽然配置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程序的执行流程如下:

  1. 分配host内存,并进行数据初始化;
  2. 分配device内存,并从host将数据拷贝到device上;
  3. 调用CUDA的核函数在device上完成指定的运算;
  4. 将device上的运算结果拷贝到host上;
  5. 释放device和host上分配的内存。

在CUDA编程中需要区分在host和device上执行的代码,对于函数需要通过特殊限定词来区分,主要的三个函数类型限定词如下:

  1. __global__:函数在device执行,仅在host上调用,这种标记的函数又称为核函数(kernel)

  2. __device__:函数在device执行,仅在device上调用

  3. __host__:函数在host执行,仅在host上调用

注意:

  • __host__是不含任何限定词时的默认情形,此时调用和执行都在host进行,相当于常规的C语言函数,因此很少单独使用这个标记
  • __host__可以和__device__组合使用,此时代码将在host和device各自编译一份

对于__global____device__,由于需要在device执行或调用,因此相比于普通的C语言函数由更多的要求:

  • 不支持递归
  • 不支持在函数体内声明静态变量
  • 不支持可变长参数

最重要的概念是__global__标记的核函数,对核函数还有更多的要求:

  • __global__限定词不能与其它限定词组合使用
  • 返回类型必须是void,因此获取返回值必须通过函数参数指针实现
  • 不能作为类的成员函数
  • ...

核函数的定义和调用示例如下:

  1. 核函数由host调用,但是由device上线程负责并行执行,一个最简单的核函数如下,这里不仅返回值是void,甚至没有函数参数,只是打印一个字符串。

    1
    2
    3
    4
    __global__ void hello_world_kernel()
    {
    printf("Hello GPU!\n");
    }

  2. 在调用核函数时,需要加上<<<grid, block>>>来指定执行核函数的线程数量,<<<grid,block>>> 表示启动一个grid个线程块,每个线程块中包含block个线程的CUDA网格,其中所有的线程一起执行核函数,示例如下,这里cudaDeviceReset()用于重置当前设备上的所有资源状态,建议在程序的最后调用这个函数,以确保释放所有由CUDA运行时分配的资源

    1
    2
    3
    4
    5
    int main()
    {
    hello_world_kernel <<<1,8>>>();
    cudaDeviceReset();
    }

  3. 每一个线程在各自执行核函数时,会分配一个唯一的线程号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
2
3
4
5
6
7
8
9
10
11
12
#include <stdio.h>

__global__ void hello_world_kernel()
{
printf("Hello GPU!\n");
}

int main()
{
hello_world_kernel <<<1,8>>>();
cudaDeviceReset();
}

编译执行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
2
3
4
5
6
7
8
Hello GPU!
Hello GPU!
Hello GPU!
Hello GPU!
Hello GPU!
Hello GPU!
Hello GPU!
Hello GPU!

这里编译必须在Visual Studio配置好的命令行环境中,否则会报错找不到cl.exe等,但是编译之后的程序运行并不需要特殊的命令行环境。

例二

第一个程序过于简单,没有返回值,下面的例子是VS2022中自带的项目,只有一个kernel.cu文件,完整代码如下

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

__global__ void addKernel(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = a[i] + b[i];
}

int main()
{
const int arraySize = 5;
const int a[arraySize] = { 1, 2, 3, 4, 5 };
const int b[arraySize] = { 10, 20, 30, 40, 50 };
int c[arraySize] = { 0 };

// Add vectors in parallel.
cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addWithCuda failed!");
return 1;
}

printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
c[0], c[1], c[2], c[3], c[4]);

// cudaDeviceReset must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaDeviceReset();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceReset failed!");
return 1;
}

return 0;
}

// Helper function for using CUDA to add vectors in parallel.
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
int *dev_a = 0;
int *dev_b = 0;
int *dev_c = 0;
cudaError_t cudaStatus;

// Choose which GPU to run on, change this on a multi-GPU system.
cudaStatus = cudaSetDevice(0);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaSetDevice failed! Do you have a CUDA-capable GPU installed?");
goto Error;
}

// Allocate GPU buffers for three vectors (two input, one output) .
cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}

cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}

cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed!");
goto Error;
}

// Copy input vectors from host memory to GPU buffers.
cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}

cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}

// Launch a kernel on the GPU with one thread for each element.
addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

// Check for any errors launching the kernel
cudaStatus = cudaGetLastError();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
goto Error;
}

// cudaDeviceSynchronize waits for the kernel to finish, and returns
// any errors encountered during the launch.
cudaStatus = cudaDeviceSynchronize();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
goto Error;
}

// Copy output vector from GPU buffer to host memory.
cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaMemcpy failed!");
goto Error;
}

Error:
cudaFree(dev_c);
cudaFree(dev_a);
cudaFree(dev_b);

return cudaStatus;
}

下面来消化理解一下这个例子。

首先关注核函数的部分,这个程序的核心功能就是让GPU来计算一个向量加法,把每一个元素的加法分配给一个线程来执行,这里的 c 是返回值向量,a 和 b 是输入向量。

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];
}

然后关注main函数部分,首先定义了向量a和向量b,以及存储结果的向量c,然后调用addWithCuda函数执行,输出计算结果c,最后释放相应资源。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
int main()
{
const int arraySize = 5;
const int a[arraySize] = { 1, 2, 3, 4, 5 };
const int b[arraySize] = { 10, 20, 30, 40, 50 };
int c[arraySize] = { 0 };

// Add vectors in parallel.
cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "addWithCuda failed!");
return 1;
}

printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
c[0], c[1], c[2], c[3], c[4]);

// cudaDeviceReset must be called before exiting in order for profiling and
// tracing tools such as Nsight and Visual Profiler to show complete traces.
cudaStatus = cudaDeviceReset();
if (cudaStatus != cudaSuccess) {
fprintf(stderr, "cudaDeviceReset failed!");
return 1;
}

return 0;
}

剩下的部分是addWithCuda函数,它就是在给GPU分配这个向量加法任务并处理调用GPU计算前后的细节:

  • cudaMalloc申请GPU内存,内存大小对应a,b,c的尺寸
  • cudaMemcpya,b,c从CPU内存拷贝到GPU内存中,得到副本dev_a,dev_b,dev_c
  • 执行核函数 addKernel<<<1, size>>>(dev_c, dev_a, dev_b);
  • 将计算结果的c从GPU内存拷贝到CPU内存中,并返回

这是CUDA官方提供的例子,代码的健壮性很高,因此包含了很多的错误处理,但是主要内容还是简单清晰的。