[YouTube] Accelerating Applications with Parallel Algorithms

[원본]  https://youtu.be/Sdjn9FOkhnA?si=M-KymEPJ2WO1wfBP

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

 

GPU 프로그래밍, 당신이 놓치고 있는 몇 가지 놀라운 진실

"GPU를 사용하면 모든 것이 그냥 빨라질 거야." 우리 중 많은 이들이 이런 기대를 품고 GPU 프로그래밍의 세계에 발을 들입니다. 하지만 막상 코드를 돌려보면 기대했던 만큼의 성능 향상이 나타나지 않거나, 때로는 오히려 더 느려지는 기묘한 현상을 마주하곤 합니다. 왜 그럴까요?

최근 NVIDIA의 공식 CUDA 강의를 살펴보면서, 저는 GPU 프로그래밍에 대한 흔한 오해를 깨부수는 몇 가지 핵심 원칙들을 발견했습니다. 이 원칙들은 단순히 기술적인 팁을 넘어, GPU 아키텍처를 근본적으로 다르게 생각하도록 만드는 놀라운 통찰을 담고 있었습니다. 이 글에서는 그중 가장 놀랍고 직관에 반하는 몇 가지 진실을 공유하고자 합니다.

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

1. GPU는 자동차가 아니라 '버스'입니다

CPU와 GPU의 가장 근본적인 차이는 무엇일까요? '자동차와 버스' 비유는 이 질문에 대한 완벽한 답을 제공하며, 실제 하드웨어 수치가 이를 뒷받침합니다.

• CPU는 자동차입니다. 단일 데이터에 대한 접근 속도(지연 시간)가 약 100ns로 매우 낮아 빠르고 민첩합니다. 하지만 한 번에 옮길 수 있는 데이터의 양(대역폭)은 약 100GB/s로 상대적으로 적습니다. 소수의 사람을 태우고 빠르게 목적지에 도착하는 스포츠카와 같습니다.
• GPU는 버스입니다. 데이터 접근 지연 시간이 약 500ns로 CPU보다 5배나 높아 출발이 굼뜹니다. 하지만 한 번 움직이기 시작하면 약 1,000GB/s(1TB/s)라는 압도적인 대역폭으로 엄청난 양의 데이터를 실어 나릅니다. 승객이 가득 찰 때 비로소 진가를 발휘하는 대용량 버스입니다.

이 비유는 실제 데이터로 증명됩니다. 한 실험에서 512개의 객체만 시뮬레이션했을 때, 단일 스레드 CPU는 GPU보다 훨씬 더 빨랐습니다. 512명 정도의 사람을 옮기는 데 굳이 거대한 버스를 동원할 필요가 없는 것과 마찬가지입니다.

객체 수가 64,000개로 늘어나자, 멀티스레드 CPU가 싱글스레드 CPU를 앞서기 시작했지만, GPU는 여전히 제 성능을 내지 못했습니다. 버스를 채우기에는 여전히 승객이 부족했던 것입니다.

하지만 시뮬레이션 대상을 2억 개의 객체로 늘리자 상황은 완전히 역전되었습니다. GPU는 CPU보다 무려 10배 더 빠른 성능을 보여주었습니다. 마침내 버스를 가득 채울 만큼 충분한 승객이 생긴 것입니다.

핵심 통찰: 따라서 GPU 프로그래밍의 첫걸음은 "내 문제가 버스를 채울 만큼 충분히 큰가?"를 자문하는 것입니다. 문제의 '규모'가 GPU의 성능을 결정하는 가장 중요한 열쇠입니다.

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

2. 코드는 저절로 GPU에서 실행되지 않습니다

많은 초보자들이 NVCC(NVIDIA CUDA 컴파일러)로 코드를 컴파일하기만 하면 마법처럼 코드가 GPU에서 실행될 것이라고 착각합니다. 하지만 현실은 그렇지 않습니다. 프로그래머는 코드가 어디서, 어떻게 실행되어야 할지 명시적으로 '지시'해야 합니다. 여기에는 두 가지 중요한 개념이 있습니다.

