想象你是一家米其林餐厅的后厨主管。后厨就是GPU，每个厨师就是一个线程。灶台边上那只冷藏总仓库，就是“全局内存”——食材全在那儿，可每走一趟要穿过整条走廊、拉开三扇门，来回得花上几分钟。而每位厨师手边的共用料理台，就是“共享内存”：肉、菜、酱料提前搬上台后，谁要谁伸手，一秒钟都不耽误。 料理台是整间后厨效率的命脉，可它也有脾气——台面只有那么大，几十个厨师同时抢同一个抽屉就得排队；洗碗工还没把上一批盘子清走，大家就得并肩干瞪眼等同步。谁摸清了料理台的规矩，谁就能把出菜速度提升上百倍。这篇文章就用这间后厨做地图，带你走一遍CUDA共享内存的每一条通路：它藏在芯片哪里、和全局内存差多少、为何会“抢抽屉”（bank conflict）、怎么用齐步哨（__syncthreads）对齐节奏、如何用“分批备菜”（tiled算法）让矩阵乘法起飞，以及声明、分配的种种门道。

岗位关联

学会共享内存，是GPU性能优化从“能跑”走向“跑得快”的分水岭，也是面试官最爱的考点。掌握它，你就能把深度学习或科学计算内核的耗时压掉一大截。

目录
  1. 可视化演示
  2. 共享内存概述
  3. 内存层次对比
  4. Bank冲突机制
  5. 线程同步机制
  6. Tiled分块算法
  7. 声明与分配
  8. 优化实践技巧
  9. AI 应用场景
  10. Python 演示
  11. 常见误区
  12. 面试题

可视化演示

展示CUDA共享内存中多线程访问Bank的两种情况，无冲突时并行访问高速，bank冲突时串行化导致延迟增加。

本节配有交互动效（进入视口后自动加载，约 3KB）。

无冲突模式 每个线程访问不同Bank，全并行执行 共享内存延迟仅1cycle 并行度 = 线程数

共享内存概述

后厨最贵的不是灶台，而是那块料理台。它贴着灶台砌成，食材搬上去后厨师一伸手就够得着；可台面就那么大，放满一筐土豆就放不下十斤牛肉。CUDA里的共享内存正是GPU这块“料理台”：它物理上建在流多处理器（SM，GPU里的一个计算单元组，相当于一间后厨）芯片内部，和线程块里的所有线程“共用”。

共享内存属于“片上内存”（on-chip），离计算单元只有一臂距离。全局内存则要穿过总线去显存里取，一趟来回要几百纳秒；共享内存十几个时钟周期就能完成，速度差了两个数量级。需要反复读的数据，放进共享内存后，每次访问都像从口袋里掏东西一样快。

它有三个硬特性：其一，位于SM内部的SRAM中；其二，对所属线程块内所有线程可见，块与块之间完全隔离——A块厨师够不到B块的料理台；其三，生命周期跟着线程块走，块结束内存即释放。最要紧的是容量极小，默认每块约48KB，放不下几张大矩阵。

后厨的直觉告诉我们：凡是会被反复使用的食材，一定要提前搬上料理台。GPU计算也一样——只要访问模式允许，把数据从全局内存“运”进共享内存，让线程们围着它干活，省下的时间就是几百倍。

TglobalTshared≈400 ns1.5 ns≈267\frac{T_{global}}{T_{shared}} \approx \frac{400\,\text{ns}}{1.5\,\text{ns}} \approx 267
全局内存访问比共享内存慢约267倍，这就是为什么“能用料理台就别跑仓库”。
  1. 全局访存延迟线程读全局内存平均要等约400纳秒，相当于厨师跑一趟马路对面的总仓库。
    Tglobal≈400 nsT_{global} \approx 400\,\text{ns}
  2. 共享访存延迟读共享内存只要约1.5纳秒，相当于手伸向旁边的料理台。
    Tshared≈1.5 nsT_{shared} \approx 1.5\,\text{ns}
  3. 求比值两者相除得到约267倍，量化了“料理台”相对“仓库”的速度优势。
    TglobalTshared≈267\frac{T_{global}}{T_{shared}} \approx 267

内存层次对比

厨房里不止有仓库和料理台，每位厨师口袋里还别着一套随身刀具——那叫寄存器（register）。整个后厨的存储体系从近到远分四层：口袋（寄存器）→ 料理台（共享内存）→ 过道搁板（L2缓存）→ 马路对面总仓库（全局内存）。越靠手边越快，越往外走容量越大、速度也越慢。

