5. 性能指南-优快云博客

本文链接：https://blog.youkuaiyun.com/qq_62704693/article/details/143352882

博客补充：CUDA C++ 编程指南学习_c++ cuda编程-优快云博客https://blog.youkuaiyun.com/qq_62704693/article/details/141225395?spm=1001.2014.3001.5501

5.1. 整体性能优化策略

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

最大化并行执行以实现最大利用率;
优化内存使用以实现最大内存吞吐量;
优化指令使用以实现最大的指令吞吐量;
最大限度地减少内存抖动。

哪些策略将为应用程序的特定部分产生最佳性能增益，具体取决于该部分的性能限制器;例如，优化主要受内存访问限制的内核的指令使用不会产生任何显著的性能提升。因此，应通过测量和监控性能限制器（例如使用 CUDA 分析器）来不断指导优化工作。此外，将特定内核的浮点运算吞吐量或内存吞吐量（以更有意义的为准）与设备的相应峰值理论吞吐量进行比较，可以表明内核还有很大的改进空间。

5.2. 最大化利用率

为了最大限度地提高利用率，应用程序的结构应尽可能多地公开并行性，并有效地将这种并行性映射到系统的各个组件，以使它们在大部分时间都保持忙碌。

5.2.1. 应用程序级别

概括地说，应用程序应该通过使用异步函数调用和流来最大化主机、设备和将主机连接到设备的总线之间的并行执行，如异步并发执行中所述。它应该为每个处理器分配它最擅长的工作类型：将串行工作负载分配给主机;将工作负载并行到设备。

对于并行工作负载，在算法中由于某些线程需要同步才能相互共享数据而导致并行性中断的点，有两种情况：这些线程属于同一个块，在这种情况下，它们应该通过同一内核调用中的共享内存使用和共享数据，或者它们属于不同的块，在这种情况下，它们必须使用两个单独的内核调用通过全局内存共享数据，一个用于写入全局内存，一个用于读取全局内存。第二种情况不太理想，因为它增加了额外内核调用和全局内存流量的开销。因此，应通过将算法映射到 CUDA 编程模型来最大限度地减少其发生，以便尽可能在单个线程块内执行需要线程间通信的计算。__syncthreads()

5.2.2. 设备级别

在较低级别，应用程序应最大化设备多处理器之间的并行执行。

多个内核可以在一个设备上并发执行，因此也可以通过使用流来启用足够的内核来并发执行，从而实现最大利用率，如异步并发执行中所述。

5.2.3. 多处理器级别

在更低的级别上，应用程序应该最大化 multiprocessor 中各个功能单元之间的并行执行。

如硬件多线程中所述，GPU 多处理器主要依赖于线程级并行性来最大限度地利用其功能单元。因此，利用率与常驻 warp 的数量直接相关。在每个 INSTRUCTION 发出时，warp 调度器都会选择一条准备执行的 INSTRUCTION。此指令可以是同一 warp 的另一个独立指令，利用指令级并行性，或者更常见的是另一个 warp 的指令，利用线程级并行性。如果选择了 ready to execute 指令，则会将其发送到 warp 的活动线程。warp 准备执行下一条指令所需的 clock cycles 数称为 latency，当所有 warp schedulers 在该 latency 期间的每个 clock cycles总是有一些指令要为某个 warp 发出时，或者换句话说，当 latency 完全 “隐藏” 时，就可以实现完全利用。隐藏 L clock cycles 的 latency 所需的指令数量取决于这些指令各自的吞吐量（有关各种算术指令的吞吐量，请参阅 Arithmetic Instructions）。如果我们假设指令具有最大吞吐量，则它等于：

4L 对于计算能力为 5.x、6.1、6.2、7.x 和 8.x 的器件，因为对于这些器件，多处理器在一个 clock cycle上为每个 warp 发出一条指令，一次四个 warp，如 Compute Capabilities 中所述。
2L 对于计算能力为 6.0 的设备，因为对于这些设备，每个周期发出的两条指令是针对两个不同 warp 的一条指令。

warp 没有准备好执行其下一条指令的最常见原因是该指令的 input 操作数尚不可用。

