4. Hardware Implementation

NVIDIA GPU架构是围绕一系列可扩展的多线程流式多处理器（SM）构建的。当主机CPU上的CUDA程序调用内核网格时，网格的块将被枚举并分配给具有可用执行能力的多处理器。线程块的线程在一个多处理器上并发执行，并且多个线程块可以在一个多处理器上并发执行。当线程块终止时，在空出的多处理器上启动新的块。

多处理器被设计为同时执行数百个线程。为了管理如此大量的线程，它采用了一种称为SIMT（单指令多线程）的独特体系结构，该体系结构在SIMT体系结构中进行了描述。这些指令是流水线式的，在单个线程中利用指令级并行性，以及通过硬件多线程中详细介绍的同步硬件多线程实现的广泛线程级并行性。与CPU内核不同，它们是按顺序发出的，没有分支预测或推测执行。

SIMT体系结构和硬件多线程描述了所有设备通用的流多处理器的体系结构特征。计算能力5.x、计算能力6.x和计算能力7.x分别提供计算能力5.x、6.x和7.x的设备的细节。

4.1. SIMT Architecture

多处理器以32个并行线程为一组（称为warps）创建、管理、调度和执行线程。组成warp的各个线程一起从相同的程序地址开始，但它们有自己的指令地址计数器和寄存器状态，因此可以自由地分支和独立执行。经纱一词起源于织造，这是第一种平行线技术。半扭曲是扭曲的第一半或第二半。四分之一经线是经线的第一、第二、第三或第四个四分之一。
 

当一个多处理器有一个或多个线程块要执行时，它会将这些线程块划分为多个warp，每个warp都由warp调度器调度执行。将块划分为wrap的方式总是相同的;每一经线包含连续的、递增的线ID的线，其中第一经线包含线0。线程层次描述了线程ID与块中线程索引的关系。

一个warp一次执行一条公共指令，因此当warp的所有32个线程都同意它们的执行路径时，实现了完全的效率。如果warp的线程通过依赖于数据的条件分支而发散，则warp执行所采用的每个分支路径，禁用不在该路径上的线程。分支发散仅发生在warp内。不同的warp独立地执行，而不管它们是执行公共的还是不相交的代码路径。

SIMT体系结构类似于SIMD（单指令多数据）向量组织，因为单指令控制多个处理元件。一个关键区别是SIMD向量组织向软件公开SIMD宽度，而SIMD指令指定单个线程的执行和分支行为。与SIMD向量机相比，SIMT使程序员能够为独立的标量线程编写线程级并行代码，并为协调线程编写数据并行代码。出于正确性的目的，程序员基本上可以忽略SIMT行为;然而，通过注意代码很少要求扭曲中的线程发散，可以实现实质性的性能改进。实际上，这类似于传统代码中缓存行的作用：在设计正确性时，可以安全地忽略缓存线大小，但在设计峰值性能时，必须在代码结构中加以考虑。另一方面，矢量架构要求软件将负载合并到矢量中，并手动管理差异。

在NVIDIA Volta之前，warps使用一个程序计数器，该计数器在warp中的所有32个线程之间共享，并使用一个活动掩码指定warp的活动线程。因此，来自不同区域或不同执行状态的相同warp的线程无法相互发送信号或交换数据，并且需要细粒度共享由锁或互斥体保护的数据的算法很容易导致死锁，这取决于争用线程来自哪个warp。

从NVIDIA Volta架构开始，独立线程调度允许线程之间的完全并发，而不考虑扭曲。通过独立线程调度，GPU可以维护每个线程的执行状态，包括程序计数器和调用堆栈，并可以在每个线程的粒度上执行，以便更好地利用执行资源或允许一个线程等待另一个线程生成数据。调度优化器确定如何将来自同一个warp的活动线程分组到SIMT单元中。这保留了以前NVIDIA GPU中SIMT执行的高吞吐量，但具有更大的灵活性：线程现在可以以亚翘曲粒度发散和再会聚。