从CUDA视角看，这四层的“可见范围”各不相同：寄存器和局部内存只属于单个线程私有；共享内存属于线程块；L2和全局内存所有线程共享。还有一级L1缓存，它和共享内存住在同一个SM上，在新款GPU里两者共享同一片物理硬件，可以通过配置按比例拆分。

延迟和带宽的数据直接决定策略：寄存器1个时钟周期；共享内存约1.5纳秒；L2要几十纳秒；全局内存约400纳秒。更惊人的是聚合带宽——一个SM内共享内存总带宽可达约19TB/s，而全局内存约0.8TB/s，差了二十多倍。算法需要高频重读数据时，用共享内存“以空间换延迟”非常划算。

所以后厨的老规矩是：备菜切墩全在料理台上完成，只有“进货”和“上新菜”才用仓库。CUDA优化本质上就是一个搬运问题——把高频访问的数据从慢层挪到快层，让稀缺的全局带宽留给真正绕不开的数据流。

BsharedBglobal≈19 TB/s0.8 TB/s≈24\frac{B_{shared}}{B_{global}} \approx \frac{19\,\text{TB/s}}{0.8\,\text{TB/s}} \approx 24
共享内存的聚合带宽约为全局内存的24倍，频率越高、复用越多的数据越应该往料理台上放。
  1. 共享内存聚合带宽所有bank并行工作时的总吞吐约19TB/s，如同32个抽屉同时出菜。
    Bshared≈19 TB/sB_{shared} \approx 19\,\text{TB/s}
  2. 全局内存带宽显存接口能提供的总带宽约0.8TB/s，如同仓库门口只有一条窄传送带。
    Bglobal≈0.8 TB/sB_{global} \approx 0.8\,\text{TB/s}
  3. 带宽之比两者相除约等于24，说明料理台传送食材的吞吐能力远胜仓库。
    BsharedBglobal≈24\frac{B_{shared}}{B_{global}} \approx 24

Bank冲突机制

料理台上有32个小抽屉（bank），每个抽屉同一时刻只能服务一个厨师。硬件规定抽屉号 = （地址字节数 ÷ 4） 对32取余。如果两个厨师恰好同时把手伸向同一个抽屉，后面那个只能等着——这就是“抽屉打架”，CUDA称之为 bank conflict。

一个线程束（warp，GPU调度的最小单位，32个线程一伍）执行一条共享内存访问指令时，硬件会把它拆成若干“事务”（transaction）。如果32个线程访问的地址落在32个不同抽屉，一次事务全部办完；如果全都挤进同一个抽屉，就要拆成32次串行事务，慢32倍。冲突越严重，事务拆得越多，等效带宽越低。

典型冲突场景：声明 tile[32][32]，按列访问 tile[tid][0] 时行距正好32个整数，每个元素都映射到同一个抽屉，产生32路冲突。破解方法是“填充”（padding）：把数组声明成 tile[32][33]，让列间距变成33，抽屉号自然错开。

抢抽屉是个概率问题，更是个设计问题。把数组弄“歪”一格（padding），让每位厨师各开一屉，后勤效率立刻拉满——这是共享内存优化的第一把金钥匙。

Taccess=w×Tbank,w=max⁡b∈[0,32)#{t:bank(t)=b}T_{access} = w \times T_{bank}, \qquad w = \max_{b \in [0,32)} \#\{t : \text{bank}(t) = b\}
一次共享内存访问的耗时等于“最挤的那个抽屉挤了几个线程”乘以单次访问时间；全部线程都挤同一个抽屉时 w=32，慢了32倍。
  1. 访问总耗时T_{access} = w \times T_{bank}，表示整个warp完成一次访问要串行排队w轮。
    Taccess=w×TbankT_{access} = w \times T_{bank}
  2. 冲突路数ww 等于所有bank中“被最多线程同时命中”的那个抽屉的人数。
    w=max⁡b∈[0,32)#{t:bank(t)=b}w = \max_{b \in [0,32)} \#\{t : \text{bank}(t) = b\}
  3. bank映射函数\text{bank}(t) = \lfloor \text{addr}(t)/4 \rfloor \bmod 32，由地址除以4再取余32得到抽屉号。
    bank(t)=⌊addr(t)/4⌋ mod 32\text{bank}(t) = \lfloor \text{addr}(t)/4 \rfloor \bmod 32
  4. 意义w=1 时无冲突，w=32 时是最坏情况，硬件被迫把请求拆成32个事务。
    1≤w≤321 \le w \le 32

