【CUDA编程】系列博客参考NVIDIA官方文档“CUDA C++ Programming Guide（v12.6）”。
本文为原创文章，未经本人允许，禁止转载。转载请注明出处。

1.Overall Performance Optimization Strategies

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

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

对于应用程序的特定部分，哪种策略能带来最佳的性能提升取决于该部分的性能限制因素。例如，针对一个主要受内存访问限制的kernel优化指令使用，通常不会带来显著的性能提升。因此，优化工作应该始终通过测量和监控性能限制因素来指导，例如使用CUDA性能分析工具（CUDA profiler）。此外，将特定kernel的浮点操作吞吐量或内存吞吐量（根据实际情况选择更相关的指标）与device的理论峰值吞吐量进行比较，可以指示该kernel还有多大的优化空间。

2.Maximize Utilization

为了最大化利用率，应用程序的结构设计应尽可能多地利用并行性，并将这种并行性高效地映射到系统的各个组件上，从而使这些组件大部分时间都能保持忙碌状态。

2.1.Application Level

high level的来看，应用程序应通过使用异步函数调用和stream（详见：Asynchronous Concurrent Execution），在host、device以及连接host和device的主线（bus）之间最大化并行执行。同时，应用程序应将各处理器最擅长的任务分配给它们：串行任务分配给host，并行任务分配给device。

对于并行任务，因为某些线程需要同步以便彼此共享数据，从而导致并行性被破坏，这包含两种情况：

如果这些线程属于同一个线程块，它们应该使用__syncthreads()并通过共享内存在同一个kernel调用中共享数据。
如果这些线程属于不同的线程块，它们必须通过全局内存共享数据，这需要两次单独的kernel调用：一次用于写入全局内存，另一次用于从全局内存中读取数据。

第二种情况效率较低，因为它增加了额外的kernel调用和全局内存访问的开销。因此，应尽量通过将算法映射到CUDA编程模型，确保需要线程间通信的计算尽可能在单个线程块内完成，以减少这种情况的发生。

2.2.Device Level

从low level来说，应用程序应在device的多处理器之间最大化并行执行。

多个kernel可以在同一device上并发执行，因此可以通过使用stream来启用足够多的kernel实现并发，从而达到最大化利用率。

2.3.Multiprocessor Level

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

正如在Hardware Multithreading中所描述的那样，GPU多处理器主要依靠线程级并行性来最大化其功能单元的利用率。因此，利用率直接与驻留的warp数量相关。在每个指令发出时，warp调度器会选择一条准备好执行的指令。这条指令可以是同一个warp中另一条独立指令（利用指令级并行性），更常见的是另一个warp中的一条指令（利用线程级并行性）。如果选中的指令已准备好执行，则它会被分配到warp的活跃线程中。

这里简单再解释下warp调度器。一个SM可能有多个warp调度器，比如在TU102 GPU芯片中，一个SM有4个warp调度器，这意味着在一个时钟周期中，这4个warp调度器最多可以同时给4个warp分配指令。

一个warp准备好执行其下一条指令所需的时钟周期数称为延迟（latency）。如果在延迟期间的每个时钟周期中，所有warp调度器都有要发出的指令且刚好也都有准备好的warp可以分配，那就可以实现完全利用率，换句话说，延迟被“隐藏”了。要隐藏$L$个时钟周期的延迟所需的指令数量取决于这些指令相应的吞吐量（有关各种算术指令的吞吐量，请见第4.1部分）。如果假设指令具有最大吞吐量，则公式如下：

对于计算能力为5.x、6.1、6.2、7.x和8.x的device，由于这些device的多处理器在每个时钟周期最多可以同时向4个warp发送指令，所以隐藏延迟需要$4L$个指令。
对于计算能力为6.0的device，由于这些device的多处理器在每个时钟周期最多可以同时向2个warp发送指令，所以隐藏延迟需要$2L$个指令。

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

