CUDA编程: CUDA内存管理（二）

CUDA 编程

摘要

详细讲解共享内存存储体、存储体冲突、访问模式以及内存填充等知识点。

CUDA 共享内存

GPU 的物理内存可以分为

  • 板载内存
  • 片上内存

全局内存就是板载内存，有较高的延时；共享内存就是较小的片上内存 ，有较低的延时。共享有比全局内存更高的带宽，可以把它当作一个可编程的缓存。共享内存通常的用途有

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

共享内存（Shared Memory, SMEM），在 GPU 中的位置如下图所示

每个 SM 都有一个小的内存池，这个内存池被当前正在该 SM 上执行的线程块中的所有线程所共享。SMEM 使同一个线程块中的线程能够互相协作，以重用片上数据，并可以大大降低核函数所需的全局内存带宽。由于 SMEM 中的内容是由应用程序显式管理的，所以是可编程管理的缓存

上图可以看到 SMEM 不需要经过 L1，相比 DRAM，延迟低 20~30 倍，带宽为 DRAM 的 10倍

在每个线程块被执行时会分配给它一些 SMEM，线程块执行完毕后 SMEM 释放，线程块和它的 SMEM 有相同的生命周期。每个线程束对 SMEM 的访问请求分为以下几种情况

  • 最好的情况是当前线程束中的每个线程都访问一个不冲突的共享内存，一个事务完成整个线程束的访问

  • 最坏的情况是有冲突访问，每个线程束的 32 个线程需要不同的 32 个事务来完成

  • 如果线程束内 32 个线程访问 SMEM 中的同一个地址，那么一个线程访问完后以广播的形式告诉其它线程

一个 SM 上的所有的正在执行的线程块共会划分有限的 SMEM 资源，所以核函数使用的共享内存越多，那么处于并发活跃状态的线程块就越少

下面将围绕避免 SMEM 中多个事务访问冲突的问题展开讨论

共享内存分配

可以动态的或静态的声明使用共享内存的变量。共享内存变量在核函数中声明，作用域就只在核函数中，在核函数外声明，对所有核函数来说作用域都是全局的，我们可以声明一维，二维和三维的共享内存数组

使用__shared__修饰符来声明共享内存变量，下面声明了使用共享内存的一维，二维和三维浮点数组

1
2
3
__shared__ float a[*];
__shared__ float b[*][*];
__shared__ float c[*][*][*];

这里的 * 必须是一个编译时确定的数字，不能是变量，如果共享内存的大小在编译时是未知的，也就是动态声明一个共享内存数组，使用extern关键字

1
extern __shared__ int d[];

并将所需的大小按字节数作为三重括号内的第三个参数，isize为数组的中的元素个数

1
kernel<<<grid, block, isize * sizeof(float)>>>(...);

注意这里的动态声明只支持一维数组

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

优化内存性能的着重指标就是

  • 延时
  • 带宽

上面提到，共享内存隐藏了全局内存延迟，并且大大提高了带宽，所以了解共享内存的原理和特性会让我们更为清晰地使用共享内存

共享内存存储体 (bank)

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

存储体冲突（bank conflict）

如果线程束对共享内存有操作，且在个存储体上只请求访问最多一次，那么就由一个内存事务来完成，如果线程束在任意一个存储体上请求访问大于一次，就会由多个内存事务来完成，称为存储体冲突（bank conflict）

bank conflict 会导致请求被重复执行，GPU 会将存储体冲突的请求访问分割到尽可能多的独立的无冲突事务中，而独立内存事务的数量会直接影响内存带宽

线程束访问共享内存时有以下三种模式

  • 并行访问，多地址访问多存储体，带宽利用率最高

    如下图所示是最完美的情况，线程束中每个线程对应一个存储体

    如下图所示为不规则的访问模式，并行却不冲突，带宽利用率也是最高

    如下图所示同样为不规则的访问模式，但如果线程访问的是同一个存储体中相同的地址，广播访问就不会冲突，如果线程访问的是同一个存储体中不同的地址，就会产生冲突

  • 串行访问，多地址访问同一存储体，就会有对应 32 个线程的 32 个事务，带宽利用率最差

  • 广播访问，单一地址读取单一存储体，线程束中所有的线程都读取同一存储体中相同的地址。一个内存事务执行后，那么被访问的字就会被广播到所有请求的线程中。虽然只有一个内存事务，但只有一小部分字节被读取，所以带宽利用率很差