线程同步机制

厨师们围着料理台备菜，最怕什么？怕你菜还没切完，旁边的人已经开火把你的半成品下锅了。后厨的办法是“齐步走”：厨师长吹一声哨（__syncthreads()），所有厨师必须停下手里的活，等全厨房的人都到齐，才允许继续做后面的事。

__syncthreads() 是线程块内的“栅栏”：它保证被它隔开的两段代码有严格先后顺序，前面所有线程对共享内存的写操作，在哨声响过后对块内所有线程可见。没有这道栅栏，线程A还没写完，线程B就读了，结果就是数据竞争（race condition），每次运行结果都可能不同。

吹哨也有代价：每一次调用都要让块内最慢的线程拖住所有人，频繁同步显著降低吞吐。更危险的是，如果同步出现在分支里，只有部分线程执行它，另一部分先跑走了——大家永远等不齐，程序直接死锁（deadlock）。因此__syncthreads()必须保证被块内所有线程执行到同一行。

哨声是安全的保障，也是节奏的枷锁。好的内核像一支默契的乐队，鼓手、吉他手各干各的，只在间奏那一拍对齐；高手会把同步次数压到最少，让每位厨师在两次对齐之间最大限度地自由发挥。

Wi≺S≺Rj(∀i,j∈block)W_i \prec S \prec R_j \qquad (\forall i, j \in \text{block})
任意线程i的写、同步栅栏、任意线程j的读之间构成严格的“先于”关系，这就是可见性的数学保证。
  1. 写事件W_i 表示线程 i 对共享内存某个地址的写入操作。
    WiW_i
  2. 同步栅栏S 表示 __syncthreads() 这个对齐点，所有线程必须都到达才放行。
    SS
  3. 读事件R_j 表示线程 j 稍后对该地址的读取操作。
    RjR_j
  4. 先行关系\prec 表示“发生在……之前”，构成 W→S→R 的链，保证读完的一定是写好的。
    Wi≺S≺RjW_i \prec S \prec R_j

Tiled分块算法

后厨出套餐时最聪明的做法，不是来一桌做一桌，而是提前把下一批客人要的食材按“桌”分好，一趟推车全部搬上料理台，再集中切配。这种“分批发货”的思路，就是CUDA里的 tiled（分块）算法。

矩阵乘法 C = A×B 是最经典的例子。朴素做法里，计算每个输出元素都要读A的一行和B的一列，总共发生 2N³ 次全局访存。分块做法是：把矩阵切成 T×T 的小块，每个线程块负责一块输出，需要的数据先搬进共享内存，再在“料理台”上做乘加运算，算完一个k分块再换下一批。

关键在于“复用”：A的一块元素载入共享内存后，会被块内多个线程反复读取；B同理。全局访存从 2N³ 降到 2N³/T，T越大节省越可观。经典实现里每个线程还配合寄存器做累加，用“双缓冲”流水线式地边加载边算；遇到矩阵边长不是T的整数倍时，还需要边界分支处理越界的读写。

一盘菜切好只给一个人吃太浪费——让它在料理台上被翻来覆去地用上好几遍，进货次数就少了。Tiled算法的核心哲学就是：宁可多搬几趟大盘子，也决不零敲碎打地反复跑仓库。

Rreuse=NnaiveNtiled=2N32N3/T=TR_{reuse} = \frac{N_{naive}}{N_{tiled}} = \frac{2N^3}{2N^3 / T} = T
把矩阵切成T×T的小块后，每个从全局内存搬到共享内存的元素平均被复用了T次，全局访存因此降为原来的1/T。
  1. 朴素做法的访存次数每个输出元素要读A和B各一次，一共 N³ 次乘法、2N³ 次全局读。
    Nnaive=2N3N_{naive} = 2N^3
  2. 分块后的访存次数每块数据只从全局内存搬一次，随后在共享内存里反复使用，总次数变为 2N³/T。
    Ntiled=2N3TN_{tiled} = \frac{2N^3}{T}
  3. 复用倍数两者相除得到 T，这就是分块算法加速的根源——一份食材喂饱T张嘴。
    Rreuse=NnaiveNtiled=TR_{reuse} = \frac{N_{naive}}{N_{tiled}} = T