1. 실행 공간 지정자 (Execution Specifier) __host__와 __device__ 같은 키워드가 여기에 해당합니다. 이것은 컴파일 시점에 NVCC 컴파일러에게 보내는 지시입니다. "이 함수를 위한 CPU용 기계어 코드를 생성해라" (__host__) 또는 "GPU용 기계어 코드를 생성해라" (__device__) 라고 알려주는 것입니다. 즉, __device__는 GPU에서 실행될 코드의 '설계도'를 준비하라는 지시일 뿐, 실제로 공장을 가동하라는 명령은 아닙니다.
2. 실행 정책 (Execution Policy) thrust::host와 thrust::device 같은 파라미터가 여기에 해당합니다. 이것은 런타임 시점에 Thrust와 같은 라이브러리에게 보내는 지시입니다. "컴파일된 코드를 실제로 CPU에서 실행할지, GPU에서 실행할지"를 결정하는, 즉 공장을 가동하라는 실제 실행 명령입니다.

만약 어떤 함수를 __device__로만 지정해놓고 thrust::host 정책으로 실행하려고 하면 컴파일 오류가 발생합니다. CPU에서 실행할 수 있는 기계어 코드(설계도)가 생성되지 않았기 때문입니다. 이처럼 두 가지 개념은 반드시 함께 가야 하며, 컴파일 시점의 유연성과 런타임 시점의 결정을 분리하여 프로그래머에게 더 많은 제어권을 주기 위해 존재합니다.

핵심 통찰: GPU 프로그래밍은 단순히 코드를 옮기는 것이 아니라, 컴파일러와 런타임 모두에게 작업의 위치와 방식을 명확히 '지시'하는 정교한 과정입니다.

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

3. 최고의 최적화: '임시 데이터'를 만들지 않는 기술

두 배열의 각 요소 간 온도 차이 중 최댓값을 구하는 간단한 문제를 생각해 봅시다. 순진한 접근법은 다음과 같습니다.

1. 두 배열(A, B)의 차이를 계산하여 새로운 '임시 배열'(C)에 저장합니다. (transform 연산)
2. 이 '임시 배열'(C)에서 최댓값을 찾습니다. (reduce 연산)

이 방식은 논리적으로는 맞지만, GPU에서는 치명적입니다. transform 연산에서 각 요소에 대해 2번의 읽기(배열 A, B)와 1번의 쓰기(임시 배열 C)가 발생하고, reduce 연산에서 다시 1번의 읽기(임시 배열 C)가 발생하여 총 3N번의 읽기와 1N번의 쓰기가 일어납니다. 불필요한 메모리 할당과 데이터 이동으로 엄청난 대역폭을 낭비하는 셈입니다.

이때 등장하는 것이 '팬시 이터레이터(Fancy Iterators)'라는 강력한 개념입니다. 팬시 이터레이터의 핵심은 **"데이터를 임시 배열에 저장하는 대신, 필요할 때마다 '즉석에서(on the fly)' 값을 계산하여 보여주는 가상의 뷰(view)를 만드는 것"**입니다.

작동 방식은 이렇습니다. 먼저 zip_iterator가 두 배열(A, B)의 같은 위치에 있는 원소(예: A[i], B[i])를 하나의 튜플 (A[i], B[i])처럼 논리적으로 묶어줍니다. 그다음 transform_iterator가 이 튜플을 입력받아 즉석에서 abs(A[i] - B[i]) 연산을 수행하여 그 결과값만 reduce 함수에 전달합니다. reduce 함수는 이 가상의 뷰를 마치 실제 배열처럼 순회하며 최댓값을 찾습니다. 이 과정에서 어떤 임시 배열도 생성되지 않으며, 메모리 접근은 오직 원본 배열을 읽는 2N번의 읽기로 줄어듭니다.

이 기법의 효과는 놀라울 정도입니다. 소스 컨텍스트의 한 예시에서 이 기법을 적용했을 때의 결과는 다음과 같습니다.