如果开发人员对以前的硬件架构的warp-synchronicity 2进行假设，独立线程调度可能会导致参与执行代码的线程集与预期的线程集完全不同。特别是，任何翘曲同步代码（如无同步，翘曲内减少）应重新审查，以确保与NVIDIA Volta和超越兼容。有关详细信息，请参见计算能力7.x。
注意

4.2. Hardware Multithreading

多处理器处理的每个warp的执行上下文（程序计数器、寄存器等）在warp的整个生命周期内都在片上维护。因此，从一个执行上下文切换到另一个执行上下文没有成本，并且在每个指令发布时间，曲速调度器选择具有准备好执行其下一指令的线程的曲速（曲速的活动线程），并且将指令发布到那些线程。

特别地，每个多处理器具有在扭曲之间划分的一组32位寄存器，以及在线程块之间划分的并行数据高速缓存或共享存储器。

对于给定内核，可以驻留在多处理器上并一起处理的块和Wrap的数量取决于内核使用的寄存器和共享内存的数量以及多处理器上可用的寄存器和共享内存的数量。每个多处理器还有最大驻留块数和最大驻留扭曲数。这些限制以及多处理器上可用的寄存器和共享内存的数量是设备计算能力的函数，在“计算能力”中给出。如果每个多处理器没有足够的寄存器或共享内存来处理至少一个块，内核将无法启动。
块中的扭曲总数如下所示：

• T is the number of threads per block,

• Wsize is the warp size, which is equal to 32,

• ceil(x, y) is equal to x rounded up to the nearest multiple of y.

分配给块的寄存器总数和共享内存总量记录在CUDA工具包提供的CUDA占用计算器中。

5. Performance Guidelines

5.1. Overall Performance Optimization Strategies

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

• 最大化并行执行以实现最大利用率;

• 优化内存使用以实现最大内存吞吐量;

• 优化指令使用以实现最大指令吞吐量;

• 最小化内存抖动。

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

5.2. Maximize Utilization

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

在高级别上，应用程序应通过使用异步函数调用和流（如中所述），最大限度地提高主机、设备以及连接主机和设备的总线之间的并行执行 异步并发执行.它应该为每个处理器分配它最擅长的工作类型：主机的串行工作负载;并行工作负载分配给设备。

对于并行工作负载，在算法中由于一些线程需要同步以便彼此共享数据而中断并行性的点处，存在两种情况：这些线程要么属于同一个块，在这种情况下，它们应该使用__syncthreads（）并在同一个内核调用中通过共享内存共享数据;要么属于不同的块，在这种情况下，它们必须使用两个单独的内核调用（一个用于写入全局内存，一个用于从全局内存阅读）通过全局内存共享数据。第二种情况的最优性要差得多，因为它增加了额外的内核调用和全局内存流量的开销。因此，应通过将算法映射到CUDA编程模型来最大限度地减少这种情况的发生，映射方式应确保需要线程间通信的计算尽可能在单个线程块中执行。

5.2.2. Device Level

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

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

5.2.3. Multiprocessor Level

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

如中所述 硬件多线程，GPU多处理器主要依赖于线程级并行性来最大化其功能单元的利用率。因此，利用率与驻留经纱的数量直接相关。在每个指令发布时间，翘曲调度器选择准备好执行的指令。该指令可以是同一个warp的另一个独立指令，利用指令级并行性;或者更常见的是，该指令可以是另一个warp的指令，利用线程级并行性。如果选择了准备执行的指令，则将其发送到 活动的经线的线。一个翘曲准备好执行其下一条指令所需的时钟周期数称为等待时间，当所有翘曲调度器在该等待时间段内的每个时钟周期总是有一些指令要为一些翘曲发出时，或者换句话说，当等待时间完全“隐藏”时，就实现了充分利用.隐藏L个时钟周期的延迟所需的指令数取决于这些指令各自的吞吐量（请参见 算术指令用于各种算术指令的吞吐量）。如果我们假设指令具有最大吞吐量，则它等于

