（五）共享内存和常量内存

本篇笔记参考如下：

https://blog.csdn.net/D1557329860/article/details/143813183?ops_request_misc=&request_id=&biz_id=102&utm_term=共享内存&utm_medium=distribute.pc_search_result.none-task-blog-2~all~sobaiduweb~default-2-143813183.142

通过安排全局内存访问模式，我们学会了如何实现良好的性能并且避免了浪费事务。未对齐的内存访问是没有问题的，因为现代的GPU硬件都有一级缓存，但在跨全局内存的非合并内存访问，仍然会导致带宽利用率不会达到最佳标准。根据算法性质和相应的访问模式，非合并访问可能是无法避免的。然而，在许多情况下，使用共享内存来提高全局内存合并访问是有可能的。

5.1 CUDA共享内存概述

GPU中有两种类型的内存：

  • 板载内存
  • 片上内存

全局内存是较大的板载内存，具有相对较高的延迟。共享内存是较小的片上内存，具有相对较低的延迟，并且共享内存可以提供比全局内存高得多的带宽。

共享内存通常的用途有：

  1. 块内线程通信的通道
  2. 用于全局内存数据的可编程管理的缓存
  3. 高速暂存存储器，用于转换数据以优化全局内存访问模式

5.1.1 共享内存

共享内存（shared memory，SMEM）是GPU的一个关键部件。物理上，每个SM都有一个小的低延迟内存池，这个内存池被当前正在该SM上执行的线程块中的所有线程所共享。

共享内存使同一个线程块中的线程能够互相协作，便于重用片上数据，并可以大大降低核函数所需的全局内存带宽。由于共享内存中的内容是由应用程序显式管理的，所以它通常被描述为可编程管理的缓存。

注意：相较于二级缓存和全局内存，共享内存和一级缓存在物理上更接近SM。因此，共享内存相较于全局内存而言，延迟要低大约20～30倍，而带宽高其大约10倍。L1 硬件缓存和共享内存在物理上是共享同一块片上存储资源，如果不显式使用共享内存，这部分物理存储空间通常会被更多地分配给 L1 缓存，以提升全局内存访问的效率。

当每个线程块开始执行时，会分配给它一定数量的共享内存。这个共享内存的地址空间被线程块中所有的线程共享。

每个线程束发出共享内存访问请求。在理想的情况下，每个被线程束共享内存访问的请求在一个事务中完成。最坏的情况下，每个共享内存的请求在32个不同的事务中顺序执行。如果多个线程访问共享内存中的同一个字，一个线程读取该字后，通过多播把它发送给其他线程。

共享内存被SM中的所有常驻线程块划分，因此，共享内存是限制设备并行性的关键资源。一个核函数使用的共享内存越多，处于并发活跃状态的线程块就越少。

由于在CUDA中允许手动管理共享内存，所以通过在数据布局上提供更多的细粒度控制和改善片上数据的移动，使得对应用程序代码进行优化变得更简单了。

5.1.2 共享内存分配

在CUDA的源代码文件中，共享内存可以被声明为一个本地的CUDA核函数或是一个全局的CUDA核函数。CUDA支持一维、二维和三维共享内存数组的声明。

共享内存变量用下列修饰符进行声明：

1
__shared__

例如：

1
__shared__ float tile[size_y][size_x];

如果在核函数中进行声明，那么这个变量的作用域就局限在该内核中。如果在文件的任何核函数外进行声明，那么这个变量的作用域对所有核函数来说都是全局的。

如果共享内存的大小在编译时是未知的，那么可以用extern关键字声明一个未知大小的数组。

例如：

1
extern __shared__ int tile[];

5.1.3 共享内存存储体和访问模式

共享内存可以用来隐藏全局内存延迟和带宽对性能的影响。要想充分理解这些资源，了解共享内存是如何被安排的，对其将会有所帮助。

5.1.3.1 内存存储体

为了获得高内存带宽，共享内存被分为32个同样大小的内存模型，它们被称为存储体，它们可以被同时访问。。有32个存储体是因为在一个线程束中有32个线程。共享内存是一个一维地址空间。根据GPU的计算能力，共享内存的地址在不同模式下会映射到不同的存储体中。

5.1.3.2 存储体冲突

当线程束发出共享内存请求时，有以下3种典型的模式：

  • 并行访问：多个地址访问多个存储体
  • 串行访问：多个地址访问同一个存储体
  • 广播访问：单一地址读取单一存储体

理想状态（并行）： 如果一个 Warp 中的 32 个线程，分别访问 32 个不同的 Bank（每个 Bank 只负责一个线程的请求），那么 GPU 只需要发起 1 次内存事务 就能搞定。这就像 32 个人去 32 个窗口办事，大家互不干扰，速度最快。

每个线程访问一个32位字。因为每个线程访问不同存储体中的地址，所以没有存储体冲突。在不规则的随机访问模式下。因为每个线程访问不同的存储体，所以也没有存储体冲突。

冲突状态（串行）： 如果多个线程同时请求访问同一个 Bank 里的不同地址，就会发生存储体冲突 (Bank Conflict)。

后果： 该操作无法一次完成，必须拆分成多个内存事务依次进行。

比喻： 就像多个人挤在同一个窗口排队，后面的人必须等前面的人办完，这会导致原本并行的操作变成了串行，内存带宽利用率自然就大幅下降。

5.1.3.3 访问模式

无论是哪种架构，其核心逻辑都是将“连续的字节地址”映射到“循环的 32 个存储体”中。

  • Fermi (计算能力 2.x): 存储体宽度为 4 字节（32位）。
    • 公式： 存储体索引=(字节地址÷4)(mod32)存储体索引 = (\text{字节地址} \div 4) \pmod{32}
    • 理解： 物理上，地址每隔 4 个字节就换到一个新的 Bank。如果你按 4 字节（如 int 或 float）连续访问，正好一个 Warp 的 32 个线程对应 32 个不同的 Bank，实现完美并行。
  • Kepler (计算能力 3.x): 存储体宽度升级为 8 字节（64位）。
    • 公式： 存储体索引=(字节地址÷8)(mod32)存储体索引 = (\text{字节地址} \div 8) \pmod{32}
    • 理解： 地址每隔 8 个字节才换 Bank。这意味着一个 8 字节的 Bank 里其实“装”下了两个 32 位的字。

Fermi 的局限： 一个 Bank 在一个周期内只能吐出 32 位数据。如果两个线程访问同一个 Bank 的不同位置，必须排队。

Kepler 的优势（32位模式）： 虽然 Word 0 和 Word 32 都映射在 Bank 0，但因为 Kepler 的 Bank 物理带宽是 64位/周期，它能一次性把这 64 位数据都读出来。

即使两个线程分别要 Word 0 和 Word 32，硬件可以一次完成读取并分发，不再产生冲突。

对于50系列的Blackwell 架构，共享内存的存储体宽度（Bank Width）并未升级，依然维持在 32位（4字节）。

尽管 Bank 宽度没变，但 Blackwell 在共享内存相关的其他维度进行了重大增强：

  • SM 共享内存容量提升：每个 SM 的 L1/共享内存组合容量从 Hopper 的 192 KB 提升至 228 KB（增加了约 19%）。
  • 单块分配限制放宽：计算能力 10.0 的设备（如 B200）允许单个线程块寻址高达 227 KB 的共享内存。
  • 引入张量内存 (TMEM)：这是最重要的改变。Blackwell 额外增加了 256 KB 的 TMEM，专门服务于 Tensor Core。TMEM 的读取带宽高达 16 TB/s，极大地减轻了共享内存和寄存器的压力。

TMEM：

  • 物理容量： 每个 SM 配备 256 KB 的 TMEM，大小与该 SM 的寄存器文件（Register File）完全一致。
  • 二维布局： 逻辑上组织为 128 行（Lanes）× 512 列，每个单元格为 4 字节（32 位）。
  • 物理位置： 相比共享内存（SMEM），TMEM 在物理上与 Tensor Core 耦合得更紧密，提供了极高的访问带宽。

5.1.3.4 内存填充

内存填充是避免存储体冲突的一种方法。。假设只有5个共享内存存储体。如果所有线程访问bank 0的不同地址，那么会发生一个五向的存储体冲突。

如果所有线程访问bank 0的不同地址，那么会发生一个五向的存储体冲突。解决这种存储体冲突的一个方法是在每N个元素之后添加一个字，这里的N是存储体的数量。这就改变了从字到存储体的映射。

5.1.4 配置共享内存量

每个SM都有64 KB的片上内存。共享内存和一级缓存共享该硬件资源。CUDA为配置一级缓存和共享内存的大小提供了两种方法：

  • 按设备进行配置
  • 按核函数进行配置

使用下述的运行时函数，可以为在设备上启动的核函数配置一级缓存和共享内存的大小：

1
cudaError_t cudaDeviceSetCacheConfig(cudaFuncCache cacheConfig);

