4-1性能指南

本文参考：

• NVIDIA CUDA C++ Programming Guide  第 8 章，基于硬件架构讲述了优化方法论
• NVIDIA CUDA C best practice
• NV 官方 Profiling Guide 给出了基于 NCU 的 profile 指南 本章也会分三节来讲
• Professional CUDA C Programming 第 章

5 性能指南

5.1 性能优化策略总览

性能优化围绕四个基本策略：

Note

• 最大化并行执行以实现最大利用率；
• 优化内存使用，实现最大内存吞吐量；
• 优化指令使用，实现最大指令吞吐量；
• 尽量减少内存抖动。

对于应用程序的特定部分，哪种策略能带来最佳的性能提升，取决于该部分的性能限制因素。例如，对于一个主要受内存访问限制的内核，优化其指令使用不会带来任何显著的性能提升。因此，优化工作应始终以对性能限制因素的测量和监控为指导，例如可以使用 CUDA 分析器。此外，将特定内核的浮点运算吞吐量或内存吞吐量（取更有意义的那个）与设备相应的理论峰值吞吐量进行比较，能表明该内核还有多少改进空间。

5.2 最大化使用率

为了最大限度地提高利用率，应用程序的结构应能使其展现出尽可能多的并行性，并将这种并行性有效地映射到系统的各个组件，使它们大部分时间都处于忙碌状态。

Application Level & Device Level

在宏观层次上，应用程序应该使用异步函数调用和 Asynchronous Concurrent Execution 中描述的 Stream 来最大化主机、设备、以及将主机与设备相连的总线之间的并行执行。同时应用程序应该为不同的处理器分配它最擅长的工作类型：将串行任务发送给主机；将并行任务发送给设备。

对于并行任务来说，在算法中，由于某些线程需要同步以便彼此共享数据，因此并行性被破坏，这种有两种情况：第一种情况是这些线程属于同一个块，在这种情况下，应使用 __syncthreads () 进行同步并在同一个内核调用中通过共享内存来共享数据；另一种情况是它们属于不同的块，在这种情况下，它们必须通过全局内存来共享数据，并使用两个单独的内核调用，一个调用用于向全局内存写入，另一个调用用于从全局内存读取。第二种情况的性能会不太理想，因为它增加了额外内核调用以及全局内存移动的开销。因此，应该将算法映射到 CUDA 编程模型，使需要线程间通信的计算尽可能在单个线程块内执行，从而最大限度地减少第二种情况的发生。

在次低的层面上，应用程序应该最大化设备中 SM 之间的并行执行。

多个内核可以在一个设备上并发执行，因此也可以通过使用 Stream 来启用足够多的内核来实现设备的利用率最大化，如 Asynchronous Concurrent Execution 中所述。

这两个层面是相对顶层的，主要是尽量利用 kernel 的异步并发执行。

Multiprocessor Level

在更低的层次上，应用程序应该最大化 SM 内不同功能单元之间的并行执行。

延迟隐藏

如 Hardware Multithreading 所述，GPU SM 主要依靠线程级并行性来最大限度地利用其功能单元。因此，利用率与驻留线程束的数量直接相关。在每个指令发射时刻，线程束调度器会选择一条准备好执行的指令。这条指令可以是同一线程束中的另一条独立指令（利用指令级并行），或者更常见的是另一个线程束中的指令（利用线程级并行）。

若某条待执行指令被选中，该指令会被发射到对应 warp 的 active 线程中。warp 为执行下一条指令而等待就绪所需的时钟周期数，被称为延迟（latency）；而要实现完全利用率，需满足：在这段延迟期间的每个时钟周期内，所有线程束调度器（warp scheduler）始终有可发射给某个线程束的指令 —— 换句话说，就是延迟被完全 “隐藏（hidden）”。

对具有 $L$ 个时钟周期的时延进行隐藏所需的指令数量取决于指令自身的吞吐量（有关各种算术指令的吞吐量，请参见 CUDA C++ Best Practices Guide）。如果假设指令具有最大吞吐量，则需要的指令数量等于：

• 对于计算能力 5.x、6.1、6.2、7.x 和 8.x 的设备来说，该值为 $4L$。因为对于这些设备，SM 在一个时钟周期内，可同时向 4 个线程束各发射一条指令。如 CUDA C++ Best Practices Guide 中所述。
• 对于计算能力 6.0 的设备，该值为 $2L$。这是因为这类设备每个时钟周期发射的两条指令，分别对应两个不同的线程束（即一个时钟周期内可向 2 个线程束各发射一条指令）。

Warp 未准备好执行下一条指令的最常见原因是该指令的输入操作数尚不可用。

如果所有输入操作数都是寄存器操作数，则时延是由寄存器依赖导致的，也就是：一些输入操作数由一些先前尚未完成的指令来写入。在这种情况下，时延等于前一条指令的执行时间，Warp 调度器在此期间会需要调度其他 Warp 的指令。执行时间因指令而异。在计算能力 7.x 的设备上，大多数算术指令通常需要 4 个时钟周期。这意味着每个 SM 需要 16 个活动 warp（4 个周期，4 个 warp 调度器）来隐藏算术指令延迟（假设 Warp 以最大吞吐量执行指令，否则不需要这么多 Warp）。如果各个 Warp 表现出指令级并行性，也就是在它们的指令流中有多个相互独立指令，则需要更少的 warp，这是因为多个独立指令可以从一个 Warp 中连续发出。

以计算能力 7.x 的设备为例（如 Volta 架构）：

• 大多数算术指令的执行时间（即延迟）为4 个时钟周期。
• 每个多处理器有4 个线程束调度器，每个调度器每个时钟周期可向一个线程束发射指令。

