[YouTube] Implementing New Algorithm with CUDA Kernels

[원본] https://youtu.be/kTWoGCSugB4?si=bWa1GRMvITXv1VGm

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

 

직접 짜는 CUDA 커널, 생각보다 훨씬 깊고 위험하다: 꼭 알아야 할 5가지 교훈

서론: 최대 성능을 향한 유혹

CUDA 프로그래밍을 시작하면 Thrust나 CUB 같은 강력한 라이브러리들의 편리함에 금방 익숙해집니다. 대부분의 일반적인 병렬 처리 작업은 이 라이브러리들이 마법처럼 해결해 주죠. 하지만 어느 순간, 라이브러리가 제공하지 않는 매우 특수한 알고리즘을 구현해야 하는 상황과 마주하게 됩니다. 이때 "나만의 CUDA 커널을 직접 작성해야겠다"는 생각이 고개를 듭니다. 직접 짠 커널은 GPU의 모든 성능을 남김없이 끌어내는 궁극적인 방법처럼 보입니다.

하지만 이 길은 생각보다 훨씬 깊고 위험합니다. 단순히 코드를 GPU에서 실행시키는 것을 넘어, 수많은 스레드가 어떻게 상호작용하고 메모리에 접근하는지에 대한 근본적인 이해가 필요하기 때문입니다. 편리한 라이브러리의 추상화 계층 뒤에 숨겨져 있던 복잡한 현실과 마주하는 순간이죠. 이 글에서는 직접 커널을 작성할 때 반드시 알아야 할 5가지 핵심 교훈을 통해, 당신의 코드가 처참한 성능을 내거나 예상치 못한 오류를 일으키는 것을 막아줄 놀라운 사실들을 파헤쳐 보겠습니다.

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

1. '병렬'이라는 이름의 함정: 스레드 하나로는 10,000배 느려진다

CUDA 커널을 처음 작성할 때 저지르는 가장 흔하고 치명적인 실수는 바로 '병렬'이라는 단어에 속는 것입니다. GPU에서 함수를 실행하기만 하면 마법처럼 모든 것이 빨라질 것이라고 기대하지만, 현실은 정반대일 수 있습니다.

단순히 커널을 호출하는 것만으로는 충분하지 않다

처음 커널을 실행하면, 기본적으로 단 하나의 스레드가 생성되어 커널 내부의 모든 계산을 순차적으로 처리하게 됩니다. 이는 GPU의 수천 개 코어를 놀게 놔두고 단 하나의 코어만 사용하는 것과 같습니다. 이 접근 방식의 결과는 충격적입니다. 소스 컨텍스트에 따르면, 이렇게 작성된 커널은 CUB와 같은 최적화된 라이브러리를 사용했을 때보다 "대략 10,000배" 느립니다.

단순히 실행할 스레드 수를 늘리는 것만으로는 문제가 해결되지 않습니다. 각 스레드가 자신이 처리해야 할 데이터가 무엇인지 식별하고, 서로 다른 작업을 수행하도록 만들어야 합니다. threadIdx.x와 같은 내장 변수를 사용해 각 스레드에 고유한 ID를 부여하고, 이를 기반으로 작업 범위를 나누는 것이 진정한 병렬 처리의 첫걸음입니다. 예를 들어, 한 스레드가 셀 하나를 처리하고 나면, 다음 작업으로 넘어갈 때 전체 스레드 수를 자신의 현재 작업 인덱스에 더합니다. 이렇게 하면 0번 스레드(ID 0)는 셀 0, 2, 4를 차례로 처리하고, 1번 스레드(ID 1)는 셀 1, 3, 5를 처리하는 식으로 작업을 효율적으로 분배할 수 있습니다.

2. 보이지 않는 벽: 스레드는 무한하지 않으며, 경계를 넘어선다

수백만 개의 데이터를 처리하기 위해 수백만 개의 스레드를 한 번에 생성하려고 하면, CUDA는 곧바로 당신을 막아설 것입니다. 스레드를 관리하는 데에는 명확한 물리적, 논리적 한계가 존재하기 때문입니다.

