[YouTube] Asynchrony and CUDA Streams

[원본] https://youtu.be/pyW9St8uM8w?si=sjfE8KZFbsHMnuzt

Google NotebookLM으로 작성된 후 일부 수정된 글입니다.

당신의 GPU 코드는 생각보다 느립니다: CUDA 성능을 2배 끌어올린 4가지 놀라운 발견

서론: 시작하며

GPU를 활용하여 복잡한 계산을 수행하는 코드를 성공적으로 작성하고 실행했을 때의 만족감은 큽니다. 하지만 코드가 작동한다는 사실에 안주하는 동안, 당신의 시스템은 가진 잠재력의 절반도 발휘하지 못하고 있을지 모릅니다. 가장 흔하면서도 비효율적인 상황을 떠올려보십시오: GPU가 열심히 다음 시뮬레이션 단계를 계산하는 동안, CPU는 그저 작업이 끝나기만을 기다리며 아무것도 하지 않습니다. 이는 시스템 자원의 명백한 낭비입니다.

이 글의 목적은 단순히 코드를 몇 줄 수정하여 조금 더 빠르게 만드는 방법을 알려주는 것이 아닙니다. 성능 저하의 근본적인 원인을 깊이 파헤치고, NVIDIA 개발자들이 실제 최적화 과정에서 발견한 놀랍고 때로는 직관에 반하는 교훈들을 공유하는 데 있습니다. 지금부터 4가지 핵심적인 발견을 통해 코드에 숨어있는 비효율을 걷어내고 시스템의 잠재력을 최대한으로 끌어내는 여정을 시작하겠습니다.

--------------------------------------------------------------------------------

본문: CUDA 성능 최적화의 4가지 교훈

1. CPU를 놀리지 마세요: 비동기 처리로 숨겨진 성능을 찾다

성능 최적화의 첫 단추는 '동기(Synchronous)'와 '비동기(Asynchronous)' 작업의 근본적인 차이를 이해하는 것에서 시작합니다. thrust::tabulate와 같은 동기 함수는 CPU가 GPU에 계산을 요청하면 GPU의 작업이 완전히 끝날 때까지 멈춰서 기다리는 '블로킹(blocking)' 방식으로 작동합니다. GPU가 계산하는 동안 CPU는 아무것도 하지 않으므로 귀중한 자원이 낭비됩니다. 반면, CUB 라이브러리의 cub::transform과 같은 비동기 함수는 다릅니다. CPU가 GPU에 작업을 요청하는 즉시 제어권을 돌려받는 '논블로킹(non-blocking)' 방식으로, CPU는 GPU가 계산을 수행하는 동안 다른 유용한 작업을 처리할 수 있습니다.

이러한 비동기 방식의 힘은 cub::transform을 사용할 때 명확해집니다. GPU가 다음 시뮬레이션 단계를 계산하는 동안, CPU는 이전 단계에서 계산된 결과 데이터를 디스크에 쓰는 작업을 동시에 수행할 수 있습니다. 각자 다른 데이터를 사용하기 때문에 두 작업은 서로 의존하지 않아 완벽한 병렬 처리가 가능합니다. 이 간단한 접근법만으로도 전체 시스템 활용률이 극적으로 향상되어 잠재적으로 **"2배의 속도 향상"**을 기대할 수 있습니다.

GPU가 복잡한 연산을 수행하는 동안 CPU는 그저 기다리기만 합니다. 이 유휴 시간을 활용하는 것이 성능 최적화의 첫걸음입니다. 비동기 처리는 CPU와 GPU가 각자의 작업을 동시에 수행하는 진정한 병렬 처리를 가능하게 합니다.

그러나 CPU와 GPU가 동시에 작동하기 시작하자, 우리는 새로운 문제에 직면했습니다. 대체 어디서 시간이 낭비되고 있는지 코드만으로는 알 수 없게 된 것입니다.

2. 보이지 않으면 최적화할 수 없습니다: Nsight Systems로 병목 현상 가시화하기