因此，要完全 “隐藏 4 个时钟周期的延迟”，需要让 4 个调度器在这 4 个周期内始终有指令可发射：

• 每个周期，4 个调度器各处理 1 个线程束 → 1 个周期需要 4 个活跃线程束。
• 4 个周期就需要 4 × 4 = 16 个活跃线程束（每个线程束负责在某个周期被某个调度器调度）。

这样，当某个线程束因寄存器依赖等待时，其他 15 个线程束可以轮流被调度，让硬件始终处于忙碌状态，从而抵消延迟的影响。

若部分输入操作数位于片外内存（off-chip memory）中，则延迟会高得多：通常为数百百数百个时钟周期。在如此高的延迟周期下，要保持线程束调度器处于忙碌状态所需的线程束数量，取决于内核代码及其指令级并行性的程度。

算术强度和延时隐藏

一般而言，如果无片外内存操作数的指令（多数情况下为算术指令）数量与有片外内存操作数的指令数量之比（该比例通常称为程序的算术强度 arithmetic intensity）较低，则需要更多的线程束。

Warp 未准备好执行其下一条指令的另一个原因是由于其处于某个内存栅栏 (memory fence)（Memory Fence Functions）或同步点 (synchronization point)（Synchronization Functions）处等待。

同步点（synchronization point）可能会迫使多处理器（multiprocessor）进入空闲状态 —— 这是因为随着越来越多的线程束（warp）等待同一线程块（block）内的其他线程束完成 “同步点之前的指令执行”，多处理器的计算资源会逐渐闲置。在这种情况下，若多处理器中同时驻留多个线程块（multiple resident blocks），则有助于减少空闲时间。这是因为来自不同线程块的线程束，无需在同步点处相互等待（同步点的约束仅作用于同一线程块内的线程束）。

SM 的并行

对于一个内核调用来说，驻留在每个 SM 上的 Block 和 Warp 的数量取决于调用的执行配置（Execution Configuration）、SM 的内存资源以及内核的资源需求，如 Hardware Multithreading 中所述。使用 --ptxas-options=-v 选项编译时，编译器会报告寄存器和共享内存的使用情况。一个块所需的共享内存总量等于静态分配的共享内存量和动态分配的共享内存量之和。

内核使用的寄存器数量会对驻留 Warp 的数量产生重大影响。例如，对于计算能力为 6.x 的设备，如果内核使用 64 个寄存器并且每个块有 512 个线程并少量的共享内存，那么在多处理器上可以常驻两个块（即 32 个 warp），因为当前配置需要 $2\times 512\times 64$ 个寄存器，这个值与 SM 上可用的寄存器数量匹配。但是一旦内核多使用一个寄存器，就只能在 SM 上常驻一个块（即 16 个 warp），因为两个块需要 $2\times 512\times 65$ 个寄存器，这比 SM 上可用的寄存器的数量多。因此，编译器会在尽量减少寄存器使用的同时，将寄存器溢出（见 Device Memory Accesses）和指令数量控制在最低限度。寄存器使用量可通过 maxrregcount 编译选项、Launch Bounds 中所述的 __launch_bounds__() 限定符，或 Maximum Number of Registers per Thread 中所述的 __maxnreg__() 限定符进行控制。

寄存器文件按 32 位寄存器进行组织。因此，存储在寄存器中的每个变量都需要至少一个 32 位寄存器，例如，double 变量使用了两个 32 位寄存器。

对于特定的内核调用，执行配置（execution configuration）对性能的影响通常取决于内核代码，因此建议通过实验进行验证。应用程序还可根据寄存器文件大小和共享内存大小来参数化执行配置 —— 这些大小取决于设备的计算能力，同时也与设备的多处理器数量和内存带宽相关，而所有这些信息都可通过运行时查询（参见参考手册）。

每个线程块的线程数量应选择为线程束大小（warp size）的整数倍，以尽可能避免因线程束未填满（under-populated warps）而浪费计算资源。

SM 级并行的核心

1. GPU 多处理器的核心设计思想是通过大量线程级并行（TLP）来掩盖指令执行和内存访问的延迟。当一个 warp（线程束）因等待数据（如寄存器依赖或内存访问）而无法执行时，调度器会立刻切换到其他就绪的 warp，让计算单元始终保持忙碌。这种 “延迟隐藏” 能力的关键是足够多的驻留 warp（即同时在多处理器上活跃的 warp）。
2. Warp 未就虚的瓶颈在于数据依赖、内存访问、同步。
3. 而驻留的 Warp 会受调用的配置即 block 维度、硬件配置即算力限制，需要结合分析与实验确定。

Occupancy Calculator

有几个 API 函数来帮助程序员根据寄存器和共享内存的大小要求选择适合的线程块大小。

• 占用计算 (Occupancy Calculator) API，cudaOccupancyMaxActiveBlocksPerMultiprocessor，可以根据内核的块大小和共享内存使用情况提供占用预测。此函数根据每个 SM 的并发线程块数量报告占用情况。
  • 注意，此值可以转换为其他指标。该值乘以每个块的 Warp 数可得每个 SM 的并发 Warp 数；进一步，将并发 Warp 数除以每个 SM 的最大 Warp 数可得占用率的百分比。
• 基于占用率的启动配置 API，cudaOccupancyMaxPotentialBlockSize 和 cudaOccupancyMaxPotentialBlockSizeVariableSMem，启发式计算实现 SM-level 的最大占用率所对应的执行配置。
• 占用率计算器 API（即 cudaOccupancyMaxActiveClusters）可根据内核的集群大小、块大小和共享内存使用情况，提供占用率预测。该函数会以系统中 GPU 上特定大小的 “最大活动集群数量” 来表示占用率。