스레드 블록의 한계

무작정 스레드 수를 늘리려 할 때 "invalid configure arguments"라는 암호 같은 오류를 마주하게 됩니다. 이는 CUDA의 스레드 구조를 이해하지 못했기 때문입니다. CUDA에서 스레드는 '스레드 블록(thread block)'이라는 단위로 그룹화됩니다. 그리고 하나의 스레드 블록은 최대 1,024개의 스레드만 가질 수 있습니다.

따라서 수백만 개의 스레드를 실행하려면, 여러 개의 스레드 블록으로 구성된 '그리드(grid)'를 사용해야 합니다. 처리할 문제의 크기를 블록 크기로 나누어 필요한 블록의 개수를 계산하는 것이죠. 이때, 문제 크기가 블록 크기의 배수가 아닐 경우 마지막 몇 개의 요소가 누락될 수 있습니다. 이를 방지하기 위해 CUDA는 정수 나눗셈 후 올림 처리를 해주는 cuda::seal_div 같은 헬퍼 함수를 제공하여 정확한 블록 수를 계산하도록 돕습니다.

경계 검사의 중요성

스레드 블록을 사용하더라도 또 다른 위험이 존재합니다. 바로 '경계 초과 접근(out-of-bound access)' 문제입니다. 예를 들어, 처리할 데이터가 6개인데, 편의상 4개의 스레드를 가진 블록 2개를 할당했다고 가정해 봅시다. 총 8개의 스레드가 생성되지만 실제 데이터는 6개뿐입니다. 앞의 6개 스레드는 정상적으로 작업을 수행하지만, 마지막 2개의 스레드는 존재하지 않는 메모리 위치에 접근하려다 오류를 일으키게 됩니다.

이 문제를 해결하려면 커널 코드 내에서 반드시 경계 검사(boundary check)를 추가해야 합니다.

// 각 스레드는 자신의 전역 ID를 계산한 후,
int global_thread_id = blockIdx.x * blockDim.x + threadIdx.x;

// 자신이 처리할 데이터가 실제 문제의 크기 내에 있는지 확인해야 합니다.
if (global_thread_id < problem_size) {
    // 실제 작업 수행
}

여기서 blockIdx.x는 현재 블록의 고유 ID, blockDim.x는 블록 당 스레드 수, threadIdx.x는 블록 내 스레드의 로컬 ID를 의미합니다. 이 간단한 if 문 하나가 수많은 메모리 접근 오류로부터 프로그램을 보호해 줍니다. 만약 이런 문제를 스스로 찾기 어렵다면, NVIDIA의 compute-sanitizer와 같은 도구가 어느 파일, 어느 라인에서 잘못된 메모리 접근이 발생했는지 정확히 알려주어 큰 도움이 될 수 있습니다.

3. 선량한 스레드들의 배신: 데이터 레이스와 원자적 연산의 역설

수백만 개의 스레드가 각자 독립적으로 작업을 처리할 때는 문제가 없지만, 이들이 하나의 공유된 데이터를 동시에 수정하려고 할 때 재앙이 시작됩니다. 이를 '데이터 레이스(Data Race)'라고 부릅니다.

온도 데이터로 히스토그램을 만드는 예제를 생각해 봅시다. 여러 스레드가 동시에 특정 온도가 속한 히스토그램 '빈(bin)'의 값을 1씩 증가시키려고 합니다. 이때 다음과 같은 문제가 발생합니다.

두 개의 스레드가 히스토그램의 특정 빈(bin)에 접근합니다. 현재 빈의 값은 0입니다. 첫 번째 스레드가 값 0을 읽습니다. 거의 동시에 두 번째 스레드도 값 0을 읽습니다. 첫 번째 스레드는 0에 1을 더해 1이라는 결과를 계산하고, 이 값을 다시 메모리에 씁니다. 잠시 후, 두 번째 스레드 역시 자신이 읽었던 0에 1을 더해 1이라는 결과를 계산하고, 이 값을 메모리에 덮어씁니다. 결과적으로 두 개의 스레드가 값을 증가시켰음에도 불구하고 최종 결과는 2가 아닌 1이 됩니다. 수백만 개의 스레드가 이런 작업을 반복하면, 수많은 업데이트가 누락되어 히스토그램이 거의 비어있는 것처럼 보이게 됩니다.

