CUDA 병렬 프로그래밍(Parallel Programming) 행렬 연산

우선 행렬(Matrix) 연산을 하기 전에 CUDA에서 어떤 방식으로 여러 개의 쓰레드를 돌리는지를 알아야한다. CUDA에선 BLOCK과 GRID로 쓰레드 그룹을 관리한다.

BLOCK에서는 여러 개의 THREAD를 (x, y, z) 즉, 3차원 이하로 가지고 있을 수 있고, GRID는 여러 BLOCK을 2차원 이하(x, y) 로 가지고 있을 수 있다. 그림으로 나타내면 다음과 같다. 그리고 한 BLOCK에는 1024개가 넘는 THREAD를 가지고 있을 수 없다.

그리고 이와 같은 BLOCK과 GRID를 정의한 다음에 Kernel Function 즉 Device에서 돌아갈 함수랑 같이 쓰여야한다.

__global__ void kernelFunc();

dim3 DimGrid(10, 10); //100 thread blocks
dim3 DimBlock(4, 5, 6); // 120 thread per Blocks

kernelFunc<<< DimGrid, DimBlock>>>();

위와 같이 C++에서 템플릿을 넘기듯이 사용하면 된다. dim3 자료형 같은 경우 3개 미만으로 parameter를 넘기는 것도 가능하다.

위와 같이 정의하면 각 thread는 고유한 blockIdx와 threadIdx를 가진다.

위 두 개의 고유한 값을 이용해서 각 thread는 자기의 고유한 위치를 알 수 있다. 그리고 이렇게 정의된 Blocks of Threads는 CUDA 자체 Queue를 통해서 SMs(Streaming multiprocessors)에 할당되어 CUDA가 계산을 해준다.

행렬 덧셈 구현

#include 

__global__ void addKernel(int *c, const int *a, const int *b) {
    int x = threadIdx.x;
    int y = threadIdx.y;
    int i = y * (blockDim.x) + x; // 1차원 인덱스로 변환

    c[i] = a[i] + b[i];
}

int main() {
    //host data
    const int HEIGHT = 5;
    const int WIDTH = 5;

    int a[HEIGHT][WIDTH];
    int b[HEIGHT][WIDTH];
    int c[HEIGHT][WIDTH];

    // 배열 초기화
    for(int y = 0; y>>(dev_c, dev_a, dev_b);

    // device에서 host로 데이터 전송
    cudaMemcpy(c, dev_c, HEIGHT*WIDTH*sizeof(int), cudaMemcpyDeviceToHost);

    // 결과 출력
    for(int y = 0; y

행렬 곱셈 연산

행렬 곱셈 같은 경우는 O(n^3) 이다. 다음과 같이 구현해서 한 쓰레드 당 O(n)의 시간복잡도를 가지는 프로그램을 작성할 수 있다.

__global__ void mulKernel(int *c, const int *a, const int *b, int N) {
    // global index로 변환
    int ROW = blockIdx.y*blockDim.y + threadIdx.y;
    int COL = blockIdx.x*blockDim.x + threadIdx.x;

    double sum = 0.0;

    /*
        이건 thread는 메모레에 Random Access를 하기 때문에
        thread가 불필요한 계산까지도 할 수 있다. 그래서 그것을
        방직하기 위해서 매트릭스에 속하지 않은 것들은 계산하지 않기
        위해서 조건문을 걸어둔 것 이다.
    */
    if( ROW < N && COL < N) {
        for(int k = 0; k

출처

Programming Massively Parallel Processors(ECE498AL) by Davia Kirk and Win-mei W. Hwu

QuantStart: MATRIX-MATRIX MULTIPLICATION ON THE GPU WITH NVIDIA CUDA