첫 CUDA 실행: vectorAdd

vectorAdd는 CUDA의 첫 실습으로 좋다. 덧셈 자체가 중요해서가 아니라, CUDA 프로그램의 전체 흐름이 한 번에 드러나기 때문이다.

CPU 배열 준비
→ GPU 배열 할당
→ CPU에서 GPU로 복사
→ kernel launch
→ GPU에서 CPU로 복사
→ 결과 확인

Colab 확인

GPU 런타임을 켠 뒤 먼저 확인한다.

Colab GPU 확인 명령 보기

!nvidia-smi
!nvcc --version

첫 코드

전체 vecadd.cu 코드 보기

%%writefile vecadd.cu
#include <cstdio>
#include <cuda_runtime.h>

__global__ void vectorAdd(float* a, float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < n) {
        c[i] = a[i] + b[i];

        if (i < 8) {
            printf("blockIdx.x=%d threadIdx.x=%d global_i=%d\n",
                   blockIdx.x, threadIdx.x, i);
        }
    }
}

int main() {
    int n = 16;
    size_t bytes = n * sizeof(float);

    float* h_a = (float*)malloc(bytes);
    float* h_b = (float*)malloc(bytes);
    float* h_c = (float*)malloc(bytes);

    for (int i = 0; i < n; i++) {
        h_a[i] = i;
        h_b[i] = i * 10;
    }

    float* d_a;
    float* d_b;
    float* d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    int threadsPerBlock = 4;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;

    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);
    cudaDeviceSynchronize();

    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    for (int i = 0; i < n; i++) {
        printf("c[%d] = %.1f\n", i, h_c[i]);
    }

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

컴파일하고 실행한다.

컴파일/실행 명령 보기

!nvcc vecadd.cu -o vecadd
!./vecadd

이 카드에서 볼 것

아직 모든 문법을 깊게 이해하지 않아도 된다. 먼저 이 세 줄만 표시해 둔다.

int i = blockIdx.x * blockDim.x + threadIdx.x;
int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

인터랙티브 설계

h_a, h_b, h_c는 CPU 쪽 박스로 표시한다.
d_a, d_b, d_c는 GPU 쪽 박스로 표시한다.
cudaMemcpy가 실행될 때 화살표가 CPU에서 GPU 또는 GPU에서 CPU로 움직인다.
kernel이 실행될 때 c[0]부터 c[15]까지 채워지는 모습을 보여준다.

확인

이 코드는 배열 전체 합 하나를 구하는가, 아니면 elementwise add를 하는가?
threadsPerBlock = 4, n = 16이면 block은 몇 개 필요한가?
printf 출력 순서가 항상 정렬되어 있다고 기대하면 안 되는 이유는 무엇인가?