CUDA 编程指南 13.4.2 · 非官方中文译本

原书编排 · 13.4.2

完整目录

章节编号与顺序保持原书编排。

CUDA 编程指南
1. CUDA 简介
1.1. 简介
1.1.1. 图形处理器
1.1.2. 使用 GPU 的优势
1.1.3. 快速入门
1.2. 编程模型
1.2.1. 异构系统
1.2.2. GPU 硬件模型
1.2.2.1. 线程块与网格
1.2.2.1.1. 线程块簇
1.2.2.2. 线程束与 SIMT
1.2.2.3. CUDA 中的分块编程
1.2.2.3.1. 数组与分块
1.2.2.3.2. 分块空间与数据移动
1.2.2.3.3. 分块操作
1.2.2.3.4. 与 SIMT 编程的关系
1.2.3. GPU 内存
1.2.3.1. 异构系统中的 DRAM 内存
1.2.3.2. 局部性域：
1.2.3.3. GPU 中的片上内存
1.2.3.3.1. 缓存
1.2.3.4. 统一内存
1.3. CUDA 平台
1.3.1. 计算能力与流式多处理器（SM）版本
1.3.2. CUDA Toolkit 与 NVIDIA 驱动程序
1.3.2.1. CUDA 运行时 API 与 CUDA 驱动程序 API
1.3.3. 并行线程执行（PTX）
1.3.4. Cubins 与 Fatbins
1.3.4.1. 二进制兼容性
1.3.4.2. PTX 兼容性
1.3.4.3. 即时编译
1.3.4.4. 二进制最终化
2. 使用 CUDA 进行 GPU 编程
2.1. CUDA C++ 入门
2.1.1. 使用 NVCC 编译
2.1.2. 核函数
2.1.2.1. 指定核函数
2.1.2.2. 启动核函数
2.1.2.2.1. 三重尖括号记法
2.1.2.3. 线程和网格索引内建函数
2.1.2.3.1. 边界检查
2.1.3. GPU 计算中的内存
2.1.3.1. 统一内存
2.1.3.2. 显式内存管理
2.1.3.3. 内存管理与应用程序性能
2.1.4. 同步 CPU 与 GPU
2.1.5. 整合所有内容
2.1.6. 运行时初始化
2.1.7. CUDA 中的错误检查
2.1.7.1. 错误状态
2.1.7.2. 异步错误
2.1.7.3. CUDA_LOG_FILE
2.1.8. 设备函数与主机函数
2.1.9. 变量说明符
2.1.9.1. 检测设备编译
2.1.10. 线程块簇
2.1.10.1. 使用三重尖括号表示法进行带簇的启动
2.2. CUDA Python 简介
2.2.1. CUDA Python 生态系统
2.2.1.1. 在 Python 中使用 CUDA 库
2.2.1.2. 本章范围
2.2.1.3. 环境设置
2.2.1.4. 运行 CUDA Python 应用程序
2.2.2. Python 中的 SIMT 核函数
2.2.2.1. 指定核函数
2.2.2.2. 启动核函数
2.2.2.2.1. 核函数启动
2.2.2.2.2. 多维网格和线程块
2.2.2.3. 线程和网格索引内建函数
2.2.3. GPU 计算中的内存
2.2.3.1. 在 GPU 上实例化数组
2.2.3.2. 在主机内存与 GPU 内存之间复制数组
2.2.3.3. ndarray 对象类型
2.2.4. 同步 CPU 与 GPU
2.2.5. 综合运用
2.2.6. CUDA Python 中的错误检查
2.3. 编写 SIMT 核函数
2.3.1. SIMT 基础
2.3.2. 线程层次结构
2.3.2.1. 线程块同步
2.3.3. GPU 设备内存空间
2.3.3.1. 全局内存
2.3.3.2. 共享内存
2.3.3.2.1. 共享内存的静态分配
2.3.3.2.2. 共享内存的动态分配
2.3.3.3. 寄存器
2.3.3.4. 局部内存
2.3.3.5. 常量内存
2.3.3.6. 缓存
2.3.3.7. 纹理和表面内存
2.3.3.8. 分布式共享内存
2.3.4. 内存性能
2.3.4.1. 合并的全局内存访问
2.3.4.1.1. 使用全局内存的矩阵转置示例
2.3.4.2. 共享内存访问模式
2.3.4.2.1. 使用共享内存的矩阵转置示例
2.3.4.2.2. 共享内存存储体冲突
2.3.5. 原子操作
2.3.5.1. 类似于 C++ std::atomic 的原子操作
2.3.5.2. Python 中的内存原子操作
2.3.6. 协作组
2.3.7. 核函数启动与占用率
2.4. 编写分块核函数
2.4.1. 核函数和函数声明
2.4.2. 启动核函数
2.4.2.1. 网格大小设置模式
2.4.3. 查询块的位置
2.4.4. 创建分块
2.4.5. 编译时常量
2.4.5.1. Python Constant[T]
2.4.5.2. C++ integral_constant 和 _ic 字面量
2.4.6. 加载和存储分块
2.4.6.1. 分块空间加载与存储
2.4.6.1.1. 分区视图加载与存储
2.4.6.1.2. Python 单次调用加载与存储
2.4.6.1.3. 分块空间边界处理
2.4.6.2. 聚集与散布
2.4.6.2.1. 聚集与散布的边界处理
2.4.7. 控制流
2.4.7.1. 循环
2.4.7.2. 条件语句
2.4.8. 逐元素算术运算与广播
2.4.8.1. 广播
2.4.8.2. 算术运算符
2.4.9. 分块原语
2.4.9.1. 矩阵乘法
2.4.9.2. 归约与扫描
2.4.9.3. 转置与置换
2.4.9.4. 逐元素选择
2.4.9.5. 数学函数
2.4.10. 原子内存操作
2.4.10.1. 块间争用
2.4.10.2. 块内争用
2.4.10.3. 支持的原子操作
2.4.11. 优化提示
2.4.11.1. C++： cutile::hint 属性
2.4.11.2. Python：装饰器参数与调用点关键字
2.4.11.3. 提示种类
2.4.11.3.1. 每个簇的 CTA 数量
2.4.11.3.2. 占用率
2.4.11.3.3. 内存访问延迟
2.4.11.3.4. 允许 TMA
2.4.12. C++ 性能提示
2.4.12.1. 对内存中的数组使用 __restrict__ 指针
2.4.12.2. 将数组指针标记为按 16 字节对齐
2.4.12.3. 内存访问优先使用 ct::partition_view
2.4.12.4. 对有界循环使用 ct::irange
2.5. 异步执行
2.5.1. 什么是异步并发执行？
2.5.2. CUDA 流
2.5.2.1. 创建和销毁 CUDA 流
2.5.2.2. 在 CUDA 流中启动核函数
2.5.2.3. 在 CUDA 流中发起内存传输
2.5.2.4. 流同步
2.5.3. CUDA 事件
2.5.3.1. 创建和销毁 CUDA 事件
2.5.3.2. 向 CUDA 流中插入事件
2.5.3.3. 对 CUDA 流中的操作计时
2.5.3.4. 检查 CUDA 事件的状态
2.5.4. 来自流的回调函数
2.5.4.1. 使用 cudaStreamAddCallback()
2.5.4.2. 异步错误处理
2.5.5. CUDA 流顺序
2.5.6. 阻塞流、非阻塞流和默认流
2.5.6.1. 传统默认流
2.5.6.2. 每线程默认流
2.5.7. 显式同步
2.5.8. 隐式同步
2.5.9. 其他主题与高级主题
2.5.9.1. 流优先级设置
2.5.9.2. 通过流捕获介绍 CUDA 图
2.5.10. 异步执行小结
2.6. 统一内存与系统内存
2.6.1. 统一虚拟地址空间
2.6.2. 统一内存
2.6.2.1. 统一内存范式
2.6.2.1.1. 统一内存范式：代码示例
2.6.2.2. 统一内存功能的完整支持
2.6.2.2.1. 具有硬件相干性的完整统一内存
2.6.2.2.2. HMM：具有软件相干性的完整统一内存
2.6.2.3. 有限的统一内存支持
2.6.2.4. 内存建议和预取
2.6.3. 页锁定主机内存
2.6.3.1. 映射内存
2.6.3.1.1. cudaMallocHost 和 cudaHostAlloc
2.6.3.1.2. cudaHostRegister
2.6.3.1.3. 统一内存与映射内存的比较
2.6.4. 总结
2.7. NVCC：NVIDIA CUDA 编译器
2.7.1. CUDA 源文件和头文件
2.7.2. NVCC 编译工作流
2.7.3. NVCC 基本用法
2.7.3.1. NVCC 的 PTX 和 Cubin 生成
2.7.3.2. 主机代码编译注意事项
2.7.3.3. GPU 代码的分离编译
2.7.4. 常用编译器选项
2.7.4.1. 语言特性
2.7.4.2. 调试选项
2.7.4.3. 优化选项
2.7.4.4. 链接时优化（LTO）
2.7.4.5. 性能分析选项
2.7.4.6. Fatbin 压缩
2.7.4.7. 编译器性能控制
3. 高级 CUDA
3.1. 高级 CUDA API 和特性
3.1.1. cudaLaunchKernelEx
3.1.2. 启动簇：
3.1.2.1. 使用 cudaLaunchKernelEx 以簇方式启动
3.1.2.2. 将线程块视为簇
3.1.3. 关于流和事件的更多内容
3.1.3.1. 流优先级
3.1.3.2. 显式同步
3.1.3.3. 隐式同步
3.1.4. 程序化依赖核函数启动
3.1.5. 批量内存传输
3.1.6. 环境变量
3.2. 高级核函数编程
3.2.1. 使用 PTX
3.2.2. 硬件实现
3.2.2.1. SIMT 执行模型
3.2.2.1.1. 独立线程调度
3.2.2.2. 硬件多线程
3.2.2.3. 异步执行特性
3.2.2.3.1. 异步线程与异步代理
3.2.3. 线程作用域
3.2.4. 高级同步原语
3.2.4.1. 限定作用域的原子操作
3.2.4.1.1. 线程作用域与内存顺序
3.2.4.1.2. 性能注意事项
3.2.4.2. 异步屏障
3.2.4.3. 流水线
3.2.5. 异步数据复制
3.2.6. 配置 L1/共享内存的平衡
3.3. CUDA 驱动程序 API
3.3.1. 上下文
3.3.2. 模块
3.3.3. 核函数执行
3.3.4. 运行时 API 与驱动程序 API 之间的互操作性
3.4. 多 GPU 系统编程
3.4.1. 多设备上下文与执行管理
3.4.1.1. 设备枚举
3.4.1.2. 设备选择
3.4.1.3. 多设备流、事件和内存复制行为
3.4.2. 多设备对等传输和内存访问
3.4.2.1. 对等内存传输
3.4.2.2. 对等内存访问
3.4.2.3. 对等内存一致性
3.4.2.4. 多设备托管内存
3.4.2.5. 主机 IOMMU 硬件、PCI 访问控制服务和虚拟机
3.4.2.6. 计算互连传输
3.5. CUDA 功能概览
3.5.1. 提高核函数性能
3.5.1.1. 异步屏障
3.5.1.2. 异步数据复制与张量内存加速器（TMA）
3.5.1.3. 流水线
3.5.1.4. 通过簇启动控制进行工作窃取
3.5.2. 改善延迟
3.5.2.1. 绿色上下文
3.5.2.2. 局部性域
3.5.2.3. 流有序内存分配
3.5.2.4. CUDA 图
3.5.2.5. 程序化依赖启动
3.5.2.6. 延迟加载
3.5.3. 功能性特性
3.5.3.1. 扩展 GPU 内存
3.5.3.2. 动态并行
3.5.4. CUDA 互操作性
3.5.4.1. CUDA 与其他 API 的互操作性
3.5.4.2. 进程间通信
3.5.5. 细粒度控制
3.5.5.1. 虚拟内存管理
3.5.5.2. 计算互连传输
3.5.5.3. 驱动入口点访问
3.5.5.4. 错误日志管理
4. CUDA 特性
4.1. 统一内存
4.1.1. 完全支持 CUDA 统一内存的设备上的统一内存
4.1.1.1. 统一内存：深入示例
4.1.1.1.1. 以文件为后备的统一内存
4.1.1.1.2. 使用统一内存进行进程间通信（IPC）
4.1.1.2. 性能调优
4.1.1.2.1. 内存分页与页大小
4.1.1.2.1.1. 选择合适的页大小
4.1.1.2.1.2. CPU 和 GPU 页表：硬件相干性与软件相干性
4.1.1.2.1.3. 混合使用硬件相干性和软件相干性
4.1.1.2.2. 从主机直接访问统一内存
4.1.1.2.3. 主机原生原子操作
4.1.1.2.4. 原子访问和同步原语
4.1.1.2.5. 使用统一内存时 Memcpy()/Memset() 的行为
4.1.1.2.6. 统一内存的内存分配器概述
4.1.1.2.7. 局部性域和统一内存
4.1.1.2.8. 访问计数器迁移
4.1.1.2.9. 避免从 CPU 频繁写入驻留于 GPU 的内存
4.1.1.2.10. 利用对系统内存的异步访问
4.1.2. 仅支持 CUDA 托管内存的设备上的统一内存
4.1.3. Windows、WSL 和 Tegra 上的统一内存
4.1.3.1. 多 GPU
4.1.3.2. 相干性和并发性
4.1.3.3. 与流关联的统一内存
4.1.3.3.1. 流回调
4.1.3.3.2. 与流关联的托管内存允许更细粒度的控制
4.1.3.3.3. 关于多线程主机程序的一个更详细的示例
4.1.3.3.4. 与流关联的统一内存的数据移动
4.1.4. 性能提示
4.1.4.1. 数据预取
4.1.4.2. 数据使用提示
4.1.4.3. 内存丢弃
4.1.4.4. 查询托管内存的数据使用属性
4.1.4.5. GPU 内存超额使用
4.2. CUDA 图
4.2.1. 图结构
4.2.1.1. 节点类型
4.2.1.2. 边数据
4.2.2. 构建和运行图
4.2.2.1. 图的创建
4.2.2.1.1. 图 API
4.2.2.1.2. 流捕获
4.2.2.1.2.1. 跨流依赖关系与事件
4.2.2.1.2.2. 禁止的操作与未处理的操作
4.2.2.1.2.3. 失效
4.2.2.1.2.4. 捕获内省
4.2.2.1.3. 综合运用
4.2.2.2. 图的实例化
4.2.2.3. 图的执行
4.2.3. 更新已实例化的图
4.2.3.1. 整图更新
4.2.3.2. 单个节点更新
4.2.3.3. 单个节点启用
4.2.3.4. 图更新限制
4.2.4. 条件图节点
4.2.4.1. 条件句柄
4.2.4.2. 条件节点主体图要求
4.2.4.3. 条件 IF 节点
4.2.4.4. 条件 WHILE 节点
4.2.4.5. 条件 SWITCH 节点
4.2.5. 图内存节点
4.2.5.1. 简介
4.2.5.2. API 基础
4.2.5.2.1. 图节点 API
4.2.5.2.2. 流捕获
4.2.5.2.3. 在执行分配的图之外访问和释放图内存
4.2.5.2.4. cudaGraphInstantiateFlagAutoFreeOnLaunch
4.2.5.2.5. 子图中的内存节点
4.2.5.3. 优化的内存复用
4.2.5.3.1. 图内地址复用
4.2.5.3.2. 物理内存管理与共享
4.2.5.4. 性能注意事项
4.2.5.4.1. 首次启动 / cudaGraphUpload
4.2.5.5. 物理内存占用
4.2.5.6. 对等访问
4.2.5.6.1. 通过图节点 API 进行对等访问
4.2.5.6.2. 通过流捕获进行对等访问
4.2.6. 设备图启动
4.2.6.1. 设备图创建
4.2.6.1.1. 设备图要求
4.2.6.1.2. 设备图上传
4.2.6.1.3. 设备图更新
4.2.6.2. 设备启动
4.2.6.2.1. 即发即弃启动
4.2.6.2.1.1. 图执行环境
4.2.6.2.2. 尾部启动
4.2.6.2.2.1. 尾部自启动
4.2.6.2.3. 同级启动
4.2.7. 使用图 API
4.2.8. CUDA 用户对象
4.3. 流有序内存分配器
4.3.1. 简介
4.3.2. 内存管理
4.3.2.1. 分配内存
4.3.2.2. 释放内存
4.3.3. 内存池
4.3.3.1. 默认/隐式池
4.3.3.2. 显式池
4.3.3.3. 用于多 GPU 支持的设备可访问性
4.3.3.4. 为 IPC 启用内存池
4.3.3.4.1. 创建和共享 IPC 内存池
4.3.3.4.2. 在导入进程中设置访问权限
4.3.3.4.3. 从导出的池中创建和共享内存分配
4.3.3.4.4. IPC 导出池的限制
4.3.3.4.5. IPC 导入池的限制
4.3.4. 最佳实践与调优
4.3.4.1. 查询支持情况
4.3.4.2. 物理页缓存行为
4.3.4.3. 资源使用统计
4.3.4.4. 内存重用策略
4.3.4.4.1. cudaMemPoolReuseFollowEventDependencies
4.3.4.4.2. cudaMemPoolReuseAllowOpportunistic
4.3.4.4.3. cudaMemPoolReuseAllowInternalDependencies
4.3.4.4.4. 禁用重用策略
4.3.4.5. 同步 API 的操作
4.3.5. 补充说明
4.3.5.1. cudaMemcpyAsync 对当前上下文/设备的敏感性
4.3.5.2. cudaPointerGetAttributes 查询
4.3.5.3. cudaGraphAddMemsetNode
4.3.5.4. 指针属性
4.3.5.5. CPU 虚拟内存
4.4. 协作组
4.4.1. 简介
4.4.2. 协作组句柄与成员函数
4.4.3. 默认行为/无组执行
4.4.3.1. 尽早创建隐式组句柄
4.4.3.2. 仅按引用传递组句柄
4.4.4. 创建协作组
4.4.4.1. 避免组创建风险
4.4.5. 同步
4.4.5.1. 同步
4.4.5.2. 屏障
4.4.6. 集合操作
4.4.6.1. 归约
4.4.6.2. 扫描
4.4.6.3. 单线程调用
4.4.7. 异步数据移动
4.4.7.1. 异步内存复制的对齐要求
4.4.8. 大规模组
4.4.8.1. 何时使用 cudaLaunchCooperativeKernel
4.5. 编程式依赖启动与同步
4.5.1. 背景
4.5.2. API 说明
4.5.3. 在 CUDA 图中使用
4.6. 绿色上下文
4.6.1. 动机 / 何时使用
4.6.2. 绿色上下文：易用性
4.6.3. 绿色上下文：设备资源和资源描述符
4.6.4. 绿色上下文创建示例
4.6.4.1. 步骤 1：获取可用的 GPU 资源
4.6.4.2. 步骤 2：划分 SM 资源
4.6.4.3. 步骤 2（续）：添加工作队列资源
4.6.4.4. 步骤 3：创建资源描述符
4.6.4.5. 步骤 4：创建绿色上下文
4.6.5. 绿色上下文：启动工作
4.6.6. 其他执行上下文 API
4.6.7. 绿色上下文示例
4.7. 局部性域
4.7.1. 发现局部性域
4.7.2. 计算资源局部化
4.7.3. 分配局部化设备内存
4.7.3.1. 创建局部化内存池
4.7.3.2. 创建局部化 VMM 分配
4.7.4. 查询内存放置情况
4.7.5. 编程指导
4.7.5.1. 采用分叉/汇合模式的局部化核函数
4.7.5.2. CUDA 图
4.7.5.2.1. 流捕获
4.7.5.2.2. 图构建
4.7.5.3. 兼容性注意事项
4.8. 延迟加载
4.8.1. 简介
4.8.2. 变更历史
4.8.3. 延迟加载的要求
4.8.3.1. CUDA 运行时版本要求
4.8.3.2. CUDA 驱动程序版本要求
4.8.3.3. 编译器要求
4.8.3.4. 核函数要求
4.8.4. 用法
4.8.4.1. 启用与禁用
4.8.4.2. 在运行时检查是否启用了延迟加载
4.8.4.3. 在运行时强制立即加载模块
4.8.5. 潜在风险
4.8.5.1. 对核函数并发执行的影响
4.8.5.2. 大规模内存分配
4.8.5.3. 对性能测量的影响
4.9. 错误日志管理
4.9.1. 背景
4.9.2. 启用
4.9.3. 输出
4.9.4. API 说明
4.9.4.1. 回调注册
4.9.4.2. 迭代器与日志转储
4.9.4.3. 转储行为
4.9.5. 限制与已知问题
4.10. 异步屏障
4.10.1. 初始化
4.10.2. 屏障的阶段：到达、递减计数、完成与重置
4.10.2.1. 线程束纠缠
4.10.3. 显式阶段跟踪
4.10.4. 提前退出
4.10.5. 完成函数
4.10.6. 跟踪异步内存操作
4.10.7. 使用屏障的生产者-消费者模式
4.11. 流水线
4.11.1. 初始化
4.11.2. 提交工作
4.11.3. 消费工作
4.11.4. 线程束纠缠
4.11.5. 提前退出
4.11.6. 跟踪异步内存操作
4.11.7. 使用流水线的生产者-消费者模式
4.12. 异步数据复制
4.12.1. 使用 LDGSTS
4.12.1.1. 在条件代码中批量执行加载
4.12.1.2. 预取数据
4.12.1.3. 通过线程束专门化实现生产者-消费者模式
4.12.2. 使用张量内存加速器（TMA）
4.12.2.1. 使用 TMA 传输一维数组
4.12.2.1.1. 预取数据
4.12.2.2. 使用 TMA 传输多维数组
4.12.2.2.1. 在设备上编码张量映射
4.12.2.2.2. 在设备端编码和修改张量映射
4.12.2.2.3. 使用修改后的张量映射
4.12.2.2.4. 使用驱动程序 API 创建模板张量映射值
4.12.2.2.5. 共享内存存储体重排
4.12.3. 使用 STAS
4.13. 使用簇启动控制进行工作窃取
4.13.1. API 细节
4.13.1.1. 线程块取消
4.13.1.2. 线程块取消的约束
4.13.2. 示例：向量与标量乘法
4.13.2.1. 用例：线程块
4.13.2.2. 用例：线程块簇
4.14. L2 缓存控制
4.14.1. 为持久访问预留 L2 缓存
4.14.2. 持久访问的 L2 策略
4.14.3. L2 访问属性
4.14.4. L2 持久性示例
4.14.5. 将 L2 访问重置为普通状态
4.14.6. 管理 L2 预留缓存的利用情况
4.14.7. 查询 L2 缓存属性
4.14.8. 控制为持久内存访问预留的 L2 缓存大小
4.15. 内存同步域
4.15.1. 内存栅栏干扰
4.15.2. 使用域隔离流量
4.15.3. 在 CUDA 中使用域
4.16. 进程间通信
4.16.1. 使用旧版进程间通信 API 进行 IPC
4.16.2. 使用虚拟内存管理 API 进行 IPC
4.17. 虚拟内存管理
4.17.1. 预备知识
4.17.1.1. 定义
4.17.1.2. 查询支持情况
4.17.2. API 概述
4.17.3. 单播内存共享
4.17.3.1. 分配和导出
4.17.3.2. 共享和导入
4.17.3.3. 预留和映射
4.17.3.4. 访问权限
4.17.3.5. 释放内存
4.17.4. 多播内存共享
4.17.4.1. 分配多播对象
4.17.4.2. 向多播对象添加设备
4.17.4.3. 将内存绑定到多播对象
4.17.4.4. 使用多播映射
4.17.5. 高级配置
4.17.5.1. 内存类型
4.17.5.2. 可压缩内存
4.17.5.3. 虚拟别名支持
4.17.5.4. 用于 IPC 的操作系统特定句柄详情
4.18. 计算互连传输
4.18.1. 先决条件与范围
4.18.2. 预备知识
4.18.2.1. 定义
4.18.2.2. 查询支持情况
4.18.3. API 概述
4.18.4. 验证互连组成员资格
4.18.5. 创建端点
4.18.6. 逻辑端点 id
4.18.6.1. 限制与对齐
4.18.7. 共享端点
4.18.7.1. 确认就绪状态
4.18.8. 绑定内存
4.18.9. 互连网络（fabric）操作
4.18.9.1. 逐线程操作与线程束集体操作
4.18.10. 完成状态与错误报告
4.18.10.1. 配置屏障到达计数与事务计数
4.18.11. 清理
4.19. 扩展 GPU 内存
4.19.1. 预备知识
4.19.1.1. EGM 平台：系统拓扑
4.19.1.2. 插槽标识符：它们是什么？如何访问它们？
4.19.1.3. 分配器与 EGM 支持
4.19.1.4. 对现有 API 的内存管理扩展
4.19.2. 使用 EGM 接口
4.19.2.1. 单节点、单 GPU
4.19.2.2. 单节点、多 GPU
4.19.2.2.1. 使用 VMM API
4.19.2.2.2. 使用 CUDA 内存池
4.19.2.3. 多节点、多 GPU
4.20. CUDA 动态并行
4.20.1. 简介
4.20.1.1. 概述
4.20.2. 执行环境
4.20.2.1. 父网格与子网格
4.20.2.2. CUDA 原语的作用域
4.20.2.3. 流和事件
4.20.2.4. 顺序与并发
4.20.3. 内存相干性与一致性
4.20.3.1. 全局内存
4.20.3.2. 映射内存
4.20.3.3. 共享内存与局部内存
4.20.3.4. 局部内存
4.20.3.4.1. 纹理内存
4.20.4. 编程接口
4.20.4.1. 基础
4.20.4.2. CDP 的 C++ 语言接口
4.20.4.2.1. 设备端核函数启动
4.20.4.2.1.1. 启动是异步的
4.20.4.2.1.2. 启动环境配置
4.20.4.2.2. 事件
4.20.4.2.3. 同步
4.20.4.2.4. 设备管理
4.20.5. 编程指南
4.20.5.1. 性能
4.20.5.1.1. 启用动态并行的核函数开销
4.20.5.2. 实现的约束与限制
4.20.5.2.1. 运行时
4.20.5.2.1.1. 内存占用
4.20.5.2.1.2. 待处理的核函数启动
4.20.5.3. 兼容性与互操作性
4.20.6. 从 PTX 进行设备端启动
4.20.6.1. 核函数启动 API
4.20.6.1.1. cudaLaunchDevice
4.20.6.1.2. cudaGetParameterBuffer
4.20.6.2. 参数缓冲区布局
4.21. CUDA 与 API 的互操作性
4.21.1. 图形互操作性
4.21.1.1. OpenGL 互操作性
4.21.1.2. Direct3D 互操作性
4.21.1.3. 可扩展链接接口（SLI）配置中的互操作性
4.21.2. 外部资源互操作性
4.21.2.1. Vulkan 互操作性
4.21.2.1.1. 设置 Vulkan 设备
4.21.2.1.2. 使用匹配的设备 UUID 初始化 CUDA
4.21.2.1.3. 导出 Vulkan 内存对象
4.21.2.1.4. 导出 Vulkan 同步对象
4.21.2.1.5. 导入内存对象
4.21.2.1.6. 将缓冲区映射到已导入的内存对象
4.21.2.1.7. 将 mipmap 数组映射到已导入的内存对象
4.21.2.1.8. 导入同步对象
4.21.2.1.9. 向已导入的同步对象发出信号/等待已导入的同步对象
4.21.2.1.10. OpenGL 互操作性
4.21.2.2. Direct3D 互操作性
4.21.2.2.1. 匹配设备 LUID
4.21.2.2.2. 导入内存对象
4.21.2.2.3. 将缓冲区映射到导入的内存对象
4.21.2.2.4. 将多级渐远纹理数组映射到导入的内存对象
4.21.2.2.5. 导入同步对象
4.21.2.2.6. 向导入的同步对象发出信号或等待这些对象
4.21.2.3. NVIDIA 软件通信接口互操作性（NVSCI）
4.21.2.3.1. 导入内存对象
4.21.2.3.2. 将缓冲区映射到导入的内存对象
4.21.2.3.3. 将多级渐远纹理数组映射到导入的内存对象
4.21.2.3.4. 导入同步对象
4.21.2.3.5. 向导入的同步对象发出信号或等待这些对象
4.22. 驱动程序入口点访问
4.22.1. 简介
4.22.2. 驱动程序函数类型定义
4.22.3. 获取驱动程序函数
4.22.3.1. 使用驱动程序 API
4.22.3.2. 使用运行时 API
4.22.3.3. 获取每线程默认流版本
4.22.3.4. 访问新的 CUDA 功能
4.22.4. cuGetProcAddress 使用准则
4.22.4.1. 运行时 API 使用准则
4.22.5. 确定 cuGetProcAddress 失败的原因
5. 技术附录
5.1. 计算能力
5.1.1. 获取 GPU 计算能力
5.1.2. 功能可用性
5.1.2.1. 架构特定功能
5.1.2.2. 系列特定功能
5.1.2.3. 功能集编译器目标
5.1.3. 功能和技术规格
5.2. CUDA 环境变量
5.2.1. 设备枚举和属性
5.2.1.1. CUDA_VISIBLE_DEVICES
5.2.1.2. CUDA_DEVICE_ORDER
5.2.1.3. CUDA_MANAGED_FORCE_DEVICE_ALLOC
5.2.2. JIT 编译
5.2.2.1. CUDA_CACHE_DISABLE
5.2.2.2. CUDA_CACHE_PATH
5.2.2.3. CUDA_CACHE_MAXSIZE
5.2.2.4. CUDA_FORCE_PTX_JIT 和 CUDA_FORCE_JIT
5.2.2.5. CUDA_DISABLE_PTX_JIT 和 CUDA_DISABLE_JIT
5.2.2.6. CUDA_FORCE_PRELOAD_LIBRARIES
5.2.3. 执行
5.2.3.1. CUDA_LAUNCH_BLOCKING
5.2.3.2. CUDA_DEVICE_MAX_CONNECTIONS
5.2.3.3. CUDA_DEVICE_MAX_COPY_CONNECTIONS
5.2.3.4. CUDA_SCALE_LAUNCH_QUEUES
5.2.3.5. CUDA_GRAPHS_USE_NODE_PRIORITY
5.2.3.6. CUDA_DEVICE_WAITS_ON_EXCEPTION
5.2.3.7. CUDA_DEVICE_DEFAULT_PERSISTING_L2_CACHE_PERCENTAGE_LIMIT
5.2.3.8. CUDA_DISABLE_PERF_BOOST
5.2.3.9. CUDA_PREFER_SPINLOCKS_ON_HOST
5.2.3.10. CUDA_AUTO_BOOST [[已弃用]]
5.2.4. 模块加载
5.2.4.1. CUDA_MODULE_LOADING
5.2.4.2. CUDA_MODULE_DATA_LOADING
5.2.4.3. CUDA_BINARY_LOADER_THREAD_COUNT
5.2.4.4. CUDA_DEVICE_DEFAULT_STACK_SIZE_MAX
5.2.5. CUDA 错误日志管理
5.2.5.1. CUDA_LOG_FILE
5.3. C++ 语言支持
5.3.1. C++11 语言特性
5.3.2. C++14 语言特性
5.3.3. C++17 语言特性
5.3.4. C++20 语言特性
5.3.5. C++23 语言特性
5.3.6. CUDA C++ 标准库
5.3.7. C 标准库函数
5.3.7.1. clock() 和 clock64()
5.3.7.2. printf()
5.3.7.3. memcpy() 和 memset()
5.3.7.4. malloc() 和 free()
5.3.7.5. alloca()
5.3.8. Lambda 表达式
5.3.8.1. Lambda 表达式与 __global__ 函数参数
5.3.8.2. 扩展 lambda
5.3.8.3. 扩展 lambda 类型特征
5.3.8.4. 扩展 lambda 的限制
5.3.8.5. 主机-设备 lambda 优化注意事项
5.3.8.6. *this 按值捕获
5.3.8.7. 实参依赖查找（ADL）
5.3.9. 多态函数包装器
5.3.10. C/C++ 语言限制
5.3.10.1. 不支持的特性
5.3.10.2. 命名空间保留
5.3.10.3. 指针和内存地址
5.3.10.4. 变量
5.3.10.4.1. 局部变量
5.3.10.4.2. const 限定的变量
5.3.10.4.3. volatile 限定的变量
5.3.10.4.4. static 变量
5.3.10.4.5. extern 变量
5.3.10.5. 函数
5.3.10.5.1. 递归
5.3.10.5.2. 外部链接
5.3.10.5.3. 形式参数
5.3.10.5.4. __global__ 函数参数
5.3.10.5.5. __global__ 函数实参传递
5.3.10.6. 类
5.3.10.6.1. 类类型变量
5.3.10.6.2. 数据成员
5.3.10.6.3. 函数成员
5.3.10.6.4. 隐式声明的函数和非虚的显式预置函数
5.3.10.6.5. 多态类
5.3.10.6.6. Windows 特有的类布局
5.3.10.7. 模板
5.3.10.8. 分块代码中的限制
5.3.11. C++11 限制
5.3.11.1. inline 命名空间
5.3.11.2. inline 无名命名空间
5.3.11.3. constexpr 函数
5.3.11.4. constexpr 变量
5.3.11.5. __global__ 可变参数模板
5.3.11.6. 预置函数 = default
5.3.11.7. [cuda::]std::initializer_list
5.3.11.8. [cuda::]std::move , [cuda::]std::forward
5.3.12. C++14 限制
5.3.12.1. 具有推导返回类型的函数
5.3.12.2. 变量模板
5.3.13. C++17 限制
5.3.13.1. inline 变量
5.3.13.2. 结构化绑定
5.3.14. C++20 限制
5.3.14.1. 三路比较运算符
5.3.14.2. consteval 函数
5.3.15. C++23 限制
5.3.15.1. 相等运算符（P2468R2）
5.4. C/C++ 语言扩展
5.4.1. 函数和变量注解
5.4.1.1. 执行空间说明符
5.4.1.2. 内存空间说明符
5.4.1.2.1. __shared__ 内存
5.4.1.2.2. __managed__ 内存
5.4.1.2.3. __tile__ 变量
5.4.1.3. 内联说明符
5.4.1.4. __restrict__ 指针
5.4.1.5. __grid_constant__ 参数
5.4.1.6. 注解汇总
5.4.2. 内建类型和变量
5.4.2.1. 主机编译器类型扩展
5.4.2.2. 内建变量
5.4.2.3. 内建类型
5.4.3. 核函数配置
5.4.3.1. 线程块簇
5.4.3.2. 启动边界
5.4.3.3. 每个线程的最大寄存器数量
5.4.4. 同步原语
5.4.4.1. 线程块同步函数
5.4.4.2. 线程束同步函数
5.4.4.3. 内存栅栏函数
5.4.5. 原子函数
5.4.5.1. 旧式原子函数
5.4.5.1.1. atomicAdd()
5.4.5.1.2. atomicSub()
5.4.5.1.3. atomicInc()
5.4.5.1.4. atomicDec()
5.4.5.1.5. atomicAnd()
5.4.5.1.6. atomicOr()
5.4.5.1.7. atomicXor()
5.4.5.1.8. atomicMin()
5.4.5.1.9. atomicMax()
5.4.5.1.10. atomicExch()
5.4.5.1.11. atomicCAS()
5.4.5.2. 内建原子函数
5.4.5.2.1. __nv_atomic_fetch_add() , __nv_atomic_add()
5.4.5.2.2. __nv_atomic_fetch_sub() , __nv_atomic_sub()
5.4.5.2.3. __nv_atomic_fetch_and() , __nv_atomic_and()
5.4.5.2.4. __nv_atomic_fetch_or() , __nv_atomic_or()
5.4.5.2.5. __nv_atomic_fetch_xor() , __nv_atomic_xor()
5.4.5.2.6. __nv_atomic_fetch_min() , __nv_atomic_min()
5.4.5.2.7. __nv_atomic_fetch_max() , __nv_atomic_max()
5.4.5.2.8. __nv_atomic_exchange() , __nv_atomic_exchange_n()
5.4.5.2.9. __nv_atomic_compare_exchange() , __nv_atomic_compare_exchange_n()
5.4.5.2.10. __nv_atomic_load() , __nv_atomic_load_n()
5.4.5.2.11. __nv_atomic_store() , __nv_atomic_store_n()
5.4.5.2.12. __nv_atomic_thread_fence()
5.4.6. 线程束函数
5.4.6.1. 线程束活跃掩码
5.4.6.2. 线程束投票函数
5.4.6.3. 线程束匹配函数
5.4.6.4. 线程束归约函数
5.4.6.5. 线程束洗牌函数
5.4.6.6. 线程束 __sync 内建函数约束
5.4.7. CUDA 特定宏
5.4.7.1. __CUDA_ARCH__
5.4.7.2. __CUDA_ARCH_SPECIFIC__ 和 __CUDA_ARCH_FAMILY_SPECIFIC__
5.4.7.3. CUDA 特性测试宏
5.4.7.4. __nv_pure__ 属性
5.4.8. CUDA 专用函数
5.4.8.1. 地址空间谓词函数
5.4.8.2. 地址空间转换函数
5.4.8.3. 低级加载和存储函数
5.4.8.4. __trap()
5.4.8.5. __nanosleep()
5.4.8.6. 动态规划扩展（DPX）指令
5.4.9. 编译器优化提示
5.4.9.1. #pragma unroll
5.4.9.2. __builtin_assume_aligned()
5.4.9.3. __builtin_assume() 和 __assume()
5.4.9.4. __builtin_constant_p()
5.4.9.5. __builtin_expect()
5.4.9.6. __builtin_unreachable()
5.4.9.7. 自定义 ABI pragma 指令
5.4.9.8. MMA 吞吐量 pragma 指令
5.4.10. 调试与诊断
5.4.10.1. 断言
5.4.10.2. 断点函数
5.4.10.3. 诊断 pragma 指令
5.4.11. 线程束矩阵函数
5.4.11.1. 说明
5.4.11.2. 替代浮点格式
5.4.11.3. 双精度
5.4.11.4. 子字节运算
5.4.11.5. 限制
5.4.11.6. 元素类型和矩阵大小
5.4.11.7. 示例
5.5. 浮点计算
5.5.1. 浮点数简介
5.5.1.1. 浮点格式
5.5.1.2. 正规值与次正规值
5.5.1.3. 特殊值
5.5.1.4. 结合律
5.5.1.5. 融合乘加（FMA）
5.5.1.6. 点积示例
5.5.1.7. 舍入
5.5.1.8. 关于主机/设备计算准确性的说明
5.5.2. 浮点数据类型
5.5.3. CUDA 与 IEEE-754 的符合性
5.5.4. CUDA 与 C/C++ 的符合性
5.5.5. 浮点功能的提供方式
5.5.6. 内置算术运算符
5.5.7. CUDA C++ 数学标准库函数
5.5.7.1. 基本运算
5.5.7.2. 指数函数
5.5.7.3. 幂函数
5.5.7.4. 三角函数
5.5.7.5. 双曲函数
5.5.7.6. 误差函数与伽马函数
5.5.7.7. 舍入至最近整数的浮点运算
5.5.7.8. 浮点数处理函数
5.5.7.9. 分类与比较
5.5.8. 非标准 CUDA 数学函数
5.5.9. 内建函数
5.5.9.1. 基本内建函数
5.5.9.2. 仅限单精度的内建函数
5.5.9.3. --use_fast_math 的影响
5.5.10. 参考文献
5.6. 设备可调用的 API 和内建函数
5.6.1. 内存屏障原语接口
5.6.1.1. 数据类型
5.6.1.2. 内存屏障原语 API
5.6.2. 流水线原语接口
5.6.2.1. memcpy_async 原语
5.6.2.2. 提交原语
5.6.2.3. 等待原语
5.6.2.4. 到达屏障原语
5.6.3. 协作组 API
5.6.3.1. cooperative_groups.h
5.6.3.1.1. thread_block 类
5.6.3.1.2. cluster_group 类
5.6.3.1.3. grid_group 类
5.6.3.1.4. thread_block_tile 类
5.6.3.1.5. coalesced_group 类
5.6.3.2. cooperative_groups/async.h
5.6.3.2.1. memcpy_async
5.6.3.2.2. wait 和 wait_prior
5.6.3.3. cooperative_groups/partition.h
5.6.3.3.1. tiled_partition
5.6.3.3.2. labeled_partition
5.6.3.3.3. binary_partition
5.6.3.4. cooperative_groups/reduce.h
5.6.3.4.1. Reduce 运算符
5.6.3.4.2. reduce
5.6.3.5. cooperative_groups/scan.h
5.6.3.5.1. inclusive_scan 和 exclusive_scan
5.6.3.6. cooperative_groups/sync.h
5.6.3.6.1. barrier_arrive 和 barrier_wait
5.6.3.6.2. sync
5.6.3.6.2.1. 网格同步
5.6.4. CUDA 设备运行时
5.6.4.1. 在 CUDA 代码中包含设备运行时 API
5.6.4.2. CUDA 设备运行时中的内存
5.6.4.2.1. 配置选项
5.6.4.2.2. 分配与生存期
5.6.4.2.2.1. 内存声明
5.6.4.2.2.1.1. 设备内存和常量内存
5.6.4.2.2.1.2. 纹理和表面
5.6.4.2.2.1.3. 共享内存变量声明
5.6.4.2.2.1.4. 常量内存
5.6.4.2.2.1.5. 符号地址
5.6.4.3. SM ID 和线程束 ID
5.6.4.4. 启动设置 API
5.6.4.5. 设备管理
5.6.4.6. API 参考
5.6.4.7. API 错误和启动失败
5.6.4.8. 设备运行时流
5.6.4.8.1. 隐式（NULL）流
5.6.4.8.2. 即发即弃流
5.6.4.8.3. 尾部启动流
5.6.4.9. ECC 错误
5.7. CUDA C++ 内存模型
5.7.1. 线程作用域
5.7.1.1. 作用域关系
5.7.2. 同步原语
5.7.3. 原子性
5.7.4. 数据竞争
5.7.5. 示例：消息传递
5.8. CUDA C++ 执行模型
5.8.1. 主机线程
5.8.2. 设备线程
5.8.3. CUDA API
5.8.3.1. 依赖项
6. 声明
6.1. 声明
6.2. OpenCL
6.3. 商标

原著 © 2007 至 2026 NVIDIA Corporation 及其关联公司。保留所有权利。
中文译本仅作学习参考，非 NVIDIA 官方出版物。技术含义有疑问时请核对NVIDIA 英文原文。
关于本译本

输入关键词搜索全书。