이제까지 우리는 CUDA의 기본 문법, 메모리 계층, 비동기 스트림, 2D 데이터 처리 등 다양한 기초 개념을 다뤄왔습니다. 하지만 이 모든 것이 현업이나 프로젝트에서 활용되기 위해서는 체계적인 빌드 환경과 현대적 C++ 기능과의 자연스러운 결합이 중요합니다. 이번 글에서는 CMake를 통한 빌드 시스템 설정 방법과 C++20/23 기능(Coroutine, Concept, Constexpr, Lambda 개선 등)을 CUDA 코드와 융합하여 더 깔끔하고 유지보수하기 쉬운 코드를 만드는 방법을 살펴보겠습니다.
CMake로 CUDA 빌드 환경 현대적으로 관리하기
기본 구조 복습
이전 글에서 간단히 소개했듯이, CMake는 CUDA를 별도의 언어로 인식하고, project(... LANGUAGES CUDA CXX) 형태로 프로젝트를 정의하면 NVCC와 C++ 컴파일러를 함께 다룰 수 있습니다.
예시 CMakeLists.txt:
cmake_minimum_required(VERSION 3.20 FATAL_ERROR)
project(my_cuda_project LANGUAGES CXX CUDA)
set(CMAKE_CXX_STANDARD 20)
set(CMAKE_CUDA_STANDARD 20)
# 빌드 타입, 최적화 옵션 설정
if(NOT CMAKE_BUILD_TYPE)
set(CMAKE_BUILD_TYPE Release)
endif()
# CUDA 아키텍처 설정 (예: sm_70)
set(CMAKE_CUDA_ARCHITECTURES 70)
add_executable(my_app main.cu utils.cu)
위와 같이 하면 CMake가 CUDA 소스(.cu) 파일을 NVCC로 처리하고, C++20 표준과 CUDA 표준을 모두 적용합니다. 또한 CMAKE_CUDA_ARCHITECTURES를 통해 특정 GPU 아키텍처에 맞춰 최적화된 코드를 생성하도록 설정할 수 있습니다.
CUDA 전용 타겟 속성 활용
CMake는 CUDA 타겟에 대해 전용 속성을 지정할 수 있습니다.
target_compile_options(my_app PRIVATE $<$<COMPILE_LANGUAGE:CUDA>:-Xptxas=-v>)
이런 식으로 CUDA 관련 전용 컴파일 옵션을 설정할 수 있습니다. 이를 통해 디버깅, 프로파일링 옵션을 손쉽게 활성화하거나, 특정 최적화 플래그를 붙일 수 있습니다.
Modern C++ 기능을 CUDA 코드에 적용하기
C++20/23에서는 다양한 언어 기능들이 추가되며, 이들은 GPU 코드에도 점진적으로 적용 가능합니다. CUDA는 C++ 표준을 점진적으로 지원해왔으며, 새로운 NVCC나 Clang CUDA 컴파일러와의 조합을 통해 람다 표현식, 콘셉트, constexpr 등을 GPU 코드에도 활용할 수 있습니다.
람다 표현식과 CUDA
람다(lambda) 표현식은 간결한 코드 작성에 유용합니다. 예를 들어, 간단한 Host 코드를 GPU 호출 사이사이에 정리하는데 람다를 써서 가독성을 높일 수 있습니다.
auto init_data = [&](std::vector<float>& v){
for (auto& x : v) x = static_cast<float>(rand()) / RAND_MAX;
};
std::vector<float> h_a(N), h_b(N), h_c(N);
init_data(h_a);
init_data(h_b);
또한 CUDA 11부터는 GPU 내부에서도 람다 표현식을 지원하는 방향으로 발전하고 있어, 향후 커널 내 람다 활용도 점차 쉬워질 것입니다(다만 모든 C++20 기능이 완전히 지원되는 것은 아니므로 최신 문서 확인 필요).
Concepts(개념) 활용하기
Concept는 템플릿 인자를 제약하는 기능입니다. GPU 코드에서 템플릿 커널을 짤 때, 특정 타입이 GPU에서 사용할 수 있는 타입인지, 혹은 특정 연산을 지원하는지 명확히 제약할 수 있습니다.
template<typename T>
concept GpuArithmetic = requires(T a, T b) {
{ a + b } -> std::convertible_to<T>;
};
__global__ void vector_add_kernel(const T* a, const T* b, T* c, int N) requires GpuArithmetic<T> {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < N) {
c[idx] = a[idx] + b[idx];
}
}
위 예제는 개념적이며, 실제로는 CUDA 컴파일러 지원 상태를 확인해야 하지만, 이런 식으로 GPU에서 템플릿 타입을 안전하게 활용할 때 Concept가 가독성과 안정성을 높여줄 수 있습니다.
Constexpr 함수와 GPU 코드
CUDA 커널이나 디바이스 함수(__device__ 함수)에도 constexpr 함수 적용이 가능해지면서, GPU 코드 컴파일 시점에 상수로 계산할 수 있는 부분을 미리 처리할 수 있습니다.
__device__ __host__ constexpr float scale_value(float x) {
return x * 2.0f;
}
이렇게 하면 Host와 Device 양쪽에서 활용 가능한 constexpr 함수를 정의할 수 있고, GPU 커널 호출 시 상수 표현 최적화가 가능해지는 경우도 있습니다.
Before vs After: 전통적 빌드 스크립트 vs 현대적 CMake + Modern C++
Before:
- Makefile 수작업 및 NVCC 옵션 일일이 지정
- C++11 수준의 코드 작성(람다/auto 정도만 활용), 제한된 표현력
- GPU 코드와 Host 코드가 명확히 구분되지 않아 유지보수성 감소
After:
- CMake를 통한 자동화된 빌드, GPU 아키텍처 별 코드생성 옵션 간편 관리
- Modern C++ 기능(C++20/23)으로 제약 조건(Concept), 상수 표현(Constexpr), 람다 활용으로 가독성 향상
- 유지보수성과 확장성이 크게 증가
이러한 변화는 팀 개발, 대규모 프로젝트, 장기 유지보수에 특히 빛을 발합니다.
다음 글 예고
다음 글(#8)에서는 디버깅과 간단한 성능 프로파일링 방법을 다룹니다. 현대적 C++와 CMake를 갖추고, 기본 최적화를 살짝 맛본 지금 시점에서, 실제로 GPU 코드가 예상대로 동작하는지 확인하고 성능 병목 구간을 찾아내는 방법을 배워보겠습니다.