GPU의 연산/메모리 계층과 루프라인 모델을 바탕으로 메모리 바운드·컴퓨트 바운드 체제, 오버헤드, 연산자 퓨전과 타일링, 코얼레스드 로드와 동기화, 뱅크/워프 구조와 뱅크 충돌, 산술 집약도 향상, 점유율과 지연 은닉, 스레드 발산 회피, 양자화 등 성능 핵심 개념을 정리합니다.
마지막 업데이트: 2025-06-18
GPU가 어떻게 작동하는지 감을 잡으려고 노력해 왔다. 온라인에서 많은 자료를 읽었는데, 특히 아래 글들이 큰 도움이 되었다:
이 글은 위 자료들에서 배운 다양한 사실들을 모아 정리한다.
감사의 말: Alex McKinney에게 독립 스레드 스케줄링에 대한 코멘트에 감사한다.
목차
GPU의 설계는 주 메모리에 접근하는 속도보다 연산 속도가 훨씬 빠르기 때문에 불균형을 낳는다. 예를 들어 NVIDIA A100 GPU는 초당 19.5조회의 32비트 부동소수점 연산(TFLOPS)을 수행할 수 있지만, 메모리 대역폭은 약 1.5 테라바이트/초(TB/s)에 불과하다. 4바이트 숫자 하나를 읽는 데 걸리는 동안, GPU는 50회 이상의 계산을 수행할 수 있다.
아래는 NVIDIA A100 GPU의 연산과 메모리 계층을 도식화한 그림이다. 본문에서 인용하는 FLOPS/s와 TB/s 수치는 A100에 한정한다.
+---------------------------------------------------------------------------------+
| Global Memory (VRAM) |
| (~40 GB, ~1.5 TB/s on A100) |
+----------------------------------------+----------------------------------------+
| (Slow off-chip bus)
+----------------------------------------v----------------------------------------+
| Streaming Multiprocessor (SM) |
| (1 of 108 SMs on an A100, each ~(19.5/108) TFLOPS) |
| (2048 threads, 64 warps, 32 blocks) |
| +-----------------------------------------------------------------------------+ |
| | Shared Memory (SRAM) / L1 Cache |
| | (~192 KB on-chip workbench, 19.5 TB/s) |
| +-----------------------------------------------------------------------------+ |
| | Register File (~256 KB, ? TB/s) |
| +-----------------------------------------------------------------------------+ |
| | | |
| | //-- A "Block" of threads runs on one SM --// | |
| | +--------------------------+ +------------------------+ | |
| | | Warp 0 (32 thr) | | Warp 1 (32 thr) | ... (up to 32 warps)| |
| | | +----------------------+ | +----------------------+ | | |
| | | | Thread 0 Registers | | | Thread 32 Registers | | | |
| | | | [reg0: float] | | | [reg0: float] | | | |
| | | | [reg1: float] ... | | | [reg1: float] ... | | | |
| | | +----------------------+ | +----------------------+ | | |
| | +--------------------------+ +------------------------+ | |
| | | |
+---------------------------------------------------------------------------------+
이 다이어그램은 성능 계층을 보여준다.1 전역 메모리(VRAM)는 모든 데이터가 처음에 존재하는 크고 느린 오프칩 메모리 풀이다. 스트리밍 멀티프로세서(SM)는 GPU의 연산 단위다. 동작하려면 느린 버스를 통해 데이터를 가져와야 한다. 이를 완화하기 위해 각 SM에는 19.5 TB/s의 대역폭을 갖는 빠른 온칩 공유 메모리(SRAM)가 있다.2 프로그래머는 이를 수동으로 관리하는 캐시처럼 사용한다.
스레드는 가장 작은 실행 단위다. 각 스레드는 즉시 계산에 사용할 값을 담는 개인 레지스터 집합을 갖고 있으며, 접근 속도는 ?? TB/s 이상이다.3 하드웨어는 스레드를 32개 단위의 워프로 묶는다. 이 글은 단순화된 모델인 락스텝 실행을 사용해 성능을 분석하는데, 여기서 워프의 32개 스레드는 동일한 명령을 동시에 실행한다고 가정한다.4 A100에서 하나의 SM은 최대 64 워프를 수용할 수 있다. 프로그래머는 스레드를 블록으로 묶는데, 이는 단일 SM에서 실행됨이 보장되는 스레드 격자다. 블록은 1, 2, 3차원일 수 있다. 단순화를 위해, 본문에서는 총 스레드 수가 하드웨어 제한인 1024를 넘지 않는 정사각형 2차원 블록 BLOCK_DIM x BLOCK_DIM에 초점을 맞춘다. 블록 내 모든 스레드는 동일한 온칩 공유 메모리에 접근한다.
호스트(CPU)가 실행을 요청해 다수의 GPU 스레드가 병렬로 실행하는 함수를 커널이라고 한다. 커널의 성능은 메모리 대역폭 또는 연산 처리량 중 하나에 의해 제한된다. 이 두 한계가 성능 체제를 규정한다.
연산이 메모리 바운드라는 것은 런타임이 전역 메모리에서 SM으로 데이터를 전송하는 속도에 의해 좌우된다는 뜻이다. 예를 들어 원소별 덧셈 y = x + 1 같은 연산은 읽은 각 원소에 대해 수행하는 FLOPs가 매우 작다. SM은 대부분의 시간을 데이터 대기 때문에 유휴 상태로 보낸다.
연산이 컴퓨트 바운드라는 것은 런타임이 SM의 산술 처리 속도에 의해 좌우된다는 뜻이다. 대규모 행렬 곱셈이 대표적 예다. 일단 데이터가 SM으로 로드되면 막대한 수의 계산이 수행되고, 이 동안 메모리 버스는 유휴 상태다.
산술 집약도(Arithmetic Intensity, AI)는 체제를 결정하는 정량 지표다. 이는 작업량 대비 메모리 트래픽의 비율이다.
Arithmetic Intensity = Total FLOPs / Total Bytes Accessed
루프라인 모델에서 Total Bytes Accessed는 전역 메모리(HBM)와 온칩 SM 사이에서 전송된 데이터를 뜻한다. 이 모델은 커널 성능을 주 병목인 느린 오프칩 메모리 버스 관점에서 평가하기 때문이다. 공유 메모리에서 레지스터로의 온칩 트래픽 등은 이 계산에 포함하지 않는다.
루프라인 모델은 커널이 달성 가능한 성능(초당 FLOPs)을 AI의 함수로 그린다. 두 개의 "지붕"은 GPU의 물리적 한계다.
^ 성능 (TFLOPS)
|
| 메모리-바운드 영역 ¦ 컴퓨트-바운드 영역
| ¦
| /¦---------------------- <-- 피크 연산 (~19.5 TFLOPS)
| / ¦
| / ¦
| 피크 전역 메모리 /<--¦------ 비효율적 연산 지붕 (예: 스칼라 연산, 초월함수)
| 대역폭 (~1.5 / ¦
| TB/s) / ¦
| / ¦
+---------------------¦---------------------------> 산술 집약도 (FLOPs/Byte)
^
¦
하드웨어 리지 포인트 (~13)
커널의 성능은 다음과 같이 결정된다:
Runtime = Bytes_Accessed / Memory_Bandwidth. 따라서 성능은 Performance = Total_FLOPs / Runtime = AI * Memory_Bandwidth. 로그-로그 플롯에서 이는 대각선이다.Performance = Peak_Compute_FLOPs. 이는 수평선이다.커널의 실제 성능은 두 값의 최소값이다. 리지 포인트는 두 성능 천장이 교차하는 AI다. A100의 경우 19.5 TFLOPS / 1.5 TB/s ≈ 13 FLOPs/Byte. 커널이 컴퓨트 바운드가 되려면 이 AI를 넘어야 한다. AI가 13보다 낮으면 메모리-바운드 영역, 높으면 컴퓨트-바운드 영역에서 동작한다. 최적화의 목표는 AI를 증가시켜 커널의 작동점을 오른쪽으로 이동시키고, 성능을 끌어올려 연산 지붕에 닿게 하는 것이다.
"피크 연산"인 19.5 TFLOPS는 텐서 코어 행렬 곱과 같은 고도로 최적화된 명령과 충분한 전력 제한에서만 달성 가능한 이상적 값이다. 연산이 컴퓨트 바운드여도 이 피크에 한참 못 미칠 수 있다. 예컨대 AI가 높지만 스칼라 산술이나 sin, exp 같은 복잡한 초월함수가 지배적인 커널은 해당 느린 명령의 처리량에 의해 제한된다. 이 경우 도표처럼 더 낮은 효과적 "지붕"이 생긴다. AI를 높이는 것만으로는 충분치 않으며, 수행하는 FLOPs 자체가 효율적이어야 한다.
AI를 높이는 기본 전략은 일단 SM의 빠른 온칩 메모리에 데이터를 로드한 뒤 그 데이터를 최대한 재사용하는 것이다. 아래는 스레드가 전역 메모리에서 자신의 개인 레지스터로 직접 데이터를 읽는 단순화된 모델이다. 이 분석은 최소 필요 데이터 전송량을 계산한다. 실제 메모리 트래픽은 접근 패턴에 따라 달라지며, 이는 뒤에서 논의한다.
C = A@B를 계산한다고 하자. 모든 행렬은 N x N이고 4바이트 실수를 사용한다.
전략 1: 한 스레드가 C의 한 원소 C[i,j]를 계산
C[i,j]를 계산하려면 스레드는 N개의 곱셈-덧셈을 수행한다. 즉 2*N FLOPs.2*N개의 실수, 8*N 바이트.(2*N FLOPs) / (8*N Bytes) = 0.25 FLOPs/Byte.이 AI는 낮다. 커널은 메모리 바운드다.
전략 2: 한 스레드가 C의 2x2 타일을 계산
2x2 타일(C[i,j], C[i,j+1], C[i+1,j], C[i+1,j+1])을 계산한다고 하자.
4개 원소 × 2*N FLOPs/원소 = 8*N FLOPs.A[i,:], A[i+1,:])과 B의 두 열(B[:,j], B[:,j+1])을 읽는다. 총 4*N개의 실수, 16*N 바이트.(8*N FLOPs) / (16*N Bytes) = 0.5 FLOPs/Byte.이 AI 값들은 A100의 리지 포인트 약 13 FLOPs/Byte에 한참 못 미친다. 이 단순 레지스터 전용 모델만으로는 행렬 곱을 컴퓨트 바운드로 만들 수 없다.5 높은 AI를 달성하는 핵심은 블록 내 스레드들이 협력해 A와 B의 훨씬 큰 타일을 공유 메모리(SRAM)에 로드하는 것이다. 이 공유 데이터를 함께 사용하면, 1024개의 스레드로 구성된 블록이 AI 13을 넘길 수 있다. 자세한 메커니즘은 공유 메모리 섹션에서 다룬다.
성능은 호스트 측 오버헤드에 의해 제한될 수도 있다. 이는 GPU에 작업을 준비시키는 데 CPU(호스트)가 소비하는 시간으로, 예컨대 파이썬 인터프리터나 프레임워크 디스패치 과정의 시간이다.
애플리케이션의 GPU 커널이 너무 작거나 너무 많으면 오버헤드 바운드가 된다. GPU는 각 작은 작업을 빨리 끝내고 다음 명령을 내릴 CPU를 기다리며 유휴 상태가 된다. 런타임은 GPU에 일을 충분히 빠르게 공급하지 못하는 CPU 때문에 지배된다.
현대 프레임워크는 비동기 실행으로 이를 완화한다. 호스트는 각 명령이 끝나기를 기다리지 않고도 GPU에 명령 스트림을 큐잉할 수 있다. 개별 GPU 연산이 충분히 크다면, 호스트는 "앞서 달릴" 수 있고, 한 커널의 런치 오버헤드는 이전 커널의 실행 동안 가려진다.
이후 내용에서는 커널이 충분히 커서 오버헤드가 주된 제한 요인이 아니라고 가정하고, 메모리와 연산에 집중한다.6
커널이 충분히 커서 런치 오버헤드가 무시할 수 있을 때, 성능은 GPU의 두 물리적 한계(메모리 대역폭과 연산 처리량)에 의해 좌우된다. 따라서 커널의 성능을 높이는 것은 루프라인 모델에서 작동점을 위쪽과 오른쪽으로 밀어 올리는 일이다. 이를 위한 기본 전략은 두 가지다.
각 전략을 차례로 다룬다.
y = relu(x + 1) 같은 단순 연산의 체인이 흔하다. 각 연산(add, relu)은 산술 집약도가 매우 낮아 메모리 바운드다. 이를 별도의 GPU 커널로 순차 실행하는 것은 비효율적이다. 이 시퀀스를 최적화하는 기본 전략이 연산자 퓨전이다.
문제는 중간 메모리 트래픽이다. 퓨전하지 않은 y = relu(x + 1) 실행을 생각해 보자:
add): 전역 메모리에서 텐서 x 전체를 읽는다. tmp = x + 1을 계산한다. 중간 텐서 tmp 전체를 전역 메모리에 쓴다.relu): 전역 메모리에서 텐서 tmp 전체를 읽는다. y = relu(tmp)를 계산한다. 최종 텐서 y를 전역 메모리에 쓴다.이 접근은 낭비가 심하다. 두 번의 커널 런치 오버헤드를 발생시키고, 중간 tmp 텐서 때문에 느린 전역 메모리 왕복을 강제한다.
퓨전은 이 단계들을 하나의 더 효율적인 GPU 커널로 결합한다. Triton이나 torch.compile의 Inductor 백엔드 같은 JIT 컴파일러가 이 변환을 자동으로 수행할 수 있다.
퓨전된 커널에서는:
x의 한 원소를 자신의 레지스터로 읽는다.tmp = x + 1, 이어서 y = relu(tmp) 등 모든 계산을 그 빠른 레지스터 안에서 수행한다.y만 전역 메모리에 쓴다.# 비퓨전(개념)
def unfused_add_relu(x):
tmp = torch.add(x, 1) # HBM에서 x를 읽고, tmp를 HBM에 씀
y = torch.relu(tmp) # HBM에서 tmp를 읽고, y를 HBM에 씀
return y
# 퓨전(개념)
@torch.compile
def fused_add_relu(x):
# 컴파일러가 둘을 하나의 커널로 퓨전함.
# x+1의 중간 결과는 결코 HBM에 닿지 않음.
return torch.relu(x + 1)
중간 텐서 tmp는 일시적이 되어 전역 메모리에 물리화되지 않는다. 이는 메모리 트래픽을 절반으로 줄이고(한 번의 x 읽기, 한 번의 y 쓰기), 두 번째 커널의 런치 오버헤드를 제거한다.
C=A@B의 레지스터 전용 모델은 산술 집약도 0.25 FLOPs/Byte를 주었는데, 이는 A100의 리지 포인트 ~13에 한참 못 미친다. 이는 단일 스레드가 2*N 값을 읽어 2*N FLOPs를 수행하고 데이터를 한 번 쓰고 버리기 때문이다. 데이터 재사용을 늘려 컴퓨트 바운드가 되려면, 블록 내 스레드들이 협력해 입력 행렬의 큰 타일을 SM의 빠른 온칩 공유 메모리에 로드해야 한다.
이 협력의 논리는 행렬 곱의 분해에 기초한다. 한 원소 C[i,j]의 계산은 k 차원에 대한 합으로 C[i,j] = sum_k A[i,k] B[k,j]이다. 이 합은 타일에 대한 부분합들의 합으로 나눌 수 있다. 정사각형 타일의 경우, 내부 k 차원을 외부 차원과 일치하는 BLOCK_DIM 크기의 타일로 쪼갠다. 수식은 다음과 같다:
타일링 알고리즘은 외부 합의 항(하나의 부분 곱셈)을 반복마다 하나씩 계산한다. 한 블록의 스레드가 하나의 출력 C_tile을 계산하기 위해 k 차원을 따라 A와 B의 타일을 로드하고, 온칩에서 곱을 계산해 누적한다. 이는 로드, 동기화, 연산의 3단계 패턴으로 달성된다.
# 하나의 스레드 블록이 하나의 출력 타일 C_tile을 계산하는 개념적 알고리즘.
# C_tile은 예: C[block_row_start:end, block_col_start:end]에 해당.
# 블록의 각 스레드는 자신의 레지스터에 C_tile의 일부를 보유. 0으로 초기화.
thread_private_C_accumulator = zeros(...)
# k-차원을 따라 A와 B의 타일을 순회.
# 각 반복은 위 합에서 하나의 부분 곱을 계산.
for k_tile_idx in range(NUM_K_TILES):
# 단계 1: 로드
# 블록 내 모든 스레드가 협력해 느린 전역 메모리에서 빠른 공유 메모리로
# A의 타일 하나와 B의 타일 하나를 로드.
A_tile = load_A_tile_from_global_mem(k_tile_idx)
B_tile = load_B_tile_from_global_mem(k_tile_idx)
# 단계 2: 동기화
# 어떤 스레드도 계산을 시작하기 전에 모든 스레드가 로드를 마칠 때까지 대기.
# A_tile과 B_tile이 완전히 채워졌음을 보장.
__syncthreads()
# 단계 3: 연산
# 각 스레드는 온칩 matmul의 자신의 부분을 계산.
# A_tile과 B_tile의 데이터는 공유 메모리에서 광범위하게 재사용됨.
thread_private_C_accumulator += on_chip_matmul_piece(A_tile, B_tile)
# 다음 타일을 로드하기 전에 모든 스레드가 계산을 마칠 때까지 대기.
__syncthreads()
# 루프가 끝나면, 최종 누적 결과를 전역 메모리에 기록.
write_C_tile_to_global_mem(thread_private_C_accumulator)
이제 3단계 로드-동기화-연산 패턴의 메커니즘을 살펴본다.
첫 단계는 느린 전역 메모리(HBM)에서 빠른 온칩 공유 메모리(SRAM)로 A와 B의 타일을 로드한다. 목표는 가능한 최대 메모리 대역폭으로 이 전송을 수행하는 것이다. 이를 위해 코얼레스드(coalesced) 메모리 접근이 필요하다. 워프의 32개 스레드가 하나의 연속된 128바이트 HBM 블록을 한 번의 트랜잭션으로 접근할 때, 그 접근은 코얼레스드라고 한다.
이를 달성하려면 커널이 스레드 인덱스를 메모리 주소에 매핑해야 한다. BLOCK_DIM x BLOCK_DIM 스레드 블록이 동일 크기의 데이터 타일을 로드할 때 흔한 매핑은 스레드 (tx, ty)가 전역 메모리의 A[global_row + ty, global_k + tx]를 공유 메모리 A_tile[ty, tx]에 쓰도록 하는 것이다. 여기서 BLOCK_DIM은 32다.
ty는 고정이고 tx가 0에서 31까지인 단일 워프를 생각하자.
(0, ty)는 A[global_row + ty, global_k + 0]을 읽는다.(1, ty)는 A[global_row + ty, global_k + 1]을 읽는다.(31, ty)는 A[global_row + ty, global_k + 31]을 읽는다.행 우선(row-major) 저장을 가정하면, 이 스레드들은 32개의 연속된 4바이트 실수, 즉 연속된 128바이트 구간을 접근한다. 이는 완벽한 코얼레스드 읽기다. 전체 32x32 타일은 블록의 각 워프마다 한 번씩, 총 32번의 코얼레스드 읽기로 로드된다.
Thread Block (32x32) Global Memory (HBM)
(One row of A's tile)
+--------------------+
| Warp 0 (ty=0) | ----> [A_ij, A_i,j+1, ..., A_i,j+31] (128 bytes)
| (tx = 0..31) | (One coalesced memory transaction)
+--------------------+
| Warp 1 (ty=1) | ----> [A_i+1,j, ..., A_i+1,j+31] (128 bytes)
+--------------------+
| ... |
+--------------------+
| Warp 31 (ty=31) | ----> [A_i+31,j, ..., A_i+31,j+31] (128 bytes)
+--------------------+
이 로드는 벡터화된 접근으로 더 효율적으로 만들 수 있다. 코얼레스드 읽기의 물리적 메모리 트랜잭션은 어차피 HBM에서 전체 128바이트를 가져온다. 차이는 SM이 이 데이터를 요청하는 방식이다.
스칼라 로드의 경우, 워프는 32개의 32비트 로드 명령을 발행해야 한다. 벡터화된 로드의 경우, 8개의 더 넓은 128비트 로드 명령만 발행한다. 이는 SM이 클록 사이클당 발행할 수 있는 명령 슬롯 수가 제한적이기 때문이다. 32개의 좁은 명령보다 8개의 넓은 명령으로 데이터를 요청하는 것이 하드웨어 자원을 덜 소모한다. 이로써 메모리 컨트롤러가 폭넓은 요청의 연속 흐름으로 지속적으로 바빠지게 되어, SM 측 병목을 줄여 활용 메모리 대역폭을 높인다.
벡터화된 접근은 디바이스 코드에서 포인터 캐스팅(예: float*에서 float4*로)을 통해 활성화되며, 이는 메모리가 벡터 크기에 정렬(aligned)되어 있음을 컴파일러에 보장하는 행위다.
이 벡터화된 로드의 효율성은 메모리 정렬에 의존한다. 단일 float4 명령은 16바이트 벡터를 로드한다. 4바이트 실수 행렬의 경우 이 벡터에는 원소 4개가 담긴다. 하드웨어는 메모리 주소가 16의 배수일 때만 이 명령을 효율적으로 실행한다. 이는 행렬의 내부 차원 K(열의 수)가 4의 배수여야 함을 의미한다. K가 4의 배수가 아니면, 행의 경계가 16바이트 세그먼트와 어긋난다.
4바이트 실수와 16바이트 세그먼리를 사용하는 메모리 시스템을 생각해 보자.
Memory: |<--- 16B --->|<--- 16B --->|
[Seg 0 ][Seg 1 ]
Row 0: [e0 e1 e2 e3 | e4 e5 e6 e7] (A float4 load for e0-e3 is aligned)
Row 1: [e0 e1 e2 e3 | e4 e5 e6 e7] (A float4 load for e0-e3 is aligned)
Memory: |<--- 16B --->|<--- 16B --->|<--- 16B --->|
[Seg 0 ][Seg 1 ][Seg 2 ]
Row 0: [e0 e1 e2 e3 e4 e5 e6]
Row 1: [e0 e1 e2 e3 e4 e5 e6] (A float4 load for Row 1's e0-e3 spans Seg 0 and Seg 1)
이러한 미정렬은 하드웨어가 더 복잡하고 느린 로드 연산을 발행하도록 강제하여 메모리 대역폭을 떨어뜨린다.7
중요: 이 행 단위 전략은 행렬 A에는 코얼레스드 접근을 제공하지만, 행렬 B에는 요구되는 접근 패턴이 상충한다.
행 우선 행렬에서 열을 직접 로드하는 것은 비코얼레스드, 즉 보폭(stride)이 큰 접근이며 HBM 트랜잭션을 직렬화한다. 따라서 해법은 B 타일을 HBM에서는 행 단위 코얼레스드로 읽고, 공유 메모리에 쓸 때 재배열하는 것이다. 이 재배열의 구조는 공유 메모리의 물리적 뱅크 구조에 의해 결정된다.
__syncthreads() 호출은 배리어로 작동한다. 블록 내 어떤 스레드도 이 지점에 도달할 때까지 다른 스레드가 진행하지 않는다. 이는 연산 단계가 시작되기 전에 A_tile과 B_tile이 공유 메모리에 완전히 로드되어 있음을 보장한다.8
공유 메모리는 스트리밍 멀티프로세서(SM)에 위치한 물리 자원이다. 스레드 블록이 SM에서 실행되도록 스케줄되면, 그 블록은 해당 SM의 공유 메모리 일부를 독점적으로 할당받는다.
공유 메모리는 동일 크기의 32개 독립 메모리 모듈, 즉 뱅크로 물리적으로 분할되어 있다. 이 뱅크들은 병렬로 메모리 요청을 처리할 수 있다. 이 숫자는 임의가 아니며 워프 크기와 맞춰져 있다. 워프는 락스텝으로 명령을 실행하는 32개 스레드로 구성되며, 메모리 접근의 기본 단위다. 32개 뱅크는 단일 워프로부터의 32개 메모리 요청을 한 클록 사이클에 병렬로 처리하도록 설계되었는데, 이는 요청들이 서로 다른 뱅크를 겨냥할 때에 한한다.
주소는 4바이트 워드 단위로 뱅크에 교차 배치(interleaved)된다.
bank 0: [word 0, word 32, word 64, ...]
bank 1: [word 1, word 33, word 65, ...]
...
bank 31: [word 31, word 63, word 95, ...]
주어진 워드 주소의 뱅크는 bank_id = address % 32로 결정된다.
공유 메모리의 전체 대역폭을 달성하려면, 워프의 32개 스레드가 32개의 서로 다른 뱅크에 해당하는 워드를 접근해야 한다. 서로 다른 스레드가 같은 뱅크에 매핑되는 서로 다른 주소를 접근하면 뱅크 충돌이 발생한다. 하드웨어는 이러한 요청을 직렬화해 처리하며, 대역폭이 떨어진다. 모든 스레드가 같은 주소를 읽는 브로드캐스트는 빠르고 충돌이 없는 연산이다.
이 점이 행렬 곱에서 문제를 만든다. 공유 메모리에 행 우선 레이아웃으로 저장된 BLOCK_DIM x BLOCK_DIM 타일을 생각하자(BLOCK_DIM=32). tile[row, col]의 주소는 row * 32 + col이다.
A_tile[fixed_row, t]를 t = 0..31에 대해 접근한다. 각 주소는 fixed_row * 32 + t이고, 각 스레드 t의 뱅크는 (fixed_row * 32 + t) % 32 = t % 32다. t가 스레드마다 유일하므로, 32개의 유일한 뱅크를 접근한다. 충돌 없는 최대 대역폭 접근이다.B_tile[t, fixed_col]을 t = 0..31에 대해 접근한다. 각 주소는 t * 32 + fixed_col이고, 각 스레드 t의 뱅크는 (t * 32 + fixed_col) % 32 = fixed_col % 32다. 32개 스레드가 모두 동일한 뱅크를 겨냥한다. 이는 32방향 뱅크 충돌을 일으켜 접근을 직렬화한다.해법은 공유 메모리 내에서 B_tile을 전치된 레이아웃으로 저장하는 것이다.
# 로드 단계에서 스레드 (tx, ty)가 수행하는 동작
# A는 그대로 로드, B는 로드시 즉시 전치하여 저장
A_tile[ty, tx] = A_global[global_row + ty, global_k + tx]
B_tile[tx, ty] = B_global[global_k + ty, global_j + tx] # 인덱스를 바꿔 씀
이 "로드-전치" 기법은 온칩 계산을 바꾼다. 부분 곱의 한 원소 계산은 더 이상 A_tile의 한 행과 B_tile의 한 열 간의 내적이 아니다. 대신 온칩 B_tile이 전치되어 있으므로 수식은 다음과 같이 된다:
이 공식에서는, 고정된 i에 대해 서로 다른 j 값을 계산하는 워프가 A_tile의 한 행과 온칩 B_tile의 한 행을 접근한다. 둘 모두 충돌 없는 접근 패턴이다. 이 단일 전략으로 HBM 코얼레스드 요구와 SRAM 뱅크 충돌 문제를 동시에 해결한다.
로드-전치 동작 (스레드 tx, ty)
HBM에서는 행 단위로 읽고, SRAM에는 열 단위로 씀
Global Memory (HBM) Shared Memory (SRAM)
+-------------------------+ +-----------------------+
| B[k_base+ty, j_base+tx] | -----> | B_tile[tx, ty] |
+-------------------------+ +-----------------------+
결과: HBM 읽기는 코얼레스드, SRAM 읽기는 충돌 없음.
데이터가 공유 메모리에 대기하면, 블록은 연산을 수행한다. 목표는 이 빠른 온칩 메모리에서 데이터 재사용을 극대화하는 것이다. 온칩 연산을 구성하는 두 가지 전략을 분석한다.
전략 1: 한 스레드가 한 출력 계산
가장 단순한 접근은 한 출력 원소를 한 스레드에 매핑하는 것이다. BLOCK_DIM x BLOCK_DIM 스레드 블록이 TILE_DIM x TILE_DIM 데이터 타일을 계산하며, 여기서 BLOCK_DIM과 TILE_DIM은 같다. 이 전략은 공유 메모리 캐싱을 도입하는 Boehm의 글의 "커널 3"과 개념적으로 유사하다.9 블록당 1024 스레드라는 하드웨어 한계로 인해 BLOCK_DIM은 최대 32로 제한된다. 스레드 (tx, ty)는 단일 출력 원소 C_partial[ty, tx]를 담당한다.
# BLOCK_DIM = TILE_DIM인 단일 스레드 (tx, ty)의 동작
c_accumulator = 0.0
for k in range(TILE_DIM):
c_accumulator += A_tile[ty, k] * B_tile[tx, k]
이 전략의 산술 집약도는 TILE_DIM / 4다.
2 * TILE_DIM^3 FLOPs를 수행한다.8 * TILE_DIM^2 바이트다.(2 * TILE_DIM^3) / (8 * TILE_DIM^2) = TILE_DIM / 4 FLOPs/Byte.TILE_DIM이 32로 제한되므로 최대 AI는 32 / 4 = 8이다. 이는 A100의 리지 포인트 ~13을 넘지 못한다. 커널은 여전히 메모리 바운드다.
전략 2: 한 스레드가 여러 출력을 계산
AI를 높이려면 스레드 수를 늘리지 않고 TILE_DIM을 늘려야 한다. 이를 위해 데이터 타일 크기와 스레드 블록 크기를 분리한다. 각 스레드에 더 많은 일을 할당한다. 이 전략은 Boehm의 글의 "커널 5"가 겨냥하는 바와 같다.
16x16 스레드 블록(BLOCK_DIM = 16, 256 스레드)이 64x64 데이터 타일(TILE_DIM = 64)을 계산할 수 있다. 이제 각 스레드는 출력의 4x4 서브타일을 계산한다. 단, TILE_DIM=64가 공유 메모리 용량을 넘지 않아야 한다.10
# 스레드가 4x4 출력 서브타일을 계산
# TILE_DIM = 64, BLOCK_DIM = 16
c_regs = [[0.0] * 4 for _ in range(4)]
a_regs = [0.0] * 4
b_regs = [0.0] * 4
for k in range(TILE_DIM):
# A_tile과 B_tile의 얇은 슬라이스를 레지스터로 로드
for i in range(4): a_regs[i] = A_tile[thread_row*4 + i, k]
for j in range(4): b_regs[j] = B_tile[thread_col*4 + j, k]
# 레지스터에서 외적을 계산하여 c_regs에 누적
for i in range(4):
for j in range(4):
c_regs[i][j] += a_regs[i] * b_regs[j]
AI 계산은 여전히 TILE_DIM / 4다. TILE_DIM = 64이면 AI는 64 / 4 = 16 FLOPs/Byte로, A100의 리지 포인트를 넘는다. 커널은 이제 컴퓨트 바운드다.
컴퓨트 바운드 커널의 런타임은 SM의 산술 처리량으로 제한된다. 이는 높은 절대 성능을 보장하지는 않는다. 예컨대 스페셜라이즈드 하드웨어(텐서 코어11) 대신 스칼라 FP32 연산을 사용하거나, 전력 제한 때문에 GPU가 피크 클록 이하로 동작하면 컴퓨트 바운드여도 느릴 수 있다.
위 코드의 내부 루프는 더 최적화할 수 있다. 스레드는 A_tile에서 네 개의 개별 float 값을 a_regs에 로드한다. 대신 단일 명령으로 16바이트 float4 벡터를 로드할 수 있다. 공유 메모리에서의 이 벡터화된 로드는 온칩 데이터 이동에 필요한 명령 수를 줄여 연산 단계의 효율을 높인다. 이는 Boehm의 글의 "커널 6"에서 사용된 온칩 벡터화 정련과 상응한다.
마지막 고려: 타일 양자화(quantization)
행렬 차원이 타일 크기의 배수가 아니면, 커널은 쓸모없는 계산을 수행하는 여분의 블록을 런치한다.
M x N 행렬을 TILE_M x TILE_N 타일로 덮기 위해 GPU는 ceil(M/TILE_M) x ceil(N/TILE_N) 스레드 블록 격자를 런치한다. 65x65 행렬을 32x32 타일로 타일링하려면 ceil(65/32) x ceil(65/32) = 3x3 블록 격자가 필요하다. 커널의 로직은 고정되어 있으며, 각 블록은 전체 32x32 타일의 산술을 수행하도록 프로그램되어 있다.
Columns 0-31 Columns 32-63 Columns 64-95
+-----------------+-----------------+-----------------+
R 0 | | | |
o-31| Block 0,0 | Block 0,1 | Block 0,2 |
w | (Full work) | (Full work) | (Wasted work) |
s | | | |
+-----------------+-----------------+-----------------+
R 32| | | |
o-63| Block 1,0 | Block 1,1 | Block 1,2 |
w | (Full work) | (Full work) | (Wasted work) |
s | | | |
+-----------------+-----------------+-----------------+
R 64| | | |
o-95| Block 2,0 | Block 2,1 | Block 2,2 |
w | (Wasted work) | (Wasted work) | (Wasted work) |
s | | | |
+-----------------+-----------------+-----------------+
NVIDIA에 따르면, “라이브러리는 어떤 타일도 유효하지 않은 메모리 접근을 수행하지 않도록 보장하지만, 모든 타일은 동일한 양의 수학 연산을 수행한다.” 이렇게 되는 이유에 대한 나의 이해(틀렸다면 알려주기 바람): 경계 블록들이 낭비 계산을 수행하는 것은 커널이 데이터를 명시적으로 패딩하기 때문이다. 행렬 경계 바깥 원소를 로드하도록 할당된 스레드는 가드 조건으로 실제 로드를 하지 못하게 막는다. 대신, 온칩 공유 메모리 타일에서 자신의 위치에 0을 쓴다. 산술 루프는 줄어들지 않는다. 커널의 로직은 타일 전반에 균일하다. 워프의 모든 스레드는 동일한 곱셈-덧셈 명령을 실행한다. 패딩된 0에 해당하는 데이터를 가진 스레드는 명령을 실행하되, C += A * 0처럼 무의미한 계산을 수행할 뿐이다. 하드웨어 자원은 사용되지만, 작업은 폐기된다.
커널을 컴퓨트 바운드로 만들었다. 이제 성능은 온칩 산술 속도로 제한된다. 하지만 하드웨어의 다른 측면을 관리해 더 빠르게 만들 수 있다. 다음은 그중 세 가지다. 다른 것들도 있지만, 아직 충분히 숙련되지 못해 다루지 않는다. 자세한 내용은 Boehm의 글을 보라.
워프는 전역 메모리 읽기 같은 장지연 명령을 실행할 때 정지(stall)한다. 데이터가 도착할 때까지 다음 명령을 실행할 수 없으며, 이는 수백 클록 사이클이 걸릴 수 있다. 이 동안 정지한 워프만 있다면 SM의 연산 유닛은 유휴 상태가 된다.
SM은 다른 작업을 실행해 이 지연을 숨긴다. 여러 스레드 블록을 동시에 수용하여 상주 워프 풀을 만든다. 한 워프가 정지하면 SM의 하드웨어 스케줄러는 즉시 이 풀에서 실행 준비가 된 다른 워프로 전환한다. 이를 지연 은닉(latency hiding)이라고 한다.
+-------------------------------------------------------------------+
| Streaming Multiprocessor (SM) |
| |
| [Block A] [Block B] |
| - Warp A1 (Ready) - Warp B1 (Ready) |
| - Warp A2 (Stalled -> waiting on HBM) |
| | | |
| +------------------v------------------+ |
| [ Pool of Ready-to-Run Warps ] |
| [ A1, B1 ] |
| | |
| +-------v-------+ |
| | SM Scheduler | --> [Execute instructions] |
| +---------------+ |
| |
+-------------------------------------------------------------------+
점유율(Occupancy)은 SM에서 활성 워프 수를 SM이 수용 가능한 최대 워프 수로 나눈 비율이다. 높은 점유율은 스케줄러에 더 큰 워프 풀을 제공해, 매 사이클 실행 가능한 워프를 찾을 가능성을 높여 연산 유닛을 바쁘게 유지한다.
이는 블록당 사용하는 자원과 SM에 동시에 상주시킬 수 있는 블록 수 사이의 트레이드오프로 이어진다. 두 극단을 다음과 같이 시각화할 수 있다:
+------------------------------------+ +----------------------------------------------+
| SM with High AI, Low Occupancy | | SM with Low AI, High Occupancy |
| | | |
| +--------------------------------+ | | +----------+ +-----------+ +-----------+ |
| | Block 0 (uses 64KB SMEM) | | | | Block 0 | | Block 1 | ... | Block N | |
| | TILE_DIM=128 -> High AI | | | | (8KB SMEM) | (8KB SMEM)| | (8KB SMEM)| |
| +--------------------------------+ | | +----------+ +-----------+ +-----------+ |
| | | |
| --> Low # of resident blocks. | | --> High # of resident blocks. |
| --> Small pool of warps for | | --> Large pool of warps for |
| latency hiding. | | latency hiding. |
+------------------------------------+ +----------------------------------------------+
커널의 자원 사용을 조율하여 높은 AI의 이점과 충분한 점유율의 필요 사이 균형을 잡는다. 이를 위한 주요 레버는 스레드 블록 차원(BLOCK_DIM), 블록당 할당하는 공유 메모리 양(즉 TILE_DIM에 의해 결정), 스레드당 사용하는 레지스터 수다.12
워프 내 스레드들이 분기(if-else)의 결과에 합의하지 못하면 스레드 발산(thread divergence)이 발생한다.13 이 경우 하드웨어는 서로 다른 코드 경로를 직렬로 실행한다. 먼저 if 경로를 택한 스레드들이 그 코드를 실행하는 동안 나머지는 비활성화되고, 다음 사이클에는 반대로 동작한다.
# 32개 스레드의 워프가 if 문을 만났다고 하자:
if (thread_id < 16)
# Path A
else
# Path B
Execution Timeline:
Time ->
+------------------------------------------------------------------+
| Warp Execution |
| |
| Cycle 1: Path A is executed. |
| - Threads 0-15: Active, execute Path A code. |
| - Threads 16-31: Inactive, masked off. |
| |
| Cycle 2: Path B is executed. |
| - Threads 0-15: Inactive, masked off. |
| - Threads 16-31: Active, execute Path B code. |
| |
| Result: Two cycles are required instead of one. |
| Effective throughput is halved. |
+------------------------------------------------------------------+
이러한 직렬화는 발산 코드의 실행 시간을 두 배로 늘려 워프의 유효 처리량을 절반으로 만든다. 성능 크리티컬 구간에서는 if-else 대신 min, max 같은 프리미티브를 사용해 분기 없는 코드를 작성해 이 비용을 피한다.
양자화는 텐서 원소의 정밀도를 FP32에서 FP16이나 BF16 등으로 낮춘다. 이는 두 가지 효과가 있다. 첫째, 원소 하나를 저장하는 데 필요한 메모리를, 예컨대 2배 줄인다. 따라서 전역 메모리에서 공유 메모리로 초당 옮길 수 있는 원소 수가 두 배가 된다. 이는 AI를 2배 높인다.
둘째, A100 같은 GPU는 저정밀 원소에서 더 빠르게 동작할 수 있다. 예를 들어 A100에서는 특정 FP16 연산에서 312 TFLOPS가 가능하지만, FP32 연산은 19.5 TFLOPS로 제한된다. 따라서 이론상 계산 속도를 최대 16배까지 올릴 수 있다.
양자화는 따라서 루프라인 플롯에서 작동점을 위쪽과 오른쪽으로 이동시킬 수 있다.
이 글에서 배운 바에 따르면, 전력 제한이 클록 속도에 영향을 주므로 이런 “피크” 수치는 감안하고 봐야 한다.↩
공유 메모리의 (피크) 19.5 TB/s 수치는 다음과 같이 유도된다: . 읽기와 쓰기가 독립 포트를 사용하므로 어떤 의미에서는 대역폭이 두 배가 된다는 점에 주의.↩
이 수치는 NVIDIA가 공개하지 않는다고 한다. 어차피 핵심 요지는 이것이 가장 빠르게 읽고 쓸 수 있는 메모리라는 점이므로 자세히 추적하지 않았다.↩
최신 GPU(예: A100)는 실제로 독립 스레드 스케줄링(ITS)을 지원한다. ITS에 대한 좋은 소개는 이 문서의 14페이지 이후를 보라. 특히 저자들은 ITS에서 “각 스레드는 고유한 프로그램 카운터를 갖기 때문에, 이론적으로 각 스레드는 다음에 수행하려는 고유한 명령을 저장할 수 있다. 스레드의 실행은 여전히 워프로 일어나며, 같은 사이클에 서로 다른 명령을 수행할 수는 없다. 그러나 워프는 이제 그 안의 스레드들이 현재 보유한 서로 다른 프로그램 카운터 중 어느 것이든에서 진행하도록 스케줄될 수 있다. 더 나아가 ITS는 ‘진행 보장’을 제공한다: 여러 사이클에 걸쳐 워프 내 스레드들이 유지하는 모든 개별 프로그램 카운터가 결국 방문된다. 이는 가령 실행이 분기되어 두 갈래가 생긴 경우, 둘 다 sooner or later 실행될 것이 보장됨을 뜻한다.”고 쓴다. ITS는 오래된 아키텍처에 비해 분기 코드의 정확성을 조금 더 쉽게 보장하지만, 여전히 워프가 가능한 한 많이 락스텝으로 동작하도록 코드를 작성해 최대 병렬 레인을 활용하는 것이 바람직하다.↩
또한 단일 스레드의 레지스터 집합에 행렬의 전체 열을 로드할 수 있다고 가정했다. 각 SM의 레지스터 파일이 256 KB이므로 큰 행렬에서는 불가능하다.↩
오버헤드에 대한 더 많은 내용은 Making Deep Learning Go Brrrr From First Principles을 보라.↩
한편 Horace He가 지적했듯이), 이 효과는 모델의 어휘 크기를 패딩하면 성능이 향상되는 이유를 설명한다.↩
독립 스레드 스케줄링에서는 워프 내 암묵적 동기화에 의존할 수 없다. 예컨대 어떤 스레드도 다른 스레드가 공유 메모리에 데이터를 쓰기 전에 이를 읽지 않도록 보장하려면 __syncthreads() 같은 명시적 배리어가 필요하다.↩
이 글은 HBM 코얼레싱(예: "커널 2")과 SRAM 뱅크 충돌(Boehm은 "커널 6"에서 행렬 A에 대해 다룸)을 온칩에서 B 타일을 전치함으로써 동시에 해결한다. 따라서 공유 메모리를 도입하지만 여전히 뱅크 충돌이 있는 Boehm의 "커널 3"보다 약간 더 진보된 전략이다.↩
두 개의 64x64 타일(4바이트 실수)은 2 * 64 * 64 * 4 = 32768 바이트(32 KB)의 공유 메모리를 요구한다. 이는 기본 설정에서 A100의 블록당 48 KB 한도 내에 들어간다. 커널이 더 많은 공유 메모리를 요청(최대 100 KB)할 수도 있지만, 그러면 SM에서 동시에 실행 가능한 스레드 블록 수가 줄어 점유율과 트레이드오프가 된다.↩
텐서 코어는 NVIDIA GPU의 특수 하드웨어 유닛으로, 행렬-곱-누산(D = A@B + C) 연산을 실행한다. Ampere GPU에서는 FP32 입력을 먼저 TensorFloat-32 (TF32) 형식으로 반올림해 가속한다. TF32는 FP32와 동일한 8비트 지수를 사용하지만 10비트 가수만 사용하며, 이는 FP16의 정밀도와 같다. 내부 누산은 FP32로 수행된다. 이 과정은 딥러닝 프레임워크의 일부 연산에 기본 적용되며, 정밀도 저하의 대가로 속도를 높인다. 최대 처리량을 위해 텐서 코어는 FP16, BF16, INT8 등에서도 동작한다. 그런데 손실형인데도 TF32를 TF32라고 부를 수 있는 걸까
SM에는 유한한 물리적 레지스터 파일(예: A100의 65,536개 32비트 레지스터)이 있다. 블록이 요구하는 총 레지스터 수는 (블록당 스레드 수) * (스레드당 레지스터 수)다. SM은 자신의 레지스터 파일과 공유 메모리 용량 안에 들어오는 만큼의 블록만 동시 수용할 수 있다. 따라서 스레드당 더 많은 레지스터를 사용하면 상주 가능한 블록 수가 줄어 점유율이 낮아진다. 컴파일러는 레지스터를 워프 단위로 고정 크기 청크로 할당하므로, 예컨대 스레드당 33개를 요청하면 40개가 할당되어 자원 계산에 더 악영향을 줄 수 있다.↩
독립 스레드 스케줄링이 발산의 비용을 없애지는 않는다. 각 스레드가 다음 명령을 가리키는 자신만의 프로그램 카운터(PC)를 갖지만, 워프는 여전히 사이클당 단일 주소에서만 명령을 발행할 수 있는 스케줄러에 의해 관리된다. 워프 내 PC가 분기되면, 스케줄러는 각 고유 경로를 직렬로 실행해야 하며 그동안 다른 스레드는 유휴 상태가 된다. 따라서 본문에 있는 단순화된 if-else 모델은 여전히 스레드 발산을 설명하는 정확한 방식이며, 직렬화는 워프의 병렬성을 약화시킨다.↩