CUDA Techniques to Maximize Compute and Instruction Throughput [S72685]

Ben Pinzone, Compute Developer Technology Engineer
David Clark, Compute Developer Technology Engineer
GTC 2025, March 17th, 2025

目录

• 目标 (Goals
• 议程 (Agenda
• GPU 线程层级结构、SIMT、线程束分化 (Warp divergence
• 线程束调度器 (Warp scheduler) 和内核性能分析概览
• 延迟隐藏 (Latency hiding) 和提升指令吞吐量
• 减少指令数量并使吞吐量有效
• Tensor Core 总结
• 问答

目标 (Goals)

• 回顾 GPU 编程概念。

• 简要说明这些概念在性能分析工具（Nsight Compute）中的体现。

  • 您将看到许多 Nsight Compute 的截图 - 这是 Nvidia 的 CUDA 内核性能分析工具。

• 概述通用优化策略，并使用具体示例。

  • 这些是想法，不是硬性规定！您的实际效果可能会有所不同。
  • 非简单复杂度的代码更改应由性能分析器驱动。

• 有些要点会快速提及或仅引用其他资源，目的在于提高认知。

议程 (Agenda)

• GPU 线程层级结构、SIMT 和线程束分化 (Warp divergence)。
• 线程束调度器 (Warp scheduler) 和内核性能分析概览。
• 延迟隐藏 (Latency hiding) 和提升指令吞吐量。
• 减少指令数量并使吞吐量有效。
• Tensor Core 总结。

GPU 线程层级结构、SIMT、线程束分化 (Warp divergence)

---

GPU 概览

NVIDIA H100 SXM

该图展示了 NVIDIA H100 SXM GPU 的整体架构，主要组件包括：
- 132 个 SMs (Streaming Multiprocessors) 和 第四代 Tensor Cores。
- PCIe Gen 5：128 GB/s 双向带宽。
- 50 MB L2 缓存。
- 80 GB HBM3：3.3 TB/s 双向带宽。
- 第四代 NVLink：900 GB/s 双向带宽。

流式多处理器 (Streaming Multiprocessor - SM)

Hopper 架构

Hopper 架构的 SM 包含以下特性：
- SM 拥有 4 个子分区 (sub-partitions)。
- 128 个 FP32 单元。
- 64 个 FP64 单元。
- 64 个 INT32 单元。
- 4 个混合精度 Tensor Cores。
- 16 个特殊功能单元 (special function units, 用于超越函数)。
- 4 个线程束调度器 (warp schedulers)。
- 32 个加载/存储 (LD/ST) 单元。
- 64k 个 32-bit 寄存器。
- 256 KiB 统一 L1 数据缓存和共享内存。
- Tensor 内存加速器 (Tensor Memory Accelerator - TMA)。

线程层级结构：网格 (Grid) 与线程块 (Blocks)

• 一个 CUDA 内核以一个线程块网格 (grid of thread blocks) 的形式启动，这些线程块是完全独立的。
• 线程块在 SMs 上执行。
  • 多个线程块可以并发地驻留在同一个 SM 上。
  • 线程块不会迁移。
  • 每个线程块可以按任意顺序被调度到任何可用的 SMs 上，可以是并发或串行。

线程层级结构：集群 (Clusters)

• 从 Hopper GPU 开始，CUDA 在线程层级结构中引入了一个可选的层级，称为线程块集群 (Thread Block Clusters)。
• 一个集群内的线程块保证可以被并发调度，并支持跨多个 SMs 的线程之间进行高效的协作和数据共享。

线程层级结构：线程束 (Warps)

<blockquote>

"编织，第一种并行技术" - CUDA C++ 编程指南

</blockquote>

• 1 个线程束 (warp) = 32 个线程。
• 在运行时，一个线程块 (block) 被划分为多个线程束以进行 SIMT 执行。
• 一个线程束内的线程拥有连续的线程 ID。
• 一个线程块中的线程束总数定义为：
  • Ceil(每个块的线程数 / 线程束大小)

SIMT 架构

单指令，多线程 (Single-Instruction, Multiple-Thread)

在共享内存中进行工作排队 (Work Queueing in shared memory)

对于某些工作负载，其昂贵的计算由一个轻量级的检查来守护。一个简单的实现可能会因为并非所有线程都能通过该检查而遭受高度的分歧。

解决方案：

• 当一个线程找到一个可以深入处理的地方时，它会将其添加到队列中，然后继续前进。
• 偶尔，所有线程会同时工作以清空队列。
• 注意： 完成侦察（scouting）的线程会留下来帮助清空队列。

该过程分为两个阶段：
1. 侦察阶段 (Scouting phase): 线程块中的线程并行地寻找工作。当某些线程（图中以绿色'x'标记）找到有效工作时，它们会将工作项添加到共享内存中的队列。
2. 处理阶段 (Process phase): 当队列中累积的工作项达到一定数量（Invoke process size）时，整个线程块会切换到处理阶段，共同从队列中取出并执行这些工作项。这确保了在处理阶段所有线程都在执行相同的任务，从而避免了分歧。

线程束调度器 (Warp scheduler) 和内核性能分析概览

Warp调度器统计数据：性能分析的思维模型

调度器运行的可视化

以下几页通过一个简化的、逐周期的示例，构建了一个用于理解 Warp 调度器行为和相关性能指标的思维模型。

Warp 状态定义:
- Unused (未使用): 空闲的 Warp 插槽。
- Active (活动): Warp 已被分配到该插槽。
- Stalled (停滞): 活动的 Warp 正在等待某个依赖项（例如，内存读取、指令执行完成），当前周期无法被调度。
- Eligible (符合条件): 活动的 Warp 已准备好执行下一条指令。
- Selected (已选择): 在当前周期，从所有符合条件的 Warp 中被调度器选中并发射指令的 Warp。

周期 N:
在一个给定的周期（Cycle N）中，调度器从所有符合条件的 Warp（图中浅绿色块）中选择一个（图中带斜线的浅绿色块）来发射（Issue）其指令。

周期 N+1:
调度器选择下一个符合条件的 Warp（插槽4）。上一个周期被选择的 Warp（插槽3）现在变为停滞状态，因为它正在等待指令完成，因此不符合被调度的条件。（当然，如果它有其他独立的指令可以执行，情况可能会有所不同）。

周期 N+2:
位于插槽4的 Warp 执行完毕并退出。此时，没有任何符合条件的 Warp，因此调度器的发射单元（Issue Slot）在本周期处于空闲状态，未发射任何指令。

周期 N+3:
新的 Warp 被调度到之前空闲的插槽中。调度器从新的符合条件的 Warp 中选择一个（插槽2）并发射指令。

聚合性能指标

基于以上4个周期的示例，我们可以计算出一系列聚合的性能指标。

1. cycles_active (活动周期数): 4
- 观察到的

延迟绑定内核：调度器统计

每个调度器一个 warp 的情况下的度量（报告中的数据，已四舍五入）：

• warps_active (活跃 warps): 1.00
• warps_stalled (停滞 warps): 0.87
• warps_eligible (合格 warps): 0.13
• warps_selected (已选 warps): 0.13

上图显示了调度器发出指令的摘要。每个调度器维护一个可以发出指令的 warp 池。理论上，池的上限受启动配置的限制。在每个周期，调度器检查池中已分配 warp 的状态（Active Warps）。未停滞的 warp（Eligible Warps）准备好发出它们的下一条指令。调度器从合格 warp 集合中选择一个 warp 来发出一或多条指令（Issued Warps）。在没有合格 warp 的周期中，发出槽被跳过，没有指令发出。许多被跳过的发出槽表明存在延迟隐藏问题。

• 问题：发出槽利用率不足
  • 每个调度器每个周期能够发出一条

延迟隐藏 (Latency hiding) 和提升指令吞吐量

延迟绑定内核：改进方案

GPU通过并行处理大量任务（in-flight work）来掩盖延迟。增加Warp的数量可以隐藏延迟。

下图展示了通过增加Warp数量来填充指令流水线，避免了因单个Warp停滞（Stalled）而导致的“无Warp分发”（No warp issuing）的空闲周期。当Warp 0停滞时，调度器可以选择Warp 1或Warp 2来执行，从而保持硬件的繁忙状态。

利特尔定律（Little's Law）在此处的应用描述了我们需要多少在执行中的指令（instructions in flight）才能避免暴露延迟。

延迟绑定内核：通过增加指令级并行性改进

另一种隐藏延迟的方法是增加指令级并行性（Instruction Level Parallelism, ILP）。

如下面的代码和图表所示，通过使用float2代替float，在同一个循环迭代中执行两个独立的操作（result.x += 3.14f; 和 result.y += 3.14f;）。这使得即使在一个Warp内部，当一条指令等待时（例如，等待前一指令的结果），硬件也可以调度执行另一条独立的指令，从而减少停滞，提高执行效率。

关于停滞（Stalling）的底线

1. 如果SM或内存系统资源已经繁忙，则无需担心停滞或未使用的分发槽。

   • 更频繁地分发指令没有帮助，因为资源已经饱和。

2. 否则，你的程序就是延迟绑定的。需要为硬件提供更多的并发工作。可以尝试以下方法：

   • 更频繁地分发指令。
   • 减少停滞的频率。
   • 在停滞期间处理其他事务。
   • 减少停滞的持续时间（例如，使用延迟更低的指令）。

占用率（Occupancy）：是什么？如何确定？

• 占用率是在一个SM上可以并发活动的最大Warp数量。

  • 设备（Device）：GPU的计算能力所决定的理论最大值。
  • 可实现（Achievable）：取决于内核实现和编译器。
  • 已实现（Achieved）：主要取决于网格大小或工作负载行为，如工作负载不平衡。

  

• CUDA内核的可实现占用率将受到以下至少一个因素的限制：

  • SM资源分配：

    • 例如：共享内存和寄存器必须在线程之间进行分区。

  • 块大小。

  • 其他硬件因素：
    • 例如：每个SM的最大块数、每个SM的最大Warp数等。
    • 具体信息在https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#technical-specifications-per-compute-capability中有概述。

<blockquote>

可以使用NVIDIA Nsight Compute分析CUDA内核的占用率。

</blockquote>

占用率详解

占用率可以用一个“发牌”的比喻来理解：线程块（Blocks）就像牌，它们有特定的资源需求（如共享内存、寄存器）。发牌官（SM）在有足够资源服务它们的情况下，将线程块并发地分发出去。

• 上图案例：给定块大小为128个线程（即4个Warp），所需的共享内存、寄存器和各种硬件限制。

  • 第一个耗尽的资源是共享内存。
  • 每个SM实现了8个Warp的占用率。

• 下一步，我们可以通过两种方式提高占用率：

  • 重新编程内核以使用更少的共享内存。
  • 将块大小减小到3个Warp，即96个线程。（注意这会浪费24KB的共享内存）。
    • 通过分配更小/更精细的块，每个（新）块需要70 * (3/4) = 52.5 KB。因此，我们可以分配164 / 52.5 = 3个块。
    • 现在每个SM实现了9个Warp的占用率。这是一个小的努力带来的改进。

占用率限制器：寄存器

• 寄存器使用情况：使用--ptxas-options=-v编译可以报告每个线程的寄存器数量。

• 最大寄存器数可以手动设置：

  • 在编译时，对每个文件使用nvcc的-maxrregcount标志。
  • 对每个内核，使用__launch_bounds__或__maxrreg__限定符。

• Hopper架构每个SM拥有64K（65536）个寄存器。

  • 寄存器以256个为单位的固定大小块进行分配。

• 示例：

  • 内核每个线程使用63个寄存器。
  • 每个Warp使用的寄存器数 = 63 * 32 = 2016个。
  • 每个Warp分配的寄存器数 = 2048个（向上取整到分配块大小）。
  • 每个SM可实现的活动Warp数 = 65536 / 2048 = 32个。
  • 占用率 = 32 / 64 * 100 = 50%。
  • Hopper架构每个SM最多支持64个Warp。

超出寄存器限制

• 如果编译器需要的寄存器数量超过了设备或用户指定的限制，它会将多余的寄存器溢出（spill）到本地内存（local memory）。
  • 本地内存是位于设备内存中的线程私有存储空间，并在L1和L2缓存中进行缓存。
  • 编译器也可能出于其他原因使用本地内存。

<blockquote>

优化技巧
有限的本地内存使用可能对性能有益。

</blockquote>

• 如果数据保留在L1缓存中，访问速度会很快，并且在大多数情况下可以提高占用率；然而，这需要进一步的调查。

下图显示了大量的本地内存请求（3.70M Req），这些请求被发送到L1/TEX缓存。

如何在NCU中定位寄存器压力

在NVIDIA Nsight Compute (NCU) 中，可以通过查看源代码视图和SASS（汇编代码）视图中的“Live Registers”（活跃寄存器）列来识别寄存器压力大的

减少指令数量并使吞吐量有效

数学运算的成本层级

使用尽可能轻量级的工具。数学运算的成本从高到低排列如下：

• 红色 (最高成本): 三角函数 (Trig)、平方根 (sqrt) 等。

• 黄色: 除法 (Division)、取模 (mod) 运算符。

  • 硬件中没有除法或取模指令。根据除数的属性，可能会转化为数百条指令。
  • 通过 constexpr 进行的除法/取模通常没有问题。

• 浅绿色: 乘法 (Multiply)、加法 (add)、减法 (subtract)。

• 深绿色 (最低成本): 融合乘加 (Fused multiply adds)。

仅使用你的应用所需的精度

• 研究较低的精度是否足以满足您的用例。
• 注意隐式类型转换为 double！在您的数值字面量上使用 .f 后缀以避免这种情况。

• 利用快速数学优化 (--use_fast_math)。

  • 可以通过 Single Precision Intrinsics 实现增量采用，例如 __cosf(), __expf(), __frcp_{0}, __fsqrt_ 等。
  • 单精度三角 API 函数可以使用一些双精度指令和局部内存。

• 下图展示了默认数学 API 路径与 __expf 或快速数学路径的对比，以及 CUDA 编程指南中关于内在函数误差范围的信息。

代数优化

静态考量

• 观察代码循环内部：能否将乘以常量的操作移出循环？等等。
• 将除数移到比较运算符的另一侧。
• 如果内核中存在按运行时常量进行的除法：请在主机端计算倒数，然后传递到内核。在内核中通过乘以倒数来实现。
• 对于在编译时已知或值范围有限的任何变量，使用模板参数。运行时编译可以进一步推动这一点。

运行时考量

• 你是否比编译器更了解情况？关于：
  • 表达式可能产生的值的范围？
  • 表达式在后续表达式中是如何概念性地使用的？
  • 关键阈值是什么？
  • 条件结果何时会改变？
  • 你如何利用这些信息？

小改动可能产生大影响

无符号整数溢出是已定义行为，编译时需要考虑这一点，可能导致额外的指令。有符号整数溢出是未定义行为，这为编译器生成更快的代码提供了更大的灵活性。

优化技巧：使用有符号整数而不是无符号整数作为循环计数器。

下面的代码和图表展示了使用 unsigned int 和 int 作为循环计数器的性能差异。

示例：线段相交

给定一个线段数组，计算有多少交点发生。

• 完整源代码可在 Accelerated Computing Hub 上获取。

计算受限于除法慢速路径

与C语言除法运算符相关的SASS（汇编）调用了一个注入的辅助SASS（函数）。

• SASS中的FCHK指令之后，CALL.REL.NOINC指令会跳转到慢速路径，而BRANCH指令则会跳过CALL指令（快速路径）。
• 棘手细节：如何知道程序执行快速路径或慢速路径的频率？
  • 这取决于运行时的浮点数范围。
  • 需要查看FCHK指令的谓词依赖关系。

线段相交示例：可能的优化

1. 修改源码以避免除法。
2. 改变算法以避免除法。
3. 执行包围盒预检查。
4. 在不同的浮点数范围上操作以避免慢速路径。

优化选项 1: 修改源码以避免除法

观察除法结果的使用方式。

• 我们只关心结果是否在 [0, 1] 范围内。
• 因此，我们无需执行除法，而是比较分子和分母的符号和大小。

优化选项 2: 采用不同算法以避免除法

新算法：线段分割二维平面。

• 该方法通过计算叉积来判断线段端点相对于另一条线段的位置，从而确定是否存在交点。
• 如下图所示，通过判断点相对于另一条线段的位置关系来判断相交情况。

优化选项 3: 添加包围盒预检查

这是一种轻量级的预检查：在运行任何相交测试之前，检查两个线段的包围盒是否重叠。

• 这会增加指令，但可能会引起线程束发散 (divergence)。是否值得这样做需要权衡。

优化选项 4: 设法改变你的输入数据

• 设法将输入的浮点数变得更小。

  • 这样，快速除法路径将被采用，因为FCHK指令不会标记溢出风险。

• 这种方法依赖于具体应用，可能无法实现。

线段相交结果 (A100)

下表展示了在不同浮点数量级、不同实现和优化策略下的运行时间。

• 对于大数量级浮点数 (1e19): 基线版本（DIV）走的是慢速路径。修改源码（DIV_FREE_ND）和新算法（DIV_FREE_CLRS）能带来显著提升 (2.54x - 2.67x)。包围盒检查在这种数据上没有帮助，反而更糟。
• 对于中等数量级浮点数 (1e18): 基线版本走的是快速路径。修改源码和新算法也能带来少量提升 (1.05x - 1.10x)。
• 对于小数据且相交频繁的情况: 包围盒检查（BBox）开始显示出优势 (1.43x - 1.48x)。

优化多项式求值

本节讨论在编译时已知常数的情况下，重复计算高阶多项式的优化方法。

替换幂函数调用

通过将 pow 函数调用替换为运行中的指数计算来优化。初始基线使用 pow 函数，运行时间为 40,862 us。通过手动展开计算，避免调用通用函数，性能得到显著提升。

• 基线 (Baseline): 40,862 us, 1x 加速比
• 使用 FP32 字面量 (FP32 literals): 6,198 us, 6.6x 加速比
• 避免通用函数 (Avoiding general purpose functions): 454.8 us, 90x 加速比

优化技巧：尽可能优先使用更快、更专业的数学函数，而不是更慢、更通用的函数。

---

采用霍纳（Horner）方法

使用霍纳方法可以减少指令数量，进一步优化多项式求值。该方法通过重构表达式来减少乘法操作。

• 霍纳方法 (Horner's method): 319.1 us, 128x 加速比

优化技巧：检查表达式是否可以被重构以产生更少的指令。

---

使用融合乘加（Fused Multiply-Add）指令

通过使用融合乘加（FMA）指令，可以将乘法和加法操作合并为一条指令，从而提高性能。这可以通过编译器选项 fmad=true 或使用内部函数（intrinsics）如 __fmaf_rn 来实现。

• 融合乘加 (Fused multiple adds): 170.6 us, 240x 加速比

优化技巧：在可能的情况下使用融合乘加指令。

---

增加指令级并行度（ILP）

霍纳方法会产生很少的指令级并行性（ILP），因为它创建了一个长的依赖指令链。当需要 ILP 时，可以采用 Estrin 方案。Estrin 方案可以用于添加指令序列化，而开销很小，因为它将多项式分解为两个子多项式，然后用霍纳方法分别求值。

下表比较了在不同并行度（通过每个 SM 的线程块数来体现）下霍nor方法和Estrin方案的性能。
- 在高并行度（8个线程块/SM）下，两种方案性能相近。
- 在低并行度（1个线程块/SM）下，Estrin方案（192 us）由于其更高的ILP，性能优于霍纳方法（326 us）。

优化技巧：如果需要指令级并行度（ILP），请打破依赖指令链。

---

Tensor Core 总结

矩阵乘法：深度学习中的关键操作

矩阵乘法是深度学习中的核心运算。在全连接层中，输入（inputs）乘以权重（weights）得到输出（outputs），这本质上是一个矩阵乘法操作。批处理维度（batch dimension）使输入成为一个二维矩阵。

Tensor Core 是专门用于矩阵乘法运算的硬件流水线。

---

示例：HMMA 1688

HMMA 1688 是一个 FP16 16x8x8 的 Tensor Core 操作示例。

• 单条 HMMA 1688 指令可计算 128 个结果（相当于 1024 次 FMA 浮点运算）。
• 运算公式为 D = A x B + C。

• 矩阵维度：

  • A: 16 x 8 FP16 输入
  • B: 8 x 8 FP16 输入
  • C/D: 16 x 8 FP16 或 FP32 输出

• 整个 Warp（32个线程）参与计算。

• 计算值存储在寄存器中。

---

Tensor Core 的历史与特性

Tensor Core 随着 NVIDIA GPU 架构的演进而不断发展，在数据中心和消费级 GPU 上均有部署。

• Volta (sm_70): 仅用于数据中心 GPU。支持 FP16，8线程宽度，每个 SM 每时钟周期可执行 1024 次 FP16 Flops。
• Turing (sm_75): 扩展到游戏和通用数据中心 GPU。支持 FP16 和 INT8/4/1，采用 warp-wide 操作，每个 SM 每时钟周期 1024 FP16 Flops。

• Ampere (sm_80, sm_86/87):

  • GA100 (sm_80): 数据中心 GPU。增加了 BF16, TF32, FP64 支持，引入了稀疏计算（Sparsity），每个 SM 每时钟周期 2048 FP16 Flops。
  • GA10x (sm_86/87): 通用 GPU。支持与 GA100 类似的数据类型和稀疏性，每个 SM 每时钟周期 1024 FP16 Flops。

• Hopper (sm_90a): 数据中心 GPU。引入 FP8 支持和共享内存（Shmem）输入，采用 4-warp-wide 操作，性能提升至每个 SM 每时钟周期 4096 FP16 Flops。

• Ada (sm_89): 通用 GPU。引入 FP8 支持，采用 warp-wide 操作，每个 SM 每时钟周期 1024 FP16 Flops。
• Blackwell (sm_100a, sm_120):
  • GB100 (sm_100a): 数据中心 GPU。引入 FP4/FP6 支持和块级缩放（Block-scaling FPx），采用 4 warps x 2 CTA 异步 TMEM + Shmem，性能提升至每个 SM 每时钟周期 8192 FP16 Flops。
  • GB20x (sm_120): 通用 GPU。引入 FP4/FP6 支持，每个 SM 每时钟周期 1024 FP16 Flops。

注：峰值 FP16 Flops 数据均基于非稀疏计算。

---

GPU 特定指令

Tensor Core 指令通常是不可移植的，除非它们基于 Ampere SuperMMA API。为了更好的硬件抽象，推荐使用更高层次的 API，如 CUTLASS。

指令集演进：

• Ampere: mma PTX
• Hopper: PTX wgmma
• Blackwell: PTX TCgen5
• 从 Ampere 到 Ada 和 Blackwell 的通用 GPU 分支，mma PTX 指令扩展以支持新的数据类型。

---

Tensor Core 的使用者（Providers）

开发者可以在不同抽象层次上使用 Tensor Core，从低级 API 到高级 API。

• 低级 API (Low level APIs)

  • 优点: 易于获得硬件的理论峰值性能（SOL performance）。
  • 缺点: 编程难度大，代码与特定 GPU 强相关。
  • 示例: PTX 指令，如 wmma (warp-level), mma (warp-wide), wgmma (warp-group, Hopper), tcgen05.mma (Blackwell)。

• 高级 API (High level APIs)

  • 优点: 易于编程，提供更好的硬件抽象。
  • 缺点: 较难达到理论峰值性能。
  • 示例:
    • 高级框架: PyTorch, TensorFlow, OAI Triton 等。
    • 库: CUBLAS, CUDNN, TRT。
    • 抽象库: CUTLASS，它提供了从高级核函数到底层 PTX 包装器的桥梁，是 NVIDIA 官方推荐的库。

Tensor Core APIs

编写您自己的 Tensor Core 内核

• 在大多数情况下，CUTLASS 应该可以满足需求，并提供硬件抽象和出色的性能。

  • https://github.com/NVIDIA/cutlass

• 查看相关的 CUTLASS 演讲：

  • 使用 CUTLASS 编程 Blackwell Tensor Cores [S72720]
  • 使用 CUTLASS 4.0 在 Python 中启用 Tensor Core 编程 [S74639]

• 如果您想使用 PTX，请查阅 PTX 文档或 CUTLASS 代码。

  • 适用于 Turing、Ampere 及所有“较小”GPU 的 PTX mma 指令
  • Hopper 上的 wgmma 指令
  • GB100 Blackwell 上的 tcgen05.mma 指令
  • Blackwell TensorCore 和内存子系统应用优化 [CWE72552]

问答

+ GTC 2025 更多 CUDA 开发者会议

• 通用 CUDA

  • S72571 - CUDA 究竟是什么？
  • S72897 - 如何编写 CUDA 程序：并行编程版

• CUDA Python

  • S72450 - 加速 Python：社区与生态系统之旅
  • S72448 - CUDA Python 开发者工具箱
  • S72449 - 编写 CUDA 内核的 1001 种方法
  • S74639 - 使用 CUTLASS 4.0 在 Python 中启用 Tensor Core 编程

• CUDA C++

  • S72574 - 飞速构建 CUDA 软件
  • S72572 - CUDA C++ 开发者工具箱
  • S72575 - 您应该如何编写 CUDA C++ 内核

• 开发者工具

  • S72527 - 比您想象的更简单——使用智能开发者工具调试和优化 CUDA

• 与专家交流

  • CWE72433 - CUDA 开发者最佳实践
  • CWE73310 - 使用 NVIDIA CUDA 编译器工具链进行高效的 GPGPU 编程
  • CWE72393 - 您的开发者工具箱里有什么？CUDA 与图形分析、优化和调试工具
  • CWE75584 - 与《Programming Massively Parallel Processors》作者 Wen-mei Hwu 博士交流

• 多 GPU 编程

  • S72576 - 多 GPU 扩展入门：分布式库
  • S72570 - 深入多 GPU 扩展：基于任务的运行时
  • S72578 - 高级多 GPU 扩展：通信库

• 性能优化

  • S72681 - 最大化内存带宽和隐藏延迟的 CUDA 技术
  • S72685 - 最大化计算和指令吞吐量的 CUDA 技术
  • S72686 - 最大化并发和系统利用率的 CUDA 技术
  • S72687 - 从 Grace Hopper 获得最佳性能

更多信息请访问：http://nvidia.com/gtc/sessions/cuda-developer