CUDA C编程入门笔记（1）

异构计算 = CPU（控制者）+ GPU（加速者）协同工作

“为任务选择适合它的架构”。

费脑子的事（逻辑、分支预测、IO读取）给 CPU。
费体力的事（大规模矩阵运算、像素处理）给 GPU。

用GPU输出：`Hello world!`#

准备工作#

检查CUDA编译器是否正确安装，以及路径：

1
which nvcc

通常结果是/usr/local/cuda/bin/nvcc。

检查GPU：

1
ls -l /dev/nv*

我这里的输出结果是：

1
crw-rw-rw- 1 root root     195,   0 11月 24 15:55 /dev/nvidia0
2
crw-rw-rw- 1 root root     195,   1 11月 24 15:55 /dev/nvidia1
3
crw-rw-rw- 1 root root     195,   2 11月 24 15:55 /dev/nvidia2
4
crw-rw-rw- 1 root root     195,   3 11月 24 15:55 /dev/nvidia3
5
crw-rw-rw- 1 root root     195, 255 11月 24 15:55 /dev/nvidiactl
6
crw-rw-rw- 1 root root     235,   0 11月 24 16:10 /dev/nvidia-uvm
7
crw-rw-rw- 1 root root     235,   1 11月 24 16:10 /dev/nvidia-uvm-tools
8

9
/dev/nvidia-caps:
10
total 0
11
cr-------- 1 root root 238, 1 11月 24 15:55 nvidia-cap1
12
cr--r--r-- 1 root root 238, 2 11月 24 15:55 nvidia-cap2

/dev/nvidia0 到 /dev/nvidia3表示有4张独立显卡。

第一个代码#

1
#include <stdio.h>
2

3
__global__ void hello(void){
4
    printf("Hello world from GPU !\n");
5
}
6

7
int main(){
8
    printf("Hello world from CPU !\n");
9
    hello<<<1,10>>> ();
10
    cudaDeviceReset();
11
    return 0;
12
}

我保存到test.cu中，然后运行nvcc test.cu -o test编译，./test来运行。运行结果：

1
$ ./test
2
Hello world from CPU !
3
Hello world from GPU !
4
Hello world from GPU !
5
Hello world from GPU !
6
Hello world from GPU !
7
Hello world from GPU !
8
Hello world from GPU !
9
Hello world from GPU !
10
Hello world from GPU !
11
Hello world from GPU !
12
Hello world from GPU !

可以看见，输出了一次CPU，十次GPU。

我的理解是，.cu文件，实际上是一个C程序，只是通过CUDA调用了GPU。而nvcc是一个编译器驱动，调用gcc编译CPU代码、以及NVDIA的自己的编译器来翻译GPU语言。

在CUDA编程里面，没有 __global__ 的函数，是普通C函数，比如main函数。

加上 __global__ 的函数（也叫 Kernel 核函数），由CPU来调用，通过GPU来执行。

因此前面CPU只输出一次。

1
int main(){
2
    printf("Hello world from CPU !\n");  // CPU运行，只有一个主线程，所以只打印一次
3
    // ...
4
}

后面的：

1
int main(){
2
    printf("Hello world from CPU !\n");
3
    hello<<<1,10>>> ();  // 启动1个线程块，每个块10个线程同时跑，所以打印10次
4
    // ...
5
}

三重尖括号意味着从主线程到设备端代码的调用，格式是函数名<<< 组的数量, 每组的人数 >>>(参数);

块和线程#

前面这段代码，如果运行hello<<< 1, 10 >>>和hello<<< 2, 5 >>>，效果应该是一样的。

但是其实应该是不一样的。

用下面这段修改后的代码，打印出块号blockIdx.x和线程号threadIdx.x。

1
#include <stdio.h>
2

3
__global__ void hello(void){
4
    printf("Block %d,\t Thread %d\n", blockIdx.x, threadIdx.x);
5
}
6

7
int main(){
8
    printf("Case : <<< 1, 10 >>>\n");
9
    hello<<<1,10>>> ();
10
    cudaDeviceSynchronize();
11

12
    printf("Case : <<< 2, 5 >>>\n");
13
    hello<<<2, 5>>> ();
14
    cudaDeviceSynchronize();
15
    return 0;
16
}

输出结果：

