CUDA内存合并优化实操：极速搞定8倍性能提升

各位跨境电商、游戏、支付、贸易和广告行业的朋友们，大家好！作为一名深耕行业多年的老兵，今天咱们来聊聊GPU性能优化里的一个核心秘密——CUDA全局内存访问。别看这技术有点硬核，但它直接关系到你的程序跑得快不快、效率高不高，是咱们提升算力、降本增效的“武功秘籍”之一。

掘金GPU高性能：CUDA全局内存访问优化实战

CUDA设备上的全局内存，就好比咱们GPU系统里一块非常重要的“公共存储区”。它不仅容量大，而且主机（CPU）和所有在GPU上运行的线程都能读写它。咱们今天的重点，就是要手把手教大家如何高效利用这块公共区域，让GPU的性能飙起来。

全局内存，你了解多少？

在CUDA设备上，内存种类可不少，每种都有它独特的“脾气”——比如访问范围、生命周期和缓存策略等等。而全局内存，也就是大家常说的“设备内存（device memory）”，它可是CUDA设备上的“主干道”。它安家在设备的DRAM里，和咱们CPU系统里的RAM差不多。之所以叫“全局”，是因为它适用范围广，主机和所有在核函数网格（kernel grid）里的线程都能访问和修改它。

全局内存的声明方式有两种。一种是你在全局作用域里，用 __device__ 这个关键词来静态声明。另一种呢，就是咱们更常用、更灵活的方式，通过CUDA运行时API来动态分配，比如 cudaMalloc() 或 cudaMallocManaged()。数据从主机传到设备，用 cudaMemcpy()，用完了要记得 cudaFree() 释放掉，这些动态分配的内存在你手动释放前会一直存在。

当然，除了传统方式，咱们也可以用统一内存（Unified Memory）来分配和释放全局内存，这部分内容比较复杂，牵扯到数据如何在主机和设备之间“自由穿梭”，以后有机会咱们再深入讲解。今天，咱们就先聚焦在全局内存的使用，以及它对咱们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的执行模型。咱们之前聊过，线程们会聚集成线程块（thread blocks），这些线程块会被分配到设备上的多处理器（multiprocessors）去执行。而在执行过程中，还有更精细的线程分组，那就是“Warp（线程束）”。GPU上的多处理器，是以SIMT（单指令多线程）模式来执行每个Warp里的指令的。目前所有支持CUDA的GPU，Warp的大小都是32个线程。

划重点了，朋友们！当Warp里的不同线程访问全局内存时，它们访问的内存地址之间有什么关系，这可是个大讲究！这种内存访问模式，直接决定了内存访问的效率，进而影响咱们整个程序的性能表现。

全局内存是以32字节为单位进行内存事务处理的。当一个CUDA线程请求全局内存数据时，同一个Warp里所有线程的内存访问会被“合并”成最少数量的内存事务。需要多少个内存事务，就取决于每个线程访问的数据大小，以及这些内存地址在线程间的分散程度。

高效访问示例：合并式访问

下面这段代码就展示了一个理想场景：Warp里连续的线程访问连续的4字节数据元素。这会创造出一个最优的内存访问模式。一个Warp发出的所有加载请求，可以由四个32字节的内存扇区来满足，这样就能最大化地利用内存带宽。大家看图1，每个线程是如何访问连续内存中4字节数据的，一目了然。

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

合并式内存访问模式，显示Warp中的线程如何访问连续的128字节内存块，分为四个32字节的扇区。

图1. 合并式内存访问模式：Warp中的线程（箭头所示）访问连续的128字节内存块，分成四个32字节的扇区。

低效访问示例：散列式访问

反过来，如果线程们以很大的步长（stride）访问内存，那每次内存事务都会抓取远超实际需要的数据。每个线程只请求4字节数据，内存却要从全局内存中抓取一整个32字节的扇区，大部分数据都白白浪费了。图2就给大家展示了这种不理想的模式。

