共享内存概述_block块共享内存-优快云博客

从本章往后将要讲解

共享内存概述

共享内存的生命周期

从本章往后将要讲解

数据在共享内存中的安排
二维共享内存到线性全局内存的索引转换
解决不同访问模式中的存储体中的冲突
在共享内存中缓存数据以减少对全局内存的访问
使用共享内存避免非合并全局内存的访问
常量缓存和只读缓存之间的差异
线程束洗牌指令编程

明白一个核心就知道为什么共享内存如此的重要

矩阵转置，无论是以行读还是列读，也就是朴素转置，我们都无法避免一个问题，绝对有一边是无法对齐合并的，那针对这种场景，是开发当中占绝大多数的，所以共享内存就是来解决这个场景的，如果都合并，你不使用共享内存也是可以达到很大的带宽，借助l1缓存l2缓存

随机索引查表 / 稀疏访问
图像邻域计算 / 模板卷积（这个的话就是每一行连续，但是上下相隔很远）
结构体数组（AoS）单字段访问
多维数组非首维度遍历

核心点就是无法按照内存存储的顺序进行访问

大量真实算法的访问逻辑，天然就要求跨步、离散地访问全局内存，这种非合并不是代码写得不好，而是算法本身的数学逻辑决定的，无法从根源上消除；
这种场景下，共享内存就是性能优化的核心手段 —— 把「全局内存的非合并低效访问」，转化为「合并读写全局内存 + 共享内存内自由访问」，用极低的共享内存开销，换取全局内存带宽的大幅提升。

共享内存概述

nvidia-ada-gpu-architecture.pdf
┌───────────────────────────────────────────────────────┐
│                      SM (Ada)                         │
├───────────────────────────────────────────────────────┤
│  [寄存器文件] 256 KB (独立)                            │
│                                                       │
│  [L1 / 共享内存] 128 KB (共享，可配置划分)              │
│  │  └── (内部划分：共享内存 + L1 数据缓存)              │
│  │  └── **L1 数据缓存同时承担以下角色：**               │
│  │      1. 普通 L1 缓存 (全局内存、局部内存的读写)      │
│  │      2. 纹理缓存 (纹理请求的缓存，逻辑通路)          │
│  │      3. 只读缓存 (通过 __ldg() 触发的只读缓存通路)   │
│         4. 常量缓存                                   │
│                              │
│  [指令缓存] 32 KB (独立 SRAM)                          │
└───────────────────────────────────────────────────────┘
这个是ai根据白皮书总结出来的，但问了多个ai都有出入，所以具体的大家可以看白皮书

因为白皮书当中没有提到常量纹理和只读，单单提到了一个指令缓存，所以其余的缓存可能已经并入l1缓存了，具体的大家可以先看一下白皮书，因为博主问了多个ai总结都有出入

但是不妨碍接下来的讲解，后续如果博主找到相关资料或者能够验证再说

共享内存的作用

线程块内高速协同通信，支撑并行算法
转化全局内存非合并访问，根治带宽浪费
缓存复用数据，大幅减少全局内存访问次数
缓解寄存器压力，替代寄存器溢出

本章我们研究两个例子：

归约核函数
矩阵转置核函数

共享内存的生命周期

从全局内存过来的全都要经过L2缓存，相比之下l1缓存和共享内存更接近cuda核心，读取效率高于l2，带宽高，延迟低

共享内存的生命周期：随着 Block 分配和释放

共享内存是 Block 独占的，它的生命周期和 Block 绑定。当一个 Block 被全局调度器（GigaThread Engine）分配到某个 SM 上时，SM 的硬件资源管理器就会把 Block 请求的共享内存大小，从 SM 的共享内存池里一次性划拨出来给这个 Block。直到 Block 的所有线程都执行完毕、退出 SM，这块共享内存才会被释放回池子里。

同一个 SM 上可以同时驻留多个 Block，每个 Block 都有自己独立的一份共享内存，彼此互不可见。如果一个 Block 用了 32KB 共享内存，SM 总共有约 100KB，那么这个 SM 最多能同时容纳 3 个这样的 Block。如果 Block 只用了 16KB，SM 就能同时容纳 6 个。这就是共享内存对占用率的限制——每个 Block 用的共享内存越多，SM 能同时容纳的 Block 越少。

共享内存block内共享，block之间是互相独立的，如果单个block的占用共享内存越大，整体的占有率就越少

共享内存的读取分配

1. 静态声明（编译时决定大小）

这是最简单、最常用的方式。在核函数内部，用 __shared__ 关键字直接声明一个固定大小的数组。大小必须是一个编译期常量，不能是变量。
__global__ void kernel_static() {
    __shared__ float tile[32][32];  // 大小固定为 32×32×4 = 4096 字节
    // ...
}
特点：