声明与分配

料理台怎么搭，有两种方式。一种是装修时就在图纸上画好格子（静态共享内存），尺寸写死在代码里；一种是开工前根据当天菜单临时拼桌（动态共享内存），台面大小由内核启动参数临时指定。

静态声明写法：__shared__ float tile[32][33];，在核函数体内写死大小，编译期就分配好，简单直观。动态声明写法：extern __shared__ float smem[];，不带大小，由内核启动时“第三个配置参数”传入字节数；如果需要的动态共享内存超过默认上限，还要先用 cudaFuncSetAttribute 显式放开。

动态内存的大小在运行时才确定，因此常用来适配不同输入规模。注意三点：第一，extern 声明的不定长数组必须放在所有静态共享内存声明的最后；第二，需要多个动态数组时要手动算偏移，让第二个数组从第一个的末尾接上；第三，默认每线程块共享内存上限48KB，超出需显式配置。静态适合大小恒定的场景，动态适合输入规模波动、需要弹性适配的场景。

静态格子省心但死板，动态拼桌灵活但要多算一把尺子。后厨高手两手都会：窗口期固定的菜用静态，菜量随客流波动的用动态——选择的标准永远是“这张台面怎么铺最不浪费”。

offseti=∑j=0i−1sizej\text{offset}_i = \sum_{j=0}^{i-1} \text{size}_j
当你需要把一块动态共享内存切成多个数组时，第i个数组的起始偏移量等于它前面所有数组大小之和。
  1. 偏移量含义\text{offset}_i 表示第 i 个数组在共享内存中的起始位置（单位：字节或元素数）。
    offseti\text{offset}_i
  2. 前缀累加\sum_{j=0}^{i-1} \text{size}_j 把前面 i 个数组的大小全部加起来。
    ∑j=0i−1sizej\sum_{j=0}^{i-1} \text{size}_j
  3. 指针切分得到偏移后，用 float* a = smem; float* b = smem + offset_1; 把一整块内存切成两段使用。
    b=smem+offset1b = \text{smem} + \text{offset}_1

优化实践技巧

把散落的技巧连成一套后厨SOP。第一，抽屉不打架：访问模式尽量让地址连续错开，必要时做padding，或者利用“广播”特性——所有线程读同一个地址时硬件一次搞定，不会冲突。第二，台面别铺太满：共享内存用得太多，能同时进厨房的厨师（线程）就变少，这叫做占用率（occupancy）下降。

占用率的公式是：SM能并发的线程数 = 活动块数 × 每块线程数。而活动块数受三类资源共同限制——共享内存、寄存器、线程上限，三者取最小。例如每块用48KB共享内存，而SM总共只有48KB，那就只能驻留1个块，其余块全在门外排队，隐藏延迟的能力大幅缩水。

具体调优建议：tile大小常取16×16或32×32，在复用率和占用率之间找平衡；同步只在真正有数据依赖的位置加；让每个线程连续访问相邻地址（合并访问），配合共享内存做转置也能减少全局访存；用 cudaFuncSetAttribute 放开大共享内存限制时，要实测性能——大台面挤走厨师，有时反而更慢。

后厨经营的终极问题：台面、人手、菜量三者怎么配比？共享内存优化不是无脑堆料，它是在带宽、容量、延迟与并发之间反复称重。掌握了这台天平，你的内核就能像米其林后厨一样，出菜快、翻台高。

O=Bactive×TblockTmax,Bactive=min⁡(Bsmem,Breg,Blimit)O = \frac{B_{active} \times T_{block}}{T_{max}}, \qquad B_{active} = \min(B_{smem}, B_{reg}, B_{limit})
占用率等于“能同时干活的线程数”除以SM理论最大线程数；决定它的活动块数，是共享内存、寄存器、线程上限这三块短板中的最短者。
  1. 占用率定义O 表示SM上活跃线程占理论容量的比例，越高越能隐藏访存延迟。
    O=Bactive×TblockTmaxO = \frac{B_{active} \times T_{block}}{T_{max}}
  2. 活动块数B_{active} 是当前资源约束下能同时驻留的线程块数。
    Bactive=min⁡(Bsmem,Breg,Blimit)B_{active} = \min(B_{smem}, B_{reg}, B_{limit})
  3. 共享内存限制\lfloor S_{SM} / S_{block} \rfloor 表示把SM的共享内存总量除以每块用量，能得到最多放几个块。
    Bsmem=⌊SSMSblock⌋B_{smem} = \left\lfloor \frac{S_{SM}}{S_{block}} \right\rfloor
  4. 取最小值三个候选块数中取最小的那个，短板决定并发上限，如同木桶效应。
    min⁡(Bsmem,Breg,Blimit)\min(B_{smem}, B_{reg}, B_{limit})

