如何在CUDA中执行多个矩阵乘法？

Question

我有一个方阵矩阵,int *M[10];以便M[i]找到i-th矩阵的第一个元素.我想将所有矩阵乘以M[i]另一个矩阵N,以便我接收一个方阵矩阵int *P[10]作为输出.

我看到了不同的可能性:

1. 确定计算M[i]不同线程的不同元素; 例如,我有大小的10矩阵,4x4以便涉及的线程数量160; 如何使用CUDA实现这种方法？
2. 在上面的例子的框架中,创建一个复合矩阵大小40x40(即,收集10,4x4大小矩阵在一起)和使用40x40线程; 但这种方法似乎需要更多时间; 我正在尝试使用矩阵数组,但我认为我做错了; 我如何在10矩阵中使用这种方法？如何在内核函数中编写代码？

这就是我正在尝试的;

void GPU_Multi(int *M[2], int *N, int *P[2], size_t width)
{

    int *devM[2];
    int *devN[2];
    int *devP[2];
    size_t allocasize =sizeof(int) *width*width;

    for(int i = 0 ; i < 10 ; i ++ ) 
    {
        cudaMalloc((void**)&devM[ i ], allocasize );
        cudaMalloc((void**)&devP[ i ], allocasize ); 
    }

    cudaMalloc((void**)&devN, allocasize );

    for(int i = 0 ; i < 10 ; i ++ ) {

        cudaMemcpy(devM[ i ],M[ i ], allocasize , cudaMemcpyHostToDevice);
        cudaMemcpy(devN, N, allocasize , cudaMemcpyHostToDevice);
        dim3 block(width*2, width*2);
        dim3 grid(1,1,1);
        Kernel_Function<<<grid, block>>>  (devM[2], devN, devP[2],width);

        for(int i = 0 ; i < 10 ; i ++ ) 
        {
            cudaMemcpy(P[ i ], P[ i ], allocatesize, cudaMemcpyDeviceToHost);
            cudaFree(devM[ i ]);   
            cudaFree(devP[ i ]);
        }

    }

Answer 1

从上面的评论和Robert Crovella的答案中可以看出,有不同的可能方法.每种方法可以更好地适用于不同的情况,即,对于不同数量N的矩阵来乘以和不同的矩阵维度MxM.我在下面总结一下:

1. 如果N小而且M很大,也许最好的方法是使用cublas<t>gemm从主机代码调用;
2. 如果N是中等且M适中的,并且如果计算能力至少3.5可用的设备可用,则很可能使用动态并行性,即创建线程的线程网格N并cublas<t>gemm从内核中启动a ; 也许这种方法会因为大N或M由于需要大量线程而失败;
3. 如果N规模很大且M很小,那么与罗伯特克罗维拉相关联的cuBLAS批量方法可能会引起人们的兴趣.
4. 同样,如果N规模很大而且M很小,那么基于cuBLAS流的方法值得一试,罗伯特的评论也提到了这一点;
5. 如果N很大并且M非常小,使用线程网格N线程的方法,每个"手动"计算优化的矩阵乘法可能是吸引人的; 例如,如果必须为矩阵构造矩阵乘法算法4x4,则可以根据用于乘以4×4矩阵的基本乘法的数量来优化由每个线程执行的矩阵乘法.

Answer 2

我认为通过使用专门为此目的设计的CUBLAS批量gemm函数(执行大量"相对较小的"矩阵 - 矩阵乘法运算),可能会实现最快的性能.

即使您希望将矩阵数组(M[])乘以单个矩阵(N),批量gemm函数也需要您传递一个矩阵数组N(即N[]),这在您的情况下都是相同的.

编辑:现在我已经通过一个例子,我似乎很清楚,通过修改下面的例子,我们可以传递一个N矩阵并让GPU_Multi函数简单地将单个N矩阵发送到设备,同时传递一个指针数组因为N,即d_Narray在下面的例子中,所有指针指向设备上的相同N矩阵.

这是一个完全工作的批量GEMM示例:

#include <stdio.h>
#include <cuda_runtime.h>
#include <cublas_v2.h>
#include <assert.h>

#define ROWM 4
#define COLM 3
#define COLN 5

#define cudaCheckErrors(msg) \
    do { \
        cudaError_t __err = cudaGetLastError(); \
        if (__err != cudaSuccess) { \
            fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
                msg, cudaGetErrorString(__err), \
                __FILE__, __LINE__); \
            fprintf(stderr, "*** FAILED - ABORTING\n"); \
            exit(1); \
        } \
    } while (0)


