CUDA C编程权威指南-第五章：共享内存和常量内存

系列导航：导读 | 上一篇：第4章 全局内存 | 下一篇：第6章 流和并发

系列第 5 篇。第 4 章讲了全局内存的对齐与合并、矩阵转置里的读写取舍；这一章把数据「搬到片上」：共享内存的布局与 bank 冲突、用共享内存做矩阵转置与归约、常量内存和线程束洗牌怎么用，把「怎么访问全局内存」和「怎么在片上复用、重组数据」串成一条线。

---

前言：从全局内存到片上内存

第4章指出：当算法不得不以非合并方式访问全局内存时（例如矩阵转置中按列读或按列写），带宽利用率会下降。需要留意的是：在许多情况下，可以通过共享内存来改善全局内存的合并访问——先把数据以合并方式从全局内存加载到共享内存，在共享内存中按需重组或复用，再以合并方式写回全局内存。这样既减少了全局内存事务数，又利用了片上存储的低延迟与高带宽。

此外，共享内存与 L1 缓存在 Fermi 等架构上共享同一块片上存储（如 64 KB），程序员可通过运行时 API 调整二者比例；常量内存则提供只读、带专用缓存的存储，当线程束内所有线程访问同一地址时能获得广播优势。本章将系统覆盖：在共享内存中数据如何排布、什么是 bank 与 bank 冲突、如何用 padding 消除冲突、二维共享内存与线性全局内存的索引转换、矩阵转置与并行归约的共享内存实现、常量内存与只读缓存的区别，以及线程束洗牌指令（warp shuffle）的入门用法。

---

一、本章在全书中的位置与学习目标

1.1 为什么要学「共享内存和常量内存」

第 4 章从全局内存回答了「如何对齐与合并、为何矩阵转置难以双端合并」；但性能瓶颈的突破往往在于把热数据搬到片上、按需复用。这一章会说明：大多数 GPU 应用程序受内存带宽限制，而共享内存是 SM 上的片上可编程存储器，延迟远低于全局内存、带宽更高；在许多情况下，可以通过共享内存来改善全局内存的合并访问。因此，理解共享内存的布局与 bank 冲突、掌握用共享内存做 tile 与归约、区分常量内存与只读缓存的适用场景，并会用线程束洗牌在 warp 内无共享内存参与地交换数据，是写出高性能核函数的关键一步；本章与第 3 章（执行模型、归约）、第 4 章（全局内存、矩阵转置）紧密衔接。

1.2 学完本章，你应该能回答

学习目标	检验方式
理解共享内存的作用域、生命周期及静态/动态分配方式	能写出 __shared__ 与 extern __shared__ 的用法，并说明动态共享内存的第三个参数含义
理解 __syncthreads() 的语义及竞争条件	能解释「先加载到共享内存 → __syncthreads → 再计算」的模式，以及分支内误用的危险
掌握** bank 与 bank 冲突**的定义及 bank 索引公式	能判断「连续 32 个 float」「同一列」「步长 2」等访问是否冲突
掌握用 padding 消除二维共享内存按列访问的 bank 冲突	能写出 tile[32][33] 的声明及线性索引 row*(BLOCK_DIM+1)+col，并说明为何 +1 能消除冲突
掌握共享内存矩阵转置的流程与索引转换	能写出读入 tile[ty][tx]、写出 tile[tx][ty] 及 idx_in、idx_out 的推导
掌握共享内存归约的流程（加载→块内多轮归约→写回）	能与第 3 章全局内存归约对比，说明为何共享内存版本减少全局访问次数
理解常量内存的适用场景及 cudaMemcpyToSymbol 用法	能说明同址广播、不同址串行化及 64KB 限制；会写主机端拷贝与核函数内只读使用
区分常量缓存与只读缓存（__ldg）	能说出数据来源、容量与最佳访问模式的区别
理解线程束洗牌的概念及 __shfl_down_sync 等用法	能写出 warp 内归约片段，并说明 mask 与「无需共享内存」的优势

1.3 博客阅读导图（本章架构）