• 4L for devices of compute capability 5.x, 6.1, 6.2, 7.x and 8.x since for these devices, a multiprocessor issues one instruction per warp over one clock cycle for four warps at a time, as mentioned in Compute Capabilities.
  对于计算能力5.x、6.1、6.2、7.x和8.x的设备，因为对于这些设备，多处理器在一个时钟周期上针对一次四个线程束每个线程束发出一个指令，如在图4A和4B中所提到的。 计算能力.

• 2L for devices of compute capability 6.0 since for these devices, the two instructions issued every cycle are one instruction for two different warps.
  对于计算能力为6.0的设备，这是因为对于这些设备，每个周期发出的两个指令是用于两个不同扭曲的一个指令。

一个warp没有准备好执行它的下一条指令的最常见的原因是指令的输入操作数还不可用

如果所有输入操作数都是寄存器，则延迟由寄存器相关性引起，即，一些输入操作数由一些尚未完成执行的先前指令写入。在这种情况下，等待时间等于前一指令的执行时间，并且Wrap调度器必须在该时间期间调度其它Wrap的指令。执行时间因指令而异。在计算能力为7.x的器件上，对于大多数算术指令，它通常为4个时钟周期。这意味着每个多处理器需要16个活动的扭曲（4个周期，4个Wrap调度器）来隐藏算术指令延迟（假设扭曲以最大吞吐量执行指令，否则需要较少的Wrap）。如果各个warp表现出指令级并行性，即在它们的指令流中具有多个独立指令，则需要较少的warp，因为来自单个warp的多个独立指令可以背对背地发出。

如果某个输入操作数驻留在片外存储器中，则延迟要高得多：通常为数百个时钟周期。在如此高的延迟时间段内，使warp调度器保持忙碌所需的warp数量取决于内核代码及其指令级并行度。一般来说，如果不具有芯片外存储器操作数的指令的数目的比率（即，大多数时间为算术指令）与具有芯片外存储器操作数的指令的数目的比率较低（此比率通常称为程序的算术强度）

warp未准备好执行下一条指令的另一个原因是它正在某个内存栅栏（内存栅栏函数）或同步点（同步函数）处等待。同步点可以迫使多处理器空闲，因为越来越多的warp等待同一块中的其他warp在同步点之前完成指令的执行。在这种情况下，每个多处理器拥有多个驻留块有助于减少空闲，因为来自不同块的扭曲不需要在同步点彼此等待。

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

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

内核使用的寄存器数量对驻留warp的数量有很大的影响。例如，对于计算能力为6.x的设备，如果内核使用64个寄存器并且每个块具有512个线程并且需要非常少的共享存储器，则两个块（即，32次扭曲）可以驻留在多处理器上，因为它们需要2 × 512 × 64个寄存器，这与多处理器上可用的寄存器的数量完全匹配。但是一旦内核再使用一个寄存器，就只有一个块（即，16次扭曲）可以驻留，因为两个块将需要2 × 512 × 65个寄存器，这比多处理器上可用的寄存器更多。因此，编译器尝试在保持寄存器溢出的同时最小化寄存器使用（请参见 设备内存访问）和指令的数量减少到最小。可以使用maxrregcount编译器选项或启动边界控制寄存器的使用，如中所述 发射界限.

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

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

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

5.2.3.1. Occupancy Calculator

有几个API函数可以帮助程序员根据寄存器和共享内存要求选择 choosing thread block size and cluster size

• 占用计算器API cudaOccupancyMaxActiveBlocksPerMultiprocessor可以基于内核的块大小和共享存储器使用来提供占用预测。此函数以每个多处理器的并发线程块数量报告占用情况。

  请注意，此值可以转换为其他指标。乘以每个块的扭曲数得到每个多处理器的并发扭曲数;进一步将并发warp除以每个多处理器的最大warp，给出了百分比形式的占用率。
• 基于占用率的启动配置器API cudaOccupancyMaxPotentialBlockSize和cudaOccupancyMaxPotentialBlockSizeVariableSMem启发式地计算实现最大多处理器级别占用率的执行配置。
• 占用率计算器API cudaOccupancyMaxActiveClusters可以基于内核的簇大小、块大小和共享存储器使用来提供占用率预测。此函数报告系统中GPU上给定大小的最大活动群集数的占用率。

