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 (또는 이미 수행 한 다른 라이브러리)을 사용하는 것이 좋습니다.