使用CUDA進行矩陣乘法，執行時間很長

Question

我是CUDA的新手，我一直試圖弄清楚我在這裏做錯了什麼。 CUDA所花費的時間不僅僅是使用CPU來乘以一個矩陣。如果我做錯了，請告訴我。 這裏是我的代碼：

#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 
#include <stdio.h> 
#include <stdlib.h> 
#include <cstdlib> 
#include <assert.h> 
#include <time.h> 
#define size 100 // Matrix size 
#define cols size // Matrix width 
#define rows size // Matrix height 

void checkCUDAError(const char *msg) 
{ 
    cudaError_t err = cudaGetLastError(); 
    if(cudaSuccess != err) 
    { 
     fprintf(stderr, "Cuda error: %s: %s.\n", msg, cudaGetErrorString(err)); 
     exit(EXIT_FAILURE); 
    }       
} 
__global__ void matrixMul(int *A, int *B, int *C) 
{ 
    int bx = blockIdx.x; // Block index 
    int tx = threadIdx.x; // Thread index 
    int ts = blockDim.x; // number of threads 
    // Declaration of the shared memory C element 
    extern __shared__ int c_element_sum[]; 
    c_element_sum[tx] = A[tx+((bx/ts)*ts)] * B[(bx%ts)+(tx*ts)]; 

    //Block until all threads in the block have written their data to shared mem 
    __syncthreads(); 

    int sum; 
    for(int i=0; i<ts; i++){ 
     if(i==0){ 
      sum=c_element_sum[i]; 
     } 
     else{ 
      sum+=c_element_sum[i]; 
     } 
    } 
    C[bx] = sum; 

} 


///////////////////////////////////////////////////////// 
// Program main 
///////////////////////////////////////////////////////// 

