💻 4장. CUDA 프로그래밍 기본

🎯 학습 목표

이 장에서는 CUDA 프로그래밍의 핵심 실습 내용을 단계별로 익히며,
실제 코드 작성, 실행, 메모리 처리, 동기화, 성능 측정 및 기본 최적화 기법을 배웁니다.

4.1 CUDA 코드 작성 및 실행 방법

🧾 기본 코드 예제

#include <stdio.h>

__global__ void helloFromGPU() {
    printf("Hello from GPU thread %d\n", threadIdx.x);
}

int main() {
    helloFromGPU<<<1, 4>>>();            // 1블록, 4스레드
    cudaDeviceSynchronize();             // GPU 연산 완료 대기
    return 0;
}

⚙️ 빌드 및 실행

nvcc hello.cu -o hello
./hello

✅ 실행 결과

Hello from GPU thread 0  
Hello from GPU thread 1  
Hello from GPU thread 2  
Hello from GPU thread 3

4.2 커널 함수 정의 및 실행

__global__ void square(float* out, float* in) {
    int idx = threadIdx.x;
    out[idx] = in[idx] * in[idx];
}

실행 예시

square<<<1, 256>>>(d_out, d_in);  // 1개의 블록, 256개의 스레드

4.3 스레드 인덱싱 및 블록 인덱싱

기본 인덱싱

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

2차원 인덱싱 예시

int x = threadIdx.x + blockIdx.x * blockDim.x;
int y = threadIdx.y + blockIdx.y * blockDim.y;

변수	설명
threadIdx.x	블록 내 스레드 번호
blockIdx.x	그리드 내 블록 번호
blockDim.x	블록 내 스레드 수

4.4 메모리 할당 및 데이터 전송

필수 함수

함수	설명
cudaMalloc()	GPU 메모리 할당
cudaMemcpy()	CPU ↔ GPU 간 데이터 복사
cudaFree()	GPU 메모리 해제

예제 코드

float *d_in, *d_out;
int size = N * sizeof(float);
cudaMalloc(&d_in, size);
cudaMemcpy(d_in, h_in, size, cudaMemcpyHostToDevice);
cudaMemcpy(h_out, d_out, size, cudaMemcpyDeviceToHost);
cudaFree(d_in); cudaFree(d_out);

📐 메모리 정렬 최적화 팁 (Memory Alignment)

성능 저하 원인

스레드들이 비연속적인 메모리 위치에 접근하면 coalescing이 되지 않아 성능 저하 발생

✅ 최적화 전략

전략	설명
연속 주소 접근	스레드들이 연속된 메모리 위치를 접근해야 함
벡터 타입 활용	float4, int4 등 자연스럽게 정렬된 타입 사용
패딩 추가	구조체 내 멤버 간 정렬 보장
공유 메모리 bank conflict 방지	공유 메모리 접근 시 주소 정렬 필요

예시

struct __align__(16) Vec4 {
    float x, y, z, w;
};

💡 메모리 접근 패턴이 정렬되어 있으면 전역 메모리 접근 속도 최대 10배 이상 차이 날 수 있습니다.

4.5 동기화 및 실행 시간 측정

동기화

cudaDeviceSynchronize(); // GPU 연산 완료 대기

실행 시간 측정 (CUDA Event API)

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);

myKernel<<<grid, block>>>();

cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms = 0;
cudaEventElapsedTime(&ms, start, stop);
printf("Execution Time: %.3f ms\n", ms);

🛡️ 에러 처리 매크로

#define CUDA_CHECK(call) { \
    cudaError_t err = call; \
    if (err != cudaSuccess) { \
        fprintf(stderr, "CUDA Error: %s\n", cudaGetErrorString(err)); \
        exit(EXIT_FAILURE); \
    } \
}

💡 모든 cudaMalloc, cudaMemcpy, 커널 실행 후에는 반드시 에러 체크 필요

✅ 종합 예제: 벡터 덧셈

__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 = 512;
    size_t size = N * sizeof(float);

    float *h_a = (float*)malloc(size);
    float *h_b = (float*)malloc(size);
    float *h_c = (float*)malloc(size);
    for (int i = 0; i < N; i++) {
        h_a[i] = i;
        h_b[i] = i * 2;
    }

    float *d_a, *d_b, *d_c;
    CUDA_CHECK(cudaMalloc(&d_a, size));
    CUDA_CHECK(cudaMalloc(&d_b, size));
    CUDA_CHECK(cudaMalloc(&d_c, size));

    CUDA_CHECK(cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice));
    CUDA_CHECK(cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice));

    int threads = 256;
    int blocks = (N + threads - 1) / threads;

    vectorAdd<<<blocks, threads>>>(d_a, d_b, d_c, N);
    CUDA_CHECK(cudaGetLastError());

    CUDA_CHECK(cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost));
    printf("h_c[0] = %.1f\n", h_c[0]);  // 결과: 0.0

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

🧯 디버깅 팁

도구	설명
cudaGetLastError()	커널 실행 오류 확인
cuda-memcheck	메모리 접근 오류 탐지
printf()	커널 내부 디버깅용
NVIDIA Nsight	성능 분석 및 디버깅 지원 GUI 도구

🌐 실습 환경 안내

환경	특징
로컬 GPU 환경	가장 안정적이고 권장
Google Colab (Pro)	제한된 GPU 환경에서 실습 가능
Docker + NVIDIA 컨테이너	독립된 CUDA 개발 환경 구성 가능
Visual Studio + Nsight	Windows 기반 개발/디버깅 최적화 IDE

📌 최종 요약표

항목	내용
작성 및 실행	.cu 파일, nvcc 컴파일, <<<>>> 커널 실행
메모리 처리	cudaMalloc, cudaMemcpy, cudaFree
스레드 인덱싱	threadIdx, blockIdx, blockDim 조합
성능 측정	cudaEventElapsedTime() 사용
에러 처리	CUDA_CHECK, cudaGetLastError()
정렬 최적화	연속 접근, 벡터 타입, 구조체 정렬 고려
디버깅	cuda-memcheck, Nsight, printf 등 활용