__global__ void uncoalesced_access(float* input, float* output, int n) {
  int tid = blockIdx.x * blockDim.x + threadIdx.x;
  if (tid < n) {
    // 以32（128字节）的步长访问，通过取模操作保持在边界内
    int scattered_index = (tid * 32) % n;
    output[tid] = input[scattered_index] * 2.0f;

散列式内存访问模式，显示每个线程（向下箭头）在单独的32字节内存扇区中访问数据。

图2. 散列式内存访问模式：每个线程（箭头所示）访问的数据都在一个单独的32字节内存扇区。

用NVIDIA Nsight Compute (NCU) “透视”内存访问

光说不练假把式！咱们现在就用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字节数据，并推测这可能是由于线程间的步长（stride）造成的。这完全符合咱们预设的“好坏内存性能”对比场景。

为了进一步深挖，咱们可以看看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_write.sum                            sector      8,240,680
dram__sectors_read.sum                             sector      67,110,368
--------------------------------------------------- ----------- ------------

对比这两段输出，咱们能看到两个核函数在 dram__sectors_read.sum （读取扇区总数）这个指标上有巨大的差异。咱们的核函数都是读取一个数组，然后写回同一个数组，所以理论上读取和写入的数据量应该差不多。但在 uncoalesced_access 这种非合并访问情况下，读取扇区数和写入扇区数之间竟然相差了足足8倍！这足以说明问题所在。

接着，咱们来分析一级缓存（L1 Cache）的行为。咱们可以用这条命令：

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字节数据块）和“请求”（Warp发起的内存事务）这两个概念，能为咱们揭示内存合并行为提供宝贵的洞察。扇区与请求的比率，清晰地反映了代码利用内存系统的效率。

如果咱们只关心这两个指标，可以用下面的命令单独收集它们：

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字节数据，所以一个Warp的每次请求，就需要抓取32个扇区。虽然内存系统抓取了32个扇区（共1024字节），但每个线程只需要它对应的扇区中的4字节。这种8倍的效率差异，对GPU性能有着深远的影响，因为内存带宽往往是限制GPU核函数最终性能的瓶颈。

步长访问的深远影响

现在，咱们再来看看步长（stride）对内存带宽到底有多大的影响。在CUDA内存访问模式中，步长指的是一个Warp中的线程访问连续内存位置之间的距离（以数组元素或字节为单位衡量）。像上面那两种核函数，当访问步长不同时，其带宽测量结果如图3所示。这里不是为了展示能达到的最大带宽，而是想直观地告诉大家，一个简单的核函数在全局内存访问步长发生变化时，它的带宽会如何变化。
GH200上步长从0到31的带宽与步长关系图，显示数值递减。

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

从图中可以清晰地看到，当步长较大时，有效带宽会变得很差，这完全符合咱们的预期。当Warp里的线程访问物理内存中相距很远的地址时，硬件就无法高效地合并这些访问。

多维数组访问：如何优化“矩阵操作”

接下来，咱们聊聊多维数组，也就是矩阵，在内存访问上的门道。想要达到最佳性能，实现合并式内存访问，关键在于让连续的线程访问数组中连续的元素，这和咱们处理一维数组的思路是一样的。

当咱们在CUDA核函数中使用2维或3维线程块时，线程的布局是线性的，其中X索引（threadIdx.x）变化最快，然后是Y（threadIdx.y），最后是Z（threadIdx.z）。举个例子，如果咱们有一个大小为(4,2)的2D线程块，那么线程的顺序将是：(0,0)(1,0)(2,0)(3,0), (0,1)(1,1)(2,1)(3,1)。在访问像矩阵这样的2D数据时，通常会使用2D线程块。

当咱们用2D线程块来访问一个矩阵（通常以1D数组的形式存储）时，考虑到C++是按行主序（row-major form）存储2D数据的，行访问就是连续的。如果我们能让连续的线程访问一行中连续的内存位置，那么这些访问就会非常高效（合并式）；反之，列访问则会效率低下（步长式、非合并式）。

新媒网跨境认为， 由于Warp内连续的threadIdx.x值应该访问连续的内存元素以实现合并，因此具有相同threadIdx.y值的线程应该访问矩阵的一行。这能确保Warp中的线程在访问矩阵元素时，遵循自然的行主序内存布局，从而实现高效的合并式内存事务，最大化内存带宽利用率。

下面这个合并式核函数（coalesced_matrix_access），就是通过巧妙地将线程索引映射到矩阵坐标，实现了高效的合并访问。在这里，每个块的X维度（threadIdx.x）被分配给列索引。这意味着当Warp中连续的线程增加它们的threadIdx.x时，它们访问的是矩阵中连续的列，但始终保持在同一行（如图4所示）。由于行主序将连续的内存位置存储为同一行内的元素，所以跨行访问能让Warp中的每个线程访问连续的内存位置。

__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;  // 行主序 ⇒ 合并访问
    matrix[idx] = matrix[idx] * 2.0f + 1.0f;
  }
}

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

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

而对于下面这个非合并式核函数（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;  // 列主序 ⇒ 非合并访问
    matrix[idx] = matrix[idx] * 2.0f + 1.0f;
  }
}

这里为了说明问题，核函数特意把行主序存储的矩阵，当成列主序来处理，使用了 col * height + row 的索引计算方式。这意味着当Warp中连续的线程增加它们的threadIdx.x（增加列索引）时，它们访问的是在列主序布局中连续的元素，但在行主序内存布局中却是步长式的。由于数据实际是按行主序物理存储的，却以列主序的索引方式访问，结果导致连续的线程访问的内存位置相距 height 个元素，从而形成了很大的步长模式，彻底打乱了GPU将这些访问合并成高效事务的能力（如图5所示）。这种存储顺序与访问模式的不匹配，会导致全局内存带宽利用率极低。
非合并式2D访问，显示2D线程块如何映射到2D矩阵，以及如何映射到矩阵所在的线性内存。连续的线程访问连续的列元素，这些元素在内存中是不连续的。

图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的“每请求扇区比”则暴露了其非合并式内存访问的问题，步长式的访问模式迫使内存子系统抓取远超必要的扇区数量，才能满足相同的内存请求。

总结与展望

朋友们，高效利用GPU内存，是咱们获得最佳性能的关键之一。要实现全局内存的最佳性能，合并式内存访问就是咱们的“杀手锏”。务必最小化全局内存的步长访问，并且养成好习惯，总是用NVIDIA Nsight Compute来对你的GPU核函数进行性能分析，确保内存访问是合并的。只有这样，咱们才能最大程度地压榨出GPU代码的性能潜力！

风险前瞻与时效提醒：

当前技术发展日新月异，咱们做跨境的，对这些前沿技术一定要保持敏感度。本教程所涉内容基于2025年的主流技术和工具（例如NCU），未来可能会有迭代更新。例如，GPU架构和CUDA版本升级都可能带来新的内存管理机制和优化策略。所以，大家在实际操作时，一定要留意官方（如NVIDIA）发布的最新文档和工具版本，确保咱们的知识库和实战技能始终保持最新，才能在这场激烈的“海外淘金”竞争中立于不败之地。

致谢：

本篇文章是在NVIDIA公司Mark Harris先生2013年原版文章基础上更新而来，在此表示感谢。

新媒网（公号: 新媒网跨境发布），是一个专业的跨境电商、游戏、支付、贸易和广告社区平台，为百万跨境人传递最新的海外淘金精准资讯情报。

本文来源：新媒网 https://nmedialink.com/posts/cuda-mem-coalescing-opt-8x-boost.html