이 문제의 해결책은 '원자적 연산(Atomic Operation)'입니다. 원자적 연산은 '읽기-수정-쓰기' 과정을 누구도 중간에 끼어들 수 없는 단일 연산으로 묶어버립니다. CUDA에서는 cuda::std::atomic_ref를 사용해 특정 메모리 위치에 fetch_add()와 같은 원자적 연산을 수행할 수 있습니다.

하지만 여기에 역설이 존재합니다. 원자적 연산은 데이터의 정확성을 보장하는 대신, 본질적으로 연산을 '직렬화(serialization)'합니다. 수백만 개의 스레드가 하나의 메모리 위치에 접근하기 위해 길게 줄을 서서 차례를 기다리는 것과 같습니다. 이로 인해 GPU의 병렬 처리 능력이 무력화되고 심각한 성능 저하가 발생합니다. 이것이 바로 데이터 레이스를 해결한 히스토그램 커널의 성능이 처참할 정도로 나빴던 이유입니다.

4. 메모리의 비밀 계층: shared 메모리가 성능을 4배 끌어올리는 마법

원자적 연산이 초래하는 직렬화 병목을 어떻게 해결할 수 있을까요? 해답은 모든 스레드가 하나의 전역 메모리(Global Memory)에 집중하는 것을 막는 데 있습니다. '사유화(privatization)' 전략이 바로 그것입니다.

전역 메모리 대신 로컬 메모리를 활용하라

이 전략의 핵심은 각 스레드 블록이 자신만의 '작은 로컬 히스토그램'을 갖게 하는 것입니다. 스레드들은 멀리 있는 전역 히스토그램 대신, 자신들의 블록에 할당된 이 로컬 히스토그램을 먼저 업데이트합니다. 블록 내 모든 스레드가 작업을 마치면, 그 결과를 한 번에 모아 전역 히스토그램에 반영합니다. 이 접근법은 원자적 연산의 수를 극적으로 줄입니다. 초기 접근법에서는 4백만 개의 원자적 연산이 전역 메모리의 한 지점을 향해 줄을 섰지만, 사유화 전략을 통해 이 병목은 각 스레드 블록당 하나씩, 총 16,000개의 원자적 연산으로 감소합니다.

하지만 이것만으로는 부족합니다. 원자적 연산의 범위를 제한하여 성능을 한 단계 더 끌어올릴 수 있습니다. CUDA는 원자적 연산이 동기화를 보장해야 하는 스레드의 범위를 지정하는 thread_scope를 제공합니다. 각 블록의 로컬 히스토그램은 오직 해당 블록 내의 스레드들만 접근하므로, 원자적 연산의 범위를 cuda::thread_scope_block으로 제한할 수 있습니다. 이 간단한 변경만으로도 커널의 성능은 6GB/s에서 100GB/s로 급상승했습니다.

여기서 멈추지 않고, 로컬 히스토그램을 어디에 저장하는지 최적화하여 성능을 극대화할 수 있습니다. GPU의 각 SM(Streaming Multiprocessor)에는 전역 메모리보다 훨씬 빠른, L1 캐시와 유사한 속도를 내는 특별한 메모리 공간이 존재합니다. 바로 **'공유 메모리(Shared Memory)'**입니다.

__shared__ 키워드를 사용해 이 메모리에 블록별 로컬 히스토그램을 할당하기 전, 반드시 초기화 과정이 필요합니다. 블록 내의 스레드들을 사용해 로컬 히스토그램의 모든 빈(bin)을 0으로 설정한 뒤, 모든 스레드가 초기화된 상태를 볼 수 있도록 sync_threads()를 호출해야 합니다. 이 준비가 끝나면, 스레드들은 전역 메모리에 접근할 필요 없이 L1 캐시 수준의 속도로 데이터를 읽고 쓸 수 있습니다. 이 기법을 적용하자, 히스토그램 커널의 성능은 100GB/s에서 400GB/s로 무려 4배나 향상되었습니다.