AI 应用场景

LLM推理：Tiled Softmax 加速注意力计算

阶段：模型推理优化 / 内核编写

解决的问题：Softmax 需要先算出整行的最大值和指数和，才能归一化。朴素实现把整行 score 存在全局内存，带宽和显存开销大、延迟高。用 tiled 算法把 score 分块搬进共享内存，边走边维护滚动最大值和累加和，最后一次性归一化，完全省掉全局中间缓冲。

具体案例：在单卡部署 LLaMA-7B 时，Flash-Attention 风格的内核把 QKᵀ 的 score tile 放进共享内存，避免物化整块注意力矩阵。TensorRT-LLM 与 vLLM 中这类内核让首 token 延迟降低 30% 以上，长序列场景收益更明显。

科学计算：GEMM 矩阵乘法内核加速

阶段：HPC 内核开发 / 算子库编写

解决的问题：深度学习训练和科学计算里矩阵乘法是绝对热点。朴素 CUDA 实现每次乘加都访问全局内存，带宽被白白耗尽。用共享内存做 tiled 分块，A、B 的小块提前搬进料理台，复用 T 次后再换新块，全局访存降为 1/T。

具体案例：在 4090 上实现 4096×4096 的 FP16 GEMM，朴素版本约 200 GFLOPS，采用 16×16 tiled + 双缓冲后可达 15 TFLOPS 以上，逼近 cuBLAS 的七成性能。许多私有算子库正是以此为起点迭代优化。

实时视觉：图像卷积的 Halos 处理

阶段：端侧部署 / 视觉算法优化

解决的问题：3×3 或 5×5 卷积中，每个输出像素要读周围邻域，相邻线程的大量读取重叠。把输入图像切成 tile 搬进共享内存，连同边缘一圈“光晕”（halo）一起加载，块内所有线程都在料理台上取数，全局访存大幅减少。

具体案例：在自动驾驶的 1080P 图像预处理链路中，把 5×5 高斯模糊与 Sobel 边缘检测合并成一个共享内存内核，端到端耗时从 4.2ms 降到 1.8ms，为后续检测模型争取了宝贵的帧时间预算。

Python 演示

python56 行
import numpy as np try:    import pycuda.driver as cuda    import pycuda.autoinit    from pycuda.compiler import SourceModuleexcept ImportError:    raise SystemExit("[!] 请先安装 PyCUDA: pip install pycuda") KERNEL = r"""__global__ void block_reduce(float* in, float* out, int n) {    __shared__ float tile[256];    int tid = threadIdx.x;    int gid = blockIdx.x * blockDim.x + tid;     float local_sum = 0.0f;    for (int i = gid; i < n; i += gridDim.x * blockDim.x) {        local_sum += in[i];    }    tile[tid] = local_sum;    __syncthreads();     for (int s = blockDim.x / 2; s > 0; s >>= 1) {        if (tid < s) {            tile[tid] += tile[tid + s];        }        __syncthreads();    }    if (tid == 0) {        out[blockIdx.x] = tile[0];    }}""" def main():    n = 1 << 20    h_in = np.random.randn(n).astype(np.float32)    h_out = np.zeros(256, dtype=np.float32)     d_in = cuda.mem_alloc(h_in.nbytes)    d_out = cuda.mem_alloc(h_out.nbytes)    cuda.memcpy_htod(d_in, h_in)     mod = SourceModule(KERNEL)    block_reduce = mod.get_function("block_reduce")    block_reduce(d_in, d_out, np.int32(n), block=(256, 1, 1), grid=(256, 1))     cuda.memcpy_dtoh(h_out, d_out)    gpu_sum = float(np.sum(h_out))    cpu_sum = float(np.sum(h_in))    print("GPU 归约和:", gpu_sum)    print("CPU 对照和:", cpu_sum)    print("误差小于 1e-2:", abs(gpu_sum - cpu_sum) < 1e-2) if __name__ == "__main__":    main()
目的：用 PyCUDA 实现一个“共享内存块内归约”内核——每个线程先把一段数据的局部和写进共享内存，再用树形合并把 256 个局部和归约为 1 个，最后每个块输出一个值，全程演示了 __syncthreads 与无 bank 冲突的 stride 访问模式。
  1. __shared__ float tile[256];在 GPU 的“料理台”上切出 256 格共享内存，块内线程共用。
  2. tile[tid] = local_sum;每个厨师先把自己那份半成品放上台面。
  3. __syncthreads();哨声对齐，保证所有线程的局部和都放好，才能开始合并。
  4. tile[tid] += tile[tid + s];相邻厨师两两配对累加，stride 模式天然无 bank 冲突。
  5. out[blockIdx.x] = tile[0];每桌领班把最终结果端出料理台，写入全局内存。