如果所有输入操作数都来自寄存器，那么延迟是由寄存器依赖性（register dependencies）引起的。也就是说，这些输入操作数中的某些值是由之前的指令写入的，而这些指令尚未完成执行。在这种情况下，延迟等于之前指令的执行时间，warp调度器必须在此期间调度其他warp的指令来掩盖这种延迟。指令的执行时间取决于具体的指令类型。在计算能力7.x的device上，大多数算术指令的执行时间通常为4个时钟周期。这意味着，要隐藏算术指令的延迟，每个多处理器需要有16个活跃的warp（4个时钟周期$\times$4个warp调度器），这里假设指令具有最大吞吐量，否则需要的warp数量会减少。如果某个warp本身具有指令级并行性（即指令流中有多条独立指令），那么所需的warp数量会减少，因为来自同一warp的多条独立指令可以连续发出，而不需要切换到其他warp。

如果某些输入操作数存储在off-chip memory中，那么延迟会显著提高，通常达到数百个时钟周期。在这样的高延迟期间，为了让warp调度器始终保持忙碌，需要的warp数量取决于kernel代码以及代码中指令级并行性的程度。我们将非off-chip memory的指令（比如算术指令）与off-chip memory指令的数量比值称为程序的算术强度（arithmetic intensity），通常来说，算术强度越低，就需要更多的warp来保持调度器的忙碌状态。

warp未准备好执行其下一条指令的另一个原因是它正在等待某个内存屏障（memory fence）或同步点（synchronization point）。同步点可能会导致多处理器逐渐进入空闲状态，因为越来越多的warp可能需要等待同一线程块中其他warp完成同步点之前的指令。在这种情况下，让每个多处理器同时驻留多个线程块（resident blocks）可以帮助减少空闲时间，因为来自不同线程块的warp不需要在同步点处相互等待。

对于一次kernel调用，每个多处理器上驻留的线程块和warp数量取决于以下因素：

调用时的执行配置。
多处理器的内存资源。
kernel的资源需求（见：Hardware Multithreading）。

在使用--ptxas-options=-v编译选项编译代码时，编译器会报告寄存器和共享内存的使用情况。

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

一个kernel使用的寄存器数量会对驻留warp的数量产生显著影响。例如，对于计算能力为6.x的device，如果一个kernel使用了64个寄存器，且每个线程块包含512个线程，并且需要的共享内存很少，那么多处理器上可以驻留两个线程块（即32个warp），因为它们需要的寄存器数量为$2 \times 512 \times 64$，正好等于多处理器可用的寄存器总数。但是，如果kernel再多使用一个寄存器（即使用65个寄存器），那么只能驻留一个线程块（即16个warp），因为两个线程块将需要$2 \times 512 \times 65$个寄存器，这超出了多处理器可用的寄存器数量。因此，编译器会尝试在减少寄存器使用的同时，尽量避免寄存器溢出（详见第3.2部分）并减少指令数量。这里简单解释下，单个线程使用的寄存器越多，其执行效率可能就越高，但寄存器使用过多会减少驻留warp的数量，从而降低并发度，影响硬件资源的利用率。如果发生寄存器溢出，则超出部分的数据会被溢出到device内存（比如全局内存），访问device内存的延迟非常高，可能会显著降低性能。可以通过使用maxrregcount编译器选项、__launch_bounds__()限定符或__maxnreg__()限定符来控制寄存器的使用。

寄存器文件以32位寄存器为单位。因此，每个变量至少需要一个32位寄存器。例如，一个double类型的变量需要占用两个32位寄存器。

对于一个kernel调用，执行配置（execution configuration）对性能的影响通常取决于kernel代码本身。因此，推荐通过实验来确定最佳配置。此外，应用程序还可以根据以下参数动态调整执行配置：

寄存器文件大小和共享内存大小，这些参数取决于device的计算能力。
多处理器数量和device的内存带宽。