如果所有 input 操作数都是 registers，则延迟是由 register dependencies 引起的，即一些 input 操作数是由一些尚未完成执行的先前指令写入的。在这种情况下，延迟等于前一条指令的执行时间，并且 warp 调度器必须在该时间内调度其他 warp 的指令。执行时间因指令而异。在计算能力为 7.x 的设备上，对于大多数算术指令，通常为 4 个 clock cycles。这意味着每个多处理器需要 16 个活动 warps（4 个周期，4 个 warp 调度器）来隐藏算术指令延迟（假设 warp 以最大吞吐量执行指令，否则需要更少的 warp）。如果单个 warp 表现出指令级并行性，即在其指令流中有多个独立的指令，则需要更少的 warp，因为来自单个 warp 的多个独立指令可以背靠背发出。

如果一些 input 操作数驻留在片外 memory中，则 latency 要高得多：通常为数百个 clock cycles。在如此高延迟期间，使 warp 调度程序保持忙碌所需的 warp 数量取决于内核代码及其指令级并行度。通常，如果没有片外存储器操作数的指令数（即大多数时候是算术指令）与具有片外存储器操作数的指令数之比较低（这个比率通常称为程序的算术强度），则需要更多的 warps。

warp 没有准备好执行下一条指令的另一个原因是它正在等待某个内存围栏（Memory Fence Functions）或同步点（Synchronization Functions）。同步点可以强制多处理器空闲，因为越来越多的 warp 等待同一块中的其他 warp 在同步点之前完成指令的执行。在这种情况下，每个多处理器拥有多个常驻块有助于减少空闲，因为来自不同块的 warp 不需要在同步点相互等待。

对于给定内核调用，驻留在每个多处理器上的块和 warp 的数量取决于调用的执行配置（Execution Configuration）、多处理器的内存资源以及内核的资源要求，如硬件多线程中所述。使用该选项进行编译时，编译器会报告 Register 和 shared memory 使用情况。--ptxas-options=-v

块所需的共享内存总量等于静态分配的共享内存量和动态分配的共享内存量之和。

kernel 使用的 registers 数量对 resident warp 的数量有很大影响。例如，对于计算能力为 6.x 的设备，如果内核使用 64 registers 并且每个块有 512 个线程并且需要很少的共享内存，那么两个块（即 32 个 warp）可以驻留在多处理器上，因为它们需要 2x512x64 registers，这与 multiprocessor 上可用的 registers 数量完全匹配。但是一旦内核多使用一个 register，就只有一个 block（即 16 个 warps）可以 resident 的，因为两个块需要 2x512x65 寄存器，这比多处理器上可用的寄存器多。因此，编译器会尝试最小化 register usage 同时将 register spilling （请参阅 Device Memory Accesses）和指令数量保持在最低限度。寄存器使用情况可以是使用 compiler 选项、Launch Bounds 中描述的限定符或 Maximum Number of Registers per Thread 中描述的限定符进行控制。maxrregcount __launch_bounds__() __maxnreg__()

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

执行配置对给定内核调用性能的影响通常取决于内核代码。因此，建议进行实验。应用程序还可以根据 register 文件大小和共享内存大小参数化执行配置，这取决于器件的计算能力，以及 multiprocessors 的数量和器件的内存带宽，所有这些都可以使用 runtime 进行查询（参见参考手册）。

每个块的线程数应选择 warp 大小的倍数，以避免尽可能浪费计算资源，因为 warp 填充不足。

5.2.3.1. 占用率计算器

存在多个 API 函数，可帮助程序员根据 register 和共享内存要求选择线程块大小和集群大小。

占用率计算器 API 可以根据内核的区块大小和共享内存使用情况提供占用率预测。此函数根据每个多处理器的并发线程块数来报告占用率。cudaOccupancyMaxActiveBlocksPerMultiprocessor
- 请注意，此值可以转换为其他量度。乘以每个块的 warp 数得到每个多处理器的并发 warp 数;进一步将并发 WARPS 除以每个多处理器的最大 WARPS 得到占用率百分比。
基于占用率的启动配置器 API 和启发式计算实现最大多处理器级别占用率的执行配置。cudaOccupancyMaxPotentialBlockSize cudaOccupancyMaxPotentialBlockSizeVariableSMem
占用率计算器 API 可以根据内核的集群大小、数据块大小和共享内存使用情况提供占用率预测。此函数根据系统中存在的 GPU 上给定大小的最大活动集群数来报告占用率。cudaOccupancyMaxActiveClusters

