Сравнение производительности кастомного ядра cuda, cublas и Cutensor

Я провел следующие тесты CUDA для сравнения показателей производительности умножения (квадратных) матриц, работающих на Ubuntu 24.04 с картой графического процессора Quadro T1000 Mobile с вычислительными возможностями 7.5 (Arch=SM_75) и драйвером графического процессора nvidia-driver. -535 .

Примечание

• cuda_runtime и cublas_v2 устанавливаются с помощью apt-get install nvidia-cuda-dev, который предоставляется репозиторием archive.ubuntu.com/ubuntu благородный/multiverse.
• тарбол Cutensor извлекается и устанавливается соответствующим образом.

Конфигурация

Test              Lib                     Cores
----              ---                     ----
mat_mul_custom    cuda_runtime v12.0.1    cuda-cores
mat_mul_blas      cublas_v2 v12.0.1       cuda-cores
mat_mul_tensor    cutensor v2.0.2.4       dedicated-cuda-cores (non-tensor-cores)
n/a               cutensor                tensor-cores

Mat_mul_test.cu

#include <cuda_runtime.h>
#define STRIDE 1024
#define SUB_STRIDE 32

__global__ void mat_mul_custom_d(float*, float*, float*);
__global__ void mat_mul_custom_d(float* A, float* B, float* C) {
    int by = blockIdx.y, bx = blockIdx.x, ty = threadIdx.y, tx = threadIdx.x;
    int aBegin = by * SUB_STRIDE * STRIDE + ty * STRIDE + tx,
        aEnd = aBegin + STRIDE, bBegin = SUB_STRIDE * bx + ty * STRIDE + tx,
        bStep = SUB_STRIDE * STRIDE;
    float sC = 0;
    for (int a = aBegin, b = bBegin; a < aEnd; a += SUB_STRIDE, b += bStep) {
        __shared__ float As[SUB_STRIDE][SUB_STRIDE], Bs[SUB_STRIDE][SUB_STRIDE];
        As[ty][tx] = A[a];
        Bs[ty][tx] = B[b];
        __syncthreads();
        for (int k = 0; k < SUB_STRIDE; ++k) sC += As[ty][k] * Bs[k][tx];
        __syncthreads();
    }
    C[by * SUB_STRIDE * STRIDE + SUB_STRIDE * bx + ty * STRIDE + tx] = sC;
}

void mat_mul_custom(float*, float*, float*);
void mat_mul_custom(float* A, float* B, float* C) {
    dim3 block(SUB_STRIDE, SUB_STRIDE);
    dim3 grid(STRIDE / SUB_STRIDE, STRIDE / SUB_STRIDE);
    mat_mul_custom_d<<<grid, block>>>(A, B, C);
}

#include <cublas_v2.h>
#define ALPHA 1
#define BETA 0

void mat_mul_blas(float*, float*, float*);
void mat_mul_blas(float* A, float* B, float* C) {
    const float alpha = ALPHA, beta = BETA;
    cublasHandle_t cublasH;
    cublasCreate(&cublasH);
    cublasSgemm(cublasH, CUBLAS_OP_N, CUBLAS_OP_N, STRIDE, STRIDE, STRIDE,
                &alpha, A, STRIDE, B, STRIDE, &beta, C, STRIDE);
}

#include <cutensor.h>
#define N_MODES 2
#define K_ALIGNMENT 128

void mat_mul_tensor(float*, float*, float*);
void mat_mul_tensor(float* A, float* B, float* C) {
    const float alpha = ALPHA, beta = BETA;
    const int modeC[] = {'i', 'j'}, modeA[] = {'i', 'k'}, modeB[] = {'k', 'j'};
    const int64_t extentC[] = {STRIDE, STRIDE}, extentA[] = {STRIDE, STRIDE},
                  extentB[] = {STRIDE, STRIDE};
    cutensorHandle_t handle;
    cutensorCreate(&handle);
    cutensorTensorDescriptor_t descA, descB, descC;
    cutensorCreateTensorDescriptor(handle, &descA, N_MODES, extentA, 0,
                                   CUTENSOR_R_32F, K_ALIGNMENT);
    cutensorCreateTensorDescriptor(handle, &descB, N_MODES, extentB, 0,
                                   CUTENSOR_R_32F, K_ALIGNMENT);
    cutensorCreateTensorDescriptor(handle, &descC, N_MODES, extentC, 0,
                                   CUTENSOR_R_32F, K_ALIGNMENT);
    cutensorOperationDescriptor_t desc;
    cutensorCreateContraction(handle, &desc, descA, modeA, CUTENSOR_OP_IDENTITY,
                              descB, modeB, CUTENSOR_OP_IDENTITY, descC, modeC,
                              CUTENSOR_OP_IDENTITY, descC, modeC,
                              CUTENSOR_COMPUTE_DESC_32F);
    cutensorPlanPreference_t planPref;
    cutensorCreatePlanPreference(handle, &planPref, CUTENSOR_ALGO_DEFAULT,
                                 CUTENSOR_JIT_MODE_NONE);
    cutensorPlan_t plan;
    cutensorCreatePlan(handle, &plan, desc, planPref, 0);
    cutensorContract(handle, plan, (void*)&alpha, A, B, (void*)&beta, C, C, 0,
                     0, 0);
}