以下代码示例展示了计算 MyKernel 的占用率。然后，通过计算并发 Warp 数与每个 SM 的最大 Warp 之间的比值输出 occupancy level。

// Device code
__global__ void MyKernel(int *d, int *a, int *b)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    d[idx] = a[idx] * b[idx];
}
 
// Host code
int main()
{
    int numBlocks;        // Occupancy in terms of active blocks
    int blockSize = 32;
 
    // These variables are used to convert occupancy to warps
    int device;
    cudaDeviceProp prop;
    int activeWarps;
    int maxWarps;
 
    cudaGetDevice(&device);
    cudaGetDeviceProperties(&prop, device);
 
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &numBlocks,
        MyKernel,
        blockSize,
        0);
 
    activeWarps = numBlocks * blockSize / prop.warpSize;
    maxWarps = prop.maxThreadsPerMultiProcessor / prop.warpSize;
 
    std::cout << "Occupancy: " << (double)activeWarps / maxWarps * 100 << "%" << std::endl;
 
    return 0;
}

下面的代码示例展示了根据用户输入配置了一个基于占用率的内核启动 MyKernel。

// Device code
__global__ void MyKernel(int *array, int arrayCount)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < arrayCount) {
        array[idx] *= array[idx];
    }
}
 
// Host code
int launchMyKernel(int *array, int arrayCount)
{
    int blockSize;      // The launch configurator returned block size
    int minGridSize;    // The minimum grid size needed to achieve the
                        // maximum occupancy for a full device
                        // launch
    int gridSize;       // The actual grid size needed, based on input
                        // size
 
    cudaOccupancyMaxPotentialBlockSize(
        &minGridSize,
        &blockSize,
        (void*)MyKernel,
        0,
        arrayCount);
 
    // Round up according to array size
    gridSize = (arrayCount + blockSize - 1) / blockSize;
 
    MyKernel<<<gridSize, blockSize>>>(array, arrayCount);
    cudaDeviceSynchronize();
 
    // If interested, the occupancy can be calculated with
    // cudaOccupancyMaxActiveBlocksPerMultiprocessor
 
    return 0;
}

以下代码示例展示了如何使用集群占用率 API 来查找特定大小的最大活动集群数量。下面的示例代码计算了集群大小为 2 且每个块 128 个线程的占用率。

从计算能力 9.0 开始，集群大小为 8 具有向前兼容性，但在 GPU 硬件或 MIG 配置过小而无法支持 8 个多处理器的情况下，最大集群大小会相应减小。不过，建议用户在启动集群内核前查询最大集群大小。可以使用 cudaOccupancyMaxPotentialClusterSize API 查询最大集群大小。

{
  cudaLaunchConfig_t config = {0};
  config.gridDim = number_of_blocks;
  config.blockDim = 128; // threads_per_block = 128
  config.dynamicSmemBytes = dynamic_shared_memory_size;
 
  cudaLaunchAttribute attribute[1];
  attribute[0].id = cudaLaunchAttributeClusterDimension;
  attribute[0].val.clusterDim.x = 2; // cluster_size = 2
  attribute[0].val.clusterDim.y = 1;
  attribute[0].val.clusterDim.z = 1;
  config.attrs = attribute;
  config.numAttrs = 1;
 
  int max_cluster_size = 0;
  cudaOccupancyMaxPotentialClusterSize(&max_cluster_size, (void *)kernel, &config);
 
  int max_active_clusters = 0;
  cudaOccupancyMaxActiveClusters(&max_active_clusters, (void *)kernel, &config);
 
  std::cout << "Max Active Clusters of size 2: " << max_active_clusters << std::endl;
}

CUDA Nsight Compute 用户界面还在 <CUDA_Toolkit_Path>/include/cuda_occupancy.h 中提供了一个独立的占用率计算器和启动配置器实现，适用于任何无法依赖 CUDA 软件栈的使用场景。Nsight Compute 版本的占用率计算器尤其适合作为学习工具，它能可视化展示影响占用率的参数（块大小、每个线程的寄存器数量、每个线程的共享内存）发生变化时所产生的影响。

5.3 最大化存储吞吐量

要最大限度提高应用程序的整体内存吞吐量，第一步是尽量减少低带宽的数据传输。

1. 这意味着要减少主机（Host）与设备（Device）之间的数据传输（如 Data Transfer between Host and Device），因为这类传输的带宽远低于全局内存（Global Memory）与设备之间的数据传输带宽。
2. 这同时也意味着，要通过最大限度利用片上内存（On-chip Memory）来减少全局内存与设备之间的数据传输：片上内存包括共享内存（Shared Memory）和高速缓存（即，计算能力 2.x 及以上的设备所具备的 L1 缓存与 L2 缓存，以及所有设备均具备的纹理缓存（Texture Cache）和常量缓存（Constant Cache））。

共享内存 Shared memory是一种用户可管理的缓存：应用程序可以显式分配它和访问它。如 CUDA Runtime 所示，典型的程序模式是将来自设备内存的数据暂存到共享内存中；换句话说，对于一个块中的每个线程：

1. 从设备内存（Device Memory）中将数据加载到共享内存（Shared Memory）；
2. 与线程块（Block）中的所有其他线程进行同步，确保每个线程都能安全读取由其他线程填充到共享内存中的数据；
3. 在共享内存中对数据进行处理；
4. 若有必要，再次进行同步，以确保共享内存已更新为处理后的结果；
5. 将结果写回设备内存。

对于某些应用程序（例如，全局内存访问模式是数据依赖的），使用传统的硬件管理缓存的方法会更好地利用数据局部性，如 Compute Capability 7.x, Compute Capability 8.x and Compute Capability 9.0 中所述，对于计算能力 3.x、7.x 和 8.x 的设备，它们用于 L1 和共享内存的片上存储器数量一致，并可针对每个内核调用的情况对 L1 以及共享内存的使用量进行针对性配置。