常见误区

误区我写内核时故意把共享内存数组开得很大，反正容量用满就不亏，性能肯定更好。

正解共享内存是 SM 上的稀缺资源，用得多会挤压能同时驻留的线程块数，拉低占用率，反而变慢。就像料理台铺满食材，厨房就挤不进更多厨师了。最优策略是“刚好够用”：tile 大小取 16×16 或 32×32，需要大容量时再用 cudaFuncSetAttribute 显式放开，并实测验证收益。

为何常见：优化方向常被表述为“用共享内存减少全局访存”，初学者很容易把“多用”当成“好用”，忽略了资源置换的成本。

误区加了同步以后线程就不会读到我还没写完的数据了，我想在循环里加几次就加几次。

正解它是“全块参与”的硬栅栏——必须块内所有线程都执行到这一行才算通过。如果放在分支里、或两个分支的同步次数不一致，先到的线程会永远等不到没来的线程，直接死锁。它保证共享内存写的可见性，但每次都会让全员等最慢的人，所以只应在真正有数据依赖的位置使用。

为何常见：“等一等”的直觉很柔和，容易让人忽略它其实是一个条件苛刻、代价很高的集合点。

误区只要把数组从全局内存搬到共享内存，性能就自动起飞。

正解搬数据本身有开销——全局到共享要一次拷贝和一次同步；如果数据只被读一次，或者访问模式有严重 bank conflict，共享内存可能比直接读全局还慢。能不能赚，取决于数据复用率 R 是否大于搬运成本。R=1 时搬数据纯属白忙，R=T 时才划算。

为何常见：教程总强调共享内存快两百多倍，却很少强调“搬运成本”和“复用率”这两个前提条件。

误区我的数组才 32×32 这么小，怎么会冲突？大数组才会吧。

正解bank 只有 32 个，任何“地址除以4再取余32”落在同一个抽屉的访问都会冲突，跟数组大小无关。小数组按列访问 tile[tid][0] 时，行距正好 32 个整数，每个元素反而精确命中同一个 bank，形成最典型的 32 路冲突。

为何常见：把 bank 理解成“大容量内存的分页”，以为容量大才容易撞车，实际上它是固定的 32 个抽屉做模运算，跟容量无关。

误区动态共享内存很简单，把 `__shared__ float s[256]` 里的 256 换成变量 N 不就行了？

正解静态共享内存只允许编译期常量大小，写成变量会编译失败。动态必须用 extern __shared__ float s[]; 声明不定长数组，并在内核启动的第三个配置参数里传入字节数；多数组还要手动算偏移。两种声明看起来只差一个 extern，实际机制完全不同。

为何常见：两种声明在代码上只差一个 extern，初学者容易误以为动态只是把常量换成变量，忽略了启动参数这个关键环节。

面试题

Q1. 共享内存在物理上位于 GPU 的哪个位置？

单选题 · easy

  • A.显存中的高速缓存区
  • B.流多处理器片上SRAM← 正确答案
  • C.主板上的系统内存条
  • D.全局内存的L2缓存

答案：B。共享内存像后厨的料理台，建在 SM 芯片内部的 SRAM 上，离计算单元一臂之遥。显存、系统内存、L2 都不对——L2 虽在片上但是另一层缓存，全局内存更是住在显存里。

Q2. 下列哪种数据最适合放进共享内存？

