앞선 첫 글에서는 개발 환경을 설정하고, GPU 상에서 “Hello, GPU!” 문자열을 출력하는 예제를 통해 CUDA 코드가 동작하는 것을 확인했습니다. 이제 두 번째 단계로, Host 코드와 Device 코드의 역할을 파악해보겠습니다. 이 개념을 확실하게 알아야 앞으로 진행될 벡터 덧셈이나 메모리 전송, 블록/스레드 개념을 자연스럽게 이해할 수 있습니다.
Host 코드와 Device 코드란?
- Host 코드(호스트 코드): CPU에서 실행되는 코드
- 우리가 평소에 작성하는 일반적인 C++ 코드와 크게 다르지 않습니다.
- 메모리 할당, 입출력, CUDA 커널 호출 등의 작업을 담당합니다.
- Device 코드(디바이스 코드): GPU에서 실행되는 코드
- 병렬 계산을 수행하는 부분입니다.
- “커널 함수(Kernel Function)”로 불리는 함수 형태로 작성하며, __global__ 키워드를 통해 정의합니다.
- GPU 상에서 수천~수백만 개의 스레드를 동시에 실행할 수 있으므로 대용량 데이터 처리에 강력한 성능을 발휘합니다.
이 두 코드 영역은 논리적으로는 한 프로젝트 안에 섞여 있지만, 실제 실행 시에는 다른 프로세서(CPU vs GPU)에서 동작합니다. 따라서 CPU와 GPU 간에 필요한 데이터 이동, 커널 런칭 시 스레드/블록 수 지정 등의 절차가 필요합니다.
코드 예제: Host와 Device 코드 구분하기
지난 예제 main.cu를 다시 살펴보면, 아래와 같이 구분할 수 있습니다.
#include <cstdio>
// Device 코드: GPU 커널 함수
__global__ void hello_kernel() {
printf("Hello from the GPU!\n");
}
// Host 코드: main 함수는 CPU에서 실행
int main() {
// Host 코드에서 Device 코드를 호출(런칭)
hello_kernel<<<1,1>>>();
cudaDeviceSynchronize(); // GPU 연산이 끝날 때까지 대기
return 0;
}
여기서 hello_kernel 함수가 GPU에서 실행되는 커널 함수(__global__ 키워드)이고, main() 함수는 CPU에서 돌아가는 Host 코드입니다. Host 코드는 GPU 커널을 호출한 뒤 cudaDeviceSynchronize()로 GPU 연산이 끝날 때까지 기다립니다.
키워드 정리
- __global__ : GPU에서 실행되는 커널 함수를 정의할 때 사용.
- __device__ : GPU 디바이스에서만 실행되는 함수를 정의할 때 사용 (호스트에서 직접 호출 불가).
- __host__ : 기본적으로 CPU에서 실행되는 함수를 명시할 때 사용 (사실 별도로 안 붙여도 CPU용이지만, 혼용 코드에서 명시적으로 사용할 수 있음).
앞으로 다양한 예제를 통해 __global__, __device__, __host__ 키워드를 적절히 활용하는 방법을 익히게 될 것입니다.
GPU 커널 호출의 구조 이해하기
CUDA에서 GPU 커널을 호출할 때는 특이한 문법을 사용합니다.
kernel_function<<<블록수, 스레드수_per_블록>>>(arguments...);
- <<<...>>> 이라는 특수한 문법은 CUDA 런치 파라미터를 설정하는 부분입니다.
- “블록(Block)”과 “스레드(Thread)” 개수를 지정함으로써 병렬 계산 단위를 정의합니다.
앞으로 벡터 연산 같은 예제를 통해 “어떻게 스레드와 블록을 개념화하고, 실행 단위를 분할하느냐”를 살펴보겠습니다. 아직은 “아, GPU에서 함수를 호출할 때는 이렇게 특이한 구문을 쓰는구나” 정도로 이해하면 충분합니다.
Host와 Device 간 데이터 흐름
Host와 Device 코드는 각자 독립적인 메모리 공간을 가지고 있습니다. CPU 메모리(RAM)와 GPU 메모리(Global Memory)는 분리되어 있기 때문에, GPU에서 연산을 하려면 다음과 같은 절차가 필요합니다.
- Host 메모리에 있는 데이터를 GPU 메모리로 복사(cudaMemcpy 등)
- GPU 커널을 실행하며 GPU 메모리에 있는 데이터를 처리
- 연산 결과를 다시 Host 메모리로 가져오기(cudaMemcpy 등)
이 흐름을 이해하는 것이 매우 중요합니다. 왜냐하면 GPU 성능을 제대로 활용하려면 계산만 빠른 게 아니라, 데이터 이동도 효율적으로 관리해야 하기 때문입니다.
다음 글(#3)에서 메모리 할당(cudaMalloc)과 데이터 복사(cudaMemcpy)를 다루며 이 개념을 구체화할 예정입니다.
예제: Host와 Device 코드가 함께하는 간단한 연산
아직 벡터 덧셈까지 가지 않고, 아주 간단한 형태의 예제를 살짝 맛봅시다. 예를 들어, GPU에서 간단히 숫자 한 개를 증가시키는 연산을 해보겠습니다.
#include <cstdio>
__global__ void increment_kernel(int* d_val) {
// 모든 스레드가 d_val[0]을 1 증가
d_val[0] += 1;
}
int main() {
int h_val = 10; // Host 메모리에 있는 정수
int* d_val = nullptr; // Device 메모리 포인터
// Device 메모리 할당
cudaMalloc(&d_val, sizeof(int));
// Host -> Device 복사
cudaMemcpy(d_val, &h_val, sizeof(int), cudaMemcpyHostToDevice);
// GPU 커널 런칭(블록 1개, 스레드 1개)
increment_kernel<<<1,1>>>(d_val);
// GPU 완료 대기
cudaDeviceSynchronize();
// Device -> Host 복사 (결과 가져오기)
cudaMemcpy(&h_val, d_val, sizeof(int), cudaMemcpyDeviceToHost);
printf("Result: %d\n", h_val); // 원래 10에서 +1 증가 -> 11 예상
cudaFree(d_val);
return 0;
}
위 예제는 아직 “구조”만 파악하기 위한 것입니다. 여기서 중요한 포인트:
- Host 코드(main)에서 Device 메모리 할당(cudaMalloc)
- Host에서 Device로 데이터 전송(cudaMemcpyHostToDevice)
- GPU 커널 호출(<<<...>>>)로 Device에서 연산 수행
- 다시 Device에서 Host로 결과 전송(cudaMemcpyDeviceToHost)
이 과정을 통해 Host와 Device 코드가 함께 협력하는 큰 그림을 그릴 수 있습니다. 이후 글에서 이런 패턴이 더 복잡하고 효율적인 형태로 발전할 것입니다.
Before vs After: 개념 이해의 중요성
이전(아무 개념 없이)은 “GPU 코드를 어떻게 돌리나?” 정도의 관점이었다면,
Before:
- Host/Device 개념을 정확히 이해하지 못하고, 그냥 GPU 커널 호출이 마법처럼 돌아가는 줄 앎.
- 데이터 복사, 메모리 관리 절차의 필요성을 직관적으로 느끼지 못함.
After:
- Host와 Device가 별개 메모리 공간, 별개 실행 단위임을 명확히 이해.
- 커널 실행, 데이터 전송 흐름을 머릿속에 그림으로 그릴 수 있음.
- 이후 다룰 메모리 전송, 스레드/블록 개념, 최적화 기법을 배울 토대 형성.
이러한 이해를 바탕으로 앞으로 더 구조적인 코드 작성과 성능 최적화를 고민할 수 있게 됩니다.
다음 글 예고
다음 글(#3)에서는 실제로 Host <-> Device 간 데이터 전송하는 법, 메모리 할당(cudaMalloc), 복사(cudaMemcpy), 그리고 커널 런칭 인자를 통해 스레드를 확장하는 기본 개념을 배울 예정입니다. 이를 통해 간단한 벡터 덧셈 예제를 시도해보며, 진짜 “GPU 병렬 계산” 맛보기로 넘어갑니다.
유용한 링크 & 리소스
- NVIDIA CUDA 프로그래밍 가이드 - Host/Device 코드 개념 및 CUDA 키워드 설명
- CUDA Runtime API 문서 - cudaMalloc, cudaMemcpy 등 함수 활용법
- NVIDIA CUDA 샘플 코드 레포지토리 - 다양한 예제 코드 확인 가능
- Modern C++ Features 정리 - 최신 C++ 기능을 숙지해두면 CUDA 코드에도 적용할 수 있음