以下代码示例计算 MyKernel 的占用率。然后，它报告占用率以及每个多处理器的并发 warp 与最大 warp 之间的比率。

#include <stdio.h>
#include <cuda_runtime.h>
#include <iostream>
// 设备代码
__global__ void MyKernel(int *d, int *a, int *b)
{
    // 计算全局线程索引
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    
    // 执行元素级乘法并将结果存储在数组d中
    d[idx] = a[idx] * b[idx];
}

// 主机代码
int main()
{
    int numBlocks;        // 以占用率表示的活动块数量
    int blockSize = 32;   // 每个块中的线程数量

    // 这些变量用于将占用率转换为warp数量
    int device;           // 用于保存设备ID的变量
    cudaDeviceProp prop;  // 用于保存设备属性的结构
    int activeWarps;      // 根据占用率计算的活动warp数量
    int maxWarps;         // 每个多处理器支持的最大warp数量

    // 获取当前设备ID
    cudaGetDevice(&device);
    
    // 获取当前设备的属性
    cudaGetDeviceProperties(&prop, device);

    // 计算MyKernel每个多处理器的最大活动块数
    cudaOccupancyMaxActiveBlocksPerMultiprocessor(
        &numBlocks,        // 输出：活动块数量
        MyKernel,          // 要分析的内核函数
        blockSize,         // 每个块中的线程数量
        0);                // 共享内存大小（0表示不使用共享内存）

    // 根据活动块的数量计算活动warp的数量
    activeWarps = numBlocks * blockSize / prop.warpSize;

    // 计算设备可以支持的最大warp数量
    maxWarps = prop.maxThreadsPerMultiProcessor / prop.warpSize;

    // 输出占用率百分比
    std::cout << "占用率: " << (double)activeWarps / maxWarps * 100 << "%" << std::endl;

    return 0; // 表示成功结束
}

运行结果：

(base) wkj@ubuntu-GPU2:~/wkj$ ./occupancy
占用率: 50%

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

#include <iostream>
#include <cuda_runtime.h>

// 设备代码
__global__ void MyKernel(int *array, int arrayCount)
{
    // 计算全局线程索引
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    // 确保索引在数组范围内
    if (idx < arrayCount) {
        // 将数组元素平方
        array[idx] *= array[idx];
    }
}

// 主机代码
int launchMyKernel(int *array, int arrayCount)
{
    int blockSize;      // 启动配置器返回的块大小
    int minGridSize;    // 为了实现满设备占用率所需的最小网格大小
    int gridSize;       // 基于输入大小所需的实际网格大小

    // 计算潜在的最大活动块大小和最小网格大小
    cudaOccupancyMaxPotentialBlockSize(
        &minGridSize,     // 输出：最小网格大小
        &blockSize,       // 输出：块大小
        (void*)MyKernel,  // 要分析的内核函数
        0,                // 共享内存大小（0表示不使用共享内存）
        arrayCount);      // 输入数组的大小

    // 根据数组大小向上取整计算网格大小
    gridSize = (arrayCount + blockSize - 1) / blockSize;

    // 启动内核
    MyKernel<<<gridSize, blockSize>>>(array, arrayCount);
    
    // 等待设备完成所有先前的任务
    cudaDeviceSynchronize();

    return 0; // 表示成功结束
}

// 主函数
int main()
{
    const int arrayCount = 1024; // 数组大小
    int *h_array = new int[arrayCount]; // 主机数组
    int *d_array; // 设备数组

    // 初始化主机数组
    for (int i = 0; i < arrayCount; i++) {
        h_array[i] = i + 1; // 示例数据
    }

    // 分配设备内存
    cudaMalloc((void**)&d_array, arrayCount * sizeof(int));
    
    // 将主机数组复制到设备数组
    cudaMemcpy(d_array, h_array, arrayCount * sizeof(int), cudaMemcpyHostToDevice);

    // 启动内核
    launchMyKernel(d_array, arrayCount);

    // 将结果从设备复制回主机
    cudaMemcpy(h_array, d_array, arrayCount * sizeof(int), cudaMemcpyDeviceToHost);

    // 打印结果
    for (int i = 0; i < arrayCount; i++) {
        std::cout << h_array[i] << " "; // 打印每个元素的平方
    }
    std::cout << std::endl;

    // 释放设备和主机内存
    cudaFree(d_array);
    delete[] h_array;

    return 0; // 表示成功结束
}

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