可以通过传递不同的枚举值来改变分配比例：

配置常量 描述 适用场景
cudaFuncCachePreferNone 默认设置 无特殊偏好，由驱动程序自动平衡。
cudaFuncCachePreferShared 偏好共享内存 算法显式使用了大量 __shared__ 变量，需要更大的共享空间。
cudaFuncCachePreferL1 偏好 L1 缓存 算法未显式使用共享内存，或者存在大量全局内存重复读取。
cudaFuncCachePreferEqual 均等分配 L1 缓存和共享内存分配相同大小的空间（部分架构支持）。

如果需要执行一个核函数，它可自由地选择不同的配置。每个核函数的配置可以覆盖设备范围的设置，也可以使用以下运行时函数进行设置：

1
cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCacheca cheConfig);

核函数使用的这种配置是由核函数指针func指定的。启动一个不同优先级的内核比启动有最近优先级设置的内核更可能会导致隐式设备同步。

GPU缓存的行为比CPU缓存的行为更难以理解。GPU使用不同的启发式算法删除数据。在GPU上，数百个线程共享相同的一级缓存，数千个线程共享相同的二级缓存。因此，数据删除在GPU上可能会发生得更频繁而且更不可预知。

5.1.5 同步

并行线程间的同步是所有并行计算语言的重要机制。

同步的两个基本方法如下所示：

  • 障碍
  • 内存栅栏

5.1.5.1 弱排序内存模型

现代的内存架构有一个宽松的内存模型。这意味着，内存访问不一定按照它们在程序中出现的顺序进行执行。

如果指令之间是相互独立的，线程从不同内存中读取数据的顺序和读指令在程序中出现的顺序不一定相同。

5.1.5.2 显式障碍

在CUDA中，障碍只能在同一线程块的线程间执行。在核函数中，可以通过调用下面的函数来指定一个障碍点：

1
__syncthreads();

syncthreads作为一个障碍点来发挥作用，它要求块中的线程必须等待直到所有线程都到达该点。syncthreads还确保在障碍点之前，被这些线程访问的所有全局和共享内存对同一块中的所有线程都可见。__syncthreads用于协调同一块中线程间的通信。当块中的某些线程访问共享内存或全局内存中的同一地址时，会有潜在问题（写后读、读后写、写后写），这将导致在那些内存位置产生未定义的应用程序行为和未定义的状态。

5.1.5.3 内存栅栏

内存栅栏的功能可确保栅栏前的任何内存写操作对栅栏后的其他线程都是可见的。

1)我们通过以下固有函数可以在线程块内创建内存栅栏：

1
void __threadfence_block();

__threadfence_block保证了栅栏前被调用线程产生的对共享内存和全局内存的所有写操作对栅栏后同一块中的其他线程都是可见的。

2)使用下面的固有函数来创建网格级内存栅栏：

1
void __threadfence();

__threadfence挂起调用的线程，直到全局内存中的所有写操作对相同网格内的所有线程都是可见的。

3)使用下面的函数可以跨系统（包括主机和设备）设置内存栅栏：

1
void __threadfence_system();

__threadfence_system挂起调用的线程，以确保该线程对全局内存、锁页主机内存和其他设备内存中的所有写操作对全部设备中的线程和主机线程是可见的。

5.1.5.4 Volatile修饰符

在全局或共享内存中使用volatile修饰符声明一个变量，可以防止编译器优化，编译器优化可能会将数据暂时缓存在寄存器或本地内存中。当使用volatile修饰符时，编译器假定任何其他线程在任何时间都可以更改或使用该变量的值。因此，这个变量的任何引用都会直接被编译到全局内存读指令或全局内存写指令中，它们都会忽略缓存。

5.2 共享内存的数据布局

为了全面了解如何有效地使用共享内存，本节将使用共享内存研究几个简单的例子:

  • 方阵与矩阵数组
  • 行主序与列主序访问
  • 静态与动态共享内存的声明
  • 文件范围与内核范围的共享内存
  • 内存填充与无内存填充

5.2.1 方形共享内存

使用共享内存可以直接缓存具有方形维度的全局数据。方形矩阵的简单维度可以很容易从二维线程索引中计算出一维内存偏移。

使用下面的语句静态声明一个二维共享内存变量：

1
__shared__ int tile[N][N];

下图显示了一个共享内存块，它在每个维度有32个元素，且按行主序进行存储。上部的图显示了一维数据布局的实际排列。

在同一个线程束中若有访问独立存储体的线程，则它是最优的。相同线程束中的线程可由连续的threadIdx.x值来确定。属于不同存储体的共享内存元素也可以通过字偏移进行连续存储。

5.2.1.1 行主序访问和列主序访问

在例子中网格有一个二维线程块，块中每个维度包含32个可用的线程。可以使用下面的宏来定义块维度：

1
2
#define BDIMX 32
#define BDIMY 32

使用下面的宏来定义核函数的执行配置：

1
2
dim3 block (BDIMX, BDIMY);
dim3 grid	(1, 1);

核函数有两个简单操作：

  • 将全局线程索引按行主序写入到一个二维共享内存数组中
  • 从共享内存中按行主序读取这些值并将它们存储到全局内存中

可以用如下方法静态声明一个二维共享内存数组：

1
__shared__ int tile[BDIMY][BDIMX];

接下来，需要为每个线程计算全局线程索引，它是根据其二维线程ID进行计算的。因为只有一个线程块将被启动，该索引转换可以被简化为：

1
unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

idx用于模拟从输入矩阵中读取值。基于线程全局ID的写入位置，存储idx的值到输出数组，将允许可视化核函数的访问模式。

将全局线程索引按行主序顺序写入共享内存块，可以按如下方式进行：

1
tile[threadIdx.y][threadIdx.x] = idx;

一旦达到同步点（使用syncthreads函数），所有线程必须将存储的数据送到共享内存块中，这样就可以按行主序从共享内存给全局内存赋值。

1
out[idx] = tile[threadIdx.x][threadIdx.y];