下面的代码示例计算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);
    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个线程的集群的占用率

群集大小为8是向前兼容的起始计算能力9.0，但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计算用户界面还在<CUDA_Toolkit_Path>/include/cuda_occupancy.h中提供了独立的占用计算器和启动配置器实施，用于无法依赖CUDA软件堆栈的任何用例。Nsight Compute版本的占用率计算器作为一种学习工具特别有用，它可以直观地显示影响占用率的参数（块大小、每个线程的寄存器和每个线程的共享内存）的更改所产生的影响。

5.3. Maximize Memory Throughput

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

这意味着尽量减少主机和设备之间的数据传输，如主机和设备之间的数据传输中所述，因为这些数据传输的带宽比全局内存和设备之间的数据传输低得多。

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

共享内存相当于用户管理的缓存： 应用程序显式分配和访问它。CUDA运行时典型的编程模式是将来自设备存储器的数据分级到共享存储器中;换句话说，使块的每个线程：

• 将数据从设备存储器加载到共享存储器，

• 与块的所有其他线程同步，以便每个线程可以安全地读取由不同线程填充的共享内存位置，

• Process the data in shared memory,处理共享内存中的数据，

• 有必要，请再次同步，以确保共享内存已用结果更新，

• Write the results back to device memory.将结果写回设备存储器。

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

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

5.3.1. Data Transfer between Host and Device

应用程序应尽量减少主机和设备之间的数据传输。实现这一点的一种方法是将更多的代码从主机移动到设备，即使这意味着运行的内核没有暴露足够的并行性来在设备上以最高效率执行。中间数据结构可以在设备存储器中创建、由设备操作、以及在从未被主机映射或复制到主机存储器的情况下被销毁。

此外，由于与每次传输相关联的开销，将许多小的传输批处理成单个大的传输总是比单独进行每次传输执行得更好。

在具有前端总线的系统上，通过使用页锁定主机内存（如中所述），可以提高主机与设备之间的数据传输性能 页锁定主机内存.

此外，当使用映射页锁定内存（映射内存）时，不需要分配任何设备内存，也不需要在设备和主机内存之间显式复制数据。每次内核访问映射内存时，都会隐式执行数据传输。为了获得最佳性能，这些内存访问必须像访问全局内存一样合并（请参阅设备内存访问）。假设它们是一致的，并且映射内存只读写一次，那么在设备和主机内存之间使用映射页锁定内存而不是显式副本可能会提高性能。

在设备内存和主机内存在物理上相同的集成系统上，主机内存和设备内存之间的任何复制都是多余的，应该使用映射页锁定内存。应用程序可以通过检查集成设备属性（参见设备枚举）是否等于1来查询设备是否为integrated。

5.3.2. Device Memory Accesses

访问可寻址存储器的指令（即，全局、局部、共享、常量或纹理存储器）可能需要根据存储器地址在warp内的线程上的分布而被多次重新发布。这种分配方式如何影响指令吞吐量取决于每种类型的内存，以下各节对此进行了说明。例如，对于全局存储器，作为一般规则，地址越分散，吞吐量减少得越多。

Global Memory全局内存

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

当 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的最佳访问模式

• 使用满足以下大小和对齐要求部分中详细描述的大小和对齐要求的数据类型，

• 在某些情况下填充数据，例如，在访问下面的二维数组一节中所述的二维数组时。

Size and Alignment Requirement

全局存储器指令支持阅读或写大小等于1、2、4、8或16字节的字。当且仅当数据类型的大小为1、2、4、8或16字节且数据自然对齐（即，其地址是该大小倍数）

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

对齐要求将自动满足 内置向量类型.

For structures, the size and alignment requirements can be enforced by the compiler using the alignment specifiers__align__(8) or __align__(16), such as                   

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

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

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

