Skip to content

Commit ebaa3de

Browse files
archibatearchibate
committed
addcuda1
1 parent e389522 commit ebaa3de

File tree

5 files changed

+343
-27
lines changed

5 files changed

+343
-27
lines changed

‎slides/moderncuda/CMakeLists.txt

+19-3
Original file line numberDiff line numberDiff line change
@@ -2,10 +2,26 @@ cmake_minimum_required(VERSION 3.12)
22

33
set(CMAKE_CXX_STANDARD 20)
44
set(CMAKE_CUDA_STANDARD 20)
5-
set(CMAKE_CUDA_SEPARABLE_COMPILATION ON)
5+
set(CMAKE_CUDA_SEPARABLE_COMPILATION OFF)
66
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr --expt-extended-lambda")
7-
if (NOT DEFINED CMAKE_CUDA_ARCHITECTURES AND CMAKE_VERSION VERSION_GREATER_EQUAL 3.24)
8-
set(CMAKE_CUDA_ARCHITECTURES native)
7+
8+
if (NOT DEFINED CMAKE_CUDA_ARCHITECTURES)
9+
# try to set CMAKE_CUDA_ARCHITECTURES to native card version
10+
if (CMAKE_VERSION VERSION_GREATER_EQUAL 3.24)
11+
set(CMAKE_CUDA_ARCHITECTURES native)
12+
else()
13+
message(STATUS "Running __nvcc_device_query for CMAKE_CUDA_ARCHITECTURES")
14+
execute_process(COMMAND "__nvcc_device_query"
15+
RESULT_VARIABLE NVCC_DEVICE_QUERY_RESULT
16+
OUTPUT_VARIABLE CMAKE_CUDA_ARCHITECTURES
17+
OUTPUT_STRIP_TRAILING_WHITESPACE)
18+
if (NVCC_DEVICE_QUERY_RESULT EQUAL 0)
19+
message(STATUS "CMAKE_CUDA_ARCHITECTURES set to ${CMAKE_CUDA_ARCHITECTURES}")
20+
else()
21+
message(WARNING "CMAKE_CUDA_ARCHITECTURES not defined and __nvcc_device_query failed.")
22+
unset(CMAKE_CUDA_ARCHITECTURES)
23+
endif()
24+
endif()
925
endif()
1026

1127
project(moderncuda LANGUAGES CXX CUDA)

‎slides/moderncuda/README.md

+280-20
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,22 @@
11
# 现代 C++ 的 CUDA 编程
22

3-
## 安装 NVIDIA 驱动
3+
参考资料：https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html
4+
5+
## 配置 CUDA 开发环境
6+
7+
硬件方面建议使用至少 GTX 1060 以上显卡，但是更老的显卡也可以运行。
8+
9+
软件方面则可以尽可能最新，以获得 CUDA C++20 支持，我安装的版本是 CUDA 12.5。
10+
11+
以下仅演示 Arch Linux 中安装 CUDA 的方法，因为 Arch Linux 官方源中就自带 `nvidia` 驱动和 `cuda` 包，而且开箱即用，其他发行版请自行如法炮制。
12+
13+
Wendous 用户可能在安装完后遇到“找不到 cuxxx.dll”报错，说明你需要拷贝 CUDA 安装目录下的所有 DLL 到 `C:\\Windows\\System32`。
14+
15+
WSL 用户要注意，WSL 环境和真正的 Linux 相差甚远。很多 Linux 下的教程，你会发现在 WSL 里复刻不出来。这是 WSL 的 bug，应该汇报去让微软统一修复，而不是让教程的作者零零散散一个个代它擦屁股。建议直接在 Wendous 本地安装 CUDA 反而比伺候 WSL 随机拉的 bug 省力。
16+
17+
Ubuntu 用户可能考虑卸载 Ubuntu，因为 Ubuntu 源中的版本永不更新。想要安装新出的软件都非常困难，基本只能安装到五六年前的古董软件，要么只能从网上下 deb 包，和 Wendous 一个软耸样。所有官方 apt 源中包的版本从 Ubuntu 发布那一天就定死了，永远不会更新了。这是为了起夜级服务器安全稳定的需要，对于个人电脑而言却只是白白阻碍我们学习，Arch Linux 这样的滚动更新的发行版才更适合个人桌面用户。
18+
19+
### 安装 NVIDIA 驱动
420

