使用共享内存也能帮助避免对未合并的全局内存的访问，矩阵转置就是一个典型的例子：读操作被自然合并，但写操作是按照交叉访问的。其中交叉访问是全局内存中最糟糕的访问模式，因为它浪费总线带宽，在共享内存的帮助下，可以先在共享内存中进行转置操作，然后再对全局内存进行合并写操作。接下来将介绍使用多个线程块对基于交叉的全局内存访问重新排序到合并访问。
基 准 转 置 内 核
  作为基准，下面的核函数是一个仅使用全局内存的矩阵转置的朴素实现：

__global__ void naiveGmem(float *out, float *in, const int nx, const int ny)
{
   unsigned int ix = blockIdx.x * blockDim.x + threadIdx.x;
   unsigned int iy = blockIdx.y * blockDim.y + threadIdx.y;
   
   if(ix < nx && iy < ny)
   {
      out[ix * ny + iy] = in[iy * nx + ix];
   }
}

使 用 共 享 内 存 的 矩 阵 转 置
  为了避免交叉全局访问，可以使用二维共享内存来缓存原始矩形的数据。从二维共享内存中读取的一列可以被转移到转置矩阵行中，它被存储在全局内存中。虽然朴素实现将导致共享内存存储体冲突，但是这个结果将比未合并的全局访问好得多。下图显示了在矩阵转置中是如何使用共享内存的。

  下面的核函数实现了使用共享内存的矩阵转置。它可以被看作前面所讨论的setRowReadCol函数的扩展，这两个核函数的差别在于setRowReadCol使用一个线程块处理输入矩阵的单块转置，而transposeSmem扩展了转置操作，使用了多个线程块和多个数据块：

//存储操作是无冲突的，但是加载操作显示有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];
}

__global__ void transposeSmem(float *out, float *in, int nx, int ny)
{
   __shared__ float tile[BDIMY][BDIMX];
   
   unsigned int ix, iy, ti, to;
   //首先，基于线程索引和块索引计算原始矩阵坐标
   ix = blockIdx.x * blockDim.x + threadIdx.x;
   iy = blockIdx.y * blockDim.y + threadIdx.y;
   
   //然后计算全局内存的索引
   ti = iy * nx + ix;
   
   unsigned int bidx, irow, icol;

   //计算转置矩阵的坐标
   bidx = threadIdx.y * blockDim.x + threadIdx.x;
   //两个新的变量icol和irow被引入以替代threadIdx
   irow = bidx / blockDim.y;
   icol = bidx % blockDim.y;
   
   ix = blockIdx.y * blockDim.x + icol;
   iy = blockIdx.x * blockDim.y + irow;
   //计算用于存储转置矩阵的全局内存索引
   to = iy * ny + ix;
   
   if(ix < nx && iy < ny)
   {
      //利用计算出的偏移量，线程中的线程束可以从全局内存中连续读取，
      //并对二维共享内存数组tile的行进行写入
      tile[threadIdx.y][threadIdx.x] = in[ti];
      __synthreads();
      //全局内存的读操作是合并的，同时共享内存存储体中的写操作没有冲突，
      //所以线程中的线程束可以从共享内存中tile中读取一列，并连续写入到全局内存中
      out[to] = tile[icol][irow]; 
   }
}

使 用 填 充 共 享 内 存 的 矩 阵 转 置
  通过给二维共享内存数组tile中的每一行添加列填充，可以将原矩阵相同列中的数据元素均匀的划分到共享内存存储体中。需要填充的列数取决于设备的计算能力和线程块的大小。对于一个大小为32×16的线程块被测试内核来说，在Tesla K40中必须增加两列填充，在Tesla M2090中必须增加一列填充。在Tesla K40中，下面的语句声明了填充的共享内存：__shared__ float tile[BDIMY][BDIMX + 2];

使 用 展 开 的 矩 阵 转 置
  下面的核函数展开了两个数据块的同时设置：每个线程现在转置了一个数据块跨越的两个数据元素。这种转化的目标是通过创造更多的同时加载和存储以提高设备内存带宽利用率。

__global__ void transposeSmemUnrollPad(float *out, float *in, const int nx, const int ny)
{
   //添加列填充的一维共享内存数组tile被静态声明
   __shared__ float tile[BDIMY*(BDIMX * 2 + IPAD)];
   
   unsigned int ix, iy, ti, to;
   //对于一个给定的线程，计算输入矩阵的坐标和用于存储输入矩阵的全局内存数组的索引
   ix = 2 * blockIdx.x * blockDim.x + threadIdx.x;
   iy = blockIdx.y * blockDim.y + threadIdx.y;
   

   ti = iy * nx + ix;
   
   unsigned int bidx, irow, icol;

   //共享内存转置块中的新线程索引计算如下：
   bidx = threadIdx.y * blockDim.x + threadIdx.x;

   irow = bidx / blockDim.y;
   icol = bidx % blockDim.y;
 
   ix2 = blockIdx.y * blockDim.x + icol;
   iy2 = 2 * blockIdx.x * blockDim.y + irow;

   to = iy2 * ny + ix2;
   
   if(ix + BlockDim.x < nx && iy < ny)
   {
     //由于共享内存数组tile是一维的，所以必须将二维线程索引转换为一维共享内存索引
     unsigned int row_idx = threadIdx.y * (blockDim.x * 2 + IPAD) + threadIdx.x;
     tile[row_idx] = in[ti];
     tile[row_idx + BDIMX] = in[ti + BDIMX];
     __synthreads();
     unsigned int col_idx = icol * (blockDim.x * 2 + IPAD) + irow'
     out[to] = tile[col_idx];
     out[to + ny * BDIMX] = tile[col_idx + BDIMX];
   }
}

增 大 并 行 性
  一个简单有效的优化技术是调整线程块的维度，以找出最佳的执行配置。下图总结了在Tesla K40上各种线程块配置的测试效果。块大小为16×16时展示出了最好的性能，因为它有更多的并发线程块，从而有更好的设备并行性。