GPU 성능 이해하기: 루프라인 모델

기본 행렬 곱셈을 구현했으니, 이런 궁금증이 생길 수 있습니다: 우리 커널은 실제로 얼마나 잘 동작하고 있을까? GPU의 연산 능력에 의해 제한되는 걸까, 아니면 다른 무언가가 발목을 잡고 있는 걸까?

루프라인 모델(역주: 루프라인은 “상한선“이라는 뜻으로, 성능이 넘을 수 없는 한계를 지붕 선에 비유한 이름입니다)은 GPU 최적화의 나침반입니다. 커널의 성능을 제한하는 하드웨어 병목이 무엇인지 알려주고, 가장 효과적인 최적화 방향으로 안내합니다. 감으로 개선하는 대신, 루프라인 모델이 정확히 어디에 집중해야 하는지 보여줍니다.

1. 모든 GPU 커널의 두 가지 성능 상한

모든 GPU 커널은 두 가지 근본적인 제약 아래에서 동작합니다:

• 연산 상한(compute ceiling) – 코어가 부동소수점 연산을 얼마나 빠르게 수행할 수 있는가 (최대 FLOPs/s)
• 메모리 상한(memory ceiling) – 메모리 시스템이 코어에 데이터를 얼마나 빠르게 공급할 수 있는가 (최대 bytes/s)

어떤 상한이 커널을 제약하는지 파악하는 것이 최적화 전략의 핵심입니다. 루프라인 모델은 두 가지 핵심 지표를 그래프로 표현하여 이 관계를 시각화합니다:

X축: 산술 강도(Arithmetic Intensity) – 데이터 1바이트당 수행하는 연산량

\[\Large I = \frac{\text{Total FLOPs}}{\text{Total Bytes from Memory}} \quad [\text{FLOP/B}]\]

Y축: 실측 성능(Sustained Performance) – 커널이 실제로 달성하는 속도

\[\Large P_{\text{sustained}} = \frac{\text{Total FLOPs}}{\text{Elapsed Time}} \quad [\text{GFLOP/s}]\]

두 개의 “상한(roof)“이 달성 가능한 성능의 상한을 정합니다:

임계 강도(critical intensity)

\[\Large I^* = \frac{P_{\text{peak}}}{B_{\text{peak}}}\]

는 커널이 메모리 바운드(\(I < I^\ast\) )에서 연산 바운드(\(I > I^\ast\) )로 전환되는 지점입니다.

2. 하드웨어 예시: NVIDIA A100 사양

이론을 NVIDIA A100의 구체적인 숫자로 확인해 보겠습니다:

최대 FP32 처리량 \[\Large P_{\text{peak}} = 19.5 \text{ TFLOP/s} = 19{,}500 \text{ GFLOP/s}\]

최대 HBM2 대역폭 \[\Large B_{\text{peak}} = 1{,}555 \text{ GB/s}\]

임계 강도 \[\Large I^* = \frac{19{,}500}{1{,}555} \approx 12.5 \text{ FLOP/B}\]

출처: NVIDIA A100 Tensor Core GPU Architecture

이는 산술 강도가 12.5 FLOP/B 미만인 커널은 메모리 바운드, 그 이상인 커널은 연산 바운드임을 뜻합니다.

3. 행렬 곱셈 구현의 시각화

아래 애니메이션은 이 퍼즐의 구현들이 A100의 루프라인 모델에 어떻게 대응하는지 보여줍니다:

이 시각화는 이 퍼즐에서 거치게 될 최적화 과정을 보여줍니다:

1. 하드웨어 제약 – 빨간색 메모리 상한과 파란색 연산 상한이 성능 한계를 정의
2. 출발점 – 기본 구현(주황색 점)이 메모리 상한 위에 위치
3. 최적화 목표 – 공유 메모리 버전(청록색 점)으로 산술 강도가 개선됨
4. 궁극적 목표 – 금색 화살표는 커널이 연산 바운드가 되는 임계 강도 지점을 가리킴

4. 기본 구현 분석

이전 섹션의 기본 커널이 왜 이런 성능을 보이는지 살펴보겠습니다. \(2 \times 2\) 행렬 곱셈의 경우:

출력 원소당 연산량: \(\text{SIZE} + (\text{SIZE}-1) = 3 \text{ FLOPs }\)

각 원소에는 \(\text{SIZE}\) 회의 곱셈과 \(\text{SIZE} - 1\) 회의 덧셈이 필요합니다: \[C_{00} = A_{00} \cdot B_{00} + A_{01} \cdot B_{10}\] \(\text{SIZE} = 2\) 일 때 곱셈 2회 + 덧셈 1회 = 3 FLOPs

출력 원소당 메모리 접근:

• 행렬 A의 행: \(2 \times 4 = 8\) bytes (FP32)
• 행렬 B의 열: \(2 \times 4 = 8\) bytes (FP32)
• 합계: 출력 원소당 \(16\) bytes

산술 강도: \[\Large I_{\text{naive}} = \frac{3 \text{ FLOPs}}{16 \text{ bytes}} = 0.1875 \text{ FLOP/B}\]