비동기 코드는 강력하지만, CPU와 GPU에서 동시에 여러 작업이 일어나기 때문에 코드만 보고서는 어느 지점에서 시간이 지연되는지, 즉 성능 병목 지점을 찾거나 디버깅하기가 매우 어렵습니다. 추측에 기반한 최적화는 시간 낭비일 뿐입니다.

이때 필요한 것이 바로 NVIDIA의 프로파일링 도구인 **'Nsight Systems'**입니다. 이 도구는 CPU와 GPU에서 일어나는 모든 활동을 시간의 흐름에 따라 정밀하게 시각화한 타임라인을 제공합니다. 이를 통해 우리는 CPU가 언제 유휴 상태에 있는지, GPU 커널은 언제 실행되는지, 그리고 두 작업이 얼마나 효율적으로 중첩되는지를 한눈에 파악할 수 있습니다. 복잡한 애플리케이션에서는 **'NVTX (NVIDIA Tools Extensions API)'**를 사용하여 분석의 깊이를 더할 수 있습니다. 코드의 특정 구간에 copy, compute, write와 같은 사용자 정의 레이블(range)을 추가하면, 이 레이블이 Nsight Systems 타임라인에 그대로 표시되어 복잡한 타임라인을 코드와 직접 연결하여 훨씬 쉽고 직관적으로 분석할 수 있습니다.

추측에 기반한 최적화는 시간 낭비일 뿐입니다. Nsight Systems와 같은 프로파일링 도구는 CPU와 GPU의 상호작용을 눈으로 직접 확인하게 해 주어, 데이터에 기반한 정확한 성능 개선을 가능하게 합니다.

Nsight Systems를 통해 우리의 코드를 들여다보자, 우리는 또 다른 미스터리와 마주하게 됩니다. 다음 단계로 나아갈 준비가 되셨나요?

3. '비동기' 호출만으로는 부족합니다: CUDA 스트림의 숨겨진 규칙

우리는 cudaMemcpyAsync와 같은 비동기 함수를 사용하여 메모리 복사와 GPU 연산을 분리했습니다. 당연히 두 작업이 중첩되어 동시에 실행될 것이라 기대했습니다. 하지만 Nsight Systems의 타임라인은 충격적인 사실을 보여주었습니다. 두 작업은 여전히 순차적으로 실행되며 전혀 중첩되지 않았습니다. 왜일까요?

원인은 CUDA의 '기본 스트림(default stream)' 동작 방식에 있습니다. 우리가 별도의 스트림을 지정하지 않으면, 모든 비동기 작업은 GPU 내의 단 하나의 작업 큐, 즉 '기본 스트림'에 순서대로 쌓입니다. GPU는 이 큐에 들어온 순서에 따라 작업을 하나씩 처리할 뿐입니다. CPU 입장에서는 비동기적으로 명령을 내렸을지라도, GPU 입장에서는 모든 작업을 동기적으로 처리하는 것과 같은 결과가 나타나는 것입니다. 이 문제를 해결하기 위해 **'CUDA 스트림'**을 직접 생성하고 사용해야 합니다. 예를 들어, 메모리 복사 작업은 '복사 스트림'에 할당하고, GPU 연산 작업은 '연산 스트림'에 할당하는 것입니다. 서로 다른 스트림에 속한 독립적인 작업들은 GPU가 순서에 상관없이 동시에 실행할 수 있어 비로소 우리가 의도했던 진정한 병렬 처리가 가능해집니다.

GPU에게 비동기 명령을 내리는 것만으로는 충분하지 않습니다. 별개의 'CUDA 스트림'을 사용해 작업을 명확히 분리해야만, GPU는 비로소 여러 작업을 동시에 처리하며 진정한 병렬성의 힘을 발휘합니다.

이제 비동기 함수와 별도의 스트림이라는 두 가지 무기를 모두 갖췄습니다. 하지만 Nsight Systems를 다시 확인했을 때, 우리는 최적화 여정의 가장 깊은 곳에 숨겨진 마지막 비밀과 마주하게 됩니다.

4. 진짜 병목은 메모리였습니다: '고정 메모리'가 모든 것을 바꾸는 이유