简单直接，大小在编译时就完全确定。

受限于每个 Block 最大共享内存上限（通过 cudaDeviceProp.sharedMemPerBlock 查询，通常约 48 KB 或 100 KB）。

如果超出硬件限制，编译器会直接报错。

2. 动态声明（运行时决定大小）

当共享内存的大小依赖输入数据或运行环境时，就需要动态声明。它分两步完成：

第一步：在核函数内用 extern __shared__ 声明
__global__ void kernel_dynamic(int *input, int *output, int N) {
    extern __shared__ float cache[];  // 大小由主机端启动时指定
    // ...
}
关键限制：一个核函数只能声明一个动态共享内存数组。如果需要多个数组，必须手动用指针偏移来划分。

第二步：在主机端启动核函数时，通过 <<< >>> 的第三个参数指定大小
int sharedMemSize = block.x * sizeof(float);  // 例如每个线程一个 float
kernel_dynamic<<<grid, block, sharedMemSize>>>(d_input, d_output, N);
<<<grid, block, sharedMemSize>>> 中的第三个参数就是动态共享内存的总字节数。

如果动态声明需要多个数组：

必须手动用指针偏移来划分：
__global__ void kernel_dynamic(int *input, int *output, int N) {
    extern __shared__ float raw_memory[];

    // 手动划分：前 N 个 float 给 A，后 N 个给 B
    float *A = raw_memory;
    float *B = &raw_memory[N];

    // 现在可以像使用独立数组一样使用 A 和 B
    // ...
}
注意：必须确保主机端传入的 sharedMemSize 等于你手动划分的所有数组的总字节数。

动态共享内存的典型应用场景：

处理变长数据：输入数据量不确定，需要根据实际大小分配共享内存。

提高通用性：同一个核函数可以处理不同规模的 tile，而不必为每个 tile 大小写一个版本。

节省编译时间：避免为不同大小生成多个核函数实例。

设备限制未知：在运行时根据 GPU 的实际共享内存上限（cudaDeviceProp.sharedMemPerBlock）动态决定大小。

3. 调整 L1 缓存与共享内存的硬件划分

共享内存和 L1 缓存共用同一块片上 SRAM，你可以通过 API 改变两者的大小比例，以获得更多共享内存容量。这不是分配共享内存本身，而是扩大共享内存的总容量上限。

方式一：针对单个核函数设置
cudaFuncSetCacheConfig(my_kernel, cudaFuncCachePreferShared);
// 在此之后启动 my_kernel，就会采用优先共享内存的划分方案
my_kernel<<<grid, block>>>(...);
可选配置：

配置选项

含义

cudaFuncCachePreferNone

默认配置，不主动偏好

cudaFuncCachePreferShared

偏好共享内存，让共享内存尽可能大，L1 尽可能小

cudaFuncCachePreferL1

偏好 L1 缓存，让 L1 尽可能大，共享内存尽可能小

cudaFuncCachePreferEqual

均分，两者各占一半

方式二：针对整个设备设置
cudaDeviceSetCacheConfig(cudaFuncCachePreferShared); 
// 影响之后所有核函数
//如何查询调整后的实际容量？
//用 cudaDeviceGetAttribute 在运行时查询：
int sharedSize, l1Size;
cudaDeviceGetAttribute(&sharedSize, cudaDevAttrMaxSharedMemoryPerBlockOptin, 0);
cudaDeviceGetAttribute(&l1Size, cudaDevAttrMaxL1CacheSize, 0);
printf("最大可选共享内存: %d bytes\n", sharedSize); // 如 102400 (100KB)
printf("L1 缓存大小: %d bytes\n", l1Size);        // 如 28672 (28KB)
sharedSize 就是你设置为 PreferShared 后每个 Block 可用的最大共享内存（约 100KB）。

这个值也等于 cudaDeviceProp.sharedMemPerBlockOptin。

总结：三种“分配”的区别与联系

层面

方法

决定时机

灵活性

关键语法/API

声明共享内存

静态声明

编译时

低，固定大小

__shared__ int arr[1024];

指定共享内存

动态声明

运行时

高，大小可变

extern __shared__ float arr[]; + <<<grid, block, bytes>>>

调整容量上限

L1/共享划分

运行时

控制总量

cudaFuncSetCacheConfig / cudaDeviceSetCacheConfig

根据你的需求选择合适的方式。如果需要共享内存大小依赖输入数据，就用动态声明；如果算法固定，静态声明更简洁。如果应用对共享内存需求极大，再考虑调整划分比例。