阅读非自然对齐的8字节或16字节字会产生不正确的结果（偏离几个字），因此必须特别注意保持这些类型的任何值或值数组的起始地址对齐。这可能容易被忽视的典型情况是当使用一些自定义全局存储器分配方案时，由此多个阵列的分配（具有对cudaMalloc()或cuMemAlloc()的多次调用）被分配被划分成多个阵列的单个大存储器块所取代，在这种情况下，每个阵列的起始地址从块的起始地址偏移。

Two-Dimensional Arrays

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

BaseAddress + width * ty + tx

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

特别是，这意味着，如果实际分配的宽度向上舍入到该大小的最接近倍数，并相应地填充其行，则访问宽度不是该大小倍数的数组的效率会高得多。参考手册中描述的cudaMallocPitch()和cuMemAllocPitch()函数以及相关的存储器复制函数使程序员能够编写非硬件相关的代码来分配符合这些约束的数组。

Local Memory本地存储器

本地内存访问仅适用于提到的某些自动变量 变量内存空间说明符.编译器可能放置在本地内存中的自动变量有：

• Arrays for which it cannot determine that they are indexed with constant quantities,
  无法确定其是否用常量索引的数组，

• Large structures or arrays that would consume too much register space,
  会消耗太多寄存器空间的大型结构或数组，

• Any variable if the kernel uses more registers than available (this is also known as register spilling).
  内核使用的寄存器多于可用寄存器时的任何变量（这也称为寄存器溢出）。

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

本地内存空间驻留在设备内存中，因此本地内存访问与全局内存访问具有相同的高延迟和低带宽，并且遵循相同的内存合并要求，如中所述 设备内存访问.然而，本地存储器被组织成使得连续的32位字被连续的线程ID访问。因此，只要一个warp中的所有线程访问相同的相对地址（例如，数组变量中的相同索引，结构变量中的相同成员），访问就被完全合并。

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

Shared Memory共享内存

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

为了实现高带宽，共享内存被划分为大小相等的内存模块（称为存储体），可以同时访问这些模块。因此，由落入n个不同存储体中的n个地址构成的任何存储器读或写请求都可以同时得到服务，从而产生n倍于单个模块带宽的总带宽。

然而，如果存储器请求的两个地址落在同一存储体中，则存在存储体冲突，并且访问必须串行化。硬件将具有存储体冲突的存储器请求拆分为所需数量的独立无冲突请求，从而以等于独立存储器请求数量的因子降低吞吐量。如果单独的存储器请求的数量是η，则初始存储器请求被称为引起η路存储体冲突。

为了获得最大性能，因此理解存储器地址如何映射到存储器组以便调度存储器请求从而最小化组冲突是重要的。计算 能力5.x 、 计算能力6.x 、 计算能力7.x、计算能力8.x和计算能力9.0分别针对计算能力5.x、6.x、7.x、8.x和9.0的设备进行了描述。

Constant Memory 常量内存

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

然后，请求被分成与初始请求中存在的不同存储器地址一样多的单独请求，从而以等于单独请求的数目的因子来降低吞吐量。

然后，在高速缓存命中的情况下，以常量高速缓存的吞吐量来服务所产生的请求，否则以设备存储器的吞吐量来服务所产生的请求。

Texture and Surface Memory纹理和表面内存

纹理和表面内存空间驻留在设备内存中，并缓存在纹理缓存中，因此纹理提取或表面读取仅在缓存未命中时才需要从设备内存读取一次内存，否则只需从纹理缓存读取一次内存。纹理缓存针对2D空间局部性进行了优化，因此读取2D中靠近的纹理或表面地址的相同扭曲的线程将实现最佳性能。而且，它被设计用于具有恒定延迟的流提取;高速缓存命中减少DRAM带宽需求，但不减少提取等待时间。

通过纹理或表面提取阅读设备内存具有一些优点，使其成为从全局或常量内存读取设备内存的有利替代方法：

• If the memory reads do not follow the access patterns that global or constant memory reads must follow to get good performance, higher bandwidth can be achieved providing that there is locality in the texture fetches or surface reads;
  如果存储器读取不遵循全局或恒定存储器读取必须遵循以获得良好性能的存取模式，那么可实现较高带宽，前提是在纹理提取或表面读取中存在局部性;

