LeetCUDA-学习记录-矩阵转置(mat-transpose)-CSDN博客

本文链接：https://blog.csdn.net/CCDRAG/article/details/148098986

LeetCUDA-学习记录

前言

本系列文章主要是记录作者本人学习LeetCUDA项目的笔记以及作者本人对该项目中对于一系列算子设计的思路进行理解分析与优化分析，有所错误望大佬们指正~

一、矩阵转置(mat-transpose)

项目设计思路

按照项目中给定的最优方案，主要分为两种写法，第一种是以行主序为主的，通过float4向量加大计算延迟从而隐藏指令延迟以及共享内存降低访存耗时，第二种是以列主序为主的，优化思路同行主序一样采用float4向量和共享内存使用

行主序： 按照一行一行的数据从左至右，从上到下的顺序进行读取
列主序： 按照一列一列的数据从上到下，从左至右的顺序进行读取

这里我只针对项目中行主序进行分析。
在mat-transpose kernel设计当中，项目的mat_transpose_f32x4_shared_bcf_col2row_kernel(float4向量化，共享内存，去bank conflict)版本 中通过对于二维shared memory数组在一维上增加一个PAD=1,加大一维大小，从而实现线程读取bank错位来解决**mat_transpose_f32x4_shared_col2row_kernel(float4向量化版本，共享内存)**中的bank问题。个人理解如下：
在这里插入图片描述
但是只是将4路降为了2路。并没有完全去除bank conflict，通过nsight compute分析可以看到

测试获取问题

No add PAD–f32x4_shared_col2row_kernel(float4向量化版本，共享内存)版本

在这里插入图片描述

add PAD–f32x4_shared_bcf_col2row_kernel(float4向量化，共享内存，去bank conflict)版本

在这里插入图片描述
因此我想尽可能消除所有的bank conflict。于是我设计了这样的一个核函数，同样有float4向量和共享内存优化，同时完全去除了bank conflict。

个人解决方案

my_f32x4_shared_col2row_kernel

__global__ void mat_transpose_my_f32x4_shared_col2row2d_kernel(float *x, float *y, const int row, const int col){
    const int global_x = blockIdx.x * blockDim.x + threadIdx.x;
    const int global_y = blockIdx.y * blockDim.y + threadIdx.y;
    const int local_x = threadIdx.x;
    const int local_y = threadIdx.y;
    __shared__ float tile[WARP_SIZE_S][WARP_SIZE_S * 4];
    if( global_x * 4 + 3 < col && global_y < row){
        float4 x_val = FLOAT4(x[global_x * 4 + global_y * col]);
        FLOAT4(tile[local_y][local_x * 4]) = FLOAT4(x_val);
        __syncthreads();
        float4 smem_val;
        const int tx = local_x + (local_y % 4) * WARP_SIZE_S; // 0 ~ 64
        const int ty = local_y / 4; // 0 ~ 4
        smem_val.x = tile[ty * 4 + 0][tx];
        smem_val.y = tile[ty * 4 + 1][tx];
        smem_val.z = tile[ty * 4 + 2][tx];
        smem_val.w = tile[ty * 4 + 3][tx];

        const int out_x = 4 * blockIdx.x * blockDim.x + tx;
        const int out_y = blockIdx.y * blockDim.y + ty * 4;
        FLOAT4(y[out_x * row + out_y]) = FLOAT4(smem_val);
    }
}