配置选项	含义
`cudaFuncCachePreferNone`	默认配置，不主动偏好
`cudaFuncCachePreferShared`	偏好共享内存，让共享内存尽可能大，L1 尽可能小
`cudaFuncCachePreferL1`	偏好 L1 缓存，让 L1 尽可能大，共享内存尽可能小
`cudaFuncCachePreferEqual`	均分，两者各占一半

层面	方法	决定时机	灵活性	关键语法/API
声明共享内存	静态声明	编译时	低，固定大小	`__shared__ int arr[1024];`
指定共享内存	动态声明	运行时	高，大小可变	`extern __shared__ float arr[];` + `<<<grid, block, bytes>>>`
调整容量上限	L1/共享划分	运行时	控制总量	`cudaFuncSetCacheConfig` / `cudaDeviceSetCacheConfig`

共享内存存储体

了解完分配要了解其特性，对于全局内存的带宽和延迟，如何采用共享内存去优化，学会优化的前提是理解其物理特性，之前就简单的提到过，但是这里将详细的从头到尾介绍

共享内存物理上被分为 32 个宽度为 4 字节的 Bank。

也就是你可以看成32个通道，为什么是32？这个和warp是对应的32线程

如果你的线程访问共享内存，如果32线程都访问不同的bank，那就不会冲突，一个事务就可以，如果多个线程访问同一个 Bank 的不同地址，就发生 Bank Conflict，请求被串行化，导致多个事务才能拿到数据，这样带宽利用率低

注意：32个bank通道是SM级别的，也就是所有的warp共享32个，只有32个

L1缓存不是按照bank的设计，而是按照我们之前所讲的sector+cacheLine设计的

为什么不是 Bank 结构？ 因为 L1 缓存服务的请求来源更复杂——不同的 Warp、不同的地址、不同的时间。它的目标是提高命中率、减少延迟，而不是支持 32 个线程同时无冲突地访问不同地址。用 Bank 结构反而会增加硬件复杂度和功耗。

那 L1 缓存内部怎么处理并行请求？ L1 缓存控制器内部有多条独立的流水线通道，可以同时处理来自不同 LSU 的请求。这和共享内存的 Bank 不是同一个概念——Bank 是存储本身的并行访问能力，而 L1 的流水线并行是请求处理能力。

设想一下场景：对于之前的转置，如果你的共享内存还设计成L1缓存那样，那合并不了，带宽就很小，为什么采用bank？核心就是因为我压根不需要合并，我只要我的数据落在不同的bank，我就可以一个周期访问数据，那带宽就得到极大的提升，这就是为什么要设计出来bank。如果没有 Bank 结构，共享内存内部同一时间只能处理一个请求。那对于访问，这 32 个线程的请求就会全部串行，共享内存的优势就荡然无存了。Bank 结构，就是为了在“无法合并”的情况下，仍然提供最大并行能力。

简单说，Bank 和合并访问是相反的策略。合并访问是“把多个请求捏成一个”，Bank 是“提供多个通道同时处理多个请求”。

注意：

线程束访问共享内存的时候有下面3种经典模式：

并行访问，多地址访问多存储体
串行访问，多地址访问同一存储体
广播访问，单一地址读取单一存储体（访问同一个地址是可以广播的，这不会冲突）

冲突是指：一个bank里面是有多个不同的地址的，如果你访问的是一个bank里面的不同地址就会冲突，如果是一个bank里面相同的地址就可以广播

以下说的就是访问同一个bank里面的不同地址导致的bank conflict

流程：

Warp Scheduler 发射 LDS（Load Shared）或 STS（Store Shared）指令给 LSU。

LSU 收集 32 个线程的地址，但不进行合并（因为共享内存不走 L1/L2 缓存，是直接访问 SRAM）。

LSU 通过交叉开关（Crossbar）直接访问共享内存的 32 个 Bank。

如果发生 Bank Conflict，请求被串行化；如果没有冲突，32 个请求在一个时钟周期内同时完成。

LSU 把读到的数据写入各线程的寄存器（Load），或把寄存器数据写入共享内存（Store）。

Warp Scheduler：厨师长，负责指挥：“配菜员，去仓库拿 32 份胡萝卜！”

LSU：专职配菜员，负责执行命令：跑去仓库（缓存/DRAM），把胡萝卜拿回来，分给 32 个厨师（CUDA 核心）。