공유 메모리에 쓴 내용이 다른 스레드에게 보이도록 보장하고, 모든 스레드가 특정 지점에서 작업을 마칠 때까지 기다리게 하려면 sync_threads()와 같은 동기화 장벽이 필수적입니다. 이는 스레드 블록 내에서의 질서를 유지하는 핵심적인 도구입니다.

5. 결국 다시 라이브러리로: 커널 안의 작은 조력자, CUB

직접 커널을 작성한다고 해서 모든 것을 바닥부터 재발명해야 하는 것은 아닙니다. 때로는 전문가들이 만들어 놓은 최적화된 도구를 현명하게 활용하는 것이 더 나은 결과를 가져옵니다.

CUDA는 커널 내에서 직접 호출하여 사용할 수 있는 고도로 최적화된 알고리즘 모음인 '협력 그룹(Cooperative Groups)' 라이브러리 CUB를 제공합니다. 여기서 '협력'이란, 알고리즘이 여러 스레드에 의해 동시에 호출되고 이들이 집단적으로 실행하는 모델을 의미합니다. 이는 한 스레드가 호출하고 실행하는 직렬(serial) 알고리즘이나, 한 스레드가 호출하면 백그라운드에서 다수의 스레드가 실행하는 Thrust와 같은 병렬(parallel) 알고리즘과는 구별됩니다.

앞서 우리가 공유 메모리와 동기화를 사용해 수작업으로 구현했던 블록-레벨 히스토그램 로직은, CUB의 BlockHistogram을 사용하면 훨씬 더 간단하고 효율적으로 대체할 수 있습니다. CUB 협력 알고리즘의 일반적인 사용 패턴은 다음과 같습니다.

1. 알고리즘의 템플릿 구조체를 인스턴스화합니다 (예: cub::BlockHistogram<...> ).
2. 알고리즘이 필요로 하는 임시 저장 공간을 __shared__ 메모리에 할당합니다.
3. 할당된 공유 메모리 참조를 전달하여 알고리즘 객체를 생성합니다.
4. 블록 내 모든 스레드에서 알고리즘의 멤버 함수(예: Histogram())를 호출합니다.

이는 매우 현명한 접근 방식입니다. 즉, 비즈니스의 핵심 로직처럼 특수한 부분은 직접 커널로 구현하되, 리덕션(reduction), 스캔(scan), 정렬(sort)과 같이 널리 사용되는 병렬 처리 패턴은 CUB와 같은 라이브러리에 맡기는 것입니다. 이를 통해 개발 시간을 단축하고, 전문가 수준으로 최적화된 성능을 보장받을 수 있습니다.

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

결론: 하드웨어와의 대화를 시작하며

지금까지 우리는 5가지 교훈을 통해 직접 CUDA 커널을 작성하는 것이 단순히 코드를 옮기는 작업이 아님을 확인했습니다.

1. 단일 스레드의 함정을 피하고 진정한 병렬 처리를 구현해야 합니다.
2. 스레드 블록의 한계를 이해하고 경계 검사를 통해 메모리 오류를 막아야 합니다.
3. 데이터 레이스를 인지하고 원자적 연산의 성능 역설을 극복해야 합니다.
4. 원자적 연산의 범위를 제한하고 공유 메모리를 활용해 성능을 극대화해야 합니다.
5. CUB와 같은 라이브러리를 커널 내에서 활용하여 생산성과 성능을 모두 잡아야 합니다.

결국 CUDA 커널 프로그래밍은 코드를 작성하는 행위를 넘어, 스레드 그룹이 어떻게 동작하고, 메모리 계층 구조가 어떻게 구성되어 있으며, 동기화가 왜 필요한지 등 하드웨어의 동작 원리를 깊이 이해하고 소통하는 과정입니다. 당신이 다음에 마주할 성능 최적화 문제는 어디에 숨어있을 것이라고 생각하십니까?