群集大小为 8 是向前兼容的起始计算能力 9.0，但在 GPU 硬件或 MIG 配置上，它们太小而无法支持 8 个多处理器，在这种情况下，最大群集大小将减小。但建议用户在启动集群内核之前查询最大集群大小。可以使用 API 查询最大集群大小。cudaOccupancyMaxPotentialClusterSize

{
  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 计算用户界面还为任何无法依赖 CUDA 软件堆栈的使用案例提供独立的占用计算器和启动配置器实现。占用率计算器的 Nsight Compute 版本作为学习工具特别有用，它可以可视化影响占用率的参数（块大小、每个线程的寄存器数和每个线程的共享内存）的变化的影响。<CUDA_Toolkit_Path>/include/cuda_occupancy.h

5.3. 最大化内存吞吐量

最大化应用程序的整体内存吞吐量的第一步是最小化低带宽的数据传输。

这意味着最大限度地减少主机和设备之间的数据传输，如主机和设备之间的数据传输中所述，因为它们的带宽比全局内存和设备之间的数据传输低得多。

这也意味着通过最大限度地利用片上内存来最大限度地减少全局内存和设备之间的数据传输：共享内存和缓存（即，计算能力为 2.x 及更高版本的设备上可用的 L1 缓存和 L2 缓存，所有设备上都可用的纹理缓存和常量缓存）。

共享内存等同于用户管理的缓存：应用程序显式分配和访问它。如 CUDA 运行时所示，典型的编程模式是将来自设备内存的数据暂存到共享内存中;换句话说，要拥有块的每个线程：

将数据从设备内存加载到共享内存，
与块的所有其他线程同步，以便每个线程都可以安全地读取由不同线程填充的共享内存位置。
处理共享内存中的数据，
如有必要，请再次同步，以确保共享内存已使用结果进行更新。
将结果写回设备内存。

对于某些应用程序（例如，全局内存访问模式依赖于数据），传统的硬件管理高速缓存更适合利用数据局部性。如计算能力 7.x、计算能力 8.x 和计算能力 9.0 中所述，对于计算能力 7.x、8.x 和 9.0 的设备，L1 和共享内存使用相同的片上内存，并且专用于 L1 和共享内存的量可针对每个内核调用进行配置。

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

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

应用程序应努力最大程度地减少主机和设备之间的数据传输。实现此目的的一种方法是将更多代码从主机移动到设备，即使这意味着运行的内核没有公开足够的并行度，无法在设备上高效执行。中间数据结构可以在设备内存中创建，由设备操作，并在不被主机映射或复制到主机内存的情况下销毁。

此外，由于与每次传输相关的开销，将许多小型传输批处理为单个大型传输始终比单独进行每个传输效果更好。

在具有前端总线的系统上，通过使用页面锁定主机内存可实现主机和设备之间数据传输的更高性能，如页面锁定主机内存中所述。

此外，当使用 mapped page-locked memory （Mapped Memory）时，无需分配任何器件内存，也无需在器件和主机存储器之间显式复制数据。每次内核访问 Map 内存时，都会隐式执行数据传输。为了获得最佳性能，这些内存访问必须与对全局内存的访问一样合并（请参阅Device Memory Accesses）。假设它们是，并且映射的内存只被读取或写入一次，那么在设备和主机内存之间使用映射的页面锁定内存而不是显式副本可以提高性能。

在 device memory 和 host memory 物理相同的集成系统上，host 和 device memory 之间的任何复制都是多余的，应该使用 mapped pagelocked memory。应用程序可以通过检查集成设备属性（请参阅设备枚举）是否等于 1 来查询设备。integrated

5.3.2. 设备内存访问

访问可寻址内存（即全局、本地、共享、常量或纹理内存）的指令可能需要多次重新发出，具体取决于内存地址在 warp 内线程之间的分布。分布如何以这种方式影响 instruction 吞吐量特定于每种类型的内存，并在以下部分中介绍。例如，对于全局内存，作为一般规则，地址越分散，吞吐量越低。

全局内存

全局内存驻留在器件内存中，器件内存通过 32 字节、64 字节或 128 字节的内存事务访问。这些 memory transactions 必须自然对齐：只有与其大小对齐的 32 字节、64 字节或 128 字节的器件内存段（即，其第一个地址是其大小的倍数）才能被 memory transactions 读取或写入。

当 warp 执行访问全局内存的指令时，它会根据每个线程访问的字的大小和内存地址在线程之间的分布，将 warp 中线程的内存访问合并到一个或多个这些内存事务中。一般来说，需要的事务越多，除了线程访问的字之外，传输的未使用的字就越多，从而相应地降低了指令吞吐量。例如，如果为每个线程的 4 字节访问生成 32 字节的内存事务，则吞吐量除以 8。

需要多少事务以及最终影响的吞吐量量因设备的计算能力而异。计算能力 5.x、计算能力 6.x、计算能力 7.x、计算能力 8.x 和计算能力 9.0 提供了有关如何为各种计算能力处理全局内存访问的更多详细信息。

因此，为了最大化全局内存吞吐量，必须通过以下方式最大化合并：

遵循基于计算能力 5.x、计算能力 6.x、计算能力 7.x、计算能力 8.x 和计算能力 9.0 的最佳访问模式
使用满足以下大小和对齐要求部分中详述的大小和对齐要求的数据类型，
在某些情况下填充数据，例如，在访问二维数组时，如下面的 Two-Dimensional Arrays 部分所述。

尺寸和对齐要求

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

如果未满足此大小和对齐要求，则访问将编译为多个指令，这些指令具有交错访问模式，从而阻止这些指令完全合并。因此，建议对驻留在全局内存中的数据使用满足此要求的类型。

内置向量类型会自动满足对齐要求。

对于结构，编译器可以使用对齐说明符（例如__align__(8) or __align__(16)

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

或

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

驻留在全局内存中或由驱动程序或运行时 API 的内存分配例程之一返回的变量的任何地址始终与至少 256 字节对齐。

读取非自然对齐的 8 字节或 16 字节字会产生不正确的结果（相差几个字），因此必须特别小心保持这些类型的任何值或值数组的起始地址对齐。一个可能很容易被忽视的典型情况是使用一些自定义的全局内存分配方案，其中多个数组的分配（多次调用 or ）被分区为多个数组的单个大内存块的分配所取代，在这种情况下，每个数组的起始地址都与块的起始地址偏移。cudaMalloc()cuMemAlloc()

二维数组

常见的全局内存访问模式是，当索引的每个线程使用以下地址来访问宽度为 2D 数组的一个元素时，该数组位于类型的地址处（满足最大化利用率中描述的要求）：(tx,ty) width BaseAddress type* type

BaseAddress + width * ty + tx

要使这些访问完全合并，thread 块的宽度和数组的宽度都必须是 warp 大小的倍数。

特别是，这意味着如果数组的宽度不是此大小的倍数，并且实际分配的宽度四舍五入到此大小的最接近倍数，并且其行相应地填充，则访问该数组的效率将大大提高。参考手册中描述的 and 函数和相关的内存复制函数使程序员能够编写不依赖于硬件的代码来分配符合这些约束的数组。cudaMallocPitch()cuMemAllocPitch()

本地内存

本地内存访问仅针对某些自动变量发生，如 Variable Memory Space Specifiers 中所述。编译器可能放置在本地内存中的自动变量是：

无法确定它们是否使用常量进行索引的数组，
会占用太多寄存器空间的大型结构体或数组，
如果 kernel 使用的 registers 多于可用 registers，则为任何变量（这也称为 register spilling）。

检查 PTX 汇编代码（通过使用 or 选项进行编译获得）将判断变量是否在第一个编译阶段被放置在本地内存中，因为它将使用助记符声明并使用和助记符进行访问。即使没有，后续编译阶段可能仍会做出其他决定，尽管如果他们发现它为目标架构占用了太多的寄存器空间：检查 cubin 对象 using 将判断是否是这种情况。此外，在使用选项进行编译时，编译器会报告每个内核（）的总本地内存使用量。请注意，某些数学函数具有可能访问本地内存的实现路径。-ptx -keep .local ld.local st.local cuobjdump lmem --ptxas-options=-v

本地内存空间驻留在设备内存中，因此本地内存访问与全局内存访问具有相同的高延迟和低带宽，并且受与设备内存访问中所述的内存合并要求相同。但是，本地内存的组织方式是连续的 32 位字由连续的线程 ID 访问。因此，只要 warp 中的所有线程访问相同的相对地址（例如，数组变量中的相同索引，结构变量中的相同成员），访问就会完全合并。

在计算能力 5.x 及更高版本的设备上，本地内存访问始终以与全局内存访问相同的方式缓存在 L2 中（请参阅计算能力 5.x 和计算能力 6.x）。

共享内存

由于共享内存是片上存储器，因此与本地或全局存储器相比，它具有更高的带宽和更低的延迟。

为了实现高带宽，共享内存被划分为大小相等的内存模块，称为 bank，可以同时访问。因此，由属于 n 个不同内存条的 n 个地址发出的任何内存读取或写入请求都可以同时处理，从而产生比单个模块带宽高 n 倍的总带宽。

但是，如果内存请求的两个地址位于同一个内存 bank，则存在 bank 冲突，并且必须序列化访问。硬件根据需要将具有存储区冲突的内存请求拆分为多个单独的无冲突请求，从而将吞吐量降低一个系数，其降低量等于单独内存请求的数量。如果单独内存请求的数量为 n，则称初始内存请求会导致 n 路组冲突。

因此，为了获得最佳性能，了解内存地址如何映射到内存 bank 非常重要，以便调度内存请求，从而最大限度地减少 bank 冲突。计算能力 5.x、计算能力 6.x、计算能力 7.x、计算能力 8.x 和计算能力 9.0 分别在计算能力 5.x、计算能力 6.x、计算能力 7.x、计算能力 8.x 和计算能力 9.0 中进行了介绍。

常量内存

常量内存空间驻留在设备内存中，并缓存在常量缓存中。

然后，请求被拆分为与初始请求中不同的内存地址一样多的单独请求，从而将吞吐量降低一个等于单独请求数量的系数。

然后，在缓存命中的情况下，以常量缓存的吞吐量为生成的请求提供服务，否则以设备内存的吞吐量为结果请求提供服务。

纹理和表面内存

纹理和表面内存空间驻留在设备内存中，并缓存在纹理缓存中，因此纹理获取或表面读取仅在缓存未命中时消耗从设备内存中读取一次内存，否则只需从纹理缓存中读取一次内存。纹理缓存针对 2D 空间局部性进行了优化，因此读取在 2D 中彼此靠近的纹理或表面地址的相同 warp 线程将获得最佳性能。此外，它还专为具有恒定延迟的流式提取而设计;缓存命中会减少 DRAM 带宽需求，但不会减少获取延迟。

通过纹理或表面获取读取设备内存具有一些好处，使其成为从全局内存或常量内存读取设备内存的有利替代方案：

如果内存读取不遵循全局或常量内存读取必须遵循的访问模式才能获得良好的性能，则只要纹理提取或表面读取中存在局部性，就可以实现更高的带宽;
寻址计算由专用单元在内核外部执行;
打包的数据可以在单个操作中广播到单独的变量;
可以选择将 8 位和 16 位整数输入数据转换为 [0.0， 1.0] 或 [-1.0， 1.0] 范围内的 32 位浮点值（请参阅纹理内存）。

5.4. 最大化指令吞吐量

为了最大限度地提高指令吞吐量，应用程序应该：

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

在本节中，吞吐量以每个 multiprocessor每个 clock cycle 的 operations 数给出。对于 warp 大小 32，一条指令对应 32 次操作，因此如果 N 是每个 clock cycle的操作数，则指令吞吐量为每 clock cycles N/32 个指令。

所有吞吐量都适用于一个多处理器。它们必须乘以设备中 multiprocessor 的数量，才能获得整个设备的吞吐量。

5.4.1. 算术指令

下表提供了各种计算能力的设备的硬件中本机支持的算术指令的吞吐量。

表 4 本机算术指令的吞吐量。（Number of Results per Clock Cycle per Multiprocessor）
计算能力	5.0, 5.2	5.3	6.0	6.1	6.2	7.x	8.0	8.6	8.9	9.0
16 位浮点加、乘、加	不适用	256	128	2	256	128	2563		128	256
32 位浮点加、乘、加	128		64	128		64		128
64 位浮点加、乘、加	4		32	4		325	32	2		64
32 位浮点倒数、倒数平方根、以 2 为底的对数（）、以 2 为底的指数（）、正弦（）、余弦（`__log2f` `exp2f` `__sinf` `__cosf`)	32		16	32		16
32 位整数加法、扩展精度加法、减法、扩展精度减法	128		64	128		64
32 位整数乘法、乘法加法、扩展精度乘法加法	多个指示。					646
24 位整数乘法（`__[u]mul24`)	多个指示。
32 位整数移位	64		32	64
比较、最小值、最大值	64		32	64
32 位整数位反向	64		32	64		16
位域提取/插入	64		32	64		多个指令。			64
32 位按位 AND、OR、XOR	128		64	128		64
前导零计数，最高有效无符号位	32		16	32		16
人口计数	32		16	32		16
Warp Shuffle （扭曲随机播放）	32					328	32
Warp Reduce（变形减少）	多个指示。						16
扭曲投票	64
绝对差值之和	64		32	64
SIMD 视频说明`vabsdiff2`	多个指示。
SIMD 视频说明`vabsdiff4`	多个指示。					64
所有其他 SIMD 视频说明	多个指示。
从 8 位和 16 位整数类型到 32 位整数类型的类型转换	32		16	32		64
与 64 位类型之间的类型转换	4		16	4		1610	16	2	2	16
所有其他类型转换	32		16	32		16
16 位 DPX	多个指示。									128
32 位 DPX	多个指示。									64

