内存访问模式

目录

前言

存储时的对齐

对齐与合并访问

全局内存读取

全局内存写入

结构体数组和数组结构体

总结

前言

本文将通过线程束的视角了解数据如何流动，以更加基础的视角去优化性能，如何通过优化全局内存的访问模式来榨干显存带宽

存储时的对齐

int *d_data = nullptr; cudaMalloc((void**)&d_data, sizeof(int) * n); 也就是说，返回的地址一定满足： address % 256 == 0 比如返回的地址可能是： 0x700000000 （十六进制，末尾两位都是0，说明至少256字节对齐） 0x700000100 0x700000200 而不会是： 0x700000001 （这种地址cudaMalloc不会返回） 0x700000003

256字节对齐是最低保证，实际上很多情况下是512字节甚至更大粒度的对齐，取决于驱动实现。

为什么是256字节？不是其他的字节呢？？？

GDDR6（GraphicsDoubleDataRate6）是第六代图形专用双倍数据速率同步动态随机存储器。显存位宽是128位，也就是拥有128条满足GDDR6的数据线，显存控制器是协调这 128 条线按照 GDDR6 协议工作的硬件逻辑。

突发传输（Burst Transfer）就是显存控制器在一次操作中，连续多个时钟周期不间断地搬数据，把多个 128 字节的 Cache Line“一口气”搬完，而不是一个 Cache Line 一个 Cache Line 地分开搬。

因为显存控制器启动是有开销的，不如一起开销，多搬运几个cacheLine，所以可能一次搬运256字节，512字节等等，所以256是可以让你不担心一次突发会横跨两个物理块，保证一次突发的情况下这些同属于一个物理块

256字节满足所有小字节类型的对齐要求

cudaMalloc返回的地址，比如：0x700000000 0x700000000 % 256 == 0 ✓（cudaMalloc保证） 0x700000000 % 128 == 0 ✓（顺带满足） 0x700000000 % 64 == 0 ✓（顺带满足） 0x700000000 % 32 == 0 ✓（顺带满足，sector对齐） 0x700000000 % 16 == 0 ✓（顺带满足，int4对齐） 0x700000000 % 8 == 0 ✓（顺带满足，int2对齐） 0x700000000 % 4 == 0 ✓（顺带满足，int对齐）

对齐与合并访问

┌─────────────────────────────────────────────────────────────────┐ │ 显卡 (Graphics Card) │ │ ┌───────────────────────────────────────────────────────────┐ │ │ │ GPU 核心芯片 (片上 SRAM) │ │ │ │ ┌─────────────────────────────────────────────────────┐ │ │ │ │ │ SM (流多处理器) × 24 │ │ │ │ │ │ ┌─────────────────┐ ┌─────────────────┐ │ │ │ │ │ │ │ 寄存器文件 │ │ 共享内存 / L1 │ │ │ │ │ │ │ │ (256 KB/SM) │ │ (约 128 KB/SM) │ │ │ │ │ │ │ │ ● 线程私有 │ │ ● Block内共享 │ │ │ │ │ │ │ │ ● 速度最快 │ │ ● 速度极快 │ │ │ │ │ │ │ └─────────────────┘ └─────────────────┘ │ │ │ │ │ │ ┌─────────────────────────────────────────────┐ │ │ │ │ │ │ │ CUDA 核心 (128/SM) │ Tensor Core (4/SM) │ │ │ │ │ │ │ │ Warp Scheduler ×4 │ SFU, LD/ST 等 │ │ │ │ │ │ │ └─────────────────────────────────────────────┘ │ │ │ │ │ └─────────────────────────────────────────────────────┘ │ │ │ │ │ │ │ │ │ ┌───────────────────────┴─────────────────────────────┐ │ │ │ │ │ L2 缓存 (24 MB, 所有 SM 共享) │ │ │ │ │ └───────────────────────┬─────────────────────────────┘ │ │ │ │ │ │ │ │ │ ┌───────────────────────┴─────────────────────────────┐ │ │ │ │ │ 显存控制器 (Memory Controller) │ │ │ │ │ │ 与 L2 通过内部高速总线连接 │ │ │ │ │ └───────────────────────┬─────────────────────────────┘ │ │ │ └──────────────────────────┼────────────────────────────────┘ │ │ │ │ │ ┌──────────────────┼──────────────────┐ │ │ │ 128 位物理数据线 (位宽 = 128 bit) │ │ │ └──────────────────┼──────────────────┘ │ │ │ │ │ ┌──────────────────────────┴──────────────────────────────┐ │ │ │ 显存颗粒 (片外 DRAM) │ │ │ │ 总容量：8 GB GDDR6 │ │ │ │ ┌─────────────────────────────────────────────────────┐ │ │ │ │ │ 全局内存 (Global Memory) │ │ │ │ │ │ - 通过 cudaMalloc 分配 │ │ │ │ │ │ - 所有线程/SM 可读写 │ │ │ │ │ │ - 包含本地内存、常量内存数据区、纹理内存数据区 │ │ │ │ │ └─────────────────────────────────────────────────────┘ │ │ │ └──────────────────────────────────────────────────────────┘ │ │ │ │ ┌──────────────────────────────────────────────────────────┐ │ │ │ PCIe 接口 (连接 CPU) │ │ │ └──────────────────────────────────────────────────────────┘ │ └─────────────────────────────────────────────────────────────────┘

