管理内存是编写 GPU 内核时影响性能的关键因素之一。本文将为您介绍全局内存及其对性能影响的重要知识。
全局内存
CUDA 设备上存在多种类型的显存，它们在作用范围、生命周期和缓存行为方面各有不同。全局内存（也称为设备内存）是 CUDA 设备中的主要内存空间，存储于设备的 DRAM 中，其功能类似于 CPU 系统中的主存。这里的“全局”指的是其作用范围，即主机以及内核网格中的所有线程均可访问和修改该内存空间。
可以使用 `static` 声明符在全局范围内静态声明全局内存，也可以通过 CUDA 运行时 API（如 `malloc` 或 `cudaMalloc`）动态分配全局内存。可利用 `cudaMemcpy` 将数据从主机传输到设备，并通过 `cudaFree` 释放内存。这些动态分配的内存将在释放之前一直有效。
还可以通过统一内存来实现全局内存的分配与释放。全局内存的分配与释放，以及数据在设备间移动的过程涉及诸多复杂因素，我们将在后续文章中详细探讨。本文则聚焦于全局内存的使用对 CUDA 内核性能的影响。
一个典型的使用模式示例是：主机在核函数启动前分配并初始化全局内存，随后执行核函数，由CUDA线程读取数据、计算并将结果写回全局内存，最后在核函数执行完毕后，主机读取并获取结果。
示例：动态分配、传输、内核处理与资源清理
// Host allocates global memory
float* d_input;
float* d_output;
cudaMalloc(&d_input, n * sizeof(float));
cudaMalloc(&d_output, n * sizeof(float));
// Transfer data to device
cudaMemcpy(d_input, h_input, n * sizeof(float), cudaMemcpyHostToDevice);
// Call a kernel to operate on the device
someKernel<<<1024, 1024>>>(d_input, d_output, n);
// Copy the result back to the host
cudaMemcpy(h_output, d_output, n * sizeof(float), cudaMemcpyDeviceToHost);
// Cleanup
cudaFree(d_input);
cudaFree(d_output);
全局内存合并
在讨论全局内存访问性能之前，有必要进一步理解 CUDA 的执行模型。我们已经了解了如何将线程组织成线程块，并将这些线程块分配到设备的多处理器上。在实际执行过程中，线程还会被进一步细分为更小的单元，称为线程束。GPU 上的多处理器以 SIMT（单指令多线程）方式执行每个线程束的指令。目前所有支持 CUDA 的 GPU，其线程束的大小（即 SIMT 宽度）均为 32 个线程。
在 CUDA 中访问全局内存时，需要重点关注同一线程束中各线程所访问内存位置之间的关系。这种内存访问模式会直接影响内存访问的效率，进而影响应用程序的整体性能。
全局内存通过 32 字节的内存事务进行访问。当 CUDA 线程请求全局内存中的数据时，该线程束内所有线程的内存访问会被合并为尽可能少的内存事务。所需内存事务的数量取决于每个线程访问的数据大小以及各线程间内存地址的分布情况。
以下代码演示了一种场景：线程束中的连续线程访问连续的 4 字节数据元素，从而形成高效的内存访问模式。所有线程束发出的加载请求均可由内存中的四个 32 字节扇区满足，有效提升了内存带宽的利用率。图 1 展示了每个线程如何访问连续内存中的 4 字节数据元素。
__global__ void coalesced_access(float* input, float* output, int n) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < n) {
// Each thread accesses consecutive 4-byte words
output[tid] = input[tid] * 2.0f ;
}
}
相反，当线程以较大的步长访问内存时，每次内存事务所获取的数据量会远超实际所需。对于每个线程请求的 4 字节元素，系统会从全局内存中读取完整的 32 字节内存段，导致其中大部分数据并未被使用。图 2 展示了这种访问模式的一个示例。
__global__ void uncoalesced_access(float* input, float* output, int n) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < n) {
// Access with a stride of 32 (128 bytes), wrapped around to stay within bounds
int scattered_index = (tid * 32) % n;
output[tid] = input[scattered_index] * 2.0f;
我们来深入探讨如何利用 NVIDIA Nsight Compute（NCU）分析两种截然不同的 CUDA 内核的内存访问模式。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展示了上述具有不同访问步长的内核的带宽测量结果。该图的目的并非体现所能达到的极限带宽，而是为了说明在以跨步方式访问全局内存时，简单内核的带宽变化情况。
本图显示，当线程束中的线程访问物理内存中距离较远的地址时，硬件难以高效合并这些内存访问，导致预期的带宽显著降低。
多维数组
现在，我们来讨论多维数组（或矩阵）的内存访问问题。为了实现良好的性能并促进内存访问的合并，确保连续的线程访问数组中的连续元素至关重要，这一点与一维情况类似。
在 CUDA 核函数中，当使用二维或三维线程块时，线程的布局按维度依次展开：X 方向（threadIdx.x）的索引变化最快，其次是 Y 方向（threadIdx.y），最后是 Z 方向（threadIdx.z）。例如，若有一个大小为 24×24 的二维线程块，则线程的排列顺序将按照 X 方向优先、Y 方向次之的方式进行线性映射。
在访问矩阵等二维数据时，CUDA 中通常采用二维线程块。由于 C++ 以行主序方式存储二维数据，行方向上的内存访问是连续的。若能使连续的线程访问矩阵某一行中连续的内存位置，则这些内存访问将高度高效（即合并访问）；而列方向的访问则会导致跨步访问，无法合并，因而效率较低。
由于线程束中连续的 threadIdx.x 值应访问连续的内存元素以实现内存合并，因此具有相同 threadIdx.y 值的线程应访问矩阵的同一行。这种访问模式符合矩阵按行优先存储的内存布局，有助于线程束以更高效的方式完成合并内存访问，从而提升内存带宽的利用率。
对于遵循内存访问模式的合并核（coalesced_matrix_access），由于线程索引到矩阵坐标的映射方式（基于行主序存储），能够实现高效的合并内存访问。具体而言，每个线程块的 x 维度（threadIdx.x）被分配为列索引，这意味着当线程束中的连续线程递增其 threadIdx.x 时，它们会访问同一行内连续的列。由于行主序存储将同一行中的元素放置在连续的内存地址中，这种访问模式使得线程束中的每个线程恰好读取相邻的内存位置，从而实现内存的连续访问。
__global__ void coalesced_matrix_access(float* matrix, int width, int height)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < height && col < width) {
int idx = row * width + col; // row-major ⇒ coalesced
matrix[idx] = matrix[idx] * 2.0f + 1.0f;
}
}
对于下文所示的未合并内核（uncoalesced_matrix_access），其内存访问模式将导致访问效率低下。
__global__ void uncoalesced_matrix_access(float* matrix, int width, int height)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
if (row < height && col < width) {
int idx = col * height + row; // column-major ⇒ uncoalesced
matrix[idx] = matrix[idx] * 2.0f + 1.0f;
}
}
为说明这一点，核函数通过索引计算 col * height + row，人为地将行主序存储的矩阵视为列主序进行访问。这意味着，当线程束中的连续线程递增其 threadIdx.x（即列索引）时，它们所访问的元素在列主序布局下是连续的，但在实际的行主序内存布局中却是间隔排列的。由于数据在物理内存中按行主序存储，而访问时却采用列主序的索引方式，导致连续线程访问的内存地址之间相隔 height 个元素，形成大步长的访问模式，使得 GPU 无法将这些内存访问合并为高效的内存事务（如图5所示）。这种存储顺序与访问模式之间的不匹配，显著降低了全局内存的带宽利用率。
我们可以通过分析以下结果来观察这一行为：
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 个扇区，说明其内存访问呈跨步模式，导致内存子系统不得不读取远超实际需求的扇区，造成资源浪费。
总结
高效利用 GPU 显存是提升性能的关键因素之一。要实现优异的全局内存性能，关键在于实现内存访问的合并。应尽量减少对全局内存的跨步访问，并始终借助 Nsight Compute 对 GPU 内核进行分析，以确保内存访问模式的合并性。通过这种方法，可以显著提升 GPU 代码的执行效率。
致谢
这篇博文是对 NVIDIA 的 Mark Harris 于 2013 年首次发布的文章的更新版本。