线程块中的线程数量应选择为warp大小的整数倍，以尽量避免因未完全填满的warp而浪费计算资源。

2.3.1.Occupancy Calculator

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

占用率计算器（occupancy calculator）API：cudaOccupancyMaxActiveBlocksPerMultiprocessor，可以根据kernel的线程块大小和共享内存使用情况提供占用率预测。此函数以每个多处理器上的并发线程块数量为单位报告占用率。
- 需要注意的是，这个值可以转换为其他指标。与线程块中的warp数相乘可以得到每个多处理器的并发warp数量。将并发warp数量除以多处理器支持的最大warp数量，可以得出占用率的百分比形式。
基于占用率的启动配置API：cudaOccupancyMaxPotentialBlockSize和cudaOccupancyMaxPotentialBlockSizeVariableSMem，使用启发式算法计算可以实现最大多处理器占用率的执行配置。
占用率计算器API：cudaOccupancyMaxActiveClusters，可以根据cluster大小、线程块大小和共享内存使用情况提供占用率预测。该API用于计算目标device上可以并发执行的最大活跃cluster数量。

以下代码示例计算了MyKernel的占用率。它通过计算并发warp的数量与每个多处理器支持的最大warp数量的比值，来报告占用率水平。

// 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); //返回正在使用的device
    cudaGetDeviceProperties(&prop, device); //返回该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;
}

cudaOccupancyMaxActiveBlocksPerMultiprocessor函数的解释：

//函数用途：返回一个device函数的占用率。
//参数解释：
//1.numBlocks：返回的占用率，即并发的线程块数量。
//2.func：需要计算占用率的kernel函数。
//3.blockSize：kernel设置的线程块的大小。
//4.dynamicSMemSize：每个线程块预期使用的动态共享内存，单位为字节。
template < class T > __host__cudaError_t cudaOccupancyMaxActiveBlocksPerMultiprocessor (
    int *numBlocks, 
    T func, 
    int blockSize, 
    size_t dynamicSMemSize
)

以下代码示例根据用户输入，配置了一个基于占用率的kernel启动。

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

cudaOccupancyMaxPotentialBlockSize函数的解释：

//函数用途：给定一个device函数，返回能够使其达到最大潜在占用率的grid size和block size。
//1.minGridSize：返回能够使其达到最佳潜在占用率的最小grid size。
//2.blockSize：返回block size。
//3.func：需要计算的kernel函数。
//4.dynamicSMemSize：每个线程块预期使用的动态共享内存，单位为字节。
//5.blockSizeLimit：给func设置的最大线程数。0表示没有限制。
template < class T > __host__cudaError_t cudaOccupancyMaxPotentialBlockSize (
    int *minGridSize,
    int *blockSize, 
    T func, 
    size_t dynamicSMemSize, 
    int blockSizeLimit
)

以下代码示例展示了如何使用cluster占用率API来计算在给定大小的cluster中最大的活跃cluster数量。以下代码示例计算了cluster大小为2且每个线程块有128个线程时的占用率。

从计算能力9.0的device开始，cluster大小为8是向前兼容的。但在某些情况下，比如GPU硬件或MIG配置太小，从而无法支持8个多处理器时，最大cluster大小将会被减小。因此，在用户启动cluster kernel之前，推荐先查询最大cluster大小。最大cluster大小可以通过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;
}

cudaOccupancyMaxPotentialClusterSize函数的解释：

//函数用途：给定kernel函数（func）和启动配置（config），返回最大的cluster size
template < class T > __host__cudaError_t cudaOccupancyMaxPotentialClusterSize (
    int *clusterSize,
    T *func, 
    const cudaLaunchConfig_t *config
)