1
$ ./test
2
Case : <<< 1, 10 >>>
3
Block 0,         Thread 0
4
Block 0,         Thread 1
5
Block 0,         Thread 2
6
Block 0,         Thread 3
7
Block 0,         Thread 4
8
Block 0,         Thread 5
9
Block 0,         Thread 6
10
Block 0,         Thread 7
11
Block 0,         Thread 8
12
Block 0,         Thread 9
13
Case : <<< 2, 5 >>>
14
Block 0,         Thread 0
15
Block 0,         Thread 1
16
Block 0,         Thread 2
17
Block 0,         Thread 3
18
Block 0,         Thread 4
19
Block 1,         Thread 0
20
Block 1,         Thread 1
21
Block 1,         Thread 2
22
Block 1,         Thread 3
23
Block 1,         Thread 4

为什么要分块？这看上去运行hello<<< 1, 10 >>>和hello<<< 2, 5 >>>结果并没有差很多，都是多个同时运行。

应该是应该硬件限制，每个块的线程有上限。

而且同一个block内部，有共享内存。

可以提升可扩展性。让程序可以在不同的CPU上都能跑。

同步与重置#

cudaDeviceReset()、cudaDeviceSynchronize()的问题。

理论上，cudaDeviceReset是全部清空上下文，重新初始化整个环境。cudaDeviceSynchronize()是等前面的GPU干完了，不会清空上下文。

cudaDeviceSynchronize()意思是，等前面调用的GPU代码都干完了。如果没有调用这个，前面没有干完，CPU也会继续往下。因为CPU调用核函数是非阻塞的。

在前面运行这段代码的时候：

1
#include <stdio.h>
2

3
__global__ void hello(void){
4
    printf("Block %d,\t Thread %d\n", blockIdx.x, threadIdx.x);
5
}
6

7
int main(){
8
    printf("Case : <<< 1, 10 >>>\n");
9
    hello<<<1,10>>> ();
10
    cudaDeviceSynchronize();
11

12
    printf("Case : <<< 2, 5 >>>\n");
13
    hello<<<2, 5>>> ();
14
    cudaDeviceSynchronize();
15

16
    return 0;
17
}

显然可以注意到，第一次打印出来Case : <<< 1, 10 >>>后，要等一会儿（很明显），才能看见后续的一堆东西打印出来，而且是立即全部打印出来，第二次的Case : <<< 2, 5 >>>不需要等待。

第一次等待，应该是因为有一个冷启动。

如果把cudaDeviceSynchronize()全部换成cudaDeviceReset()，理论上应该是有区别的！毕竟说过了后者会清空上下文，但是前面这段运行起来，体感上没有区别，但是我感觉应该是有区别的。

用下面的代码测试一下时间：

1
#include <stdio.h>
2
#include <time.h>
3

4
__global__ void hello(void){
5
    // printf("Block %d,\t Thread %d\n", blockIdx.x, threadIdx.x);
6
}
7

8
int main(){
9
    int N = 100;
10

11
    clock_t start, end;
12
    double cost_time;
13

14
    hello<<<1, 1>>>();
15
    cudaDeviceSynchronize();
16

17
    printf("Test Synchronize: \n");
18

19
    start=clock();
20
    for (int i=0;i<N;i++){
21
        hello<<<1,1>>>();
22
        cudaDeviceSynchronize();
23
    }
24
    end=clock();
25
    cost_time=((double)(end-start))/CLOCKS_PER_SEC;
26
    printf("Synchronize cost %.4f seconds\n", cost_time);
27

28
    printf("Test Reset: \n");
29

30
    start=clock();
31
    for (int i=0;i<N;i++){
32
        hello<<<1,1>>>();
33
        cudaDeviceReset();
34
    }
35
    end=clock();
36
    cost_time=((double)(end-start))/CLOCKS_PER_SEC;
37
    printf("Reset cost %.4f seconds\n", cost_time);
38

39
    return 0;
40
}

测试结果：

1
~$ ./test
2
Test Synchronize:
3
Synchronize cost 0.0006 seconds
4
Test Reset:
5
Reset cost 11.5860 seconds

感觉还是很明显的。Reset慢很多。

CUDA编程结构#

Host (主机) = CPU + 内存。负责逻辑控制、读取文件、指挥 GPU 。

Device (设备) = GPU + 显存。专门负责并行的繁重计算。