SM 内部 CUDA 核心 │ ▼ 加载/存储单元 (LSU) ──── 缓存控制器逻辑 ──── L1 缓存 (片上 SRAM) │ ▼ L2 缓存 (片上 SRAM) │ ▼ 显存控制器 (Memory Controller) │ ▼ 显存颗粒 (片外 DRAM)

内存事务是显存控制器为满足一次或多次内存访问请求而执行的一次最小数据搬运。每次事务以固定大小的 Cache Line（通常为 128 字节）为单位，从显存（DRAM）读取或写入数据到 L2 缓存，或从 L2 缓存传输到 L1 缓存。

概念	英文	大小	操作对象	作用层级
内存事务	Memory Transaction	128 字节	DRAM ↔ L2 缓存	显存控制器级别
缓存扇区	Cache Sector	32 字节	L1 ↔ L2 缓存内部	缓存控制器级别

一个 128 字节的 Cache Line 由4 个 Sector组成，每个 Sector 32 字节。Sector 是 L1/L2 缓存的内部最小访问单位，而 Transaction 是跨层级搬运的最小单位。（一个是内部管理数据的维度，一个是横跨不同层级的搬运维度，要注意区分），这样管理内部4个sector，细粒度更细，如果其中一个sector dirty了，但是别的进来时读取别的sector还能读取，不需要重新从Dram读取，而如果L1L2内部也用transaction的话如果只是一点dirty，那其他进来读取必然也dirty了，所以每次都要从Dram重新读

注意有L1缓存控制器，是接收LSU发出来的指令看数据是不是在L1缓存当中，L2也有缓存控制器，是管理L2的，接收L1缓存控制器的命令

一、加载的完整流程（从 DRAM 到寄存器）

假设一个 Warp 的 32 个线程要读取全局内存中的 32 个float（128 字节）：

1. Warp Scheduler 发射指令：Warp Scheduler 发射一条LDG（Load Global）指令给LSU。

2. LSU 合并地址：LSU 收集这 32 个线程的地址，发现它们落在同一个 128 字节 Cache Line 内，合并成一次请求。

3. 查 L1 缓存：LSU 向L1 缓存控制器发送请求。L1 控制器检查请求的 Cache Line 是否在 L1 里：

   • 命中→ 直接从 L1 返回数据给 LSU，跳到第 6 步。

   • 缺失→ 继续第 4 步。

4. 查 L2 缓存：L1 控制器把请求转发给L2 缓存控制器。L2 控制器检查：

   • 命中→ L2 把整个 128 字节 Cache Line 返回给 L1，L1 缓存后返回数据给 LSU，跳到第 6 步。

   • 缺失→ 继续第 5 步。

5. 查 DRAM：L2 控制器把请求发给显存控制器。显存控制器向GDDR6 显存颗粒发起一次 128 字节的物理事务，搬来整个 Cache Line。数据原路返回：DRAM → 显存控制器 → L2 → L1 → LSU。

6. LSU 分发数据：LSU 拿到数据块，按每个线程的原始请求切分，把正确的 4 字节写入每个线程的目标寄存器。指令完成。

二、存储的完整流程（从寄存器到 DRAM）

