CUDA 이미지 처리 파이프라인 최적화

이 요약은 AI가 원문을 분석해 생성했습니다. 정확한 내용은 원문 기준으로 확인하세요.

TL;DR

NVIDIA는 RGB 이미지를 grayscale로 변환하고 32×32 타일별 median을 계산하는 CUDA 예제를 여섯 단계로 최적화한다. Compute Sanitizer와 CCCL의 인덱싱·mdspan API로 shared memory의 범위를 벗어난 접근을 잡고, Nsight Systems와 NVTX로 median 커널이 병목임을 확인한다. 이후 CUB의 GPU 알고리즘, 메모리 풀 기반 컨테이너, pinned host memory, 이미지별 CUDA Stream과 비동기 복사를 차례로 적용해 세 이미지 처리 시간을 6.8초에서 23밀리초로 줄인다. 저수준 최적화 없이도 최종 실행 시간이 약 300배 단축된다.

빠른 이해

새로운 점

저수준 GPU 최적화 없이 CCCL·CUB의 안전한 API와 메모리·Stream 관리만 단계적으로 바꿔 6.8초의 이미지 파이프라인을 23밀리초로 줄인 실습형 접근이다.

핵심 메커니즘

CPU의 RGB 이미지 입력을 pinned host memory에서 이미지별 CUDA Stream으로 비동기 복사하고, CUB의 DeviceTransform으로 grayscale 변환한 뒤 각 32×32 타일을 BlockRadixSort로 정렬해 median을 산출한다. device memory pool은 GPU 버퍼를 재사용하고, 여러 Stream은 복사·커널 실행을 겹치며, 결과 median은 다시 CPU로 비동기 복사된다.

핵심 수치

초기 전체 처리 시간: 6.8초- 세 이미지 처리, Nsight Systems 측정
초기 median 계산 시간: 이미지당 약 2.1초- computeMedian 커널
CUB 적용 후 median 계산: 773마이크로초- 기존 대비 2,717배 빠름
CUB 적용 후 전체 처리: 635밀리초- 초기 대비 10배 빠름
pinned memory 적용 후 전체 처리: 약 25밀리초- host-to-device 전송 10배 단축
최종 전체 처리: 약 23밀리초- 초기 6.8초 대비 약 300배 빠름

섹션별 상세

출발점: RGB 이미지 처리 파이프라인

예제는 CPU에 있는 red·green·blue 이미지를 GPU로 옮긴 뒤 grayscale로 변환하고, 각 32×32 픽셀 타일의 값을 정렬해 가운데 값을 median으로 선택한다. 마지막으로 타일별 median 배열을 GPU에서 CPU로 복사하며, 이 과정을 세 이미지에 대해 OpenMP로 병렬 실행한다. 기존 구현은 RGB 변환과 median 계산을 직접 작성한 CUDA 커널에 맡기고, 각 이미지마다 GPU 메모리를 할당한 뒤 작업이 끝나면 해제한다. 따라서 shared memory 인덱싱 오류, 비효율적인 수제 알고리즘, 반복 할당, pageable 메모리 전송, default Stream의 직렬 실행이 함께 성능과 안정성을 제한한다.

Compute Sanitizer와 CCCL로 메모리 오류 차단

기본 코드는 illegal memory access를 발생시키며, Compute Sanitizer는 computeMedian 커널의 shared memory에 대한 범위 초과 쓰기를 0_base_error_example.cu 55번째 줄과 블록·스레드 위치까지 찾아낸다. 문제의 원인은 블록 내부 shared memory에 데이터를 적재하면서 이미지 전체에서 쓰는 global index를 그대로 사용한 데 있다. CCCL의 cuda::launch API는 global 좌표와 block 좌표를 구분해 커널에 전달하고, cuda::std::span·cuda::std::mdspan은 raw pointer 대신 연속 메모리의 범위를 표현해 debug 모드에서 잘못된 접근을 assertion으로 잡는다. shared memory에도 cuda::shared_memory_mdspan을 적용하면 타일의 행·열 범위를 검사할 수 있어, 성능 최적화 전에 커널의 기능적 오류를 제거하는 흐름이 마련된다.

