이제 우리는 Host(CPU)와 Device(GPU) 코드 구조를 이해하고, 간단한 커널 호출 예제를 통해 GPU 코드가 어떻게 동작하는지 살펴봤습니다. 이번 글에서는 실제로 데이터를 GPU로 보내고, 연산한 뒤 다시 결과를 가져오는 과정을 좀 더 구체적으로 다뤄볼 예정입니다. 즉, Host <-> Device 메모리 전송(cudaMemcpy) 및 메모리 할당(cudaMalloc) 방법을 배우고, 이를 활용한 간단한 벡터 덧셈 예제를 통해 GPU 병렬 연산의 기본을 체험해보겠습니다.
Host와 Device 메모리 관리의 기본
앞서 말했듯 Host와 Device는 별개의 메모리 공간을 사용합니다. 이는 마치 서로 다른 섬에 사는 두 존재가 다리를 통해 상자를 옮기는 상황과 비슷합니다. 여기서 상자는 데이터, 다리는 cudaMemcpy 같은 API라고 할 수 있습니다.
- Host 메모리 할당: 그냥 new, malloc, std::vector 등을 통해 CPU 메모리에 변수를 확보하면 됩니다. 이건 우리가 익숙하게 해왔던 부분이죠.
- Device 메모리 할당: GPU 메모리에 공간을 확보하기 위해 cudaMalloc 함수를 사용합니다.
float* d_arr; cudaMalloc(&d_arr, size_in_bytes); - 데이터 복사(cudaMemcpy): Host 메모리에 있는 데이터를 Device 메모리로 복사하거나, 반대로 Device 결과를 Host로 가져올 때 사용합니다.
cudaMemcpy(d_arr, h_arr, size_in_bytes, cudaMemcpyHostToDevice); // H->D cudaMemcpy(h_arr, d_arr, size_in_bytes, cudaMemcpyDeviceToHost); // D->H
이렇게 Host와 Device 메모리를 분리해서 생각하면, 데이터 이동 과정이 명확해집니다.
커널 호출과 연계하기
지난 글에서 보았던 커널 런칭 문법을 다시 떠올려봅시다.
kernel_function<<<블록수, 스레드수_per_블록>>>(...);
호출하기 전에 이미 GPU 메모리에 연산 대상 데이터가 있어야 GPU가 작업을 효율적으로 수행할 수 있습니다. 반대로, 연산이 끝난 뒤에는 결과를 Host로 가져와야 유용하게 활용할 수 있습니다.
이 흐름을 단계별로 정리하면 다음과 같습니다.
- Host에서 연산 대상 데이터 준비 (예: std::vector<float> 등)
- Device 메모리에 공간 할당(cudaMalloc)
- Host → Device 데이터 복사(cudaMemcpy)
- GPU 커널 런칭(<<<...>>>)
- GPU 연산 완료 대기(cudaDeviceSynchronize())
- Device → Host 결과 복사(cudaMemcpy)
- Device 메모리 해제(cudaFree)
이제 이 과정을 벡터 덧셈 예제로 구체화해봅시다.
벡터 덧셈 예제: CPU vs GPU
우리가 C++에서 벡터 덧셈을 할 때는 단순히 반복문을 돌며 요소-wise로 더해주면 됩니다. 예를 들어:
for (int i = 0; i < N; ++i) {
c[i] = a[i] + b[i];
}
이 연산을 GPU에서 한다면, 모든 요소를 독립적으로 처리할 수 있으므로 수천 개의 스레드가 동시에 덧셈을 수행할 수 있습니다. 각각의 스레드가 벡터의 일부를 담당하게 되는 것이죠.
GPU 커널 코드 예제
__global__ void vector_add_kernel(const float* a, const float* b, float* c, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
c[idx] = a[idx] + b[idx];
}
}
- blockIdx.x, threadIdx.x는 CUDA에서 제공하는 내장 변수로, 각 스레드가 자신만의 인덱스를 계산하는 데 사용됩니다.
- if (idx < N)는 전체 스레드 수가 N보다 크더라도 인덱스 범위를 벗어나지 않도록 안전장치를 두는 것입니다.
Host 코드에서의 흐름
#include <cstdio>
#include <vector>
#include <cstdlib>
// 벡터 덧셈 커널
__global__ void vector_add_kernel(const float* a, const float* b, float* c, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
c[idx] = a[idx] + b[idx];
}
}
int main() {
int N = 100000;
size_t size = N * sizeof(float);
// Host 메모리에 벡터 준비
std::vector<float> h_a(N), h_b(N), h_c(N);
for (int i = 0; i < N; ++i) {
h_a[i] = static_cast<float>(rand()) / RAND_MAX;
h_b[i] = static_cast<float>(rand()) / RAND_MAX;
}
// Device 메모리 할당
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
// Host -> Device 복사
cudaMemcpy(d_a, h_a.data(), size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b.data(), size, cudaMemcpyHostToDevice);
// 커널 런칭: 스레드와 블록 설정
int threadsPerBlock = 256;
int blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
vector_add_kernel<<<blocks, threadsPerBlock>>>(d_a, d_b, d_c, N);
// GPU 연산 완료 대기
cudaDeviceSynchronize();
// Device -> Host 복사 (결과 가져오기)
cudaMemcpy(h_c.data(), d_c, size, cudaMemcpyDeviceToHost);
// 결과 확인(간단히 첫 몇 개 출력)
for (int i = 0; i < 5; ++i) {
printf("c[%d] = %f\n", i, h_c[i]);
}
// 메모리 해제
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
return 0;
}
여기서 핵심 포인트는 다음과 같습니다.
- threadsPerBlock와 blocks로 스레드 개수를 조정해 전체 벡터를 커버.
- cudaMalloc, cudaMemcpy, cudaFree로 Host-Device 간 데이터 흐름 명확하게 관리.
- GPU 커널 내부에서 인덱스를 계산해 벡터 요소에 접근.
이 과정을 거치면 대량의 데이터를 짧은 시간에 병렬로 처리할 수 있게 됩니다.
Before vs After: 병렬 연산 준비 완료
이전까지는 단순히 “GPU에서 한 번 출력해보는 정도”였다면, 이제는 데이터 이동과 병렬 연산을 통해 Host-Device 협력 구조가 보입니다.
Before:
- GPU 코드 실행 여부만 확인 (“Hello, GPU!” 출력)
- 데이터 이동 개념이 불명확함
After:
- Host에서 데이터 준비 → Device로 전송 → GPU 병렬 연산 → 결과 회수 과정 확립
- 실제 연산 예제(벡터 덧셈)로 GPU 병렬성 체감
- 이후 스레드, 블록, 그리드 최적화, 메모리 계층 구조 이해 등으로 확장 가능한 발판 마련
다음 글 예고
다음 글(#4)에서는 스레드와 블록, 그리고 그리드 개념을 좀 더 깊게 파헤칩니다. 벡터 덧셈 예제를 통해 우리가 설정한 블록과 스레드 개수의 의미를 명확히 알고, 이를 다양한 문제에 적용하는 방법을 알아보겠습니다.
유용한 링크 & 리소스
- NVIDIA CUDA Samples: 다양한 예제 코드
- CUDA C Programming Guide: 메모리 관리 및 커널 런칭 설명
- CppReference: 최신 C++ 기능 참고
- NVIDIA Developer Forums: 질문 & 답변으로 문제 해결 도움