0

我正在编写我的第一个 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;
}

我的矩阵结构的定义：

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 上运行。

4

1 回答 1

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

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

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

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

于 2013-01-27T02:35:31.627 回答