假设同一个 Warp 要把计算结果写回全局内存：

1. Warp Scheduler 发射指令：Warp Scheduler 发射一条STG（Store Global）指令给LSU。

2. LSU 合并地址：LSU 收集 32 个线程的目标地址，如果连续就合并成一次写请求。

3. 写 L1 缓存：LSU 向L1 缓存控制器发送写请求。

   • L1 写回策略：数据先写入L1 缓存，并标记为“脏”（Dirty），不立即写到 L2 或 DRAM。只有这个脏的 Cache Line 被替换时，才一次性写回 L2。

   • L1 写穿透策略：某些 GPU 架构支持此模式，数据同时写入 L1 和 L2，确保 L2 立刻看到更新。

4. 写 L2 缓存：当 L1 需要替换掉一个脏 Cache Line 时，数据被写回L2 缓存。L2 也采用写回策略——数据先在 L2 里，标记为脏，只有被替换时才写回 DRAM。

5. 写 DRAM：当 L2 需要替换掉一个脏 Cache Line 时，数据才被写回DRAM。显存控制器向 GDDR6 显存颗粒发起一次 128 字节的物理写事务。

对于优化内存，我们主要关心以下两个特性

• 对齐内存访问：回答的是“数据的起始地址在哪里”的问题。它要求数据的首地址是对齐到某个固定值（如 128 字节）的。
• 合并内存访问：回答的是“同一个 Warp 的 32 个线程，它们的地址是否连续”的问题。它要求这些地址落在同一个 128 字节的对齐区域内。

为了最大化全局内存带宽，你写的核函数必须让同一个 Warp 的线程产生“对齐合并访问”。

两个必须同时满足的条件：

1. 对齐：这次合并访问的起始地址必须是 128 字节的整数倍。

2. 合并：同一个 Warp 的 32 个线程要访问的数据，全部落在同一个 128 字节 Cache Line 里。

假设arr[0]地址是0x100。

• 对齐合并（最优）：32 个线程访问arr[0]到arr[31]。起始地址0x100是 128 的整数倍，全部 128 字节落在同一个 Cache Line 里。一次事务，满载而归。

• 不对齐但连续（需两次事务）：32 个线程访问arr[2]到arr[33]。起始地址0x108不是 128 的整数倍，横跨了两个 Cache Line（0x100和0x180）。即使地址连续，也需要两次事务。

• 对齐但不连续（无法完美合并）：32 个线程随机或跨步访问同一个 Cache Line 内的地址。大部分请求仍需多次事务。

以上的观点就是想要说明一个问题：我们一次搬运是128字节，我们如果想要带宽打满，那就需要把128字节全部用来搬运有效数据，遗憾的是我们不能随便指定这128字节，我们需要这128字节是连续的并且需要对齐，这样一次性就能够打满带宽，提升效率，所以之前为什么要在x维度上面满足32的整数倍，就是因为一个warp是32线程为单位，这32线程是连续的，访问的地址是连续的话效率会非常高

全局内存读取

注意这里讲的是读取，也就是load，而不是store

以下是load的路径

路径	存储类型	缓存路径	程序员如何触发	性能特性
① 全局内存路径	全局内存 (cudaMalloc)	L1 → L2 → DRAM	默认 Load（无特殊修饰符）	需要合并访问，128 字节事务
② 共享内存路径	共享内存 (__shared__)	无缓存，直接访问	显式用__shared__声明	极快（SRAM），需注意Bank Conflict
③ 常量内存路径	常量内存 (__constant__)	只读常量缓存	显式用__constant__声明	适合 Warp 内统一地址广播，串行化惩罚大（联想重力加速度G常量）
④ 纹理内存路径	全局内存（通过纹理对象绑定）	只读纹理缓存	用cudaTextureObject_t绑定	适合二维空间局部性，支持硬件插值（这个暂不了解，先放）
⑤ 局部内存路径	局部内存（线程私有溢出区）	L1 → L2 → DRAM	编译器自动分配（寄存器溢出/大数组）	物理在 DRAM，访问慢，应尽量避免

