첫 CUDA 프로그램

학습 목표

• •CUDA 프로그램의 기본 구조를 익힌다
• •cudaMalloc, cudaMemcpy, cudaFree 함수를 이해한다
• •벡터 덧셈 커널을 직접 작성한다
• •커널 실행 구성을 설정한다

CUDA 프로그램 기본 구조

모든 CUDA 프로그램은 5단계 패턴을 따릅니다.

기본 패턴

1. Host 메모리 할당 및 초기화
2. Device 메모리 할당
3. Host → Device 데이터 복사
4. 커널 실행
5. Device → Host 결과 복사
6. 메모리 해제

시각적 흐름

cudaMalloc - GPU 메모리 할당

함수 시그니처

cuda
복사
cudaError_t cudaMalloc(void** devPtr, size_t size);

사용법

cuda
복사
float* d_array;  // device 포인터 (관례적으로 d_ 접두사 사용)
int n = 1000;
size_t size = n * sizeof(float);

// GPU 메모리 할당
cudaMalloc(&d_array, size);
// d_array는 이제 GPU 메모리를 가리킴

주의사항

✓ 반드시 포인터의 주소(&)를 전달
✓ 크기는 바이트 단위
✗ CPU에서 d_array 값을 직접 읽으면 안됨
  → 반드시 cudaMemcpy로 복사 후 사용

cudaMemcpy - 메모리 복사

함수 시그니처

cuda
복사
cudaError_t cudaMemcpy(void* dst, const void* src,
                       size_t count, cudaMemcpyKind kind);

복사 방향

사용 예시

cuda
복사
float h_input[1000];   // Host 배열
float* d_input;        // Device 포인터

// GPU 메모리 할당
cudaMalloc(&d_input, 1000 * sizeof(float));

// Host → Device 복사
cudaMemcpy(d_input, h_input, 1000 * sizeof(float),
           cudaMemcpyHostToDevice);

// 커널 실행 후...

// Device → Host 복사
cudaMemcpy(h_input, d_input, 1000 * sizeof(float),
           cudaMemcpyDeviceToHost);

cudaFree - GPU 메모리 해제

함수 시그니처

cuda
복사
cudaError_t cudaFree(void* devPtr);

사용법

cuda
복사
float* d_array;
cudaMalloc(&d_array, size);

// ... 사용 ...

// 메모리 해제
cudaFree(d_array);
d_array = nullptr;  // 권장: 해제 후 nullptr 설정

메모리 누수 방지

cuda
복사
// 나쁜 예: 해제하지 않으면 메모리 누수
void badFunction() {
    float* d_temp;
    cudaMalloc(&d_temp, size);
    // cudaFree 없이 함수 종료 → 메모리 누수!
}

// 좋은 예: 항상 쌍으로 사용
void goodFunction() {
    float* d_temp;
    cudaMalloc(&d_temp, size);
    // ... 작업 ...
    cudaFree(d_temp);  // 반드시 해제
}

벡터 덧셈 예제

문제 정의

두 벡터 A, B를 더해 C를 계산
C[i] = A[i] + B[i], for i = 0, 1, ..., n-1

CPU 구현 (순차)

c
복사
void vectorAddCPU(float* a, float* b, float* c, int n) {
    for (int i = 0; i < n; i++) {
        c[i] = a[i] + b[i];  // 순차 처리
    }
}
// 10,000개 원소 → 10,000번 반복

GPU 구현 (병렬)

cuda
복사
__global__ void vectorAddGPU(float* a, float* b, float* c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];  // 각 스레드가 하나의 원소 처리
    }
}
// 10,000개 원소 → 10,000개 스레드가 동시에 처리!

전역 인덱스 계산

idx = threadIdx.x + blockIdx.x * blockDim.x

예시: blockDim.x = 256, blockIdx.x = 3, threadIdx.x = 100
idx = 100 + 3 × 256 = 100 + 768 = 868
→ 이 스레드는 배열의 868번 원소를 처리

경계 검사의 중요성

cuda
복사
// 왜 if (idx < n) 검사가 필요한가?

n = 10000, blockSize = 256
gridSize = (10000 + 255) / 256 = 40 블록
총 스레드 수 = 40 × 256 = 10,240개

→ 240개의 스레드는 유효한 데이터가 없음
→ 경계 검사로 잘못된 메모리 접근 방지

완전한 벡터 덧셈 프로그램

cuda
복사
#include <cuda_runtime.h>
#include <stdio.h>
#include <stdlib.h>

// 커널 함수
__global__ void vectorAdd(float* a, float* b, float* c, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

int main() {
    int n = 10000;
    size_t size = n * sizeof(float);

    // 1. Host 메모리 할당
    float* h_a = (float*)malloc(size);
    float* h_b = (float*)malloc(size);
    float* h_c = (float*)malloc(size);

    // 2. 데이터 초기화
    for (int i = 0; i < n; i++) {
        h_a[i] = (float)i;
        h_b[i] = (float)(i * 2);
    }

    // 3. Device 메모리 할당
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, size);
    cudaMalloc(&d_b, size);
    cudaMalloc(&d_c, size);

    // 4. Host → Device 복사
    cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);

    // 5. 커널 실행 구성
    int blockSize = 256;
    int gridSize = (n + blockSize - 1) / blockSize;
    printf("Grid: %d blocks, Block: %d threads\n", gridSize, blockSize);

    // 6. 커널 실행
    vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

    // 7. Device → Host 복사
    cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);

    // 8. 결과 검증
    bool correct = true;
    for (int i = 0; i < n; i++) {
        if (h_c[i] != h_a[i] + h_b[i]) {
            correct = false;
            break;
        }
    }
    printf("Result: %s\n", correct ? "PASS" : "FAIL");

    // 9. 메모리 해제
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);

    return 0;
}

컴파일 및 실행

bash
복사
nvcc vector_add.cu -o vector_add
./vector_add

# 출력:
# Grid: 40 blocks, Block: 256 threads
# Result: PASS

커널 실행 구성

블록 크기 선택 가이드

올림 나눗셈 공식

cuda
복사
// gridSize 계산 (올림 나눗셈)
int gridSize = (n + blockSize - 1) / blockSize;

// 동일한 표현
int gridSize = ceil((float)n / blockSize);
int gridSize = (n - 1) / blockSize + 1;

에러 처리

에러 체크 매크로

cuda
복사
#define CHECK_CUDA(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA Error: %s at line %d\n", \
                   cudaGetErrorString(err), __LINE__); \
            exit(1); \
        } \
    } while(0)

// 사용 예시
CHECK_CUDA(cudaMalloc(&d_a, size));
CHECK_CUDA(cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice));

커널 에러 체크

cuda
복사
// 커널 실행 후 에러 체크
vectorAdd<<<gridSize, blockSize>>>(d_a, d_b, d_c, n);

// 방법 1: getLastError
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("Kernel error: %s\n", cudaGetErrorString(err));
}

// 방법 2: 동기화 후 체크
cudaDeviceSynchronize();
err = cudaGetLastError();

정리

CUDA 프로그램 체크리스트

□ cudaMalloc 성공 여부 확인
□ cudaMemcpy 방향 올바르게 지정
□ 커널에서 경계 검사 (idx < n)
□ 커널 실행 후 에러 체크
□ 모든 cudaMalloc에 대응하는 cudaFree
□ Host 메모리도 free

핵심 포인트: cudaMalloc → cudaMemcpy → kernel<<<>>> → cudaMemcpy → cudaFree 패턴을 숙지하세요!