单选题 · easy

  • A.只被单个线程读一次
  • B.被块内线程反复复用← 正确答案
  • C.写入后立即被丢弃
  • D.需要跨块共享的数

答案：B。共享内存的价值在于“复用”，就像一盘备好的菜喂饱多张嘴。只读一次的数据搬上台面纯属白忙；跨块共享的数据根本放不进共享内存，因为它是线程块私有的。

Q3. 32 个线程的 warp 执行一条共享内存访问，若全部命中同一 bank，会被拆成多少个事务？

单选题 · mid

  • A.1 个事务快速完成
  • B.8 个事务串行处理
  • C.32 个事务串行处理← 正确答案
  • D.取决于数组总容量

答案：C。32 个线程全挤同一个抽屉，硬件只能一次服务一个，串行拆成 32 个事务，即 32 路冲突。这与数组容量无关，只与“最挤的抽屉挤了几个人”有关。

Q4. 修复 bank 冲突最常用的“填充”做法是？

单选题 · mid

  • A.每行末尾多留一个元素← 正确答案
  • B.把数组声明为二维指针
  • C.增大整个块的线程数
  • D.减少共享内存分配量

答案：A。填充就是在每行末尾多留一个元素，如 tile[32][33]，让行距不再是 32 的整数倍，抽屉号自然错开。改指针、加线程、减容量都改变不了地址模 32 的映射关系。

Q5. 动态共享内存的正确声明方式是？

单选题 · mid

  • A.静态声明，大小固定
  • B.extern声明无大小← 正确答案
  • C.device声明于全局
  • D.主机端malloc分配

答案：B。动态共享内存必须用 extern __shared__ float s[]; 声明不定长数组，大小在启动内核时通过第三个配置参数传入。__device__ 是全局内存变量，malloc 是主机端分配，都与动态共享内存无关。

Q6. 关于 __syncthreads() 的正确说法有哪些？

多选题 · mid

  • A.只同步线程块内线程← 正确答案
  • B.同步整个GPU所有线程
  • C.分支中使用可能死锁← 正确答案
  • D.保证共享内存写可见← 正确答案

答案：ACD。它是“块内栅栏”，不是全局同步，所以 A 对 B 错；它保证栅栏前的共享内存写在栅栏后对全块可见，D 对；若放进分支，部分线程永远到不了栅栏，全员死锁，C 对。

Q7. 矩阵乘法 tiled 分块能提速的根本原因是？

单选题 · mid

  • A.减少全局内存访存量← 正确答案
  • B.增加每个线程的计算
  • C.让编译器自动向量化
  • D.降低内核启动的开销

答案：A。分块把数据搬进共享内存复用 T 次，全局访存从 2N³ 降到 2N³/T，这才是提速根源。它并没有增加计算量，也不改变启动开销，向量化更不是分块算法的功劳。

Q8. 下列哪些资源会限制 SM 的占用率？

多选题 · hard

  • A.每块的共享内存大小← 正确答案
  • B.每线程的寄存器数量← 正确答案
  • C.线程块的最大线程数← 正确答案
  • D.全局内存总带宽大小

答案：ABC。活动块数 = min（共享内存限、寄存器限、线程数限），三块短板共同决定占用率，如木桶效应。全局内存带宽影响的是访存速度，不决定 SM 上能驻留多少线程。

Q9. warp 中 32 个线程读同一共享内存地址（广播），硬件行为是？

单选题 · hard

  • A.产生32路bank冲突
  • B.一次事务广播完成← 正确答案
  • C.触发未定义的行为
  • D.自动退化成L2访问

答案：B。所有线程读同一个地址时，硬件识别出这是“广播”，把数据一次分发给全部线程，只算一个事务，不产生冲突。这就像厨师长喊一嗓子，全体厨师同时听到了。

Q10. 关于共享内存与 L1 缓存的关系，正确的说法有哪些？

多选题 · hard

  • A.默认两者占据同一片硬件← 正确答案
  • B.运行时可以按比例拆分← 正确答案
  • C.共享内存数据可自动缓存
  • D.两者都位于SM芯片内部← 正确答案

答案：ABD。在 Volta+ 架构上，L1 和共享内存共用 SM 内的同一片 SRAM，可用 cudaFuncSetAttribute 按比例拆分，所以 A、B、D 对。共享内存是用户手动管理的，不会像全局内存那样被 L1 自动缓存，C 错。