这两者通过 PCI-Express 总线连接，但它们的内存是物理隔离的。这意味着，CPU 里的变量，GPU 没法直接读；GPU 算好的结果，CPU 也没法直接拿。

NOTE
可以把 CPU 上的变量名加前缀 h_ (host)，GPU 上的变量名加前缀 d_ (device)。

因为内存不互通，所以我们往往需要：

H->D：把数据从 CPU 内存（Host）拷贝到 GPU 显存（Device）。
核函数执行：CPU 发令，GPU 开始并行计算。（这是异步的，CPU 发完令可以去干别的，也可以等 GPU 算完。）
D->H：把算好的结果从 GPU 显存拷贝回 CPU 内存。

在GPU分配空间，用的是cudaMalloc函数：

1
cudaError_t cudaMalloc (void** devPtr, size_t size)

这个是在GPU显存上，申请一个size字节的连续内存，返回的是是否分配成功，然后分配的GPU显存的指针，会被存在devPtr指向的地址中。

主机和设备之间的数据传输，用的是cudaMemcpy函数：

1
cudaError_t cudaMemcpy (void* dst, const void* src, size_t count, cudaMemcpyKind kind)

从src这个地址开始，复制count字节到dst指向的地址。

最后的kind，用来表示复制方向，有四种：

cudaMemcpyHostToHost（疑似是为了对称的，主机搬到主机，和memcpy没啥区别）
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice

然后cudaError_t是一个枚举类型，成功就是cudaSuccess，失败就是cudaErrorMemoryAllocation。

可以用这个函数，把错误代码转成可以读的错误消息：

1
char* cudaGetErrorString (cudaError_t error)

向量相加#

下面是一个简单的向量相加代码。

1
#include <stdio.h>
2

3
__global__ void kernel_add(int *d_a, int *d_b, int *d_c, int n)
4
{
5
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
6
    if (idx < n)
7
        d_c[idx] = d_a[idx] + d_b[idx];
8
}
9

10
int main()
11
{
12
    int n = 1000;
13

14
    int h_a[n + 5], h_b[n + 5], h_c[n + 5];
15
    for (int i = 0; i < n; i++)
16
        h_a[i] = i, h_b[i] = i;
17

18
    size_t sz = n * (sizeof(int));
19

20
    int *d_a, *d_b, *d_c;
21
    cudaMalloc(&d_a, sz);
22
    cudaMalloc(&d_b, sz);
23
    cudaMalloc(&d_c, sz);
24

25
    cudaMemcpy(d_a, h_a, sz, cudaMemcpyHostToDevice);
26
    cudaMemcpy(d_b, h_b, sz, cudaMemcpyHostToDevice);
27

28
    int block_size = 256; // 每个块256个线程
29
    int block_cnt = (n + block_size - 1) / block_size;
30
    // up(n/block_size)
31

32
    kernel_add<<<block_cnt, block_size>>>(d_a, d_b, d_c, n);
33

34
    cudaMemcpy(h_c, d_c, sz, cudaMemcpyDeviceToHost);
35

36
    printf("result: \n");
37
    for (int i = 0; i < 4; i++)
38
        printf("%d\t+\t%d\t=\t%d\n", h_a[i], h_b[i], h_c[i]);
39
    printf("...\n");
40
    for (int i = n - 4; i < n; i++)
41
        printf("%d\t+\t%d\t=\t%d\n", h_a[i], h_b[i], h_c[i]);
42

43
    cudaFree(d_a);
44
    cudaFree(d_b);
45
    cudaFree(d_c);
46

47
    return 0;
48
}

运行：

1
nvcc test.cu -o test
2
./test

结果：

1
result:
2
0       +       0       =       0
3
1       +       1       =       2
4
2       +       2       =       4
5
3       +       3       =       6
6
...
7
996     +       996     =       1992
8
997     +       997     =       1994
9
998     +       998     =       1996
10
999     +       999     =       1998

这里，似乎加上(void **)强转比较好，但是我编译没有报错：

1
    cudaMalloc(&d_a, sz);
2
    cudaMalloc(&d_b, sz);
3
    cudaMalloc(&d_c, sz);

加上：

1
  cudaMalloc((void**)&d_a, size);
2
  // ...

错误检查宏#

如果我们把n改成10000，然后后面增加进程数：

