이제 여기까지 왔습니다! 우리가 지금까지 다룬 내용을 바탕으로, 간단한 실전 예제를 통해 모든 개념을 종합해볼 시간입니다. 이번 글에서는 이미지 처리 예제를 선택하겠습니다. 이유는 다음과 같습니다:
- 2D 데이터 처리(블록/스레드 2D 설정) 개념 확립
- Shared Memory나 Constant Memory를 활용해볼 수 있는 기회
- 비동기 스트림으로 데이터 전송과 연산 오버랩 시도
- Modern C++ 빌드 및 언어 기능을 통한 코드 가독성 개선
- 디버깅 및 프로파일링 연습
“간단한 이미지 필터” 예제를 통해 실제 GPU 코드가 어떻게 구조화되고, 이전에 배운 개념들이 어떻게 융합되는지 살펴보겠습니다.
실전 예제 개요: 간단한 블러(Blur) 필터 적용하기
우리가 구현할 내용은 다음과 같습니다.
- Host에서 이미지를 로딩 (간단히 그레이스케일 배열로 가정)
- Device로 이미지 데이터 전송
- GPU 커널에서 3x3 블러 필터를 적용
- 결과를 다시 Host로 가져와 확인
추가적으로:
- Constant Memory에 필터 커널(예: 3x3 가우시안 블러)을 저장
- 비동기 스트림(cudaMemcpyAsync, 커널 호출) 사용
- CMake를 통한 프로젝트 구조화
- Modern C++(람다, constexpr) 일부 문법 활용
이미지는 실제 파일 IO 대신 가상 데이터를 생성하는 것으로 대체하겠습니다. 실제로는 stb_image 등 라이브러리를 사용해 Host 메모리에 이미지를 로드할 수 있습니다.
코드 예제: CMakeLists.txt (프로젝트 설정)
CMakeLists.txt (예시)
cmake_minimum_required(VERSION 3.20 FATAL_ERROR)
project(cuda_image_blur LANGUAGES CXX CUDA)
set(CMAKE_CXX_STANDARD 20)
set(CMAKE_CUDA_STANDARD 20)
# Release 빌드 타입 권장
if(NOT CMAKE_BUILD_TYPE)
set(CMAKE_BUILD_TYPE Release)
endif()
set(CMAKE_CUDA_ARCHITECTURES 70)
add_executable(blur_app main.cu)
target_compile_options(blur_app PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-Xptxas=-v>)
여기서 main.cu에 전체 코드를 작성합니다. 실제로는 파일을 나눌 수도 있지만, 단일 파일로 정리하겠습니다.
코드 예제: main.cu
#include <cstdio>
#include <vector>
#include <cstdlib>
#include <algorithm>
#include <cassert>
// 3x3 필터를 Constant Memory에 저장
__constant__ float d_filter[9];
// GPU 커널: 2D 블러 필터 적용
__global__ void blur_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) {
float sum = 0.0f;
for (int fy = -1; fy <= 1; ++fy) {
for (int fx = -1; fx <= 1; ++fx) {
int ix = min(max(x+fx,0),N-1);
int iy = min(max(y+fy,0),M-1);
float val = static_cast<float>(input[iy * N + ix]);
float fval = d_filter[(fy+1)*3 + (fx+1)];
sum += val * fval;
}
}
output[y * N + x] = static_cast<unsigned char>(sum);
}
}
int main() {
// 이미지 크기
int M = 1024;
int N = 1024;
size_t size = M * N * sizeof(unsigned char);
// Host에서 가상 이미지 데이터 생성 (그레이스케일)
std::vector<unsigned char> h_input(M*N), h_output(M*N, 0);
{
// 간단한 패턴: rand()로 채운다
for (int i = 0; i < M*N; ++i) {
h_input[i] = static_cast<unsigned char>(rand() % 256);
}
}
// 3x3 블러 필터 (가우시안 비슷하게)
float h_filter[9] = {
1/16.0f, 2/16.0f, 1/16.0f,
2/16.0f, 4/16.0f, 2/16.0f,
1/16.0f, 2/16.0f, 1/16.0f
};
// Constant Memory에 필터 복사
cudaMemcpyToSymbol(d_filter, h_filter, 9 * sizeof(float));
// Device 메모리 할당
unsigned char *d_input, *d_output;
cudaMalloc(&d_input, size);
cudaMalloc(&d_output, size);
// 비동기 스트림 생성
cudaStream_t stream;
cudaStreamCreate(&stream);
// Host -> Device 복사 비동기
cudaMemcpyAsync(d_input, h_input.data(), size, cudaMemcpyHostToDevice, stream);
// 스레드/블록 설정(16x16)
dim3 threadsPerBlock(16,16);
dim3 blocks((N+threadsPerBlock.x-1)/threadsPerBlock.x,
(M+threadsPerBlock.y-1)/threadsPerBlock.y);
// 커널 런칭 비동기
blur_kernel<<<blocks, threadsPerBlock, 0, stream>>>(d_input, d_output, M, N);
// Device -> Host 복사 비동기
cudaMemcpyAsync(h_output.data(), d_output, size, cudaMemcpyDeviceToHost, stream);
// Host에서 다른 작업 수행(가상)
auto some_host_work = [](){
// CPU 작업 시뮬레이션
for (volatile int i = 0; i < 1000000; ++i) {}
};
some_host_work();
// 스트림 동기화로 모든 GPU 작업 완료 대기
cudaStreamSynchronize(stream);
// 결과 검증 (간단히 몇 픽셀 출력)
for (int i = 0; i < 5; ++i) {
printf("output[%d] = %u\n", i, h_output[i]);
}
// 자원 해제
cudaFree(d_input);
cudaFree(d_output);
cudaStreamDestroy(stream);
return 0;
}
주요 포인트 정리
- Constant Memory: __constant__ 메모리에 필터 저장.
- 2D Grid/Block: 이미지 처리에 맞게 2D 인덱싱 활용.
- 비동기 스트림: cudaMemcpyAsync, cudaStreamCreate로 Host-Device 동작을 겹침.
- Modern C++ 문법: 간단한 람다(some_host_work) 활용.
- CMake 빌드: cmake .. && make로 간편 빌드.
디버깅 & 프로파일링 예제 활용
이번 예제를 가지고 Nsight Systems나 Nsight Compute를 돌려볼 수 있습니다.
- nsys profile -o report ./blur_app → 전체 타임라인 분석
- ncu ./blur_app → 커널 레벨 성능 분석
이를 통해 어떤 커널이 오래 걸리고, 메모리 대역폭 활용이 어떤지 등을 파악한 뒤 Shared Memory 최적화나 다른 개선을 시도해볼 수 있습니다.
Before vs After: 개념 학습 → 실전 예제
Before:
- 여러 개념을 따로따로 이해했으나, 실제 코드 예제가 부족함
After:
- Host-Device 메모리 전송, Constant Memory 활용, 2D 블록, 비동기 스트림, Modern C++ 빌드, 디버깅/프로파일링까지 모두 적용한 하나의 흐름 파악
- 향후 다른 이미지 처리나 대규모 행렬 연산, 과학적 시뮬레이션 등으로 확장 가능
이렇게 한 번에 종합한 예제를 통해, 실제 프로젝트에 CUDA와 Modern C++를 접목할 때 어떤 식으로 코드를 구조화할지 감을 잡을 수 있습니다.
다음 글 예고
다음(마지막) 글(#10)에서는 지금까지의 여정을 마무리하고, 추가로 공부할 만한 Intermediate/Advanced 주제를 소개하겠습니다. 이 시리즈를 통해 기반을 탄탄히 쌓았다면, 이제 한 단계 더 나아가 복잡한 알고리즘 최적화나 GPU 가속 라이브러리 활용, 심지어 Multi-GPU 확장까지 고려해볼 수 있을 겁니다.
유용한 링크 & 리소스
- NVIDIA CUDA Samples: 실제 이미지 처리 예제 참조
- OpenCV GPU 모듈: 실제 이미지 프로세싱 라이브러리와 결합 예제
- stb_image: 실제 이미지 로딩을 위한 헤더 라이브러리
- Modern CMake 예제: CMake 활용 패턴
반응형