关于一些比较旧的文章，讲解旧的架构的时候，可能会说L1缓存可以关闭，那是因为以前的L1缓存非常小，16~48KB，warp访问的时候导致cache thrashing严重，针对流式的数据，也就是只访问一次的数据，没必要加入缓存当中，关闭L1缓存，反而效率会更高
但是针对现代架构，对这个概念反而弱化了，针对以前关闭L1缓存的相关操作也逐渐弱化删除，现代架构编译器会根据架构特性做自己的判断，我们能做的就是调整共享内存和L1缓存的比例（关于L1缓存在对齐与合并那里已经讲的非常详细了，这里不过多赘述）

#include <cuda_runtime.h> #include <stdio.h> #include <string> #include "../freshman.hpp" using namespace std; void sumArrays(float * a,float * b,float * res,int offset,const int size){ for(int i=0,k=offset;k<size;i++,k++){ res[i]=a[k]+b[k]; } } __global__ void sumArraysGPU(float*a,float*b,float*res,int offset,int n){ int i=blockIdx.x*blockDim.x+threadIdx.x; int k=i+offset; if(k<n) res[i]=a[k]+b[k]; } int main(int argc, char* argv[]){ int dev = 0; cudaSetDevice(dev); int nElem=1<<18; int offset=0; if(argc>=2){ offset=stoi(argv[1]); } printf("Vector size:%d\n",nElem); int nByte=sizeof(float)*nElem; float *a_h=(float*)malloc(nByte); float *b_h=(float*)malloc(nByte); float *res_h=(float*)malloc(nByte); float *res_from_gpu_h=(float*)malloc(nByte); memset(res_h,0,nByte); memset(res_from_gpu_h,0,nByte); float *a_d,*b_d,*res_d; CHECK(cudaMalloc((float**)&a_d,nByte)); CHECK(cudaMalloc((float**)&b_d,nByte)); CHECK(cudaMalloc((float**)&res_d,nByte)); CHECK(cudaMemset(res_d,0,nByte)); initialData(a_h,nElem); initialData(b_h,nElem); CHECK(cudaMemcpy(a_d,a_h,nByte,cudaMemcpyHostToDevice)); CHECK(cudaMemcpy(b_d,b_h,nByte,cudaMemcpyHostToDevice)); dim3 block(1024); dim3 grid(nElem/block.x); double iStart,iElaps; iStart=efficiency(); sumArraysGPU<<<grid,block>>>(a_d,b_d,res_d,offset,nElem); cudaDeviceSynchronize(); iElaps=efficiency()-iStart; CHECK(cudaMemcpy(res_from_gpu_h,res_d,nByte,cudaMemcpyDeviceToHost)); printf("Execution configuration<<<%d,%d>>> Time elapsed %f sec --offset:%d \n",grid.x,block.x,iElaps,offset); sumArrays(a_h,b_h,res_h,offset,nElem); if(check(res_h,res_from_gpu_h,nElem)!=0){ cout<<"正确"<<endl; } cudaFree(a_d); cudaFree(b_d); cudaFree(res_d); free(a_h); free(b_h); free(res_h); free(res_from_gpu_h); return 0; }

nvcc -O3 -arch=sm_89 -Xptxas -dlcm=ca -o main main.cu #开启L1 nvcc -O3 -arch=sm_89 -Xptxas -dlcm=cg -o main main.cu #关闭L1

大家可以试一下是否还可以关闭和开启

开启的

关闭的

可以看出来关闭L1会导致所有的数据都不经过L1缓存了
但是有趣的是，如果你的offset=0的时候，因为数据是流式的访问，也就是只访问一次，硬件识别到了，所以无论你开启和关闭L1，他都不经过L1，直接经过L2

所以现代编译器的作用还是挺大的，很多理论的分析，实验出来都可能不一样，所以我们需要结合实际去分析

以下是实验分析：

实验数据是1<<18个元素，每个元素是float，4个字节，然后a数组和b数组，所以是两份

所以总的read数据量=1<<18*4*2=2MB

为此我们需要监测显存的read.sum和L1缓存的sector

dram_bytes_read.sum l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio

offset	l1tex__average_t_sectors_per_request...ratio	是否对齐？	解读
0	4.00	✅ 对齐	32 线程的连续访问恰好装满一个 128 字节 Cache Line（4 个 Sector），一次事务满载
1	5.00	❌ 不对齐	128 字节数据横跨两个 Cache Line，多跨越一个边界，需要 5 个 Sector 才能覆盖
2	5.00	❌ 不对齐	同上，offset=2 同样横跨两个 Cache Line
3	5.00	❌ 不对齐	同上，offset=3 同样横跨两个 Cache Line