CUDA Nsight Compute用户界面还在<CUDA_Toolkit_Path>/include/cuda_occupancy.h中提供了一个独立的占用率计算器和启动配置工具实现，适用于任何无法依赖CUDA软件栈的场景。Nsight Compute版本的占用率计算器特别有用，作为一种学习工具，它可以直观地显示影响占用率的参数变化（例如线程块大小、每个线程寄存器数量、每个线程的共享内存）的影响。

3.Maximize Memory Throughput

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

这意味着要尽量减少host和device之间的数据传输，因为与全局内存和device之间的数据传输相比，host和device之间的传输带宽要低得多，详见第3.1部分。

这也意味着要通过最大限度地利用on-chip memory来减少全局内存与device之间地数据传输。其中，on-chip memory包括共享内存和cache（比如在计算能力2.x及以上device中支持的L1 cache和L2 cache，以及在所有device中都支持的texture cache和constant cache）。

共享内存相当于是一种由用户管理的cache：应用程序需要显式的分配和访问它。一个典型的编程模式是将device内存中的数据调入共享内存。换句话说，让线程块中的每个线程完成以下操作：

从device内存加载数据到共享内存。
与线程块中的所有其他线程同步，以确保每个线程能够安全地读取由其他线程填充的共享内存位置。
在共享内存中处理数据。
如果有必要，再次同步以确保共享内存中的数据已被更新。
将结果写回到device内存。

kernel的内存访问吞吐量可能会因每种内存类型的访问模式不同而相差一个数量级。因此，最大化内存吞吐量的下一步是根据最佳内存访问模式（详见第3.2部分），尽可能高效的组织内存访问。这种优化对于访问全局内存尤为重要，因为与on-chip memory的带宽和算术指令吞吐量相比，全局内存的带宽较低，因此访问全局内存通常会对性能产生较大影响。

3.1.Data Transfer between Host and Device

应用程序应尽量减少host和device之间的数据传输。实现这一目标的一种方法是将更多代码从host转移到device。

此外，由于每次数据传输都会产生额外开销，因此可以通过将多次小规模数据传输合并为一次大规模传输来进行性能优化。

在具有FSB（前端总线，front-side bus）的系统中，通过使用页锁定host内存进行host与device之间的数据传输，可以实现更高的性能。

此外，当使用映射的页锁定内存时，无需分配任何device内存，也无需显式地在device和host之间拷贝数据。每次kernel访问映射内存时，数据传输会隐式进行。为了获得最佳性能，这些内存访问必须与对全局内存的访问合并（详见第3.2部分）。假设这些访问是合并的，并且映射内存只被读取或写入一次，那么相比在host和device之间进行显式的拷贝，使用映射的页锁定内存可以提升性能。

在device内存和host内存在物理上相同的集成系统中，host与device之间的任何拷贝操作都是多余的，因此应使用映射的页锁定内存。应用程序可以通过检查集成设备属性（见：Device Enumeration）是否等于1来查询设备是否integrated。

3.2.Device Memory Accesses

访问可寻址内存（即global memory、local memory、shared memory、constant memory、texture memory）的指令可能需要根据warp中线程的内存地址分布情况多次重新发出。地址分布对指令吞吐量的影响取决于每种内存的具体类型，接下来会详细描述。例如，对于全局内存，一般来说，地址分布越分散，吞吐量下降得越明显。

👉Global Memory

全局内存位于device内存中，device内存通过32字节、64字节或128字节的内存事务（memory transactions）进行访问。这些内存事务必须是自然对齐的：只有那些与事务大小对齐的32字节、64字节或128字节的device内存段（即，其起始地址是事务大小的整数倍）才能被内存事务读写。

当一个warp执行访问全局内存的指令时，它会根据每个线程访问的字（word）大小和线程内存地址的分布，将warp内线程的内存访问合并为一个或多个内存事务。一般来说，所需的内存事务越多，传输中未被线程实际访问的字越多，从而降低了指令的吞吐量。例如，如果每个线程访问4字节数据，但每次需要生成一个32字节的内存事务，那么吞吐量将被减少到原来的八分之一。