共享内存：厨房里的公共调料架，配菜员可以直接从架子上取调料（直接访问 SRAM），不需要跑仓库。
Warp调度器0     Warp调度器1     Warp调度器2     Warp调度器3
     │              │              │              │
     ▼              ▼              ▼              ▼
  warp0的         warp1的         warp2的         warp3的
  smem指令        smem指令        smem指令        smem指令
     │              │              │              │
     └──────────────┴──────────────┴──────────────┘
                          │
                          ▼
                    LSU指令队列
                    （串行出队）
                          │
                  ┌───────▼────────┐
                  │    仲裁器       │
                  │ 分析32个地址    │
                  │ 落在哪些bank    │
                  └───────┬────────┘
                          │
              ┌───────────┴───────────┐
              ▼                       ▼
       无冲突或广播              有n路冲突
       1个周期完成              n个周期完成
       bank conflict+0         bank conflict+(n-1)
              │                       │
              └───────────┬───────────┘
                          ▼
                    32个Bank的SRAM
                    （实际读写操作）
理想情况（无冲突）：
LSU0的warp请求bank0,1,2...7
LSU1的warp请求bank8,9,10...15
LSU2的warp请求bank16,17...23
LSU3的warp请求bank24,25...31

→ 四组请求访问不同bank
→ Crossbar同时建立四条独立连接
→ 一个周期内完成所有四个warp的访问
→ 四个LSU真正并行工作

仲裁器位于Crossbar的入口
每个时钟周期做两件事：

第一件事：处理同一LSU内部的bank conflict
    LSU0发来32个线程的地址
    仲裁器分析这32个地址落在哪些bank
    如果有冲突 → 拆分成多个轮次串行处理
    如果无冲突 → 一个周期完成

第二件事：处理不同LSU之间的bank竞争
    LSU0的warp要访问bank5
    LSU1的warp也要访问bank5
    → 同一个bank同一时刻只能服务一个请求
    → 仲裁器决定谁先访问
    → 另一个等待下一个周期
其实本质要管好一个warp的32线程即可，让这32个线程在访问共享内存的时候尽量在不同的bank，或者说如果冲突只要少量也可以，不要全部都冲突，那效率极差

冲突的根本原因：访问同一个bank的不同地址

所以我们要调整数据、代码、算法去规避冲突，提升性能

共享内存访问模式

架构	Bank宽度	说明
Fermi (2.x)	4字节	老架构
Kepler (3.x)	8字节	历史数据，已过时
Maxwell及以后	4字节	现代架构统一标准

#include <cuda_runtime.h>
#include "../freshman.hpp"
using namespace std;
int main()
{
    int device = 0;
    cudaSetDevice(device);
    cudaSharedMemConfig config;
    CHECK(cudaDeviceGetSharedMemConfig(&config));
    cout << "SharedMem Config value: " << config << endl;
    if (config == cudaSharedMemBankSizeFourByte){
        cout << "Mode: 4 Byte per Bank" << endl;
    }
    else if (config == cudaSharedMemBankSizeEightByte){
        cout << "Mode: 8 Byte per Bank" << endl;
    }
    else{
        cout << "cudaSharedMemBankSizeDefault" << endl;
    }
    cudaSharedMemConfig rconfig=cudaSharedMemBankSizeEightByte;
    cudaDeviceSetSharedMemConfig(rconfig);
    CHECK(cudaDeviceGetSharedMemConfig(&config));
    cout << "SharedMem Config value: " << config << endl;
    if (config == cudaSharedMemBankSizeFourByte){
        cout << "Mode: 4 Byte per Bank" << endl;
    }
    else if (config == cudaSharedMemBankSizeEightByte){
        cout << "Mode: 8 Byte per Bank" << endl;
    }
    else{
        cout << "cudaSharedMemBankSizeDefault" << endl;
    }
    return 0;
}

bank大小仍然为四字节，说明已经被锁定了

存储索引按照以下方式进行计算

Bank_Index = (Byte_Address / 4) % 32

内存填充

padding填充：https://mp.youkuaiyun.com/mp_blog/creation/editor/162274281

在这篇文章里面我已经讲的很详细，并且还有实验

简单来说就是多开一列，然后让数据映射关系不是整齐的，也就是同一列不是存储在同一个bank当中，当然本身一行也不是，为的就是访问一列的时候让其不产生bank conflict

同步

一个block被调度到SM上，然后又被划分为warp被warp schedule调度，一个warp通常是32个线程，就算不满也是按照32个来执行

那么就注定了一个block内的线程是不同步的，但是一个warp内的32线程是同步的（注意线程束分化这种情况，也就是if-else，这种会破坏天然的同步性，特性情况）

方法：

屏障同步（Barrier / 障碍）：既阻塞执行、又保证内存可见性
内存栅栏（Memory Fence）：只保证内存操作顺序可见性，不阻塞线程执行

弱排序内存模型

