이번 글에서는 지금까지 배운 스레드/블록/그리드 개념을 2차원 데이터에 적용하고, GPU의 비동기(Asynchronous) 기능을 살짝 맛보며 최적화를 향한 첫걸음을 내딛어보겠습니다. 이미지 처리나 행렬 연산 같은 2D 데이터 처리는 실제 GPU 활용 분야에서 매우 흔하고도 중요한 영역입니다. 또한 비동기 스트림을 활용하면 Host와 Device 사이의 작업을 겹치게 만들어 전체 처리 시간을 단축할 수 있다는 점을 알아볼 것입니다.
2D 데이터 처리를 위한 블록/그리드 설정
지금까지는 1차원 데이터(예: 벡터)에 대한 처리를 주로 다뤘습니다. 하지만 이미지(2D 배열)나 행렬을 다룬다면 2차원적으로 스레드와 블록을 배치하는 것이 직관적일 수 있습니다.
2D 스레드 인덱스 계산 예제
예를 들어, M x N 크기의 이미지가 있고, 각 픽셀에 대해 병렬로 연산을 하고 싶다고 합시다. 이때 블록과 스레드는 2차원으로 정의할 수 있습니다.
dim3 threadsPerBlock(16, 16);
dim3 blocks((N + threadsPerBlock.x - 1) / threadsPerBlock.x,
(M + threadsPerBlock.y - 1) / threadsPerBlock.y);
이렇게 하면 (threadsPerBlock.x, threadsPerBlock.y) 형태로 블록 내 스레드를 구성하고, (blocks.x, blocks.y) 형태로 그리드를 정의할 수 있습니다.
커널에서 인덱스를 구할 때도 2D를 활용합니다.
__global__ void process_image_kernel(const unsigned char* input, unsigned char* output, int M, int N) {
int x = blockIdx.x * blockDim.x + threadIdx.x; // 열 인덱스
int y = blockIdx.y * blockDim.y + threadIdx.y; // 행 인덱스
if (x < N && y < M) {
int idx = y * N + x;
// 여기서 픽셀을 처리. 예: grayscale 변환 등
output[idx] = 255 - input[idx]; // 단순 반전 예제
}
}
이제 각 스레드는 (x, y) 위치의 픽셀 하나를 담당하게 됩니다. 이를 통해 2D 데이터 처리 시 코드가 한층 명확해집니다.
비동기 스트림(Asynchronous Stream)이란?
지금까지는 cudaDeviceSynchronize()를 통해 Host 코드가 GPU 연산이 끝날 때까지 대기하는 방식으로 진행했습니다. 하지만 GPU 연산이 꽤 오래 걸리는 경우, Host가 그동안 아무 일도 하지 않고 기다리는 것은 비효율적입니다. **비동기 스트림(Asynchronous Stream)**을 활용하면 Host와 Device 사이의 작업을 겹쳐서(overlap) 수행할 수 있습니다.
기본 개념
- 기본적으로 모든 CUDA 호출(커널 런칭, memcpy)은 default stream(스트림 0)에서 순차적으로 실행됩니다.
- 별도의 스트림(cudaStream_t)을 생성하면, Host는 곧바로 다음 일로 넘어갈 수 있고, Device는 해당 스트림에서 커널을 진행하는 동안 Host는 다른 작업을 할 수 있습니다.
- 다양한 스트림을 사용해 Host-Device 전송과 커널 실행을 겹치는(Overlap) 전략을 취하면 전체 프로그램 실행 시간이 단축될 수 있습니다.
간단한 예제 코드 (개념적 예시)
cudaStream_t stream;
cudaStreamCreate(&stream);
int N = 100000;
size_t size = N * sizeof(float);
float *h_a, *h_b, *h_c;
float *d_a, *d_b, *d_c;
// 호스트 메모리 준비 (생략)
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
// Host->Device 복사 비동기 호출
cudaMemcpyAsync(d_a, h_a, size, cudaMemcpyHostToDevice, stream);
cudaMemcpyAsync(d_b, h_b, size, cudaMemcpyHostToDevice, stream);
// 커널 런칭 (비동기, stream에서 실행)
int threadsPerBlock = 256;
int blocks = (N + threadsPerBlock - 1) / threadsPerBlock;
vector_add_kernel<<<blocks, threadsPerBlock, 0, stream>>>(d_a, d_b, d_c, N);
// Device->Host 복사 비동기 호출
cudaMemcpyAsync(h_c, d_c, size, cudaMemcpyDeviceToHost, stream);
// 다른 Host-side 계산을 여기서 수행 가능 (GPU와 병렬)
// ... 다른 작업 ...
// 모든 작업 완료 대기
cudaStreamSynchronize(stream);
// 정리
cudaStreamDestroy(stream);
cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
여기서 cudaMemcpyAsync와 커널 런칭은 지정한 스트림(stream)에서 비동기적으로 수행됩니다. Host는 cudaStreamSynchronize(stream)를 호출하기 전까지 해당 스트림의 작업이 끝날 때까지 기다리지 않고, 자유롭게 다른 일을 할 수 있습니다.
성능 향상을 위한 첫 발걸음
비동기 스트림을 사용하면 다음과 같은 장점이 있습니다.
- 오버랩(Overlap): Host에서 다른 계산을 하는 동안 GPU에서 커널이 돌아가거나, GPU가 메모리 전송을 하는 동안 Host가 다른 작업을 수행할 수 있습니다.
- 파이프라이닝(Pipelining): 데이터 전송과 연산을 번갈아가며 수행해, GPU나 Host 어느 한 쪽도 공회전하지 않도록 할 수 있습니다.
물론 여기에는 더 깊은 최적화 기법(스트림 다중 활용, 워크로드 분할, CUDA 이벤트(Event) 활용 등)이 있지만, 지금은 단순히 “비동기 스트림을 통해 Host와 Device가 동시에 일할 수 있다”는 개념을 이해하는 것만으로도 충분합니다.
Before vs After: 동기적 대기 vs 비동기 오버랩
이전까지는 단순히 cudaDeviceSynchronize()로 모든 일이 끝날 때까지 기다렸습니다.
Before:
- Host는 GPU 연산이 끝날 때까지 대기
- 작업 순차적 진행 → CPU, GPU 중 하나는 반드시 대기 상태
After:
- 비동기 스트림 활용 시 Host와 GPU가 동시에 다른 작업 수행 가능
- 데이터 전송과 연산을 겹치면 처리 시간 단축 가능
- 향후 고성능 애플리케이션에서 필수적으로 고려할 최적화 방향 제시
다음 글 예고
다음 글(#6)에서는 더 복잡한 메모리 계층 구조(Shared Memory, Constant Memory 등)와 간단한 최적화 전략에 대해 살짝 맛볼 예정입니다. 아직 입문 단계이므로 너무 깊게 들어가지 않고, 어떤 방향으로 최적화할 수 있는지만 살펴보며, 앞으로 있을 중급/고급 주제의 예고편 같은 성격을 지닙니다.
유용한 링크 & 리소스
- NVIDIA CUDA Streams 관련 문서
- CUDA C++ Streams와 이벤트 활용 예시
- NVIDIA CUDA Samples 중 stream 관련 예제
- HPC GPU Optimization 가이드 - 고성능 GPU 계산 기법 미리보기