第 5 章  共享内存和常量内存
│
├── 二、共享内存概述（书 5.1 节）
│   ├── 为什么使用共享内存、声明与生命周期（静态/动态）
│   └── 同步与竞争条件、__syncthreads 的正确使用
│
├── 三、共享内存的布局与 Bank 冲突（书 5.2 节）★ 重难点
│   ├── 数据排布、二维与线性索引
│   ├── Bank 与 bank 冲突、典型访问模式表
│   └── 使用 padding 消除 bank 冲突
│
├── 四、使用共享内存的矩阵转置（书 5.3 节）
│   ├── 思路与索引转换（读入 tile、写出转置）
│   └── 完整核函数与有效带宽
│
├── 五、使用共享内存的并行归约（书 5.4 节）
│   ├── 思路与共享内存阶段
│   └── 完整核函数与第 3 章归约的对比
│
├── 六、常量内存（书 5.5 节）
│   ├── 常量内存与常量缓存、同址广播与不同址串行化
│   ├── cudaMemcpyToSymbol 与核函数用法
│   └── 常量缓存与只读缓存的区别
│
├── 七、线程束洗牌（书 5.6 节）
│   ├── 概念与适用场景
│   └── warp 内求和示例与 mask
│
└── 八、本章小结与重难点回顾

---

二、共享内存概述（书 5.1 节）

本节对应书中对共享内存的总体介绍：为什么需要共享内存、它与全局内存的关系，以及基本用法与同步要求。

2.1 为什么使用共享内存

这一节会说明：大多数设备端数据访问都是从全局内存开始的，且多数 GPU 应用程序受内存带宽限制。全局内存虽然容量大，但延迟高（约 400～800 周期）、带宽有限；而共享内存是 SM 上的片上可编程存储器，仅对同一线程块内的线程可见，延迟远低于全局内存（约数十周期），带宽也远高于全局内存。因此，把会被多次使用或需要在线程间共享的数据先加载到共享内存，在块内完成计算后再写回全局内存，可以显著减少对全局内存的访问次数并改善有效带宽。

共享内存的典型用途包括（业界共识）：(1) 作为可编程缓存：将全局内存中的一块数据（tile）以合并方式读入共享内存，供块内线程多次访问；(2) 重组访问模式：例如矩阵转置时，先按行合并读入 tile，在共享内存中按列写出，使写回全局内存时也是合并的；(3) 块内协作与归约：同一块内线程通过共享内存交换中间结果，并用 __syncthreads() 保证可见性。第4章已提到：矩阵转置若只靠全局内存，要么读合并写不合并，要么写合并读不合并；使用共享内存做中转即可实现读写两端都接近合并。

理解与体会：共享内存的本质是「程序员可控的片上缓存」——CPU 的 L1/L2 对程序员透明，而 GPU 把这块存储暴露出来，让我们显式决定「先把哪块数据搬进来、以什么顺序读写」。用好共享内存，就能在「合并访问」与「复杂访问模式」之间架起桥梁：全局内存只负责合并地进、合并地出；中间的重组与复用全在片上完成。

2.2 共享内存的声明与生命周期

共享内存在核函数内通过 __shared__ 修饰符声明，其作用域为线程块，生命周期与线程块一致——块内所有线程共享同一块物理存储，块执行结束后该存储被释放。常见有两种分配方式：

分配方式	语法	特点
静态共享内存	__shared__ float tile[N];	编译时确定大小，N 为常量
动态共享内存	extern __shared__ float s[];	大小在核函数启动时由第三参数指定

动态共享内存的典型用法：在主机端启动核函数时，通过执行配置的第三个参数传入每个块所需的共享内存字节数，例如：

// 主机端：每个块需要 blockDim.x * blockDim.y * sizeof(float) 字节共享内存
size_t smemSize = block.x * block.y * sizeof(float);
kernel<<<grid, block, smemSize>>>(d_in, d_out, nx, ny);

核函数内则声明为 extern __shared__ float s[];，由运行时根据 smemSize 分配。若一个核函数中同时需要多种用途的共享内存（例如一块做 tile、一块做归约缓冲），可在同一块 extern __shared__ 中通过偏移手动划分区域。需要留意的是：共享内存是有限的资源；第3章已给出不同计算能力下「每块最大共享内存」等限制，分配过多会导致每 SM 可驻留的块数减少，影响占用率。

易错点：动态共享内存的第三个参数是字节数，不是元素个数；若传 block.x * block.y 而忘记乘 sizeof(float)，会导致越界或未定义行为。另外，extern __shared__ 只能声明一次，多块共享内存需在同一块缓冲区里用偏移划分。

2.3 同步与竞争条件

块内线程通过共享内存共享数据时，必须注意写后读、读后写等竞争条件。这一节会说明：在逻辑上并行执行的线程，在物理上并非同时执行；若线程 A 写入共享内存某位置，线程 B（可能在不同线程束中）读取该位置，必须在 A 写完之后、B 再读，否则结果未定义。__syncthreads() 提供了块内栅栏：当被调用时，块内所有线程必须都到达该点，并且在该点之前的所有全局内存和共享内存访问对该块内所有线程可见之后，才能继续执行。因此，典型模式是：先由所有线程合作把全局数据加载到共享内存 → __syncthreads() → 再在共享内存上做计算或重组 → 必要时再次 __syncthreads() → 最后写回全局内存。第3章已强调：不同块之间没有同步；块间协作只能通过结束内核、再启动新内核实现。

