c++ nvcc编译CUDA程序入门示例

nvcc

nvcc是NVIDIA CUDA Compiler，用来编译host和device程序。

这里的术语：

host：指CPU及其内存
device：指GPU及其内存

使用nvcc，就可以编译CUDA程序，CUDA程序包括host代码和device代码。

在安装CUDA Toolkit后，nvcc内含其中。

注意要安装与显卡版本匹配的CUDA Toolkit。

我的nvcc版本：

% nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2019 NVIDIA Corporation
Built on Wed_Oct_23_19:24:38_PDT_2019
Cuda compilation tools, release 10.2, V10.2.89

hello, world

作为第一个示例，非“hello, world”莫属。

编写以下代码，保存为demo.cpp：

#include <iostream>
#include <cuda_runtime.h>using namespace std;int main()
{cout << "hello, world!" << endl;return 0;
}

那么，该文件可以使用以下方式编译：

g++ demo.cpp

或者：

nvcc demo.cpp

它们都生成a.out，都输出“hello, world!”。

ok，这是第一次成功使用nvcc。

加入device代码

修改代码：

#include <iostream>
#include <cuda_runtime.h>using namespace std;__global__ void demo(void)
{}int main()
{demo<<<1,1>>>();cout << "hello, world!" << endl;return 0;
}

其中，__global__是CUDA C/C++关键字，指定该段代码在device上执行，并且从host上调用。

再用这两种方式编译时，都编译失败了。使用g++编译失败是必然，使用nvcc编译报错如下：

% nvcc demo.cpp
hello.cpp: In function ‘int main()’:
hello.cpp:12:11: error: expected primary-expression before ‘<’ tokendemo<<<1,1>>>();^
hello.cpp:12:17: error: expected primary-expression before ‘>’ tokendemo<<<1,1>>>();^
hello.cpp:12:19: error: expected primary-expression before ‘)’ tokendemo<<<1,1>>>();

这是因为nvcc使用文件扩展名来确定如何处理文件的内容。

如果文件中包含CUDA语法，则文件扩展名必须为.cu，否则nvcc会将未修改的文件直接传递给主机编译器，从而导致语法错误。

所以，把demo.cpp文件名修改为demo.cu即可使用nvcc编译通过。

输出同上。

ok，这是第一次成功使用nvcc编译device代码。

做点有意义的事

上个示例中，在device上什么也没有做，这次做一下矩阵加法运算，代码如下：

#include <iostream>
#include <cuda_runtime.h>using namespace std;__global__ void add(int *a, int *b, int *c, int n)
{int index = threadIdx.x + blockIdx.x * blockDim.x;if (index < n) {c[index] = a[index] + b[index];}
}void random_ints(int* a, int n)
{for (int i = 0; i < n; ++i)a[i] = rand();
}int main()
{int n = (2048 * 2048);int threads_per_block = 512;int *a, *b, *c;int *d_a, *d_b, *d_c;int size = n * sizeof(int);cudaMalloc((void**)&d_a, size);cudaMalloc((void**)&d_b, size);cudaMalloc((void**)&d_c, size);a = (int*)malloc(size);random_ints(a, n);b = (int*)malloc(size);random_ints(b, n);c = (int*)malloc(size);cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);add<<<(n + threads_per_block - 1)/threads_per_block, threads_per_block>>>(d_a, d_b, d_c, n);cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);cudaFree(d_a);cudaFree(d_b);cudaFree(d_c);for (int i = 0; i < n; ++i) {cout << c[i] << ",";}cout << endl;cout << cudaGetErrorString(cudaGetLastError()) << endl;free(a);free(b);free(c);return 0;
}

示例虽小，五脏俱全。

流程梳理如下：

准备待处理数据
在device上分配存储空间
把数据从host拷贝到device
执行device运算
把结果从device拷贝回host
结束后释放device空间
回收host资源

更具体的细节后续博客中介绍。

小结

CUDA的入门还是比较简单的，根据这个流程做就行。

需要特别注意的是CUDA程序的文件名必须以cu后缀，这个有点反Linux，哈哈。

还有很多细节需要深刻理解，否则GPU的优势难以发挥，有点大炮打蚊子的嫌疑了。