CUDA中更快的矩阵乘法

Question

目前,我在cuda c中制作了一个神经网络程序.因为我需要操纵矩阵乘法,所以我没有将CUBLAS用于MM.我为MM使用以下代码.我想知道是否有人有一些建议让它更快,这可能非常有用,因为我需要在学习期间使用MM数百万次.谢谢.这是MakeFile:

# cuda root
_CUDA_ROOT_ = /usr/local/cuda

NVCC = nvcc
# include and lib paths
INCLUDES=-I${_CUDA_ROOT_}/include
LIB_PATH=-L${_CUDA_ROOT_}/lib64

# libraries to link against
LIB= -lcudart -lcublas
CU_SRC= main.cu
EXE=$(CU_SRC:.cu=)
#------------------------------
# Choose your gpu arch
SM = sm_35
all: $(EXE)
$(EXE): $(CU_SRC)
        $(NVCC) -arch $(SM) $(CU_SRC) -o $(EXE) $(LIB_PATH) $(LIB)

clean:
        rm -f *.o *.cu_o $(EXE)

这是MM代码:

__global__
void matrixMulti(float* A_d, float* B_d, float* C_d, int m, int k, int n)
{
    __shared__ float ds_A[TILE_WIDTH][TILE_WIDTH];
    __shared__ float ds_B[TILE_WIDTH][TILE_WIDTH];
    int col = blockIdx.x*blockDim.x + threadIdx.x;
    int row = blockIdx.y*blockDim.y + threadIdx.y;
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    float sum = 0;

    for(int t=0; t<(n-1)/TILE_WIDTH+1; t++)
    {
        if(row<m && t*TILE_WIDTH+tx<n)
            ds_A[ty][tx] = A_d[row*n + t*TILE_WIDTH+tx];
        else
            ds_A[ty][tx] = 0.0;
        if(t*TILE_WIDTH+ty<n && col<k)
            ds_B[ty][tx] = B_d[(t*TILE_WIDTH+ty)*k + col];
        else
            ds_B[ty][tx] = 0.0;
        __syncthreads();
        for(int i=0; i<TILE_WIDTH; i++)
            sum += ds_A[ty][i] * ds_B[i][tx];
        __syncthreads();
    }
    if(row<m && col<k)
        C_d[col+row*k] = sum;
}

这是代码主要部分的示例:

const int TILE_WIDTH = 32;

int main()
{
    int m, k, n;
    m = 10000, k = 10000, n = 10000;
    float *A, *B, *C;
    A = new float[m*n];
    B = new float[n*k];
    C = new float[m*k];
    float *A_d, *B_d, *C_d;
    for (int i=0; i<m*n; i++)
    {
        A[i] = 2;
    }
    for (int i=0; i<n*k; i++)
    {
        B[i] = 3;
    }
    cudaMalloc(&A_d, sizeof(float)*m*n);
    cudaMalloc(&B_d, sizeof(float)*n*k);
    cudaMalloc(&C_d, sizeof(float)*m*k);
    cudaMemcpy(A_d, A, sizeof(float)*m*n, cudaMemcpyHostToDevice);
    cudaMemcpy(B_d, B, sizeof(float)*k*n, cudaMemcpyHostToDevice);
    dim3 dimGrid((k-1)/TILE_WIDTH+1, (m-1)/TILE_WIDTH+1, 1);
    dim3 dimBlock(TILE_WIDTH, TILE_WIDTH, 1);
    matrixMulti<<<dimGrid,dimBlock>>>(A_d, B_d, C_d, m, k, n);
    cudaMemcpy(C, C_d, sizeof(float)*m*k, cudaMemcpyDeviceToHost);
    return 0;
}

Answer 1

首先,要确定这是你想要做的.在没有描述你想要做的操作的情况下,很难对此进行评论,但要注意矩阵乘法是一个n-cubed操作.如果你的操作不是那么复杂,那么你只需要使用cuBLAS就可以做得更好.

为什么是这样？cuBLAS可能会比您编写的任何内容都快,而且随着新GPU架构的推出,它将更加可维护.像GEMM这样的最佳实现将根据体系结构而有所不同,因此您现在为硬件编写的任何代码都必须针对新硬件进行重新优化.

现在,问题.您应该考虑使用许多技术来优化此代码:

1. 计算每个线程的多个输出值.这样可以减少共享内存的压力,因为可以在多个计算中使用切片数据.
2. 修复共享内存中的银行冲突.这应该由文档很好地涵盖.
3. Vectorise共享内存加载和存储.我注意到你正在编译sm_35.该架构的共享存储体每个都具有64位/时钟的带宽.加载单个浮点数只有32位,因此在没有向量化的情况下浮点数不会获得全带宽.你应该看看float2/float4类型.
4. 考虑双缓冲.将数据加载到一个共享内存块中,同时在另一个上操作 这样可以更有效地隐藏全局内存操作的高延迟,降低同步开销,并且通常可以更好地执行.它使用两倍的共享内存,因为你需要一次两个tile.

关于GPU上矩阵乘法的实现有很多论文,我建议你检查一下.你会从这些论文中获得更多的细节,而不是你会在SO上提出广泛的问题.

最后......你确定你不想使用cuBLAS吗？我不会指望获得75%的cuBLAS性能,即使这将是一个挑战.