易错点：__syncthreads() 必须被块内所有线程无差别地执行。若放在 if (tid < 32) 等分支内，则只有部分线程到达栅栏，其余线程不执行该调用，会导致死锁或未定义行为。因此，要么所有线程都进入同一分支并调用 __syncthreads()，要么在分支外、所有线程必经的路径上调用。

---

三、共享内存的布局与 Bank 冲突（书 5.2 节）

本节对应书中「数据在共享内存中如何排布」以及「bank 与 bank 冲突」的讲解，是使用共享内存时避免性能陷阱的核心。

3.1 共享内存的数据排布

共享内存在物理上被组织为线性地址空间，按字节地址连续排列。这一节会说明：从编程视角看，声明为 __shared__ float tile[32][32] 的二维数组，在内存中仍按行优先线性存储——tile[row][col] 的线性字节地址可表示为：

[
\text{字节地址} = (\text{row} \times \text{列数} + \text{col}) \times \text{每元素字节数}
]

例如 tile[row][col]（float 为 4 字节），线性索引为 row * 32 + col，对应字节地址为 (row*32+col)*4。从二维共享内存到线性全局内存的索引转换是矩阵转置等核函数中的常见操作：全局内存中矩阵按行优先存储，块内线程通常负责一块子矩阵（tile）；线程 (tx, ty) 在块内对应 tile 的 (tx, ty)，在全局矩阵中的行、列则需要加上块的起始偏移（blockIdx.y * blockDim.y、blockIdx.x * blockDim.x），从而得到全局线性索引 row * nx + col。掌握这一转换是编写正确且高效的 tile 型核函数的基础。

3.2 Bank 与 Bank 冲突

为了提供高带宽，共享内存被划分为多个 bank。这一节会说明：在 Fermi 及之后的架构中，共享内存通常被组织为 32 个 bank，每个 bank 的宽度为 4 字节；连续 4 字节字（word）依次映射到 bank 0、1、2、…、31，再循环。即：

[
\text{bank 索引} = (\text{字节地址} \div 4) \bmod 32
]

因此，若 32 个线程（一个 warp）各访问一个 float，且这些 float 的线性索引连续（如 tid 对应 s[tid]），则 32 个地址分别落在 32 个不同 bank，无冲突，可在一个周期内完成。但若 32 个线程访问同一 bank 的不同地址（例如同一列的不同行），则这些访问会串行化，产生 bank 冲突，导致多次事务、延迟增加。需要留意的是：若线程束内多个线程访问同一 bank 的不同 4 字节字，就会发生 bank 冲突；若访问的是同一 bank 内的同一地址，硬件会广播该字给所有请求线程，计为一次访问，不产生冲突。

下表归纳几种典型访问模式与是否冲突（以 32 bank、4 字节宽为例；原书配有共享内存 bank 布局示意图）：

访问模式	示例（32 线程）	Bank 情况	是否冲突
连续 32 个 float	s[threadIdx.x]	各线程不同 bank	无冲突
同一地址	s[0] 全 warp 读	同一 bank 同一地址，广播	无冲突
同一列（行优先二维）	s[row][0]，row 不同	同一 bank 不同地址	有冲突
步长为 2 的间隔访问	s[threadIdx.x * 2]	偶数 bank 重复	2 路冲突

因此，在编写使用二维共享内存的核函数时（如矩阵转置），按列访问共享内存往往会导致 bank 冲突，需要结合下文 padding 或索引重排来消除。

理解与体会：Bank 冲突的本质是「多个线程同时向同一个 bank 要不同地址」——硬件上每个 bank 一次只能服务一个地址，所以这些请求会被串行化。反之，若 32 个线程访问 32 个不同 bank，或访问同一 bank 的同一地址（硬件会广播），则无冲突。记忆窍门：连续、错开、同址往往安全；**同一逻辑列（行优先下的列）**在未 padding 时易冲突。

3.3 使用 Padding 消除 Bank 冲突