所需的内存事务数量以及最终吞吐量的影响取决于device的计算能力。

为了最大化全局内存的吞吐量，重要的是通过以下方法尽可能实现合并访问：

基于device的计算能力，选择最优的访问模式。
使用满足大小和对齐要求的数据类型（详见Size and Alignment Requirement部分）。
在某些情况下对数据进行padding，比如像Two-Dimensional Arrays部分描述的那样。

👉Size and Alignment Requirement

全局内存指令支持读取或写入大小为1字节、2字节、4字节、8字节或16字节的字（word）。只有当数据类型的大小是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为width，起始地址为BaseAddress，类型为type*（其中type符合第2部分中描述的要求）：

BaseAddress + width * ty + tx

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

这意味着如果数组的宽度不是warp大小的倍数，则它的访问效率会大大降低，此时，可以使用cudaMallocPitch()和cuMemAllocPitch()函数分配内存，这些函数会自动为每行添加适当的填充。

👉Local Memory

本地内存实际存储在全局内存中，访问速度慢于寄存器，但每个线程都有自己独立的本地内存。

本地内存访问仅发生在某些自动变量上。编译器可能会将以下类型的自动变量放置在本地内存中：

如果编译器无法确定数组的索引是常量（例如，索引值是动态计算的），它可能会将该数组存储在本地内存中。
CUDA的寄存器数量是有限的。如果某个变量（如大型结构或数组）需要大量寄存器空间，而寄存器不足，则会被放入本地内存。
寄存器溢出：如果kernel使用的寄存器数量超过了可用寄存器数量，则剩余的变量会被存储到本地内存中。

通过检查PTX汇编代码（使用-ptx或-keep选项编译生成），可以判断一个变量是否在第一个编译阶段（first compilation phases）被放置到本地内存中。如果变量被放置在本地内存中，它会通过.local指令声明，并通过ld.local和st.local指令进行访问。不过，即使在第一个编译阶段未被放置在本地内存中，后续的编译阶段仍可能因为其占用了太多的寄存器空间而将其放入本地内存，这取决于目标架构的限制：使用cuobjdump检查cubin对象来判断是否发生了这种情况。此外，当使用--ptxas-options=-v选项编译时，编译器会报告每个kernel的本地内存总使用量（lmem）。需要注意的是，一些数学函数可能有访问本地内存的实现路径。

本地内存空间位于device内存中，因此对本地内存的访问与对全局内存的访问一样，具有相同的高延迟和低带宽，并且需要满足内存合并要求。然而，本地内存的组织方式是连续的32位字被连续的线程ID访问。因此，只要warp中的所有线程访问相同的相对地址（例如，同一个数组变量的索引或同一个结构体变量的成员），访问就是完全合并的。例如，线程0访问地址A，线程1访问地址A+4，线程2访问地址A+8，这样的访问模式是完全合并的。

在计算能力5.x及以上的device中，本地内存访问总是以与全局内存访问相同的方式缓存到L2缓存中。

👉Shared Memory

因为共享内存是on-chip的，所以相比本地内存和全局内存，共享内存的带宽更高且延迟更低。

为了实现高带宽，共享内存被分成大小相等的内存模块，称为bank，它们可以被同时访问。因此，任何由n个地址组成的读或写请求，如果这些地址分别位于n个不同的bank中，那么这些请求可以被同时处理，从而获得单个模块带宽n倍的总带宽。

然而，如果一个内存请求中的两个地址落在同一个bank上，就会发生冲突（bank conflict），此时访问必须被串行化。硬件会将带有bank conflict的内存请求分解成多个独立且无冲突的请求，假设分解的请求数量为n，则称初始内存请求导致了n-way bank conflict，此时吞吐量也会下降到原来的1/n。

👉Constant Memory

常量内存空间位于device内存中，并被缓存在constant cache中。