其他指令和函数是在本机指令之上实现的。对于不同计算能力的设备，实现可能会有所不同，编译后的 native 指令数量可能会随着每个编译器版本的变化而波动。对于复杂的函数，根据 Input，可以有多个代码路径。可用于检查对象中的特定实现。cuobjdump cubin

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

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

单精度浮点除法

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

单精度浮点倒数平方根

为了保留 IEEE-754 语义，编译器只有在倒数和平方根都是近似值时才能优化成（即，使用和）。因此，建议在需要时直接调用。1.0/sqrtf() rsqrtf() -prec-div=false -prec-sqrt=false rsqrtf()

单精度浮点平方根

单精度浮点平方根实现为倒数平方根，后跟倒数而不是倒数平方根，后跟乘法，以便给出 0 和无穷大的正确结果。

正弦和余弦

sinf(x)，，，，和相应的双精度指令要昂贵得多，如果参数 x 的大小很大，则更是如此。cosf(x) tanf(x) sincosf(x)

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

快速路径用于大小足够小的参数，主要由一些乘加运算组成。慢速路径用于大小较大的参数，并且包括在整个参数范围内获得正确结果所需的冗长计算。

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

由于慢速路径比快速路径需要更多的 registers，因此已尝试通过在本地内存中存储一些中间变量来降低慢速路径中的 register 压力，这可能会因为本地内存的高延迟和带宽而影响性能（参见 Device Memory Accesses）。目前，单精度函数使用 28 字节的本地内存，双精度函数使用 44 字节的本地内存。但是，确切的金额可能会发生变化。

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