1
    int block_size = 2048; // 每个块256个线程 ，改成2048
2
    int block_cnt = (n + block_size - 1) / block_size;

这样运行是直接错的：

1
result:
2
0       +       0       =       0
3
1       +       1       =       0
4
2       +       2       =       0
5
3       +       3       =       0
6
...
7
9996    +       9996    =       0
8
9997    +       9997    =       0
9
9998    +       9998    =       0
10
9999    +       9999    =       0

因为一个block最多就1024个线程。

但是在不知道这个问题的时候，可能不太好找到问题。

可以加一个这样的：

1
#define CHECK(call) \
2
{ \
3
    const cudaError_t error = call; \
4
    if (error != cudaSuccess) \
5
    { \
6
        printf("Error: %s:%d, ", __FILE__, __LINE__); \
7
        printf("code:%d, reason: %s\n", error, cudaGetErrorString(error)); \
8
        exit(1); \
9
    } \
10
}

然后代码改成下面这样：

1
#include <stdio.h>
2

3
#define CHECK(call) \
4
{ \
5
    const cudaError_t error = call; \
6
    if (error != cudaSuccess) \
7
    { \
8
        printf("Error: %s:%d, ", __FILE__, __LINE__); \
9
        printf("code:%d, reason: %s\n", error, cudaGetErrorString(error)); \
10
        exit(1); \
11
    } \
12
}
13

14
__global__ void kernel_add(int *d_a, int *d_b, int *d_c, int n)
15
{
16
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
17
    if (idx < n)
18
        d_c[idx] = d_a[idx] + d_b[idx];
19
}
20

21
int main()
22
{
23
    int n = 10000;
24

25
    int h_a[n + 5], h_b[n + 5], h_c[n + 5];
26
    for (int i = 0; i < n; i++)
27
        h_a[i] = i, h_b[i] = i;
28

29
    size_t sz = n * (sizeof(int));
30

31
    int *d_a, *d_b, *d_c;
32
    CHECK(cudaMalloc((void**)&d_a, sz));
33
    CHECK(cudaMalloc((void**)&d_b, sz));
34
    CHECK(cudaMalloc((void**)&d_c, sz));
35

36
    CHECK(cudaMemcpy(d_a, h_a, sz, cudaMemcpyHostToDevice));
37
    CHECK(cudaMemcpy(d_b, h_b, sz, cudaMemcpyHostToDevice));
38

39
    int block_size = 2048; // 每个块256个线程
40
    int block_cnt = (n + block_size - 1) / block_size;
41
    // up(n/block_size)
42

43
    kernel_add<<<block_cnt, block_size>>>(d_a, d_b, d_c, n);
44
    CHECK(cudaGetLastError());
45

46
    CHECK(cudaMemcpy(h_c, d_c, sz, cudaMemcpyDeviceToHost));
47

48
    printf("result: \n");
49
    for (int i = 0; i < 4; i++)
50
        printf("%d\t+\t%d\t=\t%d\n", h_a[i], h_b[i], h_c[i]);
51
    printf("...\n");
52
    for (int i = n - 4; i < n; i++)
53
        printf("%d\t+\t%d\t=\t%d\n", h_a[i], h_b[i], h_c[i]);
54

55
    CHECK(cudaFree(d_a));
56
    CHECK(cudaFree(d_b));
57
    CHECK(cudaFree(d_c));
58

59
    return 0;
60
}

就可以看见报错了：

1
$ ./test
2
Error: test.cu:44, code:9, reason: invalid configuration argument

同步问题#

前面说了，核函数启动是异步的。

但是我并没有调用cudaDeviceSynchronize()，那么，会导致printf在GPU算完之前就输出吗？

不会。

因为后面有个：

1
  cudaMemcpy(h_c, d_c, sz, cudaMemcpyDeviceToHost);

这个函数以同步方式执行，因为在 cudaMemcpy 函数返回以及传输操作完成之前主机应用程序是阻塞的。

在运行这个函数的时候，CUDA会先检查GPU是否在忙碌，只有前面都忙完了，才会拷贝数据。

用GPU输出：Hello world!#

准备工作#

第一个代码#

块和线程#

同步与重置#

CUDA编程结构#

向量相加#

错误检查宏#

同步问题#

用GPU输出：`Hello world!`#