结论：注意代码里面的访问都是连续的，为什么会出现offset=0的时候会是4，原因是他们是对齐且连续的，为什么offset=1是5，甚至大家可以试一试offset=7也是5，因为不对齐但是连续，注意到了offset=8是4，因为对齐且连续

---

dram__bytes_read.sum：为什么 offset=0 反而更大？

offset	dram__bytes_read.sum	解读
0	2.40 MB	对齐，但 DRAM 搬运量反而最大
1	2.10 MB	不对齐，DRAM 搬运量反而更小
2	2.10 MB	同上
3	2.11 MB	同上

这个结果初看反直觉——为什么完美对齐的 offset=0 反而比不对齐的 offset=1,2,3 搬运了更多 DRAM 数据？

有一个写分配机制，也就是SM写数据到Dram时，如果L1缓存没有，就会去L2缓存，L2没有，就会去dram读取缓存行，然后依次同步到L2L1，接着再写入L1然后标记为dirty，表示缓存的数据新于dram，后续L1缓存淘汰再同步到L2

原因：写分配

1. 读 a[k] → load，产生DRAM read 2. 读 b[k] → load，产生DRAM read 3. 写 res[i] → store，但res[i]不在缓存里 → 触发写分配：先从DRAM把res[i]所在的cache line读进L2 → 修改L2里的值，标记dirty → 这次"为了写而产生的读"也计入dram_bytes_read！

为什么数据不应该是3MB吗？因为之前采用cudaMemset的时候，数据已经被预存在缓存当中了，甚至cudaMemcpy会先经过L2缓存再到显存，所以有部分数据也会提前被缓存，所以我们可以实验一下关掉cudaMemset，看一下是否会增加到3MB

PCIe接口 → Copy Engine → L2缓存(LTS) → 显存控制器 → GDDR6

好的，写到这里博主已经很蒙蔽了，上面的是昨天的笔记，但是这里是今天，今天去测试来看，无论是注释掉还是不注释cudaMemset，基本都是2.10MB左右了，所以昨天的2.4MB是一个异常情况，是因为GPU运行时博主开了过多的后台渲染，导致L2缓存不断的刷新

简单来说：我们要清楚两个前提，一个是写分配，一个是cudaMemset已经提前帮我们缓存到L2了，还有cudaMemcpy也会经过L2被缓存一些，所以我们只要关心最终的即可，不要再去详细的分析为什么不是3为什么不是2等等

写到这里，博主想要说明的是：对齐合并能够保证每次128字节事务都装满有效数据，我们的每一次请求都会有4个sector，而不是5个6个等等

1. 线程块 x 维度设为 32 的倍数：这是保证 Warp 内线程 ID 连续、访问地址连续的最简单方法。

2. 保证起始地址对齐

3. 每个线程处理连续地址

4. 避免跨步访问：不要用a[i * stride]这类大跨步访问。

仅仅改变一下访问模式，我们就能获得一些优化

nvidia-ada-gpu-architecture.pdf

这份是关于ada架构的白皮书，所有有关的如果觉得博主讲解有问题的，可以自行查阅相关资料，博主也是刚入门学习

• Kepler 等老架构：纹理缓存有独立于 L1 的专用 SRAM，属于 TPC 层级的较大存储，顶层架构图可以画出独立模块；
• Maxwell 及之后（包含 Ada）：NVIDIA 把纹理数据存储收拢进 SM 内部的 L1TEX，仅保留 TEX 运算单元在 SM 里，存储层面和 L1 合并，顶层芯片布局就不再单独画出纹理缓存区块。

注意纹理缓存+只读缓存+L1缓存是共用同一块SRAM的，然后再加上共享内存，也就是相当于我们之前讲的L1缓存内部再进行细分，普通L1缓存+纹理缓存+只读缓存，这三者的处理逻辑不一样但是SRAM是同一块

