第8章 占用率调优、线程束效率与指令级并行 (Occupancy Tuning, Warp Efficiency, and Instruction-Level Parallelism)¶
现代GPU加速工作负载正在将硬件推向极限。像Blackwell这样的多芯片GPU通过10 TB/s的NV-HBI链路连接多个光刻限制的芯片，并将L2缓存增加到126 MB。这些硬件设计选择显著改变了内存与计算的权衡以及占用率的最佳点。这使得分析和优化比以往任何时候都更加关键。在内存优化基础知识的基础上，我们现在转向高级延迟隐藏技术和吞吐量增强，旨在充分利用现代GPU的全部能力。
Modern GPU-accelerated workloads are pushing hardware to its limits. Multi-die GPUs like Blackwell connect multiple reticle-limited dies with a 10 TB/s NV-HBI link and increase L2 to 126 MB. These hardware design choices materially change memory-vs-compute tradeoffs and occupancy sweet spots. This makes profiling and optimization even more critical than ever. Building on the fundamentals of memory optimizations, we now turn to advanced latency-hiding techniques and throughput enhancements designed to fully leverage the full power of modern GPUs.
我们将专注于识别性能瓶颈，然后应用一套系统的优化策略逐一消除它们。本章的关键主题包括调优占用率、优化线程束效率以及增加指令级并行性。
We will focus on identifying performance bottlenecks and then applying a systematic set of optimization strategies to eliminate them one by one. Key themes in this chapter include tuning occupancy, optimizing warp efficiency, and increasing instruction-level parallelism.
在本章结束时，您将能够识别GPU利用率不足的根本原因——以及应用正确的优化组合。我们还将为您准备更高级的技术，如内核融合以及使用CUDA Graphs和CUDA流等原语的流水线操作，这些将在后续章节中介绍。
By the end of the chapter, you will be able to identify root causes of GPU underutilization-as well as apply the right combination of optimizations. We will also prepare you for more advanced techniques like kernel fusion and pipelining with primitives like CUDA Graphs and CUDA streams, which we cover in subsequent chapters.
虽然我们专注于CUDA C++等高级语言和PyTorch等AI框架，但分析和调优的原则适用于从软件栈一直到硬件的所有层级。因此，理解底层硬件性能对于诊断高级抽象难以完全解决的瓶颈仍然至关重要。
While our focus is higher-level languages like CUDA C++ and AI frameworks like PyTorch, the principles of profiling and tuning apply at all levels of the stack right down to the hardware. As such, understanding low-level hardware performance remains critical for diagnosing bottlenecks that higher-level abstractions make difficult to fully resolve.
分析和诊断GPU瓶颈 (Profiling and Diagnosing GPU Bottlenecks)¶
在优化之前，我们必须首先识别代码中的瓶颈，以确定哪个硬件或软件资源限制了我们的性能。现代NVIDIA GPU非常复杂，性能下降可能来自许多来源，包括内存带宽、内存延迟、指令吞吐量、同步开销、并行性不足、主机-设备传输延迟等。
Before optimizing, we must first identify the bottlenecks in our code to determine which hardware or software resource is limiting our performance. Modern NVIDIA GPUs are complex, and a slowdown could come from many sources, including memory bandwidth, memory latency, instruction throughput, synchronization overheads, insufficient parallelism, host-device transfer delays, and more.
NVIDIA的分析生态系统包括Nsight Systems（命令行界面nsys）和Nsight Compute（命令行界面ncu）。Nsight Systems捕获CPU线程、GPU内核和内存传输的系统级时间线。它还可以捕获Python回溯和Python采样。
NVIDIA's profiling ecosystem includes Nsight Systems (command-line interface nsys) and Nsight Compute (command-line interface ncu). Nsight Systems captures a system-level timeline of CPU threads, GPU kernels, and memory transfers. It can also capture Python backtraces and Python sampling.
结合PyTorch分析器和各种可视化工具，Nsight Systems和Nsight Compute可以帮助您诊断内核性能瓶颈、分析屋顶线图，并测量迭代优化工作的有效性。
Combined with PyTorch profiler and various visualization tools, Nsight Systems and Nsight Compute can help you diagnose kernel performance bottlenecks, analyze roofline plots, and measure the effectiveness of your iterative optimization efforts.
Nsight Systems时间线视图 (Nsight Systems Timeline View)¶
Nsight Systems时间线视图帮助精确定位并发问题、传输开销和空闲时段。例如，运行以下代码生成详细的时间线，显示内核启动重叠、CPU准备间隙、数据传输时序和NVTX标记范围：
The Nsight Systems timeline view helps pinpoint concurrency issues, transfer overhead, and idle periods. For example, run the following code to produce a detailed timeline showing kernel launch overlaps, CPU preparation gaps, data transfer timings, and NVTX-marked ranges: nsys profile --trace=... --capture-range=... --force-overwrite=true
此外，Nsight Systems GUI让您交互式地检查时间线。它可视化CPU线程、GPU内核，甚至用户定义的NVTX范围，具有缩放和平移功能用于详细分析（见图8-1）。
In addition, the Nsight Systems GUI lets you interactively inspect the timeline. It visualizes CPU threads, GPU kernels, and even user-defined NVTX ranges with zoom and pan features for detailed analysis (see Figure 8-1).
请记住，NVTX注释对于复杂应用程序至关重要。在代码中使用NVTX范围标记感兴趣的区域。然后您可以使用Nsight Systems捕获范围配置文件。虽然CUDA分析器Start和Stop API支持捕获控制，但NVTX范围是框架工作流的推荐机制。例如，您可以使用PyTorch中的torch.cuda.nvtx提供的NVTX range push/pop来标记"前向传播"和"反向传播"等阶段。这使Nsight Systems时间线更加可解释，因为分析器可以捕获性能关键的迭代并清楚地划分关键计算段。
Remember that NVTX annotations are essential for complex applications. Use NVTX ranges in your code to mark regions of interest. You can then use Nsight Systems to capture range profiles. And while the CUDA profiler Start and Stop API is supported for capture control, NVTX ranges are the recommended mechanism for framework workflows. For instance, you can use NVTX range push/pop, available using torch.cuda.nvtx in PyTorch, to label phases such as "forward pass" and "backprop." This makes the Nsight Systems timeline far more interpretable since the profiler can capture the performance-critical iterations and clearly delineate key computation segments.
Figure 8-1. Nsight Systems interactive UI (source: https://oreil.ly/YEiWS)
分析和调优数据流水线 (Profiling and Tuning the Data Pipeline)¶
如前所述，Nsight Systems提供整个流水线的高级视图，包括数据加载和处理步骤。例如，如果您看到GPU空闲而CPU忙于准备数据或运行数据增强（在训练循环中很常见），瓶颈可能不是GPU内核，而是数据流水线本身。
As mentioned earlier, Nsight Systems provides a high-level view of the entire pipeline, including the data loading and processing steps. For instance, if you see the GPU idle while the CPU is busy preparing data or running augmentation, common in training loops, the bottleneck might not be the GPU kernel but the data pipeline itself.
对于瓶颈是真实代表性数据集上的数据流水线的场景，您可以调优数据加载器线程数、使用双缓冲将CPU预处理与GPU计算重叠，或将更多预处理移至GPU。
For scenarios in which the bottleneck is the data pipeline on a realistic and representative dataset, you can tune the number of data loader threads, overlap CPU preprocessing with GPU compute using double buffering, or move more preprocessing onto the GPU.
始终确保您感知的"GPU性能问题"实际上不是由内核执行上游或下游的内容（如数据加载）引起的。
Always ensure that what you perceive as a "GPU performance issue" is not actually caused by something upstream or downstream of the kernel execution like data loading.
有了来自分析的清晰图景，我们现在可以着手解决特定的瓶颈。接下来的章节将详细介绍核心优化技术。我们将从占用率调优开始，因为充足的线程束供应是隐藏延迟和最大化吞吐量的基础。
With a clear picture from profiling, we can now proceed to address specific bottlenecks. The next sections cover the core optimization techniques in detail. We'll start with occupancy tuning since a sufficient supply of warps is fundamental to hiding latency and maximizing throughput.
Nsight Compute和屋顶线分析 (Nsight Compute and Roofline Analysis)¶
另一方面，Nsight Compute（ncu）是一个为单个内核收集深度指标的分析器。例如，它跟踪实现的占用率、每周期发出的线程束指令、内存吞吐量（GB/s）、执行单元利用率等。这些指标共同描绘了系统性能配置文件的完整图景。
Nsight Compute (ncu), on the other hand, is a profiler that collects in-depth metrics for individual kernels. For instance, it tracks achieved occupancy, issued warp instructions per cycle, memory throughput (GB/s), utilization of execution units, and many others. Together, these paint a complete picture of your system's performance profile.
这些指标被组织成内存、计算和吞吐量等部分。Nsight Compute的自动分析规则会标记低内存利用率和分支分歧等低效问题--甚至提供优化提示。这些内置规则会针对新的GPU架构持续更新。它们可以根据FLOPS每字节比率、停顿原因等指标快速突出您是内存受限、延迟受限等。
These metrics are organized into sections like memory, compute, and throughput. Nsight Compute's automated analysis rules flag inefficiencies like low memory utilization and branch divergence-and even provide optimization hints. These built-in rules are updated continuously for new GPU architectures. They can quickly highlight if you are memory bound, latency bound, etc., based on metrics like FLOPS per byte ratios, stall reasons, etc.
Nsight Compute还包括屋顶线分析部分以及自动指导。这可以绘制内核实现的FLOPS相对于硬件屋顶线的图表--甚至突出您是否接近内存带宽或计算限制。
Nsight Compute includes a Roofline analysis section as well as automated guidance. This can plot your kernel's achieved FLOPS against hardware rooflines-and even highlight if you are near the memory bandwidth or compute limits.
请记住，内存与计算的区别使用屋顶线模型进行量化，该模型将内核性能相对于内存带宽和计算吞吐量的硬件上限进行绘制。Nsight Compute现在直接提供屋顶线分析，显示每个内核相对于峰值的实现GFLOPS及其算术强度（每字节FLOPS）。落在内存屋顶线以下的内核表示内存受限行为，而接近计算屋顶线的内核则是ALU受限的。使用Nsight Compute中的屋顶线部分直接获取算术强度（FLOPs/byte）和各精度的FLOPs。将其与硬件的理论峰值FLOPS每字节进行比较，您可以看到内核运行在屋顶线下方多远。示例屋顶线图如图8-2所示。
Remember that the memory-versus-compute distinction is quantified using the Roofline model, which plots kernel performance against hardware ceilings for memory bandwidth and compute throughput. Nsight Compute now directly provides Roofline analysis, showing each kernel's achieved GFLOPS relative to peak and its arithmetic intensity (FLOPS per byte). A kernel falling below the memory roof indicates memory-bound behavior, while one near the compute roof is ALU bound. Use the Roofline section in Nsight Compute to obtain arithmetic intensity (FLOPs/byte) and FLOPs by precision directly. Compare this to the hardware's theoretical peak FLOPS per byte, you can see how far below the roofline the kernel is operating. An example Roofline chart is shown in Figure 8-2.
Figure 8-2. Roofline chart shown in the Nsight Compute UI (https://oreil.ly/wUbIz)
通常有效的做法是首先使用Nsight Systems在时间线上找到热点内核或瓶颈操作。然后，您可以使用Nsight Compute深入分析每个内核，进行更细粒度的分析和诊断。这种两步工作流程，从系统级视图到内核级深入分析，是处理复杂GPU性能调试和调优的常见方法。
It's often effective to first use Nsight Systems to find hot kernels or bottleneck operations on the timeline. Then, you can zoom into each of these kernels with Nsight Compute to perform a more fine-grained analysis and diagnosis. This two-step workflow, moving from a system-wide view to a kernel-level deep dive is a common approach to handling complex GPU performance debugging and tuning.
PyTorch分析器和可视化工具 (PyTorch Profiler and Visualization Tools)¶
使用PyTorch等高级框架时，torch.profiler API可以在模型训练/推理期间收集类似的性能指标。PyTorch分析器使用Kineto库在底层实际执行数据收集。
When using high-level frameworks like PyTorch, the torch.profiler API can collect similar performance metrics during model training/inference. The PyTorch profiler uses the Kineto library to actually perform the data collection under the hood.
Kineto与CUDA的性能工具接口（CUPTI）后端集成，以捕获操作级执行时间、GPU内核启动、内存复制和硬件计数器指标。它还与Linux perf集成以记录CPU事件。Kineto将所有这些信息合并成一个连贯的、按时间排序的跟踪，可以使用PyTorch分析器UI、Nsight Systems GUI或仅使用Chrome浏览器（例如Chrome跟踪格式）进行可视化。示例Chrome跟踪可视化如图8-3所示。
Kineto integrates with CUDA's Performance Tools Interface (CUPTI) backend to capture operator-wise execution times, GPU kernel launches, memory copies, and hardware counter metrics. It also integrates with Linux perf to record CPU events. Kineto merges all of this information into a coherent, time-ordered trace, which can be visualized using the PyTorch profiler UI, Nsight Systems GUI, or just a Chrome browser (e.g., Chrome tracing format). An example Chrome tracing visualization is shown in Figure 8-3.
Figure 8-3. Chrome tracing visualization generated by the PyTorch profiler
PyTorch分析器允许您直接识别模型代码中的瓶颈操作。例如，您可以使用torch.profiler.profile(..., with_flops=True, profile_memory=True)来分析训练循环，以记录内存使用情况并估计支持的操作符（如矩阵乘法）的FLOPs。（注意：这些是操作级别的基于公式的估计，而不是每内核硬件计数器。）这种集成使弥合PyTorch模型代码与底层CUDA性能分析之间的差距变得更加容易。
PyTorch profiler allows you to identify bottleneck operations in your model code directly. For example, you can profile a training loop with torch.profiler.profile(..., with_flops=True, profile_memory=True) to record memory usage and to estimate FLOPs for supported operators such as matrix multiplication. (Note: These are formula-based estimates at the operator level rather than per-kernel hardware counters.) Such integration makes it easier to bridge the gap between PyTorch model code and low-level CUDA performance analysis.
现代版本的Nsight Systems可以收集Python回溯采样，还提供PyTorch专注模式，支持Python调用栈采样和PyTorch域以实现更相关的跟踪。这有助于将框架活动与系统时间线关联起来。Nsight Compute关联到CUDA C或C++源代码以及PTX或SASS。您可以使用-lineinfo编译设备代码以启用源代码行映射。要将模型代码与内核关联，请使用Python中的torch.cuda.nvtx提供的NVTX范围。此外，现代版本的Nsight Compute提供源代码视图，包括指令混合以及吞吐量细分，有助于精确定位受停顿和吞吐量限制影响的源代码行。
Modern versions of Nsight Systems can collect Python backtrace sampling and also provide a PyTorch-focused mode, which supports Python call-stack sampling and a PyTorch domain for more correlated tracing." This helps correlate framework activity with the system timeline. Nsight Compute correlates to CUDA C or C++ source and PTX or SASS. You can compile device code with -lineinfo to enable source line mapping. To correlate model code with kernels, use NVTX ranges from Python using torch.cuda.nvtx. In addition, modern versions of Nsight Compute provide a source view, which includes instruction mix-as well as a throughput breakdown that helps pinpoint the source code lines impacted by stalls and throughput limitations.
传统上，缺乏捕获Python信息使PyTorch开发者不愿意使用Nsight Systems/Compute。然而，如果您正在使用PyTorch/Python，值得重新审视这些工具，以提供与系统其余部分更全面和相关的分析。
The lack of capturing Python information traditionally discouraged PyTorch developers from using Nsight Systems/Compute. However, if you are using PyTorch/Python, it's worth revisiting these tools to provide a more holistic and correlated analysis with the rest of the system.
本节概述了诊断GPU瓶颈的工作流程，包括如何解释关键分析器指标以及它们对瓶颈类型的暗示。我们将专注于Nsight Compute进行深度内核分析，包括线程束停顿和内存吞吐量--以及Nsight Systems进行高级应用程序分析，包括并发和CPU-GPU重叠。
This section outlines a workflow for diagnosing GPU bottlenecks, including how to interpret key profiler metrics and what they imply about the bottleneck type. We will focus on Nsight Compute for deep-dive kernel analysis, including warp stalls and memory throughput-and Nsight Systems for high-level application profiling, including concurrency and CPU-GPU overlap.
分析器引导的分析 (Profiler-Guided Analysis)¶
关键是利用所有这些洞察来指导您的下一步，因为不同的瓶颈需要不同的优化。例如，您可以利用Nsight Compute的引导分析规则，它可能会标记"Memory Bound: L2 transactions per FLOP high"或"Compute Bound: issue stall reasons indicate pipe contention."这些规则会针对每个架构更新，包括Blackwell。它们可以快速确认瓶颈是在内存吞吐量、指令调度、延迟隐藏等方面。这可以将您引导到最相关的指标并改善解决瓶颈的时间。
The key is to use all of these insights to guide your next steps since different bottlenecks require different optimizations. For instance, you can take advantage of Nsight Compute's guided analysis rules, which might flag "Memory Bound: L2 transactions per FLOP high" or "Compute Bound: issue stall reasons indicate pipe contention." These rules are updated for each architecture, including Blackwell. They can quickly confirm whether a bottleneck is in memory throughput, instruction dispatch, latency hiding, etc. This can direct you to the most relevant metrics and improve your bottleneck time-to-resolution.
Nsight Systems和Nsight Compute不断更新到最新的GPU架构，以帮助性能工程师在将工作负载迁移到更新的GPU代际时诊断和修复性能问题。
Nsight Systems and Nsight Compute are constantly updated to the latest GPU architectures to assist performance engineers in diagnosing and fixing performance issues when migrating workloads to newer GPU generations.
使用Nsight Compute分析线程束停顿原因 (Analyzing Warp Stall Reasons with Nsight Compute)¶
如第6章所述，NVIDIA GPU以SIMT方式执行线程束或32个线程组。如果线程束频繁暂停而不是发出指令，分析器会对原因进行分类。
As discussed in Chapter 6, NVIDIA GPUs execute warps, or groups of 32 threads, in SIMT fashion. If warps are frequently pausing instead of issuing instructions, the profiler categorizes the reasons why.
Nsight Compute中最有洞察力的视图之一是内核的线程束状态统计细分。这有时称为线程束停顿原因。通过检查这些停顿原因，您通常可以精确定位限制因素。
One of the most insightful views in Nsight Compute is the Warp State Statistics breakdown for a kernel. This is sometimes called the warp stall reasons. By examining these stall reasons, you can often pinpoint the limiting factor.
有几种不同类型的线程束停顿：内存相关、执行依赖、执行争用以及其他如纹理缓存相关停顿。让我们逐一看看这些。
There are a few different types of warp stalls: memory-related, execution dependency, execution contention, and others like texture-cache-related stalls. Let's take a look at each of these.
内存相关停顿 (Memory-Related Stalls)¶
如果内核在等待全局内存加载，Nsight Compute会报告高百分比的"Stall: Long Scoreboard"周期。记分牌跟踪每个线程束的未完成内存请求，因此Long Scoreboard表示线程束频繁等待全局DRAM加载的高延迟，如图8-4所示。
If a kernel is waiting on global memory loads, Nsight Compute reports a high percentage of "Stall: Long Scoreboard" cycles. The scoreboard tracks each warp's outstanding memory requests, so Long Scoreboard indicates warps frequently waiting on the high latency of global DRAM loads, as shown in Figure 8-4.
Figure 8-4. Long Scoreboard stall caused by waiting on high-latency global memory accesses (e.g., waiting for data fetched from device memory into registers, or for spilled local memory to be written to/from device memory)
类似地，Short Scoreboard停顿是由等待共享内存和寄存器之间的内存传输引起的。如图8-5所示。
Similarly, a Short Scoreboard stall is caused by waiting on memory transfers between shared memory and the registers. This is shown in Figure 8-5.
Figure 8-5. Short Scoreboard stall caused by high-latency data transfers between shared memory and the registers
类似地，标记为"Stall: Memory Throttle"的指标意味着加载/存储流水线已饱和，因此无法发出额外的内存请求，因为硬件的内存队列已满。"Stall: Not Selected"意味着虽然线程束有资格发出，但它们必须等待空闲的内存事务槽才能继续。每当这些内存相关停顿原因主导您的停顿配置时，这是内核内存受限的明显迹象。
Similarly, metrics labeled "Stall: Memory Throttle" mean the load/store pipelines are saturated so that no additional memory requests can be issued because the hardware's memory queues are full. "Stall: Not Selected" means that although warps are eligible to issue, they must wait for free memory transaction slots before they can proceed. Whenever these memory-related stall reasons dominate your stall profile, it is a clear sign the kernel is memory bound.
使用Nsight Compute的源代码视图来突出显示受记分牌依赖和其他硬件限制（例如特殊功能单元吞吐量等）影响的代码。
Use Nsight Compute's source code view to highlight code impacted by scoreboard dependencies and other hardware limitations (e.g., special function unit throughput, etc.).
执行依赖停顿 (Execution-Dependency Stalls)¶
高比例的"Stall: Exec Dependency"意味着线程束经常等待先前指令的结果。例如，一条指令依赖于尚未完成的先前指令的输出。这通常是每个线程内指令级并行性不足的迹象--或者是ALU延迟瓶颈。在简单的依赖算术操作序列中，后续指令必须等待前面的指令完成。这导致线程束变为空闲。如果Exec Dependency停顿占主导，内核可能是在等待计算而延迟受限。换句话说，每个线程的指令流水线由于顺序依赖关系而不够忙碌。
A high fraction of "Stall: Exec Dependency" means that warps are often waiting on the results of previous instructions. For instance, one instruction depends on the output of a prior instruction that has not yet completed. This is usually a sign of insufficient instruction-level parallelism within each thread-or an ALU latency bottleneck. In a simple sequence of dependent arithmetic operations, later instructions must wait for earlier ones to finish. This causes the warp to go idle. If Exec Dependency stalls dominate, the kernel is likely latency bound waiting on compute. In other words, each thread's instruction pipeline isn't busy enough due to sequential dependencies.
执行单元争用 (Execution Unit Contention)¶
如果您看到显著的"Stall: Compute Unit Busy"--或者FP32 CUDA核心或Tensor Core的非常高活跃利用率--内核可能是计算受限的。在这种情况下，数学单元（FP32/FP64 ALU、Tensor Core等）已饱和。具体来说，线程束准备执行更多指令，但执行单元无法更快地服务它们。这通常表现为接近峰值的"ALU pipe busy"指标，并可能与高功耗相关。某些GPU上的相关停顿原因是"Stall: Math Pipe Throttle"，这意味着线程束正在等待，因为数学流水线在每个周期都完全占用。
If you see significant "Stall: Compute Unit Busy"-or simply very high active utilization of the FP32 CUDA cores or Tensor Cores-the kernel is likely compute bound. In this case, the math units (FP32/FP64 ALUs, Tensor Cores, etc.) are saturated. Specifically, warps are ready to execute more instructions, but the execution units can't service them any faster. This often manifests as near-peak "ALU pipe busy" metrics and can correlate with high power usage. A related stall reason on some GPUs is "Stall: Math Pipe Throttle," which means that warps are waiting because the math pipelines are fully occupied on every cycle.
其他停顿原因 (Other Stall Reasons)¶
虽然前面的停顿原因是最常见的罪魁祸首，但还有许多其他不太常见的停顿类别，包括指令获取停顿、纹理单元停顿、同步停顿等。例如，"Stall: Memory Dependency"类似于Long Scoreboard，其中操作正在等待先前的内存操作。"Stall: Texture Throttle"表示纹理缓存子系统的瓶颈。"No Eligible"或"Idle"表示线程束调度器在给定周期没有准备发出的线程束。这通常是由于先前的同步--或者仅仅是启动的线程束不够。
While the previous stall reasons are the most common culprits, there are many other less common stall categories, including instruction fetch stalls, texture unit stalls, synchronization stalls, etc. For instance, "Stall: Memory Dependency" is similar to Long Scoreboard in which an operation is waiting on a prior memory operation. "Stall: Texture Throttle" indicates a bottleneck in the texture caching subsystem. "No Eligible" or "Idle" indicates the warp scheduler had no warps ready to issue at a given cycle. This is often due to a prior synchronization-or simply not enough warps launched.
在实践中，您不需要记住每个停顿原因。相反，查看停顿细分的最大部分。如果内存相关停顿占主导，则是内存受限。如果执行依赖占主导，则是ALU延迟受限。
In practice, you don't need to memorize every stall reason. Instead, look at the largest portions of the stall breakdown. If memory-related stalls dominate, it's memory bound. If execution dependency dominates, it's latency bound on ALU.
如果几乎没有停顿且流水线活跃度接近100%，则可能是计算受限。如果线程束经常"未被选中"或空闲，可能表示并行工作不足或占用率问题。通过检查内核的线程束停顿细分，您可以缩小瓶颈类别。
If there are almost no stalls and near 100% pipeline active, it's likely compute bound. If warps are often "not selected" or idle, it could indicate insufficient parallel work or an occupancy issue. By examining the warp stall breakdown for your kernel, you can narrow down the bottleneck category.
Nsight Compute为每个内核启动提供这些停顿指标。确保分析具有代表性的工作负载。小型测试内核可能不会表现出与实际应用程序相同的停顿配置。在得出结论之前，始终从实际运行中收集数据。
Nsight Compute provides these stall metrics per kernel launch. Make sure to profile representative workloads. A tiny test kernel might not exhibit the same stall profile as your real application. Always collect data from realistic runs before drawing conclusions.
表8-1显示了常见线程束停顿原因、其典型解释和可能的优化方法的列表。
Table 8-1 shows a list of common warp stall reasons, their typical interpretation, and possible optimization approaches.
表8-1. 常见线程束停顿原因和优化提示
| 线程束停顿原因 | 含义/原因 | 潜在优化 |
|---|---|---|
| 执行依赖 | 线程束在等待先前的依赖指令。这通常表示线程内指令级并行性（ILP）不足。 | 增加ILP（在同一线程中做独立工作）。展开循环或重新排序指令，使长延迟操作（如乘法或复杂数学）有其他工作可以重叠。如果ILP已最大化但仍然停顿，依赖更多线程束（占用率）来隐藏延迟。 |
| 内存依赖（Long Scoreboard停顿） | 线程束在等待内存加载（"长延迟"）完成后才能继续。线程束中没有其他工作可以在数据到达之前继续。 | 通过有更多准备运行的线程束或更高的占用率来隐藏内存延迟，以便其他线程束可以在此线程束等待时运行。或使用异步内存预取将数据复制到共享内存，然后继续计算。这将使内存操作与计算重叠。为最小化延迟，确保内存访问高效，使用正确的合并，并利用正确的缓存层次结构。在现代GPU上，您应该使用张量内存加速器（TMA）进行批量多维复制，以便内存移动以低线程开销与计算重叠。 |
| 同步（屏障） | 线程束在__syncthreads()或其他同步处等待，空闲直到所有线程到达屏障。通常表示频繁同步或负载不平衡（某些线程束较早到达并等待）。 |
减少不必要的同步。重新检查算法以找到合并阶段的方法或使用细粒度同步。确保工作均匀分布，使线程束大致同时到达屏障。在某些情况下，使用更新的同步原语（例如线程束级同步或集群同步）来限制同步范围。 |
| 指令获取/发出 | 线程束在等待获取下一条指令（可能是指令缓存未命中或流水线问题）或未发出因为所需的执行流水线繁忙。这可能发生在非常大的内核或某些流水线争用场景下。 | 如果指令缓存未命中是问题，考虑减少内核大小（拆分内核或避免过度展开导致代码大小膨胀）。如果是流水线问题（一种类型的指令使功能单元饱和），尝试混合指令类型或再次使用ILP来利用不同的流水线。 |
| 未选中（调度器） | 线程束准备就绪但在该周期未被选中发出（调度器选择了另一个线程束）。这通常意味着有其他线程束可用，这个线程束只是等待轮次。这不是依赖停顿。通常表示其他准备就绪的线程束在该周期被选中发出，这在健康的占用率下是预期的。 | 通常不是问题--意味着GPU有其他工作要做，并在该周期选择了不同的线程束。如果您看到高百分比的"未选中"，这意味着高占用率正在发挥作用（隐藏延迟）。除非它表示不平衡（一个线程束占用调度器；在罕见情况下，您可能使用调度器提示或yield），否则不需要采取行动。 |
Table 8-1. Common warp stall reasons and optimization hints
通过解释这些停顿原因，您可以决定采用哪种优化。例如，高内存停顿时间意味着您应该尝试通过更多线程束、更好的内存访问模式或异步预取来隐藏延迟。
By interpreting these stall reasons, you can decide which optimization to pursue. For example, high memory stall time means you should try to hide latency with more warps, better memory access patterns, or asynchronous prefetch.
看到高执行依赖停顿表明您应该增加ILP或重新排列代码。频繁的屏障停顿意味着您可能应该重新设计同步。这一分析步骤将您引导到最有效的优化，而不是盲目地调整一切。
Seeing high execution dependency stalls suggests that you should increase ILP or rearrange code. Frequent barrier stalls mean you should probably redesign synchronization. This analysis step points you to the most effective optimizations instead of blindly tuning everything.
检查实现的占用率和GPU利用率 (Inspecting Achieved Occupancy and GPU Utilization)¶
分析器报告的另一个关键指标是实现的占用率，即执行期间每个SM上占用的硬件线程槽（线程束）的平均比例。例如，如果GPU支持每个流式多处理器（SM）64个线程束，实现的占用率为30%，那么每个SM平均有19个线程束处于活跃状态。
Another key metric reported by profilers is achieved occupancy, or the average fraction of hardware thread slots, warps, that were occupied on each SM during execution. For example, if a GPU supports 64 warps per streaming multiprocessor (SM) and achieved occupancy is 30%, then on average 19 warps were active per SM.
低实现的占用率通常表示利用率不足，因为您无法用少量活跃线程束隐藏足够的延迟。另一方面，如果您以接近最大的占用率运行，但仍然没有看到预期的性能提升，其他瓶颈可能是原因。这些包括内存带宽限制、低效的指令流和次优的内存访问模式。在这些非理想情况下，简单地添加更多线程无济于事。相反，您应该调查以改善内存合并、增加算术强度并优化线程束级效率。我们将在接下来的章节中介绍这些技术。
Low achieved occupancy often signals underutilization since you can't hide enough latency with only a few active warps. On the other hand, if you're running at near-maximum occupancy, but still not seeing the expected performance benefits, other bottlenecks are likely the cause. These include memory bandwidth limitations, inefficient instruction streams, and suboptimal memory-access patterns. In these nonideal cases, simply adding more threads won't help. Instead, you should investigate to improve memory coalescing, increase arithmetic intensity, and optimize warp‐level efficiency. We'll cover these techniques in the upcoming sections.
Nsight Compute还报告占用率限制因素，包括哪个资源正在限制理论占用率。例如，"Limited by max registers per thread"表示您的内核的寄存器使用阻止每个SM调度更多线程束。
Nsight Compute also reports occupancy limiters, including which resource is constraining the theoretical occupancy. For example, "Limited by max registers per thread" means your kernel's register usage is preventing more warps from being scheduled per SM.
然而，"Limited by shared memory per block"表示内核的每个块共享内存分配是占用率的瓶颈。而"Limited by thread count"表示启动配置本身、网格或块大小没有请求足够的线程来填满GPU。
However, "Limited by shared memory per block" means the kernel's shared memory allocation per block is the bottleneck for occupancy. And "Limited by thread count" means the launch configuration itself, grid, or block size didn't request enough threads to fill the GPU.
这些限制因素每一个都暗示不同的修复方法，我们将在下一节详细讨论。例如，如果寄存器是限制因素，您可能会减少寄存器使用或使用__launch_bounds__来允许更多块。如果共享内存是限制因素，您可以尝试使用更小的分块和每个块更少的共享内存。
Each of these limiters hints at different fixes, which we will discuss in more detail in the next section. For instance, if registers are the limiter, you might reduce register usage or use launch_bounds to allow more blocks. If shared memory is the limiter, you can try to use smaller tiles and less shared memory per block.
超过某一点，增加占用率可能不会产生加速。非常低的占用率（例如10%-20%）会由于糟糕的延迟隐藏而损害性能。另一方面，如果内存带宽和指令依赖等其他因素成为瓶颈，将占用率推到100%并不总是有益的。
Beyond a certain point, increasing occupancy may not produce a speedup. Very low occupancy (e.g., 10%-20%) will hurt performance due to poor latency hiding. On the other hand, pushing occupancy to 100% isn't always beneficial if other factors like memory bandwidth and instruction dependencies become the bottleneck.
因此，您应该检查占用率之外的硬件利用率。Nsight Compute报告诸如实现的内存带宽（GB/s）、实现的FLOPS（TFLOPS）、每周期指令数（IPC）、发射槽利用率和其他资源利用率统计等指标。这些数字将显示您的内核与GPU物理硬件限制的接近程度。
As such, you should examine hardware utilization beyond just occupancy. Nsight Compute reports metrics such as achieved memory bandwidth in GB/s, achieved FLOPS in TFLOPS, instructions per cycle (IPC), issue-slot utilization, and other resource utilization statistics. These numbers will show how close your kernel is to the GPU's physical hardware limits.
例如，如果您的内核ALU利用率低而内存吞吐量达到峰值的95%，您几乎可以确定是内存带宽受限。相反，如果ALU利用率接近最大值但内存吞吐量仍然适中，则内核是计算受限的。在这种情况下，您只能通过增加算术吞吐量来获得速度--通常通过切换到更低精度类型（FP16、FP8或FP4）并将工作转移到GPU更快的Tensor Core。
For example, if your kernel's ALU utilization is low while its memory throughput is at 95% of peak, you are almost certainly memory-bandwidth bound. Conversely, if ALU utilization is near its maximum but memory throughput remains modest, the kernel is compute bound. In this case, you'll gain speed only by increasing arithmetic throughput-typically by switching to lower-precision types (FP16, FP8, or FP4) and moving work onto the faster Tensor Cores of the GPUs.
如果ALU利用率和内存吞吐量都很低，您的内核可能正在经历长延迟操作、同步开销或仅仅是并行工作不足。这可能表示低指令级并行性（将在接下来的章节中讨论）--或者您没有启动足够的线程来充分利用GPU。
If both ALU utilization and memory throughput are low, your kernel may be experiencing long-latency operations, synchronization overhead, or simply insufficient parallel work. This could indicate low instruction-level parallelism (discussed in an upcoming section)-or that you haven't launched enough threads to fully utilize the GPU.
您可以使用屋顶线模型来构建此分析，该模型将内核的FLOPS相对于其算术强度（每字节内存访问的FLOPS）进行绘制。屋顶线定义了计算屋顶（GPU可以维持的最大FLOPS）和内存屋顶（最大内存带宽）。如果您的内核的FLOP/byte比率低于硬件的计算与内存比率，您是内存受限的，因为您无法足够快地提供数据。如果比率很高但实际FLOPS仍远低于峰值，内核可能是延迟受限或缺乏足够的ILP来使计算单元饱和。
You can frame this analysis using the Roofline model, which plots a kernel's FLOPS against its arithmetic intensity (FLOPS per byte of memory accessed). The roofline defines a compute roof (the maximum FLOPS the GPU can sustain) and a memory roof (the maximum memory bandwidth). If your kernel's FLOP/byte ratio falls below the hardware's compute-to-memory ratio, you are memory bound because you cannot supply data fast enough. If the ratio is high but actual FLOPS remains far below peak, the kernel may be latency bound or lacking sufficient ILP to saturate the compute units.
随着每一代新GPU，内存带宽适度增加，但计算能力增长得更快。因此，内核往往会随着时间的推移变得更加内存受限。
With each new GPU generation, memory bandwidth increases modestly, but compute capacity grows much faster. As such, kernels tend to become more memory bound over time.
在实践中，始终比较两个关键数字：您的内核的内存吞吐量与硬件的峰值内存带宽--以及您的内核的计算吞吐量与硬件的峰值FLOPS。这些比较将告诉您下一次优化应该专注于内存访问、计算工作还是并行性。让我们接下来看看这些。
In practice, always compare two key figures: your kernel's memory throughput versus the hardware's peak memory bandwidth-as well as your kernel's compute throughput versus the hardware's peak FLOPS. Those comparisons will tell you whether your next optimization should focus on memory access, compute work, or parallelism. Let's look at each of these next.
内核内存吞吐量与峰值HBM内存带宽 (Kernel Memory Throughput Versus Peak HBM Memory Bandwidth)¶
Nsight Compute报告您的内核实现多少GB/s。如果您的内核显示接近峰值的内存带宽利用率，执行更多计算不会有帮助。您必须通过减少每次操作的内存流量来增加算术强度。您可以通过使用Tensor Core的降低精度、通过并发隐藏更多延迟以及使用内核融合来增加算术强度，我们稍后会介绍。这些技术通过减少正在传输的中间数据大小来帮助减少全局内存流量。
Nsight Compute reports how many GB/s your kernel achieves. If your kernel is showing near-peak memory bandwidth utilization, performing more computations won't help. You would have to increase arithmetic intensity by reducing memory traffic per operation. You can increase arithmetic intensity by using reduced precision with Tensor Cores, hiding more latency with concurrency, and using kernel fusion, which we'll cover in a bit. These techniques help reduce global memory traffic by decreasing the size of intermediate data that is being transferred.
例如，如果内核达到GPU内存带宽的约80%或更多，它可能是内存受限的，因为几乎没有剩余空间。但是，请注意Blackwell GPU具有相对较大的126 MB L2缓存。其双芯片设计使用高带宽10 TB/s互连NVIDIA高带宽接口（NV-HBI），因此两个GPU芯片在内存访问方面表现为一个。因此，许多以前在旧GPU上内存受限的内核可能会在更新的GPU世代上更好地利用缓存。
For instance, if a kernel reaches ~80% or more of GPU memory bandwidth, it's likely memory bound since there's little headroom left. However, note that Blackwell GPUs have a relatively large 126 MB L2 cache. Its dual-die design uses a high-bandwidth 10 TB/s interconnect called NVIDIA High Bandwidth Interface (NV-HBI), so the two GPU dies act as one for memory access purposes. As a result, many kernels that were memory bound on older GPUs might better exploit caches on newer GPU generations.
您可以使用Nsight Compute的内存图表查看有多少流量流向L2与DRAM。高L2命中率可以缓解全局内存瓶颈，这意味着内核可能实际上是计算受限的，尽管执行了大量内存访问。片上缓存正在处理大量内存访问。
You can use Nsight Compute's memory chart to see how much traffic goes to L2 versus DRAM. A high L2 hit rate could alleviate global memory bottlenecks, which means the kernel might actually be compute-limited despite performing heavy memory accesses. The on-chip cache is servicing a lot of the memory accesses.
在Blackwell上，您可以控制L2数据持久性以保持关键工作集驻留。
On Blackwell, you can control L2 data persistence to keep critical working sets resident.
内核计算吞吐量与峰值GPU FLOPS (Kernel Compute Throughput Versus Peak GPU FLOPS)¶
低实现的FLOPS可能由低占用率或指令级停顿引起。例如，如果平均只有30%的线程束活跃--或者流水线由于内存等待经常空闲--内核将远低于计算屋顶。您可以使用Nsight Compute的占用率部分和源计数器来精确定位这些问题。确保内核启动足够的线程来填满GPU，达到您设备的每SM驻留线程束限制（例如，每SM 64个驻留线程束）。
Low achieved FLOPS can be caused by low occupancy or instruction-level stalls. For example, if only 30% of warps are active on average-or if pipelines are often idle due to memory waits-the kernel will sit far below the compute roof. You can use Nsight Compute's Occupancy section and Source Counters to pinpoint these issues. Ensure the kernel launches enough threads to fill the GPU, up to the per-SM resident warps limit for your device (e.g., 64 resident warps per SM.)
只有同一SM内的线程束必须共享资源--它们可以相互隐藏延迟。不同SM上的线程束没有交互。因此，实现的占用率是按SM测量的。
Only warps within the same SM must share resources-and they can hide one another's latency. Warps on different SMs have no interaction. As such, achieved occupancy is measured per SM.
此外，检查指令发射效率和吞吐量指标。Blackwell扩展到每设备许多SM，因此内核级别的利用率不足可能转化为大的聚合损失。（注意：双芯片封装本身不会增加每SM发射率。）如果实现的FLOPS为中等到高但内存吞吐量低，内核可能主要是计算密集型但受指令依赖限制。您可以通过检查Nsight Compute中的"Exec Dependency"停顿指标--以及任何其他停顿原因来确认。
We'll cover occupancy in a bit. Additionally, examine instruction issue efficiency and throughput metrics. Blackwell scales to many SMs per device, so underutilization at the kernel level can translate to large aggregate losses. (Note: The dual-die packaging does not, by itself, increase the per-SM issue rate.) If the achieved FLOPS are moderate to high but memory throughput is low, the kernel might be mostly compute-focused but limited by instruction dependencies. You can confirm by checking the "Exec Dependency" stall metric-as well as any other stall reasons in Nsight Compute.
如果计算吞吐量接近峰值，那么您确实是计算受限的。在这种情况下，您已经优化了内核的内存访问模式--或者您已经在使用低精度Tensor Core来达到如此高的FLOPS。
If compute throughput is near peak, then you're indeed compute bound. In this case, you've either optimized the kernel's memory access patterns-or you're already using low-precision Tensor Cores to achieve such high FLOPS.
持续接近峰值的计算利用率有时会在现代GPU上调用功率管理计算限制器以保持其健康。在分析、基准测试和调优时，请务必考虑到这一点。查看是否受到功率限制的最简单方法是使用以下命令实时监控强制功率限制以及任何"HW Slowdown"标志：
Sustained near-peak compute utilization can sometimes invoke power-management compute limiters on modern GPUs to keep them healthy. Be sure to account for this when profiling, benchmarking, and tuning. The easiest way to see if you're power limited is to monitor enforced power limits and any "HW Slowdown" flags in real time with the following command:
nvidia-smi \
--query-gpu=\
power.draw,clocks.current.sm,clocks.current.memory,\
clocks_event_reasons.active \
--format=csv -l 1
此命令每1秒打印一行，包含当前功耗、图形/内存时钟和节流原因。有了这些信息，您可以精确定位GPU何时达到功率上限并降频。
This command prints a line every 1 second with the current power draw, graphics/memory clocks, and throttling reasons. With this information, you can pinpoint when the GPU hits its power cap and downclocks.
您还可以使用NVIDIA管理库（NVML）API，它提供对CUDA C++ nvmlDeviceGetPowerUsage()和nvmlDeviceGetEnforcedPowerLimit() API的编程访问--以及等效的NVML Python API。这些非常适合自定义脚本或与监控系统集成。
You can also use the NVIDIA Management Library (NVML) API, which provides programmatic access to the CUDA C++ nvmlDeviceGetPowerUsage() and nvmlDeviceGetEnforcedPowerLimit() APIs-as well as equivalent NVML Python APIs. These are great for custom scripts or integration with monitoring systems.
简而言之，结合使用占用率、线程束停顿原因、内存吞吐量和计算吞吐量指标来诊断瓶颈。确保您看到高占用率、高线程束效率和执行单元（包括Tensor Core）的平衡使用。
In short, use a combination of occupancy, warp stall reasons, memory throughput, and compute throughput metrics to diagnose bottlenecks. Make sure you see high occupancy, high warp efficiency, and balanced use of execution units, including Tensor Cores.
迭代分析和确定内核瓶颈 (Iteratively Profiling and Determining the Kernel Bottleneck)¶
GPU可能因四个根本不同的原因而停顿：利用率不足、延迟受限、内存受限和计算受限。这些状态相关，但重要的是独立理解每一个，以便选择正确的优化来追求。通常，修复一个瓶颈会揭示另一个。
GPUs can stall for four fundamentally different reasons: underutilized, latency bound, memory bound, and compute bound. These states are related, but it's important to understand each independently so you can choose the right optimization to pursue. Often, fixing one bottleneck reveals another.
利用率不足发生在您根本没有启动足够的线程或工作时。在这种情况下，FLOPS和内存带宽都保持较低，执行时间线有空闲间隙。一旦增加并行性，您会发现您的线程束现在正在停顿并等待内存加载。
Underutilization occurs when you simply haven't launched enough threads or work. In this case, both FLOPS and memory bandwidth stay low, and the execution timeline has idle gaps. Once you increase parallelism, you'll find that your warps are now stalling and waiting on memory loads.
一旦更充分利用GPU，您现在可以区分延迟受限和内存受限。延迟受限的内核每秒发出的字节数远少于硬件可以提供的，因为单个内存加载正在使线程束停顿。修复方法是通过增加占用率、使用更多ILP、预取和流水线来增加内存-计算重叠。
Once you're more fully utilizing the GPU, you can now distinguish between latency bound and memory bound. A latency-bound kernel issues far fewer bytes per second than the hardware could deliver because individual memory loads are stalling warps. The fix is to increase memory-compute overlap by increasing occupancy, using more ILP, prefetching, and pipelining.
相比之下，内存受限的内核正在使DRAM带宽饱和，但您的ALU处于空闲状态，不是因为停顿，而是因为由于内存流水线饱和，每秒根本没有更多数据可以获取。在这种情况下，您必须通过分块、融合、利用缓存（L1/纹理）或降低精度来减少内存流量，从而提高算术强度。
In contrast, a memory-bound kernel is saturating DRAM bandwidth, but your ALUs are idle not because of stalls, but because there's simply no more data to fetch per second due to memory pipeline saturation. In this case, you must reduce memory traffic-and thus increase arithmetic intensity-by tiling, fusing, exploiting caches (L1/texture), or lowering precision.
如果这些修复都不能产生更多速度，您已经进入计算受限领域，其中GPU的算术流水线（例如ALU和Tensor Core）是限制因素。在这里，您可以通过展开和软件流水线重叠独立指令来增加每线程ILP。在现代GPU上，统一核心不能在同一时钟周期执行INT32和FP32指令。换句话说，混合INT32和FP32工作负载不会在同一周期从同一核心执行两种类型。因此，可实现的发射率取决于您的指令混合。
If none of these fixes produce more speed, you've entered the compute-bound regime, where the GPU's arithmetic pipelines (e.g., ALUs and Tensor Cores) are the limiting factor. Here, you can increase per-thread ILP by unrolling and software pipelining to overlap independent instructions. On modern GPUs, the unified cores cannot execute INT32 and FP32 instructions in the same clock cycle. In other words, mixed INT32 and FP32 workloads won't execute both types from the same core in the same cycle. As such, the achievable issue rate depends on your instruction mix.
在优化时，您经常会按以下顺序经历这些状态：利用率不足→延迟受限→内存受限→计算受限。每次修复后，您可能会遇到新的瓶颈并应用相应的策略。在优化GPU代码时，您应该遵循如下描述的结构化方法。
When optimizing, you'll often progress through these states in this order: underutilized → latency bound → memory bound → compute bound. After each fix, you may encounter a new bottleneck and apply the corresponding strategy. When optimizing GPU code, you should follow a structured approach as described next.
首先，您应该分析。接下来，您可以识别瓶颈。您可以使用Nsight Compute获取内核级指标（例如线程束停顿、实现的占用率、内存与计算利用率），使用Nsight Systems获取应用程序级时间线（例如并发、空闲间隙）。
First, you should profile. Next, you can identify the bottleneck. You can use Nsight Compute for kernel-level metrics (e.g., warp stalls, achieved occupancy, memory vs. compute utilization) and Nsight Systems for application-level timelines (e.g., concurrency, idle gaps).
一旦识别了瓶颈，您可以确定内核是内存受限、计算受限、延迟受限还是仅仅是利用率不足。当GPU没有发出足够的工作时，它是利用率不足的。表8-2总结了这四种状态，包括常见的分析指标和补救措施。
Once you've identified the bottleneck, you can determine whether the kernel is memory bound, compute bound, latency bound, or simply underutilized. The GPU is underutilized when it's not issuing enough work. Table 8-2 summarizes these four states, including common profiler metrics and remedies.
表8-2. 内存受限与延迟受限与计算受限与GPU利用率不足
| 限制因素 | 描述 | 分析器指标 | 补救措施 |
|---|---|---|---|
| 内存受限 | 您正在移动尽可能多的数据--接近峰值DRAM带宽--但每字节没有足够的工作来充分利用ALU。 | 高内存带宽利用率接近峰值，低FLOPS。 | 增加算术强度（例如分块、融合），改善合并和缓存。 |
| 计算受限 | 您已经隐藏了内存延迟，不再使内存带宽饱和。现在ALU（例如CUDA核心和Tensor Core）是瓶颈。 | 高FLOPS接近GPU峰值，低内存利用率。 | 利用更多ILP（例如双发射、循环展开），使用专用单元（例如FP16/FP8/FP4/Tensor Core），减少依赖，融合工作，利用低精度或稀疏性。 |
| 延迟受限 | 您没有维持足够的并发工作来隐藏单个加载/存储延迟，因此线程束在等待数据时停顿。 | 低实现的带宽远低于峰值，高"stall-on-scoreboard"或"not selected"百分比。 | 提高占用率，添加ILP（例如展开、多个累加器），内核内流水线，软件预取。 |
| GPU利用率不足 | 您没有充分利用SM或启动足够的工作--内存和计算资源都保持空闲。 | 低占用率和低实现带宽，低FLOPS，时间线显示间隙或稀疏内核活动。 | 增加问题大小或批量工作，启动更多线程/块，融合任务，使用持久内核（第10章）或流（第11章）。 |
Table 8-2. Memory bound versus latency bound versus compute bound versus underutilizing the GPU
内存受限的内核是性能受内存吞吐量限制的内核--具体来说，如果GPU的全局内存无法足够快地向计算单元提供数据。在这种情况下，您的内核实现的FLOPS将接近内存带宽设置的屋顶线。如果您有大量线程但无法更快移动数据，就会发生这种情况。因此，您正面临内存带宽上限。
A memory-bound kernel is one where performance is limited by memory throughput-specifically, if the GPU's global memory cannot feed data to the compute units fast enough. In this case, your kernel's achieved FLOPS will be close to the roofline set by memory bandwidth. This happens if you have many threads but cannot move data any faster. As a result, you're facing a memory bandwidth ceiling.
相反，计算受限的内核使GPU的算术单元（ALU FP32 CUDA核心或低精度Tensor Core）饱和。如果内核接近核心的峰值FLOPS屋顶线，这是显而易见的。分析指标如实现的占用率、内存利用率和执行依赖停顿可以帮助确认此分类。
Conversely, a compute-bound kernel saturates the GPU's arithmetic units (ALU FP32 CUDA cores or lower-precision Tensor Cores). This is evident if the kernel is near the peak FLOPS roofline of the cores. Profiler metrics like achieved occupancy, memory utilization, and execution dependency stalls can help confirm this classification.
另一方面，当内核延迟受限时，每个线程花费大量时间等待单个内存加载而不是做有用的工作。实际上，这意味着当线程束发出全局内存加载以获取A[idx]时，该线程束中的所有32个线程将停顿直到该获取完成。如果代码立即发出另一个依赖加载或计算，线程束只是为每次加载空闲数百个周期。
On the other hand, when a kernel is latency bound, each thread spends a lot of time waiting on individual memory loads rather than doing useful work. In practice, this means that when a warp issues a global memory load to fetch A[idx], all 32 threads in that warp will stall until that fetch completes. If the code immediately issues another dependent load or computation, the warp just sits idle for hundreds of cycles per load.
GPU的线程束调度器可能在线程束延迟受限时切换到其他线程束，但如果每个线程束的结构相同（例如，一次加载→等待→计算→写入），很少有其他工作可以填补这些空闲周期。因此，内核从来没有足够的独立操作来隐藏长延迟DRAM访问的延迟。
The GPU's warp scheduler might switch to other warps when a warp is latency bound, but if every warp has the same structure (e.g., load→wait→compute→write, one at a time), there's little other work to fill those idle cycles. As a result, the kernel never has enough independent operations to hide the latency of long-latency DRAM accesses.
要摆脱延迟受限的情况，您需要给GPU多个操作来重叠。一种方法是通过启动更多线程和线程束来增加占用率。这样，当一个线程束在其加载上停顿时，另一个线程束准备运行。
To escape the latency-bound situation, you need to give the GPU multiple operations to overlap. One way is to increase occupancy by launching more threads and warps. This way, when one warp is stalled on its load, another warp is ready to run.
但同样重要的是增加每个线程束的ILP。我们将在本章后面更详细地介绍ILP。在高层次上，如果每个线程在做任何算术之前连续发出两个或更多独立加载--通过加载A[idx]和B[idx]--GPU可以在硬件中开始重叠这些加载。
But equally important is increasing ILP within each warp. We'll cover ILP in more detail later in this chapter. At a high level, if each thread issues two or more independent loads back-to-back before doing any arithmetic-by loading A[idx] and B[idx]-the GPU can start overlapping these loads in hardware.
在这种情况下，当第一个加载等待DRAM时，加载二（以及三和四等）已经可以在进行中。一旦这些加载中的任何一个返回，依赖的算术就可以执行。这将与剩余的待处理加载进一步重叠。
In this case, while the first load is waiting on DRAM, load two (and three and four, etc.) can already be in flight. As soon as any of these loads return, dependent arithmetic can execute. This will further overlap with the remaining pending loads.
通过增加ILP，每个线程和线程束总是有足够的工作在进行中，以便DRAM延迟被其他未完成的操作掩盖。净效果是将延迟受限的内核转变为保持SM流水线忙碌的内核。这将增加吞吐量和整体内核性能。当GPU利用率不足时，您没有启动足够的工作来保持SM和内存流水线忙碌，因此许多资源保持空闲。在这种情况下，通过增加批量大小或启动更多线程来增加工作量会有所帮助。
By increasing ILP, each thread and warp always has enough work in flight so that the DRAM latency is masked by other outstanding operations. The net effect is to transform a latency‐bound kernel into one that keeps the SM's pipelines busy. This will increase throughput and overall kernel performance. When the GPU is underutilized, you're not launching enough work to keep SMs and memory pipelines busy, so many resources remain idle. In this case, it will help to increase the amount of work by increasing the batch size or launching more threads.
优化内核 (Optimizing the Kernel)¶
在实践中，性能调优应遵循清晰的逐步工作流程。首先，识别您的内核处于哪种状态：内存受限、延迟受限、计算受限或利用率不足。接下来，应用相应的优化。如果您的内核是内存受限的，专注于减少和隐藏内存流量。您可以通过改善合并、提高占用率、增加ILP以及通过缓存或分块引入数据重用来做到这一点。当内核延迟受限时，意味着单个加载或指令延迟占主导，确保有足够的独立工作在进行中。您可以通过每个线程发出多个非依赖加载/操作或增加整体占用率来实现，以便调度器总是有准备运行的线程束。如果内核是计算受限的（即，ALU或Tensor Core饱和而内存空闲），切换到低精度算术（FP16/FP8/FP4）、将工作卸载到Tensor Core或将更多操作融合在一起可以提高算术吞吐量。最后，如果GPU利用率不足，占用率低且频繁空闲周期，只需启动更多线程或块（以便所有SM都有工作）通常足以让硬件忙碌，然后再应用更深入的优化。
In practice, performance tuning should follow a clear, step‐by‐step workflow. First, identify which regime your kernel occupies: memory bound, latency bound, compute bound, or underutilized. Next, apply the corresponding optimizations. If your kernel is memory bound, concentrate on reducing and hiding memory traffic. You can do this by improving coalescing, raising occupancy, increasing ILP, and introducing data reuse through caching or tiling. When the kernel is latency bound, meaning individual load or instruction latencies dominate, ensure there is enough independent work in flight. You can achieve this by issuing multiple nondependent loads/operations per thread or by increasing overall occupancy so that the scheduler always has warps ready to run. If the kernel is compute bound (i.e., ALUs or Tensor Cores are saturated while memory is idle), switching to lower-precision arithmetic (FP16/FP8/FP4), offloading work to Tensor Cores, or fusing more operations together can increase arithmetic throughput. Finally, if the GPU is underutilized, with low occupancy and frequent idle cycles, simply launching more threads or blocks (so that all SMs have work) is often enough to keep the hardware busy before applying deeper optimizations.
以下是帮助您将GPU性能调优视为科学过程的高级优化技术列表。我们将在接下来的几章中深入讨论这些，但它们每个都应包括识别限制因素、应用适当的优化和测量结果：
The following is a list of high-level optimization techniques to help you approach GPU performance tuning as a scientific process. We'll discuss these in depth in upcoming chapters, but each should involve identifying the limiting factor, applying the appropriate optimization, and measuring the results:
将内存受限工作负载转换为计算受限
Convert memory-bound workloads to compute bound
如果内存受限（低算术强度），按如下方式增加数据重用和每次启动的工作：应用分块以使用快速共享内存并减少冗余访问，融合内核以避免不必要的内存往返，确保内存访问已优化（按照第6章的合并、避免bank冲突），并考虑使用压缩或低精度来移动更少的数据。
If memory bound (low arithmetic intensity), increase data reuse and work per launch as follows: apply tiling to use fast shared memory and reduce redundant accesses, fuse kernels to avoid unnecessary memory round trips, ensure memory accesses are optimized (coalescing, avoiding bank conflicts per Chapter 6), and consider using compression or lower precision to move less data.
进一步优化计算受限工作负载
Further optimize compute-bound workloads
如果计算受限（例如，ALU利用率高但由于依赖未达到峰值），按如下方式增加有效指令吞吐量：使用ILP技术（例如展开循环、多个累加器以重叠独立操作），检查分支分歧（见第6章）并尝试重新组织工作以减少它，因为分歧浪费ALU周期，转移到Tensor Core和低精度以提高计算上限。如果内核处于计算屋顶，看看降低精度（FP32 → FP16 → FP8 → FP4）是否可以提供进一步加速。
If compute bound (e.g., high ALU utilization but not at peak due to dependencies), increase effective instruction throughput as follows: use ILP techniques (e.g., unroll loops, multiple accumulators to overlap independent operations), check for branch divergence (see Chapter 6) and try to reorganize work to reduce it since divergence wastes ALU cycles, move to Tensor Cores and lower precision to raise the compute ceiling. If the kernel is at the compute roof, see if lowering precision (FP32 → FP16 → FP8 → FP4) can provide further speedups.
为延迟受限工作负载增加并行性
Increase parallelism for latency-bound workloads
如果延迟受限（例如，频繁线程束停顿且没有足够的并行性来隐藏延迟），按如下方式在各级别增加并发：如果可能，启动更多线程/块直到延迟被隐藏，确保寄存器/共享内存不会过度限制占用率（例如，占用率调优和平衡资源使用），使用异步复制在每个线程/线程束内重叠内存和计算（例如，内核内流水线），使用多个流重叠独立任务或重叠复制与计算（例如，第11章描述的内核间并发）。如果启动开销是问题（例如，许多小内核），考虑使用协作组合并它们或简单地合并它们的代码（例如，内核融合）。
If latency bound (e.g., frequent warp stalls and not enough parallelism to hide latency), increase concurrency at all levels as follows: launch more threads/blocks if possible until latency is hidden, ensure registers/shared memory aren't overly limiting occupancy (e.g., occupancy tuning and balancing resource usage), use async copies to overlap memory and compute within each thread/warp (e.g., intra-kernel pipelining), use multiple streams to overlap independent tasks or overlap copies with compute (e.g., inter-kernel concurrency described in Chapter 11). If launch overhead is an issue (e.g., many small kernels), consider using cooperative groups to combine them or simply merge their code (e.g., kernel fusion).
增加GPU利用率
Increase GPU utilization
如果GPU利用率不足（例如，低SM活跃度、低占用率），确保启动足够的工作来使用所有SM。具体来说，确保您的内核以足够大的网格大小启动以充分利用GPU。一个常见的错误是启动与元素一样多的线程，但忘记每个线程只做少量工作。有时您需要多次传递或每个元素更多线程。
If the GPU is underutilized (e.g., low SM activity, low occupancy), ensure you're launching enough work to use all SMs. Specifically, ensure your kernel is launched with a large enough grid size to fully utilize the GPU. A common mistake is launching as many threads as elements, but forgetting that each thread does only a small amount of work. Sometimes you need multiple passes or more threads per element.
减少同步和主机端（CPU）停顿
Reduce synchronization and host-side (CPU) stalls
您应该删除任何不必要的cudaDeviceSynchronize()或阻塞GPU的主机端等待。如果工作负载本质上很小，考虑将其与其他工作批处理或并发运行多个实例（例如，CUDA流）。您还可以使用CUDA Graphs（见第12章）预定义并高效启动许多小内核的执行图。
You should remove any unnecessary cudaDeviceSynchronize() or host-side waits that block the GPU. If the workload is inherently small, consider batching it with other work or running multiple instances concurrently (e.g., CUDA streams). You can also use CUDA Graphs (see Chapter 12) to prerecord and efficiently launch execution graphs of many small kernels.
利用专用硬件和降低/混合精度
Leverage specialized hardware and reduced/mixed precision
Blackwell Tensor Core在PTX中公开第五代MMA指令为tcgen05.mma和相关加载/存储（例如，tcgen05.ld和tcgen05.st）。Tensor Core加速微缩放格式，包括MXFP8、MXFP4（OCP MX格式）和NVIDIA的NVFP4格式。它们还支持块缩放矩阵乘法（K分组），当存在缩放元数据时库会自动选择。TMEM和TMA支撑这些高吞吐量数据路径。我们将在本章后面和第9章更详细地讨论这些技术。
Blackwell Tensor Cores expose fifth-generation MMA instructions in PTX as tcgen05.mma and related load/store (e.g., tcgen05.ld and tcgen05.st). Tensor Cores accelerate microscaling formats including MXFP8, MXFP4 (OCP MX formats), and NVIDIA's NVFP4 format. They also support block-scaled matrix multiplication (K-grouped) that libraries automatically select when scaling metadata is present. TMEM and TMA underpin these high-throughput data paths. We'll discuss these techniques in more detail later in this chapter and in Chapter 9.
验证和迭代
Validate and iterate
每次优化后，重新分析。确认目标停顿或指标改善（例如，分块后内存停顿减少，调整块大小后实现的占用率上升，使用CUDA流后"SM吞吐量%"增加等）。还要观察总运行时间改善。有时一个瓶颈掩盖另一个；您可能修复内存带宽只是接下来变成计算受限（这很好--如果需要再解决那个）。
After each optimization, reprofile. Confirm that the target stalls or metrics improve (e.g., memory stalls reduce after tiling, achieved occupancy rises after adjusting block size, "SM throughput %" increases after using CUDA streams, etc.). Also observe overall runtime improvement. Sometimes one bottleneck masks another; you might fix memory bandwidth only to become compute bound next (which is good-solve that too if needed).
保持正确性和可接受的精度
Maintain correctness and acceptable accuracy
使用低精度或新并行策略时，使用断言或与参考结果的比较进行测试。确保加速不会以精度为代价，除非这对应用程序是可接受的。通常，像FP16甚至FP8这样的技术经过仔细验证，对AI模型的精度影响可以忽略不计。
When using lower precision or new parallel strategies, test with assertions or comparisons against reference results. Ensure that speedups don't come at the cost of accuracy, unless that's acceptable for the application. Often, techniques like FP16 or even FP8, when carefully validated, have negligible accuracy impact on AI models.
调优占用率 (Tuning Occupancy)¶
请记住第6章中提到，占用率是SM上活跃线程束与SM上可能活跃的最大线程束数的比率。低占用率（< 50%）意味着平均只有一半的可能线程束是活跃的。这可能表示您的内核受资源限制--例如每个线程块的寄存器或共享内存--而不仅仅是可用并行性。
Remember from Chapter 6 that occupancy is the ratio of active warps on the SM to the maximum number of warps that could be active on the SM. Low occupancy (< 50%) means that, on average, half of the possible warps were active. This might indicate that your kernel is limited by resources such as registers or shared memory per thread block-rather than just available parallelism.
占用率是衡量相对于硬件最大容量，SM上有多少线程或线程束处于活跃状态的指标。更高的占用率或更多进行中的线程束，使GPU能够更好地隐藏延迟。这是因为当一个线程束在等待内存加载停顿时，调度器可以快速切换到运行另一个线程束。
Occupancy is the measure of how many threads, or warps, are active on an SM relative to the hardware's maximum capacity. Higher occupancy, or more in-flight warps, allows the GPU to better hide latency. This is because when one warp is stalled waiting on a memory load, the scheduler can quickly switch to running another warp.
每个SM有更多线程束通常意味着GPU的流水线保持更忙碌，等待内存的周期更少。占用率调优是调整内核启动参数和资源使用（例如寄存器和共享内存）以最大化每个SM上有用并行性的实践，如图8-6所示。
More warps per SM generally means the GPU's pipelines stay busier, with fewer cycles wasted waiting on memory. Occupancy tuning is the practice of adjusting kernel launch parameters and resource usage (e.g., registers and shared memory) to maximize useful parallelism on each SM, as shown in Figure 8-6.
Figure 8-6. Tuning occupancy by balancing resource usage (e.g., registers per thread and shared memory) with parallelism (e.g., number of threads and warps)
请记住，占用率调优的目标是保持足够多的线程束活跃，以充分利用SM的流水线并隐藏长延迟操作。在理想情况下，您将实现100%占用率，填满所有可用的线程束槽。例如，每个Blackwell B200（计算能力10.0）SM有64个线程束槽或2,048个线程可用。
Remember that the goal of occupancy tuning is to keep enough warps active to fully utilize the SM's pipelines and hide long‐latency operations. In an ideal case, you would achieve 100% occupancy, filling all available warp slots. For instance, there are 64 warp slots, or 2,048 threads, available in each Blackwell B200 (compute capability 10.0) SM.
每个SM 64个线程束的限制在现代数据中心GPU架构如Ampere、Hopper和Blackwell中保持不变--即使整体GPU核心数量增加了。硬件性能改进来自更多SM、更大缓存、多芯片等。您可以使用Nsight Compute的占用率分析来确认目标设备的确切限制。有趣的是，NVIDIA RTX PRO 6000和Spark DGX（GB10超级芯片）支持更高的计算能力（12.x），但每个SM只允许48个线程束（1,536个线程）。
This per-SM limit of 64 warps has remained the same for modern datacenter GPU architectures like Ampere, Hopper, and Blackwell-even as the overall GPU core counts have increased. The hardware-performance improvements come from more SMs, larger caches, multidie, etc. You can use Nsight Compute's Occupancy analysis to confirm the exact limits on your target device. Interestingly, the NVIDIA RTX PRO 6000 and Spark DGX (GB10 superchip) support a higher compute capability (12.x), but only allow 48 warps (1,536 threads) per SM.
许多现实世界的内核在较低占用率下仍然表现良好--特别是如果它们的内存或算术延迟已经很小。或者如果它们利用像Tensor Core这样的高吞吐量单元，不需要那么多活跃线程束就能保持流水线忙碌。例如，计算受限的内核可能仅以50%占用率就达到峰值性能，因为每个线程束在没有等待的情况下做了大量工作。相比之下，内存受限的内核通常受益于高占用率，因为一些线程束可以在SM上运行，而其他线程束因等待内存传输而停顿。
Many real‐world kernels still perform well at lower occupancy-especially if their memory or arithmetic latencies are already small. Or if they leverage high‐throughput units like Tensor Cores that keep the pipelines busy without needing as many active warps. For instance, a compute-bound kernel might achieve peak performance with only 50% occupancy because each warp is doing lots of work without waiting. In contrast, a memory-bound kernel often benefits from high occupancy since some warps can run on the SM while others are stalled waiting for memory transfers.
为您的工作负载找到合适的占用率 (Find the Right Occupancy for Your Workload)¶
在实践中，有效的占用率调优通常在达到某一点后产生收益递减。例如，如果内核严重内存受限，从10%提高到50%的实现占用率可能会带来巨大提升，因为现在您有足够的线程束来覆盖延迟。但从50%提高到100%可能只带来很小的进一步收益，因为其他因素开始占主导，如缓存未命中、内存带宽饱和等。分析有助于确定最佳占用率。例如，您可以评估每周期合格线程束和每个调度器活跃线程束等指标。这些洞察有多少线程束准备发出与硬件可以处理多少。
In practice, effective occupancy tuning often produces diminishing returns after a certain point. If a kernel is severely memory bound, for instance, going from 10% to 50% achieved occupancy might give a huge boost because now you have enough warps to cover latency. But going from 50% to 100% might give only a small further gain since other factors start to dominate, such as cache misses, memory bandwidth saturation, etc. Profiling helps determine the optimal occupancy. For example, you can evaluate metrics like warps eligible per cycle and active warps per scheduler. These give insight into how many warps are ready to issue versus how many the hardware can handle.
每周期合格线程束报告每个周期处于"准备运行"状态且没有未完成数据或依赖停顿的线程束平均数量。每个调度器活跃线程束（通常等于SM上的调度器数量）是每个周期可以发出指令的最大线程束数。
Warps eligible per cycle reports the average number of warps that are in the "ready to run" state each cycle, without outstanding data or dependency stalls. Active warps per scheduler (typically equal to the number of schedulers on the SM) is the maximum number of warps that can issue an instruction each cycle.
如果每周期合格线程束低于每个调度器活跃线程束，GPU经常耗尽准备好的线程束。在这种情况下，当一个线程束在内存或长延迟指令上停顿时，没有另一个准备好的线程束可以切换。这表明您需要更多并发，形式为更高占用率或ILP来隐藏延迟。
If warps eligible per cycle is lower than active warps per scheduler, the GPU is often running out of ready warps. In this case, when one warp stalls on memory or a long-latency instruction, there's no other ready warp to switch to. This indicates you need more concurrency in the form of higher occupancy or ILP to hide latency.
相反，如果每周期合格线程束达到或超过调度器限制但内核仍然运行缓慢，这意味着您有足够多的线程束准备就绪，但由于其他停顿（如内存带宽饱和或执行依赖）无法发出。在这种情况下，最好专注于通过更好的合并和异步复制来隐藏内存延迟--或通过展开独立工作来增加ILP。这比简单添加更多线程是更好的方法。
Conversely, if warps eligible per cycle meets or exceeds the scheduler limit but the kernel still runs slowly, it means you have plenty of warps ready but they can't issue due to other stalls (like memory bandwidth saturation or execution dependencies). In this case, it's better to focus on hiding memory latency through better coalescing and async copies-or increasing ILP by unrolling independent work. This is a better approach than simply adding more threads.
如果您提高占用率并看到"Stall: Not Selected"或空闲百分比下降，内存流水线更频繁地忙碌，您已成功改善了占用率。如果占用率很高但Long Scoreboard仍然占主导，您可能需要其他技术，如改善内存访问模式或重叠计算。
If you raise occupancy and see "Stall: Not Selected" or idle percentages drop, and the memory pipeline is busy more often, you've successfully improved occupancy. If occupancy is high but Long Scoreboard still dominates, you may need other techniques like improving memory access patterns or overlapping computation.
在实践中，一旦达到中等占用率（例如60%-70%），收益将开始递减。因此，通常更有效的是追求更好的内存局部性、更高的ILP，以及使用共享内存和寄存器等片上内存来缓存数据。
In practice, once you reach moderate occupancy (e.g., 60%-70%), returns will start to diminish. As such, it's often more effective to pursue better memory locality, higher ILP, and using on-chip memories like shared memory and registers to cache data.
虽然最大化占用率确保许多线程束可用于运行，但如果在等待内存或执行分歧代码，这些线程束可能仍然空闲。因此，在实现合理的占用率（例如50%-70%）后，专注于线程束效率和延迟隐藏，而不是痴迷于100%占用率。
While maximizing occupancy ensures many warps are available to run, those warps might still be idle if they're waiting on memory or executing divergent code. So after achieving reasonable occupancy (e.g., 50%-70%), focus on warp efficiency and latency hiding rather than obsessing over 100% occupancy.
占用率是达到目的（隐藏延迟）的手段，而不是目的本身。一旦您有足够的线程束来保持GPU忙碌，其他优化将带来更好的回报。
Occupancy is a means to an end (hiding latency), not the end goal itself. Once you have enough warps to keep the GPU busy, other optimizations will give better returns.
占用率调优技术 (Techniques for Occupancy Tuning)¶
当分析表明占用率是限制因素时，有几个简单的步骤可以改善占用率：通过启动更多线程增加并行性、调整块大小、减少每线程资源使用，以及使用__launch_bounds__或占用率API。让我们在这里讨论每一个：
There are a few straightforward steps to improve occupancy when profiling suggests that it's the limiting factor: increase parallelism by launching more threads, adjust the block size, reduce per-thread resource usage, and use launch_bounds or the occupancy API. Let's discuss each of these here:
增加并行性（启动更多线程）
Increase parallelism (launch more threads)
改善占用率最简单的方法是启动更多工作，如果当前内核启动还没有使用所有SM。例如，如果您每个SM只启动20个线程束而GPU可以支持64个，增加网格大小或每块线程数可以提高占用率（假设有足够的数据要处理）。
The simplest way to improve occupancy is to launch more work if your current kernel launch isn't already using all SMs. For instance, if you launch only 20 warps per SM while the GPU can support 64, increasing the grid size or threads per block can raise occupancy (assuming there's enough data to process).
确保通过启动至少与SM数量相同的块来利用所有SM--通常还要更多，因为块在每个SM上并行执行。如果您的GPU显示低"SM Active Cycles"因为不是所有SM都有工作，扩大工作负载或批量大小。然而，仅仅使用更多线程是不够的--特别是如果每个线程使用大量资源如寄存器和共享内存。
Make sure to utilize all SMs by launching at least as many blocks as there are SMs-and usually more, since blocks execute in parallel on each SM. If your GPU shows low "SM Active Cycles" because not all SMs have work, scale up the workload or batch size. However, simply using more threads isn't enough-especially if each thread uses a lot of resources like registers and shared memory.
调整块大小
Adjust block size
有时每块线程数或线程块大小会限制占用率。非常大的线程块（例如1,024个线程）可能使用如此多的寄存器或共享内存，以至于一次只能在一个SM上容纳一个块。这导致低占用率。相比之下，使用中等大小的块（例如128-256个线程）允许多个块同时驻留在一个SM上，这增加了可能活跃线程束的总数。
Sometimes the threads per block, or thread block size, limits occupancy. A very large thread block (e.g., 1,024 threads) might use so many registers or shared memory that only one block can fit on an SM at a time. This leads to low occupancy. In contrast, using a medium-sized block (e.g., 128-256 threads) allows multiple blocks to be resident on an SM simultaneously, which increases the total number of warps that could be active.
最佳块大小可能因内核而异。关键是平衡块大小与资源使用。较小的块单独使用较少资源，因此允许更多块并行。
The optimal block size can vary from kernel to kernel. The key is to balance block size with resource usage. Smaller blocks use fewer resources individually, so they allow more blocks to run in parallel.
您可以使用内置的Nsight Compute占用率计算器和CUDA占用率API来试验不同的启动配置，帮助找到最大化占用率而不产生其他开销的最佳点。
You can use the built-in Nsight Compute Occupancy Calculator and the CUDA Occupancy API to experiment with different launch configurations, helping find the sweet spot that maximizes occupancy without incurring other overheads.
始终用实际计时实验验证占用率计算器建议的最佳块大小，因为有时略低占用率的配置由于更少的寄存器溢出或更好的内存合并而产生更高的吞吐量。
Always validate the Occupancy Calculator's suggested bestBlockSize with actual timing experiments, as sometimes a configuration with a bit lower occupancy produces higher throughput due to less register spilling or better memory coalescing.
减少每线程寄存器和共享内存使用
Reduce per-thread register and shared-memory usage
每个线程消耗寄存器--可能还有共享内存。如果内核每个线程使用太多寄存器或共享内存，编译器必须减少SM上可以活跃的线程或线程束数量。否则，它将超过SM的总寄存器文件或共享内存分配。要解决寄存器使用问题，您可以重构代码以使用更少的活跃变量，从而使用更少的寄存器，释放寄存器容量。这允许更多线程束同时驻留。类似地，您可以将-maxrregcount=<N>传递给编译器以限制每线程分配的寄存器数量。
Each thread consumes registers-and likely shared memory. If a kernel uses too many registers or shared memory per thread, the compiler must reduce how many threads or warps can be active on an SM. Otherwise, it will exceed the SM's total register file or shared-memory allocation. To address register usage, you can refactor your code to use fewer live variables, and therefore fewer registers, to free up register capacity. This allows more warps to be resident simultaneously. Similarly, you can pass -maxrregcount=
to the compiler to limit the number of registers allocated per thread.
当您限制每线程寄存器数量时，编译器被迫将每个线程放入更少的寄存器。这允许硬件在SM上调度额外的线程束--从而提高占用率。
When you limit the per-thread register count, the compiler is forced to fit each thread into fewer registers. This allows the hardware to schedule additional warps on the SM-thus raising occupancy.
当然，如果您过于激进地限制寄存器，编译器会将多余变量溢出到本地内存，这会损害性能。因此，您应该找到最大化占用率而不产生过度溢出的最小寄存器限制。
Of course, if you limit registers too aggressively, the compiler will spill excess variables to local memory, which hurts performance. So you should find the minimum register limit that maximizes occupancy without causing excessive spilling.
类似地，如果每个块使用大量共享内存（如大分块），共享内存将限制SM上可以容纳多少块。通过优化共享内存占用、仅存储所需内容并在可能时使用片上缓存，您可以为额外块释放容量。本质上，您可以通过每个线程/块使用更精简的资源来增加占用率。
Similarly, shared memory can limit how many blocks can fit on an SM if each block uses a lot of shared memory (like large tiles). By optimizing shared memory footprint, storing only what's needed, and using on-chip caches where possible, you can free up capacity for additional blocks. Essentially, you can increase occupancy by using leaner resources per thread/block.
但请注意，过度减少寄存器或共享内存可能会降低单线程性能。这是一个权衡。分析器的占用率限制器读数会告诉您寄存器或共享内存是否是瓶颈。这将帮助指导您的优化工作。
But be careful: reducing registers or shared memory too much can hurt single-thread performance. It's a trade-off. The profiler's occupancy limiter readings will tell you if registers or shared memory are the bottleneck. This will help guide your optimization efforts.
例如，在Blackwell上，如果您激进地展开循环，可能会看到"Limited by max registers per thread"。在这种情况下，考虑使用更少的展开迭代或拆分工作。这是因为Blackwell有255个每线程寄存器限制。占用率报告有助于量化这些权衡。
For example, on Blackwell, if you aggressively unroll loops, you might see "Limited by max registers per thread." In this case, consider using fewer unroll iterations or splitting up the work. This is because Blackwell has a 255 per-thread register limit. Occupancy reports help quantify these trade-offs.
使用__launch_bounds__
Use
__launch_bounds__
在某些场景中，您可能有意限制占用率或引导编译器针对特定占用率进行优化。CUDA允许您使用__launch_bounds__注解在内核代码中设置启动边界。这向编译器提示内核的预期使用模式和启动配置。
In some scenarios, you might intentionally limit occupancy or guide the compiler to optimize for a specific occupancy. CUDA allows you to set launch bounds in your kernel code using the launch_bounds annotation. This hints to the compiler about the kernel's expected usage patterns and launch configuration.
如果您分析发现性能在约50%占用率而不是100%时达到峰值，这通常意味着每个线程使用大量寄存器或共享内存。这将进一步限制占用率。在这种情况下，您可以使用__launch_bounds__引导编译器。
If you profile and find that performance peaks at around 50% occupancy instead of 100%, this often means each thread uses a lot of registers or shared memory. This will further limit occupancy. In this case, you can use launch_bounds to guide the compiler.
例如，您可以使用__launch_bounds__限制每块线程数并指定每SM最小块数，以允许更多线程束驻留在SM上。本质上，您告诉编译器用一些每线程寄存器/共享内存资源使用换取更高的线程束数量。
For example, you can use launch_bounds to limit the threads per block and specify a minimum number of blocks per SM, allowing more warps to be resident on the SM. Essentially, you're telling the compiler to trade some per-thread register/shared memory resource usage for higher warp counts.
当您的内核延迟受限时，这将改善吞吐量，因为更多线程束将隐藏内存延迟和依赖停顿。这与纯计算受限内核形成对比，后者不会从此类优化中获得相同收益。
This will improve throughput when your kernel is latency bound since more warps will hide memory latency and dependency stalls. This is in contrast to a purely compute-bound kernel, which wouldn't benefit as much from such optimizations.
使用CUDA占用率API
Use the CUDA Occupancy API
除了__launch_bounds__注解外，CUDA的占用率API函数（例如cudaOccupancyMaxPotentialBlockSize()和cudaOccupancyMaxActiveBlocksPerMultiprocessor()）让您在运行时确定给定内核实际寄存器和共享内存需求下产生最高占用率的块大小。
Beyond the launch_bounds annotation, CUDA's Occupancy API functions (e.g., cudaOccupancyMaxPotentialBlockSize() and cudaOccupancyMaxActiveBlocksPerMultiprocessor()) let you determine at runtime the block size that yields the highest occupancy given your kernel's actual register and shared memory requirements.
这些函数提供每块线程数和每SM块数的推荐值。这样，您不必猜测哪种配置最有效。在实践中，您可能使用占用率API找到候选块大小，然后用__launch_bounds__提示微调以锁定特定占用率水平。
These functions provide recommended values for threads per block and blocks per SM. This way, you don't have to guess which configuration is most effective. In practice, you might use the Occupancy API to find candidate block sizes, then fine-tune with launch_bounds hints to lock in a specific occupancy level.
另一种方法是使用CUDA Graphs在确定最佳配置后高效启动预定义工作负载。虽然不直接改变占用率，但CUDA Graphs可以减少每次启动开销，这在增加块数量时对占用率有益。我们将在第12章介绍CUDA Graphs，在第13章和第14章介绍PyTorch对CUDA Graphs的使用，但在此背景下值得一提。
Another approach is to use CUDA Graphs to launch predefined workloads efficiently once you've determined the optimal configuration. While not directly changing occupancy, CUDA Graphs can reduce per-launch overhead, which is beneficial for occupancy when you increase the number of blocks. We'll cover CUDA Graphs in Chapter 12 and PyTorch's use of CUDA Graphs in Chapters 13 and 14, but it's worth mentioning them in this context.
简而言之，建议使用多阶段方法，通过分析系统发现理想占用率范围--然后使用编译器和运行时提示实现理想占用率。目标是高效调整线程块大小，保持足够多的线程束进行中，而不过度分配寄存器和共享内存等稀缺片上资源。接下来，让我们仔细看看使用__launch_bounds__、占用率API和PyTorch调优占用率。
In short, it's recommended to use a multiphased approach by profiling your system to discover the ideal occupancy range-and then using compiler and runtime hints to achieve the ideal occupancy. The goal is to size your thread blocks efficiently to keep enough warps in flight without over allocating scarce on‐chip resources like registers and shared memory. Next, let's take a closer look at tuning occupancy with launch_bounds, the Occupancy API, and PyTorch.
优化占用率的编译器提示 (Compiler Hints to Optimize Occupancy)¶
我们可以通过使用CUDA的__launch_bounds__内核注解引导编译器优化占用率。此注解让我们为内核指定两个参数：我们想要启动的每块最大线程数（例如256）和我们希望在每个SM上保持驻留的最小线程块数（例如4），如下所示：
We can guide the compiler to optimize for occupancy by using CUDA's launch_bounds kernel annotation. This annotation lets us specify two parameters for a kernel: the maximum number of threads per block that we want to launch (e.g., 256) and the minimum number of thread blocks we want to keep resident on each SM (e.g., 4), as shown here: global launch_bounds(256, 8) void myKernel(...) { /* ... */ } Here, we promise never to launch myKernel with more than 256 threads per block. And we request that at least 8 blocks be active on each SM.
在这里，我们承诺永远不会以超过256个每块线程启动myKernel。并且我们请求GPU尝试保持每个SM至少8个块活跃。这些提示影响编译器的寄存器分配和内联决策。
Here, we promise to never launch myKernel with more than 256 threads per block. And we ask the GPU to try to keep at least 8 blocks active per SM. These hints influence the compiler's register allocation and inlining decisions.
具体来说，__launch_bounds__将限制每个线程的寄存器使用，使得最多256个线程可以放入一个块--并且每个SM至少可以有8个块活跃。这是每个SM总共2,048个活跃线程（8个块 x 每块256个线程 = 2,048个线程）。这填满了B200等设备的线程槽，其每SM限制为2,048个线程（例如Blackwell）。在实践中，__launch_bounds__可能导致编译器限制每线程寄存器使用并限制展开/内联。这避免了溢出并允许更高占用率。因此，我们实际上是用一些每线程性能（例如不使用每个最后的寄存器或最大展开）换取更稳定的线程束吞吐量，通过保持更多线程束进行中。
Specifically, launch_bounds will limit per-thread register usage so that at most 256 threads can fit in a block-and at least 8 blocks can be active per SM. This is a total of 2,048 active threads per SM (8 blocks x 256 threads per block = 2,048 threads). This fills up the thread slots on devices like the B200, which has a limit of 2,048 threads per SM (e.g., Blackwell). In practice, launch_bounds can cause the compiler to limit per-thread register usage and limit unrolling/inlining. This avoids spilling and allows for higher occupancy. As a result, we're effectively trading some per-thread performance (e.g., not using every last register or maximum unrolling) for more steady warp throughput by keeping more warps in flight.
增加的占用率必须与增加的每线程资源平衡。您要避免通过强制太多线程导致寄存器溢出。这会导致缓慢的内存访问，因为它们用完寄存器并溢出到本地内存（由全局HBM支持）。找到最佳点通常需要实验。Nsight Compute的"每线程寄存器"和"占用率"指标可以在此指导您。
Increased occupancy must be balanced with increased per-thread resources. You want to avoid register spilling by forcing too many threads. This causes slow memory access because they run out of registers and spill to local memory (backed by global HBM.) Finding the sweet spot often requires experimentation. Nsight Compute's "Registers per Thread" and "Occupancy" metrics can guide you here.
使用占用率API确定最佳启动配置 (Determine Optimal Launch Configuration with the Occupancy API)¶
您可以在运行时使用CUDA占用率API确定最佳启动配置，包括cudaOccupancyMaxActiveBlocksPerMultiprocessor()和cudaOccupancyMaxPotentialBlockSize()。例如，cudaOccupancyMaxPotentialBlockSize()将计算给定内核产生最佳占用率的块大小，考虑其寄存器和共享内存使用，如下所示：
You can determine an optimal launch configuration at runtime using the CUDA Occupancy API, including cudaOccupancyMaxActiveBlocksPerMultiprocessor() and cudaOccupancyMaxPotentialBlockSize(). For instance, cudaOccupancyMaxPotentialBlockSize() will calculate the block size that produces the optimal occupancy for a given kernel, considering its register and shared memory usage as shown here:
int minGridSize = 0, bestBlockSize = 0;
cudaOccupancyMaxPotentialBlockSize(
&minGridSize, &bestBlockSize,
myKernel,
/* dynamicSmemBytes = */ 0,
/* blockSizeLimit = */ 0 );
// bestBlockSize contains number of threads per block that maximizes occupancy
myKernel<<<minGridSize, bestBlockSize>>>(...);
此API计算考虑内核资源使用情况下可能最大化占用率的每块线程数。然后我们可以使用建议的bestBlockSize和minGridSize进行内核启动。
This API calculates the number of threads per block that would likely maximize occupancy, considering the kernel's resource usage. We can then launch the kernel with the suggested bestBlockSize and minGridSize.
在实践中，编译器的启发式方法通常相当不错。但在需要时使用__launch_bounds__和占用率API可以给您显式控制。当您知道内核可以用一些每线程资源使用换取更多活跃线程束时使用它们。这有助于防止由于重型线程而导致SM占用不足。
In practice, the compiler's heuristics are often quite good. But using launch_bounds and the Occupancy API when needed can give you explicit control. Use them when you know the kernel can trade some per-thread resource usage for more active warps. This helps prevent underpopulating the SM due to heavyweight threads.
使用PyTorch调优占用率 (Tuning Occupancy with PyTorch)¶
对于编写高级代码的PyTorch用户，您通常不直接管理占用率--PyTorch的CUDA内核和库在底层处理启动参数。像矩阵乘法、卷积等操作已经调优为有效使用GPU。然而，理解占用率仍然很有价值。如果您编写自定义CUDA内核作为PyTorch扩展，同样的原则适用：启动足够的线程/块来利用GPU，选择合理的块大小，避免使用过多寄存器或共享内存每线程（如果它会限制占用率）。
For PyTorch users writing high-level code, you don't usually manage occupancy directly-PyTorch's CUDA kernels and libraries handle launch parameters under the hood. Operations like matrix multiplies, convolutions, etc., are already tuned to use the GPU effectively. However, understanding occupancy can still be valuable. If you write custom CUDA kernels as PyTorch extensions, the same principles apply: launch enough threads/blocks to utilize the GPU, choose reasonable block sizes, and avoid using too many registers or shared memory per thread if it would limit occupancy.
即使在Python级别，也有一些事情需要注意。如果您使用PyTorch并注意到GPU利用率不足（例如，在分析跟踪中您看到只有一小部分SM活跃），这可能是因为非常小的张量操作没有启动很多线程。
Even at the Python level, there are a few things to watch for. If you use PyTorch and notice underutilization (e.g., in profiler traces you see only a fraction of SMs active), it could be because very small tensor operations aren't launching many threads.
极小的工作负载可能无法很好地扩展到大型GPU。在这种情况下，尝试批处理或组合操作，使每次启动做更多工作。PyTorch的图编译器可以使用torch.compile融合某些操作。这增加了每次内核启动的工作量--从而改善了占用率和效率。
Tiny workloads might not scale well to large GPUs. In such cases, try batching or combining operations so that each launch does more work. PyTorch's graph compiler can fuse certain operations using torch.compile. This increases the amount of work per kernel launch-thus improving occupancy and efficiency.
在底层，torch.compile生成融合的GPU内核。它通过将许多小操作合并为一个更大的内核来实现。在可能的情况下，利用PyTorch编译器，因为它在某些情况下可以达到手工编写CUDA内核的效率。对于稳定形状的长时间训练/推理，优先使用torch.compile(mode="max-autotune")，对于小批量、图友好的循环，使用mode="reduce-overhead"。两者在有利可图时都启用CUDA Graphs。如果不希望使用图，也可以使用mode="max-autotune-no-cudagraphs"。我们将在第13章和第14章深入讨论PyTorch编译器及其不同的模式选项。
Under the hood, torch.compile generates fused GPU kernels. It does this by merging many small operations into one larger kernel. Whenever possible, leverage the PyTorch compiler, as it can achieve the efficiency of a manually written CUDA kernel in some cases. Prefer
torch.compile(mode="max-autotune")for long-running training/inference on stable shapes, ormode="reduce-overhead"for small-batch, graph-friendly loops. Both enable CUDA Graphs when profitable. You can also usemode="max-autotune-no-cudagraphs"if graphs are undesirable. We dive deep into the PyTorch compiler and its different mode options in Chapters 13 and 14.
值得注意的是，PyTorch的内置内核通常在内部使用最佳实践。例如，PyTorch中的归约操作和逐元素操作使用启动配置器实现，根据设备和张量大小选择最佳块大小和块数。
It's worth noting that PyTorch's built-in kernels often use best practices internally. For example, reduction operations and element-wise operations in PyTorch are implemented with launch configurators that choose optimal block sizes and numbers of blocks based on device and tensor sizes.
如果您深入研究PyTorch的C++ CUDA代码，您会看到限制块大小并为某些算法每个SM启动多个块的逻辑。这本质上是自动化的占用率调优。为减少优化器开销，在可用时启用融合实现。例如，使用torch.optim.AdamW(..., fused=True)并与自动混合精度（AMP）配对。
If you dig into PyTorch's C++ CUDA code, you'll see logic that limits block sizes and launches multiple blocks per SM for certain algorithms. This is essentially automated occupancy tuning. To reduce optimizer overhead, enable fused implementations when available. For instance, use torch.optim.AdamW(..., fused=True) and pair with automatic mixed precision (AMP).
PyTorch用户的要点是优先使用已经优化的高级库和操作。这比编写自定义内核更可取。如果您确实需要自定义内核，应用我们讨论的相同占用率原则。
The takeaway for PyTorch users is to prefer high-level libraries and operations that are already optimized. This is preferable to writing custom kernels. If you do need a custom kernel, apply the same occupancy principles we discussed.
现在我们了解了如何用足够的线程保持GPU忙碌，接下来我们转向使每个线程束更高效。高占用率意味着如果每个线程束在浪费周期，那也没什么意义。这引出我们对线程束级效率优化的讨论，包括线程束分歧和线程束内线程协作。
Now that we understand how to keep the GPU busy with enough threads, we next turn to making each warp more efficient. High occupancy means little if each warp is wasting cycles. This leads us to our discussion of warp-level efficiency optimizations, including warp divergence and intrawarp thread cooperation.
提高线程束执行效率（线程束分歧） (Improving Warp Execution Efficiency (Warp Divergence))¶
即使由于良好的占用率而有大量线程束可用，每个线程束的内部效率也很重要。线程束执行效率衡量线程束通道平均活跃的比例。当线程束内的线程不都做相同的工作时，就会出现线程束级低效，这称为线程束分歧。另一个原因是线程正在等待数据加载。具体来说，低线程束执行效率值通常显示高分歧分支计数、低谓词效率，以及可能大量的内存相关停顿。这意味着线程束线程由于内核中的if/else语句或数据依赖导致的分支分歧而空闲。
Even with plenty of warps available, thanks to good occupancy, each warp's internal efficiency matters. Warp execution efficiency measures the fraction of warp lanes that are active on average. Warp-level inefficiencies arise when threads within a warp do not all do the same work, called warp divergence. Another cause is if the threads are waiting on data loads. Specifically, a low warp execution efficiency value typically shows a high divergent branch count, low predicate efficiency, and possibly a high number of memory-related stalls. This means that the warp threads are idle due to branching divergence caused by if/else statements in the kernel or by data-dependencies.
在这种情况下，您应该尝试重构代码以避免低效。要提高线程束执行效率，您应该尝试改善内存合并、最小化线程分歧，并使用线程束级内建函数进行最佳的线程束内（线程级）通信。
In this case, you should try to restructure your code to avoid the inefficiencies. To improve warp execution efficiency, you should try to improve memory coalescing, minimize thread divergence, and use warp-level intrinsics for optimal intrawarp (thread-level) communication.
与更高级别的算法更改相比，这些技术通常产生适度的加速（例如5%-30%），但在高度优化的代码中，每个周期都很重要，它们可能至关重要。它们也补充了其他优化，在优化占用率和改善内存访问后，您仍然可以通过确保每个线程束尽可能高效地执行来挤出额外的性能。
These techniques often produce modest speedups (say, 5%-30%) compared to higher-level algorithmic changes, but they can be crucial in highly optimized code where every cycle counts. They also complement other optimizations in that, after you've optimized occupancy and improved memory access, you can still squeeze out extra performance by ensuring each warp executes as efficiently as possible.
线程束分歧的原因 (Causes of Warp Divergence)¶
如第7章所述，线程束分歧或分支分歧发生在同一线程束中的线程采取不同的控制流路径时。例如，考虑内核中的if/else，其中线程束中的某些线程执行if子句，其他线程执行else。在SIMT执行中，线程束必须串行执行两条路径：首先，所有采取if分支的线程执行，而该线程束中的其他线程被屏蔽（空闲）。然后采取else的线程执行，而第一组空闲，如图8-7所示。
As mentioned in Chapter 7, warp divergence, or branch divergence, occurs when threads in the same warp take different control-flow paths. For example, consider an if/else inside a kernel where some threads in a warp execute the if clause and others execute the else. In SIMT execution, the warp must execute both paths serially: first, all threads that take the if branch execute while the others in that warp are masked off (idle). Then the threads that took else execute while the first group idles, as shown in Figure 8-7.
Figure 8-7. Divergent versus nondivergent warp execution
在分歧部分期间，实际上只有一半或某个比例的线程束在做有用的工作。这降低了整体吞吐量。在if和else之间50/50分配的情况下，线程束的活跃利用率在该部分代码中下降到50%。如果分配是1个线程对31个线程，那么31/32的线程将在其中一个子分支中空闲。
During the divergent section, effectively only half or some fraction of the warp is doing useful work. This reduces overall throughput. In a 50/50 split between if and else, the warp's active utilization drops to 50% for that section of code. If the split is 1 thread versus 31 threads, then 31/32 of the threads will be idle in one of the subbranches.
如果您的内核包含多个分歧点，或者每个分歧分支承载更重的工作，消除这些分支可以复合这些收益。在理想情况下（例如50/50分配分支），消除该分歧分支可以将近吞吐量翻倍，达到约2倍加速。消除多个重度分歧路径可以复合这些收益。
If your kernel contains multiple divergence points, or if each divergent branch carries heavier work, eliminating these branches can compound these gains. In the ideal case (e.g., 50/50 split branch), eliminating that divergent branch can nearly double throughput, achieving around a 2x speedup. Eliminating multiple heavily divergent paths can compound these gains.
整体效果是线程束分歧导致某些GPU核心空闲。这增加了总指令计数，因为每个分支路径由不同的线程子集串行执行。因此，最小化线程束内分歧是线程束级效率的关键。
The overall effect is that warp divergence leaves some GPU cores idle. This increases the total instruction count since each branch path is executed serially by different subsets of threads. As such, minimizing divergence within warps is key to warp-level efficiency.
避免线程束分歧的技术 (Techniques to Avoid Warp Divergence)¶
有多种最佳实践可以最小化线程束分歧，包括重构条件和将分支分离到多个内核中。线程束分歧只影响线程束内的线程。不同线程束中的线程不会相互导致停顿。此外，您可以使用线程束一致分支、线程束投票内建函数和谓词执行。让我们讨论每一个：
There are various best practices to minimize warp divergence, including restructuring conditions and separating branches into multiple kernels. Warp divergence only affects threads within a warp. Threads in different warps do not cause each other to stall. Additionally, you can use warp-unanimous branches, warp-vote intrinsics, and predication. Let's discuss each of these:
重构条件
Restructure conditions
尽可能组织您的计算，使同一线程束中的线程遵循相同的执行路径。这可能涉及将分歧条件移到内部循环之外的更高级别--或分离到单独的内核启动中。您还可以对数据进行排序或分组，使每个线程束处理更同质的情况。
Where possible, organize your computation so that threads in the same warp follow the same execution path. This might involve moving divergent conditions to a higher level outside inner loops-or separating into separate kernel launches. You can also sort or group data so that each warp processes more homogeneous cases.
例如，如果您有一个值数组，并且希望以不同方式处理负值和非负值，朴素方法可能在内核内部放置一个if(x<0)，每个元素都会分歧。更聪明的方法是对数据进行分区，使得一个内核处理所有负值，另一个内核处理非负值。通过与线程束对齐数据，您减少单个线程束必须分割其执行的机会。
For example, if you have an array of values and want to process negative and nonnegative values differently, a naive approach might put an if(x<0) inside the kernel, diverging on every element. A smarter approach is to partition the data so that one kernel handles all negatives and another handles nonnegatives. By aligning data with warps, you reduce the chance that a single warp has to split its execution.
分离到多个内核
Separate into multiple kernels
另一种方法是将工作分离到多个内核启动中，使得一个内核处理if情况，另一个处理else情况。您可以使用前缀和或压缩将线程分配到不同的内核。这以启动更多内核和添加逻辑在内核之间分发数据为代价避免分歧。如果分歧是一个大问题且分歧部分的指令计数足够大，这可能是值得的。
Another approach is to separate work into multiple kernel launches so that one kernel handles the if case and another handles the else case. You can use prefix sum or compaction to assign threads to different kernels. This avoids divergence at the cost of launching more kernels and adding logic to distribute data between them. If divergence is a big problem and the instruction count in the divergent sections is large enough, this can be worthwhile.
将条件重写为线程束一致的
Rewrite conditions as warp-unanimous
在某些情况下，您可以将条件重写为线程束一致的。这意味着线程束中的所有32个线程要么都满足条件，要么都不满足。一个常见的技巧是在条件语句中使用线程束索引。例如，您首先将线程束ID计算为int warpId = threadIdx.x / 32;。然后一半的线程束执行任务A，另一半执行任务B。为此，您编写：if (warpId % 2 == 0) { /* 任务A */ } else { /* 任务B */ }。
In some cases, you can rewrite conditions to be warp-unanimous. This means that all 32 threads in a warp either all satisfy the condition or all do not. A common trick is to use the warp index in the conditional statement. For example, you first compute the warp ID as int warpId = threadIdx.x / 32;. Then half the warps perform task A, and the other half perform task B. To do this, you write: if (warpId % 2 == 0) { /* Task A / } else { / Task B */ }.
在这种情况下，给定线程束中的所有线程要么一起进入任务A，要么进入任务B，没有线程束分歧，因为一个线程束执行一个分支，而另一个线程束执行另一个分支。但在单个线程束内，所有线程在要执行的分支上达成一致。
In this case, all threads in a given warp will either go into Task A together or into Task B, with no warp divergence since one warp executes one branch and another warp executes the other. But within a single warp, all threads agree on which branch to execute.
轻微的开销是每个线程束仍然评估分支条件，但由于它们都同意它，每个线程束只执行其中一个分支。这种技术本质上是以一些灵活性为代价--因为您正在约束工作如何划分--来获得线程束一致性。
The slight overhead is that each warp still evaluates the branch condition, but since they all agree on it, each warp only executes one of the branches. This technique essentially trades some flexibility-since you're constraining how work is divided-for warp uniformity.
当您可以将工作划分为32个线程的粗粒度增量时，请使用线程束一致分支。
Use warp-unanimous branches when you can divide work into coarse increments of 32 threads.
利用线程束投票并行算法
Utilize warp-vote parallel algorithms
线程束内建函数（__ballot_sync、__any_sync、__all_sync）、协作组的warp.ballot/any/all以及设备端投票掩码（%WarpVote）让线程束集体决定或投票哪些通道需要"特殊"工作。然后线程束动态委托、重新分区并将该工作压缩到一个（或几个）通道中，而不是分歧所有32个线程。这避免了每通道分支分歧，但引入了一些您应该分析影响的潜在负载不平衡权衡。我们稍后将更详细地介绍线程束内建函数。
Warp intrinsics (__ballot_sync, __any_sync, __all_sync), cooperative-groups' warp.ballot/any/all, and device-side vote masks (%WarpVote) let a warp collectively decide, or vote, which lanes need "special" work. The warp then dynamically delegates, repartitions, and compacts that work into one (or a few) lane(s) instead of diverging all 32 threads. This avoids per-lane branch divergence but introduces some potential load-imbalance trade-offs that you should profile for impact. We'll cover warp intrinsics in more detail in a bit.
谓词执行短通道
Predicate short lanes
CUDA编译器有时会对短条件代码进行谓词执行。谓词执行意味着编译器将if转换为每个线程的布尔掩码--并为所有线程执行两条路径。但是，它只提交适合每个线程路径的结果。
CUDA compilers will sometimes predicate short conditional code. Predication means the compiler turns the if into a per-thread boolean mask-and executes both paths for all threads. However, it only commits the result that fits each thread's path.
谓词执行以每个线程做额外工作为代价避免分歧分支。当分支非常短且分歧很高时，这是有益的。作为程序员，您可以通过编写无分支构造来鼓励谓词执行，包括?:三元运算符或简单条件的位运算技巧。例如，考虑这个使用if/else语句的朴素实现：
Predication avoids divergent branches at the cost of doing extra work per thread. It's beneficial when branches are very short and divergence is high. As a programmer, you can encourage predication by writing branchless constructs, including the ?: ternary operator or bitwise tricks for simple conditions. For example, consider this naive implementation using an if/else statement:
相反，您可以编写以下内容，为所有线程计算f(x)和g(x)，但将每个结果乘以cond以选择匹配条件的结果：
Instead, you can write the following, which computes both f(x) and g(x) for all threads, but multiplies each result by cond to select the matching result for the condition:
在这里，没有分支，因此没有分歧，但我们确实为每个线程做了额外的工作，因为两个函数仍然运行，包括对于给定情况不需要的那个。因此，只有当额外工作比分歧成本更便宜时，谓词执行才值得。
Here, there is no branching, so no divergence, but we do extra work per thread since both functions still run, including the one not needed for a given case. As such, predication is only worthwhile if the extra work is cheaper than the divergence cost.
在实践中，您应该分析这些优化的效果。具体来说，如果使用谓词执行，检查"predicated-off threads"和整体指令计数等指标。像Nsight Compute这样的分析器会显示谓词执行是否正在减少线程束停顿--或者是否正在执行降低有效吞吐量的不必要工作。
In practice, you should profile the effects of these optimizations. Specifically, check metrics like "predicated-off threads" and overall instruction count if using predication. Profilers like Nsight Compute will show if predication is reducing warp stalls-or if it's executing unnecessary work that lowers effective throughput.
通常，谓词执行适用于非常短、简单的分支，每个只有几条指令--特别是如果许多线程束会由于给定条件而分歧。因此，如果分支涉及大量工作--或者只有少数线程分歧--谓词执行实际上可能更糟。请谨慎使用此技术--并且仅用于小型条件工作负载。
Generally, predication is good for very short, simple branches with only a few instructions each-especially if many warps would diverge due to a given condition. As such, if branches involve heavy work-or only a few threads diverge-predication can actually be worse. Use this technique judiciously-and only for small conditional workloads.
CUDA编译器通常会自动为您应用简单的谓词执行，但对于性能关键的内核，仍然值得分析和调优。
CUDA compilers will often apply simple predication automatically for you, but it's still worth profiling and tuning for performance-critical kernels.
简而言之，最小化线程束分歧需要仔细的算法和数据组织。尽可能重构您的问题，使每个线程束遵循统一的执行路径。这可能意味着拆分内核以在单独的启动中处理不同情况，重新排列数据使每个线程束处理相似的项目，或将分歧条件从内部循环中拉出来。
In short, minimizing warp divergence requires careful algorithm and data organization. Whenever possible, restructure your problem so that each warp follows a uniform execution path. This might mean splitting the kernel to handle different cases in separate launches, rearranging the data so that each warp processes similar items, or pulling divergent conditions out of inner loops.
当分歧不可避免时（例如，具有固有变化的每线程工作的树或图工作负载），保持分歧部分尽可能短和罕见，以便线程束快速重新汇聚。您还可以利用线程束级内建函数如__any_sync、__ballot_sync和其他投票原语让线程束一起在条件上达成一致。您还可以尝试将工作压缩到更少的通道中，而不是让所有32个通道分支。
When divergence is unavoidable (e.g., tree or graph workloads with inherently varying per-thread work), keep the divergent sections as short and rare as possible so that warps reconverge quickly. You can also leverage warp-level intrinsics like __any_sync, __ballot_sync, and other vote primitives to let a warp agree on conditions together. You can also try compressing work into fewer lanes instead of having all 32 lanes branch.
总之，最小化线程束分歧需要仔细的算法和数据组织。在可能的情况下，重构您的问题，使每个线程束遵循统一的执行路径。这可能意味着拆分内核以在单独的启动中处理不同情况，重新排列数据使每个线程束处理相似的项目，或将分歧条件从内部循环中拉出来。当分歧不可避免时（例如，具有固有变化的每线程工作的树或图工作负载），保持分歧代码尽可能短，并考虑使用谓词执行或线程束投票内建函数来减少其影响。
In summary, minimizing warp divergence requires careful algorithm and data organization. Whenever possible, restructure your problem so that each warp follows a uniform execution path. This might mean splitting the kernel to handle different cases in separate launches, rearranging the data so that each warp processes similar items, or pulling divergent conditions out of inner loops. When divergence is unavoidable (e.g., tree or graph workloads with inherently varying per-thread work), keep the divergent code as short as possible and consider using predication or warp-vote intrinsics to reduce its impact.
如果您看到低线程束执行效率，检查内核的分支。将条件逻辑重构到单独的内核中或使用线程束级原语（例如ballot sync）可能是有益的，以便更有效地处理分歧，我们接下来将介绍。
If you see low warp execution efficiency, check your kernel's branches. It might be beneficial to refactor conditional logic into separate kernels or use warp-level primitives (e.g., ballot sync) to handle divergence more efficiently, which we'll cover next.
使用谓词执行最小化分歧 (Using Predication to Minimize Divergence)¶
让我们展示一个使用谓词执行来分析和消除线程束分歧和分支的例子。考虑一个使用if (x[i] > 0) y[i] = x[i]; else y[i] = 0;对数组进行阈值处理的内核。如果一半值是正数，一半不是，那么大多数线程束会有一些线程采取if分支，一些采取else分支。线程束将首先为条件为真的线程执行if分支指令。它只是屏蔽掉其他线程。然后它再次运行并为剩余线程执行else分支。所以这花了两组指令来完成一组逻辑操作。这将效率降低了50%。
Let's show an example using predication to profile and eliminate warp divergence and branches. Consider a kernel that thresholds an array using if (x[i] > 0) y[i] = x[i]; else y[i] = 0;. If half the values are positive and half are not, then most warps will have some threads take the if branch and some take the else branch. The warp will first execute the if branch instructions for the threads when the condition is true. It will simply mask out the other threads. It then runs again and executes the else branch for the remaining threads. So it took two sets of instructions to do one logical operation. This reduces efficiency by 50%.
之前示例（CUDA C++）：在这个例子中，如果线程束中的某些线程满足X[i] > threshold而其他不满足，线程束将分歧。这是一个会导致线程束分支分歧的内核的明显例子：
Before example (CUDA C++). In this example, if some threads in a warp satisfy X[i] > threshold and others do not, the warp will diverge. This is a clear example of a kernel that will cause warp branch divergence:
// threshold_naive.cu
__global__ void threshold_naive(const float* X, float* Y,
float threshold, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
if (X[i] > threshold) {
Y[i] = X[i]; // branch 1
} else {
Y[i] = 0.0f; // branch 2
}
}
}
这导致线程束执行两次，一个接一个串行。一次执行将为一个线程子集运行赋值Y[i]=X[i]，另一次执行将为另一个子集运行Y[i]=0。线程束执行效率将很低--如果一半线程采取每条路径，大约为50%。
This results in the warp executing twice, one after the other serially. One execution will run the assignment Y[i]=X[i] for one subset of threads, and another execution will run Y[i]=0 for the other subset. The warp execution efficiency will be low-around 50% if half the threads take each path.
之后示例（CUDA C++）：这是一个使用谓词执行的减少分歧方法。在这个版本中，我们使用三元运算符，编译器可能将其转换为谓词移动指令（基于条件的SEL/MOV）而不是实际分支：
After example (CUDA C++). This is a divergence-reducing approach using predication. In this version, we use a ternary operator, which the compiler might convert to a predicated move instruction (SEL/MOV based on condition) instead of an actual branch:
// threshold_predicated.cu
__global__ void threshold_predicated(const float* X, float* Y,
float threshold, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
float x = X[i];
// Use a conditional move or multiplication by boolean
float val = (x > threshold) ? x : 0.0f;
Y[i] = val;
}
}
在这种情况下，所有线程执行相同的计算val然后存储它的指令序列。这避免了线程束分歧，因为控制流在线程束间是统一的。但条件为假的线程只会设置val=0。结果是线程束执行效率保持高水平，因为所有线程遵循一条路径。
In simple cases, the CUDA NVCC compiler likely generated a predicated move for the ternary operator, as expected. In PTX/assembly, you would see the PTX @p predicate syntax to guard the write without splitting into separate warp paths.
之后示例（PyTorch）
After example (PyTorch). In PyTorch, the threshold operation can be done with vectorized operations as follows:
在PyTorch中，阈值操作可以用向量化操作完成，如下所示：
# threshold_op.py
import torch
X = torch.randn(N, device='cuda')
Y = torch.maximum(X, torch.zeros_like(X)) # equivalent to Y = X > 0 ? X : 0
torch.cuda.synchronize()
The torch.maximum with 0 will execute on the GPU without branching since it uses elementwise max, which is implemented in a vectorized manner. Libraries like PyTorch ensure these elementwise ops are divergence-free at the warp level by using predication or bitwise tricks under the hood:
带有0的torch.maximum将在GPU上执行而不分支，因为它使用逐元素max，这是以向量化方式实现的。像PyTorch这样的库通过在底层使用谓词执行或位运算技巧确保这些逐元素操作在线程束级别是无分歧的：
# jit_threshold_op.py
import torch
# In PyTorch, we can compile and fuse this operation
# for even higher throughput
@torch.compile()
def threshold_op(X):
return torch.maximum(X, torch.zeros_like(X))
X = torch.randn(N, device='cuda')
Y = threshold_op(X)
torch.cuda.synchronize()
正如您将在第13章和第14章看到的，torch.compile使用TorchInductor将许多逐元素操作融合到内核中，这更好地匹配了手动CUDA C++示例的性能。但是，torch.compile不保证最佳占用率或ILP，因此您可能仍需要进一步分析和调优性能。
As you will see in Chapters 13 and 14, torch.compile uses TorchInductor to fuse many pointwise operations into kernels, which better match the performance of the manual CUDA C++ example. However, torch.compile does not guarantee optimal occupancy or ILP, so you may still need to profile and tune performance further.
在底层，这些类型的逐元素函数编译为GPU上的单指令多数据（SIMD）风格的谓词执行或位选择操作。这确保线程束中的每个线程遵循相同的指令流，并避免分歧控制流的串行化惩罚。总的来说，减少线程束分歧提高了线程束的效率，并在分歧是主要问题的情况下产生了显著的加速。表8-3显示了通过用谓词操作替换分支来减少线程束分歧的结果。
Under the hood, these types of elementwise functions compile down to single-instruction, multiple data (SIMD)‐style predication or bitwise select operations on GPUs. This ensures that every thread in a warp follows the same instruction stream and avoids the serialization penalties of divergent control flow. Overall, reducing warp divergence improves the warp's efficiency and yields substantial speedups in which divergence was a major issue. Table 8-3 shows the results of reducing warp divergence by replacing branches with predicated operations.
表8-3. 消除线程束分歧的分析结果
| 指标 | 之前 | 之后 |
|---|---|---|
| 内核执行时间 | 30 ms | 15 ms（-50%） |
| 动态指令计数 | 600 M条指令 | 300 M条指令（-50%） |
| 平均线程束分支解析停顿延迟 | 200周期 | 100周期（-50%） |
| 线程束执行效率 | 50% | 99% |
| 被谓词关闭线程（%） | 50% | 0% |
Note: The numeric values in all metrics tables are illustrative to explain the concepts. For actual benchmark results on different GPU architectures, see the GitHub repository.
通过用单个谓词max()操作替换双路径分支，我们在所有关键Nsight Compute指标上看到了显著改善。gpu__time_elapsed.avg测量的内核执行时间从30 ms下降到15 ms，有效地使吞吐量翻倍。线程束执行效率攀升至近100%，因为每个线程束中的所有线程都保持活跃进行有用的工作。分析器进一步报告0%的被谓词关闭比率，表明在谓词max()方法下没有通道被屏蔽掉。这确认几乎所有的重新汇聚开销已被消除。
By replacing the two‐path branches with a single predicated max() operation, we saw dramatic improvements across all key Nsight Compute metrics. The kernel execution time measured by gpu__time_elapsed.avg dropped from 30 ms to 15 ms, effectively doubling throughput. Warp execution efficiency climbed to nearly 100% since all threads in each warp stayed active with useful work. And the profiler further reports a 0% predicated-off ratio, indicating that no lanes were ever masked off under the predicated max() approach. This confirms that almost all reconvergence overhead has been eliminated.
同时，动态指令计数（smsp__inst_executed.sum）下降了50%，从6亿条指令降至3亿条指令，因为每个线程束不再花费周期串行执行分支的两边。平均线程束分支解析停顿延迟（smsp__average_warp_latency_issue_stalled_branch_resolving）也减半，从200周期降至100周期。
At the same time, the dynamic instruction count (smsp__inst_executed.sum) dropped by 50%, from 600 million instructions to 300 million instructions, since each warp no longer spends cycles serially executing both sides of branches. The average warp branch resolution stall latency (smsp__average_warp_latency_issue_stalled_branch_resolving) also halved, from 200 cycles to 100 cycles.
如果您的内核包含多个分歧点，或者每个分歧分支承载更重的工作，移除这些分支可以复合这些收益--可能每个移除的分支产生超过2倍的加速。
If your kernel contains multiple divergence points or if each divergent branch carries heavier work, removing those branches can compound these gains-potentially yielding more than a 2x speedup per branch eliminated.
请注意，谓词执行不是免费的。某些线程仍然计算值（例如val = x），它们的结果从未被使用。如果每个分支承载大量工作，为每个线程计算两边实际上可能比轻微分歧成本更高。
Note that predication isn't free. Some threads still compute values (e.g., val = x) and their results are never used. If each branch carries substantial work, computing both sides for every thread can actually cost more than a mild divergence.
在像我们的阈值示例这样的简单情况下，谓词执行获胜，但您应该始终进行基准测试。在Nsight Compute的线程束执行效率指标下尝试基于分支和谓词执行的版本。如果效率低且每个分支很轻，谓词执行可能会有帮助。
In simple cases like our threshold example, predication wins, but you should always benchmark. Try both branch-based and predicated versions under Nsight Compute's warp execution efficiency metric. If efficiency is low and each branch is light, predication will likely help.
如果分支很重，允许一些分歧--或使用单独的内核--可能是更好的路径。最终，最小化分歧对SIMT性能至关重要，因此尽可能构建算法以保持线程束在单一路径上，并将任何剩余的分歧逻辑隔离到其自己的内核或线程束大小的区域中。
If branches are heavy, allowing some divergence-or using separate kernels-may be the better path. Ultimately, minimizing divergence is critical for SIMT performance, so structure your algorithms to keep warps on a single path whenever possible and isolate any remaining divergent logic into its own kernel or warp-sized region.
CUDA编译器通常会自动为您应用简单的谓词执行，但对于性能关键的内核，仍然值得手工调优和分析。
CUDA compilers will often apply simple predication automatically for you, but it's still worth hand-tuning and profiling for performance-critical kernels.
使用线程束内建函数进行高效的线程束内通信 (Efficient Intrawarp Communication with Warp Intrinsics)¶
当线程束内的线程必须共享数据来计算归约/聚合时，例如，您可以使用线程束洗牌内建函数，包括__shfl_sync和__shfl_down_sync（在第6章介绍），直接通过寄存器交换值。这与将它们暂存在共享内存中并调用__syncthreads()形成对比，后者会对性能产生负面影响。与产生额外L1/L2流量并需要全块屏障的共享内存方法不同，洗牌只在寄存器之间移动数据，并在每条指令隐式同步线程束中的32个通道。这只增加几个时钟周期的延迟。
When threads within a warp must share data to compute a reduction/aggregation, for instance, you can use warp shuffle intrinsics, including __shfl_sync and __shfl_down_sync (introduced in Chapter 6) to exchange values directly through registers. This is in contrast to staging them in shared memory and calling __syncthreads(), which negatively impacts performance. Unlike the shared memory approach that generates extra L1/L2 traffic and requires a full-block barrier, shuffles move data only between registers and implicitly synchronize the 32 lanes in a warp with each instruction. This adds only a few clock cycles of latency.
有关完整讨论，包括性能比较、代码示例以及如何实现线程束级归约或前缀和，请参见第6章。
For a full discussion, including performance comparisons, code examples, and how to implement warp-level reductions or prefix sums, see Chapter 6.
如果您的协作跨越多个线程束，您仍然必须使用块级或网格级同步（例如__syncthreads()）或使用协作组（在第10章介绍）。但对于纯线程束内通信，洗牌几乎总是最快的选择。
If your collaboration spans multiple warps, you must still use block-level or grid-level synchronization (e.g., __syncthreads()) or use cooperative groups (introduced in Chapter 10). But for pure intrawarp communication, shuffles are almost always the fastest option.
简而言之，如果您的算法要求同一线程束中的线程共享中间结果或协调计算，请使用__shfl_sync和相关线程束级选项。只有在协作真正跨越多个线程束时才使用块级共享内存和__syncthreads()。
In short, if your algorithm requires threads in the same warp to share intermediate results or coordinate computations, use __shfl_sync and related warp-level options. Only use block-level shared memory and __syncthreads() when collaboration truly spans multiple warps.
PyTorch线程束级效率的考虑 (PyTorch Considerations for Warp-Level Efficiency)¶
在纯Python级别，PyTorch不暴露线程束内建函数或允许您直接控制线程束级执行。然而，作为PyTorch的最终用户，您间接受益于这些优化，因为PyTorch的许多内部CUDA内核使用线程束级技巧。例如，torch.sum(x, dim=0)的实现，它对张量的一个维度求和，当被归约的维度很小时（例如在32个元素内），通常会使用线程束级归约。在这种情况下，每个线程束的线程使用洗牌指令协作产生部分和，而不分歧或使用共享内存。此逻辑在库内部实现，因此您不必自己实现。
At the pure Python level, PyTorch doesn't expose warp intrinsics or allow you to directly control warp-level execution. However, as an end user of PyTorch, you indirectly benefit from these optimizations because many of PyTorch's internal CUDA kernels use warp-level tricks. For instance, the implementation of torch.sum(x, dim=0), which sums across a dimension of a tensor, will often use warp-level reductions when the dimension being reduced is small (e.g., within 32 elements). In such cases, each warp's threads use shuffle instructions to cooperatively produce a partial sum without diverging or using shared memory. This logic is implemented internally in the library, so you don't have to implement it yourself.
如果您编写自定义CUDA内核作为PyTorch扩展，您可以--也应该--在可能时在设备代码中使用这些线程束级优化。例如，如果您的操作需要对线程束内的值求和或在线程之间交换数据，优先使用像__shfl_sync这样的内建函数而不是朴素的共享内存方法，以增加速度并减少开销。还要注意任何自定义CUDA代码中的线程束分歧。
If you write custom CUDA kernels as PyTorch extensions, you can-and should-use these warp-level optimizations in your device code when possible. For example, if your operation needs to sum values within a warp or exchange data between threads, prefer using intrinsics like __shfl_sync over naive shared memory approaches to increase speed and reduce overhead. Also be mindful of warp divergence in any custom CUDA code.
如果您发现自己使用每元素if语句，您可以改为使用torch.where在整个张量上将条件表达为张量级操作。这允许库更有效地处理条件--可能以向量化和线程束一致的方式。
If you find yourself using per-element if statements, you can instead use torch.where to express the condition as a tensor-level operation across the whole tensor. This allows the library to handle the condition more efficiently-potentially in a vectorized and warp-consistent manner.
PyTorch的JIT编译器和图执行器（在第10章讨论）可以融合逐元素操作。这通过每个内核做更多工作--以及每移动字节内存做更多工作来增加线程束效率和算术强度。
PyTorch's JIT compiler and graph executor (discussed in Chapter 10) can fuse element-wise operations. This increases warp efficiency and arithmetic intensity by doing more work per kernel-and more work per byte of memory moved.
当多个操作融合到一个内核中时，线程束在单次传递中处理更多工作。虽然内核融合不能神奇地消除算法固有的分歧，但它确实意味着每个线程束做更多有用的工作。因此，任何分歧开销被分摊到更多计算中。
When multiple operations are fused into one kernel, warps process more work in a single pass. While kernel fusion can't magically eliminate divergence inherent to the algorithm, it does mean each warp does more useful work. As such, any divergence overhead is amortized across more computation.
例如，如果您有一个修正线性单元（ReLU）操作后跟一个abs()操作，融合内核可以在一次传递中处理两者。因此，即使有分支，它也在每个线程束一起处理两个操作。
For example, if you have a rectified linear unit (ReLU) operation followed by an abs() operation, a fused kernel can handle both in one pass. So even with branching, it processes both operations together in each warp.
作为CUDA程序员，您应该避免线程束内分歧，并在线程束内的任何集体操作中使用线程束内建函数。这确保线程束中的所有32个线程保持忙碌并行地做有用工作。
As a CUDA programmer, you should avoid intrawarp divergence and use warp intrinsics for any collective operations within a warp. This ensures all 32 threads in a warp stay busy doing useful work in parallel.
作为PyTorch用户，请注意库已经在底层为您做这些。您的角色是以允许利用这些优化的方式编写计算。例如，使用内置张量操作（用CUDA C++优化）而不是编写带有每元素条件的显式Python循环。如果您使用自定义CUDA扩展，应用我们讨论的相同最佳实践。
As a PyTorch user, be aware that the library already does this for you under the hood. Your role is to write computations in a way that allows these optimizations to be leveraged. For example, use built-in tensor operations (optimized with CUDA C++) rather than writing explicit Python loops with per-element conditions. If you use custom CUDA extensions, apply the same best practices we discussed.
每个线程束的32个通道协同工作是理想状态。通过最小化分歧并为任何需要的通信使用快速内建函数，您确保线程束执行效率。这些技术通常产生适度但有意义的加速--可能是1.1倍到1.3倍的改善。这在紧凑的内核中可能是显著的。
Having all 32 lanes of a warp work together is the ideal state. By minimizing divergence and using fast intrinsics for any needed communication, you ensure warp execution efficiency. These techniques typically yield modest but meaningful speedups-likely 1.1x to 1.3x improvements. This can be significant in tight kernels.
这些优化也为指令级并行奠定了基础，它允许每个线程做更多有用的工作。让我们接下来介绍这个。
These optimizations also lay the groundwork for instruction-level parallelism, which allows each thread to do more useful work. Let's cover that next.
暴露指令级并行性 (Exposing Instruction-Level Parallelism)¶
正如我们在占用率讨论中看到的，并发运行许多线程束让SM的调度器从任何在长延迟操作（如全局内存加载）上停顿的线程束切换开。除了启动足够的线程外，我们还可以在每个线程束内利用ILP，使单个线程束不需要等待一条指令完成后再发出下一条。
As we saw in the occupancy discussion, running many warps concurrently lets the SM's scheduler switch away from any warp that stalls on a long‐ latency operation such as a global‐memory load. In addition to launching enough threads, we can also exploit ILP within each warp so that a single warp does not need to wait for one instruction to complete before issuing the next. You can rearrange or unroll your code so that each thread issues multiple independent operations (e.g., memory loads and arithmetic instructions) before any result is consumed. This way, the GPU keeps its execution units busy while earlier instructions are still pending.
您可以重新排列或展开代码，使每个线程在消耗结果之前发出多个独立操作（例如内存加载和算术指令）。这样，GPU在早期指令仍在待处理时保持其执行单元忙碌。
You can rearrange or unroll your code so that each thread issues multiple independent operations (e.g., memory loads and arithmetic instructions) before any result is consumed. This way, the GPU keeps its execution units busy while earlier instructions are still pending.
利用ILP允许单个线程束背靠背发出某些独立指令，这改善了延迟隐藏。例如，一个线程可能加载数据，然后在该加载完成时执行不相关的算术。图8-8显示了每个周期重叠多条指令的示例。
Figure 8-8. Overlapping with ILP
通过展开循环体，您将曾经每次迭代的"加载→乘法→存储"转变为在循环返回之前加载和乘法多个元素的序列。例如，代替：
By unrolling your loop body, you turn what was once "load → multiply → store" each iteration into a sequence that loads and multiplies multiple elements before looping back. For example, instead of:
for (int i = 0; i < N; ++i) {
float ai = a[i]; // load
float bi = b[i]; // load
sum += ai * bi; // multiply after both loads complete
}
您可以展开，使每个循环迭代发出两个独立的乘法操作而不是一个，如下一代码块所示。这给硬件提供了在第一个乘法仍在进行时执行第二个乘法的机会--或者在等待其操作数从内存加载时。这适当地重叠并隐藏延迟。结果是更高的每周期指令数（IPC）和更好的延迟隐藏，如下所示通过并行计算两个乘法ai0*bi0和ai1*bi1：
By unrolling your loop body, you turn what was once "load → multiply → store" each iteration into a sequence that loads and multiplies multiple elements before looping back. For example, instead of: for (int i = 0; i < N; ++i) { float ai = a[i]; // load float bi = b[i]; // load sum += ai * bi; // multiply after both loads complete } You can unroll such that each loop iteration issues two independent multiply operations instead of one, as shown in the next code block. This gives the hardware an opportunity to execute the second multiply while the first is still in progress-or while waiting for its operands to load from memory. This appropriately overlaps and hides latency. The result is higher instructions per cycle (IPC) and better latency hiding, as shown next by computing two multiplications ai0*bi0 and ai1*bi1 in parallel:
for (int i = 0; i + 1 < N; i += 2) {
float ai0 = a[i]; // load a[i]
float bi0 = b[i]; // load b[i]
float ai1 = a[i + 1]; // load a[i+1]
float bi1 = b[i + 1]; // load b[i+1]
sum += ai0 * bi0 + ai1 * bi1;
}
在这里，我们在做任何乘法之前加载成对的值（a[i]、b[i]和a[i+1]、b[i+1]）。展开暴露了每个循环迭代两个独立的乘法指令。这样，GPU可以重叠这些算术操作并隐藏每个操作数加载的延迟。
ILP不直接改变算术强度（每字节FLOPS）。然而，它通过重叠计算与内存或计算到计算的依赖关系来提高整体吞吐量（FLOPS）。我们显式展开以增加每线程的独立工作并帮助双发射。换句话说，ILP通过时间复用独立工作来提高吞吐量--而不是通过做额外工作。
ILP doesn't directly change arithmetic intensity (FLOPS per byte). However, it improves overall throughput (FLOPS) by overlapping computation with memory or compute-to-compute dependencies. We explicitly unroll to increase per-thread independent work and help dual issue. In other words, ILP improves throughput by time-multiplexing independent work-not by doing extra work.
一个常见的误解是增加ILP会执行更多操作。然而，它不会--它只是让GPU的多个功能单元保持忙碌。ILP只有在您的程序由于延迟而有空闲发出槽时才有帮助。如果您的内核已经在每个周期完全利用执行单元（例如，没有停顿的紧密计算绑定循环），增加ILP实际上不会提高性能。
A common misconception is that increasing ILP will perform more operations. However, it doesn't-it just keeps the GPU's multiple functional units busy. ILP helps only if your program has idle issue slots due to latency. If your kernel is already fully utilizing the execution units on every cycle (e.g., a tight compute-bound loop with no stalls), increasing ILP won't actually increase performance.
要鼓励编译器增加或暴露ILP，您可以在短循环上使用#pragma unroll或调整-maxrregcount以允许编译器分配更多寄存器来保存中间值。这显式确认您希望编译器增加寄存器使用并可能降低占用率。这样，ILP补充了占用率。高占用率确保有许多线程束可以在一个停顿时切换到，ILP确保单个线程束可以发出尽可能多的独立指令来填充流水线并隐藏延迟。这很像CPU中的超标量和乱序执行，只是它由编译器对CUDA代码的调度显式指导。
To encourage the compiler to increase, or expose, ILP, you can use #pragma unroll on short loops or tune ‐maxrregcount to allow the compiler to allocate more registers for holding intermediate values. This explicitly acknowledges that you want the compiler to increase register usage and potentially reduce occupancy. In this way, ILP complements occupancy. High occupancy ensures there are many warps to switch to when one stalls, and ILP makes sure that a single warp can issue as many independent instructions as possible to fill the pipeline and hide latency. This is much like superscalar and out-of-order execution in CPUs, except that it's guided explicitly by the compiler's scheduling of your CUDA code.
线程束调度和双发射指令 (Warp Scheduling and Dual Issue Instructions)¶
每个SM包含多个线程束调度器。例如，Blackwell SM有四个线程束调度器。每个调度器每个周期可以发出最多两条独立指令，如图8-9所示。
Each SM contains multiple warp schedulers. For example, Blackwell SMs have four warp schedulers. Each scheduler can issue up to two independent instructions per cycle, as shown in Figure 8-9.
Figure 8-9. Each warp scheduler can dispatch up to two instructions per cycle called "dual issue"
在实践中，这意味着单个线程束可以在连续周期内重叠多个独立操作--有时甚至在同一周期内，如果它们针对不同的流水线，如一个数学、一个内存或一个特殊功能单元（SFU）流水线。这意味着多条指令可以同时在进行中并通过流水线推进。这种重叠就是我们之前讨论的ILP。
In practice, this means that a single warp can overlap multiple independent operations across successive cycles-and sometimes even in the same cycle if they target different pipelines such as one math, one memory, or one special-function unit (SFU) pipeline. This means that multiple instructions can be in flight and progressing through the pipeline simultaneously. This overlap is the ILP that we discussed earlier.
在Blackwell上，FP32和INT32流水线已合并为一组统一的CUDA核心。因此，它们在给定时钟周期内只能执行FP32或INT32指令，而不能同时执行两者。每个核心每个时钟周期必须选择一种数据类型。因此，任何依赖在同一周期内双发射INT和FP的ILP将不再有效，因为Blackwell在每个核心上每个周期必须选择其中之一。因此，混合INT32和FP32指令流不再受益于单核心双发射。混合流应该利用线程束/SM并发，而不是依赖每核心双发射。因此，您应该优先使用ILP让两个独立的数学和内存指令（或两个独立的数学指令如FP32和FP16）在进行中双发射。这将增加指令发射效率。
On Blackwell, the FP32 and INT32 pipelines have been merged into a single set of unified CUDA cores. As such, they can only execute either FP32 or INT32 instructions in a given cycle-not both at once. Each core must pick one data type each cycle. As a result, any ILP that depended on dual-issuing an INT and an FP in the same cycle will no longer work since Blackwell must choose one or the other per cycle on each core. As such, mixed INT32 and FP32 instruction streams no longer benefit from dual issue on a single core. Mixed streams should instead exploit warp/SM concurrency-and not rely on per-core dual issue. As such, you should prefer dual issue with two independent math and memory instructions (or two independent math instructions such as FP32 and FP16) in flight using ILP. This will increase instruction-issue efficiency.
ILP不要求两条指令在同一周期物理触发。相反，它确保当一条指令（例如长延迟全局内存加载）仍在待处理时，线程束可以在下一个周期立即发出另一条独立指令（例如算术操作）。由于这种重叠，当数据从内存返回时，算术工作已经取得进展。这有效地隐藏了加载的延迟。在Blackwell上，由于流水线技术，每个线程束可以有多条指令在进行中，调度器在连续周期内向不同的执行单元发出。在实践中，ILP表现为稳定的独立发出流，而不是每个线程束固定数量的并发指令。
ILP does not require that two instructions physically fire in the same cycle. Instead, it makes sure that while one instruction (e.g., long-latency global memory load) is still pending, the warp can immediately issue another independent instruction (e.g., arithmetic operation) on the next cycle. As a result of this overlap, by the time the data returns from memory, the arithmetic work has already made progress. This effectively hides the load's latency. On Blackwell, multiple instructions can be in flight per warp thanks to pipelining, with the scheduler issuing to different execution units on successive cycles. In practice, ILP manifests as a steady stream of independent issues, not a fixed number of concurrent instructions per warp.
考虑内核中每个线程加载两个数组元素--然后将每个乘以不同的常数再组合结果的场景。第一个元素的乘法操作可以在第二个元素的加载操作仍在进行时执行。
Consider a kernel where each thread loads two array elements-and then multiplies each by a different constant before combining the results. The multiplication for the first element can execute while the load for the second element is still in progress.
这种重叠模式保持线程束的执行单元忙碌而不是等待。这提高了整体吞吐量而不改变每个元素发生多少加载和乘法。这是增加ILP的具体示例：
This overlapping pattern keeps the warp's execution units busy instead of waiting. This improves overall throughput without changing how many loads and multiplies happen per element. This is a concrete example of increasing ILP:
// Launch configuration: e.g., 256 threads per block, enough
// blocks to cover N
__global__ void independentOps(const float *a, const float *b,
float *out, int N) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N) {
float x = a[i];
float y = b[i];
// Two independent operations (no dependency between u and v):
float u = x * x;
float v = y * y;
// Dependent operation that uses both results:
float sum = u + v;
out[i] = sqrtf(sum);
}
}
在这个内核中，每个线程从不同数组加载两个值x和y并计算两个结果u和v。u = x*x和v = y*y的计算彼此独立，因为它们都不使用另一个的结果。
In this kernel, each thread loads two values x and y from different arrays and computes two results u and v. The calculations of u = x*x and v = y*y are independent of each other since neither uses the result of the other.
以这种方式构建代码，而不是sum = x*x + y*y，给编译器提供了安排指令序列以暴露ILP并增加性能的机会。具体来说，它可以发出u = x*x的指令，然后在下一个周期在u完成之前发出v = y*y。
Structuring the code this way, instead of sum = x*x + y*y, gives the compiler an opportunity to schedule the instruction sequence to expose ILP and increase performance. Specifically, it can issue the instruction for u = x*x and then issue v = y*y on the next cycle before u completes.
当第一个乘法正在运行--或等待流水线槽时，第二个乘法可以在另一个算术单元中并行运行。当需要做sum = u + v时，可能两个乘法都已经完成。因此线程束有多条指令在进行中，这有助于隐藏这两个乘法操作的延迟。
While the first multiply is running-or waiting for a pipeline slot-the second multiply can run in parallel in another arithmetic unit. When it's time to do sum = u + v, both multiplies may have already completed. So the warp has multiple instructions in flight, which helps hide the latency of those two multiply operations.
相比之下，考虑以下代码，其中v的计算被放在使用u之后。在这个场景中，我们会创建一个限制ILP的不必要依赖链：
In contrast, consider the following code where the computation of v is placed after the use of u. In this scenario, we would create an unnecessary dependency chain that limits ILP:
内核计算u = x * x并立即使用该结果形成temp = u + 1.0f。只有在计算temp之后才发出v = y * y，这不依赖于u或temp。因此，它放在temp之后迫使GPU等到u + 1.0f指令完成才能发出v的乘法。
The kernel computes u = x * x and immediately uses that result to form temp = u + 1.0f. Only after computing temp does it issue v = y * y, which doesn't depend on u or temp. As such, it places after temp and forces the GPU to wait until the u + 1.0f instruction completes before it can issue the multiply for v.
实际上，您创建了一个串行依赖链：首先是x * x → u，然后u + 1.0f → temp，只有然后y * y → v。在硬件正在执行（或等待）u + 1.0f操作期间，它不能开始y * y，即使v在数学上是独立的。这种排序人为地限制了指令级并行性，因为一个独立的乘法必须停在一个依赖指令之后。
In effect, you've created a serial dependency chain: first x * x → u, then u + 1.0f → temp, and only then y * y → v. During the time the hardware is executing (or waiting for) the u + 1.0f operation, it cannot start y * y, even though v is mathematically independent. This ordering artificially limits instruction-level parallelism because an independent multiply has to stall after a dependent instruction.
相比之下，考虑以下重新排序和优化的代码：
In contrast, consider the following reordered and optimized code:
// Optimized ordering with better ILP:
float u = x * x; // Independent multiply on x
float v = y * y; // Independent multiply on y (issues immediately after u)
float temp = u + 1.0f; // Only now use u
// ... use v or temp as needed ...
在这里，我们重新排序代码，使x * x和y * y背靠背发生（即float u = x * x; float v = y * y; float temp = u + 1.0f;）。在这种情况下，GPU可以在连续周期发出两个独立的乘法而不等待u + 1.0f加法完成。
Here, we reordered the code so that x * x and y * y happen back-to-back (i.e., float u = x * x; float v = y * y; float temp = u + 1.0f;). In this case, the GPU can issue two independent multiplies in consecutive cycles without waiting for the u + 1.0f addition to complete.
这意味着当x * x的结果可能仍在"待处理"流水线槽中时，硬件已经可以开始执行y * y。
This means that while the result of x * x may still be in a "pending" pipeline slot, the hardware can already start executing y * y.
只有当这两个乘法都计算出结果后，GPU才执行temp = u + 1.0f。
Only after both multiplies have computed their results does the GPU execute temp = u + 1.0f.
换句话说，您现在有两条乘法同时在进行。这填充了执行单元并隐藏了延迟。这种改进的ILP，尽快发出v = y * y，确保y上的乘法与x上的乘法和后面的加法重叠--而不是被迫等待。
In other words, you now have two multiplies in flight simultaneously. This fills the execution units and hides latency. This improved ILP, issuing v = y * y as soon as possible, ensures that the multiply on y overlaps with the multiply on x and the subsequent addition-instead of being forced to wait.
因此，线程束花费更少的周期空闲，这直接转化为更高的吞吐量。
As a result, the warp spends fewer cycles idle, which directly translates to higher throughput.
GPU依赖编译器，有时依赖程序员的提示来调度独立指令。在现代CUDA编译器上，像前面这样的简单模式通常被检测并最优调度。但您可以通过适当构建代码或使用pragma编译器指令来鼓励ILP。
GPUs rely on the compiler, and sometimes on programmer hints, to schedule independent instructions. On modern CUDA compilers, simple patterns like the preceding ones are often detected and optimally scheduled. But you can encourage ILP by structuring your code appropriately or using pragma compiler directives.
这个例子很简单，但它展示了一个想法，即如果可以的话，不要在线程内串行化独立工作。如果一个线程需要加载两个数组并对两者做数学运算，并行执行或交错执行比以串行方式一个接一个执行更好。
This example is simple, but it illustrates the idea that if you can, don't serialize independent work within a thread. If a thread needs to load two arrays and do math on both, executing in parallel or interleaved is better than doing it serially one after the other.
ILP和占用率 (ILP and Occupancy)¶
在现代GPU上，每个SM的最大驻留线程束是64个线程束或2,048个线程（2,048个线程 = 每线程束32个线程 x 64个线程束）。线程束调度器可以在依赖关系允许时每周期发出多条指令。INT32和FP32核心是统一的，在给定时钟内作为FP32或INT32运行，而不是同时运行两者。每个SM的寄存器文件是256 KB，即64K个32位寄存器。这个大寄存器文件对于在不溢出的情况下维持ILP很重要，我们稍后会看到。
On modern GPUs, the maximum resident warps per SM is 64 warps, or 2,048 threads (2,048 threads = 32 threads per warp x 64 warps). The warp schedulers can issue multiple instructions per cycle when dependencies permit. The INT32 and FP32 cores are unified and operate as either FP32 or INT32 in a given clock rather than both at once. The register file per SM is 256 KB, which is 64K 32-bit registers. This large register file is important for sustaining ILP without spilling, as we'll see in a bit.
如果您编写的内核几乎没有ILP，使得每个线程一次只发出一个操作，您通常需要约1,536个活跃线程（约48个线程束或75%占用率）来饱和SM的执行单元（这些不是硬限制而是近似值）。
If you write a kernel with almost no ILP, so that each thread issues only one operation at a time, you typically need about 1,536 active threads (about 48 warps or 75% occupancy) to saturate an SM's execution units (these are not hard limits but approximations).
相比之下，暴露即使是适量的ILP也可以显著降低所需的线程数。表8-4总结了ILP如何减少饱和Blackwell SM所需的线程/占用率。
In contrast, exposing even moderate ILP can significantly reduce the number of threads required. Table 8-4 summarizes how ILP reduces the threads/occupancy needed to saturate a Blackwell SM.
表8-4. ILP如何减少饱和Blackwell SM所需的线程/占用率
Table 8-4. How ILP reduces the threads/occupancy needed to saturate a Blackwell SM
| ILP（每线程独立操作数） | ~100%利用率所需的线程/SM | 占用率（2,048线程的百分比） |
|---|---|---|
| 1（无ILP） | ~1,536个线程（48个线程束） | 75% |
| 2 | ~1,024个线程（32个线程束） | 50% |
| 3 | ~768个线程（24个线程束） | 37.5% |
| 4 | ~512个线程（16个线程束） | 25% |
ILP (independent ops/thread) Threads/SM for ~100% utilization Occupancy (% of 2,048 threads) 1 (no ILP) ~1,536 threads (48 warps) 75% 2 ~1,024 threads (32 warps) 50% 3 ~768 threads (24 warps) 37.5% 4 ~512 threads (16 warps) 25%
在这里，双向ILP，即每个线程背靠背发出两个独立操作，通常以约1,024个活跃线程（约32个线程束或50%占用率）实现全吞吐量。这是因为每个线程束在有一个操作仍待处理时保持两个操作在进行中。
Here, two-way ILP, where each thread issues two independent operations back-to-back, typically achieves full throughput with about 1,024 active threads (about 32 warps or 50% occupancy). This is because each warp keeps two operations in flight while one operation is still pending.
另一方面，三向ILP可以用约768个线程（约24个线程束或37.5%占用率）饱和。四向ILP可能只需要512个线程（约16个线程束或25%占用率）来填充流水线。这是因为每个线程束本身做四个独立操作的工作。
On the other hand, three-way ILP can saturate with about 768 threads (about 24 warps or 37.5% occupancy). Four-way ILP might need only 512 threads (about 16 warps or 25% occupancy) to fill the pipeline. This is because each warp itself does the work of four independent operations.
每个线程更多在进行中的操作让每个线程束保持数学流水线忙碌--即使在较低的线程数下。
More in-flight operations per thread lets each warp keep the math pipelines busy-even at lower thread counts.
显然，增加ILP让您可以用更少的线程束实现峰值计算吞吐量。当您的工作负载无法启动大量线程--或者当您想释放片上资源用于其他任务时，这特别有价值。重要的是要注意，您可以暴露的ILP量有实际限制。尽管Blackwell SM可以保持每个线程束许多指令在进行中，但每个线程束只能容纳有限数量的待处理指令。除此之外，编译器必须调度独立指令而不超过可用寄存器或压倒GPU的指令解码带宽。解码带宽是指令获取和解码硬件每周期可以推送到执行单元的最大指令数。一旦达到这些限制，添加更多独立操作会产生收益递减。
Clearly, increasing ILP lets you achieve peak compute throughput with fewer warps. This is especially valuable when your workload cannot launch huge numbers of threads-or when you want to free up on-chip resources for other tasks. It's important to note that there is a practical limit on how much ILP you can expose. Even though Blackwell SMs can keep many instructions in flight per warp, each warp can hold only a finite number of pending instructions. On top of that, the compiler must schedule independent instructions without exceeding available registers or overwhelming the GPU's instruction decode bandwidth. Decode bandwidth is the maximum number of instructions the instruction fetch and decode hardware can push to the execution units per cycle. Once you hit those limits, adding more independent operations yields diminishing returns.
一旦您有足够的独立工作来保持发出槽和内存系统忙碌，进一步增加ILP几乎没有实际好处。推到五向或六向ILP产生的收益很少，甚至可能由于额外的寄存器压力和指令开销而损害性能。
Once you have enough independent work to keep issue slots and the memory system busy, further increasing ILP provides little practical benefit. Pushing to five-way or six-way ILP yields little gain and may even hurt performance due to extra register pressure and instruction overhead.
在Blackwell上，理想的ILP配置通常是每个线程两到四个独立操作。此时，其调度器和缓存通常已饱和。然而，确切的点取决于内核和工作负载。使用Nsight Compute发出和停顿指标来决定在哪里停止。
On Blackwell, the ideal ILP configuration is often between two and four independent operations per thread. At this point, its schedulers and caches are typically saturated. However, the exact point is kernel and workload dependent. Use Nsight Compute issue and stall metrics to decide where to stop.
循环展开、交错和编译器提示 (Loop Unrolling, Interleaving, and Compiler Hinting)¶
要在自己的内核中利用ILP，寻找每个线程在需要任何结果之前发出多个独立算术或内存指令的机会。常见模式包括：
To exploit ILP in your own kernels, look for opportunities in which each thread issues multiple independent arithmetic or memory instructions before any one result is needed. Common patterns include the following:
展开小循环
Unrolling small loops
展开循环允许每个线程使用单独的累加器寄存器执行N次累加（例如2x或4x）。例如，考虑以下循环：
Unrolling loops allows each thread to perform N accumulations (e.g., 2x or 4x) with separate accumulator registers. For example, consider the following loop:
float sum = 0.0f;
for (int k = 0; k < 4; ++k) {
float a = A[idx * 4 + k];
sum += a * w[k]; // dependent on load a
}
您可以重写它以利用ILP，如下所示。这种转换为四个独立的加载-乘法对--每个a0..a3一个--创造了可以快速连续发出的机会：
You can rewrite it to exploit ILP as shown next. This transformation creates four independent load-multiply pairs-one for each of a0..a3-that can be issued in quick succession:
float a0 = A[idx * 4 + 0];
float a1 = A[idx * 4 + 1];
float a2 = A[idx * 4 + 2];
float a3 = A[idx * 4 + 3];
float sum0 = a0 * w[0];
float sum1 = a1 * w[1];
float sum2 = a2 * w[2];
float sum3 = a3 * w[3];
float sum = sum0 + sum1 + sum2 + sum3;
在这里，所有四个加载首先发出。当每个加载完成时，其相应的乘法可以执行并与后续加载重叠计算以隐藏内存加载的延迟。换句话说，GPU交错这些操作，使得当一对在等待内存时，另一对可以乘法。这增加了ILP并隐藏了延迟。
Here, all four loads issue first. As each load completes, its corresponding multiply can execute and overlap computation with subsequent loads to hide the memory load latency. In other words, the GPU interleaves these operations so that while one pair is waiting on memory, another can multiply. This increases ILP and hides latency.
交错独立操作
Interleaving independent operations
您可以交错使用不同数据元素的操作，如下所示：
You can interleave operations on different data elements as follows:
float x = A[idx];
float y = B[idx];
// If you wrote float u = x * c; then float v = y * d;,
// the multiplies are independent.
float u = x * c; // can start as soon as x is ready
float v = y * d; // can start as soon as y is ready, overlapping with u
在这里，您通过不让一个依赖指令阻塞下一个独立指令等来保持每个周期的执行单元忙碌。在这个交错代码中，u = x*c和v = y*d是独立指令，可以并发执行。编译器将背靠背调度它们，使得一个乘法可以在另一个乘法仍在完成时进行。这就是ILP在行动。
Here, you keep the execution units busy each cycle by not letting one dependent instruction block the next independent instruction, etc. In this interleaved code, u = x*c and v = y*d are independent instructions that can execute concurrently. The compiler will schedule them back-to-back so that one multiply can proceed while the other is still completing. This is ILP in action.
使用编译器提示
Using compiler hints
小循环上的编译器提示如#pragma unroll可以帮助编译器显式创建多个算术链。如果需要，您可以调整-maxrregcount以允许编译器为展开的变量使用更多每线程寄存器。通过允许每线程更多资源，您正在以更少线程为代价增加ILP。这自然地用更高的ILP交换一些占用率。
Compiler hints like #pragma unroll on small loops can help the compiler explicitly create multiple arithmetic chains. If needed, you can tune -maxrregcount to allow the compiler to use more per-thread registers for unrolled variables. By allowing more resources per thread, you're increasing ILP at the expense of fewer threads. This naturally trades some occupancy for higher ILP.
简而言之，要在Blackwell深度流水线上隐藏延迟并最大化吞吐量，您应该结合高占用率（例如足够的线程束/线程来隐藏停顿）和高ILP（例如每线程多条独立指令）。在这里，您确保即使当一条指令在等待（例如内存加载或长延迟算术）时，线程束也可以立即发出其他指令。这将显著增加内核实现的FLOPS。
In short, to hide latency and maximize throughput on Blackwell's deep pipelines, you should combine high occupancy (e.g., enough warps/threads to hide stalls) with high ILP (e.g., multiple independent instructions per thread). Here, you ensure that even when one instruction is waiting (e.g., memory load or long-latency arithmetic), the warp can immediately issue other instructions. This will significantly increase the FLOPS achieved by the kernel.
在大型展开上，注意Nsight Compute中的指令获取和发出停顿。过度展开会使代码大小膨胀--或使解码和分发饱和--而不会进一步增加吞吐量。
On large unrolls, watch instruction-fetch and issue stalls in Nsight Compute. Excessive unrolling can bloat code size-or saturate decode and dispatch-without further throughput gains.
分析和缓解寄存器压力 (Profiling and Mitigating Register Pressure)¶
然而，请注意寄存器压力，因为ILP通常需要更多寄存器来保存多个独立值和结果。如果您展开太多--或添加太多并行计算--可能会将寄存器使用增加到占用率显著下降的程度。一如既往，找到正确的平衡是关键。CUDA编译器在自动展开循环方面做得很好。但如果您通过手动展开太多迭代（例如使用#pragma unroll）将ILP推得太远，您可能会注意到占用率下降。
However, be mindful of register pressure since ILP typically requires more registers to hold multiple independent values and results. If you unroll too much-or add too many parallel computations-you might increase register usage to the point that occupancy falls significantly. As always, finding the right balance is key. The CUDA compiler does a good job of unrolling loops automatically. But if you push ILP too far by manually unrolling too many iterations using #pragma unroll, for instance, you might notice occupancy drop.
在这种情况下，Nsight Compute将在其占用率分析中显示"Limited by Registers"。这通常是一个明确的信号，表明您走得太远了。如果您看到这个，您应该回调展开，减少独立累加器变量的数量，或放宽启动边界。
In this case, Nsight Compute will show "Limited by Registers" in its occupancy analysis. This is usually a clear signal that you've gone too far. If you see this, you should dial back unrolling, reduce the number of independent accumulator variables, or relax launch bounds.
具体来说，您可以减少所需的MinBlocksPerSM或增加MaxThreadsPerBlock，以便每个块可以使用更多寄存器而不溢出。换句话说，允许稍微降低的占用率以避免严重溢出。这样可以缓解寄存器压力并恢复占用率。
Specifically, you can reduce MinBlocksPerSM or increase MaxThreadsPerBlock so that each block can use more registers without spilling. In other words, allow slightly lower occupancy to avoid severe spilling. This can alleviate register pressure and recover occupancy.
从分析角度来看，如果ILP有帮助，您应该看到"Stall: Exec Dependency"百分比下降，因为线程束在先前指令上等待的时间更少。您应该看到更高的每周期指令数。
From a profiling perspective, if ILP is helping, you should see the "Stall: Exec Dependency" percentage go down because warps spend less time waiting on previous instructions. You should see higher instructions per cycle.
Nsight Compute报告发出效率或IPC指标。这应该随着ILP增加而上升。例如，当您展开时，您应该看到发出效率或每周期指令数指标上升。例如，如果您看到IPC指标从每个线程束调度器1.0增加到1.8，这表明线程束现在平均每周期发出接近两条指令而不是只有一条。然而，实际值取决于指令混合和架构调度约束。
Nsight Compute reports issue efficiency or IPC metrics. This should go up as ILP increases. For example, when you unroll, you should see the issue efficiency or instructions-per-cycle metric rise. For instance, if you see the IPC metric increase from 1.0 to 1.8 per warp scheduler, it indicates that warps are now issuing close to two instructions per cycle on average instead of just one. However, the actual value depends on instruction mix and architecture scheduling constraints.
您可能还会看到您可以比以前更低的占用率实现良好的性能。理想情况是您有足够的ILP来保持每个线程束忙碌--以及足够的线程束来保持SM忙碌。在这两者之间，您正在覆盖线程束内和跨线程束的延迟。
You might also see that you can achieve good performance at lower occupancy than before. The ideal is that you have enough ILP to keep each warp busy-and enough warps to keep the SM busy. Between the two, you're covering latency within warps and across warps.
至此，我们已经解决了线程束级并行性和指令级并行性。这是保持GPU计算单元忙碌的两种方法。接下来，我们转向测量和改善算术强度。这是关于最大化每次内存访问所做的有用工作。
At this point, we've addressed warp-level parallelism and instruction-level parallelism. These are two ways to keep the GPU's compute units busy. Next, we turn to measuring and improving arithmetic intensity. This is about maximizing useful work done per memory access.
关键要点 (Key Takeaways)¶
在本章中，您看到了如何通过将工作从慢速全局内存移动到更快的片上资源和计算单元来揭示和消除GPU内核瓶颈。通过遵循分析、诊断、优化和重新分析的循环，您可以将内核从利用不足或内存受限转变为计算饱和、高吞吐量的例程。这些技术将帮助利用GPU的全部能力：
In this chapter, you saw how to uncover and eliminate GPU kernel bottlenecks by moving work from slow global memory into faster on-chip resources and compute units. By following a cycle of profiling, diagnosing, optimizing, and reprofiling, you can transform kernels from underutilized or memory‐bound into compute‐saturated, high‐throughput routines. These techniques will help utilize the full power of your GPUs:
使用Nsight和torch.profiler进行分析：从Nsight Systems开始可视化端到端时间线并揭示CPU-GPU间隙、内核重叠和NVTX范围。然后使用Nsight Compute的停顿指标、屋顶线分析和占用率报告深入各个内核。在PyTorch代码中，您可以插入分析器检测（使用基于Kineto库的torch.profiler API）将Python级操作直接映射到GPU活动。请记住，PyTorch分析器的with_flops估计是基于有限操作符集的公式。它不读取GPU硬件计数器。对于硬件计数器，在感兴趣的内核上使用Nsight Compute。
Profile with Nsight and torch.profiler：Start with Nsight Systems to visualize the end-to-end timeline and reveal CPU-GPU gaps, kernel overlaps, and NVTX ranges. Then use Nsight Compute's stall metrics, roofline analysis, and occupancy reports to drill down into individual kernels. In PyTorch code, you can insert profiler instrumentation (using the Kineto-based torch.profiler API) to map Python-level operations directly to GPU activity. Remember that PyTorch profiler's with_flops estimates are formula-based on a limited operator set. It doesn't read GPU hardware counters. For hardware counters, use Nsight Compute on kernels of interest.
调优占用率以隐藏延迟：每个SM足够的线程束对于覆盖内存和指令延迟至关重要。您可以通过代码重构或编译器提示减少每线程寄存器/共享内存使用。您还应该选择适合SM资源的块大小（例如128-256个线程）。这样做将增加内核实现的占用率到最佳范围--约50%-75%--并停止利用不足（例如Blackwell上约32-48个线程束）。
Tune occupancy to hide latency：Enough warps per SM are critical for covering memory and instruction latency. You can reduce per-thread register/shared memory usage through code refactoring or compiler hints. You should also choose block sizes that fit SM resources (e.g., 128-256 threads). Doing so will increase the kernel's achieved occupancy to an optimal range-around 50%-75%-and stop underutilization (e.g., ~32-48 warps on Blackwell).
最小化线程束分歧以提高SIMT效率：因为线程束以锁步执行，任何分支分歧意味着被屏蔽的通道空闲。重构内核使线程束中的所有线程遵循相同的路径。使用谓词操作如三元数学和torch.maximum--或使用线程束投票内建函数来压缩活跃通道。这将提高线程束执行效率并减少串行执行。
Minimize warp divergence for SIMT efficiency：Because warps execute in lockstep, any branch divergence means masked lanes sit idle. Restructure kernels so that all threads in a warp follow the same path. Use predicated operations like ternary math and torch.maximum-or use warp vote intrinsics to compress active lanes. This will improve warp execution efficiency and reduce serial execution.
识别性能状态：每个内核落入四个桶之一：利用不足、延迟受限、内存受限或计算受限。这基于FLOPS、带宽使用和停顿原因。了解哪种状态适用将有助于指导确切的优化策略。例如，高Long Scoreboard停顿意味着工作负载是延迟受限的。
Identify performance states：Every kernel falls into one of four buckets: underutilized, latency bound, memory bound, or compute bound. This is based on FLOPS, bandwidth usage, and stall reasons. Knowing which state applies will help guide the exact optimization strategy. For example, high Long Scoreboard stalls mean the workload is latency bound.
暴露ILP：通过展开循环、打破依赖链和预取数据，您暴露指令级并行性，使每个线程可以每周期发出多条独立指令。这让双向或四向ILP将填充SM所需的线程束减少一半--减少执行依赖停顿、提高IPC并降低"Stall: Exec Dependency"指标。始终分析不同的ILP深度和线程计数以找到内核的最佳点。
Expose ILP：By unrolling loops, breaking dependency chains, and prefetching data, you expose instruction-level parallelism that lets each thread issue multiple independent instructions per cycle. This lets two-way or four-way ILP cut the warps needed to fill an SM in half-reducing execution dependency stalls, raising IPC, and lowering the "Stall: Exec Dependency" metric. Always profile different ILP depths and thread counts to find the sweet spot for your kernel.
迭代优化和验证：每次更改后，包括占用率调整、ILP重构、分块或精度缩放，您应该重新分析以确认减少的内存停顿、更少的执行依赖以及更高的实现占用率或FLOPS。在降低精度时，始终使用断言或数值检查与FP32基线比较结果，以防止不可接受的精度回归。
Iterate and verify：After each change, including occupancy adjustments, ILP restructuring, tiling, or precision scaling, you should reprofile to confirm reduced memory stalls, fewer execution dependencies, and higher achieved occupancy or FLOPS. When lowering precision, always use assertions or numerical checks against an FP32 baseline to prevent unacceptable accuracy regressions.
结论 (Conclusion)¶
您看到了有效调优GPU性能需要迭代工作流程。首先，使用分析器测量。然后识别您的主要瓶颈（计算、内存、互连等）。最后，应用正确的优化并重复。
You saw that effectively tuning GPU performance requires an iterative workflow. First, measure using profilers. Then identify your primary bottlenecks (compute, memory, interconnect, etc.). Finally, apply the right optimizations and repeat.
虽然每个新GPU架构在计算和内存带宽方面带来改进，但它们也增加了复杂性。您必须不断跟上这些创新以维持峰值性能。
While every new GPU architecture brings improvements in compute and memory bandwidth, they also add complexity. You must continually keep up with these innovations to sustain peak performance.
我们介绍了如何调优占用率、避免线程束分歧和增加ILP。您还看到了如何使用分析器关联CPU、内核和跨GPU的NVLink流量。然后我们使用Nsight Compute获取每内核详细信息。
We covered how to tune occupancy, avoid warp divergence, and increase ILP. You also saw how to use profilers to correlate CPU, kernel, and cross-GPU NVLink traffic. Then we used Nsight Compute for per-kernel details.
在接下来的章节中，我们将通过专注于内核级效率来扩展这一基础，以增加算术强度。为此，您将充分利用GPU的硬件优化资源，如用于计算的Tensor Core、为Tensor Core服务的TMEM以及用于数据传输的TMA。TMA仍然是现代支持它的GPU上从全局内存到共享内存批量复制的首选方法。让我们继续推向硬件的峰值性能极限！
In the upcoming chapters, we'll extend this foundation by focusing on kernel-level efficiency to increase arithmetic intensity. To do so, you'll take full advantage of the GPU's hardware-optimized resources like Tensor Cores for compute, TMEM to serve Tensor Cores, and TMA for data movement. TMA remains the preferred method for bulk copying from global memory to shared memory on modern GPUs that support it. Let's continue pushing toward the hardware's peak performance limits!