CUDA C++ 编程指北-第四章性能提升指南

imoldpan · 2023 年9 月 9 日 14:36

Overall Performance Optimization Strategies

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

最大化并行执行以实现最大利用率；
优化内存使用以实现最大的内存吞吐量；
优化指令使用以实现最大的指令吞吐量；
最小化内存抖动。

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

Maximize Utilization

为了最大化利用，应用程序应该以尽可能多的并行方式构建，并有效地将这种并行性映射到系统的各个组件，以使它们大部分时间保持忙碌。

应用层级

从高层次来看，应用程序应该使用异步函数调用和流，如异步并发执行中所述，最大化主机、设备以及连接主机和设备的总线之间的并行执行。它应该将每种工作分配给最擅长的处理器：将串行工作负载分配给主机；将并行工作负载分配给设备。

对于并行工作负载，在算法中并行性被打破的点，因为一些线程需要同步以与其他线程共享数据，有两种情况：要么这些线程属于同一个块，这种情况下它们应该使用__syncthreads()并在同一个内核调用中通过共享内存共享数据，要么它们属于不同的块，这种情况下它们必须使用两个单独的内核调用通过全局内存共享数据，一个用于写入，一个用于从全局内存读取。第二种情况不太理想，因为它增加了额外的内核调用和全局内存流量的开销。因此，应该尽量通过将算法映射到CUDA编程模型，使需要进行线程间通信的计算尽可能地在一个线程块内进行。

设备层级

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

多个内核可以在一个设备上并发执行，因此也可以使用流来实现足够的内核并发执行，如异步并发执行中所述。

多处理器层级

在更低的层次上，应用程序应该最大化多处理器内各个功能单元之间的并行执行。

如硬件多线程中所述，GPU多处理器主要依赖线程级并行性来最大化其功能单元的利用率。因此，利用率直接与常驻线程数相关。在每个指令发出时间，一个线程调度器选择一个准备执行其下一指令的线程（即线程的活动线程），并向这些线程发出指令。

特别地，每个多处理器都有一组32位寄存器，这些寄存器在线程之间进行分区，以及一个并行数据缓存或共享内存，这些内存在线程块之间进行分区。

给定内核的每个多处理器上可以驻留和一起处理的块和线程的数量取决于多处理器的内存资源以及内核的资源需求，如硬件多线程中所述。寄存器和共享内存的使用情况在使用--ptxas-options=-v选项编译时由编译器报告。

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

内核使用的寄存器数量对常驻线程的数量有显著影响。例如，对于计算能力为6.x的设备，如果一个内核使用64个寄存器，每个块有512个线程，并且需要很少的共享内存，那么两个块（即32 warps）可以驻留在多处理器上，因为它们需要2x512x64个寄存器，这正好与多处理器上可用的寄存器数量相匹配。但是，一旦内核再多使用一个以上的寄存器，那么只有一个块（即16个warps）可以常驻，因为两个块需要2x512x65个寄存器，这比多处理器上可用的寄存器数量要多。因此，编译器试图在保持寄存器溢出（参见设备内存访问）和指令数量最少的情况下最小化寄存器使用。可以使用maxrregcount编译器选项或启动边界来控制寄存器使用，如启动边界中所述。

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

对于给定内核调用，执行配置对性能的影响通常取决于内核代码。因此，建议进行实验。应用程序还可以根据设备的计算能力、多处理器的数量和内存带宽，以及运行时可以查询的其他参数，对执行配置进行参数化（参见参考手册）。

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

Occupancy Calculator

有几个API函数可以帮助程序员根据寄存器和共享内存的需求来选择线程块大小和集群大小。

占用率计算器API，cudaOccupancyMaxActiveBlocksPerMultiprocessor，可以根据内核的块大小和共享内存使用情况提供占用率预测。此函数按每个多处理器的并发线程块数量报告占用率。
- 注意，这个值可以转换为其他指标。将其乘以每个块的线程数得到每个多处理器的并发线程数；进一步将并发线程数除以每个多处理器的最大线程数，得到占用率的百分比。
基于占用率的启动配置器APIs，cudaOccupancyMaxPotentialBlockSize 和 cudaOccupancyMaxPotentialBlockSizeVariableSMem，通过启发式方法计算达到最大多处理器级占用率的执行配置。
占用率计算器API，cudaOccupancyMaxActiveClusters，可以根据集群大小、块大小和内核的共享内存使用情况提供占用率预测。此函数按照系统中存在的给定大小的最大活动集群的数量报告占用率。

以下代码示例计算了MyKernel的占用率。然后，它使用每个多处理器的并发线程与最大线程之间的比率来报告占用率水平。

// 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;
}

以下代码示例根据用户输入配置了一个基于占用（occupancy-based）的内核启动 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;
}

以下代码示例展示了如何使用集群（Cluster）占用 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版本特别有用，可以将影响占用率的参数的更改可视化（块大小、每线程寄存器和每线程共享内存）。

CUDA C++ 编程指北-第四章 性能提升指南