根据不同内存类型的访问模式，内核访问内存的吞吐量会发生数量级程度的变化。因此，最大化内存吞吐量的下一步是根据 Device Memory Accesses 中描述的最佳内存访问模式，尽可能优化地组织内存访问方式。这种优化对于全局内存访问尤为重要，因为与可用的片上带宽和算术指令吞吐量相比，全局内存带宽较低，因此非最佳的全局内存访问方式通常会对性能产生很大影响。

5.3.1 设备与主机之间的数据传输

应用程序应尽量减少主机和设备之间的数据传输。实现这一点的一种方法是将更多代码从主机移动到设备，即使在设备中运行的内核无法提供足够的并行性以达到最优效率。中间数据结构可以在设备内存中创建，由设备操作，并在无主机映射或复制到主机内存的情况下销毁。

此外，由于每次传输都会产生相关开销，因此将多个小型传输批量整合为单次大型传输，其性能始终优于单独执行每次小型传输。

在配备前端总线（Front-Side Bus）的系统中，通过使用 “页锁定主机内存”（Page-Locked Host Memory，详见 Page-Locked Host Memory），可提升主机与设备之间数据传输的性能。

另外，当使用 “映射页锁定内存”（Mapped Memory，Mapped Memory）时，无需分配任何设备内存，也无需在设备内存与主机内存之间执行显式的数据复制操作。每当核函数（Kernel）访问映射内存时，数据传输会以隐式方式自动完成。为实现最佳性能，此类内存访问必须像访问全局内存（Global Memory）那样进行合并（见 Device Memory Accesses）。假设内存访问已满足合并条件，且映射内存仅被读取或写入一次，那么使用映射页锁定内存替代设备与主机内存间的显式复制，通常能带来性能提升。

对于设备内存和主机内存在物理上相同的集成系统来说，主机和设备内存之间的任何复制都是多余的，应使用映射锁页内存。应用程序可以通过查询集成设备属性（见 Device Enumeration）是否等于 1 来确认设备是否集成。

将更多代码从主机迁移至设备执行，中间数据也在设备内存中处理。采用锁页内存避免显式数据复制，数据传输隐式完成

5.3.2 设备内存访问

访问可寻址内存（即全局内存、本地内存、共享内存、常量内存或纹理内存）的指令，可能需要根据线程束（warp）内各线程的内存地址分布情况多次重新发射。这种地址分布对指令吞吐量的影响方式因内存类型而异，具体将在以下章节中阐述。例如，对于全局内存，一般规则是：地址分布越分散，吞吐量下降越明显。

Global Memory

全局内存（Global Memory）位于设备内存（Device Memory）中，而设备内存需通过 32 字节、64 字节或 128 字节的内存事务（Memory Transaction）进行访问。这些内存事务必须满足自然对齐（Naturally Aligned）要求：设备内存中只有那些与自身大小对齐的 32 字节、64 字节或 128 字节段（即段的起始地址是其大小的整数倍），才能通过内存事务进行读写操作。

当一个线程束（warp）执行一条访问全局内存（global memory）的指令时，会根据每个线程所访问数据的字长（word size）以及内存地址在各线程间的分布情况，将该线程束内所有线程的内存访问请求合并（coalesce）为一个或多个内存事务（memory transaction）。通常情况下，所需的内存事务数量越多，除了线程实际访问的数据字之外，需要额外传输的未使用数据字就越多，指令吞吐量（instruction throughput）也会随之降低。例如，若为每个线程 4 字节的访问请求生成一个 32 字节的内存事务，那么吞吐量会降至原来的 1/8（即除以 8）。

需要多少事务以及最终影响多少吞吐量取决于设备的计算能力。如何为各种计算能力的设备处理全局内存访问，Compute Capability 3.x、Compute Capability 5.x、Compute Capability 6.x、Compute Capability 7.x 和 Compute Capability 8.x 中提供了更多详细信息。

为了最大化全局内存吞吐量，以下方式的最大化合并非常重要：

• 遵循 Compute Capability 3.x、Compute Capability 5.x、Compute Capability 6.x、Compute Capability 7.x 和 Compute Capability 8.x 的最佳访问模式
• 使用符合下文 “大小与对齐要求”（Size and Alignment Requirement）章节中详细说明的大小及对齐要求的数据类型。
• 在部分情况下对数据进行填充（Padding），例如，在下文 “二维数组”（Two-Dimensional Arrays）章节中所述的访问二维数组时。

Size And Alignment Requirement

全局内存指令支持读取或写入大小为 1、2、4、8 或 16 字节的数据字。只有当数据类型的大小为 1、2、4、8 或 16 字节，且数据自然对齐（即其地址是该大小的整数倍）时，对全局内存中数据的任何访问（通过变量或指针）才会编译为单条全局内存指令。

如果不满足这种大小和对齐要求，访问操作会编译为多条具有交错访问模式的指令，这会阻碍这些指令实现完全的访问合并。因此，建议为驻留在全局内存中的数据使用满足此要求的数据类型。

内置向量类型（Built-in Vector Types）会自动满足对齐要求。

对于结构，大小和对齐要求可以由编译器使用对齐说明符 __align__(8) 或 __align__(16) 强制执行，例如:

struct __align__(8) {
    float x;
    float y;
};
 
struct __align__(16) {
    float x;
    float y;
    float z;
};

驻留在全局内存中，或由驱动程序或运行时 API 的某个内存分配例程返回的变量的任何地址总是对齐到至少 256 字节。

