How You Should Write a CUDA C++ Kernel

George Sifniotakis, Software Engineering Manager | March 18

目录

1. 从试错到原则化的CUDA C++内核工程
2. 问题陈述
3. 性能评测与分析
   • 3.1 内核的异步性与性能评测陷阱
   • 3.2 NVBench：一个专业的CUDA性能评测框架
   • 3.3 性能分析：光速（Speed of Light）分析法
4. CUDA Kernel 实现的演进
   • 4.1 朴素单线程内核 (Naive Single-Threaded Kernel
   • 4.2 朴素单块内核 (Naive Single-Block Kernel
   • 4.3 标准网格-块内核 (Naive Kernel
5. CUDA Kernel 编程基础
   • 5.1 如何选择块大小 (Block Size
   • 5.2 网格大小（Grid Size）的选择
   • 5.3 调试：越界访问与Compute Sanitizer
6. 拥抱现代CUDA C++
   • 6.1 使用 std::span 避免原始指针错误
   • 6.2 libcu++ 中提供的词汇类型
7. 深入理解和优化核函数性能
   • 7.1 核函数执行流程与延迟隐藏
   • 7.2 性能瓶颈与利特尔定律 (Little's Law
   • 7.3 使用 Nsight Compute 进行性能分析
   • 7.4 指令级并行 (Instruction-Level Parallelism, ILP
   • 7.5 局部内存（Local Memory）与性能陷阱
   • 7.6 编译器优化指令
   • 7.7 理解并优化占用率 (Occupancy
8. 高级优化技术
   • 8.1 向量化内存访问
   • 8.2 高效的内存移动与 PTX
9. 结论与核心要点
   • 9.1 最好的内核是你不必写的那个
   • 9.2 协同库 (Cooperative Libraries
   • 9.3 核心要点 (Takeaways

---

1. 从试错到原则化的CUDA C++内核工程

本次演讲旨在帮助开发者从一种低效的开发模式转变为一种更系统化、现代化的CUDA C++内核工程方法。

学习目标：

• 从以下状态转变：
• 低下的开发者生产力
• 凭猜测进行优化

• 不安全的代码

• 迈向以下目标：

• 系统化的开发方法
• 现代化的优化技术与最佳实践
• 使用设备端的CUDA C++库
• 强大的开发者工具

2. 问题陈述

我们以一个具体问题为例：模拟N个物体冷却至室温的过程。

• 假设室温为20°C。
• 每个物体有一个初始温度，在每个时间步后，其温度会根据某个物理模型发生变化。
• 这个过程可以抽象为一个转换操作：根据prev_temperatures（前一时刻的温度数组）计算出next_temperatures（当前时刻的温度数组）。

在CUDA生态中，可以使用thrust::transform库函数来高效实现这类操作：
thrust::transform(prev, prev + size, next, compute);

3. 性能评测与分析

3.1 内核的异步性与性能评测陷阱

在对CUDA内核进行性能测试时，我们发现了一个常见问题：CUDA内核调用是异步的。CPU将核函数调用放入CUDA流（stream）后会立即返回，而不会等待GPU执行完毕。

• 如果直接测量内核启动函数（kernel<<<...>>>）的耗时，得到的时间会非常短，因为它只包含了将任务提交到GPU的开销，而非实际执行时间。
• 这导致了与Thrust库（它会隐式地进行同步）在性能数据上的巨大差异，产生了误导性的结果。

3.2 NVBench：一个专业的CUDA性能评测框架

为了进行准确、可靠的性能评测，我们推荐使用NVBench。这是一个专为CUDA设计的C++ benchmark框架。

主要特性：

1. 简洁强大的定义方式：采用类似Google Benchmark的风格定义测试用例。
2. 原生CUDA支持：内置对CUDA事件计时器、设备、缓存刷新等的支持。
3. 提升的稳定性：通过流阻塞（stream blocking）和合理的停止标准，确保测量结果的稳定与准确。
4. 强大的参数扫描框架：可以方便地探索不同参数值（如数据大小）和静态类型对性能的影响。
5. 灵活的命令行接口：可以轻松选择要运行的benchmark、更换测试设备等。

项目地址：https://github.com/NVIDIA/nvbench

NVBench 使用示例

1. 定义Benchmark：
   一个benchmark本质上是一个函数，它接受一个nvbench::state &state参数。你需要通过NVBENCH_BENCH宏来注册这个函数。

   

2. 添加参数轴：
   使用add_..._axis系列函数（如add_int64_power_of_two_axis）来定义一个参数维度。在benchmark函数内部，通过state.get_int64("...")来获取当前迭代的参数值。

   

3. 执行与计时：
   通过state.exec([&](nvbench::launch &launch){ ... });来包裹你的CUDA内核启动代码。

   • state.exec会自动多次调用内核启动器，直到采集到具有代表性的样本。
   • 所有的设置代码（setup）都应在state.exec之前完成，这部分代码只会执行一次。

   

3.3 性能分析：光速（Speed of Light）分析法

重新评估性能

使用NVBench进行准确的性能测试后，我们发现朴素内核的性能极差，比Thrust库慢了大约47,000倍。

这引出了两个关键问题：
1. 我们怎么知道Thrust是一个好的性能参考标准？
2. 当没有任何参考实现时，我们如何判断自己的内核性能是否足够好？

光速（Speed of Light, SOL）性能分析

与其和某个参考实现进行性能比较（speedup），不如将实际性能与硬件的理论极限进行比较。这种方法被称为“光速（SOL）分析”。

• Speedup的局限：加速比无法告诉我们一个算法的性能潜力还有多大。
• SOL的优势：SOL是理论上的性能峰值。将实际性能与SOL进行比较，可以提供清晰的优化目标。

示例：
- 将一个1GB的数据块用一个线程块（thread block）进行拷贝，比用单个线程快10倍。
- 但这个性能实际上只达到了理论峰值的1/3。这说明虽然有10倍的提升，但仍有巨大的优化空间。

算术强度与Roofline模型

为了理解性能瓶颈，我们需要引入算术强度（Arithmetic Intensity）的概念。

• 定义：算术强度 = 浮点运算次数 (FLOPS) / 内存访问量 (Bytes)

Roofline模型提供了一个直观的性能分析框架。它将应用的性能上限可视化，该上限由两个因素决定：峰值内存带宽（Peak bandwidth）和峰值计算性能（Peak compute）。

• 内存密集型 (Memory bound)：应用的算术强度较低，其性能受限于内存带宽。
• 计算密集型 (Compute bound)：应用的算术强度较高，其性能受限于计算单元的处理能力。

对于我们的transform问题，其算术强度很低，因此它是一个内存密集型的应用。其性能瓶颈在于内存带宽，而非计算能力。

利用NVBench进行吞吐量分析

NVBench允许我们在benchmark中声明内存读写量，从而自动计算吞吐量和与理论峰值的差距（SOL %）。

• 使用state.add_global_memory_reads(bytes)和state.add_global_memory_writes(bytes)来声明全局内存的读写字节数。

吞吐量分析结果：

通过吞吐量分析，我们得到了更深层次的洞见：

• Thrust: 实现了高达84%的理论峰值带宽（SOL），这表明其实现已经非常接近硬件极限。
• 朴素内核: 仅实现了0%的SOL，带宽利用率极低。

结论：朴素的单线程内核之所以性能差，根本原因在于它完全没有有效利用GPU的并行内存访问能力，导致内存带宽被严重浪费。这为我们指明了下一步的优化方向：并行化。

4. CUDA Kernel 实现的演进

4.1 朴素单线程内核 (Naive Single-Threaded Kernel)

为了解决上述问题，一个最简单的想法是编写一个单线程的CUDA内核。

• 实现方式：在GPU上启动一个线程，该线程通过一个循环，顺序处理数组中的每一个元素。

代码示例：

__global__
void naive(const half *in, half *out, int size)
{
  for (int i = 0; i < size; ++i)
  {
    out[i] = compute(in[i]);
  }
}

void transform(const half *in, half *out, int size)
{
  kernel<<<1, 1, 0, stream>>>(in, out, size);
}

• 执行流程与效率问题:

  1. 一个warp调度器（warp scheduler）发出内存访问请求。
  2. 在等待数据从内存返回的过程中，该warp调度器进入停顿（stalled）状态。由于只有一个活跃的warp，整个流式多处理器（SM）都处于空闲状态，无法执行其他计算任务。
  3. 数据从内存返回后，warp调度器才能恢复执行。

• 问题: 这种方法的瓶颈在于内存延迟。计算单元在等待内存操作完成时完全闲置，极大地浪费了GPU的计算资源。

4.2 朴素单块内核 (Naive Single-Block Kernel)

为了解决单线程内核的效率问题，可以采用单线程块（single thread block）的方法。通过在一个块内启动多个线程，可以在一个SM上同时运行多个warp。

• 实现方式:
  • 在GPU上启动一个线程块。
  • 将数据元素在块内的多个线程之间进行划分。
  • 采用“网格跨步循环”（grid-stride loop）的方式，使得固定数量的线程可以处理任意数量的数据元素。

示例代码:

  
    // 定义线程数量
    const int number_of_threads = 256;
    // CUDA 内核函数
    __global__ void kernel1_half(half in, half out, int size) {
        // 获取线程在块内的索引
        int thread_index = threadIdx.x;
        // 使用跨步循环，让每个线程处理多个元素
        for (int i = thread_index; i < size; i += number_of_threads) {
            out[i] = compute(in[i]);
        }
    }
    // 在主机端启动内核
    // 启动1个线程块，每个块包含 number_of_threads 个线程
    kernel1<<<1, number_of_threads>>>(in, out, size);

• 性能提升原理 (延迟隐藏 Latency Hiding):
  当一个warp因为等待内存操作而停顿时，SM的warp调度器可以切换到另一个准备就绪的warp上继续执行计算或发起新的内存请求。这样，通过重叠计算和内存访问，可以有效隐藏内存延迟，保持计算单元的繁忙状态，从而提高SM的利用率。

• 性能对比:
  与朴素单线程（Naive 1T）内核相比，朴素单块（Naive 1B）内核的内存利用率得到了显著提升（从0.003%提升到0.059%）。
• 结论: 内存利用率有所改善。
• 不足: 即便如此，利用率仍然低于H100 GPU潜能的1%，说明还有巨大的优化空间。需要进一步利用更多的GPU资源。

4.3 标准网格-块内核 (Naive Kernel)

为了充分利用整个GPU的计算能力（而不仅仅是单个SM），标准的CUDA编程模型是启动一个由多个线程块组成的网格（a grid of blocks）。

• 核心思想:
  • 启动一个线程块的网格（grid of blocks）。
  • 为每个数据元素分配一个线程。
  • 通过 blockIdx 和 threadIdx 的组合计算全局唯一的线程ID，以映射到数据元素。

• 实现方式:
• 全局索引计算: int i = blockIdx.x * blockDim.x + threadIdx.x;
• 边界检查: if (i < size) 确保线程不会访问越界内存。
• 启动配置:
  • block_size: 每个块中的线程数。
  • grid_size: 网格中的块数，通常通过 (size + block_size - 1) / block_size 来计算，以确保所有元素都被覆盖。
  • kernel2<<<grid_size, block_size>>>(...); 启动内核。

5. CUDA Kernel 编程基础

5.1 如何选择块大小 (Block Size)

选择合适的块大小（Block Size）对于性能至关重要，但没有一个“万能”的尺寸适用于所有内核。

• 经验法则 (Rule of thumb):
  • 使用32的倍数: 块大小应为warp大小（通常是32）的倍数，以确保硬件的高效利用。
  • 默认值: 256通常是一个不错的起点和默认值。
  • 性能剖析 (Profiling): 为了获得最佳性能，需要针对具体应用和硬件进行性能剖析和调优。

5.2 网格大小（Grid Size）的选择

在启动CUDA核函数时，我们需要确定线程块的数量，即网格大小（Grid Size）。网格大小通常取决于具体问题的数据规模。为了计算所需的网格大小，可以使用 cuda::std::div_ceil 函数，这是一个向上取整的除法操作，可以确保分配足够的线程块来处理所有数据。

其计算方式如下：
grid_size = cuda::std::div_ceil(size, block_size);

其中 size 是数据元素的总数，block_size 是每个线程块中的线程数。

示例代码：

#include <cuda/std/cmath>

void kernel(...) {
  ...
}

int main() {
  int block_size = 256;
  int grid_size = cuda::std::div_ceil(size, block_size);

  kernel<<<grid_size, block_size>>>(in, out);
}

5.3 调试：越界访问与Compute Sanitizer

越界访问（Out-of-Bounds Access）问题

当使用向上取整的方式计算网格大小时，启动的线程总数可能会超过实际需要处理的数据元素数量。如果在核函数内不进行边界检查，这可能导致线程访问数组范围之外的内存，即越界访问。

如下图所示，当数据大小为6，而每个块有4个线程时，会启动2个块，总共8个线程。第二个线程块（blockIdx.x=1）中的最后两个线程（红色标记）将会访问索引为6和7的元素，这超出了数组的有效范围（0-5），从而导致错误。

产生越界访问的核函数代码示例：

__global__
void kernel(_half* in, _half* out, int size) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;

  // 缺乏边界检查，当 i >= size 时会发生越界
  out[i] = compute(in[i]);
}

使用Compute Sanitizer进行调试

越界访问等内存错误可以通过 compute-sanitizer 工具来检测。这是一个用于CUDA应用程序的内存和线程错误检测工具。

使用方法与输出示例：
通过在命令行中运行 compute-sanitizer 并指定要分析的应用程序，可以捕获运行时错误。

nvcc --extended-lambda -g -O0 rtmp/x.cu -o /tmp/x.out
compute-sanitizer /tmp/x.out

如果存在越界访问，compute-sanitizer 会生成详细的报告，指出错误的类型、位置和相关信息。

错误报告摘要：
- 错误类型: Invalid __global__ read of size 2 bytes （无效的2字节全局读取）
- 错误位置: 在核函数 _Z6kernelP6__halfS0_i 的第4行，位于块 (4,0,0) 中的线程 (192,0,0)。
- 详细信息: 地址 0x7f3271010000 越界，该地址位于最近分配的内存块（大小为20,000,000字节）之后8字节处。

6. 拥抱现代CUDA C++

6.1 使用 std::span 避免原始指针错误

std::span 是C++20中引入的一个视图类型，libcu++ 库使其可以在CUDA核函数中使用。它为连续的数据序列提供了一个非拥有（non-owning）的视图，并包含了大小信息，是原始指针的一种更安全的替代方案。

std::span 的特性：
- 包含大小信息：自带边界检查，有助于防止越界访问。
- 原始指针的更安全替代品：减少常见的指针错误。
- 轻量级视图：它不拥有数据，只是引用一段连续的内存。
- 不拥有数据：仅引用数据，不负责其生命周期管理。

通过将核函数参数从原始指针 (_half*) 和大小 (int size) 修改为 std::span<_half>，可以使代码更安全、更简洁。

代码对比：
原始指针版本:

__global__ void kernel(_half* in, _half* out, int size);
kernel<<<grid_size, block_size>>>(in_ptr, out_ptr, size);

std::span 版本:

__global__ void kernel(std::span<_half> in, std::span<_half> out);

std::span<_half> in_span(in_ptr, size);
std::span<_half> out_span(out_ptr, size);
kernel<<<grid_size, block_size>>>(in_span, out_span);

重要限制: 在 __global__ 函数（设备代码）中，不能调用 __host__ 函数的成员函数（如 operator[]）。因此 std::span 的某些主机端功能在设备上受限。

6.2 libcu++ 中提供的词汇类型

libcu++ 是NVIDIA为CUDA C++提供的标准库实现，支持在GPU上使用多种C++标准库组件。

主要类型包括：
- 复合类型 (Compound Types):
- cuda::std::pair
- cuda::std::tuple
- 可选值与变体 (Optional and Alternatives):
- cuda::std::optional
- cuda::std::variant
- 数学 (Math):
- cuda::std::complex
- cuda::std::mdspan
- 同步 (Synchronization):
- cuda::std::atomic
- cuda::std::atomic_ref
- cuda::std::atomic_flag
- CUDA 扩展 (CUDA Extensions):
- cuda::atomic
- cuda::atomic_ref
- cuda::atomic_flag

更多信息可查阅官方文档: https://nvidia.github.io/cccl/libcudacxx/

7. 深入理解和优化核函数性能

7.1 核函数执行流程与延迟隐藏

以下图示逐步展示了一个朴素核函数的执行过程，揭示了GPU如何通过线程束（Warp）调度来隐藏内存访问延迟。

1. 初始状态：两个流式多处理器（SM）上各有一个Warp准备就绪。
   

2. 内存请求：Warp发起全局内存读取请求。这些请求被发送到内存控制器，Warp进入等待状态。
   

3. Warp切换：当第一批Warp因等待内存而暂停时，SM会切换到其他就绪的Warp继续执行，以利用计算资源。
   

4. 数据返回：第一个内存请求完成，数据从内存返回。
   

5. 计算：收到数据的Warp恢复执行，进行计算操作。
   

这个过程展示了GPU的核心特性之一：延迟隐藏（Latency Hiding）。通过在等待高延迟操作（如内存访问）时切换执行其他Warp，GPU能够保持其计算单元的繁忙，从而提高整体吞吐量。

7.2 性能瓶颈与利特尔定律 (Little's Law)

对朴素核函数的性能进行分析，可以看到内存利用率仍有提升空间。

• 性能结果: 柱状图显示，某个版本的核函数（Naive 64）内存利用率（84.47%）远高于另一个版本（Naive 16，19.78%）。
• 瓶颈: 即使是表现更好的版本，内存利用率仍然低于H100 GPU理论峰值的30%。
• 问题: 如何才能更充分地利用内存带宽？

答案在于理解并应用利特尔定律（Little's Law）。

利特尔定律: concurrency = latency * throughput
- 并发度 (Concurrency): 系统中正在处理的请求数量（例如，在途的内存请求）。
- 延迟 (Latency): 处理单个请求所需的时间（例如，内存访问延迟）。
- 吞吐量 (Throughput): 单位时间内系统处理的请求数量（例如，内存带宽）。

为了达到系统的最大吞吐量（例如，饱和内存带宽），必须维持足够高的并发度，以完全覆盖内存访问的延迟。换言之，GPU需要有足够多的在途内存请求，使得在等待一个请求返回的过程中，内存总线始终在为其他请求服务。如果并发度不足，计算单元就会因等待数据而空闲，导致性能瓶颈。

下图展示了两个流式多处理器（Streaming Multiprocessors, SMs）与内存系统交互的简化模型。为了充分利用内存带宽，需要有足够多的内存请求在“飞行”中（in-flight），即已发出但尚未完成。这里的并发度（concurrency）计算为 4 * 4 = 16，表示系统需要维持16个并发的内存操作才能有效利用硬件资源。

7.3 使用 Nsight Compute 进行性能分析

Nsight Compute 是NVIDIA提供的一款强大的性能分析工具，专用于CUDA核函数（kernels）的深度分析和优化。

其主要功能包括：
- 详细的性能分析器：为您的CUDA核函数提供深入的性能数据。
- 源代码级关联：将性能指标与源代码或汇编代码逐行关联，实现精确的核函数调优。
- 引导式分析：高亮显示性能瓶颈，并提供针对性的优化建议。

通过Nsight Compute等工具，可以将底层指令的执行情况与源代码相关联，从而定位性能瓶颈。下图展示了三条汇编指令的Warp Stall（线程束暂停）采样情况。

• ld（加载）指令的暂停比例为0.91%。
• fma（乘加）指令的暂停比例高达74.74%。
• st（存储）指令的暂停比例为0.05%。

fma指令的高暂停率表明它在等待其依赖的数据（由ld指令加载）。这种由数据依赖导致的等待是性能瓶颈的主要原因。

• 指令本身不会阻塞执行：GPU硬件设计允许在一条指令等待时执行其他指令。
• 真正导致等待的是依赖性：后续指令必须等待其所需数据准备就绪。
• 利用这一特性发出多条指令：我们可以在一条长延迟指令（如内存加载）执行期间，发出其他独立的指令，从而隐藏延迟。
• 这种技术被称为指令级并行（Instruction-Level Parallelism, ILP）。

7.4 指令级并行 (Instruction-Level Parallelism, ILP)

指令级并行是一种通过重排代码指令来提高处理器资源利用率的技术。在GPU中，它常用于隐藏内存访问延迟。

以下代码示例展示了如何通过修改循环结构来实现ILP。原始代码通常将数据加载和数据计算分为两个独立的循环。

第一步：数据加载
每个线程在一个循环中加载8个元素到线程本地的数组 thread_data 中。

// ...
__half thread_data[items_per_thread];

for (int i = 0; i < items_per_thread; ++i) {
    int current_idx = idx + i * grid_size;
    if (current_idx < in.size()) {
        thread_data[i] = in[current_idx];
    }
}
// ...

第二步：数据计算和存储
在另一个独立的循环中，每个线程对加载到 thread_data 中的8个元素进行计算，并将结果写回。

// ...
for (int i = 0; i < items_per_thread; ++i) {
    int current_idx = idx + i * grid_size;
    if (current_idx < in.size()) {
        out[current_idx] = compute(thread_data[i]);
    }
}

这种分离的循环结构会导致compute操作必须等待所有load操作完成后才能开始，从而产生内存访问延迟导致的停顿。通过将加载和计算操作交错执行，可以实现指令级并行，隐藏内存延迟。

ILP 执行流程可视化

以下动画序列展示了ILP如何提高硬件利用率。

1. 初始状态：Warp（线程束）准备发出内存加载指令。此时没有在途（in flight）的内存请求。
   

2. 发出多个内存请求：Warp连续发出多个独立的内存加载指令。这些请求进入内存系统的队列，增加了在途请求的数量。
   

3. 数据返回，计算开始：当第一个内存请求的数据返回后，Warp可以立即对该数据执行计算指令。与此同时，其他的内存请求仍在处理中，内存总线保持繁忙。
   

4. 持续流水线：后续数据陆续返回，计算操作和新的内存请求可以重叠进行，形成一个高效的流水线，从而最大化地隐藏了内存延迟。
   

H200 硬件背景与ILP性能结果

现代GPU（如H200）拥有极高的内存带宽（例如4.8 TB/s）。然而，要充分利用这种带宽，必须有足够多的并发内存请求来填满硬件流水线。单个操作流的速度上限（图中示例为1.6 GB/s）远低于理论总带宽，这凸显了通过ILP和线程级并行（Thread-Level Parallelism）来饱和内存带宽的重要性。

下图比较了不同实现方式下的内存利用率。

• Naive T1: 每个线程处理1个元素，ILP机会很少，内存利用率几乎为0%。
• Naive T2: 每个线程处理2个元素，利用率提升至30.76%。
• Naive T8: 每个线程处理8个元素，提供了更多的ILP机会，内存利用率显著提升至76.27%。
• Thrust: NVIDIA官方的高度优化库，利用率达到84.42%。

结论:
- 通过增加每个线程处理的元素数量，可以有效利用指令级并行来提高内存利用率。
- 尽管手动优化取得了显著效果，但与Thrust这样的专业库相比仍有差距。
- “为什么还有差距？”这个问题暗示了可能还有其他更高级的优化技术（如更复杂的调度、共享内存使用等）可以进一步提升性能。

7.5 局部内存（Local Memory）与性能陷阱

与预期相反，为每个线程分配256个元素（使用局部内存）显著降低了性能。如下图所示，在应用了启动边界（LFB）和循环展开（U）后，性能从79.21%急剧下降到37.50%。

原因分析：

• 局部内存空间（Local Memory Space）位于设备内存（Device Memory）中，而不是片上（on-chip）的高速缓存。
• 因此，它与全局内存（Global Memory）访问具有相同的高延迟。

如下图所示，代码中声明的 thread_data 数组被分配在局部内存中。编译器会生成 st.local 和 ld.local 指令，这些指令最终会访问高延迟的设备内存。这与使用 __shared__ 关键字声明的、位于低延迟片上共享内存中的数据形成鲜明对比。

7.6 编译器优化指令

启动边界 (Launch Bounds)

在编译时，NVCC 编译器并不知道线程块（thread block）的大小，它可能大到1024。如果每个线程使用过多的寄存器，一个大的线程块可能无法装载到流式多处理器（SM）上。

我们可以通过 __launch_bounds__(BLOCK_SIZE) 来告知编译器线程块大小的上限。这会影响 NVCC 的启发式策略，允许为每个线程分配更多寄存器，从而避免寄存器溢出（register spilling，即将寄存器中的数据存回局部/全局内存）。

注意：如果内核执行时设置的每块线程数超过了其启动边界，内核将无法启动。

Pragma Unroll

默认情况下，NVCC 会展开小的循环。如果开发者不同意编译器的启发式判断，可以使用 #pragma unroll 强制展开循环，或使用 #pragma unroll 1 来禁用循环展开。

编译器指令的综合效果

通过结合使用启动边界（Launch Bounds）和强制循环展开（Pragma Unroll）等编译器指令（"Knobs"），性能得到了显著回升，从37.50%提升到了65.83%。然而，这仍然远低于最初的79.21%（LFB）和Thrust库的84.40%。这表明还有其他更关键的性能瓶颈。

7.7 理解并优化占用率 (Occupancy)

SM 资源

• 像寄存器文件（register file）这样的资源是每个 SM 独有的。
• 一个 SM 的寄存器文件通常由 64K 个 32位寄存器组成。
• 如果线程块所需的共享资源（如寄存器、共享内存）能够满足，SM 可以并发运行多个线程块。

占用率计算

占用率（Occupancy）是指一个 SM 上并发活跃的线程束（warp）数量与该 SM 支持的最大活跃线程束数量的比例。它直接受每个线程所需寄存器数量的影响。

• 高寄存器用量示例: 当 items_per_thread = 256 时，每个线程需要255个寄存器，导致整个线程块需要超过17k个寄存器。这远超 SM 的寄存器文件容量，导致 SM 一次只能容纳1个线程块。
• 低寄存器用量示例: 当 items_per_thread = 8 时，每个线程仅需13个寄存器，整个线程块需要约4k个寄存器。这样，一个 SM 可以同时容纳8个线程块。

更高的占用率意味着 SM 上有更多的线程束可以并发执行，这有助于隐藏内存访问延迟，从而提升性能。

Nsight Compute 占用率图

NVIDIA Nsight Compute 工具可以直观地展示占用率图（Occupancy Graph），该图表显示了占用率与每线程寄存器数量之间的关系。

• 该图可以帮助你快速识别限制内核占用率的瓶颈。
• 根据经验法则，优化目标通常是最大化占用率。

下图中，蓝点（items_per_thread = 8）位于高占用率区域，而绿点（items_per_thread = 256）位于低占用率区域。

Occupancy API

CUDA 提供了运行时函数来帮助计算占用率。

• 对于给定的内核和块大小，可以使用 cudaOccupancyMaxActiveBlocksPerMultiprocessor 查询每个 SM 可容纳的最大线程块数量。
• 也可以使用 cudaOccupancyMaxPotentialBlockSize 查询能够实现最大 SM 占用率的块大小。
• Thrust 库在其 transform 算法中正是使用了这个 API 来最大化占用率。

8. 高级优化技术

8.1 向量化内存访问

通过将内存访问从单个元素（如 half）合并为向量类型（如 half2），可以更有效地利用内存带宽。例如，一次性加载两个 half 类型的数据，而不是分两次加载。

向量化内存访问可以进一步提升性能。如下图所示，"Vec"版本的性能（82.18%）超过了之前的"LFB"版本（79.21%），但仍然与Thrust（84.40%）有差距，说明还有优化空间。

8.2 高效的内存移动与 PTX

为了进一步挖掘性能，可以使用异步数据移动（Asynchronous Data Movement）技术。通过 cuda::pipeline 和共享内存，可以在计算的同时，异步地将下一批数据从全局内存预取到共享内存中，从而有效隐藏数据传输的延迟。

这通常涉及直接使用 PTX（并行线程执行）指令集来实现底层的异步拷贝操作。

PTX 是 NVIDIA GPU 的一种低级中间表示语言。在 C++/CUDA 中嵌入 PTX 汇编代码（Inline PTX）有助于：
* 对生成的代码进行最大程度的控制。
* 在高级 C++ API 可用之前，提前实验新的硬件功能。

下图展示了实现异步内存拷贝的复杂 PTX 内联汇编代码。

直接编写内联 PTX 效率不高。推荐使用 cuda::ptx 等库为 C++ 提供封装，这样既能利用 PTX 的强大功能，又不会有抽象开销。

9. 结论与核心要点

9.1 最好的内核是你不必写的那个

幻灯片最后展示了 Thrust 库中 transform 操作的部分实现代码。这段代码极为复杂，融合了之前讨论的所有优化技巧，包括占用率计算、向量化、异步数据移动等。

这说明，虽然从头编写和优化 CUDA 内核是可能的，但要达到像 Thrust 这样的专业库的性能水平，需要深厚的专业知识和巨大的工程努力。因此，在大多数情况下，利用现有的、高度优化的库（如 Thrust, CUB, cuBLAS 等）是更明智、更高效的选择。

9.2 协同库 (Cooperative Libraries)

当加速库不支持您的特定用例时，编写自定义的CUDA C++核函数就变得必要。然而，这并不意味着需要从头开始实现每一个算法。CUDA提供了多种协同库，可以使您的核函数运行得更快，并缩短开发时间。

• CUB (Cooperative Unbounded Back-end)

• 在CUDA核函数内部使用的协同通用算法库。

• cuBLASDx

• 在CUDA核函数内部使用的协同线性代数函数库。

• cuFFTDx

• 在CUDA核函数内部使用的协同快速傅里叶变换库。

您可以在加速计算中心（accelerated computing hub）找到更多关于协同算法的信息。

9.3 核心要点 (Takeaways)

编写自定义CUDA核函数的时机
- 仅当您所需的应用场景未被现有加速库覆盖时。

使用“光速分析法”(Speed-of-light analysis)
- 发现算法瓶颈。
- 发现优化机会。
- 理解何时应停止优化。
- 在没有参考实现的情况下进行操作。

拥抱现代CUDA C++
- 使用现代词汇类型（如 span, mdspan 等）。
- 使用 libcu++ 以利用最新的硬件特性。
- 用原子操作（atomics）替换内部函数（intrinsics）。

使用设备端库 (device-side libraries)
- 使用CUB进行块（block）级和线程束（warp）级的通用算法。
- 使用cuBLASDx进行线性代数计算。
- 使用cuFFTDx进行快速傅里叶变换（FFT）。
- 等等。

使用开发者工具
- 使用 compute-sanitizer 进行CUDA核函数的正确性检查。
- 使用 cuda-gdb 进行CUDA核函数的调试。
- 使用 NVIDIA Nsight Systems 识别应用程序中的瓶颈。
- 使用 NVIDIA Nsight Compute 进行深入的核函数性能分析。