junyeokk
Blog
Parallel Computing·2025. 10. 20

CUDA 커널 함수 작성 방법

커널 launch 문법

CUDA 커널은 일반 함수 호출과 다른 특별한 문법을 사용한다.

c
// 일반 C/C++ 함수 호출
func_name(int param, ...);

for (int i = 0; i < SIZE; ++i) {
    func_name(param, ...);
}
c
// CUDA 커널 launch
__global__ void kernel_name(int param, ...);

kernel_name<<<1, SIZE>>>(param, ...);

<<<, >>>는 C/C++에서 사용하지 않는 연산자다. CUDA 전용 문법으로, 커널 실행 설정(execution configuration)을 지정한다.

phone call vs rocket launch
일반 함수 호출은 전화를 거는 것과 같다. 하나의 상대에게 연결된다. 커널 launch는 로켓을 발사하는 것과 같다. 수천 개의 스레드가 동시에 폭발적으로 실행된다.

<<<grid, block>>> 문법

c
kernel_name<<<gridSize, blockSize>>>(params);

첫 번째 인자 gridSize는 Grid 크기다. 몇 개의 블록을 생성할지 지정한다.

두 번째 인자 blockSize는 Block 크기다. 각 블록에 몇 개의 스레드를 생성할지 지정한다.

c
add_kernel<<<1, 8>>>(dev_c, dev_a, dev_b);

위 코드는 1개의 블록에 8개의 스레드를 실행한다. 총 8개의 스레드가 add_kernel()을 동시에 실행한다.

threadIdx: 각 스레드의 고유 번호

각 스레드는 자신의 인덱스를 알 수 있다. threadIdx라는 내장 변수로 접근한다.

c
__global__ void add_kernel(int* c, const int* a, const int* b) {
    int i = threadIdx.x;  // 각 스레드가 자신의 인덱스를 안다
    c[i] = a[i] + b[i];
}

threadIdx.x는 현재 스레드의 인덱스다. 0부터 blockSize-1까지의 값을 가진다.

c
add_kernel<<<1, SIZE>>>(dev_c, dev_a, dev_b);

이 호출은 SIZE개의 스레드를 실행한다. 각 스레드는 다른 threadIdx.x 값을 가진다.

  • 스레드 0: threadIdx.x = 0
  • 스레드 1: threadIdx.x = 1
  • 스레드 2: threadIdx.x = 2
  • ...
  • 스레드 SIZE-1: threadIdx.x = SIZE-1

각 스레드는 서로 다른 데이터를 처리한다. CPU의 for 루프가 i를 순차적으로 증가시키는 대신, GPU는 각 스레드에 고유한 i를 할당한다.

완전한 CUDA 예제

gpu-add.cu의 전체 구조다.

c
#include "./common.cpp"

// GPU 커널 함수
__global__ void add_kernel(int* c, const int* a, const int* b) {
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

int main(void) {
    // Step 1: host-side 데이터 준비
    const int SIZE = 5;
    const int a[SIZE] = { 1, 2, 3, 4, 5 };
    const int b[SIZE] = { 10, 20, 30, 40, 50 };
    int c[SIZE] = { 0 };

    // Step 2: device 메모리 할당
    int* dev_a = 0;
    int* dev_b = 0;
    int* dev_c = 0;
    cudaMalloc((void**)&dev_a, SIZE * sizeof(int));
    cudaMalloc((void**)&dev_b, SIZE * sizeof(int));
    cudaMalloc((void**)&dev_c, SIZE * sizeof(int));

    // Step 3: host → device 복사
    cudaMemcpy(dev_a, a, SIZE * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(dev_b, b, SIZE * sizeof(int), cudaMemcpyHostToDevice);

    // Step 4: 커널 실행
    add_kernel<<<1,SIZE>>>(dev_c, dev_a, dev_b);
    cudaDeviceSynchronize();

    // Step 5: device → host 복사
    cudaMemcpy(c, dev_c, SIZE * sizeof(int), cudaMemcpyDeviceToHost);

    // Step 6: 결과 출력
    printf("{1,2,3,4,5} + {10,20,30,40,50} = {11,22,33,44,55}\\n");

    // Step 7: device 메모리 해제
    cudaFree(dev_a);
    cudaFree(dev_b);
    cudaFree(dev_c);

    return 0;
}

실행 결과.

plain
CUDA: success
{1,2,3,4,5} + {10,20,30,40,50} = {11,22,33,44,55}

CPU vs GPU 실행 방식 비교

CPU 버전은 순차 실행한다.

c
for (int i = 0; i < SIZE; ++i) {
    add_kernel(i, a, b, c);
}
  • time 0: CPU가 add_kernel(0, ...)을 실행
  • time 1: CPU가 add_kernel(1, ...)을 실행
  • time 2: CPU가 add_kernel(2, ...)를 실행

GPU 버전은 병렬 실행한다.

c
add_kernel<<<1, SIZE>>>(dev_c, dev_a, dev_b);
  • time 0:GPU[core #0]부터 GPU[core #(SIZE-1)]까지 모두 동시에 실행

for 루프가 CPU의 순차 처리를, <<<, >>> 문법이 GPU의 병렬 처리를 대체한다.

CPU for 루프를 CUDA 커널로 변환

변환 과정은 다음과 같다.

1단계: 루프 body를 함수로 추출

c
for (int i = 0; i < SIZE; ++i) {
    c[i] = a[i] + b[i];
}

↓

c
void add_kernel(int idx, const int* a, const int* b, int* c) {
    int i = idx;
    c[i] = a[i] + b[i];
}

for (int i = 0; i < SIZE; ++i) {
    add_kernel(i, a, b, c);
}

2단계: __global__ 선언자 추가 및 threadIdx 사용

c
__global__ void add_kernel(int* c, const int* a, const int* b) {
    int i = threadIdx.x;  // idx 파라미터 대신
    c[i] = a[i] + b[i];
}

3단계: for 루프를 커널 launch로 교체

c
add_kernel<<<1, SIZE>>>(dev_c, dev_a, dev_b);

이제 SIZE개의 스레드가 동시에 실행한다.

threadIdx

threadIdx는 CUDA가 제공하는 내장 변수다.

c
int i = threadIdx.x;

위 코드는 CPU의 for 루프 인덱스와 동일한 역할을 한다. 차이는 CPU가 순차적으로 i를 증가시키는 반면, GPU는 각 스레드가 처음부터 서로 다른 i 값을 가진다는 점이다.

CPU의 순차 처리.

  • 시간 0: i = 0으로 c[0] 계산
  • 시간 1: i = 1로 c[1] 계산
  • 시간 2: i = 2로 c[2] 계산

GPU의 병렬 처리.

  • 시간 0: 모든 스레드가 동시에 자신의 i로 c[i] 계산
    • 스레드 0은 i = 0
    • 스레드 1은 i = 1
    • 스레드 2는 i = 2

주의사항

커널 함수는 void만 반환한다.

c
__global__ int add_kernel() { }  // 컴파일 에러!
__global__ void add_kernel() { }  // 정상

수천 개의 스레드가 동시에 실행하므로, 어느 스레드의 반환값을 받아야 할지 불명확하다. 결과는 메모리에 저장하고 cudaMemcpy()로 가져온다.

관련 문서