读取非自然对齐的 8 字节或 16 字节数据字会产生错误结果（偏移几个数据字），因此必须特别注意保持这些类型的任何值或值数组的起始地址对齐。

一种容易被忽视的典型情况是使用某些自定义全局内存分配方案时：此时，多个数组的分配（通过多次调用 cudaMalloc () 或 cuMemAlloc ()）被替换为分配单个大内存块，再将其划分为多个数组。在这种情况下，每个数组的起始地址相对于该块的起始地址存在偏移。

Two-Dimensional Arrays

一个常见的全局内存访问模式是：索引 (tx,ty) 的每个线程使用以下地址访问一个宽度为 width 的二维数组中的元素时，位于 type* 类型的地址 BaseAddress（其中 type 满足在 Maximize Utilization 中描述的要求要求）：

BaseAddress + width * ty + tx

为了让这些访问完全合并，线程块的宽度和数组的宽度都必须是 Warp 大小的整数倍。

具体而言，这意味着：如果某个数组的宽度不是上述（内存访问优化所需的）大小的整数倍，那么将其实际分配的宽度向上取整至该大小的最接近整数倍，并对数组的行进行相应填充（Padding），就能显著提升该数组的访问效率。参考手册中介绍的 cudaMallocPitch() 和 cuMemAllocPitch() 函数，以及与其配套的内存复制函数，可帮助开发者编写与硬件无关的代码，从而分配出符合这些（内存对齐与宽度）约束条件的数组。

Local Memory

本地内存访问仅发生在 Variable Memory Space Specifiers 中提到的某些自动变量中。编译器可能放置在本地内存中的变量有：

• 无法确定其索引为常量的数组；
• 会占用过多寄存器空间的大型结构体或大型数组；
• 若内核（kernel）使用的寄存器数量超过可用寄存器数量，则所有变量都会受到影响（这种情况也称为寄存器溢出，register spilling）。

通过查看 PTX 汇编代码（使用 -ptx 或 -keep 选项编译可获得），可以判断变量在编译的第一阶段是否被分配到本地内存（local memory）：若变量是通过 .local 助记符（mnemonic）声明，并通过 ld.local（本地内存读取）和 st.local（本地内存写入）助记符访问，则说明其被分配到了本地内存。

即便变量在第一阶段未被分配到本地内存，后续编译阶段若发现该变量针对目标架构（targeted architecture）会占用过多寄存器空间，仍可能改变分配策略（将其移入本地内存）：此时可通过 cuobjdump 工具查看 cubin 目标文件，以确认是否存在这种情况。此外，使用 --ptxas-options=-v 选项编译时，编译器会报告每个内核（kernel）的总本地内存使用率（标记为 lmem）。需注意，部分数学函数的实现路径中可能存在对本地内存的访问。

本地内存（local memory）空间位于设备内存（device memory）中，因此本地内存访问与全局内存（global memory）访问具有相同的高延迟和低带宽特性，并且需遵循 Device Memory Accesses 所述的相同内存合并（memory coalescing）要求。

不过，本地内存的组织方式存在特殊性：连续的 32 位数据字由连续的线程 ID（thread ID）进行访问。因此，只要一个线程束（warp）中的所有线程都访问相同的相对地址（例如，数组变量中相同的索引、结构体变量中相同的成员），本地内存访问就能实现完全合并。

在计算能力（Compute Capability）5.x 及更高版本的设备上，本地内存访问会像全局内存访问一样，始终在二级缓存（L2）中进行缓存（see Compute Capability 5.x and Compute Capability 6.x）。

Shared Memory

由于共享内存（shared memory）位于芯片上（on-chip），其带宽远高于本地内存（local memory）或全局内存（global memory），延迟也远低于后两者。

为实现高带宽，共享内存被划分为多个大小相等的内存模块，这些模块被称为 “存储体”（bank），且可被同时访问。因此，若某个内存读写请求涉及 n 个地址，且这些地址分别落在 n 个不同的 bank 中，该请求就能被并行处理，最终实现的总带宽将是单个存储体带宽的 n 倍。

然而，若一个内存请求中的两个地址落在同一个存储体中，就会产生 “存储体冲突”（bank conflict），此时访问操作必须串行执行。硬件会将存在存储体冲突的内存请求拆分为若干个无冲突的独立请求（拆分数量取决于冲突情况），这会导致吞吐量降低，降低倍数等于独立请求的数量。若独立请求的数量为 n，则称初始内存请求引发了 “n 路存储体冲突”（n-way bank conflict）。

因此，要实现最佳性能，关键在于理解内存地址如何映射到存储体，进而合理规划内存请求以最小化存储体冲突。关于这一点，针对计算能力（Compute Capability）分别为 5.x、6.x、7.x、8.x、9.0、10.0 及 12.0 的设备，相关说明已在对应计算能力版本的文档章节中给出，见 Compute Capability 5.x, Compute Capability 6.x, Compute Capability 7.x, Compute Capability 8.x, Compute Capability 9.0, Compute Capability 10.0, and Compute Capability 12.0。

Constant Memory

常量内存（constant memory）空间位于设备内存中，并被缓存到常量缓存（constant cache）中。

此时，一个请求会被拆分为与初始请求中不同内存地址数量相同的多个独立请求，吞吐量会降低，降低倍数等于独立请求的数量。

若缓存命中（cache hit），则拆分后的请求将以常量缓存的吞吐量进行处理；否则，将以设备内存的吞吐量进行处理。

Texture And Surface Memory

纹理内存（texture memory）和表面内存（surface memory）空间位于设备内存（device memory）中，并被缓存到纹理缓存（texture cache）中。因此，纹理读取（texture fetch）或表面读取（surface read）操作仅在缓存未命中（cache miss）时，才需要从设备内存读取一次数据；若缓存命中（cache hit），则只需从纹理缓存读取一次即可。