• Addressing calculations are performed outside the kernel by dedicated units;
  寻址计算由专用单元在内核外部执行;

• Packed data may be broadcast to separate variables in a single operation;
  打包的数据可以在单个操作中被广播到单独的变量;

• 8-bit and 16-bit integer input data may be optionally converted to 32 bit floating-point values in the range [0.0, 1.0] or [-1.0, 1.0] (see Texture Memory).
  8-位和16位整数输入数据可以选择转换为范围[0.0，1.0]或[-1.0，1.0]中的32位浮点值（请参见纹理内存）。

5.4. Maximize Instruction Throughput最大化指令吞吐量

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

• 最小化低吞吐量算术指令的使用; 这包括在不影响最终结果时以精度换取速度，例如使用内在函数而不是常规函数（内在函数在中列出内在函数）、单精度而不是双精度，或者将非规格化的数字刷新为零;

• 最大限度地减少控制流指令导致的发散扭曲，如中所述 控制流程说明

• 减少指令数，例如，尽可能优化同步点，如中所述 同步指令 或者通过使用如 __限制__.

在本节中，吞吐量以每个多处理器每个时钟周期的操作数表示。对于32的翘曲大小，一个指令对应于32个操作，因此如果N是每个时钟周期的操作数，则指令吞吐量是每个时钟周期的N/32个指令。

所有吞吐量都是针对一个多处理器的。它们必须乘以设备中多处理器的数量才能得到整个设备的吞吐量。

5.4.1. Arithmetic Instructions算术指令

Compute Capability	5.0, 5.2	5.3	6.0	6.1	6.2	7.x	8.0	8.6	8.9	9.0
16-bit floating-point add, multiply, multiply-add	N/A	256	128	2	256	128	2563		128	256
32-bit floating-point add, multiply, multiply-add	128		64	128		64		128
64-bit floating-point add, multiply, multiply-add	4		32	4		325	32	2		64
32-bit floating-point reciprocal, reciprocal square root, base-2 logarithm (__log2f), base 2 exponential (exp2f), sine (__sinf), cosine (__cosf)	32		16	32		16
32-bit integer add, extended-precision add, subtract, extended-precision subtract	128		64	128		64
32-bit integer multiply, multiply-add, extended-precision multiply-add	Multiple instruct.					646
24-bit integer multiply (__[u]mul24)	Multiple instruct.
32-bit integer shift	64		32	64
compare, minimum, maximum	64		32	64
32-bit integer bit reverse	64		32	64		16
Bit field extract/insert	64		32	64		Multiple Instruct.			64
32-bit bitwise AND, OR, XOR	128		64	128		64
count of leading zeros, most significant non-sign bit	32 | 16 | 32 | 16
population count	32		16	32		16
warp shuffle	32					328	32
warp reduce	Multiple instruct.						16
warp vote	64
sum of absolute difference	64		32	64
SIMD video instructions vabsdiff2	Multiple instruct.
SIMD video instructions vabsdiff4	Multiple instruct.					64
All other SIMD video instructions	Multiple instruct.
Type conversions from 8-bit and 16-bit integer to 32-bit integer types	32		16	32		64
Type conversions from and to 64-bit types	4		16	4		1610	16	2	2	16
All other type conversions	32		16	32		16
16-bit DPX	Multiple instruct.									128
32-bit DPX	Multiple instruct.									64

其它指令和功能在本地指令之上实现。对于具有不同计算能力的设备，实现可能不同，并且编译后的本地指令的数量可能随每个编译器版本而波动。对于复杂的函数，根据输入可能有多个代码路径。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&(n-1)）;编译器将执行这些转换，如果n是文字。

__brev and __popc map to a single instruction and __brevll and __popcll to a few instructions.

__[u]mul24 are legacy intrinsic functions that no longer have any reason to be used.
__[u]mul24是传统的内在函数，不再有任何理由被使用。

Half Precision Arithmetic半精度算法

