Overall Performance Optimization Strategies 整体性能优化策略
性能优化围绕四个基本策略展开：
- 最大化并行执行以实现最大利用率；
- 优化内存使用以实现最大的内存吞吐量；
- 优化指令使用以实现最大的指令吞吐量；
- 最小化内存抖动。
对于应用程序的特定部分，哪些策略会产生最佳的性能提升取决于该部分的性能限制因素；例如，优化主要受内存访问限制的内核的指令使用不会产生任何显著的性能提升。因此，优化工作应该不断地通过测量和监控性能限制因素来指导（别瞎优化！），例如使用CUDA profiler。此外，将特定内核的浮点操作吞吐量或内存吞吐量（取决于哪一个更有意义）与设备的相应峰值理论吞吐量进行比较，可以知道该内核有多少改进空间。
Maximize Utilization
为了最大化利用SM，算法应该以尽可能多的并行方式设计，并有效地将这种并行性映射到系统的各个组件，以使它们大部分时间保持忙碌。
应用层级 Application Level
从高层次来看，应用程序应该使用异步函数调用和流，如异步并发执行中所述，最大化主机、设备以及连接主机和设备的总线之间的并行执行。它应该将每种工作分配给最擅长的处理器：将串行工作负载分配给主机；将并行工作负载分配给设备。
对于并行工作负载，在算法中并行性被打破的点，因为一些线程需要同步以与其他线程共享数据，有两种情况：要么这些线程属于同一个块，这种情况下它们应该使用__syncthreads()并在同一个内核调用中通过共享内存共享数据，要么它们属于不同的块，这种情况下它们必须使用两个单独的内核调用通过全局内存共享数据，一个用于写入，一个用于从全局内存读取。第二种情况不太理想，因为它增加了额外的内核调用和全局内存流量的开销。因此，应该尽量通过将算法映射到CUDA编程模型，使需要进行线程间通信的计算尽可能地在一个线程块内进行。
设备层级 Device Level
在lower level上，应用程序应该最大化设备的多处理器之间的并行执行。
多个内核可以在一个设备上并发执行，因此也可以使用流来实现足够的内核并发执行，如异步并发执行中所述。
多处理器层级 Multiprocessor
在更低的层次上，应用程序应该最大化多处理器内各个功能单元之间的并行执行。
如硬件多线程中所述，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版本特别有用，可以将影响占用率的参数的更改可视化（块大小、每线程寄存器和每线程共享内存）。
最大化内存吞吐量
要最大化应用程序的整体内存吞吐量的第一步是尽量避免低带宽的数据传输。
这意味着要最小化host和device之间的数据传输，如主机与设备之间的数据传输中详细描述的那样，因为这些传输的带宽远低于设备与global memory之间的数据传输。
这也意味着通过最大化使用片上内存来最小化全局内存和设备之间的数据传输：共享内存和缓存（即2.x及更高计算能力的设备上可用的L1缓存和L2缓存，以及所有设备上可用的纹理缓存和常量缓存）。
共享内存相当于用户管理的缓存：应用程序需要显式地分配和访问它。如CUDA运行时中所示，典型的编程模式是将来自设备内存的数据暂存到共享内存中；换句话说，让一个块中的每个线程：
- 从设备内存加载数据到共享内存，
- 与块中的所有其他线程同步，以便每个线程可以安全地读取由不同线程填充的共享内存位置，
- 处理共享内存中的数据，
- 如果有必要，再次同步以确保共享内存已更新结果，
- 将结果写回设备内存。
对于某些应用程序（例如，全局内存访问模式是数据依赖的），传统的硬件管理缓存更适合利用数据局部性。如计算能力7.x、计算能力8.x和计算能力9.0中所述，对于计算能力为7.x、8.x和9.0的设备，相同的片上内存用于L1和共享内存，每个内核调用的L1与共享内存之间的分配方式是可配置的。
内核访问内存的吞吐量可能会因每种内存的访问模式而变化一个数量级。因此，最大化内存吞吐量的下一步是根据设备内存访问中描述的最佳内存访问模式尽可能地组织内存访问。这种优化对于全局内存访问尤为重要，因为与可用的片上带宽和算术指令吞吐量相比，全局内存带宽较低，因此非最优的全局内存访问通常对性能影响很大。
主机与设备之间的数据传输
应用程序应努力最小化主机和设备之间的数据传输。实现这一目标的一种方法是将更多的代码从主机移动到设备，即使这意味着运行的内核并没有暴露出足够的并行性来在设备上以全效率执行。可以在设备内存中创建中间数据结构，由设备操作，并销毁，而不需要由主机映射或复制到主机内存。
此外，由于与每次传输相关的开销，将许多小传输批量到一个大传输总是比单独进行每次传输性能更好。
在具有前端总线的系统上，使用页面锁定的主机内存可以实现主机和设备之间的数据传输的更高性能，如页面锁定主机内存中所述。
此外，当使用映射的页面锁定内存（映射内存）时，无需分配任何设备内存并在设备和主机内存之间显式复制数据。每次内核访问映射的内存时，都会隐式地执行数据传输。为了获得最大的性能，这些内存访问必须像访问全局内存一样合并（参见设备内存访问）。假设它们是合并的，并且映射的内存只被读取或写入一次，使用映射的页面锁定内存代替在设备和主机内存之间的显式复制可以提高性能。
在集成系统上，设备内存和主机内存在物理上是相同的，主机和设备内存之间的任何复制都是多余的，应该使用映射的页面锁定内存代替。应用程序可以通过检查集成设备属性（参见设备枚举）是否等于1来查询设备是否为集成。
设备内存访问
访问可寻址内存的指令（即，全局、局部、共享、常量或纹理内存）可能需要根据warp内的内存地址在线程之间的分布多次重新发出。这种分布如何影响指令吞吐量是特定于每种内存类型的，并在以下各节中描述。例如，对于全局内存，一般规则是，地址分散得越多，吞吐量降低得越多。
全局内存 Global Memory
全局内存位于设备内存中，设备内存通过32字节、64字节或128字节的内存事务（memory transactions）进行访问。这些内存事务（memory transactions）必须自然对齐：只有设备内存的32字节、64字节或128字节段，它们的大小对齐（即，它们的第一个地址是它们大小的倍数）可以通过内存transactions读取或写入。
当warp执行访问全局内存的指令时，它会根据每个线程访问的字大小和内存地址在线程之间的分布，将warp内的线程的内存访问合并为一个或多个这样的内存transactions。一般来说，需要的transactions越多，除了线程访问的字之外，还会传输更多的未使用的字，从而相应地降低指令吞吐量。例如，如果为每个线程的4字节访问生成一个32字节的内存事务，吞吐量就会除以8。
transactions的数量和吞吐量的影响程度随设备的计算能力而异。计算能力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字节，并且数据自然对齐（即，其地址是该大小的倍数）时，访问全局内存中的数据的任何访问（通过变量或指针）都会编译为单个全局内存指令。
如果不满足此大小和对齐要求，访问将编译为具有交错访问模式的多个指令，这些指令阻止这些指令完全合并。因此，建议使用满足此要求的类型用于位于全局内存中的数据。
内置向量类型 自动满足对齐要求。
对于结构，可以使用对齐指定符__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()) 被替换为分配一个大块的内存，该内存被划分为多个数组，这种情况下，每个数组的起始地址从块的起始地址偏移。
二维数组
常见的全局内存访问模式是每个索引为(tx,ty)的线程使用以下地址访问宽度为width的2D数组中的一个元素，该数组位于类型为type*的BaseAddress地址（其中type满足最大化利用中描述的要求）：
BaseAddress + width * ty + tx
为了使这些访问完全合并，线程块的宽度和数组的宽度都必须是warp大小的倍数。
特别地，这意味着如果一个数组的宽度不是这个大小的倍数，那么它将被更有效地访问，即使它实际上是以这个大小的最接近的倍数四舍五入的宽度分配的，并且其行相应地填充。cudaMallocPitch()和cuMemAllocPitch()函数及参考手册中描述的相关内存复制函数使程序员能够编写非硬件相关的代码来分配符合这些约束的数组。
- 具体地说，要访问宽度为
width的二维数组中的元素，其地址为BaseAddress + width * ty + tx。这里，BaseAddress是数组的起始地址，type*表示数组的数据类型指针，而type满足某些特定的要求（如在“Maximize Utilization”中描述的要求）。- 为了使内存访问完全合并（即高效地访问内存），线程块的宽度和数组的宽度都必须是warp大小的倍数。Warp是CUDA中的一个术语，表示一组并行执行的线程。
- 如果数组的宽度不是warp大小的倍数，那么为了更高效地访问它，应该将其宽度四舍五入到warp大小的最接近的倍数，并相应地填充其行。这意味着，如果原始数组的宽度不是warp大小的倍数，那么在分配内存时，应该为其分配更多的空间，以确保其宽度是warp大小的倍数。
- cudaMallocPitch()和cuMemAllocPitch()函数：这两个函数允许程序员分配满足上述约束的数组，而不依赖于硬件的特定细节。这些函数在分配内存时会考虑到最佳的内存访问模式，并返回一个适当的宽度（即四舍五入到warp大小的最接近的倍数）。与这两个函数相关的内存复制函数则允许程序员复制数据到这些特殊分配的数组中。
局部内存 local memory
只有在变量内存空间指定符中提到的一些自动变量才会发生局部内存访问。编译器可能会将以下自动变量放入局部内存：
- 它不能确定它们是用常量量索引的数组，
- 大型结构或数组会消耗太多的寄存器空间，
- 如果内核使用的寄存器比可用的寄存器多（这也称为寄存器溢出）。
PTX汇编代码的检查（通过使用-ptx或-keep选项编译获得）将告诉我们在第一编译阶段是否已将变量放入局部内存，因为它将使用.local助记符声明并使用ld.local和st.local助记符访问。即使没有，如果它们发现它为目标架构消耗了太多的寄存器空间，后续的编译阶段仍然可能决定另行处理：使用cuobjdump检查cubin对象将告诉我们是否是这种情况。此外，当使用--ptxas-options=-v选项编译时，编译器报告每个内核的总局部内存使用量（lmem）。注意，某些数学函数具有可能访问局部内存的实现路径。
局部内存空间位于设备内存中，因此局部内存访问与全局内存访问具有相同的高延迟和低带宽，并受到设备内存访问中描述的内存合并的相同要求。但是，局部内存的组织方式是连续的32位字由连续的线程ID访问。只要warp中的所有线程访问相同的相对地址（例如，在数组变量中相同的索引，结构变量中的相同成员），访问就是完全合并的。
在计算能力5.x及更高版本的设备上，局部内存访问始终以与全局内存访问相同的方式在L2中缓存（参见计算能力5.x和计算能力6.x）。
共享内存
由于它位于片上，共享内存的带宽要高得多，延迟要低得多。
为了实现高带宽，共享内存被划分为大小相等的内存模块，称为bank，这些bank可以同时访问。因此，由n地址组成的任何内存读取或写入请求，这些地址在n个不同的内存bank中，因此可以同时服务，从而产生的总带宽是单个模块的带宽的n倍。
但是，如果一个内存请求的两个地址落在同一个内存bank中，就会发生bank冲突，访问必须被序列化。硬件将带有bank冲突的内存请求分割为多个无冲突的请求，这些请求的数量与所需的数量相同，从而使吞吐量降低了一个与单独内存请求数量相等的因子。如果单独的内存请求的数量是n，则初始内存请求被称为引起n-way银行冲突。
为了获得最大的性能，因此重要的是了解内存地址如何映射到内存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位浮点值（参见纹理内存）。
提升指令处理速度
为了提升程序的指令处理速度，我们应该：
- 尽量减少使用计算速度慢的算术指令；例如在精度损失可以接受的情况下提高计算速度，可以使用内建函数（intrinsic functions，详情见Intrinsic Functions），使用单精度计算而非双精度（fp16占用更少的内存和带宽，并且硬件上执行fp16操作的速度比fp32更快），将非规格化数（denormalized numbers）刷新为零可以防止它们导致较低的算术性能。
- 按照Control Flow Instructions的指导，减少因控制流指令造成的warp不一致。分歧线程束是指在执行控制流指令时，线程束内的线程采取了不同的执行路径，这会降低执行效率；
- 减少指令的总数，比如说，根据Synchronization Instruction的建议，尽可能减少同步点，或者按照restrict的描述使用限制指针。 优化掉同步点可以减少因同步而浪费的时间，并增加程序的执行效率、使用
__restrict__关键字声明的限定指针可以告诉编译器这个指针是访问其指向的数据的唯一方式，这允许编译器进行更多的优化；
在这一部分，吞吐量是以每个多处理器每个时钟周期能完成的操作数量来衡量的。如果warp的大小是32，那么一个指令相当于32个操作。如果每个时钟周期有N个操作，那么指令吞吐量就是N除以32个指令每时钟周期。
这里的所有吞吐量数据都是基于单个多处理器的。要得到整个设备的吞吐量数据，需要把这个数字乘以设备里多处理器的数量。
算术指令
下表提供了不同计算能力设备硬件所支持的算术指令的处理速度数据。
因为格式问题，可以查看原文：CUDA C++ Programming Guide
其他指令和功能是基于原生指令实现的。对不同计算能力的设备，其实现可能各不相同，并且在每个编译器版本中，编译后的原生指令数量也可能发生变化。对于一些复杂的函数，根据输入的不同，可能存在多个代码执行路径。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用户手册中找到。
单精度浮点数除法
__fdividef(x, y)（参见Intrinsic Functions）比普通的除法运算符提供更快的单精度浮点数除法。
单精度浮点数的倒数平方根
为了保持IEEE-754的语义，只有在倒数和平方根都是近似的情况下（即，使用-prec-div=false和-prec-sqrt=false），编译器才会把1.0/sqrtf()优化成rsqrtf()。因此，建议在需要的地方直接调用rsqrtf()。
单精度浮点数平方根
为了得到0和无穷大的正确结果，单精度浮点数平方根是通过先计算倒数平方根，然后再取倒数来实现的，而不是先计算倒数平方根然后再乘以一个数来实现的。
正弦和余弦
sinf(x)，cosf(x)，tanf(x)，sincosf(x)以及相应的双精度指令的计算成本要高得多，特别是当参数x的绝对值很大时。
更具体地说，参数简化代码（可以参见Mathematical Functions来查看具体实现）包括两个被称为快速路径和慢速路径的代码执行路径。
快速路径用于那些绝对值较小的参数，并主要包括一些乘法和加法操作。而慢速路径则用于那些绝对值较大的参数，并涉及到一系列复杂的计算，以在整个参数范围内得到正确的结果。
目前，对于绝对值小于105615.0f的参数，单精度函数会选择快速路径；而对于绝对值小于2147483648.0的参数，双精度函数会选择快速路径。
由于慢速路径需要比快速路径更多的寄存器，因此已经尝试通过将一些中间变量存储在本地内存中来减少慢速路径的寄存器压力。这可能会因本地内存的高延迟和带宽而影响性能（详见Device Memory Accesses）。目前，单精度函数使用了28字节的本地内存，而双精度函数使用了44字节的本地内存。然而，这个数字可能会发生变化。
由于慢速路径中的复杂计算和本地内存的使用，当需要慢速路径时，这些三角函数的吞吐量比需要快速路径时要低一个数量级。
整数算术
整数的除法和模运算非常消耗资源，因为它们会被编译成多达20个指令。在某些情况下，它们可以被位运算替换：如果n是2的幂，那么（i/n）等价于（i>>log2(n)），（i%n）等价于（i&(n-1））；如果n是一个常数，编译器将进行这些转换。
__brev和__popc分别对应一个指令，而__brevll和__popcll对应几个指令。
__[u]mul24是一些已经没必要再使用的过时内建函数。
半精度算术
为了实现16位精度浮点数的加、乘和乘加运算的良好性能，建议使用half2数据类型表示half精度，使用__nv_bfloat162表示__nv_bfloat16精度。然后，可以使用向量内建函数（例如，__hadd2，__hsub2，__hmul2，__hfma2）在一个指令中完成两个操作。使用half2或__nv_bfloat162代替两个half或__nv_bfloat16的调用也可能提高其他内建函数的性能，比如warp shuffle。
提供了__halves2half2内建函数，用于将两个half精度值转换为half2数据类型。
还提供了__halves2bfloat162内建函数，用于将两个__nv_bfloat精度值转换为__nv_bfloat162数据类型。
类型转换
有时，编译器可能需要插入转换指令，从而引入额外的执行周期。这主要出现在以下情况：
- 一些作用于
char或short类型变量的函数，这些函数的操作数通常需要转换为int。 - 双精度浮点常数（即那些没有任何类型后缀的常数）被用作单精度浮点计算的输入（按照C/C++标准的要求）。
为了避免最后一种情况，可以使用单精度浮点常数，这些常数带有f后缀，比如3.141592653589793f，1.0f，0.5f。
控制流指令
控制流指令如if、switch、do、for和while可能对有效指令的吞吐量产生重大影响，原因是这些指令可能使同一“warp”（一种线程组）内的线程走向不同的执行路径，导致路径分歧。一旦发生这种情况，不同的执行路径就必须一个接一个地执行，从而增加了该“warp”需要执行的指令总数。
为了在控制流依赖于线程ID的情况下实现最佳性能，应当编写控制条件以尽量减少路径分歧的“warp”数量。由于“warp”在“block”（线程块）内的分布是有规律的（正如SIMT架构中提及的），所以这是可行的。一个简单的例子是控制条件只依赖于（threadIdx / warpSize），其中“warpSize”是“warp”的大小。在此情况下，由于控制条件与“warp”完全对齐，所以没有“warp”会发生路径分歧。
有时，编译器可能会展开循环，或者通过使用分支预测来优化短if或switch块，从而防止任何“warp”发生分歧。程序员还可以通过使用#pragma unroll指令来控制循环展开（参见#pragma unroll）。
使用分支预测时，控制条件相关的指令都不会被跳过。相反，每条指令都与一个线程相关的条件代码或谓词关联，这些条件代码或谓词会根据控制条件被设为真或假。尽管所有指令都已被安排执行，但实际上只有那些具有真谓词的指令才会被执行。具有假谓词的指令不会产生结果，也不会计算地址或读取操作数。
同步指令
对于计算能力为6.0的设备，__syncthreads()的吞吐量是每个时钟周期32个操作，对于计算能力为7.x和8.x的设备是每个时钟周期16个操作，对于计算能力为5.x、6.1和6.2的设备是每个时钟周期64个操作。
这里的吞吐量是指每个时钟周期内
__syncthreads()函数可以同步的操作数。不同计算能力的GPU设备，其硬件架构和资源分配有所不同，因此__syncthreads()的效率也有所不同。
需要注意的是，__syncthreads()可能会影响性能，原因是它强制多处理器进入空闲状态，具体详见设备内存访问。
避免内存频繁申请和释放
不断分配和释放内存的应用程序在运行过程中，随着时间的推移，内存分配的速度会逐渐变慢，直至达到一个限制。这种现象通常是因为内存释放回操作系统后，系统会按照其自身的机制使用这些内存，导致分配内存的速度下降。为了获得更好的性能表现，我们有以下建议：
- 根据当前问题的规模来分配内存。不要试图通过
cudaMalloc/cudaMallocHost/cuMemCreate分配所有可用的内存，因为这样会立刻使内存占用，且阻止其他应用使用这部分内存。这不仅会增加操作系统调度的压力，还可能完全阻止使用相同GPU的其他应用程序的运行。 - 尽量在应用程序开始时就按需分配适量的内存，并只在应用程序不使用内存时才释放它。减少应用程序中
cudaMalloc和cudaFree调用的次数，特别是在对性能要求较高的代码区域。 - 如果应用程序无法分配足够的设备内存，可以考虑使用
cudaMallocHost或cudaMallocManaged这类可能性能较差但可以使应用程序继续运行的内存类型作为备选方案。 - 对于支持此功能的平台，
cudaMallocManaged允许内存的超额订阅。只要启用了正确的cudaMemAdvise策略，cudaMallocManaged就可以保持与cudaMalloc相当的性能。而且，cudaMallocManaged不会强制内存驻留，直到需要或被预取，从而减轻了操作系统调度的压力，并更好地支持多租户使用场景。