纹理缓存针对二维空间局部性（2D spatial locality）进行了优化，因此当同一个线程束（warp）中的线程读取纹理或表面内存中二维空间上相邻的地址时，能实现最佳性能。此外，纹理缓存专为具有固定延迟（constant latency）的流式读取（streaming fetches）设计：缓存命中可降低对 DRAM（动态随机存取存储器）带宽的需求，但不会减少读取延迟。

通过纹理读取或表面读取的方式访问设备内存，相比从全局内存（global memory）或常量内存（constant memory）访问设备内存，具有一些独特优势，使其成为更优选择：

• 若内存读取无法遵循全局内存或常量内存为实现高性能而要求的访问模式，只要纹理读取或表面读取具有局部性（locality），仍可实现更高带宽；
• 地址计算由专用单元在核函数（kernel）外部执行；
• 打包数据（packed data）可在单次操作中广播（broadcast）到多个独立变量；
• 8 位和 16 位整数输入数据可按需转换为范围在 [0.0, 1.0] 或 [-1.0, 1.0] 之间的 32 位浮点值（详见 Texture Memory）。

5.4 最大化指令吞吐量

见 Best practice

See the CUDA C++ Best Practices Guide for more details on optimizing instruction throughput.

5.5 最小化内存抖动

频繁执行内存分配与释放操作的应用程序，可能会发现分配调用的速度会随时间逐渐变慢，直至达到某个瓶颈。通常而言，这种现象是正常的 —— 因为释放的内存会归还给操作系统供其调度使用，这一过程本身就会带来性能影响。为在该方面实现最佳性能，我们建议采取以下措施：

1. 内存分配大小应与手头问题需求匹配。不要尝试通过 cudaMalloc/cudaMallocHost/cuMemCreate 分配所有可用内存，这种操作会强制内存立即处于 “驻留”（resident）状态，导致其他应用程序无法使用这部分内存。这不仅会增加操作系统调度器的压力，甚至可能完全阻止其他使用同一 GPU 的应用程序运行。
2. 建议在应用程序启动初期，就以合适的大小分配所需内存；仅当内存确实不再被应用程序使用时，才执行释放操作。尽量减少应用程序中 cudaMalloc 与 cudaFree 的调用次数，尤其是在性能关键区域（performance-critical regions）。
3. 若应用程序无法分配足够的设备内存，可考虑改用其他类型的内存分配方式，例如 cudaMallocHost（主机锁定内存）或 cudaMallocManaged（统一内存）。这些方式的性能可能略逊于设备内存，但能确保应用程序继续正常运行。
4. 在支持该特性的平台上，cudaMallocManaged 允许 “内存超分配”（oversubscription，即分配的内存大小可超过 GPU 实际物理内存）；且通过启用正确的 cudaMemAdvise 策略，应用程序即使使用统一内存，也能保留 cudaMalloc（设备内存分配）的大部分（甚至全部）性能。此外，cudaMallocManaged 不会强制内存立即驻留，仅在内存被实际使用或预取（prefetched）时才会使其驻留，这能降低操作系统调度器的整体压力，更适合 “多租户”（multi-tenant）使用场景（即多个应用程序共享同一 GPU 资源）。

Best Practice

为了最大化指令吞吐量，应用程序应该：

• 尽量减少使用低吞吐量的算术指令；包括在不影响最终结果的情况下通过精度换取速度，例如使用内部函数而不是常规函数（内部函数在 内部函数 中列出），使用单精度而不是双精度，或者将非规范化数字赋值为零；
• 最大限度地减少由控制流指令引起的 Warp 发散，如 控制流指令 中所述
• 减少指令的数量，例如，尽可能优化同步点（如 同步指令 中所述）或使用受限指针（如 restrict 中所述）。

在本节中，吞吐量以每 SM 的每时钟周期的操作数目给出。对于 32 大小的 Warp，一条指令对应 32 次操作，因此若 N 是每个时钟周期的操作数量，则指令吞吐量为 N/32 每时钟周期/指令。

所有吞吐量都是针对单个 SM 的。整个设备的吞吐量必须乘以设备中的所有 SM 的数量。

5.4.1 算数指令

如下图所示

其他指令和函数是在 Native 指令之上实现的。计算能力不同的设备可能有不同的实现，编译后的 Native 指令的数量可能会随着编译器版本的不同而变化。对于复杂的函数，可能有多个代码路径，具体取决于输入。cuobjdump 可用于检查 cubin 对象中的特定实现。

一些函数的实现在 CUDA 头文件（math_functions.h、device_functions.h、…）中很容易获得。

通常，使用 -ftz=true 编译的代码（非规范化数字刷新为零）往往比使用 -ftz=false 编译的代码具有更高的性能。类似地，使用 -prec-div=false（不太精确的除法）编译的代码往往比使用 -prec-div=true 编译的代码具有更高的性能，使用 -prec-sqrt=false（不太精确的平方根）编译的代码往往比使用 -prec-sqrt=true 编译的代码具有更高的性能。nvcc 用户手册更详细地描述了这些编译标志。

Single-Precision Floating-Point Division

__fdividef(x, y)（参见 内部函数）提供了比除法运算符更快的单精度浮点除法。

Single-Precision Floating-Point Reciprocal Square Root

为了保留 IEEE-754 语义，编译器会将 1.0 / sqrtf() 优化为 rsqrtf()，这种情况只发生在倒数与平方根都相近时（即 -prec-div=false 和 -prec-sqrt=false）。因此，如果需要时，建议直接调用 rsqrtf()。