什么叫弱排序，就是内存访问不一定严格按照代码书写顺序执行和可见。

对于多线程并发的情况，如果线程A先写入x，再写y，但是在线程B看到的是y，此次才是x，这种情况就是弱排序，不保证顺序

1. 写操作：不是乱序执行，而是编译器可能缓存数据

对于同一个线程内，有数据依赖的内存访问，硬件和编译器保证顺序。对于没有依赖的访问，编译器可能把数据临时缓存在寄存器里，直到“合适的时候”才真正写入内存。

最常见的情况是编译器优化：你写了两行赋值，先写 a[0] 再写 b[0]。编译器发现这两者没有依赖，可能把 b[0] 的结果先写回显存，而把 a[0] 的值暂留在寄存器里稍后再写。不是硬件写操作乱序了，而是编译器改变了实际物理写入的时间点。

2. 读操作：指令级并行可能导致重排

假设你写了两条读指令，分别读 a[0] 和 b[0]。这两者没有依赖，硬件可以同时发射这两条指令，或者后发射的先返回数据。这是为了填满流水线、提高吞吐量，是性能优化，不是Bug。 这和CPU的乱序执行是一个道理。

3. 跨线程可见性：核心在于同步点，而不是执行顺序

在没有同步点的情况下，线程 A 对全局内存的写入，什么时候、以什么顺序被线程 B 看到，是无法保证的。因为数据可能还在 A 的 L1 缓存里，还没写回 L2 或显存。__syncthreads 不仅是一个障碍同步，它还包含一个内存栅栏，强制所有之前的写入对 Block 内其他线程可见。这就是为什么在共享内存的写入和读出之间必须加 __syncthreads——没有它，即使你的代码写入了数据，其他线程也读不到。

为了性能，这是根本原因：

如果要保证强一致，每次写入都得立刻刷新到所有缓存，带宽开销极大；
弱排序让硬件可以自由地批量刷新缓存、乱序执行，最大化吞吐量；
GPU 的设计哲学是 "吞吐量优先，一致性按需保证"—— 平时放开跑，需要顺序的时候用同步指令自己建。

为了可控，必须使用同步技术

显式屏障/障碍

CUDA中，障碍点设置在核函数中，注意这个指令只能在核函数中调用，并只对同一线程块内线程有效。
void __syncthreads();
__syncthreads()作为一个障碍点，他保证在同一线程块内所有线程没到达此障碍点时，不能继续向下执行。
同一线程块内此障碍点之前的所有全局内存，共享内存操作，对后面的线程都是可见的。
这个也就能解决同一线程块内，内存竞争的问题，同步，保证先后顺序，不会混乱。
避免死锁情况出现，比如下面这种情况，就会导致内核死锁：
if (threadID % 2 == 0) {
    __syncthreads();
} else {
    __syncthreads();
}
//因为偶数线程的__syncthreads()导致一直等待奇数线程，连if都出不去，死锁了
只能解决一个块内的线程同步，想做块之间的，只能通过核函数的执行和结束来进行块之间的同步。（把要同步的地方作为核函数的结束，来隐式的同步线程块）

内存栅栏

内存栅栏能保证栅栏前的内核内存写操作对栅栏后的其他线程都是可见的，有以下三种栅栏：块，设备，系统。

线程块内：保证同一block中的其他线程对于栅栏前的内存写操作可见
void __threadfence_block();
设备：保证整个SM的所有线程对于栅栏前的内存写操作可见
void __threadfence();
系统：保证整个系统可见
void __threadfence_system();
栅栏的作用就是保证栅栏之前的所有写入内存的操作不会停留在寄存器，会刷新到对应的地方让其他可见，解决的是可见性问题

假设这个是warp0
data[0] = 1;      // 写入 A
__threadfence();  // 栅栏：强制 A 在 B 之前完成
ready = 1;        // 写入 B
这个是warp1
if (ready == 1) {
    use(data[0]);  // 希望读到 Warp 0 写入的 1
}
那么栅栏的作用是啥？

是warp能够读取的时候，data肯定写到了全局内存当中，一定能读到新值，所以栅栏的作用就是保证你在ready读取的时候，data一定是新值，就是类似一个栅栏，不允许重排

所以，栅栏的价值在于：它消除了“看到 ready=1 却读不到 data[0]=1”这种不一致的情况。

要保证所有的线程同一时间看到最新值，这个是__syncthreads()去做的

栅栏：只保证“我写入的东西，在某个时间点之后对别人可见”，但不能保证“别人什么时候去看”。

障碍：强制所有线程在同一个点集合，确保所有人都完成了写入，然后所有人再一起去看。