对于行优先的二维共享内存，有一种常见情况：对于行优先的二维共享内存 tile[BLOCK_DIM][BLOCK_DIM]，若线程按列读取（即 tile[row][col] 中 col 固定、row 由线程变化），则同一 warp 内不同 row 对应的线性索引相差「列数」的整数倍；当列数为 32 的约数时，不同行会落在同一 bank，产生冲突。Padding 的做法是：在每行末尾多分配一列（或若干列），使「逻辑列数」变为 33（或 32+1），这样按列访问时，相邻行对应的线性索引相差 33 个 4 字节字，其 bank 索引相差 (33 \bmod 32 = 1)，从而 32 行会分散到 32 个不同 bank，消除冲突。

声明示例（典型写法）：

#define BLOCK_DIM 32
// 无 padding：tile[32][32]，按列访问可能 32 路冲突
// 有 padding：tile[32][33]，按列访问时 bank 分散
__shared__ float tile[BLOCK_DIM][BLOCK_DIM + 1];  // +1 即 padding 一列

线性索引公式变为：tile[row][col] 对应线性索引 row * (BLOCK_DIM + 1) + col。这样，从全局内存「按行」合并读入时，写入共享内存的列下标仍为 col，行下标为 row，对应共享内存中的位置 tile[row][col]；当从共享内存「按列」读出并写回全局内存时，读的是 tile[col][row]（转置），由于 padding，同一 warp 内不同 col 对应的 tile[col][row] 会落在不同 bank。矩阵转置核函数中会使用带 padding 的共享内存，下面第四节给出完整实现。

易错点：使用 padding 后，若你在代码里手动计算线性索引（例如做一维化处理），必须用 row * (BLOCK_DIM + 1) + col，不能再用 row * BLOCK_DIM + col，否则会访问错误位置或越界。

---

四、使用共享内存的矩阵转置（书 5.3 节）

本节对应书中「用共享内存避免非合并的全局内存访问」的矩阵转置案例：先合并读入 tile，在共享内存中转置，再合并写出。

4.1 思路与索引转换

第4章已说明：若直接对全局内存做转置，读按行则写按列（非合并），写按行则读按列（非合并）。共享内存方案：每个线程块负责矩阵中的一个 tile（例如 32×32）；(1) 块内线程按行从全局内存把该 tile 读入共享内存（合并读）；(2) __syncthreads() 确保 tile 写满；(3) 块内线程按列从共享内存读出（对应转置后的行），写回全局内存的转置位置（合并写）。这样，全局内存的读和写都是合并的。

二维共享内存到线性全局内存的索引（要点；原书通常配有 tile 与全局矩阵的索引对应图）：设块维度为 (BLOCK_DIM, BLOCK_DIM)，线程 (tx, ty) 在块内对应 tile 的 (tx, ty)。读入时：全局输入矩阵 in 的行、列为 row = blockIdx.y * BLOCK_DIM + ty，col = blockIdx.x * BLOCK_DIM + tx，线性索引 idx_in = row * nx + col。写入时：转置后输出矩阵 out 中，该元素应位于「原列作行、原行作列」，即 out_col = row，out_row = col，线性索引 idx_out = out_row * ny + out_col = col * ny + row。共享内存中，读入阶段 tile[ty][tx] = in[idx_in]；写出阶段从共享内存读 tile[tx][ty]（行列互换即转置），写 out[idx_out] = tile[tx][ty]。若使用 padding，共享内存声明为 tile[BLOCK_DIM][BLOCK_DIM+1]，则读入时仍为 tile[ty][tx]，写出时读 tile[tx][ty]，此时按「列」读共享内存（第一维为 tx，第二维为 ty），因 padding 而避免 bank 冲突。

理解与体会：矩阵转置是「用共享内存换双端合并」的典范：第 4 章在纯全局内存下只能二选一（读合并或写合并），这里用一块 32×33 的共享内存做中转，全局内存只做「按行整块读入」和「按行整块写出」，中间在片上完成行列互换。代价是多一次 __syncthreads() 和有限的共享内存占用；收益是有效带宽的显著提升，用 nvprof 对比即可验证。

4.2 完整核函数与主函数框架

下面给出带 padding 的共享内存矩阵转置核函数及简要主函数框架（含 CHECK 与执行配置）：

#include <stdio.h>
#include <cuda_runtime.h>

#define CHECK(call)                                                            \
    do {                                                                       \
        cudaError_t err = call;                                                 \
        if (err != cudaSuccess) {                                               \
            printf("Error %s:%d, code %d, reason: %s\n", __FILE__, __LINE__,   \
                   err, cudaGetErrorString(err));                               \
            exit(1);                                                            \
        }                                                                      \
    } while (0)

#define BLOCK_DIM 32