然后，一个请求会根据初始请求中不同的内存地址被分解为多个独立请求，吞吐量就会按分解请求的数量成比例下降。

分解后的请求如果发生缓存命中（cache hit），将以constant cache的吞吐量被处理，否则，将以device内存的吞吐量被处理。

👉Texture and Surface Memory

纹理内存和表面内存位于device内存中，并被缓存在texture cache中。因此，只有在缓存未命中（cache miss）时，纹理读取（texture fetch）或表面读取（surface read）才需要从device内存进行内存读取，如果缓存命中（cache hit），则只需从texture cache读取即可。

相比通过全局内存或常量内存来读取device内存，通过纹理或表面来读取device内存有以下一些优点：

如果内存读取不能满足访问全局内存或常量内存的性能优化要求，则通过纹理读取或表面读取可以实现更高的带宽。
地址的计算在kernel外部由专门的单元完成。
打包的数据可以在单次操作中被广播到多个独立的变量。
8位或16位的整型输入数据可以选择性的转换为$[0.0,1.0]$或$[-1.0,1.0]$范围内的32位浮点值（详见：Texture Memory）。

4.Maximize Instruction Throughput

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

最小化使用吞吐量较低的算术指令。这包括在不影响最终结果的情况下，用精度换取速度，比如使用内置函数代替常规函数、使用单精度代替双精度、将非规范化数字（denormalized numbers）刷新为零等。
最小化由控制流指令（control flow instructions）导致的warp发散，详见第4.2部分。
减少指令数量，比如，尽可能地优化同步点（详见第4.3部分）或使用受限指针（restricted pointers）。

在本部分，吞吐量指的是操作吞吐量，即每个多处理器在每个时钟周期完成的操作数量。对于一个大小为32的warp，每条指令对应32个操作。因此，如果每个时钟周期可以执行$N$次操作，那么指令吞吐量为$N/32$。

所有的吞吐量评估都是基于单个多处理器的。要计算整个device的总吞吐量，需要将单个多处理器的吞吐量乘以device中多处理器的数量。

4.1.Arithmetic Instructions

下表列出了在硬件中原生支持（natively supported）的算术指令的吞吐量。

表中上标的解释：

\

\

128 for __nv_bfloat16

8 for GeForce GPUs, except for Titan GPUs

2 for compute capability 7.5 GPUs

32 for extended-precision

32 for GeForce GPUs, except for Titan GPUs

16 for compute capabilities 7.5 GPUs

8 for GeForce GPUs, except for Titan GPUs

2 for compute capabilities 7.5 GPUs

其他指令和函数都是基于这些原生指令（native instructions）实现的。对于不同计算能力的device，其实现可能会有所不同，并且在编译后生成的原生指令数量可能会随每个编译器版本而有所变化。对于复杂的函数，根据输入可能会有多个代码路径。可以使用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编译的代码性能更高。

👉Single-Precision Floating-Point Division

__fdividef(x, y)提供了比除法运算符更快的单精度浮点数除法。

👉Single-Precision Floating-Point Reciprocal Square Root

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

rsqrtf()（倒数平方根，reciprocal square root）的计算方式为$\text{rsqrt}(x) = \frac{1}{\sqrt{x}}$。

👉Single-Precision Floating-Point Square Root

基于倒数平方根计算平方根有两种方式：

先求倒数平方根，然后再求倒数：$\sqrt{x} = \frac{1}{\text{rsqrt}(x)}$
先求倒数平方根，然后再做乘法：$\sqrt{x} = x \times \text{rsqrt}(x)$

单精度浮点数的平方根计算选择了方式一，因为方式一确保了在处理0和无穷时能够得到正确的结果。

👉Sine and Cosine

sinf(x)、cosf(x)、tanf(x)、sincosf(x)以及对应的双精度指令的计算代价非常高，尤其是当参数x的数值很大时，计算代价会进一步增加。

