比较自定义 cuda 内核、cublas 和 cutensor 之间的性能

问题描述 投票：0回答：1

我进行了以下 CUDA 测试，以比较（方）矩阵乘法的性能数据，在 Ubuntu 24.04 上运行，GPU 卡为 Quadro T1000 Mobile，计算兼容性为 7.5 (arch=SM_75) 和GPU 驱动程序 nvidia-driver-535 .

配置

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");
}

测试

$ 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

被测GPU在硬件上没有
张量核心，但具有处理张量指令的硬件。 mat_mul_blas 和 mat_mul_tensor 之间的测试表现出几乎相同的性能数据。他们如预期的那样吗？内核 volta_sgemm_128x64_nn 到底执行什么？测试mat_mul_blas比mat_mul_custom快得多。仅仅是因为内核volta_sgemm_128x64_nn高度优化，还是主要是因为它们在不同类型的CUDA核心上执行？

参考



  • Nvidia 定制 cuda 内核示例

  • Nvidia cublas 示例

  • Nvidia cutensor 示例

    • 
    

	


cuda

tensor

cublas


1个回答




    
    
    

        

            

                

                    
0
投票

                    
                

            

        

        


他们如预期的那样吗？




是的，预计 cublas 执行矩阵乘法运算的速度比您展示的内核代码更快。  同样明显的是，cutensor 设计者使用的方法与 cublas 大致相同 - 事实上，很明显他们以与您的 cublas 调用类似的方式调用 cublas。 （时长基本相同，内核名称也相同）




内核volta_sgemm_128x64_nn到底执行什么？




它是针对 FP32 数据的高度优化的 gemm 内核（矩阵-矩阵乘法）（这就是您看到 
sgemm
 的原因），由 cublas 设计者编写，作为闭源 cublas 库的一部分。  因此我无法提供详细的描述。  然而，编写高度优化的 gemm 内核并非易事。  
这篇文章可能会让您了解与内核代码相比所涉及的复杂性，我们现在可能将其称为“天真的”共享内存矩阵乘法代码。



仅仅是因为内核volta_sgemm_128x64_nn经过高度优化，还是主要是因为它们在不同类型的CUDA核心上执行？




因为它是高度优化的。  它们没有使用不同的内核（SM 中的功能单元）。  您可以使用 
CUDA 二进制实用程序 获得一些证据/证明。  如果你转储每种情况下的 SASS 代码，你会发现你的内核和 cublas sgemm 内核都在使用 FFMA SASS 指令来执行矩阵乘法运算。

正如已经提到的，目前没有 CUDA GPU 为 FP32 数据提供张量核计算路径，所以我们现在可以完全放弃这个概念，这不是这里发生的情况，事实上你的 cublas 都调用并且您的 cutensor 调用正在调用 cublas 非张量核心 sgemm 例程，这进一步证明了这一点。

    

	

    

    


  

  

    
最新问题


  





  

    © www.soinside.com 2019 - 2025. All rights reserved.