typedef float mytype;
// Pi = Mi x Ni
// pr = P rows = M rows
// pc = P cols = N cols
// mc = M cols = N rows
void GPU_Multi(mytype **M, mytype **N, mytype **P, size_t pr, size_t pc, size_t mc, size_t num_mat, mytype alpha, mytype beta)
{

    mytype *devM[num_mat];
    mytype *devN[num_mat];
    mytype *devP[num_mat];
    size_t p_size =sizeof(mytype) *pr*pc;
    size_t m_size =sizeof(mytype) *pr*mc;
    size_t n_size =sizeof(mytype) *mc*pc;
    const mytype **d_Marray, **d_Narray;
    mytype **d_Parray;
    cublasHandle_t myhandle;
    cublasStatus_t cublas_result;

    for(int i = 0 ; i < num_mat; i ++ )
    {
        cudaMalloc((void**)&devM[ i ], m_size );
        cudaMalloc((void**)&devN[ i ], n_size );
        cudaMalloc((void**)&devP[ i ], p_size );
    }
    cudaMalloc((void**)&d_Marray, num_mat*sizeof(mytype *));
    cudaMalloc((void**)&d_Narray, num_mat*sizeof(mytype *));
    cudaMalloc((void**)&d_Parray, num_mat*sizeof(mytype *));
    cudaCheckErrors("cudaMalloc fail");
    for(int i = 0 ; i < num_mat; i ++ ) {

        cudaMemcpy(devM[i], M[i], m_size , cudaMemcpyHostToDevice);
        cudaMemcpy(devN[i], N[i], n_size , cudaMemcpyHostToDevice);
        cudaMemcpy(devP[i], P[i], p_size , cudaMemcpyHostToDevice);
    }
    cudaMemcpy(d_Marray, devM, num_mat*sizeof(mytype *), cudaMemcpyHostToDevice);
    cudaMemcpy(d_Narray, devN, num_mat*sizeof(mytype *), cudaMemcpyHostToDevice);
    cudaMemcpy(d_Parray, devP, num_mat*sizeof(mytype *), cudaMemcpyHostToDevice);
    cudaCheckErrors("cudaMemcpy H2D fail");
    cublas_result = cublasCreate(&myhandle);
    assert(cublas_result == CUBLAS_STATUS_SUCCESS);
    // change to    cublasDgemmBatched for double
    cublas_result = cublasSgemmBatched(myhandle, CUBLAS_OP_N, CUBLAS_OP_N, pr, pc, mc, &alpha, d_Marray, pr, d_Narray, mc, &beta, d_Parray, pr, num_mat);
    assert(cublas_result == CUBLAS_STATUS_SUCCESS);

    for(int i = 0 ; i < num_mat ; i ++ )
        {
            cudaMemcpy(P[i], devP[i], p_size, cudaMemcpyDeviceToHost);
            cudaFree(devM[i]);
            cudaFree(devN[i]);
            cudaFree(devP[i]);
        }
    cudaFree(d_Marray);
    cudaFree(d_Narray);
    cudaFree(d_Parray);
    cudaCheckErrors("cudaMemcpy D2H fail");

}

int main(){

  mytype h_M1[ROWM][COLM], h_M2[ROWM][COLM];
  mytype h_N1[COLM][COLN], h_N2[COLM][COLN];
  mytype h_P1[ROWM][COLN], h_P2[ROWM][COLN];
  mytype *h_Marray[2], *h_Narray[2], *h_Parray[2];
  for (int i = 0; i < ROWM; i++)
    for (int j = 0; j < COLM; j++){
      h_M1[i][j] = 1.0f; h_M2[i][j] = 2.0f;}
  for (int i = 0; i < COLM; i++)
    for (int j = 0; j < COLN; j++){
      h_N1[i][j] = 1.0f; h_N2[i][j] = 1.0f;}
  for (int i = 0; i < ROWM; i++)
    for (int j = 0; j < COLN; j++){
      h_P1[i][j] = 0.0f; h_P2[i][j] = 0.0f;}

  h_Marray[0] = &(h_M1[0][0]);
  h_Marray[1] = &(h_M2[0][0]);
  h_Narray[0] = &(h_N1[0][0]);
  h_Narray[1] = &(h_N2[0][0]);
  h_Parray[0] = &(h_P1[0][0]);
  h_Parray[1] = &(h_P2[0][0]);

  GPU_Multi(h_Marray, h_Narray, h_Parray, ROWM, COLN, COLM, 2, 1.0f, 0.0f);
  for (int i = 0; i < ROWM; i++)
    for (int j = 0; j < COLN; j++){
      if (h_P1[i][j] != COLM*1.0f) {printf("h_P1 mismatch at %d,%d was: %f should be: %f\n", i, j, h_P1[i][j], COLM*1.0f); return 1;}
      if (h_P2[i][j] != COLM*2.0f) {printf("h_P2 mismatch at %d,%d was: %f should be: %f\n", i, j, h_P2[i][j], COLM*2.0f); return 1;}
    }
  printf("Success!\n");
  return 0;
}