┌─────────────────────────────────────────┐ │ SM │ │ │ │ ┌──────────────────────────────────┐ │ │ │ L1TEX SRAM(128KB) │ │ │ │ 可划分：通用L1 / 共享内存 │ │ │ │ 内部复用存储：只读通路、纹理数据 │ │ │ │ 多条逻辑通路共用这一块物理内存 │ │ │ └──────────────────────────────────┘ │ │ │ │ ┌──────────────────────────────────┐ │ │ │ 常量缓存（独立SRAM，64KB） │ │ │ │ 专属广播硬件，不占用L1TEX空间 │ │ │ └──────────────────────────────────┘ │ │ │ └─────────────────────────────────────────┘

Ada 架构没有独立物理 SRAM 只读缓存。 「只读缓存」本质 =GPU LSU 一条独立的只读访存通路，数据存放于l1tex（128KB 统一 SRAM）； 优势：不受-dlcm=cg影响。 当你使用-dlcm=cg让普通ptr[i]全局 load 强制绕过 L1TEX 时，__ldg()依然可以使用 L1TEX 缓存。

只读缓存（Read-only Cache）

对应的内存：global memory里的只读数据 硬件特性： 没有广播机制（这点和常量缓存不同） 但对2D空间局部性有优化（相邻地址的缓存效率更好） 容量比常量缓存大 适合的访问模式： 每个线程读不同的地址（常量缓存做不到的） 数据在整个kernel执行期间不会被修改 使用方式： // 方式一：__ldg()内置函数 int val = __ldg(&g_data[idx]); // 方式二：const __restrict__指针 __global__ void kernel(const float* __restrict__ data){ float val = data[idx]; // 编译器自动走只读缓存路径 }

常量缓存（Constant Cache）

对应的内存：__constant__声明的常量内存（64KB） 硬件特性： 有广播机制（broadcast） 同一warp内所有线程读同一地址 → 1次读取广播给32个线程，极高效 同一warp内线程读不同地址 → 串行化，极慢（32次串行读取） 适合的访问模式： 所有线程读同一个值（比如神经网络的bias） 不适合每个线程读不同地址的情况 声明方式： __constant__ float weights[256]; cudaMemcpyToSymbol(weights, h_weights, sizeof(float)*256);

纹理缓存（Texture Cache）

现代架构（Maxwell之后）的真相： 纹理缓存在物理上已经和只读缓存合并了 两者共享同一块硬件单元 在ncu里看到的"l1tex"里的"tex"就是指这个 历史上纹理缓存的特点（现在依然保留的特性）： 针对2D空间局部性优化 支持硬件插值（bilinear interpolation） 支持边界处理（clamp/wrap/mirror） 支持坐标归一化 现代用法： 图形渲染场景：用传统的texture object API 通用计算场景：直接用__ldg()或const __restrict__，走的是同一条硬件路径

全局内存写入

这里我们要讲解的是store，注意对于LSU来说，这两条通道是独立的，也就是你完全可以load的同时也store，除非他们之间有逻辑关系，比如先写后读，那就只能等

寄存器 → LSU（合并地址）→ L1 缓存 → L2 缓存 → 显存控制器 → DRAM

1. Warp Scheduler 发射指令：发射一条STG（Store Global）指令给 LSU。

2. LSU 合并地址：收集 32 个线程的目标地址。如果地址连续且对齐，合并成一次大写入请求；如果分散，拆成多次小请求。

3. 写入 L1 缓存：数据首先写入 L1 缓存，并标记为“脏”（Dirty）。此时数据还没有到达 DRAM。

4. 写回 L2：当 L1 中这个脏的 Cache Line 被替换时，数据被写回 L2 缓存。

5. 写回 DRAM：当 L2 中这个脏的 Cache Line 被替换时，数据才最终写回 DRAM。

核心缓存策略：写回 + 写分配

（1）写回策略

写入操作不会立即穿透到 DRAM。数据在 L1/L2 中暂留，直到该 Cache Line 被替换时才一次性写回。这是为了减少 DRAM 带宽消耗——如果同一地址被多次写入，只有最后一次写回真正生效。

（2）写分配策略

当写入目标地址不在 L1/L2 缓存中时，硬件不会直接写 DRAM，而是先把目标地址所在的整个 128 字节 Cache Line从 DRAM 读到 L1/L2，然后在缓存中修改，标记为脏。这个“为写入而读”的操作就是写分配。

