优化 X_transpose*X CUDA 内核的技巧

Posted 2023-03-23

【中文标题】优化 X_transpose*X CUDA 内核的技巧

【英文标题】：Tips for optimizing X_transpose*X CUDA kernel

【发布时间】：2013-01-10 12:18:25

【问题描述】：

我正在编写我的第一个 CUDA 应用程序，并且正在编写我自己的所有内核以供练习。

在一部分中，我只是在计算 X_transpose * X。

我一直在使用 cudaMallocPitch 和 cudaMemcpy2D，我首先在设备上为 X 和 X_transpose*X 分配了足够的空间。我将 X 复制到设备，我的内核接受两个输入，即 X 矩阵，然后是写入 X_transpose * X 结果的空间。

使用分析器，内核最初需要 104 秒才能在大小为 5000x6000 的矩阵上执行。我在主机上用零填充矩阵，使其成为块大小的倍数，以避免检查内核中矩阵的边界。我使用 32 x 32 的块大小。

我做了一些更改以尝试最大化合并读/写到全局内存，这似乎有很大帮助。使用可视化分析器来分析我的代码的发布版本，内核现在需要 4.27 秒来执行。

我还没有对我的 matlab 执行进行准确的计时（只是操作 X'*X;），但它似乎大约是 3 秒。我希望我能获得比使用 CUDA 的 matlab 更好的加速。

nvidia 视觉分析器无法找到我的内核的任何问题，我希望这里的社区可以就如何让它更快地运行提出一些建议。

内核代码：

__global__ void XTXKernel(Matrix X, Matrix XTX) 

//find location in output matrix
int blockRow = blockIdx.y;
int blockCol = blockIdx.x;

int row = threadIdx.y;
int col = threadIdx.x;

Matrix XTXsub = GetSubMatrix(XTX, blockRow, blockCol);
float Cvalue = 0;

for(int m = 0; m < (X.paddedHeight / BLOCK_SIZE); ++m) 

    //Get sub-matrix
    Matrix Xsub = GetSubMatrix(X, m, blockCol);
    Matrix XTsub = GetSubMatrix(X, m, blockRow);

    __shared__ float Xs[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float XTs[BLOCK_SIZE][BLOCK_SIZE];

    //Xs[row][col] = GetElement(Xsub, row, col);
    //XTs[row][col] = GetElement(XTsub, col, row);
    Xs[row][col] = *(float*)((char*)Xsub.data + row*Xsub.pitch) + col;
    XTs[col][row] = *(float*)((char*)XTsub.data + row*XTsub.pitch) + col;

    __syncthreads();

    for(int e = 0; e < BLOCK_SIZE; ++e)
        Cvalue += Xs[e][row] * XTs[col][e];

    __syncthreads();


//write the result to the XTX matrix
//SetElement(XTXsub, row, col, Cvalue);
((float *)((char*)XTXsub.data + row*XTX.pitch) + col)[0] = Cvalue;

我的Matrix结构的定义：

struct Matrix 
matrixLocation location;
unsigned int width;             //width of matrix(# cols)
unsigned int height;            //height of matrix(# rows)
unsigned int paddedWidth;       //zero padded width
unsigned int paddedHeight;      //zero padded height
float* data;                    //pointer to linear array of data elements
size_t pitch;               //pitch in bytes, the paddedHeight*sizeof(float) for host, device determines own pitch
size_t size;                //total number of elements in the matrix
size_t paddedSize;          //total number of elements counting zero padding
;

提前感谢您的建议。

编辑：我忘了提，我在开普勒卡 GTX 670 4GB 上运行。

【问题讨论】：

cs.colostate.edu/~cs675/MatrixTranspose.pdf

【参考方案1】：

16x16 或 8x8 等较小的块大小可能会更快。 This slides 还演示了更大的非方形块/共享内存对于特定矩阵大小可能更快。 对于共享内存分配，使用[BLOCK_SIZE][BLOCK_SIZE+1] 在前导维度上添加一个虚拟元素以避免银行冲突。 尝试使用 #pragma unroll 展开内部 for 循环

另一方面，对于足够大的 A'*A，您可能不会比 matlab GPU 代码快多少。由于matlab的性能瓶颈是调用开销而不是内核性能。

cuBLAS 例程culas_gemm() 可能具有最高的矩阵乘法性能。你可以和它比较。

MAGMAroutine magma_gemm() 在某些情况下比 cuBLAS 具有更高的性能。这是一个开源项目。你也可以从他们的代码中得到一些想法。