在现代 GPU 计算中，内存访问模式往往是决定内核性能的首要因素。算法复杂度相同的两个 CUDA Kernel，仅仅因为访问显存的方式不同，执行时间可能相差数倍甚至一个数量级。本文将系统梳理 CUDA 编程中各类内存的访问特性，深入讲解合并访问（Coalesced Access）与 Bank Conflict 的核心原理，并给出配套的代码示例与优化实践。
一、Global Memory 合并访问
1.1 什么是合并访问
Global Memory 是 GPU 中容量最大但延迟最高的存储层级。现代 GPU 的 L2 Cache 以 32 字节为粒度传输数据，而一个 Warp（32 个线程）在一次内存事务中，若访问的地址恰好落在连续的 128 字节内（即满足对齐要求），则只需要一次或极少的传输请求即可完成。这种「相邻线程访问相邻地址」的模式被称为合并访问（Coalesced Access）。
相反，如果线程之间访问的地址彼此分散（Strided Access），每次只能捞出少量有效数据，导致总线利用率急剧下降，实际带宽可能仅为峰值的 10%-20%。
1.2 行主序矩阵的典型反例
下面这个 Kernel 展示了最常见的反模式：按行遍历却让 threadIdx.y 变化最快。
__global__ void uncoalesced_kernel(float* out, float* in,
int width, int height) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;
// 错误：threadIdx.y 在 x 轴上步进，导致 stride = width
if (x < width && y < height) {
out[idx] = in[idx] * 2.0f;
}
}
在这个例子中，同一个 Warp 内的线程分别跨越了 width 个 float 的距离。如果 width 为 1024，则 stride 达到 4KB，完全无法合并。
1.3 正确写法与性能对比
只要调整线程布局，让 threadIdx.x 对应内存的连续维度即可恢复合并访问。
#include <cuda_runtime.h>
#include <cuda.h>
#include <iostream>
#include <chrono>
// 反模式：stride = width
__global__ void uncoalesced_kernel(float* out, const float* in, int width, int height) {
int x = blockIdx.x * blockDim.x + threadIdx.x; // block 内连续，但网格 x 跨步大
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;
if (x < width && y < height)
out[idx] = in[idx] * 2.0f; // Warp 内 stride = width * sizeof(float)
}
// 正确模式：连续访问
__global__ void coalesced_kernel(float* out, const float* in, int width, int height) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;
if (x < width && y < height)
out[idx] = in[idx] * 2.0f; // Warp 内 x 连续，addr 连续
}
__global__ void coalesced_row_kernel(float* out, const float* in, int width, int height) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;
if (x < width && y < height)
out[idx] = in[idx] * 2.0f; // 只要 blockDim.x 对应连续维度即可
}
int main() {
const int width = 4096;
const int height = 4096;
const size_t memSize = sizeof(float) * width * height;
float *h_in = new float[width * height];
for (int i = 0; i < width * height; ++i) h_in[i] = static_cast<float>(i);
float *d_in, *d_out;
cudaMalloc(&d_in, memSize);
cudaMalloc(&d_out, memSize);
cudaMemcpy(d_in, h_in, memSize, cudaMemcpyHostToDevice);
dim3 threads(32, 8);
dim3 blocks((width + threads.x - 1) / threads.x,
(height + threads.y - 1) / threads.y);
// Warmup
coalesced_kernel<<<blocks, threads>>>(d_out, d_in, width, height);
cudaDeviceSynchronize();
// Benchmark coalesced
auto t1 = std::chrono::high_resolution_clock::now();
for (int i = 0; i < 100; ++i)
coalesced_kernel<<<blocks, threads>>>(d_out, d_in, width, height);
cudaDeviceSynchronize();
auto t2 = std::chrono::high_resolution_clock::now();
double ms = std::chrono::duration<double, std::milli>(t2 - t1).count();
std::cout << "Coalesced: " << ms << " ms\n";
cudaFree(d_in); cudaFree(d_out); delete[] h_in;
return 0;
}
实际在 Ampere 架构上测试，4096 x 4096 float 数据的合并访问通常能达到 400 GB/s 以上的有效带宽，而步长（stride）极大的模式可能只有 30-60 GB/s。核心原则就是让 threadIdx.x 映射到内存地址连续变化的最内层维度。
二、Shared Memory 与 Bank Conflict
2.1 Shared Memory 的 Bank 架构
Shared Memory 位于 SM（Streaming Multiprocessor）芯片上，延迟极低（约 20-30 周期），是寄存器与 Global Memory 之间的高速缓冲。在现代 GPU 中，Shared Memory 被逻辑划分为 32 个 Bank，每个 Bank 宽度为 4 字节。这意味着地址 addr 对应的 Bank ID 为 (addr / 4) % 32。当一个 Warp（32 线程）同时访问 32 个不同 Bank 的数据时，可以在一个周期内全部完成；如果多个线程命中了同一个 Bank，则会发生 Bank Conflict，这些访问被迫串行化。
2.2 无冲突、2-way 冲突与完全冲突
- 无冲突：线程
i恰好访问 Banki，一个周期完成 32 个读取。 - 2-way 冲突：两个线程访问同一个 Bank，需要 2 个周期。
- 32-way 冲突：所有线程访问同一个 Bank，需要 32 个周期，性能回到串行。
2.3 矩阵转置中的 Bank Conflict
矩阵转置是最典型的需要处理 Bank Conflict 的场景。在共享内存中写入转置后的数据时，如果不加处理，连续的线程会写入同一列，而同一列的元素恰好映射到同一个 Bank。
// 有 Bank Conflict 的转置：写入时 stride = 32 * 4 导致同一列命中同一 Bank
#define BDIM 32
__shared__ float tile[BDIM][BDIM];
__global__ void transpose_conflict(float* out, const float* in, int width, int height) {
int x = blockIdx.x * BDIM + threadIdx.x;
int y = blockIdx.y * BDIM + threadIdx.y;
int srcIdx = y * width + x;
int dstIdx = x * height + y;
tile[threadIdx.y][threadIdx.x] = in[srcIdx];
__syncthreads();
// Bank Conflict 发生在这里：
// threadIdx.x 连续变化时读取的是 tile[*][0], tile[*][1]...
// 但写入 Global 时并不冲突；真正的冲突是在 tile 的列访问中。
}
真正的问题出现在转置写入阶段：若按 tile[threadIdx.x][threadIdx.y] 读取，threadIdx.x 变化最快，而 threadIdx.x 的连续变化落到了列索引上，导致同一列（即同一 Bank）被多个线程争抢。
2.4 Padding 消除 Bank Conflict
最简单的做法是在 Shared Memory 数组的第二维增加一列 Padding，使得同一行的相邻元素错开 Bank。
#include <cuda_runtime.h>
#include <iostream>
#define BDIM 32
#define SROW (BDIM + 1) // Padding 后 Bank ID 不再对齐
// 无冲突版本
__global__ void transpose_pad(float* out, const float* in,
int width, int height) {
__shared__ float tile[BDIM][SROW];
int bx = blockIdx.x * BDIM;
int by = blockIdx.y * BDIM;
int x = bx + threadIdx.x;
int y = by + threadIdx.y;
// 读取原矩阵一行到 Shared Memory
tile[threadIdx.y][threadIdx.x] = in[y * width + x];
__syncthreads();
// 转置后写入：此时 x/y 在 tile 内的读取方向已翻转，
// 但由于 padding，相邻线程访问的 Bank 已错开
out[(bx + threadIdx.y) * height + (by + threadIdx.x)] =
tile[threadIdx.x][threadIdx.y];
}
// 有冲突版本（用于对比）
__global__ void transpose_npad(float* out, const float* in,
int width, int height) {
__shared__ float tile[BDIM][BDIM];
int bx = blockIdx.x * BDIM;
int by = blockIdx.y * BDIM;
int x = bx + threadIdx.x;
int y = by + threadIdx.y;
tile[threadIdx.y][threadIdx.x] = in[y * width + x];
__syncthreads();
// 32-way bank conflict
out[(bx + threadIdx.y) * height + (by + threadIdx.x)] =
tile[threadIdx.x][threadIdx.y];
}
int main() {
const int width = 1024, height = 1024;
const size_t size = sizeof(float) * width * height;
float *d_in, *d_out;
cudaMalloc(&d_in, size); cudaMalloc(&d_out, size);
dim3 blocks(width/BDIM, height/BDIM);
dim3 threads(BDIM, BDIM);
transpose_pad<<<blocks, threads>>>(d_out, d_in, width, height);
cudaDeviceSynchronize();
// 可用 nsys nvprof 对比两者带宽
cudaFree(d_in); cudaFree(d_out);
return 0;
}
通过 BDIM + 1 的 Padding，线程 i 与线程 i+1 访问的 Shared Memory 地址不再相差 32 的倍数，从而映射到不同 Bank，消除了 32-way Conflict。实测中，带 Padding 的转置通常比无 Padding 版本快 5-10 倍。
三、内存传输优化
3.1 页锁定内存（Pinned Memory）
默认情况下，cudaMemcpy 的 Host 端内存是可分页的（Pageable），数据必须先被 CUDA 驱动锁定到物理内存，再由 DMA 引擎传输到设备。这个「隐式复制」会使实测带宽远低于 PCIe 理论值。使用 cudaMallocHost 分配的页锁定内存（Pinned Memory）可直接作为 DMA 源/目标，传输效率显著提升。
float *h_pageable, *h_pinned, *d_data;
size_t bytes = 1ULL << 28; // 256 MB
// 页锁定内存分配
cudaMallocHost(&h_pinned, bytes); // 推荐
cudaMalloc(&d_data, bytes);
// 异步传输（需要 pinned host memory 才能 overlap）
cudaMemcpyAsync(d_data, h_pinned, bytes, cudaMemcpyHostToDevice, 0);
cudaStreamSynchronize(0);
cudaFreeHost(h_pinned);
cudaFree(d_data);
3.2 多流重叠 H2D、Kernel 与 D2H
CUDA Stream 允许将计算与通信任务排入不同的队列，只要硬件资源（copy engine、SM）充足，它们就可以并发执行。典型的双向流重叠模式如下：
#include <cuda_runtime.h>
#include <vector>
#define NSTREAM 4
#define CHUNK (1 << 26) // 64 MB per chunk
__global__ void apply_kernel(float* d, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) d[i] = d[i] * 2.0f + 1.0f;
}
int main() {
float *h_src, *h_dst, *d_src, *d_dst;
size_t total = NSTREAM * CHUNK * sizeof(float);
cudaMallocHost(&h_src, total);
cudaMallocHost(&h_dst, total);
cudaMalloc(&d_src, total);
cudaMalloc(&d_dst, total);
cudaStream_t streams[NSTREAM];
for (int i = 0; i < NSTREAM; ++i)
cudaStreamCreate(&streams[i]);
for (int i = 0; i < NSTREAM; ++i) {
float* h_chunk_src = h_src + i * CHUNK;
float* h_chunk_dst = h_dst + i * CHUNK;
float* d_chunk_src = d_src + i * CHUNK;
float* d_chunk_dst = d_dst + i * CHUNK;
// 1) H2D
cudaMemcpyAsync(d_chunk_src, h_chunk_src, CHUNK * sizeof(float),
cudaMemcpyHostToDevice, streams[i]);
// 2) Kernel
apply_kernel<<<(CHUNK + 255)/256, 256, 0, streams[i]>>>(d_chunk_src, CHUNK);
// 3) D2H
cudaMemcpyAsync(h_chunk_dst, d_chunk_src, CHUNK * sizeof(float),
cudaMemcpyDeviceToHost, streams[i]);
}
for (int i = 0; i < NSTREAM; ++i)
cudaStreamSynchronize(streams[i]);
for (int i = 0; i < NSTREAM; ++i)
cudaStreamDestroy(streams[i]);
cudaFreeHost(h_src); cudaFreeHost(h_dst);
cudaFree(d_src); cudaFree(d_dst);
return 0;
}
在一个支持双工的 PCIe 设备上，上述代码的有效吞吐量（H2D + Compute + D2H 的总耗时）往往比串行执行版本缩短 40%-60%。
3.3 Zero-Copy 内存
使用 cudaHostAllocMapped 标志分配的 Host 内存可以直接被 GPU 内核按指针访问，无需显式 cudaMemcpy。这种方式适合只访问一次的零星数据，但不适合大规模连续读写，因为每次访问都要走 PCIe 总线。
四、Unified Memory
4.1 自动页迁移
cudaMallocManaged 创建的统一内存（Unified Memory）允许 CPU 与 GPU 共享同一个指针。在 Pascal（SM 6.x）及更新的架构上，系统以 4KB 页为粒度自动在 Host 与 Device 之间迁移数据。如果数据被 GPU 访问后又被 CPU 访问，就会发生 Page Fault，驱动将页迁移回 Host 内存。
4.2 显式预取与访问提示
为了规避隐式迁移的开销，CUDA 提供了显式控制 API：
#include <cuda_runtime.h>
int main() {
const int n = 1 << 24;
float *data;
cudaMallocManaged(&data, n * sizeof(float));
// 初始在 CPU 初始化
for (int i = 0; i < n; ++i)
data[i] = static_cast<float>(i);
int device = 0;
// 显式预取到 GPU（避免内核调用时触发 page fault）
cudaMemPrefetchAsync(data, n * sizeof(float), device, 0);
cudaDeviceSynchronize();
// 内核执行 ...
// apply_kernel<<<(n+255)/256, 256>>>(data, n);
// 预取回 CPU
cudaMemPrefetchAsync(data, n * sizeof(float), cudaCpuDeviceId, 0);
cudaDeviceSynchronize();
// 访问提示：告诉驱动这段数据将主要被 GPU 只读访问
cudaMemAdvise(data, n * sizeof(float), cudaMemAdviseSetReadMostly, device);
// 对应还有个 cudaMemAdviseSetPreferredLocation 可指定常驻设备
cudaFree(data);
return 0;
}
4.3 适用场景
Unified Memory 在以下场景能显著降低代码复杂度：
- 数据结构包含深层指针（如图、树），手动传输需要遍历打包。
- 算法的实际数据访问模式在运行时才能确定，静态分配难以预判数据温度。
- 需要快速原型验证，先跑通正确性，再针对热点路径做显式管理优化。
在性能敏感的大规模 HPC 场景中，仍然推荐使用显式 cudaMalloc/cudaMemcpyAsync。
五、Constant Memory 与 Texture Memory
5.1 Constant Memory
Constant Memory 是片上只读缓存，容量通常为 64KB。当一个 Warp 内的所有线程访问同一个地址时，Constant Memory 可实现高效的广播；若访问地址分散，则性能急剧下降。
__constant__ float const_filter[128]; // 设备端只读
__global__ void conv_kernel(float* out, const float* in, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
float sum = 0;
for (int f = 0; f < 128; ++f)
sum += in[i + f] * const_filter[f]; // broadcast within warp
out[i] = sum;
}
}
5.2 Texture Memory
Texture Memory 同样通过只读缓存访问，但它具有硬件实现的线性插值、归一化坐标以及 2D 局部性优化。对于图像处理中的 2D 采样、查找表（LUT）等需要空间局部性的随机访问场景，Texture Cache 往往比 L1/L2 Cache 更高效。
texture<float, cudaTextureType2D, cudaReadModeElementType> texRef;
__global__ void sample_kernel(float* out, int w, int h) {
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < w && y < h)
out[y * w + x] = tex2D(texRef, x + 0.5f, y + 0.5f);
}
// Host 端绑定
void bind_texture(cudaArray* cuArray) {
texRef.normalized = false;
texRef.filterMode = cudaFilterModePoint;
texRef.addressMode[0] = cudaAddressModeClamp;
texRef.addressMode[1] = cudaAddressModeClamp;
cudaBindTextureToArray(texRef, cuArray);
}
六、Nsight 性能分析
6.1 Nsight Compute 关键指标
Nsight Compute（ncu）是 CUDA 最细粒度的分析工具。在内存优化迭代中，关注以下几个指标即可快速定位瓶颈：
- Memory Throughput (%)：Global Memory 利用率，若接近 100% 说明是显存带宽受限，需从访问模式入手。
- L2 Cache Hit Rate：若命中率低于 60%，说明数据复用不足，可考虑增加 Shared Memory 缓存或调整 tile 大小。
- Shared Memory Bank Conflicts：直接给出每执行一次的 Conflict 数量，非零即需审视 Shared Memory 索引计算。
- Memory Divergence：Warp 内事务数量超过理论最小值时，存在地址分散。
6.2 快速诊断命令
ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,\
smsp__warps_issue_stalled_no_instruction,\
l1tex__t_requests_pipe_lsu_mem_global_op_ld \
./a.out
通过比较合并前后的 t_sectors 与 t_requests 比值，可以量化合并访问效率。比值越接近 4（128B / 32 threads），说明合并越充分。
总结
CUDA 内存优化的核心始终围绕「数据在哪里」与「怎么访问」两个维度展开：
- Global Memory 务必保证 Warp 内地址连续，让
threadIdx.x映射最内层连续维度。 - Shared Memory 利用 Padding 技术消除 Bank Conflict，尤其在矩阵转置、FFT、卷积等场景中这是必做项。
- Host↔Device 传输 使用 Pinned Memory 配合多 Stream 重叠，必要时引入 Unified Memory 降低复杂数据结构的传输成本。
- 只读常量/纹理数据 根据访问模式选择 Constant Memory（广播）或 Texture Memory（2D 局部性）。
- 量化迭代 不要凭感觉优化，始终用 Nsight Compute 的内存吞吐量与 Bank Conflict 指标验证每次改动的实际效果。
将上述原则落实到代码中，能够在不改变算法复杂度的情况下，获得数倍乃至数量级的加速收益。
继续阅读
探索更多技术文章
浏览归档，发现更多关于系统设计、工具链和工程实践的内容。