Single-Precision Floating-Point Square Root

单精度浮点平方根的实现是先取倒数的平方根，然后取倒数，而不是在计算倒数平方根后跟乘法，因此它在 0 和无穷大情况下提供正确的结果。

Sine and Cosine

sinf(x)、cosf(x)、tanf(x)、sincosf(x) 和相应的双精度指令更为昂贵，同时如果参数 x 的量级很大，则更是如此。

更准确地说，参数缩减代码（参见实现的 数学函数）包括两个代码路径，分别称为快速路径和慢速路径。

快速路径用于大小足够小的参数，它基本由几个乘加运算组成。慢速路径用于量级较大的参数，包含了许多冗长的计算，因为需要在整个参数范围内获得正确结果。

目前，三角函数的参数缩减代码在单精度函数量级小于 105615.0f，双精度函数量级小于 2147483648.0 时选择快速路径。

由于慢速路径比快速路径需要更多的寄存器，因此会在本地内存中存储一些中间变量来降低慢速路径的寄存器压力，但是因为本地内存的高延迟和带宽（请参阅 设备内存访问）的原因，这可能会而影响性能。目前，单精度函数使用 28 字节的本地内存，双精度函数使用 44 字节。并且，确切的本地内存数量可能会发生变化。

由于在慢路径中需要进行冗长的计算和使用本地内存，当需要进行慢路径缩减时，与快速路径缩减相比，这些三角函数的吞吐量会低一个数量级。

Integer Arithmetic

整数除法和模运算的成本很高，因为它们最多可编译为 20 条指令。在某些情况下，可以用位运算代替除法和取模运算：如果 n 是 2 的幂，则 (i/n) 等价于 (i>>log2(n)) 并且 (i%n) 等价于 (i&(n1)); 如果 n 是字母，则编译器会执行这些转换。

__brev 和 __popc 将映射为一条指令，而 __brevll 和 __popcll 将映射为几条指令。

__[u]mul24 是遗留内部函数，在任何情况下都不应该使用。

Half Precision Arithmetic

为了实现 16 位精度浮点加、乘或乘加的良好性能，建议使用 half2 数据类型替换 half 精度，使用 __nv_bfloat162 替换 __nv_bfloat16 精度，并使用 vector intrinsics 函数（例如 __hadd2、__hsub2、__hmul2、__hfma2）在一条指令中执行两个操作。使用 half2 或 __nv_bfloat162 代替使用 half 或 __nv_bfloat16 的两个调用也可能有助于提升其他 intrinsics 函数的性能，例如 Warp Shuffles。

提供了 intrinsics 的 __halves2half2 以将两个 half 精 度值转换为 half2 数据类型。

提供了 intrinsics 的 __halves2bfloat162 以将两个 __nv_bfloat 精度值转换为 __nv_bfloat162 数据类型。

Type Conversion

有时，编译器必须插入转换指令，而这引入了额外的执行周期。具体情况如下：

• 对 char 或 short 类型进行操作的函数，其操作数通常需要转换为 int，
• （由 C/C++ 标准规定）单精度浮点计算的输入会转换为双精度浮点常量（即转换为那些没有任何类型后缀定义的常量）。

最后一种情况可以通过使用单精度浮点常量来避免，这些常量使用 f 后缀定义，例如 3.141592653589793f、1.0f、0.5f。

5.4.2 控制流指令

任何流控制指令（if、switch、do、for、while）都会导致同一个 Warp 中的线程发散（即遵循不同的执行路径），并显着影响有效指令的吞吐量。如果发生这种情况，则必须对不同的执行路径进行序列化，从而增加 Warp 执行的指令总数。

在控制流依赖线程 ID 的情况下，为了获得最佳性能，应设计控制条件来最小化 Warp 发散的数量。这是可能实现的，因为一个线程块中的 Warp 分布是确定性的，如 SIMT 架构 中所说。一个简单的例子是当控制条件仅取决于 (threadIdx / warpSize) 时，这里 warpSize 是 Warp 的大小。在这种情况下，由于控制条件与 Warp 完全对齐，因此不会出现 Warp 发散。

如下所述，有时，编译器可能会进行循环展开，或者会通过分支预测来优化短的 if 或 switch 块。在这些情况下，所有 Warp 都不会发散。程序员还可以使用#pragma unroll 指令控制循环展开（参见 #pragma unroll）。

当使用分支预测时，编译器会执行那些与控制条件相关的所有指令。同时，其中的每条指令都与 per-thread 的状态码或 predicate 相关联，这些状态码或 predicate 会根据控制条件设置为 true 或 false；尽管每个指令都会都会执行，但实际上只有 predicate 为 True 的指令被执行。predicate 为 False 的指令不会写入结果，也不计算地址或读取操作数。

5.4.3 同步指令

对于计算能力为 3.x 的设备，__syncthreads() 的吞吐量为每个时钟周期 128 次操作，对于计算能力为 6.0 的设备，每个时钟周期为 32 次操作，对于计算能力为 7.x 和 8.x 的设备，每个时钟周期为 16 次操作。对于计算能力为 5.x、6.1 和 6.2 的设备，每个时钟周期 64 次操作。

请注意，__syncthreads() 可以会强制 SM 空闲对性能产生影响，如 设备内存访问 中所述。

NCU Profiling Guide

Introduction

性能模型

Roofline Model

https://zhuanlan.zhihu.com/p/34204282 https://people.eecs.berkeley.edu/~kubitron/cs252/handouts/papers/RooflineVyNoYellow.pdf https://zhuanlan.zhihu.com/p/33693725

Roof-line Model 的意义抽象出了：一定计算量和访存量的模型在一定算力和带宽的计算平台所能达到的理论性能上限。