整数算术

整数除法和模运算成本高昂，因为它们最多可编译 20 条指令。在某些情况下，它们可以替换为按位运算：如果是 2 的幂，则（）等效于（）且等效于（）;如果为 Literals，编译器将执行这些转换。n i/n (i>>log2(n)) (i%n) i&(n-1) n

__brev和映射到单个指令和和几个指令。__popc __brevll __popcll

__[u]mul24是不再有任何理由使用的旧式内部函数。

半精度算术

为了获得 16 位精度浮点加、乘或乘加的良好性能，建议将数据类型用于精度，并将 datatype 用于精度。然后，可以使用向量内部函数（例如，，，，）在单个指令中执行两个操作。使用 OR 代替两个调用使用 OR 也可能有助于提高其他内部函数的性能，例如 warp shuffle。half2 half __nv_bfloat162 __nv_bfloat16 __hadd2 __hsub2 __hmul2 __hfma2 half2 __nv_bfloat162 half __nv_bfloat16

提供内部函数用于将两个精度值转换为数据类型。__halves2half2 half half2

提供内部函数用于将两个精度值转换为数据类型。__halves2bfloat162 __nv_bfloat __nv_bfloat162

类型转换

有时，编译器必须插入 conversion 指令，从而引入额外的执行周期。这种情况适用于：

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

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