이 산술 강도는 A100의 임계 강도에 한참 못 미치므로, 기본 커널은 심각한 메모리 바운드 상태입니다.

\[\Large I_{\text{naive}} = 0.1875 \ll I^* = 12.5\]

예상 성능: \[\Large P \approx B_{\text{peak}} \times I_{\text{naive}} = 1{,}555 \times 0.1875 \approx 292 \text{ GFLOP/s}\]

이는 GPU 연산 잠재력의 \(\frac{292}{19{,}500} \approx 1.5\%\) 에 불과합니다! 시각화에서 노란색 점이 메모리 상한 위에 놓인 것이 이를 잘 보여줍니다 — 연산 상한에는 한참 미치지 못하는 수준입니다.

5. 다음 단계: 공유 메모리 최적화

루프라인 모델이 알려주는 최적화 전략은 명확합니다: 중복 메모리 접근을 줄여 산술 강도를 높이는 것입니다. 공유 메모리 접근법이 바로 이를 달성합니다:

공유 메모리의 이점:

• 협력적 로딩: 스레드들이 함께 행렬 블록을 빠른 공유 메모리에 로드
• 데이터 재사용: 로드한 원소 하나를 여러 연산에 활용
• 전역 메모리 트래픽 감소: 느린 전역 메모리에 대한 접근 횟수 감소

산술 강도 개선 예상치: \[\Large I_{\text{shared}} = \frac{12 \text{ FLOPs}}{32 \text{ bytes}} = 0.375 \text{ FLOP/B}\]

작은 \(2 \times 2\) 규모에서는 여전히 메모리 바운드이지만, 이 2배의 산술 강도 향상은 공유 메모리 타일을 훨씬 더 많이 재사용할 수 있는 큰 행렬에서 극적인 효과를 발휘합니다.

6. 루프라인이 알려주는 최적화 전략

루프라인 모델은 현재 성능을 진단할 뿐 아니라, 최적화 방향까지 알려줍니다. 이후 퍼즐에서 살펴볼 핵심 기법은 다음과 같습니다:

각 기법은 커널을 루프라인 위에서 이동시킵니다 — 메모리 상한을 따라 위로(대역폭 활용 개선), 또는 오른쪽 연산 상한을 향해(산술 강도 향상).

비동기 연산에 대한 참고: 표준 GPU 메모리 로드(ld.global)는 이미 비동기입니다 — 워프는 로드한 데이터가 실제로 필요해질 때까지 계속 실행됩니다. cp.async(CUDA)나 copy_dram_to_sram_async(Mojo) 같은 전용 비동기 복사 명령은 여기서 한 걸음 더 나아가, 전용 복사 엔진을 사용하고 레지스터를 우회하여 자원 활용을 높입니다. 단순히 동기 연산을 비동기로 바꾸는 것과는 다릅니다.

7. 단순한 루프라인을 넘어서

다단계 메모리: 고급 루프라인은 L2 캐시, 공유 메모리, 레지스터 대역폭에 대해 별도의 상한을 포함하여 어떤 메모리 계층이 성능을 제약하는지 식별합니다.

통신 루프라인: 멀티 GPU 애플리케이션에서는 메모리 대역폭 대신 인터커넥트 대역폭(NVLink, InfiniBand)을 사용하여 스케일링 효율을 분석합니다.

전용 유닛: 최신 GPU는 고유한 성능 특성을 가진 텐서 코어를 포함하며, 별도의 루프라인 분석이 필요합니다.

8. 실전에서 루프라인 활용하기

1. 커널 프로파일링: Nsight Compute 같은 도구로 실제 FLOPs와 메모리 트래픽 측정
2. 데이터 포인트 표시: 산술 강도와 실측 성능 계산
3. 병목 식별: 메모리 바운드 커널은 메모리 상한 위에, 연산 바운드 커널은 연산 상한에 근접
4. 최적화 선택: 메모리 바운드 커널에는 대역폭 개선에, 연산 바운드 커널에는 알고리즘 변경에 집중
5. 측정과 반복: 최적화가 커널을 기대한 방향으로 이동시키는지 검증

공유 메모리 퍼즐과의 연결

다음 섹션에서는 커널을 루프라인 위로 끌어올리기 시작하는 공유 메모리 최적화를 구현합니다. 시각화에서 볼 수 있듯이, 주황색 점(기본)에서 청록색 점(공유 메모리)으로 이동하게 됩니다 — 데이터 재사용 개선을 통한 확실한 성능 향상입니다.

\(2 \times 2\) 예제에서는 연산 상한에 도달하지 못하지만, 공유 메모리가 성능에 결정적인 역할을 하는 큰 행렬에서 동일한 원리가 어떻게 확장되는지 확인할 수 있습니다. 루프라인 모델은 공유 메모리가 왜 도움이 되고 얼마나 개선을 기대할 수 있는지 이해하기 위한 이론적 토대를 제공합니다.

루프라인 모델을 이해하면 GPU 최적화가 추측에서 체계적인 엔지니어링으로 바뀝니다. 이 책의 모든 최적화 기법은 이 단순하지만 강력한 성능 모델에 대한 효과로 이해할 수 있습니다.