// 使用共享内存 + padding 的矩阵转置（读写均合并）
__global__ void transposeShared(float *in, float *out, int nx, int ny) {
    __shared__ float tile[BLOCK_DIM][BLOCK_DIM + 1];  // padding 消除 bank 冲突

    unsigned int tx = threadIdx.x;
    unsigned int ty = threadIdx.y;
    unsigned int row = blockIdx.y * blockDim.y + ty;
    unsigned int col = blockIdx.x * blockDim.x + tx;

    if (row < ny && col < nx) {
        tile[ty][tx] = in[row * nx + col];  // 按行读全局内存 → 合并
    }
    __syncthreads();

    // 转置：输出 (col, row) 对应输入 (row, col)
    unsigned int out_row = col;
    unsigned int out_col = row;
    if (out_row < nx && out_col < ny) {
        out[out_row * ny + out_col] = tile[tx][ty];  // 按列读共享内存（padding 后无冲突），按行写全局 → 合并
    }
}

int main() {
    int nx = 1024, ny = 1024;
    size_t nBytes = nx * ny * sizeof(float);

    float *h_in  = (float *)malloc(nBytes);
    float *h_out = (float *)malloc(nBytes);
    // 初始化 h_in ...

    float *d_in, *d_out;
    CHECK(cudaMalloc((void **)&d_in, nBytes));
    CHECK(cudaMalloc((void **)&d_out, nBytes));
    CHECK(cudaMemcpy(d_in, h_in, nBytes, cudaMemcpyHostToDevice));

    dim3 block(BLOCK_DIM, BLOCK_DIM);
    dim3 grid((nx + block.x - 1) / block.x, (ny + block.y - 1) / block.y);
    transposeShared<<<grid, block>>>(d_in, d_out, nx, ny);

    CHECK(cudaMemcpy(h_out, d_out, nBytes, cudaMemcpyDeviceToHost));
    // 验证转置正确性 ...
    cudaFree(d_in);
    cudaFree(d_out);
    free(h_in);
    free(h_out);
    return 0;
}

有效带宽（与第4章一致）可按下式估算：

[
\text{有效带宽（GB/s）} = \frac{(\text{读字节数} + \text{写字节数}) \times 10^{-9}}{\text{核函数执行时间（秒）}}
]

对于 (n \times n) 的 float 矩阵转置，读字节数 = 写字节数 = (n^2 \times 4)。可对比「仅全局内存的转置」与「共享内存 + padding 的转置」的 nvprof 带宽或耗时；通常共享内存版本能显著提高有效带宽，因为读写两端都实现了合并。

---

五、使用共享内存的并行归约（书 5.4 节）

本节对应书中「用共享内存缓存数据以减少全局内存访问」的并行归约：第3章已在全局内存上做了相邻/交错配对归约，本节把每块的输入先加载到共享内存，在共享内存中做多轮归约，最后只把块内结果写回全局内存，从而大幅减少全局内存访问次数。

5.1 思路与共享内存阶段

典型思路：每个线程块负责一段连续元素；块内每个线程先从全局内存读取一个（或若干个）元素到寄存器或共享内存，然后在共享内存中进行多轮成对归约（与第3章的交错配对类似），每轮后**__syncthreads()，最后将块内部分和由线程 0 写回全局内存的 g_odata[blockIdx.x]。这样，每块只对全局内存做一次块大小的加载和一次标量写回**，而不是在全局内存上就地多轮读写，显著减少全局内存事务。

共享内存中的归约与第3章 reduceInterleaved 逻辑一致：stride 从 blockDim.x/2 开始减半，每次 s[tid] += s[tid + stride]，__syncthreads()。注意：共享内存数组 s[] 需在归约前由各线程用全局数据初始化，例如 s[tid] = g_idata[blockIdx.x * blockDim.x + tid]（需处理边界，如 idx >= n 时赋 0 或单位元）。

理解与体会：第 3 章的归约是在全局内存上「就地」多轮读写，每轮都有大量全局内存访问；本章把「一块」的数据一次性搬进共享内存，在片上完成多轮归约，每块只写回一个标量。这样全局内存的读从「多轮、可能非合并」变成「一次块大小的合并加载 + 一次标量写回」，是典型的「用共享内存减少全局访问次数」的模式。

5.2 完整共享内存归约核函数

下面给出基于共享内存的整数求和归约核函数（与 reduce 系列对应），含块内加载、共享内存归约、写回：

