CUDA C++가 진짜로 뭔데

C++에 확장 셋 박은 거 — 그리고 철학 하나

처음 보면 CUDA C++는 그냥 평범한 C/C++ 같아. header 박고, 함수 쓰고, 컴파일하고. 근데 확장 세 개가 게임을 통째로 바꿔:

Function qualifier — __global__, __device__, __host__ — 함수가 어디서 돌고 누가 부를 수 있는지 결정해. __global__ 커널은 GPU에서 돌지만 launch는 CPU host 코드가 해.
Triple-angle launch — kernel<<<grid, block>>>(args); — 컴파일 타임 메타데이터인데 런타임이 GPU 스케줄러용 command buffer로 바꿔.
SIMT built-in — threadIdx.x, blockIdx.x, blockDim.x, gridDim.x — thread마다 자기 좌표 줘서 어느 작업 슬라이스 처리할지 알려줘.

나머지는 C++17 toolchain 그대로 상속받아. lambda, template, <algorithm>, host쪽엔 std::vector까지. 트위스트는 device 코드가 NVCC를 통과한다는 거. NVCC가 translation unit 갈라서 GPU 부분은 PTX/SASS back-end로 컴파일하고 하나의 binary로 다시 link해.

왜 컴파일러가 따로 있냐면 GPU는 또 다른 CPU가 아니거든. 자기 ISA 있고 (PTX는 portable intermediate, SASS는 architecture별 binary), register file 모양 다르고, memory hierarchy 다르고. NVCC는 한 source 파일로 두 컴파일된 프로그램 만들어서 link 시점에 합쳐주는 이중언어 번역가야.

철학은 SIMT (Single Instruction, Multiple Threads). 수천 개 thread가 같은 커널 본문을 lockstep으로 돌면서 자기 index에 따라 분기해. 원소 loop 돌리는 코드를 쓰는 게 아니라, thread 하나가 뭘 할지 쓰고 데이터 다 덮을 만큼 thread 띄워.

Code

Function qualifier 셋 — 같은 문법, 다른 집·cuda

// __host__ — CPU에서 돌고, CPU에서 호출. qualifier 없으면 default.
__host__ float cpu_norm(const float* a, int n) {
    float s = 0.f;
    for (int i = 0; i < n; ++i) s += a[i] * a[i];
    return std::sqrt(s);
}

// __device__ — GPU에서 돌고, GPU 코드에서만 호출 가능.
__device__ float square(float x) { return x * x; }

// __global__ — GPU에서 돌고, CPU에서만 호출 가능. 이게 커널.
__global__ void norm_kernel(const float* a, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = square(a[i]);  // __device__ helper 호출
}

Triple-angle launch — NVCC가 보는 거·cuda

int n = 1 << 20;       // 1,048,576 원소
int threads = 256;     // block당 thread (32의 배수여야 함)
int blocks  = (n + threads - 1) / threads;  // ceil 나눗셈

norm_kernel<<<blocks, threads>>>(d_a, d_out, n);
cudaDeviceSynchronize();   // GPU 끝날 때까지 대기

Exercise

NVIDIA PTX ISA reference 페이지 브라우저로 열고 목차만 훑어. 외우려고 하지 말고 — PTX가 자기만의 load/store, branch, arithmetic instruction이 있고 x86이나 ARM이랑 다르다는 것만 인지해. 이게 네 CUDA C++가 변환되는 어셈블리야. 북마크 해놔; 커널이 느릴 때 generated PTX 읽으러 다시 와야 해.

C++에 확장 셋 박은 거 — 그리고 철학 하나

Code

External links

Exercise

Progress

댓글 0