为什么需要写分配？因为写入通常只修改 Cache Line 中的部分字节（比如 4 字节），而缓存一致性管理的最小粒度是 128 字节的 Cache Line。硬件必须先拥有完整的旧版本，才能正确标记哪些部分被修改。

（3）例外：流写入

对于只写不读的数据（比如输出结果），可以用流写入策略绕过写分配。通过在编译器选项或 PTX 指令中指定，LSU 会直接发起 128 字节的写事务到 DRAM，不经过 L1/L2 缓存，避免“为写入而读取”的无用开销。

合并写入同样重要
设想一下，如果写入的地址非常分散，LSU 无法合并，必须拆成多次小写入请求。每次小写入都触发独立的缓存操作和 DRAM 事务，每次 128 字节事务只有少量有效数据，带宽利用率断崖式下跌。对于写分配不一定会增加很多，因为你在写入之前可能已经把数据加载进来过了，所以取决于你之前的操作，甚至还有跟预取器有关（这个的意思就是比如你在读取a的时候，可能因为a和res比较连续，预取器由于连续会把周围的也读取进缓存），所以跟很多种因素有关，反正就是合并肯定好

结构体数组和数组结构体

这对于学过c语言的我们并不奇怪，在c语言的结构体还有对齐的规则，如果想要了解的自行查一下资料

• 结构体数组 (AoS - Array of Structures)：先定义结构体，再创建一个由该结构体组成的数组。内存排布是交替的，比如 struct1.x, struct1.y, struct1.z, struct2.x, struct2.y, struct2.z...

• 数组结构体 (SoA - Structure of Arrays)：先定义一个包含多个数组的结构体，每个数组成员存储所有元素的同一属性。内存排布是连续的，比如所有元素的 x 值连续存放，所有 y 值连续存放，所有 z 值连续存放。

// 1. 结构体数组 (AoS) struct PointAOS { float x, y, z; }; PointAOS aos_points[1024]; // 内存：[x0,y0,z0], [x1,y1,z1]... // 2. 数组结构体 (SoA) struct PointsSOA { float x[1024]; float y[1024]; float z[1024]; }; PointsSOA soa_points; // 内存：[x0,x1,x2...], [y0,y1,y2...]...

可以看出来这两者的内存结构相差很大，假设我们的thread同时访问x，左边第一种AoS那回到上面的对齐合并中offset绝对就大于0了（这取决于你的结构体是xyz还是xy还是别的），那这样是无法完美合并，那访问的效率绝对会大大折扣

• 全局内存访问与合并访问：GPU 希望同一个 Warp 的 32 个线程能连续访问地址，一次搬完 128 字节。如果只访问x属性，SoA 方式会让所有线程的访问连续，一次读满整个缓存行，完美合并访问。而 AoS 方式下，每个线程的x之间隔着y和z，地址不连续，一个缓存行只利用了三分之一，严重浪费带宽。

• 共享内存访问与 Bank Conflict：共享内存的 32 个 Bank 喜欢线程访问落在不同 Bank 上。AoS 方式容易让相邻线程踩进同一个 Bank，导致请求排队（Bank Conflict）。SoA 方式则让线程连续访问完美分布在不同的 Bank 上，实现无冲突高效访问。

如果你的数据是AoS的话，可以读到共享内存，然后转到SoA之后，在进行后续计算会更高效

实验

