内存访问模式

目录前言存储时的对齐对齐与合并访问全局内存读取全局内存写入结构体数组和数组结构体总结前言本文将通过线程束的视角了解数据如何流动以更加基础的视角去优化性能如何通过优化全局内存的访问模式来榨干显存带宽存储时的对齐int *d_data nullptr; cudaMalloc((void**)d_data, sizeof(int) * n); 也就是说返回的地址一定满足 address % 256 0 比如返回的地址可能是 0x700000000 十六进制末尾两位都是0说明至少256字节对齐 0x700000100 0x700000200 而不会是 0x700000001 这种地址cudaMalloc不会返回 0x700000003256字节对齐是最低保证实际上很多情况下是512字节甚至更大粒度的对齐取决于驱动实现。为什么是256字节不是其他的字节呢GDDR6GraphicsDoubleDataRate6是第六代图形专用双倍数据速率同步动态随机存储器。显存位宽是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 Transaction128 字节DRAM ↔ L2 缓存显存控制器级别缓存扇区Cache Sector32 字节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 个float128 字节Warp Scheduler 发射指令Warp Scheduler 发射一条LDGLoad Global指令给LSU。LSU 合并地址LSU 收集这 32 个线程的地址发现它们落在同一个 128 字节 Cache Line 内合并成一次请求。查 L1 缓存LSU 向L1 缓存控制器发送请求。L1 控制器检查请求的 Cache Line 是否在 L1 里命中→ 直接从 L1 返回数据给 LSU跳到第 6 步。缺失→ 继续第 4 步。查 L2 缓存L1 控制器把请求转发给L2 缓存控制器。L2 控制器检查命中→ L2 把整个 128 字节 Cache Line 返回给 L1L1 缓存后返回数据给 LSU跳到第 6 步。缺失→ 继续第 5 步。查 DRAML2 控制器把请求发给显存控制器。显存控制器向GDDR6 显存颗粒发起一次 128 字节的物理事务搬来整个 Cache Line。数据原路返回DRAM → 显存控制器 → L2 → L1 → LSU。LSU 分发数据LSU 拿到数据块按每个线程的原始请求切分把正确的 4 字节写入每个线程的目标寄存器。指令完成。二、存储的完整流程从寄存器到 DRAM假设同一个 Warp 要把计算结果写回全局内存Warp Scheduler 发射指令Warp Scheduler 发射一条STGStore Global指令给LSU。LSU 合并地址LSU 收集 32 个线程的目标地址如果连续就合并成一次写请求。写 L1 缓存LSU 向L1 缓存控制器发送写请求。L1 写回策略数据先写入L1 缓存并标记为“脏”Dirty不立即写到 L2 或 DRAM。只有这个脏的 Cache Line 被替换时才一次性写回 L2。L1 写穿透策略某些 GPU 架构支持此模式数据同时写入 L1 和 L2确保 L2 立刻看到更新。写 L2 缓存当 L1 需要替换掉一个脏 Cache Line 时数据被写回L2 缓存。L2 也采用写回策略——数据先在 L2 里标记为脏只有被替换时才写回 DRAM。写 DRAM当 L2 需要替换掉一个脏 Cache Line 时数据才被写回DRAM。显存控制器向 GDDR6 显存颗粒发起一次 128 字节的物理写事务。对于优化内存我们主要关心以下两个特性对齐内存访问回答的是“数据的起始地址在哪里”的问题。它要求数据的首地址是对齐到某个固定值如 128 字节的。合并内存访问回答的是“同一个 Warp 的 32 个线程它们的地址是否连续”的问题。它要求这些地址落在同一个 128 字节的对齐区域内。为了最大化全局内存带宽你写的核函数必须让同一个 Warp 的线程产生“对齐合并访问”。两个必须同时满足的条件对齐这次合并访问的起始地址必须是 128 字节的整数倍。合并同一个 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 Line0x100和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~48KBwarp访问的时候导致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 i0,koffset;ksize;i,k){ res[i]a[k]b[k]; } } __global__ void sumArraysGPU(float*a,float*b,float*res,int offset,int n){ int iblockIdx.x*blockDim.xthreadIdx.x; int kioffset; if(kn) res[i]a[k]b[k]; } int main(int argc, char* argv[]){ int dev 0; cudaSetDevice(dev); int nElem118; int offset0; if(argc2){ offsetstoi(argv[1]); } printf(Vector size:%d\n,nElem); int nBytesizeof(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; iStartefficiency(); sumArraysGPUgrid,block(a_d,b_d,res_d,offset,nElem); cudaDeviceSynchronize(); iElapsefficiency()-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 -archsm_89 -Xptxas -dlcmca -o main main.cu #开启L1 nvcc -O3 -archsm_89 -Xptxas -dlcmcg -o main main.cu #关闭L1大家可以试一下是否还可以关闭和开启开启的关闭的可以看出来关闭L1会导致所有的数据都不经过L1缓存了但是有趣的是如果你的offset0的时候因为数据是流式的访问也就是只访问一次硬件识别到了所以无论你开启和关闭L1他都不经过L1直接经过L2所以现代编译器的作用还是挺大的很多理论的分析实验出来都可能不一样所以我们需要结合实际去分析以下是实验分析实验数据是118个元素每个元素是float4个字节然后a数组和b数组所以是两份所以总的read数据量118*4*22MB为此我们需要监测显存的read.sum和L1缓存的sectordram_bytes_read.sum l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratiooffsetl1tex__average_t_sectors_per_request...ratio是否对齐解读04.00✅ 对齐32 线程的连续访问恰好装满一个 128 字节 Cache Line4 个 Sector一次事务满载15.00❌ 不对齐128 字节数据横跨两个 Cache Line多跨越一个边界需要 5 个 Sector 才能覆盖25.00❌ 不对齐同上offset2 同样横跨两个 Cache Line35.00❌ 不对齐同上offset3 同样横跨两个 Cache Line结论注意代码里面的访问都是连续的为什么会出现offset0的时候会是4原因是他们是对齐且连续的为什么offset1是5甚至大家可以试一试offset7也是5因为不对齐但是连续注意到了offset8是4因为对齐且连续dram__bytes_read.sum为什么 offset0 反而更大offsetdram__bytes_read.sum解读02.40 MB对齐但 DRAM 搬运量反而最大12.10 MB不对齐DRAM 搬运量反而更小22.10 MB同上32.11 MB同上这个结果初看反直觉——为什么完美对齐的 offset0 反而比不对齐的 offset1,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看一下是否会增加到3MBPCIe接口 → Copy Engine → L2缓存(LTS) → 显存控制器 → GDDR6好的写到这里博主已经很蒙蔽了上面的是昨天的笔记但是这里是今天今天去测试来看无论是注释掉还是不注释cudaMemset基本都是2.10MB左右了所以昨天的2.4MB是一个异常情况是因为GPU运行时博主开了过多的后台渲染导致L2缓存不断的刷新简单来说我们要清楚两个前提一个是写分配一个是cudaMemset已经提前帮我们缓存到L2了还有cudaMemcpy也会经过L2被缓存一些所以我们只要关心最终的即可不要再去详细的分析为什么不是3为什么不是2等等写到这里博主想要说明的是对齐合并能够保证每次128字节事务都装满有效数据我们的每一次请求都会有4个sector而不是5个6个等等线程块 x 维度设为 32 的倍数这是保证 Warp 内线程 ID 连续、访问地址连续的最简单方法。保证起始地址对齐每个线程处理连续地址避免跨步访问不要用a[i * stride]这类大跨步访问。仅仅改变一下访问模式我们就能获得一些优化nvidia-ada-gpu-architecture.pdf这份是关于ada架构的白皮书所有有关的如果觉得博主讲解有问题的可以自行查阅相关资料博主也是刚入门学习Kepler 等老架构纹理缓存有独立于 L1 的专用 SRAM属于 TPC 层级的较大存储顶层架构图可以画出独立模块Maxwell 及之后包含 AdaNVIDIA 把纹理数据存储收拢进 SM 内部的 L1TEX仅保留 TEX 运算单元在 SM 里存储层面和 L1 合并顶层芯片布局就不再单独画出纹理缓存区块。注意纹理缓存只读缓存L1缓存是共用同一块SRAM的然后再加上共享内存也就是相当于我们之前讲的L1缓存内部再进行细分普通L1缓存纹理缓存只读缓存这三者的处理逻辑不一样但是SRAM是同一块┌─────────────────────────────────────────┐ │ SM │ │ │ │ ┌──────────────────────────────────┐ │ │ │ L1TEX SRAM(128KB) │ │ │ │ 可划分通用L1 / 共享内存 │ │ │ │ 内部复用存储只读通路、纹理数据 │ │ │ │ 多条逻辑通路共用这一块物理内存 │ │ │ └──────────────────────────────────┘ │ │ │ │ ┌──────────────────────────────────┐ │ │ │ 常量缓存独立SRAM64KB │ │ │ │ 专属广播硬件不占用L1TEX空间 │ │ │ └──────────────────────────────────┘ │ │ │ └─────────────────────────────────────────┘Ada 架构没有独立物理 SRAM 只读缓存。 「只读缓存」本质 GPU LSU 一条独立的只读访存通路数据存放于l1tex128KB 统一 SRAM 优势不受-dlcmcg影响。 当你使用-dlcmcg让普通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 缓存 → 显存控制器 → DRAMWarp Scheduler 发射指令发射一条STGStore Global指令给 LSU。LSU 合并地址收集 32 个线程的目标地址。如果地址连续且对齐合并成一次大写入请求如果分散拆成多次小请求。写入 L1 缓存数据首先写入 L1 缓存并标记为“脏”Dirty。此时数据还没有到达 DRAM。写回 L2当 L1 中这个脏的 Cache Line 被替换时数据被写回 L2 缓存。写回 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 (124) // 提前定义 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(); // AoSGpugrid, block(res_AoS_d, a_d, b_d, nElem); // cudaDeviceSynchronize(); // double end efficiency(); // cout AoS Elapse: end - start sec endl; double start1 efficiency(); SoAGpugrid, 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*8256字节也就是32线程跨了两个cacheLine需要8个sector才能完成所以很明显AoS每次都搬运一些垃圾值而且占带宽总结在写函数的时候我们尽量根据数据的特性安排对齐合并访问提高带宽利用率这样能够很大的提高效率还有数据如果是结构体数组我们可以转化为数组结构体有关实验出入的地方很可能跟显卡配置或者当时的显卡处于什么样的环境有关如果有讲解错误的地方欢迎大家指出每个人的实验结果可能不一样我们应该看趋势在面对真实的生产环境的时候应该用ncu/nsys