cpp

auto config = cuda::make_config(cuda::block_dims(...), cuda::grid_dims(...));
cuda::launch(stream, config, kernel_name, input)

새로운 launch API로 커널의 블록·그리드 구성을 지정하고 특정 Stream에서 실행한다.

Nsight Systems로 실제 병목 측정

오류를 제거한 뒤에는 Nsight Systems에서 함수 호출과 GPU 작업의 시간 순서를 확인하고, NVTX 범위로 전체 이미지 처리와 median 커널을 구분한다. 초기 타임라인에서 GPU 시간의 98.5%가 커널 실행에, 1.5%가 메모리 작업에 쓰였으며, 세 이미지의 전체 처리 시간은 6.8초로 나타난다. 특히 computeMedian이 이미지 하나당 약 2.1초를 차지해 RGB 변환보다 훨씬 긴 실행 시간을 만든다. 이 측정 결과는 메모리 전송보다 타일 정렬 커널을 먼저 교체해야 전체 시간에 가장 큰 변화가 생긴다는 우선순위를 정한다.

cpp

cub::DeviceTransform::Transform(
    cuda::std::make_tuple(d_image_r, d_image_g, d_image_b),
    d_image_gray,
    IMAGE_SIZE,
    [] __host__ __device__ (pixel_t r, pixel_t g, pixel_t b)
    { return static_cast<pixel_t>(0.299f * r + 0.587f * g + 0.114f * b); },
    stream);

CUB의 DeviceTransform으로 RGB 세 입력을 GPU에서 grayscale 출력으로 변환한다.

CUB로 GPU 알고리즘 교체

수제 RGB 변환 커널은 세 입력 iterator의 튜플을 GPU에서 처리하는 cub::DeviceTransform::Transform으로 바꾸고, RGB 각 채널에 0.299f·0.587f·0.114f 가중치를 곱해 grayscale 픽셀을 출력한다. median 계산에서는 스레드 블록이 global memory에서 32×32 타일을 읽은 뒤 CUB의 BlockRadixSort로 block-level 정렬을 수행하고, 정렬된 중간 위치의 값을 결과 배열에 기록한다. 이 방식은 공통 병렬 알고리즘의 최적화 구현을 재사용해 직접 작성한 정렬 로직과 단일 스레드 중심의 처리 부담을 줄인다. 그 결과 median 계산 시간은 이미지당 2.1초 수준에서 773마이크로초로 줄어 2,717배 빨라졌고, 전체 세 이미지 처리는 635밀리초로 단축됐으며 병목은 메모리 할당으로 이동한다.

cpp

cuda::device_memory_pool_ref device_resource = cuda::device_default_memory_pool(cuda::device_ref{0});
cuda::device_buffer d_image_r = cuda::make_buffer(stream, device_resource, IMAGE_SIZE, cuda::no_init);

GPU 메모리 풀을 백엔드로 사용하는 device_buffer를 생성해 반복적인 할당·해제 비용을 줄인다.

메모리 풀과 pinned memory 적용

반복되는 cudaMalloc·cudaFree는 GPU 작업 중 불필요한 할당 비용을 만들므로, CCCL의 cuda::device_buffer와 cuda::device_memory_pool_ref를 사용해 메모리를 풀에서 재사용한다. device_buffer는 범위를 벗어날 때 자동으로 해제되고, 같은 풀의 반복 할당·해제에서는 매번 전체 cudaMalloc·cudaFree 비용을 지불하지 않는다. 이 변경으로 메모리 할당 시간은 거의 사라지고 이미지 처리 시간은 2.6배 개선되며, 남은 대부분의 시간은 CPU에서 GPU로 red·green·blue 데이터를 옮기는 데 쓰인다. 기본 pageable host memory에서는 CUDA 드라이버가 임시 page-locked 배열을 만들고 한 번 더 복사하므로, cuda::host_buffer와 cuda::make_pinned_buffer로 CPU 입력을 직접 pinned memory에 배치한다. 그 결과 세 이미지의 전체 시간은 약 25밀리초로 줄어 host-to-device 전송이 10배 빨라진다.