__global__ void reduceShared(int *g_idata, int *g_odata, unsigned int n) {
    __shared__ int s[256];  // 假设 blockDim.x <= 256

    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    // 从全局内存加载到共享内存，越界则填 0
    s[tid] = (idx < n) ? g_idata[idx] : 0;
    __syncthreads();

    for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
        if (tid < stride) {
            s[tid] += s[tid + stride];
        }
        __syncthreads();
    }

    if (tid == 0) {
        g_odata[blockIdx.x] = s[0];
    }
}

还可结合循环展开与 warp 内最后阶段不调用 __syncthreads() 等优化（第3章已提及）；进一步地，线程束洗牌（第七节）可在 warp 内无共享内存参与的情况下做归约，减少共享内存带宽压力。有效带宽公式仍为：(\text{有效带宽} = (\text{读字节数} + \text{写字节数}) / \text{时间})；共享内存版本因全局内存读写的总字节数不变但访问次数与模式更优，通常能获得更高的有效带宽与更短的执行时间。

---

六、常量内存（书 5.5 节）

本节对应书中对常量内存与常量缓存的介绍，以及常量内存与只读缓存的区别。

6.1 常量内存与常量缓存

常量内存是设备上的一种只读内存，在源码中通过 __constant__ 在全局/文件作用域声明，由主机在运行时通过 cudaMemcpyToSymbol 写入，设备端核函数只能读取。这一节会说明：常量内存在硬件上配有专用的常量缓存；当线程束内所有线程访问同一地址时，该地址的一个 32 位字会从常量缓存广播给整个 warp，因此只需一次读取即可满足 32 个线程，非常适合「所有线程用同一参数」的场景（如半径、系数表、查找表首地址等）。

常量内存有容量限制（如 64 KB），且当线程束内线程访问不同地址时，不同地址的请求会串行化，性能可能不如合并的全局内存访问。因此需要留意的是：常量内存最适合所有线程读取相同常量的情况；若各线程访问不同地址，应优先考虑合并的全局内存或只读缓存（__ldg / 只读数据缓存）。典型适用场景包括：小型的查找表或系数表、所有线程共用的配置参数（如滤波半径、块大小）、以及 kernel 启动前由主机一次性写入、设备端只读的常量数据。

理解与体会：常量内存与共享内存的「同址广播」类似，但作用域是全局、且为只读。其优势仅在 warp 内所有线程访问同一地址 时发挥；一旦 32 个线程访问 32 个不同地址，会退化成串行读取，反而不如合并的全局内存。所以使用前要问：这段数据是不是「全 warp 同参数」？若是小表但各线程查不同下标，更宜用全局内存 + __ldg。

6.2 常量内存的用法：cudaMemcpyToSymbol 与核函数

主机端使用 cudaMemcpyToSymbol 将数据拷贝到设备上的 __constant__ 变量：

__constant__ float constData[256];

// 主机端：拷贝到常量内存
CHECK(cudaMemcpyToSymbol(constData, h_data, sizeof(float) * 256));

核函数中直接按只读使用：

__global__ void useConstant(float *out, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        out[idx] = constData[idx % 256];  // 若 idx%256 在 warp 内相同则更优
    }
}

当 idx % 256 在 warp 内一致时（例如同一 warp 内 idx 连续且块大小 256 的整数倍），可发挥广播优势；否则会串行化。可对比「常量内存 + 同址访问」与「全局内存合并访问」的性能差异。

6.3 常量缓存与只读缓存的区别

（对应第 4 章只读缓存）：常量缓存专用于 __constant__ 变量，粒度与广播行为如上所述。只读缓存（计算能力 3.5+）可用于全局内存的只读加载（通过 __ldg 或 const __restrict__ 指针），加载粒度为 32 字节，对分散读取常优于 L1。二者区别可归纳为：

特性	常量缓存	只读缓存（全局内存只读）
数据来源	__constant__，cudaMemcpyToSymbol	全局内存（__ldg / const restrict）
容量	有限（如 64 KB）	独立于常量内存
最佳访问	同址（warp 内同一地址）广播	可分散，32 字节粒度
适用场景	小表、全 warp 同参数	大只读数组、分散读

---

七、线程束洗牌（Warp Shuffle）（书 5.6 节）

本节对应书中「Programming with the warp shuffle instruction」：在同一 warp 内通过寄存器直接交换数据，无需共享内存，延迟更低。

7.1 概念与适用场景

线程束洗牌是计算能力 3.0 及以上支持的** warp 内寄存器级数据交换指令。这一节会说明：同一 warp 内的线程可以通过 __shfl 系列内建函数，根据源 lane 索引或偏移直接读取另一线程的寄存器值，而不经过共享内存**。这样可以在 warp 内做归约、扫描等操作时减少共享内存的 bank 冲突与带宽占用，并降低延迟。

