新特性

协程组： https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#cooperative-groups

cudaLaunchKernelEx

🔹 完整代码结构

#ifndef LAUNCH_KERNEL
#ifndef DISABLE_SM90_FEATURES
#define LAUNCH_KERNEL(config, kernel, ...) CUDA_CHECK(cudaLaunchKernelEx(config, kernel, ##__VA_ARGS__))
#else
#define LAUNCH_KERNEL(config, kernel, ...) \
do { \
    kernel<<<__num_sms, __num_threads, 0, __stream>>>(__VA_ARGS__); \
    cudaError_t e = cudaGetLastError(); \
    if (e != cudaSuccess) { \
        EPException cuda_exception("CUDA", __FILE__, __LINE__, cudaGetErrorString(e)); \
        fprintf(stderr, "%s\n", cuda_exception.what()); \
        throw cuda_exception; \
    } \
} while (0)
#endif
#endif

这段代码定义了一个宏 LAUNCH_KERNEL，它根据编译时的宏定义条件，提供了两种不同的 CUDA 内核启动方式。其目的是为了在支持新特性的设备（如计算能力 9.0+ 的 GPU）上使用更高级的启动接口，而在不支持的设备上回退到传统的 <<<>>> 语法。

• #ifndef LAUNCH_KERNEL：如果尚未定义 LAUNCH_KERNEL 宏，则进入定义流程。
• #ifndef DISABLE_SM90_FEATURES：如果 没有禁用 SM90 特性（即启用新特性），则使用 cudaLaunchKernelEx。
• 否则（即禁用了 SM90 特性或目标设备不支持），使用传统的 <<<>>> 语法。

---

🔧 两种 Launch 方式的详细解释

---

🟢 方式一：cudaLaunchKernelEx（现代方式）

CUDA_CHECK(cudaLaunchKernelEx(config, kernel, ##__VA_ARGS__))

1. cudaLaunchKernelEx

• 这是 NVIDIA 在 CUDA 12.0+ 引入的 新式内核启动 API。
• 它允许通过一个 config 结构体（cudaLaunchConfig_t 或扩展）来精细控制内核启动参数，比如：
  • 动态并行化配置
  • 预编译内核句柄
  • 分布式共享内存
  • 更灵活的资源调度
• 支持更高级的功能，特别是在 Hopper 架构（SM 9.0） 上启用新特性（如 cudaLimitMaxL2FetchGranularity、异步传输等）。

2. config 参数

• 是一个结构体，包含 .gridDim, .blockDim, .sharedMem, .stream, .kernel 等字段。
• 允许在运行时动态构造启动配置，比 <<<>>> 更灵活。

3. CUDA_CHECK(...)

• 是一个宏，用于检查 CUDA 错误（封装了 cudaGetLastError() 和 cudaDeviceSynchronize() 等）。
• 如果调用失败，会抛出异常或终止程序。

✅ 优点：

• 更现代、更安全、更可扩展。
• 支持未来 GPU 架构的新特性（尤其是 Hopper SM9.0+）。
• 可用于构建高性能、可调优的运行时系统（如深度学习框架）。

⛔ 要求：

• 需要较新的 CUDA Toolkit（12.0+）。
• 目标 GPU 计算能力 ≥ 9.0（Hopper）才能发挥全部优势。

---

🔴 方式二：传统 <<<>>> 语法 + 手动错误检查

kernel<<<__num_sms, __num_threads, 0, __stream>>>(__VA_ARGS__); \
cudaError_t e = cudaGetLastError(); \
if (e != cudaSuccess) { \
    EPException cuda_exception("CUDA", __FILE__, __LINE__, cudaGetErrorString(e)); \
    fprintf(stderr, "%s\n", cuda_exception.what()); \
    throw cuda_exception; \
}

1. <<<>>> 语法

• 传统的 CUDA 内核启动方式。
• __num_sms：网格大小（grid size），通常是 SM 数量或倍数。
• __num_threads：每个线程块的线程数（block size）。
• 0：共享内存大小。
• __stream：CUDA 流。

2. cudaGetLastError()

• 由于 <<<>>> 语法是异步的，它只负责将内核排队，并不立即执行。
• 调用后必须用 cudaGetLastError() 检查是否有语法错误（如非法的 grid/block size）。
• 注意：它不能捕获内核内部的运行时错误（需用 cudaDeviceSynchronize() + cudaGetLastError() 才能捕获）。

3. 自定义异常处理

• 使用 EPException 抛出异常（可能是某个项目自定义的异常类）。
• 打印文件名、行号、错误信息，便于调试。

✅ 优点：

• 兼容所有 CUDA 版本和 GPU 架构。
• 简单直观，广泛使用。

⛔ 缺点：

• 不支持 SM9.0+ 的新特性（如 Cooperative Groups 增强、新的调度策略等）。
• 语法不够灵活，难以动态配置。
• 错误检查不如 cudaLaunchKernelEx 完善。

---

🆚 两种方式的核心区别总结

特性	cudaLaunchKernelEx	传统 <<<>>>
引入时间	CUDA 12.0+	CUDA 1.0+
API 类型	函数式 API	语言扩展语法
灵活性	高（可通过 config 动态设置）	低（编译时确定）
支持 SM9.0+ 新特性	✅ 是	❌ 否
错误检查	通过 CUDA_CHECK 封装	手动 cudaGetLastError()
兼容性	仅 CUDA 12.0+，推荐 SM9.0+	所有平台
适用场景	高性能框架、Hopper 架构优化	兼容旧设备、简单项目

---