5.4.2. 控制流说明

任何流控制指令（，，，，）都可以通过导致同一 warp 的线程发散（即遵循不同的执行路径）来显着影响有效指令吞吐量。如果发生这种情况，则必须序列化不同的执行路径，从而增加为此 warp 执行的指令总数。if switch do for while

为了在控制流依赖于线程 ID 的情况下获得最佳性能，应编写控制条件，以最大程度地减少发散 warp 的数量。这是可能的，因为 warp 在 block 中的分布是确定性的，如 SIMT 架构中所述。一个简单的例子是当控制条件仅取决于（）时，其中是 warp 大小。在这种情况下，没有 warp 发散，因为控制条件与 warp 完全一致。threadIdx / warpSize warpSize

有时，编译器可能会展开循环，或者可能会通过使用分支谓词来优化 short 或 blocks，如下所述。在这些情况下，任何 warp 都不会发散。程序员还可以使用指令控制 loop 展开（参见 #pragma unroll）。if switch #pragma unroll

使用 branch predication 时，不会跳过其执行取决于控制条件的任何指令。相反，它们中的每一个都与一个基于控制条件设置为 true 或 false 的每线程条件代码或谓词相关联，尽管这些指令中的每一个都被安排执行，但实际上只有具有 true 谓词的指令才会被执行。带有 false 谓词的指令不会写入结果，也不会计算地址或读取操作数。