典型用法：归约时，warp 内 32 个线程先各自持有一个值，通过 __shfl_down_sync 等逐步两两相加，最终 warp 内 lane 0 得到该 warp 的部分和；多个 warp 再通过共享内存或第二次核函数合并。同步：洗牌指令需指定参与线程的掩码（如 __shfl_down_sync(0xffffffff, val, offset) 表示整个 warp 参与），以保证语义正确。其他变体包括：__shfl_up_sync（从更高 lane 索引取）、__shfl_xor_sync（按 lane 索引异或得到源 lane，常用于 butterfly 式归约），以及按源 lane 索引直接读取的 __shfl_sync，便于实现更灵活的数据交换模式。

理解与体会：Warp shuffle 把「warp 内数据交换」从共享内存搬到了寄存器间直接交换，既省共享内存带宽，又降低延迟。适合「先 warp 内归约/扫描，再块内跨 warp 合并」的两阶段设计；此时共享内存只需存每个 warp 的一个结果，而不是 32 个中间值。

7.2 简单示例：warp 内求和

下面给出 warp 内 32 个线程对寄存器 val 做求和，结果在 lane 0 的片段（完整归约还需块内跨 warp 合并）：

__device__ __forceinline__ int warpReduce(int val) {
    for (int offset = warpSize / 2; offset > 0; offset >>= 1) {
        val += __shfl_down_sync(0xffffffff, val, offset);
    }
    return val;
}

__global__ void reduceShuffle(int *g_idata, int *g_odata, unsigned int n) {
    unsigned int tid = threadIdx.x;
    unsigned int idx = blockIdx.x * blockDim.x + tid;
    int val = (idx < n) ? g_idata[idx] : 0;

    int lane = tid % warpSize;
    int warpId = tid / warpSize;

    val = warpReduce(val);  // warp 内归约

    if (lane == 0) {
        // 每个 warp 的 lane 0 写入共享内存，再在块内归约
        // 此处省略共享内存合并多 warp 的代码
    }
}

__shfl_down_sync(mask, val, offset) 表示当前线程从「相对自己向下 offset 个 lane」的线程取 val；其他变体包括 __shfl_up_sync、__shfl_xor_sync 等，用于不同的数据移动模式。使用洗牌时需注意 mask 必须与当前活跃线程一致（例如整个 warp 参与则 mask = 0xffffffff），否则未定义行为。

易错点：Shuffle 的 mask 必须正确。若内核中有分支导致部分线程不参与归约，则不应使用全 warp 的 0xffffffff，而应使用实际参与线程的掩码（如 __activemask() 或根据分支计算出的 mask），否则未参与线程的寄存器值会被错误地参与交换，导致结果错误或未定义行为。

---

八、本章小结与重难点回顾

8.1 知识小结（与书中对应）

• 共享内存：片上可编程、块内可见、低延迟高带宽；用于作可编程缓存、重组访问模式（如矩阵转置）、块内协作与归约。分静态（__shared__ T arr[N]）与动态（extern __shared__ T s[]，启动时传字节数）。
• 布局与 bank：共享内存按行优先线性排布；Fermi 及以后为 32 bank、每 bank 4 字节，bank 索引 = (字节地址/4) mod 32。同一 warp 访问同一 bank 不同地址会 bank 冲突；同址则广播无冲突。
• Padding：二维共享内存每行多一列（如 tile[32][33]）可使按列访问时 bank 分散，消除典型矩阵转置中的 bank 冲突。
• 矩阵转置：用共享内存 tile 先合并读入、__syncthreads()、再按列读出写回全局，实现读写双端合并；配合 padding 避免共享内存按列读时的 bank 冲突。
• 并行归约：将块内数据先载入共享内存，在共享内存中做多轮交错配对归约，最后写回一块一个标量；可结合 warp 内洗牌进一步减少共享内存访问。
• 常量内存：__constant__ + cudaMemcpyToSymbol，只读、专用常量缓存；warp 内同址访问时广播，性能最优；不同址会串行化。
• 常量缓存 vs 只读缓存：常量缓存面向 __constant__；只读缓存面向全局内存只读加载（__ldg），适合大数组或分散读。
• 线程束洗牌：__shfl_*_sync 在 warp 内通过寄存器交换数据，无需共享内存，适合 warp 内归约/扫描；需正确设置 mask。

8.2 重难点速查