访问模式

共享内存存储体宽度（bank widths）直接影响访问模式，也就是每个存储体（bank）在一个时钟周期内的带宽，在计算能力 1.x 的设备中 bank widths 为 2 字节（16 位），计算能力 2.x 的设备中为 4 字节（32 位wjg），计算能力 3.x 以上的设备中为 8 字节（64 位）

对于计算能力为 2.0 的设备来说，bank widths 为 32 位，如下图所示就是共享内存的存储体的访问模式，字节地址除以 4 转换为 4 字节索引，再模 32，将 4 字节索引转换为存储体索引

上面的操作对应如下公式

1
bank index = (Byte address ÷ 4 byte/bank) % 32 banks

现今的 GPU 同时支持 64 位模式和 32 位模式。如果为 64 位模式，由于 SMEM 只有 32 个 bank，所以每个 bank 中的地址会被逻辑分成两侧，每个时钟周期内的每个 bank 都有 64 位的带宽，公式中的byte/bank就是 8。如下图所示是 64 位模式的几种情况，这也解释了为什么相比 32 位模式，其更不容易引起冲突

  • 每个线程访问不同的 bank，无 bank conflict

  • 多个线程访问一个 bank 中同一侧的同一个地址，地址会广播到所有线程；两个线程访问同一个 bank，所以会无 bank conflict

  • 两个线程访问同一个存储体的同一侧，为 bank conflict

  • 同一个 bank 的左侧被两个线程同时访问了不同的地址，会导致三向的 bank conflict

在现今 GPU 的 32 位模式下，由于 GPU 的每个时钟周期都是 64 带宽，所以 bank 中 32 位的数据需要 2 个时钟周期才能凑够 64 位，这也就使得两个线程读同一个 bank 时，如果读取的两个地址索引分别在两个不同的时钟周期被传输，就不会产生冲突，例如两个线程可以读4-byte word index为 0 和 32 的两个地址。如下图所示

cuda_runtime.h 提供了如下函数设置当前共享内存访问模式

1
__host__ cudaError_t cudaDeviceSetSharedMemConfig ( cudaSharedMemConfig config )
  • config: 请求的缓存配置，枚举类型，有如下值
    • cudaSharedMemBankSizeDefault = 0: 设置 bank widths 为设备默认值
    • cudaSharedMemBankSizeFourByte = 1: 设置 bank widths 为 4 字节（32 位）
    • cudaSharedMemBankSizeEightByte = 2: 设置 bank widths 为 8 字节（64 位）

注意cudaDeviceSetSharedMemConfig函数在固定共享内存大小的设备上无作用

如下函数查询当前共享内存访问模式

1
__host__ __device__ cudaError_t cudaDeviceGetSharedMemConfig ( cudaSharedMemConfig ** pConfig )

pConfig: 返回的缓存配置，枚举类型，有如下值

  • cudaSharedMemBankSizeFourByte = 1: bank widths 为 4 字节（32 位）
  • cudaSharedMemBankSizeEightByte = 2: bank widths 为 8 字节（64 位）

在不同的核函数启动之间更改共享内存的配置，可能需要一个隐式的设备同步点，更改共享内存存储体的大小对性能有重大的影响。更大的 bank widths 可能有更高的带宽，也可能导致更多的 bank conflict，需要实验得出

下面的代码简单地使用了上面两个函数，仅供参考，这里笔主的显卡固定了共享内存大小，无法演示

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
#include <cuda_runtime.h>
#include <stdio.h>

void printSharedMemConfig(enum cudaSharedMemConfig &sharedMemConfig) {
    switch (sharedMemConfig) {
        case cudaSharedMemBankSizeDefault:
            printf("Shared memory config: Default\n");
            break;
        case cudaSharedMemBankSizeFourByte:
            printf("Shared memory config: Bank size of 4 bytes\n");
            break;
        case cudaSharedMemBankSizeEightByte:
            printf("Shared memory config: Bank size of 8 bytes\n");
            break;
    }
}

