cuda矩阵转置

CUDA编程入门—矩阵转置（六）

朴素矩阵转置

• 内存访问模式

  • 读操作：原矩阵按行访问可实现合并访问，效率较高；转置矩阵按列访问为交叉访问，是 GPU 性能最差的访问模式，但在矩阵转置中无法避免。
  • 写操作：对转置矩阵写时按列访问，是交叉访问 。

• 两种转置核函数性能

  启用一级缓存时

  • NaiveCol 优势：NaiveCol 是按列加载、按行存储，存在交叉读取的问题。但由于一级缓存的存在，即使按列读取时加载了未被此次访问用到的数据，这些数据会暂留在缓存中。当后续操作需要这些数据时，就可以直接从缓存读取，发生缓存命中，减少对全局内存的访问。而全局内存访问相对较慢，所以通过缓存命中能有效提升性能。
  • NaiveRow 劣势：NaiveRow 按行加载、按列存储，写入时是交叉访问。在 GPU 中，交叉访问是性能较差的内存访问模式，且写入操作无法利用一级缓存（写操作不在一级缓存中缓存），所以这种交叉写入的低效访问模式会显著影响其性能，导致它不如 NaiveCol 。

  禁用一级缓存时

  • 失去缓存优势：禁用一级缓存后，NaiveCol 无法再依靠缓存来减少全局内存访问。其按列读取的不合并特性会浪费带宽，使得性能受到较大影响。
  • 写入劣势凸显：NaiveRow 本身写入时的交叉访问问题依旧存在，且没有了缓存对读取操作的优化辅助（虽然其读取是合并访问，但缓存能进一步提升整体性能），性能同样受到影响。所以在禁用一级缓存的情况下，二者性能都因各自内存访问模式的缺陷受到显著影响，表现出与启用缓存时不同的性能关系 。

// 基于行的朴素转置核函数
__global__ void transposeNaiveRow(float *out, float *in, const int nx, const int ny) {
    unsigned int ix = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int iy = blockDim.y * blockIdx.y + threadIdx.y;
    if (ix < nx && iy < ny) {
        out[ix * ny + iy] = in[iy * nx + ix];
    }
}

// 基于列的朴素转置核函数
__global__ void transposeNaiveCol(float *out, float *in, const int nx, const int ny) {
    unsigned int ix = blockDim.x * blockIdx.x + threadIdx.x;
    unsigned int iy = blockDim.y * blockIdx.y + threadIdx.y;
    if (ix < nx && iy < ny) {
        out[iy * nx + ix] = in[ix * ny + iy];
    }
}

优化

展开转置

• 并行计算架构：将任务细分为子任务，让多线程并行处理。契合 GPU 的 SIMT 架构，以线程束为单位执行指令，降低指令开销，提升并行度。

• 内存层次与访问：借助共享内存缓存数据，减少对慢速全局内存的访问。通过合理组织线程访问顺序，实现内存合并，提升数据读写带宽，降低延迟。

• 指令执行与数据处理：细粒度的指令序列便于 GPU 指令调度单元灵活调度，减少等待与空闲时间。更多 CUDA 核心同时工作，提高硬件资源利用率，加速转置。

  __global__ void transposeUnrolled(int *in, int *out, int width, int height) {
      int x = blockIdx.x * blockDim.x + threadIdx.x;
      int y = blockIdx.y * blockDim.y + threadIdx.y;
      if (x < width && y < height) {
          // 假设展开4次示例
          out[x * height + y] = in[y * width + x];
          out[(x + 1) * height + y] = in[y * width + (x + 1)];
          out[(x + 2) * height + y] = in[y * width + (x + 2)];
          out[(x + 3) * height + y] = in[y * width + (x + 3)];
      }
  }
  

瘦块

• 原理：在 GPU 并行计算中，将线程块设计得更 “瘦长”，即减少线程块在一个维度上的大小，增加在另一个维度上的大小。这样能让内存访问更连续，减少内存事务数量，提高内存访问效率。尤其在矩阵转置场景下，可改善对内存的访问模式，降低数据读取和写入的开销。

调用时合理设置blockDim，比如dim3 blockDim(16, 64);（具体数值依实际情况调整），让线程块在 x 方向上较窄，y 方向上较长。

对角坐标

• 原理：改变线程块到数据块的映射方式，采用对角坐标映射，使内存访问更均匀地分布到 DRAM 的不同分区，减少分区冲突。在传统笛卡尔坐标映射中，可能出现内存请求集中在某些分区的情况，而对角坐标映射可降低这种冲突，提升内存访问效率，进而优化矩阵转置性能。

__global__ void transposeDiagonalCoordinates(int *in, int *out, int width, int height) {
    // 根据对角坐标计算索引
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int x = idx % width;
    int y = idx / width;
    if (x < width && y < height) {
        out[x * height + y] = in[y * width + x];
    }
}