更具体地说，参数约简代码（argument reduction code）包含两种代码路径，分别是快速路径（fast path）和慢速路径（slow path）。

个人注解：三角函数的计算是一个十分复杂的问题。因为三角函数呈周期性，因此我们可以通过参数约简将输入化简到函数的基本周期范围内，比如$[-\pi,\pi]$或$[0,2\pi]$等。而约简的方式以及约简后三角函数的计算有多种实现方式。

快速路径用于处理值足够小的参数，其实现本质上仅需进行几次乘加运算。慢速路径用于处理值较大的参数，其实现需要较长的计算步骤，以确保在整个参数范围内得到正确的结果。

目前，在三角函数的参数约简代码中，对于单精度函数，当参数值小于105615.0f时会选择快速路径，对于双精度函数，当参数值小于2147483648.0时会选择快速路径。

由于慢速路径需要的寄存器比快速路径多，因此尝试将一些中间变量存储在local memory中，以减少慢速路径的寄存器压力。然而，由于local memory具有高延迟和带宽限制（见第3.2部分），这可能会影响性能。目前，单精度函数使用28字节的local memory，双精度函数使用44字节的local memory，但是具体使用量可能会发生变化。

由于慢速路径中的计算步骤较多且使用了local memory，当需要进行慢速路径的约简时，与快速路径约简相比，这些三角函数的吞吐量降低了一个数量级。

👉Integer Arithmetic

整数除法和取模操作在硬件上实现复杂，代价较高，尤其是当除数或模数不是2的幂时，它们在编译时可能会被转换为多达20条指令。在某些情况下，可以用位操作代替这些运算：如果n是2的幂，那么(i/n)等同于(i>>log2(n))，(i%n)等同于(i&(n-1))；如果n是常量，编译器会自动进行这些转换。

__brev和__popc映射到单个指令，而__brevll和__popcll则映射为多个指令。

个人注解：

__brev和__brevll用于按位反转（bit-reverse），即将一个整数的二进制位顺序颠倒，其中，__brev用于32位无符号整数，__brevll用于64位无符号整数。

__popc和__popcll用于统计二进制数中1的数量（popcount），其中，__popc用于32位无符号整数，__popcll用于64位无符号整数。

__[u]mul24是过时的内置函数，现在已经没有使用的必要了。

👉Half Precision Arithmetic

为了在16位浮点数（即半精度）的加法、乘法或乘加运算中获得好的性能，建议使用half2数据类型处理half精度的操作，使用__nv_bfloat162处理__nv_bfloat16精度的操作。向量化内置函数（例如__hadd2、__hsub2、__hmul2、__hfma2）可用于在一条指令中执行两个操作。使用half2或__nv_bfloat162代替half或__nv_bfloat16还可能会提高其他内置函数的性能，比如warp shuffles。

内置函数__halves2half2可将两个half精度的值转换成half2数据类型。

内置函数__halves2bfloat162可将两个__nv_bfloat精度的值转换成__nv_bfloat162数据类型。

👉Type Conversion

有时，编译器必须插入类型转换指令，从而引入额外的执行周期。这种情况通常出现在以下场景中：

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

对于第二种情况，可以通过定义带有f后缀的单精度浮点常量来避免这种转换，例如3.141592653589793f、1.0f、0.5f等。

4.2.Control Flow Instructions

任何流控制指令（比如if、switch、do、for、while）都可能显著影响指令的有效吞吐量，因为它们可能导致同一个warp中的线程发生分支分裂。如果发生分支分裂，不同的执行路径将被串行化，从而增加该warp执行的总指令数。

为了在控制流（control flow）依赖线程ID的情况下获得最佳性能，应该以最小化warp分支分裂数量的方式编写控制条件。这是可行的，因为warp在block内的分布是确定的，详见SIMT Architecture。一个简单的例子是，当控制条件仅依赖于threadIdx / warpSize（warpSize为warp的大小）时，不会发生分支分裂，因为控制条件与warp的划分完全对齐。