对于下列核函数来说

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
__global__ void setRowReadCol(int *out)
{
    // static shared memory
    __shared__ int tile[BDIMY][BDIMX];

    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[threadIdx.y][threadIdx.x] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

到目前为止，在内核中有3个内存操作：

  • 共享内存的存储操作
  • 共享内存的加载操作
  • 全局内存的存储操作

因为相同线程束中的线程有连续的threadIdx.x值，并且可以使用threadIdx.x索引共享内存数组tile的最内层维度，所以核函数无存储体冲突。

另一方面，如果在将数据分配给共享内存块时交换threadIdx.y和threadIdx.x，线程束的内存将会按列主序访问。每个共享内存的加载和存储将导致32路存储体冲突。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
__global__ void setColReadCol (int *out)
{
    // static shared memory
    __shared__ int tile[BDIMX][BDIMY];

    // mapping from thread index to global memory index
    unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;

    // shared memory store operation
    tile[threadIdx.x][threadIdx.y] = idx;

    // wait for all threads to complete
    __syncthreads();

    // shared memory load operation
    out[idx] = tile[threadIdx.x][threadIdx.y];
}

编译并采用ncu分析checkSmemSquare.cu

1
2
nvcc -arch=sm_120 checkSmemSquare.cu -o checkSmemSquare
ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum ./checkSmemSquare

1.基础读写核函数

setRowReadRow (行存行取)

  • 逻辑：线程 (threadIdx.x,threadIdx.y)(threadIdx.x, threadIdx.y) 写入 tile[y][x] 后读取同一位置。
  • Bank 映射：由于 tile 的宽度是 32，连续的线程 xx 访问连续的内存地址，正好映射到 32 个不同的 Bank。
  • 结果：无冲突。这是最理想的访问模式，访存事务可以完全并行。

setColReadCol (列存列取)

  • 逻辑：写入和读取都采用 tile[x][y]。
  • 冲突分析：
    • 写入 (Store)：Warp 内相邻线程（xx 变化）访问的是 tile[0][y], tile[1][y], ...。由于第一维跨度是 32，这些地址都映射到同一个 Bank。
    • 读取 (Load)：同理，纵向读取也会导致所有线程挤在同一个 Bank。
  • 结果：双重 32 路冲突（Load 992, Store 992）。效率极低，访存被完全串行化。

2.坐标转置核函数

setRowReadCol (行存列取)

  • 逻辑：按行顺序写入 tile[y][x]，但按列顺序读取 tile[x][y]。
  • 冲突分析：
    • 写入时是连续的，无冲突。
    • 读取时，Warp 内的线程访问 tile[0][y], tile[1][y], ...，地址间隔为 32 个 int。根据公式 Index=(Address/4)(mod32)Index = (Address / 4) \pmod{32}，所有 32 个线程的索引结果相同。
  • 结果：32 路读取冲突（Load 992）。这是典型的转置操作性能瓶颈。

setRowReadColDyn (动态共享内存版本)

  • 逻辑：使用 extern __shared__ 定义一维数组，通过手动计算 row_idx 和 col_idx 模拟二维访问。
  • 结果：与静态版本一致。动态分配不改变硬件的 Bank 映射逻辑，依然存在严重的读取冲突。

3.优化核函数 (Padding)

setRowReadColPad (静态填充)

  • 逻辑：将共享内存声明为 tile[BDIMY][BDIMX + IPAD]，即 [32][33]。
  • 原理：
    • 每行多出一个空位，使得下一行的起始地址在 Bank 索引上向后偏移 1 位。
    • 当按列读取 tile[0][y], tile[1][y] 时，它们的物理地址不再相差 32 的整数倍，而是错开了 Bank。
  • 结果：冲突降为 0。通过极小的内存开销换取了巨大的性能提升。

setRowReadColDynPad (动态填充)

  • 逻辑：在动态一维数组索引计算中加入偏移量 (blockDim.x + IPAD)。
  • 结果：冲突降为 0。证明了通过算法逻辑模拟 Padding 同样可以规避 Blackwell 架构下的存储体冲突。

得到输出如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
setColReadCol(int *) (1, 1, 1)x(32, 32, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  -------------------------------------------------------- ----------- ------------
  Metric Name                                              Metric Unit Metric Value
  -------------------------------------------------------- ----------- ------------
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                      992
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                      992
  -------------------------------------------------------- ----------- ------------

setRowReadRow(int *) (1, 1, 1)x(32, 32, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  -------------------------------------------------------- ----------- ------------
  Metric Name                                              Metric Unit Metric Value
  -------------------------------------------------------- ----------- ------------
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                        0
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
  -------------------------------------------------------- ----------- ------------

setRowReadCol(int *) (1, 1, 1)x(32, 32, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  -------------------------------------------------------- ----------- ------------
  Metric Name                                              Metric Unit Metric Value
  -------------------------------------------------------- ----------- ------------
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                      992
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
  -------------------------------------------------------- ----------- ------------

setRowReadColDyn(int *) (1, 1, 1)x(32, 32, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  -------------------------------------------------------- ----------- ------------
  Metric Name                                              Metric Unit Metric Value
  -------------------------------------------------------- ----------- ------------
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                      992
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
  -------------------------------------------------------- ----------- ------------

setRowReadColPad(int *) (1, 1, 1)x(32, 32, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  -------------------------------------------------------- ----------- ------------
  Metric Name                                              Metric Unit Metric Value
  -------------------------------------------------------- ----------- ------------
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                        0
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
  -------------------------------------------------------- ----------- ------------

setRowReadColDynPad(int *) (1, 1, 1)x(32, 32, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  -------------------------------------------------------- ----------- ------------
  Metric Name                                              Metric Unit Metric Value
  -------------------------------------------------------- ----------- ------------
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                        0
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
  -------------------------------------------------------- ----------- ------------

分析结果得到下表

核函数名称 访问模式 (Load/Store) Load 冲突数 Store 冲突数 性能影响分析
setRowReadRow 行存 / 行取 0 0 最优：Warp内线程访问连续地址，完美映射到 32 个 Banks。
setRowReadCol 行存 / 列取 992 0 严重冲突：读取步长为 32，导致 32 个线程请求同一个 Bank，效率下降 32 倍。
setColReadCol 列存 / 列取 992 992 极差：写入和读取阶段均发生 32 路冲突，内存事务完全串行化。
setRowReadColPad 行存 / 列取 (带填充) 0 0 高效优化：通过 33 步长错开 Bank 索引，彻底消除冲突。
setRowReadColDyn 动态内存 / 列取 992 0 冲突：动态分配不改变地址映射规律，依然存在 32 路读取冲突。
setRowReadColDP 动态内存 (带填充) 0 0 高效优化：动态共享内存配合 Padding 同样可以消除冲突。

5.2.1.3 动态共享内存

可以动态声明共享内存，从而实现这些相同的核函数。可以在核函数外声明动态共享内存，使它的作用域为整个文件，也可以在核函数内声明动态共享内存，将其作用域限制在该内核之中。

动态共享内存必须被声明为一个未定大小的一维数组，因此，需要基于二维线程索引来计算内存访问索引。

因为要在这个核函数中按行主序写入，按列主序读取，所以需要保留以下两个索引：

·row_idx：根据二维线程索引计算出的一维行主序内存偏移量

·col_idx：根据二维线程索引计算出的一维列主序内存偏移量

5.2.1.4 填充静态声明的共享内存

填充数组是避免存储体冲突的一种方法。填充静态声明的共享内存很简单。

1
__shared__ int tile[BDIMY][BDIMX+1];

5.2.1.5 填充动态声明的共享内存

填充动态声明的共享内存数组更加复杂。当执行从二维线程索引到一维内存索引的索引转换时，对于每一行必须跳过一个填充的内存空间

1
2
3
4
#define IPAD  1

    unsigned int row_idx = threadIdx.y * (blockDim.x + IPAD) + threadIdx.x;
    unsigned int col_idx = threadIdx.x * (blockDim.x + IPAD) + threadIdx.y;

图中显示了这些内存索引计算，这些计算使用了一个简化的五存储体共享内存实现。

5.2.1.6 方形共享内存内核性能的比较

到目前为止，从所有执行过的内核运行时间可以看出：

  • 使用填充的内核可提高性能，因为它减少了存储体冲突
  • 带有动态声明共享内存的内核增加了少量的消耗

5.2.2 矩形共享内存

矩形共享内存是一个更普遍的二维共享内存，在矩形共享内存中数组的行与列的数量不相等。

1
__shared__ int tile[row][col];

当执行一个转置操作时，不能像在方形共享内存中一样，只是通过简单地转换来引用矩形数组的线程坐标。当使用矩形共享内存时，这样做会导致内存访问冲突。需要基于矩阵维度重新计算访问索引，以重新实现之前描述的核函数。

测试一个矩形共享内存数组，其每行有32个元素，每列有16个元素。在下面的宏中定义了维度：

1
2
#define BDIMX 32
#define BDIMY 16

矩形共享内存块被分配如下：

1
__shared__int tile[BDIMY][BDIMX];

为了简单起见，内核将被启动为只有一个网格和一个二维线程块，该线程块的大小与矩形共享内存数组相同。

1
2
dim3 block (BDIMX, BDIMY);
dim3 grid (1,1);

编译并执行代码checkSmemRectangle.cu

1
2
nvcc -arch=sm_120 checkSmemRectangle.cu -o checkSmemRectangle
ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum ./checkSmemRectangle

代码分析如下

1.基础读写核函数

setRowReadRow

  • 共享内存结构：tile[16][32]。
  • 访问逻辑：Warp 内相邻线程（threadIdx.x 变化）访问 tile[y][x]，对应内存中连续的地址。
  • 分析：由于最内层维度正好是 32，步长为 1，完全没有 Bank Conflict。

setColReadCol

  • 共享内存结构：tile[32][16]。
  • 访问逻辑：Warp 访问 tile[x][y]。
  • 冲突分析：因为最内层维度变成了 16，Warp 中相邻线程访问的地址间隔为 16 个 int。根据 Index=(Address/4)(mod32)Index = (Address/4) \pmod{32}，线程 ii 和线程 i+2i+2 会映射到同一个 Bank。
  • 结果：产生 2 路冲突。

2.转置与坐标映射核函数

setRowReadCol (行存列取)

  • 坐标转换：代码通过 idx / blockDim.y 和 idx % blockDim.y 计算转置后的索引 irow 和 icol。
  • 冲突分析：
    • 读取操作访问 tile[icol][irow]。
    • 对于相邻线程，icolicol 的变化步长为 1，但 irowirow 的变化会导致访问跨度。
  • 结果：在矩形块中，由于 blockDim.y (16) 不是 32，读取时的步长不再是 32，这会引发不同程度的 Bank Conflict。

setColReadCol2

  • 逻辑：同样使用 tile[16][32] 结构，但通过复杂的坐标变换模拟转置。
  • 分析：这演示了即便共享内存声明是“安全”的（内层为 32），如果访问索引逻辑（icol, irow）导致 Warp 内部线程请求同一个 Bank 的数据，依然会发生冲突。

3.填充优化 (Padding)

setRowReadColPad (静态填充)

  • 共享内存结构：tile[16][32 + 2]（这里 IPAD=2）。
  • 原理：
    • 在每一行末尾增加 2 个 int 的填充。
    • 这改变了每一行起始元素的 Bank 映射。由于 34(mod32)=234 \pmod{32} = 2，垂直方向上的元素现在错开了 Bank 索引。
  • 结果：有效减少或消除了列读取时的 Bank 冲突。

setRowReadColDynPad (动态填充)

  • 逻辑：在动态分配的一维数组中手动计算步长：idx = y * (32 + 2) + x。
  • 分析：这展示了在动态共享内存下，只要手动维持“列宽 + Padding”的索引逻辑，同样能达到规避冲突的效果。

得到的结果如下

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
(cuda_env) root@63320d55395d:~/cuda_learn/cuda_program/CodeSamples/chapter05# ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum ./checkSmemRectangle
==PROF== Connected to process 18308 (/root/cuda_learn/cuda_program/CodeSamples/chapter05/checkSmemRectangle)
/root/cuda_learn/cuda_program/CodeSamples/chapter05/./checkSmemRectangle at device 0: NVIDIA GeForce RTX 5060 Ti with Bank Mode:4-Byte <<< grid (1,1) block (32,16)>>>
==PROF== Profiling "setRowReadRow(int *)" - 0: 0%....50%....100% - 1 pass
==PROF== Profiling "setColReadCol(int *)" - 1: 0%....50%....100% - 1 pass
==PROF== Profiling "setColReadCol2(int *)" - 2: 0%....50%....100% - 1 pass
==PROF== Profiling "setRowReadCol(int *)" - 3: 0%....50%....100% - 1 pass
==PROF== Profiling "setRowReadColDyn(int *)" - 4: 0%....50%....100% - 1 pass
==PROF== Profiling "setRowReadColPad(int *)" - 5: 0%....50%....100% - 1 pass
==PROF== Profiling "setRowReadColDynPad(int *)" - 6: 0%....50%....100% - 1 pass
==PROF== Disconnected from process 18308
[18308] checkSmemRectangle@127.0.0.1
  setRowReadRow(int *) (1, 1, 1)x(32, 16, 1), Context 1, Stream 7, Device 0, CC 12.0
    Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                        0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
    -------------------------------------------------------- ----------- ------------

  setColReadCol(int *) (1, 1, 1)x(32, 16, 1), Context 1, Stream 7, Device 0, CC 12.0
    Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                      240
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                      240
    -------------------------------------------------------- ----------- ------------

  setColReadCol2(int *) (1, 1, 1)x(32, 16, 1), Context 1, Stream 7, Device 0, CC 12.0
    Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                      240
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                      240
    -------------------------------------------------------- ----------- ------------

  setRowReadCol(int *) (1, 1, 1)x(32, 16, 1), Context 1, Stream 7, Device 0, CC 12.0
    Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                      240
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
    -------------------------------------------------------- ----------- ------------

  setRowReadColDyn(int *) (1, 1, 1)x(32, 16, 1), Context 1, Stream 7, Device 0, CC 12.0
    Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                      240
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
    -------------------------------------------------------- ----------- ------------

  setRowReadColPad(int *) (1, 1, 1)x(32, 16, 1), Context 1, Stream 7, Device 0, CC 12.0
    Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                        0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
    -------------------------------------------------------- ----------- ------------

  setRowReadColDynPad(int *) (1, 1, 1)x(32, 16, 1), Context 1, Stream 7, Device 0, CC 12.0
    Section: Command line profiler metrics
    -------------------------------------------------------- ----------- ------------
    Metric Name                                              Metric Unit Metric Value
    -------------------------------------------------------- ----------- ------------
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum                        0
    l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum                        0
    -------------------------------------------------------- ----------- ------------
核函数名称 访问模式 Load 冲突数 Store 冲突数 冲突级别 结果分析
setRowReadRow 行存 / 行取 0 0 无 步长为 1，线程束内 32 线程对应 32 个 Banks，完美并行。
setColReadCol 列存 / 列取 240 240 2 路冲突 步长为 16，两个线程撞在一个 Bank。15×16 (warps)=24015 \times 16 \text{ (warps)} = 240。
setColReadCol2 坐标变换 240 240 2 路冲突 即便内层维度设为 32，复杂的索引逻辑依然导致了 2 路冲突。
setRowReadCol 行存 / 列取 240 0 2 路冲突 读取步长为 16，导致每个 Bank 被 2 个线程竞争。
setRowReadColDyn 动态内存 240 0 2 路冲突 动态分配不改变物理 Bank 映射，依然受步长 16 影响。
setRowReadColPad 带填充 (Pad) 0 0 无 步长变为 32+2=3432+2=34，彻底错开 Bank 索引。
setRowReadColDynPad 动态 + 填充 0 0 无 算法层面的 Padding 同样在 Blackwell 上完美生效。

5.3 减少全局内存访问

使用共享内存的主要原因之一是要缓存片上的数据，从而减少核函数中全局内存访问的次数。

  • 如何重新安排数据访问模式以避免线程束分化
  • 如何展开循环以保证有足够的操作使指令和内存带宽饱和

5.3.1 使用共享内存的并行归约

我们以reduceInteger.cu代码中的核函数reduceGmem作为基准函数

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
__global__ void reduceGmem(int *g_idata, int *g_odata, unsigned int n)
{
    // set thread ID
    unsigned int tid = threadIdx.x;
    int *idata = g_idata + blockIdx.x * blockDim.x;

    // boundary check
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx >= n) return;

    // in-place reduction in global memory
    if (blockDim.x >= 1024 && tid < 512) idata[tid] += idata[tid + 512];

    __syncthreads();

    if (blockDim.x >= 512 && tid < 256) idata[tid] += idata[tid + 256];

    __syncthreads();

    if (blockDim.x >= 256 && tid < 128) idata[tid] += idata[tid + 128];

    __syncthreads();

    if (blockDim.x >= 128 && tid < 64) idata[tid] += idata[tid + 64];

    __syncthreads();

    // unrolling warp
    if (tid < 32)
    {
        volatile int *vsmem = idata;
        vsmem[tid] += vsmem[tid + 32];
        vsmem[tid] += vsmem[tid + 16];
        vsmem[tid] += vsmem[tid +  8];
        vsmem[tid] += vsmem[tid +  4];
        vsmem[tid] += vsmem[tid +  2];
        vsmem[tid] += vsmem[tid +  1];
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = idata[0];
}

该核函数通过“分治法”将一个数据块的和合并到该块的第一个元素中：

  • 数据定位：每个线程块处理全局内存中的一段连续区域 idata。
  • 阶梯式求和：每一轮迭代将参与计算的线程数减半（512 -> 256 -> 128 …），直到最后只剩下一个结果。
  • 全局内存操作：注意该函数的命名是 reduceGmem，意味着所有的加法操作都是直接在**全局内存（Global Memory）**上进行的。

代码在每一轮归约后都调用了 __syncthreads()：由于 GPU 调度 Warp 的进度不一，必须确保当前轮次的所有线程都完成了加法写入，下一轮次的线程才能读取正确的值。

当存活的线程数减少到 32（即一个 Warp）时，代码进入了特殊的处理逻辑：在最后一个 Warp 内，SIMT 架构保证了这 32 个线程是同步执行指令的。因此，程序员可以省去昂贵的 __syncthreads() 调用，直接顺序执行加法。

1
2
3
4
5
6
7
if (tid < 32)
{
    volatile int *vsmem = idata;
    vsmem[tid] += vsmem[tid + 32];
    vsmem[tid] += vsmem[tid + 16];
    ...
}

编译代码reduceInteger.cu

1
nvcc -arch=sm_120 reduceInteger.cu -o reduceInteger

采用ncu进行分析

1
ncu --metrics gpu__time_duration.sum ./reduceInteger

得到以下结果

1
2
3
4
5
6
7
reduceGmem(int *, int *, unsigned int) (32768, 1, 1)x(128, 1, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  ---------------------- ----------- ------------
  Metric Name            Metric Unit Metric Value
  ---------------------- ----------- ------------
  gpu__time_duration.sum          us       129.76
  ---------------------- ----------- ------------

时间为129.76us。

接下来测试下面的原地归约核函数reduceSmem，核心修改部分为

1
2
smem[tid] = idata[tid];
__syncthreads();

它增加了带有共享内存的全局内存操作。这个核函数和原来的reduceGmem核函数几乎相同。然而，reduceSmem函数没有使用全局内存中的输入数组子集来执行原地归约，而是使用了共享内存数组smem。

执行得到结果如下

1
2
3
4
5
6
7
reduceSmem(int *, int *, unsigned int) (32768, 1, 1)x(128, 1, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  ---------------------- ----------- ------------
  Metric Name            Metric Unit Metric Value
  ---------------------- ----------- ------------
  gpu__time_duration.sum          us        77.15
  ---------------------- ----------- ------------

时间为77.15us。

采用以下命令测试全局内存加载和存储事务

1
ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_sectors_pipe_lsu_mem_global_op_st.sum ./reduceInteger

得到如下结果

核函数类别 内存层级 加载事务 (LD Sectors) 存储事务 (ST Sectors) 访存压力分析
reduceGmem 全局内存 2,228,224 1,081,344 极高：每一轮求和都要读写全局内存。
reduceSmem 共享内存 524,288 32,768 极大优化：仅在加载数据和保存结果时访问全局内存。

从这个结果可以看出，使用共享内存明显减少了全局内存访问。

5.3.2 使用展开的并行归约

在前面的核函数中，每个线程块处理一个数据块。在第3章中，我们可以通过一次运行多个I/O操作，展开线程块来提高内核性能。以下内核展开了4个线程块，即每个线程处理来自于4个数据块的数据元素。

可以得到代码如下

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
__global__ void reduceSmemUnroll(int *g_idata, int *g_odata, unsigned int n)
{
    // static shared memory
    __shared__ int smem[DIM];

    // set thread ID
    unsigned int tid = threadIdx.x;

    // global index, 4 blocks of input data processed at a time
    unsigned int idx = blockIdx.x * blockDim.x * 4 + threadIdx.x;

    // unrolling 4 blocks
    int tmpSum = 0;

    // boundary check
    if (idx < n)
    {
        int a1, a2, a3, a4;
        a1 = a2 = a3 = a4 = 0;
        a1 = g_idata[idx];
        if (idx + blockDim.x < n) a2 = g_idata[idx + blockDim.x];
        if (idx + 2 * blockDim.x < n) a3 = g_idata[idx + 2 * blockDim.x];
        if (idx + 3 * blockDim.x < n) a4 = g_idata[idx + 3 * blockDim.x];
        tmpSum = a1 + a2 + a3 + a4;
    }

    smem[tid] = tmpSum;
    __syncthreads();

    // in-place reduction in shared memory
    if (blockDim.x >= 1024 && tid < 512) smem[tid] += smem[tid + 512];

    __syncthreads();

    if (blockDim.x >= 512 && tid < 256)  smem[tid] += smem[tid + 256];

    __syncthreads();

    if (blockDim.x >= 256 && tid < 128)  smem[tid] += smem[tid + 128];

    __syncthreads();

    if (blockDim.x >= 128 && tid < 64)   smem[tid] += smem[tid + 64];

    __syncthreads();

    // unrolling warp
    if (tid < 32)
    {
        volatile int *vsmem = smem;
        vsmem[tid] += vsmem[tid + 32];
        vsmem[tid] += vsmem[tid + 16];
        vsmem[tid] += vsmem[tid +  8];
        vsmem[tid] += vsmem[tid +  4];
        vsmem[tid] += vsmem[tid +  2];
        vsmem[tid] += vsmem[tid +  1];
    }

    // write result for this block to global mem
    if (tid == 0) g_odata[blockIdx.x] = smem[0];
}

计算逻辑：每个线程不再只读取一个全局内存数据，而是利用 idx + n * blockDim.x 读取四个不同位置的数据（a1 到 a4），并在进入共享内存前就完成初步求和。

隐藏延迟：通过在同一个线程中发起多个独立的全局内存读取请求，GPU 能够更好地利用内存控制器带宽，隐藏全局内存的高时延。

减少 Grid 开销：由于每个线程块处理了原来 4 倍的数据，因此启动核函数所需的 Grid 规模减小了 4 倍（从 32768 降至 8192），这显著降低了硬件调度的负载。

输入命令ncu --metrics gpu__time_duration.sum ./reduceInteger我们可以看到结果

1
2
3
4
5
6
7
reduceSmemUnroll(int *, int *, unsigned int) (8192, 1, 1)x(128, 1, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  ---------------------- ----------- ------------
  Metric Name            Metric Unit Metric Value
  ---------------------- ----------- ------------
  gpu__time_duration.sum          us        41.06
  ---------------------- ----------- ------------

时间为41.06us。

输入命令ncu --metrics l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum,l1tex__t_sectors_pipe_lsu_mem_global_op_st.sum ./reduceInteger测试全局内存加载和存储事务

可以得到以下表格

核函数类别 内存层级 加载事务 (LD Sectors) 存储事务 (ST Sectors) 访存压力分析
reduceGmem 全局内存 2,228,224 1,081,344 极高：每一轮求和都要读写全局内存。
reduceSmem 共享内存 524,288 32,768 极大优化：仅在加载数据和保存结果时访问全局内存。
reduceGmemUnroll 全局内存 1,081,344 401,408 中等：通过循环展开减少了总的访存次数。
reduceSmemUnroll 共享内存 524,288 8,192 最优：进一步压缩了结果存储的事务数。

5.3.3 使用动态共享内存的并行归约

并行归约核函数还可以使用动态共享内存来执行，通过以下声明，在reduceSmem-Unroll中用动态共享内存取代静态共享内存：

1
extern __shared__ int smem[];

5.3.4 有效带宽

由于归约核函数是受内存带宽约束的，所以评估它们时所使用的适当的性能指标是有效带宽。有效带宽是在核函数的完整执行时间内I/O的数量（以字节为单位）。

5.4 合并的全局内存访问

矩阵转置就是一个典型的例子：读操作被自然合并，但写操作是按照交叉访问的。在共享内存的帮助下，可以先在共享内存中进行转置操作，然后再对全局内存进行合并写操作。

5.4.1 基准转置内核

作为基准，下面的核函数是一个仅使用全局内存的矩阵转置的朴素实现。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
#define INDEX(ROW, COL, INNER) ((ROW) * (INNER) + (COL))

__global__ void naiveGmem(float *out, float *in, const int nrows, const int ncols)
{
    // matrix coordinate (ix,iy)
    unsigned int row = blockIdx.y * blockDim.y + threadIdx.y;
    unsigned int col = blockIdx.x * blockDim.x + threadIdx.x;

    // transpose with boundary test
    if (row < nrows && col < ncols)
    {
        out[INDEX(col, row, nrows)] = in[INDEX(row, col, ncols)];
    }
}

/*
col = ix
nrows = ny
row = iy


row = iy
ncols = nx
col = ix 
*/

因为ix是这个核函数二维线程配置的最内层维度，全局内存读操作在线程束内是被合并的，而全局内存写操作在相邻线程间是交叉访问的。naiveGmem核函数的性能是一个下界。

以执行合并访问为目的的更改写操作会生成副本内核。因为读写操作将被合并，但仍执行相同数量的I/O，所以copyGmem函数将成为一个性能近似的上界：

1
2
3
4
5
6
7
8
9
10
11
12
13
__global__ void copyGmem(float *out, float *in, const int nrows, const int ncols)
{
    // matrix coordinate (ix,iy)
    unsigned int row = blockIdx.y * blockDim.y + threadIdx.y;
    unsigned int col = blockIdx.x * blockDim.x + threadIdx.x;

    // transpose with boundary test
    if (row < nrows && col < ncols)
    {
		    // NOTE this is a transpose, not a copy
        out[INDEX(row, col, ncols)] = in[INDEX(row, col, ncols)];
    }
}

对于这些测试，矩阵大小被设置为4096×4096，并且还会用到一个维度为32×16的二维线程块。

编译transposeRectangle.cu

1
nvcc -arch=sm_120 transposeRectangle.cu -o transposeRectangle

采用以下命令获取

1
ncu --metrics l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio,l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_st.ratio ./transposeRectangle

得到以下结果

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
copyGmem(float *, float *, int, int) (256, 256, 1)x(16, 16, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  -------------------------------------------------------------------- ----------- ------------
  Metric Name                                                          Metric Unit Metric Value
  -------------------------------------------------------------------- ----------- ------------
  l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio      sector            4
  l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_st.ratio      sector         4.00
  -------------------------------------------------------------------- ----------- ------------

naiveGmem(float *, float *, int, int) (256, 256, 1)x(16, 16, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  -------------------------------------------------------------------- ----------- ------------
  Metric Name                                                          Metric Unit Metric Value
  -------------------------------------------------------------------- ----------- ------------
  l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio      sector            4
  l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_st.ratio      sector           16
  -------------------------------------------------------------------- ----------- ------------

对于 copyGmem (理想状态)

  • 读取 (ld) = 4, 写入 (st) = 4.00
  • 分析：一个 Warp 有 32 个线程，每个线程处理一个 float（4 字节）。当这 32 个线程访问连续内存时，总共请求 32×4=12832 \times 4 = 128 字节。
  • 计算：128 字节/32 字节每扇区=4 个 Sectors128 \text{ 字节} / 32 \text{ 字节每扇区} = 4 \text{ 个 Sectors}。
  • 结论：这说明 copyGmem 实现了完全合并访问。硬件只需发出 4 个 Sector 请求就能满足整个 Warp 的需求，效率最高。

对于 naiveGmem (转置瓶颈)

  • 读取 (ld) = 4：读取依然是连续的（行优先读取），所以保持了高效。
  • 写入 (st) = 16：这是性能杀手。
  • 分析：由于是朴素转置，相邻线程写入的是不同的行。在你的矩阵规模下，这些地址在内存中跨度很大。
函数 写入效率 (Relative Efficiency) 带宽利用率
copyGmem 100% 接近理论带宽
naiveGmem ~25% (4/16) 极低，大量带宽浪费在无效的 Sector 传输上

5.4.2 使用共享内存的矩阵转置

为了避免交叉全局内存访问，可以使用二维共享内存来缓存原始矩阵的数据。

下面的核函数实现了使用共享内存的矩阵转置。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
__global__ void transposeSmem(float *out, float *in, int nrows, int ncols)
{
    // static shared memory
    __shared__ float tile[BDIMY][BDIMX];

    // coordinate in original matrix
    unsigned int row = blockDim.y * blockIdx.y + threadIdx.y;
    unsigned int col = blockDim.x * blockIdx.x + threadIdx.x;

    // linear global memory index for original matrix
    unsigned int offset = INDEX(row, col, ncols);

    if (row < nrows && col < ncols)
    {
      // load data from global memory to shared memory
      tile[threadIdx.y][threadIdx.x] = in[offset];
    }

    // thread index in transposed block
    unsigned int bidx, irow, icol;
    bidx = threadIdx.y * blockDim.x + threadIdx.x;
    irow = bidx / blockDim.y;
    icol = bidx % blockDim.y;

	  // NOTE - need to transpose row and col on block and thread-block level:
	  // 1. swap blocks x-y
	  // 2. swap thread x-y assignment (irow and icol calculations above)
	  // note col still has continuous threadIdx.x -> coalesced gst
	  col = blockIdx.y * blockDim.y + icol;
	  row = blockIdx.x * blockDim.x + irow;

    // linear global memory index for transposed matrix
	  // NOTE nrows is stride of result, row and col are transposed
    unsigned int transposed_offset = INDEX(row, col, nrows);
    // thread synchronization
    __syncthreads();

	  // NOTE invert sizes for write check
    if (row < ncols && col < nrows)
    {
        // store data to global memory from shared memory
        out[transposed_offset] = tile[icol][irow]; // NOTE icol,irow not irow,icol
    }
}

第一步：合并读取到共享内存

1
tile[threadIdx.y][threadIdx.x] = in[offset];

线程束（Warp）内的线程访问连续的 offset，实现合并读取。数据被暂存在 tile 中。此时 tile 里的布局和原矩阵块一致。

第二步：重新映射坐标（关键点）

这部分代码执行了“逻辑上的坐标置换”：

1
2
3
bidx = threadIdx.y * blockDim.x + threadIdx.x;
irow = bidx / blockDim.y;
icol = bidx % blockDim.y;

改变线程对 Tile 中元素的访问方式。它打乱了原有的 threadIdx.x 映射关系，目的是为了在最后写回时，让原本属于不同列的元素，由连续的线程写出到连续的内存地址。

即：为了保证写回全局内存时也是连续的（合并写入），我们让连续的线程去写转置矩阵的行。既然要写转置矩阵的行，那线程就必须从共享内存里按列取出数据。

第三步：同步与合并写入

1
2
__syncthreads(); // 必须同步，确保所有线程都完成了加载
out[transposed_offset] = tile[icol][irow];

代码解决了全局内存合并访问的问题，但它引入了一个新的潜在瓶颈：共享内存bank冲突。在 tile[icol][irow] 读取时，如果多个线程同时访问同一个 Bank 的不同行（通常共享内存分为 32 个 Bank），访问会被序列化。但两者的代价完全不在一个量级：

维度 全局内存 (Global Memory) 共享内存 (Shared Memory)
延迟 几百个时钟周期 几个时钟周期
不连续访问后果 产生大量冗余事务（如你看到的 Ratio=16） 产生排队（Bank Conflict）
优化目标 首要优化点：必须合并访问 次要优化点：尽量避免 Bank Conflict

可以发现结果

1
2
3
4
5
6
7
8
transposeSmem(float *, float *, int, int) (256, 256, 1)x(16, 16, 1), Context 1, Stream 7, Device 0, CC 12.0
  Section: Command line profiler metrics
  -------------------------------------------------------------------- ----------- ------------
  Metric Name                                                          Metric Unit Metric Value
  -------------------------------------------------------------------- ----------- ------------
  l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio      sector            4
  l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_st.ratio      sector            4
  -------------------------------------------------------------------- ----------- ------------

读取 (ld) = 4, 写入 (st) = 4.00与理想状态一致

输入以下命令

1
ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum ./transposeRectangle

可以得到以下数据

指标名称 冲突总数 严重程度 对应代码逻辑
shared_op_ld.sum (读取) 3,795,868 极高 out[...] = tile[icol][irow];
shared_op_st.sum (写入) 14,157 极低 tile[y][x] = in[offset];

shared_op_ld.sum (读取): 指的是线程从共享内存读取数据（Load）时发生的**bank冲突（Bank Conflict）**总次数。

shared_op_st.sum (写入): 指的是线程向共享内存写入数据（Store）时发生的银行冲突总次数。

读取阶段存在近 380 万次的银行冲突，GPU 的 LSU（加载/存储单元） 管线会频繁发生 Stall（停顿）。

  1. 指令延迟增加：原本 1 个时钟周期能发射的指令，现在可能需要 32 个周期才能完成。
  2. 带宽浪费：虽然全局内存的写入是合并的，但数据从共享内存“流出”的速度太慢，导致全局内存的带宽利用率无法跑满。

5.4.3 使用填充共享内存的矩阵转置

通过给二维共享内存数组tile中的每一行添加列填充，可以将原矩阵相同列中的数据元素均匀地划分到共享内存存储体中。

核函数代码如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
__global__ void transposeSmemPad(float *out, float *in, int nrows, int ncols)
{
    // static shared memory with padding
    __shared__ float tile[BDIMY][BDIMX + IPAD];

    // coordinate in original matrix
    unsigned int row = blockDim.y * blockIdx.y + threadIdx.y;
    unsigned int col = blockDim.x * blockIdx.x + threadIdx.x;

    // linear global memory index for original matrix
    unsigned int offset = INDEX(row, col, ncols);

    // thread index in transposed block
    unsigned int bidx, irow, icol;
    bidx = threadIdx.y * blockDim.x + threadIdx.x;
    irow = bidx / blockDim.y;
    icol = bidx % blockDim.y;

    // linear global memory index for transposed matrix
    unsigned int transposed_offset = INDEX(col, row, nrows);

    // transpose with boundary test
    if (row < nrows && col < ncols)
    {
        // load data from global memory to shared memory
        tile[threadIdx.y][threadIdx.x] = in[offset];

        // thread synchronization
        __syncthreads();

        // store data to global memory from shared memory
        out[transposed_offset] = tile[irow][icol];
    }
}

输入以下命令

1
ncu --metrics l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum ./transposeRectangle

可以得到以下数据与优化前对比

指标 (Metric) transposeSmem (无 Padding) transposeSmemPad (有 Padding) 变化趋势
Load 冲突总数 (ld.sum) 3,795,868 621,985 下降约 83.6%
Store 冲突总数 (st.sum) 14,157 530,764 上升约 36 倍

读取冲突 (Load) 的大幅下降

在 transposeSmem 中，读取阶段产生的 379万次 冲突是性能的瓶颈。这是因为 Warp 内的线程在纵向读取共享内存列时，所有线程都撞在了同一个 Bank 上。 而在 transposeSmemPad 中，由于使用了 tile[BDIMY][BDIMX + IPAD]，每一行的数据在 Bank 中产生了错位偏移。原本垂直对齐的元素现在分布在不同的 Bank 中，使得读取冲突骤降了 80% 以上。

写入冲突 (Store) 的异常上升

这是一个非常有趣的现象：Padding 虽然解决了读取冲突，但导致写入冲突从 1.4万 升至 53万。

  • 原因分析：这通常是因为在 transposeSmemPad 的代码逻辑中，为了简化索引或匹配 Padding 后的数组步长（Stride），写入操作不再是完美的横向连续访问，或者是因为 Padding 后的地址计算导致 Warp 内的线程在写入时产生了跨 Bank 的竞争。
  • 权衡 (Trade-off)：尽管写入冲突增加了，但由于读取侧的冲突减量（约 317 万次）远大于写入侧的增量（约 51 万次），总冲突数依然大幅度减少，整体性能通常会更好。

5.4.4 使用展开的矩阵转置

下面的核函数展开两个数据块的同时处理：每个线程现在转置了被一个数据块跨越的两个数据元素。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
__global__ void transposeSmemUnrollPadDyn (float *out, float *in, const int nrows,
        const int ncols)
{
    // dynamic shared memory
    extern __shared__ float tile[];

    unsigned int row = blockIdx.y * blockDim.y + threadIdx.y;
    unsigned int col = (2 * blockIdx.x * blockDim.x) + threadIdx.x;

    unsigned int row2 = row;
    unsigned int col2 = col + blockDim.x;

    // linear global memory index for original matrix
    unsigned int offset = INDEX(row, col, ncols);
    unsigned int offset2 = INDEX(row2, col2, ncols);

    // thread index in transposed block
    unsigned int bidx = threadIdx.y * blockDim.x + threadIdx.x;
    unsigned int irow = bidx / blockDim.y;
    unsigned int icol = bidx % blockDim.y;

    // coordinate in transposed matrix
    unsigned int transposed_offset = INDEX(col, row, nrows);
    unsigned int transposed_offset2 = INDEX(col2, row2, nrows);

    if (row < nrows && col < ncols)
    {
        tile[INDEX(threadIdx.y, threadIdx.x, BDIMX * 2 + IPAD)] = in[offset];
    }
    if (row2 < nrows && col2 < ncols)
    {
        tile[INDEX(threadIdx.y, blockDim.x + threadIdx.x, BDIMX * 2 + IPAD)] =
            in[offset2];
    }

    __syncthreads();

    if (row < nrows && col < ncols)
    {
        out[transposed_offset] = tile[INDEX(irow, icol, BDIMX * 2 + IPAD)];
    }
    if (row2 < nrows && col2 < ncols)
    {
        out[transposed_offset2] = tile[INDEX(irow, blockDim.x + icol, BDIMX * 2 + IPAD)];
    }
}

一个32×16的线程块配置与一个展开大小为（32＋32）×16的数据块一起使用。

通过展开的两块，更多的内存请求将同时处于运行状态并且读/写的吞吐量会提高。这里不再通过ncu分析，有兴趣的话大家可以自行尝试分析。

5.4.5 增大并行性

一个简单而有效的优化技术是调整线程块的维度，以找出最佳的执行配置。表中总结了在Tesla K40上各种线程块配置的测试结果。块大

小为16×16时展示出了最好的性能。可以通过调优的方式找到最优解，这也与我们组里项目的工具有关hh。

5.5 常量内存

常量内存是一种专用的内存，它用于只读数据和统一访问线程束中线程的数据。常量内存位于设备的DRAM上（和全局内存一样），并且有一个专用的片上缓存。和一级缓存和共享内存一样，从每个SM的常量缓存中读取的延迟，比直接从常量内存中读取的低得多。每个SM常量内存缓存大小的限制为64KB。

不同于学习的任何其他类型的内存而言，常量内存有一个不同的最优访问模式。在常量内存中，如果线程束中的所有线程都访问相同的位置，那么这个访问模式就是最优的。如果线程束中的线程访问不同的地址，则访问就需要串行。因此，一个常量内存读取的成本与线程束中线程读取唯一地址的数量呈线性关系。

在全局作用域中必须用以下修饰符声明常量变量：

1
__constant__

常量内存变量的生存期与应用程序的生存期相同，其对网格内的所有线程都是可访问的，并且通过运行时函数对主机可访问。

常量内存变量跨多个源文件是可见的。因为设备只能读取常量内存，所以常量内存中的值必须使用以下运行时函数进行初始化：

1
cudaError_t cudaMemcpyToSymbol(const void *symbol, const void *src, size_t count, size_t offset, cudaMemcpyKind kind)

该函数通常用于将数据从主机内存拷贝到 GPU 的**常量内存（Constant Memory）**或全局符号（device 变量）中。

枚举变量kind指定了传输方向，默认情况下，kind是cudaMemcpyHostToDevice。

5.5.1 使用常量内存实现一维模板

在数值分析中，模板计算在几何点集合上应用函数，并用输出更新单一点的值。模板是求解许多偏微分方程算法的基础。在一维中，在位置x周围的九点模板会给这些位置上的值应用一些函数：$${x―4h，x―3h，x―2h，x―h，x，x＋h，x＋2h，x＋3h，x＋4h}$$

下图展示了一个九点模板。

一个九点模板的例子是实变量函数f在点x上一阶导数的第八阶中心差分公式。在本节中该公式将被作为一个示例模板。

在一维数组中对该公式的应用是对一个数据进行并行操作，在上述模板公式的例子下，系数c0、c1、c2和c3在所有线程中都是相同的并且不会被修改。这使它们成为常量内存最优的候选，因为它们是只读的，并将呈现一个广播式的访问模式：线程束中的每个线程同时引用相同的常量内存地址。

为实现一维模板计算，要使用共享内存来缓存数据，从而减少对全局内存的冗余访问。

RADIUS定义了点x两侧点的数量，这些点被用于计算x点的值。在这个例子中，为了形成一个九点模板，RADIUS被定义为4：x两侧各有4个点加上位置x的值。下图中，在每个块的左、右边界上各需要一个RADIUS个元素的光环。

因为有限差分系数被存储在常量内存中，并且这是由主机线程准备的，所以在核函数中访问它们就像访问数组一样简单。完整的核函数如下：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
__global__ void stencil_1d(float *in, float *out, int N)
{
    // shared memory
    __shared__ float smem[BDIM + 2 * RADIUS];

    // index to global memory
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    while (idx < N)
    {

        // index to shared memory for stencil calculatioin
        int sidx = threadIdx.x + RADIUS;

        // Read data from global memory into shared memory
        smem[sidx] = in[idx];

        // read halo part to shared memory
        if (threadIdx.x < RADIUS)
        {
            smem[sidx - RADIUS] = in[idx - RADIUS];
            smem[sidx + BDIM] = in[idx + BDIM];
        }

        // Synchronize (ensure all the data is available)
        __syncthreads();

        // Apply the stencil
        float tmp = 0.0f;

#pragma unroll
        for (int i = 1; i <= RADIUS; i++)
        {
            tmp += coef[i] * (smem[sidx + i] - smem[sidx - i]);
        }

        // Store the result
        out[idx] = tmp;

        idx += gridDim.x * blockDim.x;
    }
}

在常量内存中声明coef数组:

1
__constant__ float coef[RADIUS + 1];

然后使用cudaMemcpyToSymbol的CUDA API调用从主机端初始化的常量内存：

1
2
3
4
5
void setup_coef_constant (void)
{
    const float h_coef[] = {a0, a1, a2, a3, a4};
    CHECK(cudaMemcpyToSymbol( coef, h_coef, (RADIUS + 1) * sizeof(float)));
}

5.6 线程束洗牌指令

**线程束洗牌指令（Warp Shuffle Instructions）**它是提升高并发计算性能的核心利器之一，能够让线程直接在线程束内部交换数据，而无需通过共享内存或全局内存 。在没有洗牌指令之前，线程间交换数据必须通过共享内存。这涉及“存储、同步（__syncthreads()）、读取”三个步骤。

洗牌指令使得线程束中的线程彼此之间可以直接交换数据，而不是通过共享内存或全局内存来进行的。洗牌指令比共享内存有更低的延迟，并且该指令在执行数据交换时不消耗额外的内存

计算数组总和时，利用 __shfl_down_sync 可以让一个 Warp 在几条指令内完成局部求和，避免了共享内存的Bank冲突。

因为洗牌指令在线程束中的线程之间被执行，需要介绍一下束内线程（lane）的概念。一个束内线程指的是线程束内的单一线程。线程束中的每个束内线程是[0，31]范围内束内线程索引（laneindex）的唯一标识。线程束中的每个线程都有一个唯一的束内线程索引，并且同一线程块中的多个线程可以有相同的束内线程索引（就像同一网格中的多个线程可以有相同的threadIdx.x值一样）。

在一维线程块中，对于一个给定线程的束内线程索引和线程束索引可以按以下公式进行计算：

1
2
laneID = threadIdx.x % 32;
warpID = threadIdx.x / 32;

例如，线程块中的线程1和线程33都有束内线程ID 1，但它们有不同的线程束ID。对于二维线程块，可以将二维线程坐标转换为一维线程索引，并应用前面的公式来确定束内线程和线程束的索引。

5.6.1 线程束洗牌指令的不同形式

有两组洗牌指令：一组用于整型变量，另一组用于浮点型变量。每组有4种形式的洗牌指令。在线程束内交换整型变量，其基本函数标记如下：

1
int __shfl(int var, int srcLane, int width=warpSize);

var: 当前线程想要提供的变量值。

srcLane: 源线程的索引（Lane ID）。即你想从哪个线程获取数据。

width: 参与 Shuffle 操作的范围大小。它必须是 2 的幂且小于或等于 Warp 大小（通常是 32）。如果 width 是 32，则在整个 Warp 内进行数据交换。

例如：

1
int y = shfl(x, 3, 16);

当 width = 16 时，32 个线程被分为两组：

  1. 第一组 (Lane 0 ~ 15)：
    • 这是一个独立的通信范围。
    • srcLane = 3 指向这一组内的第 3 号线程。
    • 因此，线程 0 到 15 都会读取到线程 3 的 x 值。
  2. 第二组 (Lane 16 ~ 31)：
    • 这也是一个独立的通信范围。
    • 虽然参数写的是 3，但在第二组内部，索引是相对的。
    • 它的实际来源是该组的“第 3 个线程”，即逻辑索引为 16+3=1916 + 3 = 19 的线程。
    • 因此，线程 16 到 31 都会读取到线程 19 的 x 值。

当传递给shfl的束内线程索引与线程束中所有线程的值相同时，指令从特定的束内线程到线程束中所有线程都执行线程束广播操作。

1. __shfl_up (向上偏移)

核心逻辑：当前线程从比自己索引更小的线程获取数据。

1
int __shfl_up(int var, unsigned int delta, int width = warpSize);
  • 计算方式：srcLane = laneId - delta。
  • 行为：如果计算出的索引小于当前子组的边界，则该线程获得的值保持不变（通常是它自己的 var）。
  • 应用场景：计算前缀和（Prefix Sum / Scan）。例如，每个线程都加上它左边所有线程的和。

2.__shfl_down (向下偏移)

核心逻辑：当前线程从比自己索引更大的线程获取数据。

1
int __shfl_down(int var, unsigned int delta, int width = warpSize);
  • 计算方式：srcLane = laneId + delta。
  • 行为：如果计算出的索引超出了当前子组的边界（例如 laneId + delta >= 32），则该线程的值保持不变。
  • 应用场景：并行归约（Reduction）。
    • 示例：第一步让所有线程执行 __shfl_down(val, 16)，然后相加；第二步 __shfl_down(val, 8)… 这种方式比使用共享内存快得多，因为它不需要额外的存储开销和显式的同步屏障。

3.__shfl_xor (按位异或交换)

核心逻辑：通过对当前线程的 Lane ID 进行**按位异或（XOR）**运算来确定源线程。

1
int __shfl_xor(int var, int laneMask, int width = warpSize);
  • 计算方式：srcLane = laneId ^ laneMask。
  • 行为：这是一种“对称交换”。如果线程 A 从线程 B 取数据，那么线程 B 也会从线程 A 取数据。
  • 应用场景：蝶形交换算法（Butterfly Exchange）。
    • 在快速傅里叶变换（FFT）或排序网络（Sorting Networks）中极其常用。
    • 示例：当 laneMask = 1 时，相邻的线程（0和1, 2和3…）两两交换数据。

5.6.2 线程束内的共享数据

在本节中，会介绍几个有关线程束洗牌指令的例子,洗牌指令将被应用到以下3种整数变量类型中：

·标量变量

·数组

·向量型变量

1.test_shfl_broadcast

核心逻辑：验证将单一个线程的值分发给子组内所有线程的能力。

  • 操作：通常调用 __shfl_sync(mask, var, srcLane, width)。
  • 预期结果：子组内所有线程最终都持有索引为 srcLane 的线程在执行前所拥有的 var 值。
  • 用途：模拟参数广播，例如将某个计算出的权重发送给整个 Warp。

2.test_shfl_up

核心逻辑：验证数据向高索引方向的平移。

  • 操作：调用 __shfl_up_sync(mask, var, delta, width)。
  • 预期结果：线程 ii 获得线程 i−deltai - delta 的数据。
  • 边界处理：索引小于 deltadelta 的线程（如前 deltadelta 个线程）其值保持不变。
  • 图形化理解：数据向右滑动，左侧留空。

3.test_shfl_down

核心逻辑：验证数据向低索引方向的平移。

  • 操作：调用 __shfl_down_sync(mask, var, delta, width)。
  • 预期结果：线程 ii 获得线程 i+deltai + delta 的数据。
  • 边界处理：索引加上 deltadelta 后超过子组范围的线程，其值保持不变。
  • 用途：这是实现 Warp Reduction（规约） 的标准测试，通过不断减小 deltadelta（16, 8, 4, 2, 1）来汇总数据。

4.test_shfl_wrap

核心逻辑：这通常是一个自定义测试，模拟“循环移位”。

  • 操作：在标准的 Shuffle 指令中并没有直接的 wrap 函数，这通常通过 __shfl_sync 手动计算索引实现：srcLane = (laneId + offset) % width。
  • 预期结果：数据像转盘一样移动。例如，向左移 1 位，则线程 0 获得线程 1 的值，而线程 31 获得线程 0 的值。
  • 区别：它与 up / down 的最大区别在于首尾相连，没有数据丢失。

5.test_shfl_xor

核心逻辑：验证基于位运算的对称交换。

  • 操作：调用 __shfl_xor_sync(mask, var, laneMask, width)。
  • 预期结果：线程 ii 与线程 i⊕laneMaski \oplus laneMask 交换数据。
  • 特点：如果 laneMask 为 1，则线程 (0,1), (2,3) 互换；如果为 2，则 (0,2), (1,3) 互换。
  • 用途：测试蝴蝶变换（Butterfly Exchange）逻辑，是 FFT 算法的核心。

6. test_shfl_xor_array

核心逻辑：验证如何使用 Shuffle 指令高效地交换数组或结构体。

  • 技术背景：内置的 __shfl 系列函数原生只支持基本类型（如 int, float）。
  • 操作：测试通常会使用循环，或者利用 C++ 模板将一个较大的数据结构拆分为多个 int 或 float 寄存器，逐个进行 __shfl_xor 交换，最后再拼合。
  • 预期结果：整个数组或复杂的结构体在线程对之间完成了完整的搬运。

5.6.3 使用线程束洗牌指令的并行归约

一个线程块中可能有几个线程束。对于线程束级归约来说，每个线程束执行自己的归约。每个线程不使用共享内存，而是使用寄存器存储一个从全局内存中读取的数据元素：

1
int mySum = g_idata[idx];

线程束级归约作为一个内联函数实现：

1
2
3
4
5
6
7
8
9
10
__inline__ __device__ int warpReduce(int localSum)
{
    localSum += __shfl_xor(localSum, 16);
    localSum += __shfl_xor(localSum, 8);
    localSum += __shfl_xor(localSum, 4);
    localSum += __shfl_xor(localSum, 2);
    localSum += __shfl_xor(localSum, 1);

    return localSum;
}

完整的reduce-Shfl核函数如下

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
__global__ void reduceShfl (int *g_idata, int *g_odata, unsigned int n)
{
    // shared memory for each warp sum
    __shared__ int smem[SMEMDIM];

    // boundary check
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (idx >= n) return;

    // calculate lane index and warp index
    int laneIdx = threadIdx.x % warpSize;
    int warpIdx = threadIdx.x / warpSize;

    // blcok-wide warp reduce
    int localSum = warpReduce(g_idata[idx]);

    // save warp sum to shared memory
    if (laneIdx == 0) smem[warpIdx] = localSum;

    // block synchronization
    __syncthreads();

    // last warp reduce
    if (threadIdx.x < warpSize) localSum = (threadIdx.x < SMEMDIM) ?
        smem[laneIdx] : 0;

    if (warpIdx == 0) localSum = warpReduce(localSum);

    // write result for this block to global mem
    if (threadIdx.x == 0) g_odata[blockIdx.x] = localSum;
}
  1. 基础信息准备
  • __shared__ int smem[SMEMDIM]：在线程块内开辟共享内存，用于存储每个线程束计算出的部分和。
  • laneIdx：当前线程在 Warp 内的索引（0-31）。
  • warpIdx：当前线程属于线程块中的第几个 Warp。
  1. 第一级归约：线程束内规约 (Warp-Level)
1
int localSum = warpReduce(g_idata[idx]);
  • 每个线程从全局内存读取一个数据。
  • 调用 warpReduce（通常内部实现就是您之前问到的 __shfl_down_sync 循环）。
  • 结果：执行完这一行后，每个 Warp 的 0 号线程（laneIdx == 0）都持有该 Warp 内所有元素的总和。
  1. 中间结果传递 (Warp to Shared Memory)
1
2
if (laneIdx == 0) smem[warpIdx] = localSum;
__syncthreads();
  • 每个 Warp 的 0 号线程将自己算出的部分和写入共享内存 smem 的对应位置。
  • 调用 __syncthreads() 确保线程块内所有线程都完成了写入，这样后续读取才安全。
  1. 第二级归约：跨线程束汇总 (Block-Level)
1
2
3
4
5
if (threadIdx.x < warpSize) 
    localSum = (threadIdx.x < SMEMDIM) ? smem[laneIdx] : 0;

if (warpIdx == 0) 
    localSum = warpReduce(localSum);
  • 只激活第一个 Warp（threadIdx.x < warpSize）：让第一个 Warp 的线程去读取之前存入 smem 的各 Warp 部分和。
  • 最后一次 warpReduce：第一个 Warp 再次进行一次 Warp 内归约，将所有 Warp 的部分和加在一起。
  • 结果：现在，整个线程块的 0 号线程持有了该 Block 的最终总和。
  1. 结果输出
1
if (threadIdx.x == 0) g_odata[blockIdx.x] = localSum;
  • 由 0 号线程将该线程块的归约结果写入输出数组 g_odata 的对应位置（每个 block 输出一个值）。

5.7 总结

为了获得最大的应用性能，需要有一个能显式管理的内存层次结构。在C语言中，没有直接控制数据移动的方式。在本章中，介绍了不同CUDA内存层次结构类型，如共享内存、常量内存和只读缓存。介绍了当从共享内存中引入或删除数据时如何显式控制以显著提高其性能。还介绍了常量内存和只读缓存的行为，以及如何最有效地使用它们。