目录
背景
前面我们一直在深挖cuda执行模型，一直在调整核函数的各种配置，但是我们无法忽略的一个点就是内存带宽，这一点也是影响整个效率的一重大因素，有时候甚至是决定性能的第一大瓶颈
GPU的算力是非常恐怖的，我的 RTX 4060 单精度浮点峰值约 15 TFLOPS，但是你必须要时时刻刻喂满数据，否则他们就算能够工作，没有原料，也无法开机啊。
本章就是通过不同的访存模式，然后实操，看一下对比，来使得我们的核函数能够更搞笑的运行
内存层次结构的优点
对于有cpu编程的基础，我们知道寄存器最快，再到缓存，再到内存，再到磁盘
容量越小，速度越快，造价越贵我们知道一个程序是具有局部性的，这种局部性是我们访问数据的一种模式
时间局部性：同一个内存地址，在短时间内被反复访问。
例子：归约循环中反复读写
idata[tid]。好处：这个地址的数据可以被寄存器缓存或 L1 命中，避免反复读显存。
空间局部性：附近的内存地址，在短时间内被连续访问。
例子：矩阵加法中，同一个 Warp 的 32 个线程访问同一行中连续的 32 个
float。好处：这些连续地址可以被合并成一次 128 字节事务，一次搬运全部满足。
程序有局部性，那硬件也会有局部性
空间局部性支持：GPU 的 L1/L2 缓存和显存控制器能以 128 字节粒度搬运数据。这是为“空间局部性”设计的硬件基础。即使你只请求 4 字节，硬件也会搬来整个 Cache Line。
时间局部性支持：多级缓存（L1、L2）会自动保留最近访问过的数据。如果短期内再次访问同一地址，就从缓存命中，而不是再去显存读。
程序的局部性，几乎是所有程序与生俱来的特征，远早于现代缓存硬件出现。
也就是大家编写程序的时候发现往往附件的部位会被重复的访问，所以后来硬件才慢慢催生出来局部性，去配合软件达到最佳的性能，最后为了追求高效率，设备将越来越优化局部性，而程序也会越来越局部化。
在cpu编程的时候，我们很明显的发现
寄存器：由编译器自动分配，你无法直接指定哪个变量放在寄存器。编译器通过复杂的寄存器分配算法来决定，你只能通过调整代码间接影响它。
L1/L2/L3缓存：完全由硬件自动管理。内存访问模式会触发缓存行填充和替换，你无法发出“把这个数组塞进L2缓存”的指令。你唯一能做的是通过编写具有良好局部性的代码（如连续访问、循环分块），来让缓存系统发挥最佳性能。
但是在cuda当中，你可以随意的安排，CUDA编程模型将内存层次结构更好的呈献给开发者，让我们显示的控制其行为。
共享内存：通过
__shared__关键字直接声明，你决定哪些数据进共享内存，哪些线程共享它，以及通过__syncthreads()控制同步。这是CPU编程完全没有的体验。全局内存：通过
cudaMalloc分配，你显式管理数据在主机和设备间的传输（cudaMemcpy）。寄存器：虽然你依然不能直接说“这个变量必须放寄存器”，但CUDA提供了工具（如
nvcc -Xptxas -v或ncu）让你精确查看每个线程用了多少寄存器。你可以通过调整代码（如改变Block大小、减少临时变量）来主动影响寄存器用量，从而控制SM上活跃Warp的数量（占用率）。这种反馈-调整的循环是CUDA性能优化的核心环节。所以接下来我们就逐一介绍这些内存模型，然后如何去控制，如何去配合执行模型发挥最大的性能
内存模型
我们可以把大致的内存分为
1.可编程的
2.不可编程的
正如字面意思，可编程的就是我们能够通过代码去控制的，不可编程就是硬件写死的，你只能了解它的特性然后配合它，发挥最大的性能
- 寄存器
- 共享内存
- 本地内存
- 常量内存
- 纹理内存
- 全局内存
寄存器
关于寄存器在cpu编程当中我们已经很熟悉了
寄存器是 GPU 存储层级中速度最快、离计算单元最近的存储空间。
寄存器是 SM 内部的一小块静态随机存取存储器 (SRAM)。它和 CPU 里的寄存器概念完全一样，但在 GPU 里，寄存器的数量要多得多。寄存器的物理载体是 寄存器文件 (Register File)，它直接集成在 SM 内部，紧挨着 CUDA 核心。
multiProcessorCount //SM的数量 regsPerMultiprocessor //SM中的寄存器数量#include <iostream> #include "../freshman.hpp" #include <cuda_runtime.h> using namespace std; int main(){ int count=0; CHECK(cudaGetDeviceCount(&count)); for(int i=0;i<count;i++){ //查询设备信息 cudaDeviceProp prop; CHECK(cudaGetDeviceProperties(&prop,i)); cout<<"SM数量："<<prop.multiProcessorCount<<endl; cout<<"每个流多处理器的寄存器个数："<<prop.regsPerMultiprocessor<<endl; } return 0; }
我的硬件是RTX4060：每个寄存器是32位，一个SM里面有65536个寄存器,总容量为 256KB
这些65536个寄存器共同组成一个叫寄存器文件
寄存器是否是线程独有的呢？？？
- 软件编程视角：是。在 CUDA 编程模型中，每个线程拥有完全独立、私有的寄存器上下文，线程间寄存器不可互访，逻辑上完全独占。
- 硬件实现视角：不是。物理寄存器堆是单个 SM 内所有线程共享的统一资源，以 Warp（线程束）为单位静态分配，通过硬件地址映射实现 “逻辑私有”。
从写核函数的开发者角度看，每个线程的寄存器是严格隔离的：
- 上下文独立：每个线程有自己独立的寄存器编号空间，线程 A 的
r0寄存器和线程 B 的r0寄存器完全无关，互相不能直接读写，各自保存自己的局部变量、循环计数器、函数返回地址等。- 默认分配规则：核函数内的非数组局部变量，编译器会优先分配到寄存器中，每个线程独立持有一份，这也是 GPU 能支持大量线程并行执行、互不干扰的基础。
- 执行隔离：即使同一个 Warp 内的线程发生分支分化，各自的寄存器状态也会独立保留，不会互相覆盖。
唯一例外：Warp 内洗牌指令（如
__shfl_sync）可以通过硬件通道，直接读取同一个 Warp 内其他线程的寄存器值，这是 Warp 内最快的通信方式，延迟远低于共享内存。1. 通常情况：能放寄存器就放寄存器
寄存器是每个线程最快、最私有的存储。对于核函数里的标量变量（比如
int i; float tmp;），编译器总是想方设法把它们塞进寄存器里。因为这些变量被频繁读写，放寄存器里才能保证性能。2. 溢出情况：寄存器不够了就放局部内存
但 SM 的寄存器总量是严格有限的（我的每个 SM 有 65536 个 32 位寄存器）。当核函数过于复杂——变量太多、数组太大，或者强制设定了很低的寄存器使用上限——编译器会优先保留使用频率高的变量在寄存器里，把使用频率低的溢出到本地内存（这叫做寄存器分配的启发式策略）。另外生命周期短的变量反而更容易留在寄存器里（因为它用完就释放了，不长期占用寄存器slot），生命周期长的变量更容易被溢出。。不过这个过程是你无法直接控制的，但可以从结果上看出。
3. 被优化掉的情况：变量根本不存在
如果编译器发现某个局部变量从未被使用，或者它只是一个中间结果，可以被直接优化掉，那么这个变量就不会占用任何物理存储，既不占寄存器，也不占局部内存。
4. 如何知道一个变量去了哪里？
想知道声明一个变量后它到底去了哪里，最直接的方法是使用我们在
ncu中学到的方法，去“查看结果”。方法一：编译时查看总体统计
nvcc -Xptxas -v your_kernel.cu这个命令会输出每个核函数每个线程平均使用了多少个寄存器。你可以通过反复注释掉或添加某个变量，然后看寄存器的总数量是增加还是减少，间接推断这个变量是否被放入了寄存器。
方法二：运行时查看溢出情况
用ncu查看局部内存的读写指标。# 本地内存load总量（包括命中和未命中缓存的） ncu --metrics l1tex__t_bytes_pipe_lsu_mem_local_op_ld.sum \ l1tex__t_bytes_pipe_lsu_mem_local_op_st.sum ./your_program如果计数器不为0，就说明发生了寄存器溢出，有变量被放到了局部内存。
5. 规则与例外
变量类型
通常是去寄存器吗？
备注
标量变量 (
int i; float f;)是，优先放寄存器
会被频繁读写，放寄存器最快
小型数组 (
int arr[4];)优先放寄存器，但容易溢出
比标量更容易撑爆寄存器文件
大型数组 (
int arr[256];)几乎必然溢出到局部内存
寄存器根本塞不下，只能放局部内存
编译器优化掉的变量
哪里都不去，直接消失
根本没用到的变量，不占任何存储
总结
“声明一个变量就会存到寄存器” 这个说法，在理想情况下是对的，但并不是绝对的。正确的理解是：编译器总是竭尽全力把你的变量放进寄存器，但当寄存器紧张时，它会悄悄把部分变量“溢出”到局部内存（本地内存），这会严重拖慢速度。 这正是为什么我们要学会查看寄存器用量，并尽可能优化代码、减少变量占用的根本原因。
我们要尽可能的减少溢出，因为本地内存是属于全局显存的某块区域，也就是访问效率跟全局显存一样非常低效，如果溢出了就导致你的整个核函数访问某些变量的时候会慢，整体效率就不行
实验
#include <iostream> #include "../freshman.hpp" #include <cuda_runtime.h> using namespace std; __global__ void kernel(){ unsigned int idx=blockIdx.x*blockDim.x+threadIdx.x; printf("当前线程号：%d\n",idx); } int main(){ int count=0; CHECK(cudaGetDeviceCount(&count)); for(int i=0;i<count;i++){ //查询设备信息 cudaDeviceProp prop; CHECK(cudaGetDeviceProperties(&prop,i)); cout<<"SM数量："<<prop.multiProcessorCount<<endl; cout<<"每个流多处理器的寄存器个数："<<prop.regsPerMultiprocessor<<endl; } int dev=0; cudaSetDevice(dev); dim3 block(32); dim3 grid(4); kernel<<<grid,block>>>(); cudaDeviceSynchronize(); return 0; }
8 bytes stack frame：核函数的栈帧大小是 8 字节。这个栈帧是局部内存的一部分，用来存放函数调用时的返回地址、少量临时变量等。8 字节非常小，完全正常。
0 bytes spill stores和0 bytes spill loads：这两项为 0，是性能好的直接证据！
Spill（溢出）：指编译器因为寄存器不够用，被迫把某些变量从寄存器暂时存放到局部内存（显存）里。
Spill stores：存放到局部内存的字节数。0 表示完全没有变量被溢出。
Spill loads：从局部内存加载回来的字节数。0 同上。
Used 24 registers：每个线程使用了 24 个 32 位寄存器。这是一个精确的数字。可以用它来估算 SM 的占用率。比如我的 Block 有 32*4=128个线程，一个 Block 就用24*128=3072个寄存器。SM 总共有 65536 个，理论上能同时放21个这样的 Block（但受限于线程数等其他限制，实际不可能达到）。
used 0 barriers：核函数里没有使用__syncthreads()这样的同步屏障，所以不需要硬件屏障资源。
8 bytes cumulative stack size：累计栈大小，和前面的栈帧一致。
352 bytes cmem[0]：这是另一个关键的常量内存使用量。cmem[0]是专门用来从主机端传递核函数参数到设备端的常量内存。传给核函数的参数越多、越大，这个数字就越大。352 字节相当少，说明参数很简单。为什么是24个寄存器呢，我明明只定义了一个int啊？用1个不就行了？？？
mov.u32 %r0, %ctaid.x // 存blockIdx.x mov.u32 %r1, %ntid.x // 存blockDim.x mul.lo.u32 %r2, %r0, %r1 // 存乘法结果 mov.u32 %r3, %tid.x // 存threadIdx.x add.u32 %r4, %r2, %r3 // 存最终idx printf在GPU上的实现非常重，它需要： 格式字符串的地址（1个寄存器） 参数值idx（1个寄存器） printf内部缓冲区的指针（1~2个寄存器） 写入缓冲区时的偏移量计算（多个寄存器） 内部的原子操作（需要额外寄存器保存中间状态） 函数调用的返回地址、参数传递等（几个寄存器） 编译器为了让指令流水线更高效，有时会故意多保留一些中间值在寄存器里 （而不是算完立刻丢掉），用空间换时间，这也会增加寄存器用量。如果我们把printf去掉呢？？？
直接变成了4个，那就是printf占用了绝大多数的寄存器
接下来我定义一个非常小的数组
__global__ void kernel(){ // 故意声明大量变量 int a1=threadIdx.x, a2=a1*2, a3=a2*3, a4=a3*4; int b1=a4+1, b2=b1+2, b3=b2+3, b4=b3+4; // 继续堆变量... int arr[64]; // 故意放一个大数组 for(int i=0;i<64;i++) arr[i]=i*threadIdx.x; printf("%d %d\n", arr[threadIdx.x%64], a1+a2+a3+a4+b1+b2+b3+b4); } arr[64] = 64个int = 64 * 4字节 = 256字节 = 64个寄存器 加上其他变量 a1~a4, b1~b4 = 8个寄存器 加上printf = 约20个寄存器 总计 ≈ 64 + 8 + 20 = 92个寄存器/线程
30 registers：只比原来多6个，是为了处理threadIdx.x%64这个下标计算多用了几个寄存器寄存器显然没有像我们说的那样夸张，反而是栈帧变大了，每个线程的栈帧从原来的8字节变成了272字节，因为编译器分析我们这些代码发现有些arr是只写不读，这就是为什么溢出记录为0，是因为其实栈帧是预留这些空间的，但是实际上连写都没写，因为你后面压根都不读，只是预留了空间罢了
现代编译器的优化能力非常强，很多你以为会产生开销的代码，实际上编译后根本不存在。这也是为什么性能分析必须看实际的PTX/SASS指令或者ncu运行时指标，而不能只看C++源码来推断开销。
博主测试了很多的条件，发现溢出还是不溢出压根就不是理论能够分析出来的，可能是博主的能力不足，分析的不够透彻，实际中编译器的优化是相当的严重，能优化的地方全部都优化，这导致不好理论分析，可能只能分析个上界，所以实际的程序我们需要用ncu去分析为了避免溢出，我们可以采用以下方式
首先要明确溢出的根本原因在于：
每个线程使用的寄存器*总的线程数>硬件寄存器个数
单个线程使用的寄存器>255
所有的方法都是减少单个线程使用的寄存器数量
用
__launch_bounds__精确控制// 告诉编译器：这个kernel最多用256线程/block，每个SM至少驻留2个block __global__ void __launch_bounds__(MaxThreadsPerBlock, MinBlocksPerMultiprocessor) myKernel() { // 核函数体 }
MaxThreadsPerBlock（必需）：告诉编译器“我这个核函数最多会用多少个线程来启动一个 Block”。比如你总是用block(256)，这里就填256。
MinBlocksPerMultiprocessor（可选）：告诉编译器“我希望每个 SM 上至少能驻留多少个 Block”。编译器会尽量遵守这个要求。编译器会根据线程数量从而限制每个线程使用的寄存器数量
--maxrregcount限制寄存器用量nvcc --maxrregcount=64 your_kernel.cu强制每个线程最多用64个寄存器，超出的编译器自动溢出到本地内存。
这是利用可控的寄存器去换占用率，如果每个线程的寄存器不可控，那编译器在汇编阶段就会瞎分配，比如每个线程分配的很大，那这个SM的block占有率就非常的低，那最后并发就可能很差，所以我们可以手动的去控制寄存器个数，从而让block变多，那调度的的线程就变多，那并发可能会更好，也就是隐藏延迟首选：减少局部数组 → 用shared memory替代 减少同时存活的变量 → 重构代码逻辑 次选：拆分kernel → 降低单个kernel的复杂度 调优阶段：--maxrregcount / __launch_bounds__ → 精确控制寄存器和occupancy的权衡 控制#pragma unroll的展开倍数 最后手段：接受部分溢出，用更高的occupancy来补偿spill带来的延迟溢出不一定是问题，occupancy低才是问题——如果溢出量很小、且被L1/L2缓存住了，对性能的实际影响可能微乎其微，不需要花大力气去消除。真正需要处理的是"溢出导致occupancy下降、进而导致延迟隐藏能力不足"这种情况。
最后我们要根据实际的核函数+硬件+ncu去分析，而不是草率的认为，我们要评估所有的优化手段哪个的收益比较高
本地内存
┌─────────────────────────────────────────────────────────────┐ │ GPU 芯片 (片内，SRAM) │ │ ┌───────────────────────────────────────────────────────┐ │ │ │ SM (流多处理器) │ │ │ │ ┌─────────┐ ┌─────────┐ ┌──────────────────────┐ │ │ │ │ │ 寄存器 │ │ 寄存器 │ │ 共享内存 / L1 缓存 │ │ │ │ │ │ (线程0) │ │ (线程1) │ │ (Block内共享，用户控制)│ │ │ │ │ └─────────┘ └─────────┘ └──────────────────────┘ │ │ │ └───────────────────────────────────────────────────────┘ │ │ │ │ │ ┌─────┴─────┐ │ │ │ L2 缓存 │ │ │ │ (所有SM共享)│ │ │ └─────┬─────┘ │ └──────────────────────────┼──────────────────────────────────┘ │ 跨芯片总线 ┌──────────────────────────┼──────────────────────────────────┐ │ 显存颗粒 (片外，DRAM) │ │ ┌───────────────────────┴───────────────────────────────┐ │ │ │ 全局内存 (Global Memory) │ │ │ │ - 通过 cudaMalloc 分配 │ │ │ │ - 所有线程/SM 都可读写的“公共仓库” │ │ │ │ ┌─────────────────────────────────────────────────┐ │ │ │ │ │ 本地内存 (Local Memory) │ │ │ │ │ │ - 线程独享的“私人文件柜” (在全局内存中划出) │ │ │ │ │ │ - 编译器寄存器溢出 / 大型局部数组 / 栈帧超限 │ │ │ │ │ │ - 速度慢 (走L1/L2缓存, 物理上在DRAM) │ │ │ │ │ └─────────────────────────────────────────────────┘ │ │ │ └───────────────────────────────────────────────────────┘ │ │ ┌───────────────────────────────────────────────────────┐ │ │ │ 常量内存 / 纹理内存 │ │ │ └───────────────────────────────────────────────────────┘ │ └────────────────────────────────────────────────────────────┘以下情况是会存储到本地内存的
情况一：寄存器溢出（spill）：变量太多，寄存器不够用，编译器把低优先级的变量踢到本地内存。
情况二：编译器无法确定下标的局部数组：
int arr[64]; int i = some_runtime_value; arr[i] = ...; // 下标是运行时变量，编译器无法把arr拆成64个独立寄存器 // 只能整体放进本地内存的栈帧如果下标是编译期常量（比如
arr[0]、arr[1]），编译器可以把每个元素映射成独立的寄存器；一旦下标是运行时变量，整个数组就必须放进本地内存。情况三：函数调用的栈帧：之前
8 bytes stack frame，就是printf函数调用产生的，和CPU上的函数调用栈是一样的概念。但要注意如果很小时可能会优化在寄存器，当栈帧过大才可能会优化到本地内存为什么会存储到本地内存，核心原因是这些都是线程独有的，而不是共享，所以是不会存储到L1L2缓存，L1L2缓存只能作为中间层去加速计算单元和本地内存的访问效率，而不是存储
如果访问频繁，L1会把它缓存住，实际延迟会比直接打到DRAM低很多
共享内存
L1 缓存和共享内存物理上是同一块 SRAM，并且它们都属于片上（On-Chip）存储
延迟：约1~32个cycle，比寄存器（0 cycle）慢一点，比L2（约200 cycle）和DRAM（400~800 cycle）快得多。
for(int i=0;i<count;i++){ //查询设备信息 cudaDeviceProp prop; CHECK(cudaGetDeviceProperties(&prop,i)); cout<<"SM数量："<<prop.multiProcessorCount<<endl; cout<<"每个流多处理器的寄存器个数："<<prop.regsPerMultiprocessor<<endl; cout<<"每个流多处理器允许使用的最大的shared_memory："<<prop.sharedMemPerMultiprocessor<<endl; cout<<"每个流多处理器允许block使用的最大的shared_memory："<<prop.sharedMemPerBlock<<endl; }
硬件上限是100KB，48KB
注意共享内存的生命周期是随block的，假设总的可用的shared_memory是96KB，那每个block使用32KB，那最多只有3个block驻留，如果我把每个block使用的下调到16KB，那此时就能驻留6个
声明方式：关键字__shared__
静态声明（编译期确定大小）
__global__ void kernel(){ __shared__ int smem[1024]; // 固定大小，编译期已知 __shared__ float tile[32][32]; // 二维数组也可以 smem[threadIdx.x] = ...; __syncthreads(); // 必须同步才能安全读别人写的数据 int val = smem[(threadIdx.x+1) % 1024]; }动态声明（运行时确定大小）
__global__ void kernel_dynamic() { // 声明动态共享内存，大小由 <<< >>> 第三个参数决定 extern __shared__ float dynamic_data[]; // 多个数据类型时，需从同一个数组首地址手动划分 // int* int_data = (int*)dynamic_data; // 前一部分存int // float* float_data = (float*)&int_data[需要的int数量]; // 后一部分存float }动态声明适合大小依赖运行时参数的情况，但一次只能有一个
extern __shared__数组（可以用指针偏移模拟多个数组）。除了以上的划分，还可以通过以下函数去划分大小
//针对单个核函数的调整 cudaFuncSetCacheConfig(my_kernel, cudaFuncCachePreferShared); // 可选配置： // cudaFuncCachePreferNone - 默认 // cudaFuncCachePreferShared - 优先共享内存（更大的共享内存，更小的 L1） // cudaFuncCachePreferL1 - 优先 L1（更大的 L1，更小的共享内存） // cudaFuncCachePreferEqual - 均分 //针对硬件进行调整，所有的核函数都会被影响 cudaDeviceSetCacheConfig(cudaFuncCachePreferShared); // 影响之后所有核函数 //可以通过以下去查询实际动态分配了多少 int sharedSize, l1Size; cudaDeviceGetAttribute(&sharedSize, cudaDevAttrMaxSharedMemoryPerBlockOptin, 0); cudaDeviceGetAttribute(&l1Size, cudaDevAttrMaxL1CacheSize, 0);
默认：每 Block 最多 48KB 共享内存，不调用配置函数也可以直接静态分配，只要不超过 48KB。
如果需要超过 48KB：调用
cudaFuncSetCacheConfig(kernel, cudaFuncCachePreferShared)，上限提升到 100KB。无需配置的情况：共享内存用量低于 48KB 时，完全不必调用该函数。
为什么需要使用共享内存？？？
1：避免重复的读取global memory
比如之前两个矩阵加法，我们完全可以先声明一个共享内存，然后再相加之前，先把数据从global_memory读取到shared_memory
2：block内线程间通信
思路还是一样，通过共享内存，我们可以把归约的数据放到共享内存，然后再做归约，这里共享内存是可以互相通信的，所以可以加一个同步
3：替代大型局部数组
之前讲寄存器的时候，为了避免溢出到本地内存，我们也可以把局部数组放到共享内存，减少延迟，代价：共享内存是block内共享的，用这种方式要确保不同线程的区段不重叠，而且总用量不超过每个block的共享内存上限。
Bank Conflict
Bank（存储体）是 CUDA 共享内存的硬件级并行访问单元。访问共享内存是有32个物理通道，我们叫做bank，也就是一次访问只能32个bank
假设我们__shared__ int smem[1024]，那么此时就会把smem[0]放在bank0，smem[1]放在bank1，这样交错的分布，而不是把所有的smem放在bank[0]，因为你是一次访问，如果地址全部在bank0，只能是串行的去访问，只有访问的地址分布在不同的bank，这些物理通道才能并行
每个bank是独立工作的，一个周期只能被访问一次
注意bank里面是有多个地址的，如果你的线程访问同一个地址，那是不会冲突，直接广播给所有的线程，如果你访问同一个bank里面不同地址，那此时就会冲突
Bank号=(元素在共享内存中的字节偏移量/4)mod32所以你声明的数组当中如果按照从0开始，那么smem[0]在bank0，smem[32]也在bank0，如果你恰好访问smem[0]和smem[32]，那就会触发bank conflict，一个周期是读不完的，必须两个周期，那此时效率就会大大降低
所以提升效率就是避免bank conflict
比如
// 有conflict的转置：读合并，写有conflict __shared__ float tile[32][32]; tile[threadIdx.y][threadIdx.x] = input[...]; // 读，行访问，无conflict __syncthreads(); output[...] = tile[threadIdx.x][threadIdx.y]; // 写，列访问 // 列访问时，同一列的32个线程threadIdx.x各不相同，但threadIdx.y相同 // → tile[0][threadIdx.y], tile[1][threadIdx.y]... 这一列的32个元素 // → 全部属于bank threadIdx.y → 32路conflict！对于一个warp，他们的threadIdx.x是连续增大的，然后threadIdx.y是相同的，
那么此时tile是共享内存去读input全局显存，然后是按照行，所以是无conflict的
当把共享内存写到output时，tile[0][y] tile[1][y]……此时是按照列去写的，那么刚好这一列按照前面说的映射Bank号=(元素在共享内存中的字节偏移量/4)mod32，刚好落在同一个bank里面，此时直接conflict，效率直接下降，每一列都是冲突，都是串行
解决方案：多开一列
// 解决：padding一列 __shared__ float tile[32][33]; // 多一列padding // tile[0][threadIdx.y]→bank(threadIdx.y) // tile[1][threadIdx.y]→bank(threadIdx.y+1)（因为每行多了1个元素，地址错开） // 现在列访问里每个线程打到不同bank → 无conflict多开的这一列本质是为了让映射关系：Bank号=(元素在共享内存中的字节偏移量/4)mod32
也就是第一行的第33列映射到bank0，导致第二行的第一个映射到bank1，第三行的第一个映射到bank2，全部列都是错开的，此时就可以并行的去访问32路bank，大大提升并行效率
实验
//bank conflict __global__ void shared_memory(float *input, float *output){ int ix=threadIdx.x+blockIdx.x*blockDim.x; int iy=threadIdx.y+blockIdx.y*blockDim.y; __shared__ float tile[32][32]; tile[iy][ix] = input[iy*blockDim.x+ix]; __syncthreads(); output[iy*blockDim.x+ix]=tile[ix][iy]; } __global__ void shared_memory1(float *input, float *output){ int ix = threadIdx.x + blockIdx.x * blockDim.x; int iy = threadIdx.y + blockIdx.y * blockDim.y; __shared__ float tile[32][33]; tile[iy][ix] = input[iy * blockDim.x + ix]; __syncthreads(); output[iy * blockDim.x + ix] = tile[ix][iy]; } int main(){ int dev=0; cudaSetDevice(dev); dim3 block(32,32); dim3 grid(1); int numSize=1024; float* tmp=nullptr; float* input=nullptr; float* output=nullptr; tmp=(float*)malloc(sizeof(float)*numSize); initialData(tmp,1024); //开辟显存 cudaMalloc((void**)&input,sizeof(float)*numSize); cudaMalloc((void**)&output,sizeof(float)*numSize); cudaMemcpy(input,tmp,sizeof(float)*1024,cudaMemcpyHostToDevice); double start=efficiency(); shared_memory<<<grid,block>>>(input,output); cudaDeviceSynchronize(); double end=efficiency(); cout<<"未padding的："<<end-start<<endl; cudaMemcpy(input,tmp,sizeof(float)*1024,cudaMemcpyHostToDevice); double start1=efficiency(); shared_memory1<<<grid,block>>>(input,output); cudaDeviceSynchronize(); double end1=efficiency(); cout<<"padding的："<<end1-start1<<endl; return 0; }
未padding的：0.00109291
padding的：0.0000641346
ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum ./your_program这个指标统计的是共享内存加载操作中发生 Bank Conflict 的总次数。
表示总共发生了 992 次冲突的共享内存指令发射，很明显发生了32路冲突，但是是算的31*32=992，不是32*32哦，注意整个核函数执行 1 次读取操作是不算的，算多余的
如何正确使用共享内存
__global__ void correctUsage(int *g_data, int *g_out, int n){ __shared__ int smem[1024]; int tid = threadIdx.x; int idx = blockIdx.x * blockDim.x + tid; // 步骤1：从global memory加载数据到共享内存 smem[tid] = (idx < n) ? g_data[idx] : 0; // 步骤2：同步，确保所有线程都完成加载 __syncthreads(); // 步骤3：在共享内存上做计算（可以访问其他线程加载的数据） int left = (tid > 0) ? smem[tid-1] : 0; int right = (tid < 1023)? smem[tid+1] : 0; int result = left + smem[tid] + right; // 步骤4：如果后续还有共享内存写操作，再次同步 __syncthreads(); // 步骤5：写回global memory if(idx < n) g_out[idx] = result; }比较常见的共享内存的错误
错误一：忘记同步
smem[tid] = g_data[idx]; // 忘记__syncthreads() int val = smem[(tid+1) % blockDim.x]; // 可能读到未初始化的值错误二：共享内存大小超过上限
// CC 8.9每个block最多96KB共享内存 __shared__ int arr[96*1024/4 + 1]; // 超过上限，kernel launch失败可以用这个检查：
ncu --metrics shared_load_transactions_per_request ./your_program或者直接查询：
cudaDeviceProp prop; cudaGetDeviceProperties(&prop, 0); printf("每个block最大共享内存：%zu KB\n", prop.sharedMemPerBlock/1024);错误三：跨block访问
// 错误：共享内存只在当前block可见 __shared__ int smem[1024]; // block 0的线程无法访问block 1的smem，这是硬件保证的隔离总结
第一：我们得理解共享内存的特性，比如它是SM级别的共享，SM上的所有线程共享，是属于临界资源，如果访问需要注意并发访问数据导致一致性的问题，这里可以需要线程同步syncthreads()，但是频繁使用会降低效率第二：了解硬件的上限，共享内存多大
第三：重点去理解bank conflict，避免冲突，否则会降级为串行，导致效率低下
常量内存
常量内存是 CUDA 内存模型中一段只读、带专用硬件缓存、全局可见的显存地址空间。它物理上位于板载显存（DRAM），并非片上存储，但每个 SM 配备了专属的常量缓存与广播优化机制，在特定访问模式下性能接近片上共享内存，是 GPU 轻量级只读参数的首选存储方案。
查询常量内存的容量totalConstMem
64KB，注意片上专属于常量缓存是无法查询到的，大概是8~16KB，这是硬件自主控制的，能够被我们控制的才能查询到，我们更应该关心的是如何访问同一个地址而不发生冲突
如何声明常量内存
常量内存必须在所有函数（包括
main和核函数）之外，以__constant__关键字在文件作用域内声明。它的大小是静态固定的，且必须在主机端通过专门的 API 来更新数据。//使用关键字__constant__ // 全局作用域声明，可见于同一编译单元内的所有设备端代码 __constant__ float const_data[256]; int main() { float host_data[256] = { /* 初始化数据 */ }; // 使用 cudaMemcpyToSymbol 将数据从主机拷贝到常量内存 cudaMemcpyToSymbol(const_data, host_data, sizeof(host_data)); // ... } __global__ void kernel() { int tid = threadIdx.x + blockIdx.x * blockDim.x; // 所有线程都访问 const_data[tid]，访问模式决定了性能 float val = const_data[tid]; }
生命周期：与应用程序相同，在整个程序运行期间都存在。
大小限制：总量有限，通常为 64KB。
更新方式：核函数内无法修改，只能在主机端使用
cudaMemcpyToSymbol更新。常量内存的常量就是只读不改，要改只能通过主机端，不能通过核函数内部去修改
单周期广播机制
不像共享一样有bank，它内部就像一条线，只有一个窗口，一条专用的广播总线L1L2缓存是为了普遍性，其实内部也有bank和多级互连，所以内部才能把不同的地址合并，然后缓存就能识别分发不同的数据给不同的线程，但是常量缓存不是，他有自己的特性，它是为了访问同一个地址设计的，比如重力加速度g，你把它放在常量里面，后续能够加载到缓存，所有的线程同时访问时，它能够通过广播机制统一分发给不同的线程
常量缓存的设计目标极其单一：让一个 Warp 的所有线程在同一时刻，用最快的速度获取同一个值。（也就是硬件根据需求去设计了，你只能这样顺着硬件去设计软件发挥最大性能）
场景一：广播命中（最快）
地址解析：硬件检查这 32 个线程要访问的地址。
广播命中：如果发现所有 32 个线程访问的都是同一个物理地址，硬件就不会去反复读取存储体。它会直接从常量缓存中取出这个数据值。
单周期分发：通过一个专用的广播总线，将这个值同时发送给 Warp 内的所有 32 个线程。整个过程仅需 1 个时钟周期。
场景二：缓存缺失或地址不一致（变慢）
请求串行化：如果 32 个线程访问的是不同的地址，广播机制失效。这些请求无法被同时服务，会被硬件强制拆分成多次独立的请求，然后逐个从常量缓存（如果命中）或更底层的显存中读取。
性能惩罚：如果 N 个线程访问了 N 个不同的地址，最坏的情况下需要 N 个时钟周期才能完成服务。速度直接跌落神坛，甚至不如普通全局内存。
实验
#define N 256 __constant__ float c_data[N]; // 核函数A：所有线程读同一个地址 → 完美广播，最快 __global__ void broadcast_read(float *out, int idx) { int tid = threadIdx.x + blockIdx.x * blockDim.x; out[tid] = c_data[idx]; // 所有线程用同一个 idx } // 核函数B：每个线程读不同地址 → 串行化，慢 __global__ void scatter_read(float *out) { int tid = threadIdx.x + blockIdx.x * blockDim.x; out[tid] = c_data[tid % N]; // 每个线程不同地址 }
ncu --metrics \ l1tex__t_sectors_pipe_lsu_mem_constant_op_ld_lookup_hit.sum,\ l1tex__t_sectors_pipe_lsu_mem_constant_op_ld_lookup_miss.sum \ ./your_program
..._hit.sum：常量缓存命中次数（广播命中也计入）。
..._miss.sum：常量缓存缺失次数，会触发从 L2/DRAM 搬移整个 Cache Line。注意
- 老架构（图灵 sm_7x 及更早）：每个 SM 有一块独立的、专门服务常量内存的常量缓存，硬件有独立的计数器统计它的命中 / 未命中，所以这类
lookup_hit/lookup_miss指标能正常出数。- 新架构（安培 sm_80 及以后）：NVIDIA 取消了独立的常量缓存硬件设计，常量内存的访问直接并入统一 L1 / 纹理缓存，和全局内存、纹理内存共用同一块 SRAM 和同一套缓存流水线。
所以以上两个指标在我的4060上使用不了的，大家可以试试自己的设备
总结
关于常量内存我们要清楚其单周期的广播机制，根据这个机制我们应该把所有线程统一访问的只读的地址放到常量内存中加速我们的访问，避免多线程并行下访问不同的地址导致串行，从而降低效率，如果需要多线程访问不同的地址，我们可以利用共享内存，采用bank机制就可以访问不同的地址而避免串行化
纹理内存
纹理内存（Texture Memory）不是一块独立的物理显存，而是建立在全局内存之上、带专用硬件缓存与采样加速的只读访问通道。它原本为图形渲染的纹理采样设计，在通用计算中专门优化「2D/3D 空间局部性访问、硬件插值、自动边界处理」三类场景，是 CUDA 内存层级中功能特色非常鲜明的一员。
它的物理载体和全局内存、常量内存完全相同，都是板载显存；性能优势全部来自专用的纹理缓存硬件、采样加速流水线、内置的插值 / 边界处理逻辑。
本质和常量内存有点像，是从全局内存划分出来的一部分空间，然后通过纹理缓存去加速
新架构：纹理缓存、L1 数据缓存、常量缓存被整合为统一 L1 / 纹理缓存，共享同一块片上 SRAM 池，由硬件动态调度分配。
如何声明
现代编程使用纹理对象（cudaTextureObject_t），更灵活，且不与全局命名空间冲突。主机端流程
分配全局显存，拷贝原始数据（和普通全局内存完全一致）
填充资源描述符：告诉硬件数据在哪、是什么格式、几维
填充纹理描述符：告诉硬件寻址规则、是否插值、越界怎么处理
创建纹理对象，拿到可传入核函数的句柄
将纹理对象作为参数传入核函数
核函数端流程
调用采样函数（
tex1Dfetch/tex2D等）读取数据收尾
核函数执行完成后，销毁纹理对象，释放显存
// 核函数：纹理对象作为参数传入 __global__ void texture_1d_kernel(cudaTextureObject_t tex, float* output, int len) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx >= len) return; // 一维纹理采样：tex1Dfetch(纹理对象, 整数下标) // 用法和数组几乎完全一样，底层走纹理缓存流水线 float val = tex1Dfetch<float>(tex, idx); output[idx] = val * 2.0f; } int main() { const int len = 1024; const int bytes = len * sizeof(float); // ========== 第1步：分配全局显存，准备数据 ========== float* h_input = new float[len]; for (int i = 0; i < len; i++) h_input[i] = (float)i; float* d_input = nullptr; float* d_output = nullptr; CHECK(cudaMalloc(&d_input, bytes)); CHECK(cudaMalloc(&d_output, bytes)); CHECK(cudaMemcpy(d_input, h_input, bytes, cudaMemcpyHostToDevice)); // ========== 第2步：填充资源描述符 ========== cudaResourceDesc res_desc{}; res_desc.resType = cudaResourceTypeLinear; // 线性内存（一维） res_desc.res.linear.devPtr = d_input; // 指向全局显存数据 res_desc.res.linear.desc = cudaCreateChannelDesc<float>(); // 数据格式：float res_desc.res.linear.sizeInBytes = bytes; // 数据总字节数 // ========== 第3步：填充纹理描述符 ========== cudaTextureDesc tex_desc{}; tex_desc.normalizedCoords = 0; // 非归一化坐标：直接用整数下标 tex_desc.filterMode = cudaFilterModePoint; // 最近邻采样：无插值，原样返回 tex_desc.addressMode[0] = cudaAddressModeClamp; // 越界钳位：超出范围返回边缘值 // ========== 第4步：创建纹理对象 ========== cudaTextureObject_t tex_obj = 0; CHECK(cudaCreateTextureObject(&tex_obj, &res_desc, &tex_desc, nullptr)); // ========== 第5步：启动核函数，传入纹理对象 ========== dim3 block(256); dim3 grid((len + block.x - 1) / block.x); texture_1d_kernel<<<grid, block>>>(tex_obj, d_output, len); CHECK(cudaGetLastError()); CHECK(cudaDeviceSynchronize()); // ========== 验证结果 ========== float* h_output = new float[len]; CHECK(cudaMemcpy(h_output, d_output, bytes, cudaMemcpyDeviceToHost)); cout << "前5个结果："; for (int i = 0; i < 5; i++) cout << h_output[i] << " "; cout << endl; // ========== 第7步：销毁资源 ========== CHECK(cudaDestroyTextureObject(tex_obj)); CHECK(cudaFree(d_input)); CHECK(cudaFree(d_output)); delete[] h_input; delete[] h_output; return 0; }关于纹理内存有太多的特性了，好多博主都不是很清楚，等博主了解完毕之后再详细的出一章节来细讲
现在的大致了解是，可以利用一些特性去解决图像非连续空间的访问，可以加速，而且特性似乎非常厉害，这里需要详细讲，本章先简单的理解有个纹理缓存即可
全局内存
全局内存是 GPU 显存的核心主体与最终载体，是 GPU 的 “主存”，对应 CPU 侧的内存条。
全局内存容量最大、寻址最灵活、适用范围最广，但原生延迟也最高；几乎所有 CUDA 程序的性能瓶颈，最终都会落到全局内存的访问效率上。
1.物理载体与容量
物理位置：GPU 板载显存（GDDR6X / HBM 等显存颗粒），和 CPU 主存是同级概念，不在 SM 芯片内部。
总容量：就是你显卡的显存大小（如 8GB、12GB、24GB），远大于所有片上存储的总和。
底层共性：本地内存、常量内存、纹理内存的数据本体都存放在这里，只是访问路径、缓存策略、硬件流水线不同。
2. 访问权限与生命周期
访问范围：全局可见。主机端（CPU）、设备端（GPU）所有线程、所有核函数都可以读写，是唯一支持跨设备、跨核函数传递数据的存储层级。
管理方式：完全由开发者手动管理，通过
cudaMalloc分配、cudaFree释放，生命周期和进程一致，不会随核函数结束而自动销毁。寻址方式：线性字节地址空间，和 C++ 普通堆内存逻辑完全一致，支持指针运算、随机寻址。
3. 和其他内存的物理从属关系
板载显存 = 全局内存地址空间 + 常量内存地址空间 + 本地内存分配 + 纹理内存绑定的数据区 它们共用同一块物理显存，差异只在「访问路径、缓存策略、硬件加速单元」。
物理特性
高带宽，高延迟
全局内存本身的物理延迟无法改变，优化永远围绕两个方向：
- 减少访问次数：通过片上缓存、共享内存，让数据尽量留在高速区域，少去显存。
- 提高访问效率：通过访存合并，把零散请求拼成大批次传输，榨满显存带宽。
如何声明
动态分配：cudaMalloc，这里就不举例子了
静态分配：__device__
__device__ int counter = 0; // 设备端全局变量，所有核函数都能访问 __global__ void increment() { atomicAdd(&counter, 1); // 所有线程都能修改它 }CPU 不能直接读写它，必须通过
cudaMemcpyToSymbol/cudaMemcpyFromSymbol来拷贝数据。不能过多使用：因为它占用宝贵的全局内存，且过多使用会降低代码可维护性。通常用于需要跨核函数保持状态的场景（如全局计数器、标志位）。
总结
内存类型 物理位置 作用范围与访问权限 声明方式与生命周期 最佳访问模式与优化要点 寄存器 片上 SRAM 线程私有，读/写 核函数内局部变量（编译器自动分配）；线程生命周期 速度最快。避免使用过多变量导致寄存器溢出到本地内存。 共享内存 片上 SRAM Block 内共享，读/写 __shared__静态声明 或extern __shared__动态声明；Block 生命周期避免 Bank Conflict：同一 Warp 内线程尽量访问不同 Bank 的地址。 L1 / L2 缓存 片上 SRAM 硬件自动管理，对所有线程透明 无法显式声明（L1与共享内存共用，可调整比例） 无法直接控制，但可通过合并访问、数据复用提高命中率。 常量内存 片外 DRAM 全局只读，通过专用缓存加速 __constant__全局声明，cudaMemcpyToSymbol()拷贝；应用生命周期一个 Warp 内所有线程访问同一地址（广播读），否则串行化。 纹理内存 片外 DRAM 全局只读，通过专用缓存加速 cudaTextureObject_t创建绑定；纹理对象生命周期适合二维空间局部性（图像处理），支持硬件插值与自动边界处理。 全局内存 片外 DRAM 所有线程，读/写 cudaMalloc()动态分配；由程序员控制生命周期必须满足合并访问：同一 Warp 线程访问连续、对齐的 128 字节地址块。 本地内存 片外 DRAM 线程私有，读/写 编译器自动分配（寄存器溢出/大数组）；线程生命周期 性能杀手。应通过优化代码、减少寄存器压力来极力避免使用。 __device__变量片外 DRAM 全局，同一编译单元内所有核函数可读/写 __device__全局声明；应用生命周期作用类似静态全局变量，可用于跨核函数状态保持。
cuda是尽量暴露更多的内存模型给用户使用，通过学习各种内存的特性，我们才能更好的去优化我们的程序，下一节我们将通过优化深入的感受内存的特性
扫一扫
1万+
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