521
首先确保你安装了 NVIDIA 最新驱动：
622

@@ -14,21 +30,74 @@ pacman -S nvidia
1430
nvidia-smi
1531
```
1632

33+
应该能得到：
34+
35+
```
36+
Mon Aug 26 14:09:15 2024
37+
+-----------------------------------------------------------------------------------------+
38+
| NVIDIA-SMI 555.58.02 Driver Version: 555.58.02 CUDA Version: 12.5 |
39+
|-----------------------------------------+------------------------+----------------------+
40+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
41+
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
42+
| | | MIG M. |
43+
|=========================================+========================+======================|
44+
| 0 NVIDIA GeForce RTX 4070 ... Off | 00000000:01:00.0 On | N/A |
45+
| 0% 30C P8 17W / 285W | 576MiB / 16376MiB | 41% Default |
46+
| | | N/A |
47+
+-----------------------------------------+------------------------+----------------------+
48+
49+
+-----------------------------------------------------------------------------------------+
50+
| Processes: |
51+
| GPU GI CI PID Type Process name GPU Memory |
52+
| ID ID Usage |
53+
|=========================================================================================|
54+
| 0 N/A N/A 583 G /usr/lib/Xorg 370MiB |
55+
| 0 N/A N/A 740 G xfwm4 4MiB |
56+
| 0 N/A N/A 783 G /usr/lib/firefox/firefox 133MiB |
57+
| 0 N/A N/A 4435 G obs 37MiB |
58+
+-----------------------------------------------------------------------------------------+
59+
```
60+
1761
如果不行，那就重启。
1862

19-
## 安装 CUDA
63+
### 安装 CUDA
2064

2165
然后安装 CUDA Toolkit（即 nvcc 编译器）：
2266

2367
```bash
2468
pacman -S cuda
69+
```
70+
2571
打开 `.bashrc`（如果你是 zsh 用户就打开 `.zshrc`），在末尾添加两行：
2672

2773
```bash
2874
export PATH="/opt/cuda/bin:$PATH" # 这是默认的 cuda 安装位置
2975
export NVCC_CCBIN="/usr/bin/g++-13" # Arch Linux 用户才需要这一行
3076
```
31-
## 常见问题
77+
78+
然后重启 `bash`，或者执行以下命令重载环境变量：
79+
80+
```bash
81+
source .bashrc
82+
```
83+
84+
运行以下命令测试 CUDA 编译器是否可用：
85+
86+
```bash
87+
nvcc --version
88+
```
89+
90+
应该能得到：
91+
92+
```
93+
nvcc: NVIDIA (R) Cuda compiler driver
94+
Copyright (c) 2005-2024 NVIDIA Corporation
95+
Built on Thu_Jun__6_02:18:23_PDT_2024
96+
Cuda compilation tools, release 12.5, V12.5.82
97+
Build cuda_12.5.r12.5/compiler.34385749_0
98+
```
99+
100+
### 常见问题解答
32101

33102
CMake 报错找不到 CUDA？添加环境变量：
34103

@@ -37,53 +106,244 @@ export PATH="/opt/cuda/bin:$PATH" # 这里换成你的 cuda 安装位置
37106
export NVCC_CCBIN="/usr/bin/g++-13" # 只有 Arch Linux 需要这一行
38107
```
39108

40-
Clangd 报错不认识 `-forward-unknown-to-host-compiler` 选项？
109+
IDE 使用了 Clangd 静态检查插件，报错不认识 `-forward-unknown-to-host-compiler` 选项？
41110

