解锁GPU性能：CUDA全局内存访问优化指南-开发者社区

管理内存是编写GPU内核时需要考虑的最重要的性能特征之一。本文将引导您了解关于全局内存及其性能的重要方面。

全局内存

CUDA设备上有多种类型的内存，每种都有不同的作用域、生命周期和缓存行为。全局内存（也称为设备内存）是CUDA设备上的主要内存空间。它驻留在设备的DRAM中，其功能类似于CPU系统中的RAM。“全局”一词指的是其作用域，它既可以被主机访问和修改，也可以被内核网格中的所有线程访问和修改。

全局内存可以使用__device__声明说明符在全局作用域中静态声明，或者使用CUDA运行时API（如cudaMalloc()或cudaMallocManaged()）动态分配。数据可以使用cudaMemcpy()从主机传输到设备，并使用cudaFree()释放。这些分配在被释放之前是持久存在的。

全局内存也可以通过统一内存进行分配/释放。全局内存的分配/释放以及与设备之间的数据传输是一个复杂的话题，将在后续文章中探讨。在本文中，我们将重点关注在CUDA内核中使用全局内存的性能影响。

一个典型使用模式的简单示例包括：主机在内核启动前分配并初始化全局内存；接着内核执行，CUDA线程从全局内存读取数据并将结果写回全局内存；最后在内核完成后，主机检索结果。

示例：动态分配、传输、内核执行与清理

// 主机分配全局内存float*d_input;float*d_output;cudaMalloc(&d_input,n*sizeof(float));cudaMalloc(&d_output,n*sizeof(float));// 将数据传输到设备cudaMemcpy(d_input,h_input,n*sizeof(float),cudaMemcpyHostToDevice);// 调用内核在设备上执行操作someKernel<<<1024,1024>>>(d_input,d_output,n);// 将结果复制回主机cudaMemcpy(h_output,d_output,n*sizeof(float),cudaMemcpyDeviceToHost);// 清理cudaFree(d_input);cudaFree(d_output);

全局内存合并

在深入探讨全局内存访问性能之前，我们需要细化对CUDA执行模型的理解。我们已经讨论过线程如何分组到线程块中，这些线程块被分配到设备上的多处理器。在执行过程中，线程会被更精细地分组为线程束（warp）。GPU上的多处理器以SIMT（单指令多线程）方式为每个线程束执行指令。所有当前支持CUDA的GPU的线程束大小（实际上是SIMT宽度）是32个线程。

在CUDA中访问全局内存时，您需要考虑的一个关键方面是同一线程束内不同线程所访问的内存位置之间的关系。这些内存访问的模式直接影响内存访问效率和整体应用程序性能。

全局内存通过32字节的内存事务进行访问。当CUDA线程从全局内存请求数据时，该线程束中所有线程的内存访问会被合并成最少次数的内存事务。所需内存事务的数量取决于每个线程访问的字的大小以及这些内存地址在线程间的分布情况。

以下代码演示了一个场景：线程束内的连续线程访问连续的4字节数据元素，创建了最优的内存访问模式。线程束发出的所有加载操作都可以通过内存中的四个32字节扇区来满足，这允许最有效地利用内存带宽。图1显示了每个线程如何访问内存中连续的4字节数据元素。