#include <assert.h>
#include <stdlib.h>
#define RAND_UPPER 2

void mat_mul(void (*mat_mul_impl)(float*, float*, float*), bool);
void mat_mul(void (*mat_mul_impl)(float*, float*, float*), bool column_major) {
    float *A, *B, *C;
    cudaMallocManaged(&A, sizeof(float) * STRIDE * STRIDE);
    cudaMallocManaged(&B, sizeof(float) * STRIDE * STRIDE);
    cudaMallocManaged(&C, sizeof(float) * STRIDE * STRIDE);
    for (int i = 0; i < STRIDE; i++)
        for (int k = 0; k < STRIDE; k++)
            A[i * STRIDE + k] = rand() % RAND_UPPER;
    for (int j = 0; j < STRIDE; j++)
        for (int k = 0; k < STRIDE; k++)
            B[k * STRIDE + j] = rand() % RAND_UPPER;
    if (column_major)
        mat_mul_impl(B, A, C);
    else
        mat_mul_impl(A, B, C);
    cudaDeviceSynchronize();
    for (int i = 0; i < STRIDE; i++)
        for (int j = 0; j < STRIDE; j++) {
            float res = 0;
            for (int k = 0; k < STRIDE; k++)
                res += A[i * STRIDE + k] * B[k * STRIDE + j];
            assert(res == C[i * STRIDE + j]);
        }
    cudaFree(A);
    cudaFree(B);
    cudaFree(C);
}

#include <string.h>
#include <stdio.h>

int main(int argc, char** argv) {
    if (argc == 2) {
        if (!strcmp(argv[1], "custom")) {
            printf("Test: mat_mul_custom\n");
            mat_mul(mat_mul_custom, false);
            return 0;
        }
        if (!strcmp(argv[1], "blas")) {
            printf("Test: mat_mul_blas\n");
            mat_mul(mat_mul_blas, true);
            return 0;
        }
        if (!strcmp(argv[1], "tensor")) {
            printf("Test: mat_mul_tensor\n");
            mat_mul(mat_mul_tensor, true);
            return 0;
        }
    }
    printf("Usage: nvprof ./Mat_mul_test [custom|blas|tensor]\n");
}

Тесты

$ nvcc ./Mat_mul_test.cu \
       -lcublas \ 
       -L${CUTENSOR_ROOT}/lib/12 \
       -lcutensor \
       -I${CUTENSOR_ROOT}/include \
       -o ./Mat_mul_test

$ nvprof ./Mat_mul_test custom

     Avg    Name
     ---    ----
15.352ms    mat_mul_custom_d

   Count    Total Time    Name
   -----    ----------    ----
      35    4.641ms       Gpu page fault groups

$ nvprof ./Mat_mul_test blas

 4.628ms    volta_sgemm_128x64_nn

      28    3.627ms       Gpu page fault groups

$ nvprof ./Mat_mul_test tensor

 4.879ms    volta_sgemm_128x64_nn

      31    4.441ms       Gpu page fault groups

Тестируемый графический процессор не имеет тензорных ядер в аппаратном обеспечении, но имеет аппаратное обеспечение для обработки тензорных инструкций. Тесты между mat_mul_blas и mat_mul_tensor показывают почти одинаковые показатели производительности. Соответствуют ли они ожиданиям? Что именно выполняет ядро ​​volta_sgemm_128x64_nn? Тест mat_mul_blas намного быстрее, чем mat_mul_custom. Это просто потому, что ядро ​​volta_sgemm_128x64_nn сильно оптимизировано, или главным образом потому, что они выполняются на разных типах ядер CUDA?

Ссылка

• Образец пользовательского ядра cuda Nvidia
• Образец Nvidia Cublas
• Образец кутсенора Nvidia
• Почему кублы