5.4.3. 同步指令

对于计算能力为 6.0 的器件，吞吐量为每 clock cycle 32 次操作，对于计算能力为 7.x 和 8.x 的器件，每个 clock cycle 为 16 次操作，对于计算能力为 5.x、6.1 和 6.2 的器件，每个 clock cycle 为 64 次操作。__syncthreads()

请注意，这可能会强制多处理器空闲，如 Device Memory Accesses中所述，从而影响性能。__syncthreads()

5.5. 最小化内存抖动

过于频繁地不断分配和释放内存的应用程序可能会发现，分配调用往往会随着时间的推移而变慢，直到达到限制。这通常是意料之中的，因为将内存释放回操作系统供其自己使用的性质。为了获得这方面的最佳性能，我们建议执行以下操作：

尝试根据手头的问题调整分配的大小。不要尝试使用 / / 分配所有可用内存，因为这会强制内存立即驻留并阻止其他应用程序使用该内存。这可能会给操作系统调度程序带来更大的压力，或者只是阻止使用同一 GPU 的其他应用程序完全运行。cudaMalloc cudaMallocHost cuMemCreate
尝试在应用程序的早期以适当大小的分配方式分配内存，并且仅在应用程序没有任何用途时分配内存。减少应用程序中的 + 调用次数，尤其是在性能关键型区域中。cudaMalloc cudaFree
如果应用程序无法分配足够的设备内存，请考虑回退到其他内存类型，例如或，这些类型的性能可能不那么好，但会使应用程序能够继续前进。cudaMallocHost cudaMallocManaged
对于支持该功能的平台，允许超额订阅，并且启用了正确的策略，将允许应用程序保留大部分（如果不是全部）性能。也不会强制将分配驻留，直到需要或预取它，从而减少操作系统调度程序的整体压力，并更好地支持多原则用例。cudaMallocManaged cudaMemAdvise cudaMalloc cudaMallocManaged