【CUDA 】第5章 共享内存和常量内存——5.2 共享内存的数据分布（2）

待解决的问题：5.2.2.1如何看用了多少事务？？？【大概已解决】

第五章 共享内存和常量内存

5.2 共享内存的数据分布

前提知识：ncu看shared_load_transactions_per_request指标

shared_load_transactions_per_request指标的含义：每个共享内存加载请求（Load Request）对应的实际内存事务（Transactions）数量。理想值为 1，表示每个请求只需 1 次事务；若值 >1，表明存在 内存访问冲突（Bank Conflict） 或 未合并访问（Uncoalesced Access）。

通过查找资料可知，旧的指标与ncu的对照，如下图所示：

链接: 指标与ncu的对照

找到对应shared_load_transactions的指标，值为256，如下图所示：

再用这个除总请求数应该就是结果。

所以，shared_load_transactions_per_request值应该为256/16=16。
表明每个请求需要16次事务。

5.2.2 矩形共享内存（shared mem）

矩形共享内存是更普遍的情况，行列数不相同；方形共享内存是特殊的，行列数相同。

__shared__ int tile[Row][Col];

当执行转置操作时，不能像方形一样简单的转换数组的线程坐标，这样做在矩形里面会造成内存访问冲突，需要基于矩阵维度重新计算访问索引，实现之前的核函数。

测试一个矩形共享内存数组，每行32个元素（32列），每列16个元素（16行）。
在宏定义中定义维度：

#define BDINX 32//每行32个元素（32列）
#define BDINY 16//每列16个元素（16行）__shared__ int tile[BDINY][BDINX];

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

dim3 block (BDINY, BDINX);
dim3 grid(1, 1);

5.2.2.1 行主序访问与列主序访问

把之前的核函数也在方形共享内存的情况下使用：

__global__ void setRowReadRow(int *out)
__global__ void setColReadCol(int *out)

代码复制上一节的代码即可，但是只需要上面两个核函数。

使用nvprof指标检查存储体冲突的结果。

cudaC/unit5/./5-2checkSmemRectangle at device 0: NVIDIA GeForce RTX 3090 with Bank Mode:4-Byte <<< grid (1,1) block (32,16)>>>
==PROF== Profiling "setColReadCol(int *)" - 0: 0%....50%....100% - 40 passes
==PROF== Profiling "setRowReadRow(int *)" - 1: 0%....50%....100% - 40 passes
==PROF== Disconnected from process 3732277
==PROF== Report: /home/zyn/cudaC/unit5/5-2checkSmemRectangle.ncu-rep

问题：不知道如何用ncu看使用了多少事务

书上结果说，存储体宽度为8个字，16个存储体元素安排到8个存储体中，有一个8路冲突。
实际运行：读写都16路冲突。

5.2.2.2 行主序写+列主序读

本节的核函数使用矩形共享内存数组，按行主序写+列主序读。
使用共享内存执行矩阵转置，通过最大化低延迟的加载和存储来提高性能，并且合并全局内存访问。

二维共享内存块的声明：

__shared__ int tile[BDIMY][BDIMX];

内核有3个内存操作：
①写行–每个线程束的共享内存行，避免存储体冲突
②读列–每个线程束的共享内存列，完成矩阵转置
③合并访问–写入每个线程束的全局内存行

计算正确的共享和全局内存访问的步骤：
①二维线程索引转换为一维线程索引.

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

这个一维行主序的映射可以确保全局内存访问是合并的。（全局内存访问合并：当warp访问连续的内存地址时，GPU把多个线程的访问请求合并为少数次内存事务，来最大化显存带宽利用率）

②计算转置矩阵中的新坐标。因为输出的全局内存的数据是转置过。

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

通过把全局线程ID存储到二维共享内存块来初始化共享内存块。

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

共享内存中的数据是从0到BDIMX*BDIMY-1线性存储的，因为每个线程束对共享内存执行了行主序写，因此写操作期间没有存储体冲突。

③用之前计算的坐标 访问 转置过的共享内存数据。通过交换过的irow和icol访问共享内存，用一维线程ID向全局内存写入转置数据。
线程束从共享内存的一列中读数据，并对全局内存执行合并写入。

out[idx] = tile[icol][irow];

完整核函数代码：

//按行存按列读
__global__ void setRowReadCol(int *out){__shared__ int tile[BDIMY][BDIMX];//从二维线程索引映射到————》线性内存unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;//把idx转换成转置后的坐标（row,col）unsigned int irow = idx / blockDim.y;unsigned int icol = idx % blockDim.y;tile[threadIdx.y][threadIdx.x] = idx;//按行写__syncthreads();out[idx] = tile[icol][irow];//按列读
}

问题：不知道如何用ncu看使用了多少事务

书上结果说，有一个8路冲突，矩阵中所有元素都被转置。
实际结果：读写都没有冲突。

5.2.2.3 动态声明的共享内存

动态共享内存只能被声明为一维数组，当按行写+按列读的时候，把二位线程坐标转换为——》一维共享内存索引需要一个新的索引：