int main() {
    enum cudaSharedMemConfig sharedMemConfig;
    cudaDeviceGetSharedMemConfig(&sharedMemConfig);
    printSharedMemConfig(sharedMemConfig);
    cudaDeviceSetSharedMemConfig(cudaSharedMemBankSizeEightByte);
    cudaDeviceGetSharedMemConfig(&sharedMemConfig);
    printSharedMemConfig(sharedMemConfig);
    return 0;
}

内存填充（memory padding）

内存填充是避免存储单元冲突的一种方法。假设 5 个共享内存存储单元。如果所有线程访问 bank 0 的不同地址，那么会发生一个五向的存储单元冲突。解决这种存储单元冲突的一个方法是在每 5 个元素之后添加一个填充，改变从字到存储单元的映射，以错开访问每行数据

如上图所示，由于填充，之前所有属于 bank 0 的字，现在被传播到了不同的存储单元中。 填充的内存不能用于数据存储，其唯一的作用就是移动数据元素，以便将原来属于同一个存储单元中的数据分散到不同存储单元中。这样可以使得线程块可用的总共享内存的数量减少。 填充之后还需要根据前面的公式重新计算数组索引以确保能访问到正确的数据元素。例如下面的共享内存数组

1
__shared__ int a[5][4];

我们可以更改声明以还原图例中的情况

1
__shared__ int a[5][5];

配置共享内存

每个 SM 上有 64KB 的片上内存，SMEM 和 L1 共享这 64KB，并且可以配置，CUDA 为配置 L1 和 SMEM 提供以下两种方法

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

为当前设备设置首选缓存配置

1
__host__ cudaError_t cudaDeviceSetCacheConfig ( cudaFuncCache cacheConfig )
  • cacheConfig：请求的缓存配置，枚举类型

  • cudaFuncCachePreferNone = 0: 默认函数缓存配置，无优先级

  • cudaFuncCachePreferShared = 1: 首选更大的 SMEM 和更小的 L1 缓存

  • cudaFuncCachePreferL1 = 2: 首选较大的 L1 缓存和较小的 SMEM

  • cudaFuncCachePreferEqual = 3: 首选大小相同的 L1 缓存和 SMEM

使用上面哪种更好要根据核函数使用了多少 SMEM

  • SMEM 使用较多，那么首选更大的 SMEM
  • 更多的寄存器使用，那么首选较大的 L1

另一个函数是为当前核函数设置首选缓存配置

1
__host__ cudaError_t cudaFuncSetCacheConfig ( const void* func, cudaFuncCache cacheConfig )
  • func: 设备函数指针
  • cacheConfig: 请求的缓存配置

L1 和 SMEM 虽然都在同一个片上，但是与 SMEM 的 bank 不同，L1 通过缓存行进行访问。我们可以完全控制 SMEM，但 L1 的删除工作是硬件完成的

GPU使用不同的启发式算法来处理数据。在GPU上，数百个线程共享相同的 L1，数千个线程共享有网的 L2。因此，数据处理在 GPU上可能会发生的更频繁而且更不可预知，所以使用 SMEM 不仅可以显式管理数据，还可以保证 SM 的局部性

同步

同步是并行的重要机制，其主要目的就是防止冲突。同步基本方法有

  • 障碍，是所有调用线程等待其余调用线程达到的障碍点
  • 内存栅栏，所有调用线程必须等到全部内存修改对其余线程可见时才继续进行

首先需要理解 CUDA 采用的弱排序内存模型

弱排序内存模型

CUDA 允许编译器大幅优化源代码以加速程序运行效率，这就会导致内存访问的顺序被改变，也就是说 GPU 线程在不同的内存，比如 SMEM，全局内存，锁页内存或对等设备内存中，写入数据的顺序是不一定和这些数据在源代码中访问的顺序相同。当线程的写入顺序对其他线程可见的时候，它可能和写操作被执行的实际顺序不一致。如果指令之间相互独立，线程从不同内存中读取数据和指令的顺序也不一定相同。在这种不正确情况下，为了保持内存管理的可控，必须在代码中使用障碍和内存栅栏以防止冲突