int main(int argc, char** argv) 
{ 
    //create timer. 
    clock_t t1, t2; 

    //start timer 
    t1=clock(); 

    //allocate host memory for matrices 
    unsigned int size_A = cols * rows; 
    unsigned int mem_size_A = sizeof(int) * size_A; 
    int* mA = (int*) malloc(mem_size_A); 

    unsigned int size_B = cols * rows; 
    unsigned int mem_size_B = sizeof(int) * size_B; 
    int* mB = (int*) malloc(mem_size_B); 

    unsigned int size_C = cols * rows; 
    unsigned int mem_size_C = sizeof(int) * size_C; 
    int* mC = (int*) malloc(mem_size_C); 

    //initialize host memory 
    for (int i = 0; i < size_A; ++i){ 
     mA[i] = 1; 
     mB[i] = 1; 
     mC[i] = 0; 
    } 

    // allocate device memory 
    int* d_mA; 
    int* d_mB; 
    int* d_mC; 
    cudaMalloc((void**) &d_mA, mem_size_A); 
    cudaMalloc((void**) &d_mB, mem_size_B); 
    cudaMalloc((void**) &d_mC, mem_size_C); 

    //copy host memory to device (A and B) 
    cudaMemcpy(d_mA, mA, mem_size_A, cudaMemcpyHostToDevice); 
    cudaMemcpy(d_mB, mB, mem_size_B, cudaMemcpyHostToDevice); 
    cudaMemcpy(d_mC, mC, mem_size_C, cudaMemcpyHostToDevice); 

    // setup execution parameters 
    int numThreadsPerBlock = cols; 
    int numBlocks = (cols * rows); 
    int sharedMemSize = numThreadsPerBlock * sizeof(int); 

    dim3 dimGrid(numBlocks); 
    dim3 dimBlock(numThreadsPerBlock); 

    // execute the kernel 
    matrixMul <<< dimGrid, dimBlock, sharedMemSize >>>(d_mA, d_mB, d_mC); 

    //Block until device has completed 
    cudaThreadSynchronize(); 

    // check if kernel execution generated an error 
    // Check for any CUDA errors 
    checkCUDAError("kernel invocation"); 

    //copy result from device to host 
    cudaMemcpy(mC, d_mC, mem_size_C, cudaMemcpyDeviceToHost); 

    // Check for any CUDA errors 
    checkCUDAError("memcpy"); 

    //stop timer 
    t2 = clock(); 

    //check results 
    for (int i = 0; i < size_C; ++i){ 
     assert(mC[i] == cols); 
    } 

    //clean up memory 
    free(mA); 
    free(mB); 
    free(mC); 
    cudaFree(d_mA); 
    cudaFree(d_mB); 
    cudaFree(d_mC); 

    printf("WITH CUDA - clocks: %d \n\n", t2-t1); 

    ////////////////////////////// 
    ///////// CPU ONLY ////////// 
    ///////////////////////////// 

    //create timer. 
    clock_t cpu_t1, cpu_t2; 

    //start timer 
    cpu_t1=clock(); 

    //allocate host memory for matrices 
    unsigned int cpu_size_A = cols * rows; 
    unsigned int cpu_mem_size_A = sizeof(int) * cpu_size_A; 
    int* cpu_mA = (int*) malloc(cpu_mem_size_A); 

    unsigned int cpu_size_B = cols * rows; 
    unsigned int cpu_mem_size_B = sizeof(int) * cpu_size_B; 
    int* cpu_mB = (int*) malloc(cpu_mem_size_B); 

    unsigned int cpu_size_C = cols * rows; 
    unsigned int cpu_mem_size_C = sizeof(int) * cpu_size_C; 
    int* cpu_mC = (int*) malloc(cpu_mem_size_C); 

    //initialize host memory 
    for (int i = 0; i < cpu_size_A; ++i){ 
     cpu_mA[i] = 1; 
     cpu_mB[i] = 1; 
     cpu_mC[i] = 0; 
    } 

    int ts = cols; 
    for(int bx=0; bx<(cols*rows);bx++){ 
     int sum = 0; 
     for(int tx=0; tx<cols; tx++){ 
      sum += cpu_mA[tx+((bx/ts)*ts)] * cpu_mB[(bx%ts)+(tx*ts)]; 
     } 
     cpu_mC[bx]=sum; 
    } 

    //stop timer 
    cpu_t2 = clock(); 

    //check results 
    for (int i = 0; i < cpu_size_C; ++i){ 
     assert(cpu_mC[i] == cols); 
    } 

    //clean up memory 
    free(cpu_mA); 
    free(cpu_mB); 
    free(cpu_mC); 

    printf("CPU ONLY - clocks: %d \n\n", cpu_t2-cpu_t1); 

    return 0; 
} 

Answer 1

根據您的程序，這是預期的。您的計時器看起來像是爲程序的整個執行計時，其中包括複製到設備，計算時間以及將結果複製回來。考慮到您爲程序提供的相當小的工作量（100x100矩陣），內存複製的開銷遠遠超過您在使用內核執行計算時獲得的任何計算益處。你的內核本身也不是最有效的實現。

我不認爲你在做任何事情錯誤，這只是你沒有提供GPU的足夠大的工作量，你可能會進一步優化你的內核。請注意，簡單地擴大塊的大小可能不會顯着提高CPU的性能，因爲您還將擴大內存管理時間。雖然在CUDA上編寫程序的第一個實現相對簡單，但獲得良好性能顯然更困難。使用CUDA的最有效方式是計算與內存交易的比例很高。例如，讓一系列計算密集型內核的管道在一塊數據上連續運行，只需要在開始和結束時進行主機設備複製。

如果這只是幫助您學習編寫CUDA代碼的程序，那麼這是一個很棒的步驟，深入瞭解如何優化矩陣乘法內核將在很多其他情況下爲您提供良好的服務。如果您正在編寫此內核以用於生產軟件，我建議您使用高度優化的線性代數庫CUBLAS：http://developer.nvidia.com/cublas（或其他已經爲您完成的努力的其他庫）。