reduce_by_key를 팬시 이터레이터와 함께 사용하는 것은 이제 300배 더 빠릅니다. (using fancy iterator with the reduce by key is 300x faster)

핵심 통찰: 효율적인 GPU 프로그래밍은 단순히 계산을 병렬화하는 것을 넘어, 불필요한 메모리 접근 자체를 없애는 데이터 흐름의 재설계에 달려있습니다.

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

4. '편리함'이라는 함정: 보이지 않는 데이터 이동 비용

최신 CUDA는 thrust::universal_vector와 같은 통합 메모리(Unified Memory) 기능을 제공합니다. 이는 프로그래머가 CPU 메모리와 GPU 메모리를 신경 쓰지 않고도 코드를 작성할 수 있게 해주는 매우 편리한 기능입니다. 하지만 이 편리함 뒤에는 무서운 함정이 숨어있습니다.

강의의 시뮬레이션 예제에서 기묘한 현상이 발견되었습니다. 평소 몇 밀리초밖에 걸리지 않던 시뮬레이션 단계가, 데이터를 디스크에 저장하는 작업을 수행한 직후에는 100배 이상 느려졌습니다.

원인은 통합 메모리의 '보이지 않는' 데이터 이동 때문이었습니다.

1. 디스크 저장: CPU가 파일 저장을 위해 데이터에 접근합니다. 이때 통합 메모리 시스템은 데이터가 GPU에 있다는 것을 감지하고, 자동으로 데이터를 GPU 메모리에서 CPU 메모리로 복사합니다. 이 과정은 프로그래머에게 보이지 않습니다.
2. 다음 시뮬레이션: GPU가 다음 시뮬레이션을 실행하려고 합니다. 하지만 데이터는 방금 CPU로 옮겨졌습니다. 따라서 GPU는 다시 데이터를 CPU에서 GPU 메모리로 복사해야 합니다.

바로 이 두 번째의 느린 복사 과정 때문에 성능이 급격히 저하된 것입니다. 통합 메모리는 프로그래머가 데이터 위치를 신경 쓰지 않도록 편의를 제공하지만, 그 대가로 데이터 이동의 '시점'과 '비용'에 대한 통제권을 시스템에 넘겨줍니다. 이로 인해 예측 불가능한 성능 저하가 발생할 수 있습니다.

해결책은 thrust::device_vector(GPU 전용)와 thrust::host_vector(CPU 전용)를 사용해 메모리를 명시적으로 관리하는 것입니다. thrust::copy를 통해 프로그래머가 직접 데이터 이동 시점을 제어함으로써, 이러한 예측 불가능한 성능 저하를 막고 일관된 고성능을 유지할 수 있습니다.

핵심 통찰: 최신 기술이 제공하는 '자동화'와 '편리함'이 항상 최고의 성능을 보장하지는 않습니다. 때로는 내부 동작을 이해하고 직접 제어하는 것이 중요합니다.

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

결론: 생각을 바꾸면 코드가 빨라진다

오늘 살펴본 네 가지 진실—GPU는 버스라는 점, 코드 실행은 명시적으로 지시해야 한다는 점, 최고의 최적화는 임시 데이터를 만들지 않는 것이라는 점, 그리고 편리한 통합 메모리의 숨겨진 비용—을 관통하는 하나의 아이디어가 있습니다.

바로, 효과적인 GPU 프로그래밍은 하드웨어의 힘에만 의존하는 것이 아니라, 아키텍처를 깊이 이해하고 데이터가 흐르는 방식을 근본적으로 다르게 생각하는 것에서 시작된다는 것입니다.

당신의 코드에 대해 다시 한번 생각해 보십시오. 당신의 코드에서 당연하게 여겼던 '임시 배열'이나 '자동 메모리 관리'가 사실은 성능을 발목 잡는 가장 큰 병목점은 아닐까요? 그 안에 숨어있는 100배, 혹은 300배의 성능 향상을 찾아낼 준비가 되셨습니까?