方 形 共 享 内 存
使用共享内存可以直接缓存具有方形维度的全局数据。方形矩形的简单维度可以很容易从二维线程索引中计算出一维内存偏移。下图显示了一个共享内存块，它在每个维度有32个元素，且按行主序进行存储。上部的图显示了一维数据布局的实际排序，下部的图显示了带有4字节数据元素和存储体映射的二维共享内存逻辑视图。
使用下面的语句静态声明一个二维共享内存变量：__shared__ int tile[N][N];因为这个共享内存块是方形的，所以可选择一个二维线程块访问它，在x或者y维度上通过相邻线程访问临近元素：1. tile[threadIdx.y][threadIdx.x]; 2. tile[threadIdx.x][threadIdx.y]；其中，第一存取模式将比第二存取模式呈现出更好的性能和更少的存储体冲突，因为邻近线程在最内存数组维度上访问相邻的阵列单元。
行主序访问和列主序访问
考虑一个例子，在例子中网格有一个二维线程块，块中每个维度包含32个可用的线程。代码如下：
#define BDIMX 32
#define BDIMY 32
//因为相同线程束中的线程有连续的threadIdx.x值，并且可以使用
//threadIdx.x索引共享内存数组tile的最内层维度，所以核函数无存储体冲突
__global__ void setRowReadRow(int *out)
{
__shared__ int tile[BDIMY][BDIMX];
unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
tile[threadIdx.y][threadIdx.x] = idx;
__synthreads();
out[idx] = tile[threadIdx.y][threadIdx.x];
}
//在数据分配给共享内存块时交换threadIdx.y和threadIdx.x，线程束的内存将会按列主序访问。
//每个共享内存的加载和存储将导致Fermi装置中有32路存储体冲突，导致Kepler装置中有16路存储体冲突。
__global__ void setColReadCol(int *out)
{
__shared__ int tile[BDIMX][BDIMY];
unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
tile[threadIdx.x][threadIdx.y] = idx;
__synthreads();
out[idx] = tile[threadIdx.x][threadIdx.y];
}
结果表明按行访问共享可以提高性能，因为相邻线程引用相邻字。在setRowReadRow核函数中，线程束的存储和加载请求由一个事务来完成，而相同的请求在setColReadCol核函数中由16个事务完成。
按行主序写和按列主序读
下面的核函数实现了共享内存中按行主序写入和按列主序读取，按行主序写入共享内存是将线程索引的最内层维度作为二维共享内存块的列索引来实现：tile[threadIdx.y][threadIdx.x] = idxl按列主序在共享内存中给全局内存赋值，这是在引用共享内存时交换两个线程索引实现的：out[idx] = tile[threadIdx.x][threadIdx.y];，下图显示了两个内存操作，它们使用了简化的五存储体共享内存实现。
内核代码如下：
//存储操作是无冲突的，但是加载操作显示有16路冲突
__global__ void setRowReadCol(int *out)
{
__shared__ int tile[BDIMY]BDIMX];
unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
tile[threadIdx.y][threadIdx.x] = idx;
__synthreads();
out[idx] = tile[threadIdx.x][threadIdx.y];
}
动态共享内存
可以动态声明共享内存，从而实现这些相同的核函数，可以在核函数外声明动态共享内存，使它的作用域为整个文件，也可以在核函数内声明动态共享内存，将其作用域限制在该内核之中。动态共享内存必须被声明为一个未定大小的一维数组，因此，需要基于二维线程索引来计算内存访问索引。因为要在这个核函数中按行主序写入，按列主序读取，所以需要保留以下两个索引：1.row_idx:根据二维县城索引计算出的一维行主序内存偏移量；2.col_idx:根据二维线程索引计算出的一维主序内存偏移量。
核函数如下所示
//写操作无冲突，读操作会报告一个16路径冲突
__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;
__synthreads();
out[row_idx] = tile[col_idx];
}
填充静态声明的共享内存
填充数组是避免存储体冲突的一种方法。填充静态声明的共享内存很简单，只需简单地将一列添加到二维共享内存分配中，代码如下所示：__shared__ int tile[BDIMY][BDIMX + 1];下面的核函数是setRowReadCol核函数的修改版，setRowReadCol按列主序读取时报告了16路冲突，通过在每行添加一个元素，列元素便分布在了不同的存储体中，因此读和写操作都是无冲突的。
__global__ void setRowReadColPad(int *out)
{
__shared__ int tile[BDIMY][BDIMX + IPAD];
unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
tile[threadIdx.y][threadIdx.x] = idx;
__synthreads();
out[idx] = tile[threadIdx.x][threadIdx.y];
}
填充动态声明的共享内存
填充动态声明的共享内存数组更加复杂。当执行从二维线程索引到一维内存索引的索引转换时，对于每一行必须跳过一个填充的内存空间，代码如下：
unsigned int row_idx = threadIdx.y * (blockDim.x + 1) + threadIdx.x;
unsigned int col_idx = threadIdx.x * (blockDim.x + 1) + threadIdx.y;
下图显示了这些内存索引计算，这些计算使用了一个简化的五存储体共享内存实现：
因为在以下核函数中用于存储数据的全局内存小于填充的共享内存，所以需要3个索引：一个索引用于按行主序写入共享内存，一个索引用于按列主序读取共享内存，一个索引用于未填充的全局内存的合并访问，代码如下：
__global__ void setRowReadColDynPad(int *out)
{
extern __shared__ int tile[];
unsigned int row_idx = threadIdx.y * (blockDim.x + IPAD) + threadIdx.x;
unsigned int col_idx = threadIdx.x * (blockDim.x + IPAD) + threadIdx.y;
unsigned int g_idx = threadIdx.y * blockDim.x + threadIdx.x;
tile[row_idx] = g_idx;
__synthreads();
out[g_idx] = tile[col_idx];
}
方形共享内存内核性能比较
到目前为止，从所有执行过的内核运行时间可以看出：1、使用填充的内核函数可提高性能，因为它减少了存储体冲突；2.带有动态声明共享内存的内核增加了少量消耗。
矩 形 共 享 内 存
矩形共享内存是一个更普遍的二维共享内存，在矩形共享内存中数组的行与列的数量并不相同，当执行一个转置操作时，不能像在方形共享内存中一样，只是通过简单的转换来引用矩形数组的线程坐标。当使用矩形共享内存时，这样做会导致内存访问冲突。需要基于矩阵维度重新计算访问索引，以重新实现之前描述的核函数。
一般情况下，需要测试一个矩形共享内存数组，其每行有32个元素，每列有16个元素，在下面的宏中定义了维度：
#define BDIMX 32
#define BDIMY 16
按行主序访问和列主序访问
与前面方形矩阵的按行主序访问和列主序访问类似。
按行主序写和按列主序读
接下来将实现一个核函数，该核函数使用一个矩形共享内存数组，按行主序写入共享内存，并按列主序读取共享内存。这个内核在现实的应用程序中是可用的。它使用共享内存执行矩阵转置，通过最大化低延迟的加载和存储来提高性能，合并全局内存访问。代码如下：
__global__ void setRowReadCol(int *out)
{
__shared__ int tile[BDIMY]BDIMX];
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;
__synthreads();
out[idx] = tile[icol][irow];
}
动态共享内存
因为动态共享内存只能被声明为一维数组，当按照行写入和按列读取时，将二维线程坐标转换为一维共享内存索引需要一个新的索引：unsigned int col idx = icol * blockDim.x + irow;因为icol对应于线程块中最内层的维度，所以这种转换以列主序访问共享内存，这会导致存储体冲突。核函数代码如下：
//写操作无冲突，读操作会报告一个8路径冲突，动态分配共享内存不会影响存储体冲突
__global__ void setRowReadColDyn(int *out)
{
extern __shared__ int tile[];
unsigned int idx = threadIdx.y * blockDim.x + threadIdx.x;
unsigned int irow = idx / blockDim.y;
unsigned int icol = idx % blockDim.y;
unsigned int col_idx = icol * blockDim.x + irow;
tile[idx] = idx;
__synthreads();
out[idx] = tile[col_idx];
}
填充静态声明的共享内存
对于矩形共享内存，还可以使用共享内存填充来解决存储体冲突，然而，对于Kepler设备，必须计算出需要多少填充元素。为了便于编程，使用宏定义每一行添加的填充列的数量：#define NPAD 2,填充的静态共享内存被声明如下：__shared__ int tile[BDIMY][BDIMX + NPAD];具体的核函数如下所示：
__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;
__synthreads();
out[idx] = tile[icol][irow];
}
填充动态声明的共享内存
填充技术还可以应用于动态共享内存的内核中，该内核使用矩形共享内存区域。因为填充的共享内存和全局内存大小会有所不同，所以在内核中每个线程必须保留3个索引：
1.row_idx:填充共享内存的行主序索引。使用该索引，线程束可以访问单一的矩阵行；
2.col_idx:填充共享内存的列主序索引。使用该索引，线程束可以访问单一的矩阵列；
3.g_idx:线性全局内存索引。使用该索引，线程束可以对全局内存进行合并访问。
这些索引是用以下代码计算出来的：
unsigned int g_idx = threadIdx.y * blockDim.x + threadIdx.x;
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;
unsigned int col_idx = 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;
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;
tile[row_idx] = g_idx;
__synthreads();
out[g_idx] = tile[col_idx];
}
矩形共享内存内核性能的比较
在一般情况下，核函数使用共享内存填充消除存储体冲突以提高性能，使用动态共享内存的核函数会显示有少量的消耗（与前面一样）。
1239
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