显示障碍

CUDA 中，障碍点只对同一线程块内的线程执行，且只能设置在核函数中，使用如下函数设置一个障碍点

1
void __syncthreads();

__syncthreads()作为一个障碍点，保证在同一线程块内所有线程没到达此障碍点时，不能继续向下执行，也就是阻塞 block 直至 block 内的线程全都执行到这一行。且在同一线程块内，此障碍点之前的所有全局内存，共享内存操作，对后面的线程都是可见的。

__syncthreads()也可以解决同一线程块内，内存竞争的问题，保证执行的先后顺序

此外在条件语句中使用__syncthreads()，会导致无法预料的严重情况。如下面的代码，因块中的所有线程都没有达到相同的障碍点，会直接导致内核死锁

1
2
3
4
5
if (threadID % 2 == 0) {
	__syncthreads();
} else {
	__syncthreads();
}

但是__syncthreads()的局限就在于只能解决一个块内的线程同步，不能跨线程同步，线程块会以任何顺序，并行或串行地在任何 SM 上执行，线程块这种独立的特性使得 CUDA 在任意数量的核心中都是可扩展的，如果一个 CUDA 核函数要求线程块全局同步，那么只能结束核函数的运行来隐式的同步线程块

内存栅栏

内存栅栏能保证栅栏前的内核内存写操作对栅栏后的其他线程都是可见的，根据所需范围有以下三种栅栏

  • 线程块，在 block 内创建内存栅栏，阻塞线程束直至阻塞线程束发出的写操作完成，但由于阻塞线程束本身就是单指令多线程，该指令就没什么用

    1
    2
    3
    4
    5
    6
      void __threadfence_block();

    * 网格，在 grid 内创建内存栅栏，阻塞 grid 直至 grid 内的线程发出的读写操作完成，可以实现块间同步

      ```cpp
      void __threadfence();

  • 系统，可以跨系统创建内存栅栏，挂起调用的线程，以确保该线程对全局内存、锁页主机内存和其他设备内存中的所有写操作对全部设备中的线程和主机线程是可见的

    1
    void __threadfence_system();

Volatile 修饰符

在全局或共享内存中使用 volatile 修饰符声明一个变量，阻止编译器优化，可以防止这个变量存入缓存，这个变量的任何引用都会直接被编译到全局内存中，忽略缓存。举例如下

1
volatile float vfloat;

共享内存的数据布局

这部分我们通过研究如何组织共享内存的数据布局，以达到更少的 bank conflict 和最佳的性能

方形共享内存

SMEM 可以直接缓存方形维度的全局数据，如下图所示，字节地址与存储体地址的逻辑映射图，在每个维度假设有 32 个元素，且按行主序进行存储

如下静态声明一个二维共享内存变量

1
2
#define N 32
__shared__ int a[N][N];

我们可以用两种方式访问其中一个元素

1
2
3
4
5
// 行主序
a[threadIdx.y][threadIdx.x];

//列主序
a[threadIdx.y][threadIdx.x];

行主序和列主序哪个效率更高，这取决于线程与共享内存存储体的映射关系。在一个线程束中的线程由连续的threadIdx.x来确定，也就是说，threadIdx.y对应上图中的 Row 行，threadIdx.x对应上图中的 Bank 列，而每个 bank 和线程束中的每个线程对应，邻近线程在最内层数组维度上访问相邻的阵列单元。因此，相比列主序，行主序有更好的性能和更少的 bank conflict

行主序读写和列主序读写对比

下面的程序将全局线程索引值存入二维共享内存，再从共享内存中读取这些值并存到全局内存中，对比行主序和列主序

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
#include <cuda_runtime.h>
#include <stdio.h>

#define DIM_X 32
#define DIM_Y 32

__global__ void setRowReadRow(int * out) {
    __shared__ int tile[DIM_Y][DIM_X];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.y][threadIdx.x]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.y][threadIdx.x];
}

__global__ void setColReadCol(int * out)
{
    __shared__ int tile[DIM_Y][DIM_X];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.x][threadIdx.y]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.x][threadIdx.y];
}

int main() {
    enum cudaSharedMemConfig sharedMemConfig;
    cudaDeviceGetSharedMemConfig(&sharedMemConfig);
    printf("%d\n", sharedMemConfig);
    
    int nElem = DIM_X * DIM_Y;
    int nByte = sizeof(int)*nElem;
    int * out;
    cudaMalloc((int**)&out,nByte);
    dim3 block(DIM_Y, DIM_X);
    dim3 grid(1,1);
    
    setRowReadRow<<<grid,block>>>(out);
    setColReadCol<<<grid,block>>>(out);
    
    cudaDeviceSynchronize();
    return 0;
}

在 4 字节访问模式下，因为相邻线程引用相邻字，可以看到按行访问使用的时间比列访问少了很多

使用 ncu --metrics l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum 命令可以获取核函数运行阶段的共享内存加载事务数，ncu --metrics l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum命令可以获取共享内存存储事务数，关于ncu --metrics更多参数的具体解释可以看看这篇博客

1
2
3
4
5
6
setRowReadRow(int*)
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum	32
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum	32
setColReadCol(int*)
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum	1024
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum	1024

在对共享内存存储事务数/执行共享内存的访问次数，也就是二维共享内存的行DIM_Y，这里为 32，得到每次访问共享内存时的存储事务数 32/32 = 1，不会产生 bank conflict。但是在setColReadCol中每次访问共享内存时的存储事务数为 1024/32 = 32，会有 32 路 bank conflict，对应DIM_Y=32，就是因为setRowReadRow是邻近线程在最内层数组维度上访问相邻的阵列单元

下面的核函数为按行主序写和按列主序读

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
__global__ void setRowReadCol(int * out) {
    __shared__ int tile[DIM_X][DIM_X];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.y][threadIdx.x]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.x][threadIdx.y];
}

__global__ void setColReadRow(int * out) {
    __shared__ int tile[DIM_Y][DIM_X];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.x][threadIdx.y]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.y][threadIdx.x];
}

执行上面两个核函数，可以看到冲突情况符合我们的理论，即邻近线程在最内层数组维度上访问相邻的阵列单元会减少冲突

1
2
3
4
5
6
setRowReadCol(int*)
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum	1024
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum	32
setColReadRow(int*)
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum	32
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum	1024

设矩阵 size(4,4) 执行上面的四个核函数，输出out的值可以看到setRowReadCol和setColReadRow会对数组转置，这为之后我们将会讲到的转置算法作了基础

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
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
#include <cuda_runtime.h>
#include <stdio.h>

#define DIM_X 4
#define DIM_Y 4

__global__ void setRowReadRow(int * out) {
    __shared__ int tile[DIM_Y][DIM_X];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.y][threadIdx.x]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.y][threadIdx.x];
}

__global__ void setColReadCol(int * out)
{
    __shared__ int tile[DIM_Y][DIM_X];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.x][threadIdx.y]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.x][threadIdx.y];
}

__global__ void setRowReadCol(int * out) {
    __shared__ int tile[DIM_X][DIM_X];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.y][threadIdx.x]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.x][threadIdx.y];
}

__global__ void setColReadRow(int * out) {
    __shared__ int tile[DIM_Y][DIM_X];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.x][threadIdx.y]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.y][threadIdx.x];
}

int main() {
    enum cudaSharedMemConfig sharedMemConfig;
    cudaDeviceGetSharedMemConfig(&sharedMemConfig);
    printf("%d\n", sharedMemConfig);
    
    int nElem = DIM_X * DIM_Y;
    int nByte = sizeof(int)*nElem;
    int * out_h, * out_d;
    out_h = (int*)malloc(nByte);
    cudaMalloc((int**)&out_d,nByte);
    dim3 block(DIM_Y, DIM_X);
    dim3 grid(1,1);
    
    printf("setRowReadRow: ");
    setRowReadRow<<<grid,block>>>(out_d);
    cudaMemcpy(out_h, out_d, nByte, cudaMemcpyDeviceToHost);
    for(int i=0; i<nElem; i++) {
        printf("%d ", out_h[i]);
    }
    printf("\n");

    printf("setColReadCol: ");
    setColReadCol<<<grid,block>>>(out_d);
    cudaMemcpy(out_h, out_d, nByte, cudaMemcpyDeviceToHost);
    for(int i=0; i<nElem; i++) {
        printf("%d ", out_h[i]);
    }
    printf("\n");
    printf("setRowReadCol: ");
    setRowReadCol<<<grid,block>>>(out_d);
    cudaMemcpy(out_h, out_d, nByte, cudaMemcpyDeviceToHost);
    for(int i=0; i<nElem; i++) {
        printf("%d ", out_h[i]);
    }
    printf("\n");
    printf("setColReadRow: ");
    setColReadRow<<<grid,block>>>(out_d);
    cudaMemcpy(out_h, out_d, nByte, cudaMemcpyDeviceToHost);
    for(int i=0; i<nElem; i++) {
        printf("%d ", out_h[i]);
    }
    printf("\n");
    
    cudaDeviceSynchronize();
    return 0;
}
1
2
3
4
setRowReadRow: 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 
setColReadCol: 0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 
setRowReadCol: 0 4 8 12 1 5 9 13 2 6 10 14 3 7 11 15 
setColReadRow: 0 4 8 12 1 5 9 13 2 6 10 14 3 7 11 15

下面给出动态声明版本，前面提到过，动态共享内存数组只能是一维的，且要将所需大小按字节数作为核函数三重括号内的第三个参数

1
2
3
4
5
6
7
8
9
__global__ void setRowReadColDyn(int * out) {
    extern __shared__ int tile[];
    unsigned int row_idx=threadIdx.y*blockDim.x+threadIdx.x;
    unsigned int col_idx=threadIdx.x*blockDim.y+threadIdx.y;

    tile[row_idx]=row_idx;
    __syncthreads();
    out[row_idx]=tile[col_idx];
}
1
setRowReadColDyn<<<grid, block, DIM_X * DIM_Y * sizeof(int)>>>(out);
1
2
3
setRowReadColDyn(int*)
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum	1024
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum	32

memory padding

为了解决setRowReadCol，setColReadRow核函数的 bank conflict，我们要根据具体的数据分布来填充内存，在静态声明中，只需要将填充的列添加到二维共享内存分配中就可以了，代码如下

1
2
3
4
5
6
7
8
9
#define PAD 1
__global__ void setRowReadColPad(int * out) {
    __shared__ int tile[DIM_Y][DIM_X+PAD];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;

    tile[threadIdx.y][threadIdx.x]=idx;
    __syncthreads();
    out[idx]=tile[threadIdx.y][threadIdx.x];
}
1
setRowReadColPad<<<grid,block>>>(out);

在动态声明中，由于需要执行二维线程索引到一维线程索引的转换，所以对于每一行，都要跳过填充的部分，代码如下

1
2
3
4
5
6
7
8
9
10
11
#define PAD 1
__global__ void setRowReadColDynPad(int * out) {
    extern __shared__ int tile[];
    unsigned int row_idx=threadIdx.y* (blockDim.x + PAD) +threadIdx.x;
    unsigned int col_idx=threadIdx.x* (blockDim.y + PAD) +threadIdx.y;
    unsigned int g_idx=threadIdx.y*blockDim.x +threadIdx.x;

    tile[row_idx]=g_idx;
    __syncthreads();
    out[g_idx]=tile[col_idx];
}
1
setRowReadColDynPad<<<grid, block, (DIM_X + PAD) * DIM_Y * sizeof(int)>>>(out);

使用 ncu 工具可以看到每次访问共享内存请求的事务数量为 1，无 bank conflict

1
2
3
4
5
6
setRowReadColPad(int*)
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum	32
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum	32
setRowReadColDynPad(int*)
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum	32
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum	32

矩形共享内存

矩形共享内存和方形共享内存非常相似，不同的地方在于线程索引的要先映射为一维，保证访问是合并的

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

再通过二维方式访问，这里对原矩阵按列读取，irow 和 icol 对应的是转置后矩阵中的坐标

1
2
unsigned int icol=idx%blockDim.y;
unsigned int irow=idx/blockDim.y;

下面的程序将全局线程索引值存入二维共享内存，再从共享内存中读取这些值并存到全局内存中，对比行主序和列主序

1
2
3
4
5
6
7
8
9
10
11
12
#define DIM_X_RECT 32
#define DIM_Y_RECT 16

__global__ void setRowReadColRect(int * out) {
    __shared__ int tile[DIM_Y_RECT][DIM_X_RECT];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;
    unsigned int icol=idx%blockDim.y;
    unsigned int irow=idx/blockDim.y;
    tile[threadIdx.y][threadIdx.x]=idx;
    __syncthreads();
    out[idx]=tile[icol][irow];
}

使用 ncu 工具可以看到加载操作有 256/16 = 16 路冲突，而存储操作没有冲突

1
2
3
setRowReadColRect(int*)
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum	256
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum	16

为了解决 bank conflict，下面给出 memory padding 版本的核函数，这里的PAD_RECT=2是因为将长方形矩阵一行有 16 个元素，为了满足 32 个存储体的数量，每次会访问两行的数据，所以要填充对应两行的 2 个元素以错开访问，而方形矩阵一行有 32 个元素，匹配 32 个存储体的数量，每行只需要填充一个数据即可错开访问

1
2
3
4
5
6
7
8
9
10
11
#define PAD_RECT 2

__global__ void setRowReadColRectPad(int * out) {
    __shared__ int tile[DIM_Y_RECT][DIM_X_RECT+PAD_RECT];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;
    unsigned int icol=idx%blockDim.y;
    unsigned int irow=idx/blockDim.y;
    tile[threadIdx.y][threadIdx.x]=idx;
    __syncthreads();
    out[idx]=tile[icol][irow];
}
1
2
3
setRowReadColRectPad(int*)
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum	16
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum	16

下面再给出动态声明和 memory 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
#define DIM_X_RECT 32
#define DIM_Y_RECT 16
#define PAD_RECT 2

__global__ void setRowReadColRectDyn(int * out) {
    extern __shared__ int tile[];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;
    unsigned int icol=idx%blockDim.y;
    unsigned int irow=idx/blockDim.y;
    unsigned int col_idx=icol*blockDim.x+irow;
    tile[idx]=idx;
    __syncthreads();
    out[idx]=tile[col_idx];
}

__global__ void setRowReadColRectDynPad(int * out) {
    extern __shared__ int tile[];
    unsigned int idx=threadIdx.y*blockDim.x+threadIdx.x;
    unsigned int icol=idx%blockDim.y;
    unsigned int irow=idx/blockDim.y;
    unsigned int row_idx=threadIdx.y*(PAD_RECT+blockDim.x)+threadIdx.x;
    unsigned int col_idx=icol*(PAD_RECT+blockDim.x)+irow;
    tile[row_idx]=idx;
    __syncthreads();
    out[idx]=tile[col_idx];
}
1
2
3
4
dim3 block_rect(DIM_X_RECT,DIM_Y_RECT);
dim3 grid_rect(1,1);
setRowReadColRectDyn<<<grid_rect,block_rect, DIM_X_RECT * DIM_Y_RECT*sizeof(int)>>>(out);
setRowReadColRectDynPad<<<grid_rect,block_rect, (DIM_X_RECT+PAD_RECT) * DIM_Y_RECT*sizeof(int)>>>(out);
1
2
3
4
5
6
setRowReadColRectDyn(int*)
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum	256
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum	16
setRowReadColRectDynPad(int*)
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_ld.sum	16
	l1tex__data_pipe_lsu_wavefronts_mem_shared_op_st.sum	16
- ETX   Thank you for reading -
  • Copyright: All posts on this blog except otherwise stated, All adopt CC BY-NC-ND 4.0 license agreement. Please indicate the source of reprint!