__global__voidcoalesced_access(float*input,float*output,intn){inttid=blockIdx.x*blockDim.x+threadIdx.x;if(tid<n){// 每个线程访问连续的4字节字output[tid]=input[tid]*2.0f;}}

图1. 合并内存访问模式，显示一个线程束的线程（箭头）访问连续的128字节内存块，分为四个32字节扇区。

相反，如果线程以较大的步幅访问内存，每个内存事务获取的数据量远超过所需。对于每个线程请求的每个4字节元素，都会从全局内存获取整个32字节的扇区，而大部分传输的数据未被使用。图2显示了这种模式的示例。

__global__voiduncoalesced_access(float*input,float*output,intn){inttid=blockIdx.x*blockDim.x+threadIdx.x;if(tid<n){// 以步幅32（128字节）进行访问，回绕以保持在边界内intscattered_index=(tid*32)%n;output[tid]=input[scattered_index]*2.0f;}}

图2. 非合并内存访问模式，显示每个线程（箭头）在单独的32字节内存扇区中访问数据。

让我们深入分析这两个对比鲜明的CUDA内核的内存访问模式，使用某机构的Nsight Compute（NCU）工具。NCU提供了强大的指标来量化内存访问模式。

开始分析内核时，我们通常运行：

ncu --set full --print-details=all ./a.out

此命令收集所有可用的分析部分，包括内存、指令、启动、占用率、缓存等。然而，当特别关注内存访问效率时，我们将其缩小到量化内存工作负载模式的指标。要仅隔离与内存工作负载相关的细节，以下命令更合适：

ncu --section MemoryWorkloadAnalysis_Tables --print-details=all ./a.out

此命令的输出如下所示，为清晰起见已简化。

coalesced_access(float *, float *, int) (262144, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 uncoalesced_access(float *, float *, int) (262144, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 Section: Memory Workload Analysis Tables OPT Est. Speedup: 83% The memory access pattern for global loads from DRAM might not be optimal. On average, only 4.0 of the 32 bytes transmitted per sector are utilized by each thread. This applies to the 100.0% of sectors missed in L2. This could possibly be caused by a stride between threads. Check the Source Counters section for uncoalesced global loads.

从输出中，我们可以看到NCU已经识别出uncoalesced_access内核在全局加载方面存在性能改进的空间，事实上它指出我们平均只利用了每个获取的32字节扇区中的4个字节。NCU甚至暗示“这可能是由线程间的步幅引起的”。

我们专门设置这个问题来说明良好和糟糕的内存性能，所以这并不令人惊讶。为了进一步深入，我们可以看看NCU还能提供哪些其他类型的内存分析表。

由于NCU的初始输出指出了从DRAM加载的问题，接下来我们将尝试这个命令来更深入地研究DRAM统计数据。

ncu --metrics group:memory__dram_table ./a.out

coalesced_access(float *, float *, int) (262144, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 Section: Command line profiler metrics --------------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value --------------------------------------------------- ----------- ------------ dram__bytes_read.sum Mbyte 268.44 dram__bytes_read.sum.pct_of_peak_sustained_elapsed % 46.76 dram__bytes_read.sum.per_second Gbyte/s 159.76 dram__bytes_write.sum Mbyte 248.50 dram__bytes_write.sum.pct_of_peak_sustained_elapsed % 43.28 dram__bytes_write.sum.per_second Gbyte/s 147.89 dram__sectors_read.sum sector 8,388,900 dram__sectors_write.sum sector 7,765,572 --------------------------------------------------- ----------- ------------ uncoalesced_access(float *, float *, int) (262144, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 8.9 Section: Command line profiler metrics --------------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value --------------------------------------------------- ----------- ------------ dram__bytes_read.sum Gbyte 2.15 dram__bytes_read.sum.pct_of_peak_sustained_elapsed % 84.92 dram__bytes_read.sum.per_second Gbyte/s 290.16 dram__bytes_write.sum Mbyte 263.70 dram__bytes_write.sum.pct_of_peak_sustained_elapsed % 10.43 dram__bytes_write.sum.per_second Gbyte/s 35.63 dram__sectors_read.sum sector 67,110,368 dram__sectors_write.sum sector 8,240,680 --------------------------------------------------- ----------- ------------

通过这个结果，我们可以看到两个内核的dram__sectors_read.sum输出之间存在巨大差异。我们的内核读取一个数组然后写回同一个数组，所以读取的数据量应该与写入的数据量相同，但在非合并的情况下，我们看到sectors_read与sectors_write之间存在8倍的差异。

现在让我们使用这个命令分析L1行为：

ncu --metrics group:memory__first_level_cache_table ./a.out

此命令会输出大量信息，我们在此省略了，但如果您运行它，关键是注意两个内核之间不同的指标。我们想进一步调查其中两个：l1tex_t_requests_pipe_lsu_mem_global_op_ld.sum和l1tex_t_sectors_pipe_lsu_mem_global_op_ld.sum。NCU提供了一个表格来帮助您解码这些指标收集的信息。第一个指标本质上是发出的内存请求数量，第二个指标是获取的扇区数量。

在分析GPU内核的内存效率时，扇区（从内存传输的32字节数据块）和请求（由线程束发起的内存事务）为了解内存合并行为提供了宝贵的见解。扇区与请求的比率清晰地展示了代码如何有效利用内存系统。

如果我们使用以下命令，可以仅收集这两个指标：

ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum ./a.out

我们获得的输出是：

coalesced_access(float *, float *, int) (262144, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 9.0 Section: Command line profiler metrics ----------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ----------------------------------------------- ----------- ------------ l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum 2097152 l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum sector 8388608 ----------------------------------------------- ----------- ------------ uncoalesced_access(float *, float *, int) (262144, 1, 1)x(256, 1, 1), Context 1, Stream 7, Device 0, CC 9.0 Section: Command line profiler metrics ----------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ----------------------------------------------- ----------- ------------ l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum 2097152 l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum sector 67108864 ----------------------------------------------- ----------- ------------

在合并内核中，请求与扇区的比率是1:4，这正是我们预期的。回想图1，我们展示了一个完美合并的内存事务：128字节将需要四个32字节的扇区。从内存获取的每个字节都被内核使用，实现了100%的内存带宽效率。

在非合并内核中，请求与扇区的比率是1:32，这也是我们预期的，回想图2，每个线程从不同的32字节扇区请求4个字节。因此，线程束的每个请求都需要32个扇区。虽然内存系统获取了32个扇区（总共1024字节），但每个线程只需要其各自扇区中的4个字节。

这8倍的效率差异对GPU性能有深远的影响，因为内存带宽通常决定了GPU内核的最终性能极限。有关分析的更多信息，包括内存扇区，可以在“分析指南”部分找到。

步幅访问

现在让我们看看步幅对内存带宽的影响。在CUDA内存访问模式的上下文中，步幅指的是线程束中的线程访问的连续内存位置之间的距离（以数组元素或字节为单位）。

如上所示的具有不同访问步幅的内核的带宽测量结果如图3所示。这并非旨在显示可实现的最大带宽，而只是为了展示当对全局内存的访问有步幅时，简单内核的带宽如何变化。

图3. GH200上步幅从0到31的带宽与步幅关系图，显示带宽值递减。

图表显示，对于大步幅，有效带宽很差，正如预期的那样。当线程束中的线程访问物理内存中相距较远的内存地址时，硬件无法有效地合并这些访问。

多维数组

现在我们来讨论多维数组或矩阵情况下的内存访问。为了获得最佳性能并实现合并的内存访问，连续线程访问数组中的连续元素非常重要，就像在一维情况下一样。

在CUDA内核中使用二维或三维线程块时，线程按线性排列，X索引（threadIdx.x）变化最快，然后是Y（threadIdx.y），最后是Z（threadIdx.z）。例如，如果我们有一个大小为(4,2)的二维线程块，线程的顺序将是：(0,0)(1,0)(2,0)(3,0)(0,1)(1,1)(2,1)(3,1)。

在CUDA中，当访问二维数据（如矩阵）时，通常使用二维线程块。当我们考虑使用二维线程块访问矩阵（以1D内存数组存储）时，由于C++以行主序形式存储2D数据，因此行访问是连续的。如果我们能让连续的线程连续访问行中的内存位置，那么这些访问将是高效的（合并的），而列访问则是低效的（有步幅，非合并的）。

由于线程束内连续的threadIdx.x值应该访问连续的内存元素以实现合并，具有相同threadIdx.y值的线程应该访问矩阵的一行。这确保了当线程束中的线程访问矩阵元素时，它们遵循自然的行主序内存布局，从而实现高效的合并内存事务并最大化内存带宽利用率。

对于遵循内存访问模式的内核（coalesced_matrix_access），由于线程索引如何映射到矩阵坐标（给定行主序存储顺序），因此可以实现高效的合并访问。在这里，每个块的x维度（threadIdx.x）被分配给列索引，这意味着当线程束内的连续线程增加它们的threadIdx.x时，它们访问矩阵的连续列，同时保持在相同的行内（图4）。由于行主序将连续的内存位置存储为同一行内的元素，跨行访问允许线程束中的每个线程访问连续的内存位置。

__global__voidcoalesced_matrix_access(float*matrix,intwidth,intheight){introw=blockIdx.y*blockDim.y+threadIdx.y;intcol=blockIdx.x*blockDim.x+threadIdx.x;if(row<height&&col<width){intidx=row*width+col;// 行主序 ⇒ 合并matrix[idx]=matrix[idx]*2.0f+1.0f;}}

图4. 合并的2D访问，显示2D线程块如何映射到2D矩阵，以及它如何映射到矩阵所在的线性内存。连续线程访问连续的行元素，这些元素在内存中是连续的。

对于接下来显示的非合并内核（uncoalesced_matrix_access），内存访问模式导致了低效的非合并访问。

__global__voiduncoalesced_matrix_access(float*matrix,intwidth,intheight){introw=blockIdx.y*blockDim.y+threadIdx.y;intcol=blockIdx.x*blockDim.x+threadIdx.x;if(row<height&&col<width){intidx=col*height+row;// 列主序 ⇒ 非合并matrix[idx]=matrix[idx]*2.0f+1.0f;}}

这里，为了说明这一点，内核通过使用索引计算col * height + row，人为地将行主序矩阵视为列主序。这意味着当线程束内的连续线程增加它们的threadIdx.x（增加列索引）时，它们访问的是在列主序布局中连续的元素，但在行主序内存布局中是有步幅的。由于数据以行主序物理存储，但以列主序索引访问，连续的线程最终访问相隔height个元素的内存位置，创建了一个大的步幅模式，消除了GPU将这些访问合并为高效事务的能力（图5）。这种存储顺序和访问模式之间的不匹配导致全局内存带宽利用率低下。

图5. 非合并的2D访问，显示2D线程块如何映射到2D矩阵，以及它如何映射到矩阵所在的线性内存。连续线程访问连续的列元素，这些元素在内存中是不连续的。

我们可以通过检查下面的分析结果来观察这种行为：

coalesced_matrix_access(float *, int, int) (512, 512, 1)x(32, 32, 1), Context 1, Stream 7, Device 0, CC 9.0 Section: Command line profiler metrics ----------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ----------------------------------------------- ----------- ------------ l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum 8388608 l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum sector 33554432 ----------------------------------------------- ----------- ------------ uncoalesced_matrix_access(float *, int, int) (512, 512, 1)x(32, 32, 1), Context 1, Stream 7, Device 0, CC 9.0 Section: Command line profiler metrics ----------------------------------------------- ----------- ------------ Metric Name Metric Unit Metric Value ----------------------------------------------- ----------- ------------ l1tex__t_requests_pipe_lsu_mem_global_op_ld.sum 8388608 l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum sector 268435456 ----------------------------------------------- ----------- ------------

两个内核生成相同数量的内存请求（8,388,608），但合并版本只需要33,554,432个扇区，而非合并版本需要268,435,456个扇区。这转化为合并内核的每个请求扇区比率为4，而非合并内核为32。合并内核的低比率（每个请求4个扇区）表示高效的内存合并，由于连续的访问模式，GPU可以在更少的内存扇区内满足多个线程请求。相反，非合并内核的高比率（每个请求32个扇区）表明内存访问未合并，其中步幅访问模式迫使内存子系统获取比满足相同内存请求所需更多的扇区。