42111
创建文件 `~/.config/clangd/config.yaml`：
43112

44113
```yaml
45114
CompileFlags:
46-
Add:
47-
- --cuda-gpu-arch=sm_86
48-
Remove:
115+
Add: # 要额外添加到 Clang 的 NVCC 没有的参数
116+
- --no-cuda-version-check
117+
Remove: # 移除 Clang 不认识的 NVCC 参数
49118
- -forward-unknown-to-host-compiler
50-
- --expt-relaxed-constexpr
51-
- --expt-extended-lambda
119+
- --expt-*
52120
- --generate-code=*
53121
- -arch=*
122+
- -rdc=*
123+
```
124+
125+
### 建议开启的 CMake 选项
126+
127+
#### CUDA 编译器路径
128+
129+
如果你无法搞定环境变量，也可以通过 `CMAKE_CUDA_COMPILER` 直接设置 `nvcc` 编译器的路径：
130+
131+
```cmake
132+
set(CMAKE_CUDA_COMPILER "/opt/cuda/bin/nvcc") # 这里换成你的 cuda 安装位置
54133
```
55134

56-
### 建议开启的选项
135+
不建议这样写，因为会让使用你项目的人也被迫把 CUDA 安装到这个路径去。
57136

58-
```bash
59-
nvcc --expt-relaxed-constexpr --expt-extended-lambda
137+
建议是把你的 `nvcc` 安装好后，通过 `PATH` 环境变量，`cmake` 就能找到了，不需要设置这个变量。
138+
139+
#### CUDA C++ 版本
140+
141+
CUDA 是一种基于 C++ 的领域特定语言，CUDA C++ 的版本和正规 C++ 一一对应。
142+
143+
目前最新的是 CUDA C++20，可以完全使用 C++20 特性的同时书写 CUDA 代码。
144+
145+
- 在 `__host__` 函数（未经特殊修饰的函数默认就是此类，在 CPU 端执行）中，CUDA 和普通 C++ 没有区别，任何普通 C++ 代码，都可以用 CUDA 编译器编译。
146+
- 在 `__device__` 函数（CUDA kernel，在 GPU 端执行）中，能使用的函数和类就有一定限制了：
147+
- 例如你不能在 `__device__` 函数里使用仅限 `__host__` 用的 `std::cout`（但 `printf` 可以，因为 CUDA 团队为了方便用户调试，为你做了 `printf` 的 `__device__` 版特化）。
148+
- `__device__` 中不能使用绝大多数非 `constexpr` 的 STL 容器，例如 `std::map` 等，但是在 `__host__` 侧还是可以用的！
149+
- 所有的 `constexpr` 函数也是可以使用的，例如各种 C++ 风格的数学函数如 `std::max`，`std::sin`，这些函数都是 `constexpr` 的，在 `__host__` 和 `__device__` 都能用。
150+
- 如果一个容器的成员全是 `constexpr` 的，那么他可以在 `__device__` 函数中使用。例如 `std::tuple`、`std::array` 等等，因为不涉及 I/O 和内存分配，都是可以在 `__device__` 中使用的。
151+
- 例如 C++20 增加了 constexpr-new 的支持，让 `std::vector` 和 `std::string` 变成了 `constexpr` 的容器，因此可以在 `__device__` 中使用 `std::vector`（会用到 `__device__` 版本的 `malloc` 函数，这是 CUDA 的一大特色：你可以在 kernel 内部用 `malloc` 动态分配设备内存，并且从 CUDA C++20 开始 `new` 也可以了）。
152+
- `std::variant` 现在也是 `constexpr` 的容器，也可以在 `__device__` 函数中使用了。
153+
- 异常目前还不是 `constexpr` 的，因此无法在 `__device__` 函数中使用 `try/catch/throw` 系列关键字。
154+
- 总之，随着，我们可以期待越来越多纯计算的函数和容器能在 CUDA kernel（`__device__` 环境）中使用。
155+
156+
正如 `CMAKE_CXX_STANDARD` 设置了 `.cpp` 文件所用的 C++ 版本，也可以用 `CMAKE_CUDA_STANDARD` 设置 `.cu` 文件所用的 CUDA C++ 版本。
157+
158+
```cmake
159+
set(CMAKE_CXX_STANDARD 20) # .cpp 文件采用的 C++ 版本是 C++20
160+
set(CMAKE_CUDA_STANDARD 20) # .cu 文件采用的 CUDA C++ 版本是 C++20
60161
```
61162

62-
* --expt-relaxed-constexpr: 让所有 constexpr 函数都自动带有 __host__ __device__
63-
* --expt-extended-lambda: 允许 lambda 表达式具有 __host__ 和/或 __device__ 修饰
163+
### 赋能现代 C++ 语法糖
64164

65-
### 建议开启的 CMake 选项
165+
```cmake
166+
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr --expt-extended-lambda")
167+
```
168+
169+
* `--expt-relaxed-constexpr`: 让所有 `constexpr` 函数默认自动带有 `__host__ __device__`
170+
* `--expt-extended-lambda`: 允许为 lambda 表达式指定 `__host__` 或 `__device__`
171+
172+
#### 显卡架构版本号
173+
174+
不同的显卡有不同的“架构版本号”，架构版本号必须与你的硬件匹配才能最佳状态运行，可以略低，但将不能发挥完整性能。
66175

67176
```cmake
68-
set(CMAKE_CUDA_SEPARABLE_COMPILATION ON)
69-
set(CMAKE_CUDA_ARCHITECTURES 86) # 不同的显卡有不同的“架构版本号”，架构版本号必须与你的硬件匹配才能最佳状态运行，可以略低，但将不能发挥完整性能。
70-
set(CMAKE_CUDA_ARCHITECTURES native) # 如果你有 CMake 3.24 以上，可以设定本参数，让 CMake 自动检测当前显卡，并选择准确的架构版本号。
177+
set(CMAKE_CUDA_ARCHITECTURES 86) # 表示针对 RTX 30xx 系列（Ampere 架构）生成
178+
set(CMAKE_CUDA_ARCHITECTURES native) # 如果 CMake 版本高于 3.24，该变量可以设为 "native"，让 CMake 自动检测当前显卡的架构版本号
71179
```
72180

73181
架构版本号：例如 75 表示 RTX 20xx 系列（Turing 架构）；86 表示 RTX 30xx 系列（Ampere 架构）；89 表示 RTX 40xx 系列（Ada 架构）等。
74182

75183
完整的架构版本号列表可以在 [CUDA 文档](https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html#virtual-architecture-feature-list) 中找到。
76184

77-
我的选项是：
185+
也可以运行如下命令（如果有的话）查询当前显卡的架构版本号：
186+
187+
```bash
188+
__nvcc_device_query
189+
```
190+
191+
#### 设备函数分离定义
192+
193+
默认只有 `__host__` 函数可分离声明和定义。如果你需要分离 `__device__` 函数的声明和定义，就要开启这个选项：
194+
195+
```cmake
196+
set(CMAKE_CUDA_SEPARABLE_COMPILATION ON) # 可选
197+
```
198+
199+
#### 创建 CUDA 项目
200+
201+
完成以上选项的设定后，使用 `project` 命令正式创建 CUDA C++ 项目。
202+
203+
```cmake
204+
project(这里填你的项目名 LANGUAGES CXX CUDA)
205+
```
206+
207+
> {{ icon.fun }} 我见过有人照抄代码把“这里填你的项目名”抄进去的。
208+
209+
如需在特定条件下才开启 CUDA，可以用 `enable_language()` 命令延迟 CUDA 环境在 CMake 中的初始化：
210+
211+
```cmake
212+
project(这里填你的项目名 LANGUAGES CXX)
213+
214+
...
215+
216+
option(ENABLE_CUDA "Enable CUDA" ON)
217+
218+
if (ENABLE_CUDA)
219+
enable_language(CUDA)
220+
endif()
221+
```
222+
223+
#### CMake 配置总结
224+
225+
注意！以上这些选项设定都必须在 `project()` 命令之前！否则设定了也无效。
226+
227+
因为实际上是 `project()` 命令会检测这些选项，用这些选项来找到编译器和 CUDA 版本等信息。
228+
229+
总之，我的选项是：
78230

79231
```cmake
80232
cmake_minimum_required(VERSION 3.12)
81233
82234
set(CMAKE_CXX_STANDARD 20)
83235
set(CMAKE_CUDA_STANDARD 20)
84-
set(CMAKE_CUDA_SEPARABLE_COMPILATION ON)
236+
set(CMAKE_CUDA_SEPARABLE_COMPILATION OFF)
85237
set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} --expt-relaxed-constexpr --expt-extended-lambda")
86238
if (NOT DEFINED CMAKE_CUDA_ARCHITECTURES AND CMAKE_VERSION VERSION_GREATER_EQUAL 3.24)
87239
set(CMAKE_CUDA_ARCHITECTURES native)
88240
endif()
241+
242+
project(你的项目名 LANGUAGES CXX CUDA)
243+
244+
file(GLOB sources "*.cpp" "*.cu")
245+
add_executable(${PROJECT_NAME} ${sources})
246+
target_link_libraries(${PROJECT_NAME} PRIVATE cusparse cublas)
247+
```
248+
249+
## 开始编写 CUDA
250+
251+
CUDA 有两套 API：
252+
253+
- [CUDA runtime API](https://docs.nvidia.com/cuda/cuda-runtime-api/index.html)：更加简单，兼顾性能，无需手动编译 kernel，都替你包办好了，但不够灵活。
254+
- [CUDA driver API](https://docs.nvidia.com/cuda/cuda-driver-api/index.html)：更加灵活多变，但操作繁琐，需要手动编译 kernel，适合有特殊需求的用户。
255+
256+
他们都提供了大量用于管理 CUDA 资源和内存的函数。
257+
258+
我们要学习的是比较易懂、用的也最多的 CUDA runtime API。
259+
260+
使用 `<cuda_runtime.h>` 头文件即可导入所有 CUDA runtime API 的函数和类型：
261+
262+
```cuda
263+
#include <cuda_runtime.h>
264+
```
265+
266+
虽然 CUDA 基于 C++（而不是 C 语言），支持所有 C++ 语言特性。但其 CUDA runtime API 依然是仿 C 风格的接口，可能是照顾了部分从 C 语言转过来的土木老哥，也可能是为了方便被第三方二次封装。
267+
268+
我们的课程主题是：用现代 C++ 赋能更好的 CUDA 开发，所以会对 CUDA 原生的 C 风格 API 做一些 C++ 封装，使其呈更直观易用的接口，帮助你避免出错（例如内存泄漏）。
269+
270+
### 认识 CUDA 语言
271+
272+
一份 CUDA 源码和 C/C++ 一样，是由大量的函数组成。
273+
274+
由于 GPU 编程的特殊性，GPU 代码和 CPU 代码是需要分离的，他们的指令集完全不同。
275+
276+
CPU 的责任是决定什么时候要“启动(launch)” GPU 代码，而 GPU 只专注于计算。
277+
278+
所以 CPU 的指令集中含有大量的条件和判断，而 GPU 则以计算指令为主。
279+
280+
而 C++ 中，函数就是代码，根据代码执行的位置不同，函数可以分为：
281+
282+
- GPU 端执行。
283+
- CPU 端执行。
284+
285+
为此，就需要对进行标识：
286+
287+
CUDA 中的函数分为三大类：
288+
289+
- `__host__` 函数：在 CPU 端执行，只能被 CPU 端函数调用。编译器编译时，会将其编译为 CPU 的汇编（x86 汇编），可以使用所有 C++ 标准库功能（如 `std::cout`）。
290+
- `__device__` 函数：在 GPU 端执行，只能被 GPU 端函数调用。编译器编译时，会将其编译为 GPU 的汇编（PTX 汇编），只能使用 C++ 标准库中纯计算的部分功能（如 `std::sin`）。
291+
- `__global__` 函数：也是在 GPU 端执行，类似于 `__device__`。区别在于：
292+
- `__global__` 函数可以被 CPU 端的函数（`__host__`）调用。
293+
- `__device__` 只能被 GPU 端的函数（`__device__` 或 `__global__`）调用。
294+
295+
因而 `__global__` 就像一座桥梁一样，是从 CPU 走向 GPU 的入口点。
296+
297+
> {{ icon.tip }} 但反过来，GPU 函数不能再走回 CPU。所以 `__global__` 是单向的一次性桥梁，一旦进入 GPU，就只能等整个 `__global__` 函数退出了。
298+
299+
`__global__` 就像 GPU 版的 `main` 函数一样，是所有 GPU 代码的入口点。
300+
301+
只不过这个入口点可以有很多个，CPU 可以多次提交不同的 `__global__` 函数，就像是创建了许多个不同的“进程”，分别有各自的 `main` 入口点一样。
302+
303+
这样的一次 `__global__` 调用所产生的一个“GPU 版进程”，被称作一个“网格(grid)”。
304+
305+
正如 CPU 上的单个进程由很多子线程组成一样；GPU 上的每个“网格”，由许多“块(block)”组成，“块”又进一步由许多“线程(thread)”组成。
306+
307+
> {{ icon.detail }} 实际上，GPU 的“块”才最接近 CPU 上线程概念，而 GPU 所谓的“线程”实际上对应于 CPU 的 SIMD 矢量，稍后会详细介绍。
308+
309+
因为很多时候，人们喜欢直接在 `__global__` 函数里写上所有的计算代码，而不会再定义一个 `__device__` 函数了，所有的计算核心功能都在 `__global__` 函数中。
310+
311+
“网格(grid)”是“内核(kernel)”的实例，正如 CPU 上进程是可执行文件(exe)的实例一样：同一个可执行文件可以被 Shell 多次调用产生多个进程，同一个“内核”也可以被 host 函数多次调用产生多个“网格”。
312+
313+
但是，CPU 上的每个进程只会调用一次 `main` 入口点，而 GPU 上的 `__global__` 会被调用 n 次（n 的大小在 host 函数中指定），所有启动的 n 个 `__global__` 函数互相之间是并行执行的，每个线程的入口点都是 `__global__`，因此一个“网格”含有多个“线程”。
314+
315+
```cuda
316+
#include <cuda_runtime.h>
317+
318+
/*__host__*/ void host_func() {
319+
}
320+
321+
__device__ void device_func() {
322+
}
323+
324+
__host__ __device__ void host_device_func() {
325+
}
326+
327+
constexpr void constexpr_func() {
328+
}
329+
330+
__global__ void kernel() {
331+
device_func();
332+
host_device_func();
333+
constexpr_func(); // 需开启 --expt-relaxed-constexpr
334+
auto device_lambda = [] __device__ (int i) { // 需开启 --expt-extended-lambda
335+
return i * 2;
336+
};
337+
device_lambda(1);
338+
}
339+
340+
int main() {
341+
host_func();
342+
host_device_func();
343+
constexpr_func();
344+
auto host_lambda = [] (int i) {
345+
return i * 2;
346+
};
347+
host_lambda(1);
348+
}
89349
```

0 commit comments

Comments
 (0)