이미지별 CUDA Stream으로 작업 중첩

pinned memory를 적용한 뒤에도 서로 다른 OpenMP 스레드의 작업이 default Stream 하나에서 순서대로 실행되어 GPU 병렬성이 제한된다. 각 이미지 처리 반복문에 cuda::stream을 하나씩 만들고, cuda::copy_bytes로 host-to-device·device-to-host 전송을 비동기로 실행하면 이미지별 작업 큐가 분리된다. CUB 호출과 커널은 전달받은 Stream에서 기본적으로 비동기 실행되며, 초기 pinned buffer 할당에 사용한 init_stream은 다른 Stream에서 작업을 시작하기 전에 동기화한다. 각 반복의 마지막 stream.sync()는 CPU가 결과를 읽기 전에 복사가 끝났는지 보장한다. 세 Stream에서 메모리 전송과 커널이 겹친 최종 실행 시간은 약 23밀리초로, 초기 6.8초보다 약 300배 짧아진다.

cpp

cuda::copy_bytes(stream, h_images_r[i], d_image_r);
cub::DeviceTransform::Transform(..., stream.get());
cuda::launch(stream, ...);
cuda::copy_bytes(stream, d_median, h_medians[i]);
stream.sync();

각 스레드의 Stream에서 CPU-GPU 복사와 GPU 연산을 비동기로 실행한 뒤 결과 접근 전에 동기화한다.

용어 해설

CUDA C++ Core Libraries(CCCL): — NVIDIA가 제공하는 CUDA C++용 라이브러리 모음으로, GPU 알고리즘·메모리 컨테이너·Stream·인덱싱 API를 표준화된 형태로 제공한다. 원문에서는 CUB, cuda::device_buffer, cuda::host_buffer, cuda::launch 같은 기능을 통해 CUDA 코드의 안전성·성능·유지보수성을 높이는 기반으로 활용된다.
Compute Sanitizer: — CUDA 프로그램의 메모리 접근과 기능적 정확성을 검사하는 NVIDIA 도구다. 커널 실행 중 발생한 잘못된 shared memory 접근 위치와 스레드·블록 정보를 표시해, 일반적인 CUDA 오류 메시지만으로 찾기 어려운 인덱싱 버그를 추적하는 데 쓰인다.
NVTX: — CUDA 애플리케이션의 특정 코드 구간에 이름을 붙이는 주석·범위 표시 API다. Nsight Systems 타임라인에서 함수와 커널, 메모리 전송 구간을 구분해 병목이 어디에 있는지 확인할 수 있도록 한다.
페이지 고정 메모리(Pinned Memory): — CPU 메모리의 물리적 위치를 고정해 CUDA가 직접 전송할 수 있도록 만든 메모리다. 일반 pageable 메모리를 사용할 때 드라이버가 임시 pinned 배열을 거치는 추가 복사 단계를 줄여 host-to-device 전송 시간을 단축한다.
CUDA Stream: — GPU 커널 실행과 메모리 전송을 순서대로 배치하는 작업 큐다. 이미지나 OpenMP 스레드마다 별도 Stream을 할당하고 비동기 작업을 사용하면 서로 다른 입력의 커널과 전송이 겹쳐 GPU 유휴 시간을 줄일 수 있다.

기술

CUDA
CCCL
Compute Sanitizer
cuda::launch
cuda::std::span
cuda::std::mdspan
cuda::shared_memory_mdspan
NVIDIA Nsight Systems
NVTX
CUB
cub::DeviceTransform::Transform
cub::BlockRadixSort
cuda::device_buffer
cuda::device_memory_pool_ref
cuda::host_buffer
cuda::make_pinned_buffer
cuda::stream
cuda::copy_bytes
OpenMP
Google Colab

언급된 리소스

문서원문 링크