C++高性能并行编程与优化 - 课件 - 08 CUDA 开启的 GPU 编程

• 语言: 中文（简体）
• 格式: .pptx
• 来源: github.com/archibate

摘要

文档主要介绍了C++在高性能并行编程中的应用，特别是使用CUDA进行GPU编程。内容包括GPU与CPU的性能对比、核函数的定义与调用、动态并行的实现、内存优化策略以及内联函数的使用。文档通过示例代码展示了如何在GPU上进行高效的并行计算，并强调了避免线程组分歧、跨步访问和银行冲突等优化技巧。此外，还提到了CUDA的多段编译特点和如何利用宏指令针对CPU和GPU生成不同的代码。

AI总结

### 文档总结

#### 1. **GPU 编程基础**
- **数据类型选择**：`sin` 是 `double` 类型，而 `sinf` 是 `float` 类型，使用时注意区分，避免性能损失。
- **性能测试**：
- 使用 `ticktock.h` 测试 CPU 和 GPU 的运行时间。
- 测试结果：GPU 性能远优于 CPU，`cpu_sinf: 0.213827s` vs `gpu_sinf: 0.013417s`。
- 注意：`cudaDeviceSynchronize()` 必须放在同步之后，以确保时间测量准确。

#### 2. **代码生成与条件编译**
- CUDA 具有多段编译特性，代码会被预处理多次，生成 CPU 和 GPU 的指令码。
- 使用 `#ifdef` 判断 `___CUDA_ARCH___` 宏，实现针对 GPU 和 CPU 的代码生成。
- 示例：通过条件编译输出不同的日志信息，如 GPU 和 CPU 的问候语。

#### 3. **动态并行与核函数调用**
- 从 Kepler 架构开始，GPU 核函数可以调用其他核函数，实现动态并行。
- 优势：避免数据返回 CPU 导致的同步开销，适合动态计算线程数和网格数。
- 示例：`kernel` 核函数动态调用 `another` 核函数，指定不同的线程数和网格数。

#### 4. **GPU 优化手法**
- **线程组分歧（Divergence）**：尽量保证 32 个线程走同一条分支，否则会导致性能损失。
- **延迟隐藏（Latency Hiding）**：通过足够大的 `blockDim` 调度其他线程组，避免内存等待。
- **寄存器打翻（Register Spill）**：局部变量过多会导致寄存器不足，建议减少 `blockDim`。
- **共享内存（Shared Memory）**：减少全局内存访问开销，适合多次访问的数据。
- **跨步访问（Coalesced Access）**：顺序访问共享内存，提升带宽利用率。
- **银行冲突（Bank Conflict）**：避免同一 warp 中多个线程访问共享内存的同一银行，可通过错开步长（如 33 跨步）解决。

#### 5. **内联函数**
- CUDA 提供 `__inline__` 和 `__forceinline__` 来实现函数内联，优化函数调用开销。
- 注意：`inline` 在现代 C++ 中仅是弱符号声明，与性能无关。

#### 6. **课程内容与工具**
- **下一课主题**：
- CUDA 流（cudaStream）异步编程。
- Nsight profiler 性能分析工具。
- 常用库（如 curand、cufft、cublas）。
- GPU 和 CPU 实现的红黑高斯方法。
- **前置条件**：
- 熟悉 C/C++ 编程。
- 英伟达 GTX900 及以上显卡。
- CUDA 11 及以上版本。
- CMake 3.18 及以上版本。

#### 7. **总结**
- CUDA 提供了强大的 GPU 编程能力，通过条件编译、动态并行和优化手法，可以显著提升程序性能。
- 注意数据类型、同步机制和内存访问模式，避免性能损失。

---

以上为文档的核心内容总结，语言简洁，重点突出关键信息，便于理解和查阅。