比如以下代码：

//不会发生分支分裂，性能大幅提高
if (threadIdx.x / warpSize == 0) {
    //路径A
    //线程ID 0~31属于warp0，全部执行路径A
} else {
    //路径B
    //线程ID 32~63属于warp1，全部执行路径B
}

有时，编译器可能会对循环进行展开（unroll loops），或者通过分支预测来优化短的if或switch代码块。在这些情况下，warp永远不会发生分支分裂。程序员还可以使用#pragma unroll来控制循环的展开。

循环展开是编译器的一种优化技术，它将循环体的多次迭代直接展开为连续的指令，从而减少循环控制开销：

//适用于较小的循环
//未展开的循环
for (int i = 0; i < 4; i++) {
    arr[i] = arr[i] * 2;
}

//展开后的代码
arr[0] = arr[0] * 2;
arr[1] = arr[1] * 2;
arr[2] = arr[2] * 2;
arr[3] = arr[3] * 2;

分支预测是GPU的一种优化技术，它通过在每个线程中执行分支的所有路径并屏蔽不需要的结果，来避免warp分支分裂的影响。这种方式适用于短小的if或switch代码块。这和CPU中的分支预测机制不同，CPU使用复杂的硬件逻辑来预测分支条件的结果，并提前加载预测的分支路径，如果预测正确，则直接继续执行；如果预测错误，则回滚并重新执行正确的路径。而GPU中的分支预测会同时计算所有可能的分支路径，然后根据条件屏蔽掉不需要的路径。

4.3.Synchronization Instruction

对于__syncthreads()，其吞吐量在不同计算能力的device上如下：

对于计算能力6.0的device，每个时钟周期支持32次操作。
对于计算能力7.x和8.x的device，每个时钟周期支持16次操作。
对于计算能力5.x、6.1和6.2的device，每个时钟周期支持64次操作。

注意，__syncthreads()可能会强制让多处理器空闲，从而影响性能，详见第3.2部分。

5.Minimize Memory Thrashing

如果频繁分配和释放内存，系统需要反复分配和回收内存块，这会带来显著的开销。并且如果内存分配和释放模式不连续，内存空间可能会变得零散，随着程序运行时间的增长，找到足够大的连续内存块的难度增加，导致分配时间变长。为了在这方面获得最佳性能，我们推荐以下措施：

尽量根据实际问题的需求调整分配的内存大小。不要尝试通过cudaMalloc、cudaMallocHost、cuMemCreate分配所有可用内存，因为这会强制内存立即驻留并阻止其他应用程序使用这些内存。这种行为可能会增加操作系统调度程序的压力，或者完全阻止其他使用相同GPU的应用程序运行。
尽量在应用程序的早期就分配适当大小的内存。减少在性能关键区域中cudaMalloc+cudaFree的调用次数。
如果一个应用程序无法分配足够的device内存，可以考虑使用其他类型的内存，例如cudaMallocHost或cudaMallocManaged，尽管它们的性能可能不如device内存，但可以帮助应用程序继续执行下去。
对于支持该功能的平台，cudaMallocManaged允许内存超额使用，如果启用了正确的cudaMemAdvise策略，应用程序还可以保留cudaMalloc的大部分性能。使用cudaMallocManaged，数据可以部分存储在host内存中，并在需要时自动迁移到GPU。同时，cudaMallocManaged分配的内存不会立即驻留，只有当数据被使用或明确预取（prefetch）时才会迁移到device。这降低了对操作系统调度程序的总体压力，并更好地支持multi-tenet场景（指多个用户或应用程序共享同一个GPU资源的情况）。

【CUDA编程】【25】【5.Performance Guidelines】

Overall Performance Optimization Strategies，Maximize Utilization，Maximize Memory Throughput，Maximize Instruction Throughput，Minimize Memory Thrashing