重难点	要点
共享内存 vs 全局内存	共享内存片上、块内可见、低延迟高带宽；用作可编程缓存时，全局内存只负责合并进、合并出，中间重组在片上完成。
Bank 与 bank 冲突	32 bank、每 bank 4 字节；bank 索引 = (字节地址/4) mod 32。同一 warp 访问同一 bank 不同地址→冲突；同址→广播无冲突。
Padding 消除冲突	二维行优先下按列访问易冲突；每行多一列（如 33）使相邻行线性索引差 33 字，33 mod 32=1，32 行分散到 32 个 bank。
矩阵转置索引	读入 tile[ty][tx]=in[row*nx+col]，写出 out[col*ny+row]=tile[tx][ty]；row/col 含块偏移。
__syncthreads 使用	必须被块内所有线程执行，不能放在仅部分线程进入的分支内，否则死锁或未定义行为。
常量内存适用场景	仅当 warp 内同址访问时才有广播优势；不同址会串行化，且容量有限（如 64KB）。大表或分散读宜用全局内存+__ldg。
Warp shuffle 与 mask	warp 内寄存器级交换，无需共享内存；mask 须与参与线程一致，分支导致部分线程不参与时不能用 0xffffffff。

8.3 学习思考

• 与第 4 章的衔接：第 4 章指出矩阵转置在纯全局内存下无法同时做到读合并与写合并；本章用共享内存做 tile 中转，先合并读入、在片上转置、再合并写出，是「用空间换带宽」的典型。理解这一点，就能举一反三到其他「访问模式与合并要求矛盾」的场景。
• 与第 3 章的关系：第 3 章归约在全局内存上多轮读写；本章把块内数据搬进共享内存后多轮归约，每块只写回一个标量，大幅减少全局访问次数。进一步地，warp 内用 shuffle 做第一段归约，共享内存只存每个 warp 的一个结果，可同时减少共享内存占用与 bank 冲突风险。
• 实践建议：写新内核时，若存在「先按一种顺序读、再按另一种顺序写」的需求，优先考虑共享内存 tile + __syncthreads；使用二维共享内存时，若按列访问，检查是否需 padding（如 tile[32][33]）；常量内存只留给「全 warp 同参数」的小表或配置；归约类内核可先实现共享内存版，再在 warp 内用 __shfl_down_sync 做一层优化。

下表归纳本章涉及的三种可编程内存的典型用途与性能要点（与书中描述一致）：

内存类型	作用域	典型用途	性能要点
共享内存	线程块	tile 缓存、转置/归约中转、块内协作	低延迟、高带宽；注意 bank 冲突与 padding
全局内存	全部	主数据、输入输出缓冲	对齐与合并；可配合共享内存减少访问
常量内存	全部	只读参数、小表、全 warp 同址数据	同址广播优；不同址易串行化，容量有限（如 64 KB）

---

下一章预告

在下一篇博客中，我们将进入第 6 章：流和并发：

• CUDA 流的概念与默认流
• 基于流的重叠：主机与设备、传输与计算、多流并发
• 流回调与事件（cudaEvent）
• 多 GPU 编程入门

从「单次内核与单次传输」到「用流与事件组织并发」，是进一步提升吞吐与隐藏延迟的关键一步。

---

本章自测

1. 共享内存的 bank 冲突是什么？如何用 padding 消除 32×32 tile 的写冲突？
2. 用共享内存做矩阵转置时，典型步骤是什么？为何能同时实现读合并与写合并？
3. 块内线程在共享内存上协作时，为什么必须在「加载完成」与「使用数据前」之间加 __syncthreads()？

答案与解析

1. 同一 warp 内多线程若访问同一 bank 的不同地址，会串行化（bank 冲突）。32×32 按行存时，同一列属于同一 bank；在每行后多开一列（如 32×33）做 padding，使同一列元素分布到不同 bank，可消除写冲突。
2. 步骤：块内线程按行从全局内存合并读入共享内存 → __syncthreads() → 按列从共享内存读出并写回全局转置位置（写合并）。全局侧只做「整块合并读入」和「整块合并写出」，中间在片上转置。
3. 因为线程并行执行，写共享内存的线程与读共享内存的线程可能不同步；若不加 __syncthreads()，读线程可能读到未写入的数据。栅栏保证「所有线程写完、对块内可见」后再继续。

---

系列导航：导读 | 上一篇：第4章 全局内存 | 下一篇：第6章 流和并发

---

本文为「CUDA C编程权威指南」系列博客第 5 篇，共 10 章。基于《Professional CUDA C Programming》by John Cheng, Max Grossman, Ty McKercher。