#include "../freshman.hpp" #include <cuda_runtime.h> #include <iostream> using namespace std; #define N (1<<24) // 提前定义 N // 结构体数组 (AoS) struct AoS { float x; float y; }; // 数组结构体 (SoA) struct SoA { float x[N]; float y[N]; }; void checkResult_structAoS(float* res_h, struct AoS* res_from_gpu_h, int nElem) { for (int i = 0; i < nElem; i++) if (res_h[i] != res_from_gpu_h[i].x) { printf("check fail at %d!\n", i); exit(0); } printf("result check success!\n"); } void checkResult_structSoA(float* res_h, struct SoA* res_from_gpu_h, int nElem) { for (int i = 0; i < nElem; i++) if (res_h[i] != res_from_gpu_h->x[i]) { printf("check fail at %d!\n", i); exit(0); } printf("result check success!\n"); } void sumCpu(float* res, float* a, float* b, int n) { for (int i = 0; i < n; i++) res[i] = a[i] + b[i]; } __global__ void AoSGpu(struct AoS* res, float* a, float* b, int n) { int i = threadIdx.x + blockIdx.x * blockDim.x; if (i < n) { res[i].x = a[i] + b[i]; } } __global__ void SoAGpu(struct SoA* res, float* a, float* b, int n) { int i = threadIdx.x + blockIdx.x * blockDim.x; if (i < n) { res->x[i] = a[i] + b[i]; } } int main() { int device = 0; cudaSetDevice(device); int nElem = N; int nSize = nElem * sizeof(float); int nSize_struct = nElem * sizeof(struct AoS); dim3 block(1024); dim3 grid(nElem / block.x); // 主机内存 float *a_h = (float*)malloc(nSize); float *b_h = (float*)malloc(nSize); float *res_h = (float*)malloc(nSize); struct AoS *res_from_gpu_AoS_h = (struct AoS*)malloc(nSize_struct); struct SoA *res_from_gpu_SoA_h = (struct SoA*)malloc(sizeof(struct SoA)); initialData(a_h, nElem); initialData(b_h, nElem); // 设备内存 float *a_d, *b_d; struct AoS *res_AoS_d; struct SoA *res_SoA_d; CHECK(cudaMalloc((void**)&a_d, nSize)); CHECK(cudaMalloc((void**)&b_d, nSize)); CHECK(cudaMalloc((void**)&res_AoS_d, nSize_struct)); CHECK(cudaMalloc((void**)&res_SoA_d, sizeof(struct SoA))); // 拷贝到设备 CHECK(cudaMemcpy(a_d, a_h, nSize, cudaMemcpyHostToDevice)); CHECK(cudaMemcpy(b_d, b_h, nSize, cudaMemcpyHostToDevice)); // 启动核函数 // double start = efficiency(); // AoSGpu<<<grid, block>>>(res_AoS_d, a_d, b_d, nElem); // cudaDeviceSynchronize(); // double end = efficiency(); // cout << "AoS Elapse: " << end - start << " sec" << endl; double start1 = efficiency(); SoAGpu<<<grid, block>>>(res_SoA_d, a_d, b_d, nElem); cudaDeviceSynchronize(); double end1= efficiency(); cout << "SoA Elapse: " << end1 - start1 << " sec" << endl; // 回拷结果 CHECK(cudaMemcpy(res_from_gpu_AoS_h, res_AoS_d, nSize_struct, cudaMemcpyDeviceToHost)); CHECK(cudaMemcpy(res_from_gpu_SoA_h, res_SoA_d, sizeof(struct SoA), cudaMemcpyDeviceToHost)); // 验证 sumCpu(res_h, a_h, b_h, nElem); //checkResult_structAoS(res_h, res_from_gpu_AoS_h, nElem); checkResult_structSoA(res_h, res_from_gpu_SoA_h, nElem); // 释放 free(a_h); free(b_h); free(res_h); free(res_from_gpu_AoS_h); free(res_from_gpu_SoA_h); cudaFree(a_d); cudaFree(b_d); cudaFree(res_AoS_d); cudaFree(res_SoA_d); return 0; }

• AoS 稳定值：约0.0027~0.0030 秒。

• SoA 稳定值：约0.0017~0.0023 秒，且呈逐步下降趋势（L2 缓存预热效应）。

• 加速比：SoA 比 AoS 快约1.5~1.8 倍（取稳定值 0.0028 / 0.0017 ≈ 1.65×）。

AoS的每次请求花费的sector是SoA的两倍，这也是性能拖垮的重要原因

SoA刚好占满128字节，但是AoS是32*8=256字节，也就是32线程跨了两个cacheLine，需要8个sector才能完成，所以很明显AoS每次都搬运一些垃圾值，而且占带宽

总结

在写函数的时候，我们尽量根据数据的特性，安排对齐合并访问，提高带宽利用率，这样能够很大的提高效率，还有数据如果是结构体数组，我们可以转化为数组结构体，有关实验出入的地方很可能跟显卡配置或者当时的显卡处于什么样的环境有关，如果有讲解错误的地方，欢迎大家指出，每个人的实验结果可能不一样，我们应该看趋势，在面对真实的生产环境的时候，应该用ncu/nsys