unsigned int col idx = icol * blockDim.x + irow;

icol对应线程块中最内层的访问维度，这种转换以列主序访问共享内存，会造成存储体冲突。

完整核函数代码：

//动态声明共享内存+按行写按列读————》对应主函数中内核启动必须指定共享内存大小
__global__ void setRowReadColDyn(int *out){//动态声明共享内存extern __shared__ int tile[];//只能被声明为一维数组unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;//从线程索引映射到————》全局内存索引unsigned int irow = idx / blockDim.y;//把idx转换成转置后的坐标（row,col）unsigned int icol = idx % blockDim.y;//转回共享内存idx来访问转置元素unsigned int col_idx = icol * blockDim.x + irow;//这里乘blockDim.xtile[idx] = idx;__syncthreads();out[idx] = tile[col_idx];
}

对应主函数修改：

    //动态声明共享内存+按行写按列读cudaMemset(d_C, 0, nBytes);setRowReadColDyn<<<grid, block, BDIMX * BDIMY * sizeof(int)>>>(d_C);cudaMemcpy(gpuRef, d_C, nBytes, cudaMemcpyDeviceToHost);if(iprintf) printData("set row read col   \n", gpuRef, nx*ny);

问题：不知道如何用ncu看使用了多少事务

书上结果说，写操作无冲突，读操作有一个8路冲突。动态分配共享内存不会影响存储体冲突。
实际结果：写操作无冲突，读操作有一个16路冲突。

5.2.2.4 填充静态动态声明的共享内存

矩形共享内存也可以用共享内存填充来解决存储体冲突，但是对于Kepler架构必须计算出填充多少元素。
为了便于编程，用宏来定义填充列的数量

#define NPAD 2

填充的静态共享内存被声明如下：

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

只有静态共享内存填充不同，剩下的核函数与setRowReadCol相同。

完整核函数代码：

//填充静态声明共享内存+按行写按列读
__global__ void setRowReadColPad(int *out){//静态共享内存填充__shared__ int tile[BDIMY][BDIMX + IPAD];unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;unsigned int irow = idx / blockDim.y;unsigned int icol = idx % blockDim.y;tile[threadIdx.y][threadIdx.x] = idx;//按行写__syncthreads();out[idx] = tile[icol][irow];//按列读
}

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

该内核使用矩形共享内存区域，因为填充的共享内存和全局内存大小有所不同，所以在内核中每个线程必须保留3个索引：
①row_idx：填充共享内存的行主序。可以访问单一的矩阵行
②col_idx：填充共享内存的列主序。可以访问单一的矩阵列
③g_idx：线性全局内存索引，可以对全局内存进行合并访问

//线程索引————>全局内存索引
unsigned int g_idx = threadIdx.y * blockDim.x + threadIdx.x;//idx————>转职后的的坐标（row,col）
unsigned int irow = g_idx / blockDim.y;
unsigned int icol = g_idx % blockDim.y;
unsigned int row_idx = threadIdx.y * (blockDim.x + IPAD) + threadIdx.x;//填充后的行索引//转回共享内存idx来访问转置后的元素
unsigned int col_dix = icol * (blockDim.x + IPAD) + irow;

完整核函数代码：

//填充动态声明共享内存+按行写按列读
__global__ void setRowReadColDynPad(int *out){//动态共享内存extern __shared__ int tile[];//只能被声明为一维数组unsigned int g_idx = threadIdx.y * blockDim.x + threadIdx.x;//线程索引————>全局内存索引unsigned int irow = g_idx / blockDim.y;//idx————>转职后的的坐标（row,col）unsigned int icol = g_idx % blockDim.y;unsigned int row_idx = threadIdx.y * (blockDim.x + IPAD) + threadIdx.x;//填充后的行索引unsigned int col_idx = icol * (blockDim.x + IPAD) + irow;//转回共享内存idx来访问转置后的元素tile[row_idx] = g_idx;//按行写__syncthreads();out[g_idx] = tile[col_idx];//按列读
}

对应主函数中核函数调用同上。

5.2.2.6 矩形共享内存内核性能比较

1、运行时间对比：

核函数用共享内存填充，可以消除存储体冲突提高性能（5和6）。使用动态共享内存的核函数会有一些消耗（4和6）。

2、shared_load_transactions_per_request指标计算

request指标：所有函数shared load requests和shared store requests都是16

函数名	shared_load_transactions_per_request	shared_store_transactions_per_request
setColReadCol	16	16
setRowReadRow	1	1
setRowReadCol	16	1
setRowReadColDyn	16	1
setRowReadColPad	1	1
setRowReadColDynPad	1	1

可以看到两个填充后的核函数都是1，每个请求只需 1 次事务，没有冲突。

但是值得注意的是，第二个核函数按行写按行读也没有冲突，因为它本身就是行主序，访问不同的列，所以不会出现冲突。（与书上结果不同，大概原因是因为书上用K40显卡，K40Bank宽度为4字节+无Bank冲突合并优化，RTX 4090 bank宽度为8字节+支持Bank访问合并）