Profile-Driven Optimization Profiling is the act of analyzing program performance by measuring:

• The space (memory) or time complexity of application code
• The use of particular instructions
• The frequency and duration of function calls There are three common limiters to performance for a kernel that you may encounter:
• Memory bandwidth
• Compute resources
• Instruction and memory latency

CUDA 优化

CUDA 性能优化的核心在于最大化硬件利用率和减少各种瓶颈，需从内存访问、计算效率和资源调度三个维度进行系统性优化。以下是关键技术点的深度解析：

一、内存优化（Memory Optimization）

内存访问是 GPU 性能的关键瓶颈，优化目标是减少全局内存访问次数，并提高内存带宽利用率。

1. 全局内存合并访问（Coalesced Access）

• 原理：Warp 内连续线程访问连续内存地址，触发硬件批量传输（如 A100 每次传输 128 字节）。

• 示例（非合并→合并）：

  // 非合并访问（随机索引）
  float val = d_data[threadIdx.x * 4];  // 可能导致多次内存事务
   
  // 合并访问（连续索引）
  float val = d_data[blockIdx.x * blockDim.x + threadIdx.x];  // 最优模式

• 工具检测：Nsight Compute 的 global_hit_rate 指标（理想值接近 100%）。

2. 共享内存 tiling（数据重用）

• 典型场景：矩阵乘法（将全局内存数据分块加载到共享内存）。

3. 内存带宽利用率计算

• 公式：

  实际带宽 = 内存访问量 / 内核执行时间  
  带宽利用率 = 实际带宽 / 硬件峰值带宽（如A100约1.5TB/s）

• 优化目标：利用率 > 80%。若过低，检查是否存在：

  • 非合并访问
  • 过多原子操作（原子操作带宽仅为普通内存访问的 1/4）
  • 内存访问与计算未重叠

二、计算优化（Computation Optimization）

计算优化的核心是最大化计算单元利用率，特别是Tensor Core的高效使用。

1. Tensor Core 编程（以 FP16 矩阵乘为例）

• 原理：Tensor Core 是专为矩阵运算优化的硬件单元，支持 A×B+C 的融合操作（FMA）。
• 性能对比：Tensor Core 的 FP16 计算吞吐量是普通 CUDA Core 的 8 倍以上。

3. 减少分支发散（Branch Divergence）

• 原理：同一 Warp 内线程执行不同分支时，需串行执行各分支（而非并行）。
• 优化示例：
• 工具检测：Nsight Compute 的 branch_efficiency 指标（理想值 100%）。

三、资源利用率优化（Resource Utilization）

优化目标是充分利用 GPU 硬件资源（SM、内存带宽、寄存器等），避免资源闲置。

1. 线程块调度与 Occupancy

• Occupancy 定义：实际活跃 Warp 数与最大可能 Warp 数的比值。

• 计算公式：

  每个SM的最大Warp数 = min(最大线程数/SM, 最大Warp数/SM, 最大寄存器数/SM ÷ 每个线程的寄存器数)
  Occupancy = 实际活跃Warp数 ÷ 每个SM的最大Warp数

• 优化方法：

  1. 使用 nvcc --ptxas-options=-v 查看寄存器使用情况，避免寄存器溢出（如超过 255 / 线程）。
  2. 调整线程块大小（通常 128~512 线程），平衡寄存器压力和 Occupancy。
  3. 使用 Occupancy Calculator 工具（NVIDIA 提供）预测最佳配置。

2. 计算与内存重叠（Overlap）

• 技术：使用 CUDA 流（Stream）将内存操作与计算并行。
• 性能收益：理论上可将内存传输时间隐藏在计算中，实现近 2 倍加速。

3. 共享内存 Bank 冲突优化

• 原理：共享内存分为 32 个 Bank（每个 Bank 带宽为 32 位 / 周期），若多线程访问同一 Bank 会导致串行。

• 优化方法：

  // 有Bank冲突的声明
  __shared__ float data[256];  // 线程0和线程32访问同一Bank
   
  // 无Bank冲突的声明（添加填充）
  __shared__ float data[256 + 32];  // 错开地址，避免同一Bank访问冲突

• 工具检测：Nsight Compute 的 shared_load_throughput 和 shared_store_throughput 指标。

四、性能分析工具链

1. Nsight Compute（内核级分析）

• 关键指标：

  |sm_efficiency|SM 利用率|>80%|

  |achieved_occupancy|实际 Occupancy|接近理论值|

  |gld_efficiency/gst_efficiency|全局加载 / 存储效率|>90%|

  |tensor_core_throughput|Tensor Core 使用率|>70%|

  |branch_efficiency|分支效率|接近 100%|

• 使用示例：

  ncu --set full --metrics sm_efficiency,tensor_core_throughput -o profile ./my_app

2. Nsight Systems（系统级分析）

• 关键视图：
  • Timeline：可视化 CPU/GPU 活动，检查计算与传输是否重叠。
  • CUDA API Calls：分析 API 调用耗时（如 cudaMemcpy 是否阻塞）。
  • Memory Transfers：检查 PCIe 传输带宽利用率。

五、实战优化流程

1. 基准测试：
   • 运行原始代码，记录性能指标（如吞吐量、延迟）。
2. 瓶颈识别：
   • 使用 Nsight Compute 分析内核热点，确定瓶颈类型（内存 / 计算 / 调度）。
3. 针对性优化：
   • 内存瓶颈：应用合并访问、共享内存 tiling。
   • 计算瓶颈：启用 Tensor Core、向量化编程。
   • 调度瓶颈：调整线程块大小、减少同步操作。

• NVIDIA Performance Optimization Guide
• CUDA C++ Programming Guide
• Nsight Compute Documentation