In order to achieve good performance for 16-bit precision floating-point add, multiply or multiply-add, it is recommended that the half2 datatype is used for half precision and __nv_bfloat162 be used for __nv_bfloat16 precision. Vector intrinsics (for example, __hadd2, __hsub2, __hmul2, __hfma2) can then be used to do two operations in a single instruction. Using half2 or __nv_bfloat162 in place of two calls using half or __nv_bfloat16 may also help performance of other intrinsics, such as warp shuffles.

The intrinsic __halves2half2 is provided to convert two half precision values to the half2 datatype.
提供固有的__halves2half2将两个half精度值转换为half2数据类型。

The intrinsic __halves2bfloat162 is provided to convert two __nv_bfloat precision values to the __nv_bfloat162 datatype.
提供固有的__halves2half2将两个half精度值转换为half2数据类型。

Type Conversion类型转换

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

• 操作类型为char或short的变量的函数，其操作数通常需要转换为int，

• 双精度浮点常量（即，定义为没有任何类型后缀的那些常量），用作单精度浮点计算的输入（如C/C++标准所要求的）。

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

5.4.2. Control Flow Instructions控制流程说明

任何流控制指令（if、switch、do、for、while）可通过致使相同线程束的线程发散（即，发散）而显著影响有效指令吞吐量。遵循不同的执行路径）。如果发生这种情况，不同的执行路径必须被串行化，从而增加了为该翘曲执行的指令总数。

为了在控制流依赖于线程ID的情况下获得最佳性能，应该编写控制条件，以便最小化发散扭曲的数量。这是可能的，因为如SIMT体系结构中所述，扭曲在块上的分布是确定性的。一个简单的例子是当控制条件仅取决于（threadIdx / warpSize）时，其中warpSize是扭曲大小。在这种情况下，由于控制条件与扭曲完全对齐，因此没有扭曲发散。

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

使用分支谓词时，不会跳过执行依赖于控制条件的任何指令。相反，它们中的每一个都与基于控制条件被设置为真或假的每线程条件码或谓词相关联，并且尽管这些指令中的每一个都被调度用于执行，但是实际上仅执行具有真谓词的指令。谓词为假的指令不写入结果，也不计算地址或读取操作数。

5.4.3. Synchronization Instruction同步指令

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

请注意，__syncthreads()会强制多处理器空闲，从而影响性能，详见设备内存访问。

5.5. Minimize Memory Thrashing最小化内存抖动

经常分配和释放内存的应用程序可能会发现，分配调用往往会随着时间的推移而变慢，直到达到某个限制。由于将内存释放回操作系统供其自己使用的特性，这通常是可以预料到的。为了在这方面获得最佳性能，我们建议采取以下措施：

• 试着根据手头的问题确定你的分配额。不要尝试使用cudaMalloc / cudaMallocHost / cuMemCreate分配所有可用内存，因为这会强制内存立即驻留，并阻止其他应用程序使用该内存。这会给操作系统调度程序带来更大的压力，或者只是阻止使用同一GPU的其他应用程序完全运行。

• 尽量在应用程序的早期以适当大小的分配方式分配内存，并仅在应用程序不使用内存时进行分配。减少应用程序中cudaMalloc+cudaFree调用的数量，特别是在性能关键的区域。

• 如果应用程序无法分配足够的设备内存，请考虑回退到其他内存类型（如cudaMallocHost或cudaMallocManaged），这些内存类型可能性能不佳，但可以使应用程序继续运行。

• 对于支持该功能的平台，cudaMallocManaged允许超额订阅，并且启用了正确的cudaMemAdvise策略，将允许应用程序保留cudaMalloc的大部分（如果不是全部）性能。cudaMallocManaged也不会强制一个分配在被需要或预取之前一直驻留，从而降低了操作系统调度器的整体压力，并更好地支持多原则用例。

6. CUDA-Enabled GPUs

The compute capability, number of multiprocessors, clock frequency, total amount of device memory, and other properties can be queried using the runtime (see reference manual).
计算能力、多处理器数量、时钟频率、设备内存总量和其他属性都可以使用运行时查询（参见参考手册）。