矩阵转置算子优化-使用向量化读取

#define FETCH_CFLOAT4(p) (reinterpret_cast<const float4*>(&(p))[0])
#define FETCH_FLOAT4(p) (reinterpret_cast<float4*>(&(p))[0])

template <int BLOCK_SZ>
__global__ void mat_transpose_kernel_v3_5(const float* idata, float* odata, int M, int N) {
    const int bx = blockIdx.x, by = blockIdx.y;
    const int tx = threadIdx.x, ty = threadIdx.y;

    __shared__ float sdata[BLOCK_SZ][BLOCK_SZ];

    int x = bx * BLOCK_SZ + tx * 4;
    int y = by * BLOCK_SZ + ty;

    if (x < N && y < M) {
        FETCH_FLOAT4(sdata[ty][tx * 4]) = FETCH_CFLOAT4(idata[y * N + x]);
    }
    __syncthreads();

    x = by * BLOCK_SZ + tx * 4;
    y = bx * BLOCK_SZ + ty;
    float tmp[4];
    if (x < M && y < N) {
        #pragma unroll
        for (int i = 0; i < 4; ++i) {
            tmp[i] = sdata[tx * 4 + i][ty];
        }
        FETCH_FLOAT4(odata[y * M + x]) = FETCH_FLOAT4(tmp);
    }
}

void mat_transpose_v3_5(const float* idata, float* odata, int M, int N) {
    constexpr int BLOCK_SZ = 32;
    dim3 block(BLOCK_SZ / 4, BLOCK_SZ);
    dim3 grid(Ceil(N, BLOCK_SZ), Ceil(M, BLOCK_SZ));
    mat_transpose_kernel_v3_5<BLOCK_SZ><<<grid, block>>>(idata, odata, M, N);
}

先看宏定义：
#define FETCH_FLOAT4§ (reinterpret_cast<float4*>(&p))[0]
float4 是 CUDA 内置的向量类型，包含 4 个连续的 float（类似 struct { float x, y, z, w; }）。
reinterpret_cast<float4*>(&p)：将指针 &p 强制转换为 float4* 类型（把 p 所在地址当作 float4 的起始地址）。
[0]：取这个 float4* 指向的第一个 float4 对象

#define FETCH_CFLOAT4§ (reinterpret_cast<const float4*>(&p))[0]
它是 FETCH_FLOAT4 的只读版本（指针加了 const 修饰），用于从只读内存（如输入矩阵 idata）中批量读取 float4，避免意外修改源数据

核心作用：向量化内存访问
为什么需要这种优化？
矩阵转置是典型的内存密集型操作（大量读写全局内存）。通过 float4 向量化访问：
减少内存事务次数：原本 4 次 float 读写 → 1 次 float4 读写，降低内存控制器压力。
利用 GPU 硬件对向量类型的优化：CUDA 架构对 float4 等向量类型的内存访问有专门优化（如合并访存、缓存友好性），进一步提升效率。