비동기 함수와 별도의 CUDA 스트림까지 사용했지만, Nsight Systems 타임라인에서는 여전히 메모리 복사와 GPU 연산이 중첩되지 않는 더욱 미스터리한 상황이 펼쳐졌습니다. 코드는 완벽해 보이는데 도대체 무엇이 문제일까요? 원인은 우리 코드의 로직이 아닌, 컴퓨터의 메모리가 작동하는 방식에 숨어있었습니다.

기본적으로 프로그램이 할당하는 메모리는 '페이징 가능 메모리(Pageable Memory)'입니다. 이는 운영체제가 필요에 따라 해당 메모리 데이터를 물리적 RAM에서 디스크로 옮길 수 있다는 의미입니다. 하지만 GPU는 데이터 전송 중 데이터가 갑자기 디스크로 옮겨지는 위험을 감수할 수 없습니다. 따라서 GPU는 물리적 RAM에 항상 상주하도록 '고정된(pinned)' 메모리에만 직접적이고 비동기적인 접근이 가능하다는 핵심적인 제약 사항을 가집니다.

우리가 일반적인 '페이징 가능 메모리'를 사용해 비동기 복사를 요청하면, CUDA 드라이버는 데이터를 GPU로 직접 보내지 못합니다. 대신, 내부적으로 숨겨진 **'고정 메모리 스테이징 버퍼'**로 데이터를 먼저 동기적으로 복사한 후, 그 버퍼에서 GPU로 데이터를 비동기 전송합니다. 이 숨겨진 복사 단계가 동기적으로 작동해야만 하는 이유는 명확합니다. 만약 비동기적으로 복사하는 동안 운영체제가 원본인 페이징 가능 메모리를 디스크로 옮겨버린다면, 데이터가 손상될 위험이 있기 때문입니다. CUDA 드라이버는 이 위험을 방지하기 위해 데이터가 고정 메모리로 완전히 옮겨질 때까지 모든 작업을 멈추고 기다리는 것입니다. 바로 이 '숨겨진 동기 복사' 단계 때문에 우리가 의도한 비동기 중첩 효과가 사라졌던 것입니다.

이 문제의 해결책은 놀랍도록 간단합니다. thrust::host_vector 대신, 이름에 'universal'이 붙은 것처럼 CPU와 GPU 양쪽에서 모두 접근 가능하여 코드를 단순하게 유지하면서도 처음부터 '고정 메모리'를 할당하는 thrust::universal_host_pinned_vector와 같은 컨테이너를 사용하는 것입니다. 이 간단한 타입 변경만으로 숨겨진 동기 단계를 완전히 건너뛸 수 있습니다. 그 결과, Nsight Systems 타임라인에서 마침내 메모리 복사와 GPU 연산이 완벽하게 중첩되는 모습을 확인할 수 있습니다.

가장 교묘한 성능 병목은 코드 라인이 아닌 시스템 아키텍처에 숨어 있습니다. GPU는 '고정 메모리'라는 특별한 통로로만 데이터를 비동기적으로 받을 수 있으며, 이 사실을 아는 것이 진정한 성능 최적화의 마지막 열쇠입니다.

--------------------------------------------------------------------------------

결론: 맺음말

단순한 비동기 함수 호출에서 시작하여 Nsight Systems를 통한 프로파일링, CUDA 스트림을 이용한 작업 분리, 그리고 시스템 메모리 아키텍처의 깊은 이해에 이르기까지, 우리는 숨겨진 성능 병목을 찾아 해결하는 여정을 함께했습니다. 이를 통해 배운 가장 중요한 교훈은, 성능 최적화란 단순히 더 빠른 코드를 작성하는 기술이 아니라 시스템 전체가 어떻게 상호작용하는지를 깊이 있게 이해하는 과정이라는 것입니다.

이제 당신의 코드를 다시 한번 살펴보십시오. 당신의 코드 성능을 조용히 갉아먹고 있는 또 다른 '숨겨진 가정'은 무엇일까요? 이 글에서 얻은 통찰력을 바탕으로 시스템의 잠재력을 최대한 끌어내는 새로운 발견을 시작해 보시길 바랍니다.