Scan at the Speed of Light

Jake 와 Georgii 의 강의는 L021 가 끝난 자리에서 출발한다. multi-block scan 이 동작하긴 하는데 — peak 의 절반도 안 나온다. 왜? 어떻게 더 짜낼 수 있는가? 이론적 한계는 어디인가?

강의가 깐 큰 질문 세 개.

1. “빠르다” 의 절대 기준 — speed-of-light. Nvidia 가 내부적으로 쓰는 분석 frame. 이건 § 02 의 본론.
2. scan 의 SoL 이 정확히 무엇인가 — input 한 번 read, output 한 번 write 의 memcpy. 그것보다 빠를 수는 없다(§ 03).
3. 3 알고리즘의 상한이 SoL 의 몇 % 인가 — reduce-then-scan ~50%, hierarchical ~50%, stream ~2%, decoupled lookback ~95%. 이 분석이 알고리즘 선택의 frame(§ 04–07).

SoL 의 정의 — 알고리즘이 반드시 해야 하는 minimum work / hardware 의 peak throughput. 두 종류가 있다.

• compute-bound SoL — 알고리즘의 산술 양 (FLOPs) ÷ peak FLOPS. matmul 같은 algorithm.
• memory-bound SoL — 알고리즘의 minimum 메모리 트래픽 (byte) ÷ HBM bandwidth. scan, copy, reduction 같은 algorithm.

알고리즘이 inherently memory-bound 이면 (예: scan), “반드시 read 해야 하는 byte 수 + 반드시 write 해야 하는 byte 수” ÷ HBM 가 SoL 시간의 하한.

scan 의 minimum memory 트래픽 — input N 을 한 번 읽고, output N 을 한 번 쓴다. 총 2N memory operation. 산술은 byte 수보다 훨씬 작아 무시 가능. scan 의 SoL = 2N / HBM bandwidth.

그리고 이건 정확히 memcpy 의 시간. memcpy 도 N byte read + N byte write. 같은 시간.

가장 단순한 multi-block scan 의 한 형태. “각 block 의 sum 을 먼저 다 구한다 — 그 sum 들의 prefix 를 구한다 — block 을 다시 통과하면서 prefix 를 더한다”.

1. pass 1: 각 block 이 자기 segment 의 sum 만 계산 (intermediate 저장 안 함). N read.
2. pass 2: 작은 array 의 sum 들의 scan (작아서 무시).
3. pass 3: 각 block 이 자기 segment 다시 통과 — prefix 받아서 scan + write. N read + N write.

총 메모리 트래픽 — 2N read + N write = 3N. SoL 의 2N 대비 67%. 최선 ~67% SoL.

L021 의 multi-block scan. block 별 scan 결과를 일단 메모리에 적었다, partial sum 을 또 scan, 다시 읽어서 더한다.

1. launch 1: 각 block scan. N read + N write (intermediate scan 결과 저장).
2. launch 2: partial sum array 의 scan (작아서 무시).
3. launch 3: 각 block 의 prefix add. N read + N write.

총 — 2N read + 2N write = 4N. SoL 2N 대비 50%. 강의의 figure 가 이걸 정확히 보여줌 — “이 알고리즘은 이론적으로 50% 가 한계”.

Naive 한 single-pass 시도. 한 kernel 안에서 block 0 이 자기 prefix 를 다 만들고, block 1 이 그 prefix 를 받아서 자기 것 만들고, … 직렬로. 메모리는 2N (read + write 각 한 번).

이론상 SoL 100%. 그런데 실제 측정 — SoL 의 2% 미만. 왜?

이게 강의의 가장 큰 통찰 중 하나 — “단순히 메모리 트래픽이 적은 알고리즘이 빠르지 않다”. parallelism 도 같이 봐야 한다. SoL 분석이 만능은 아니라는 자리.

stream scan 의 trick — “block 들이 자기 sum 을 빨리 publish 한 뒤, 다른 일을 하면서 prefix 를 기다린다”. NVIDIA 가 발표한 single-pass scan 의 핵심.

각 block 의 한 “status flag” — 세 상태.

한 block 의 한 round

1. local sum — 자기 segment 합 계산. flag = A, sum publish.
2. lookback — 자기 앞 block 들을 거꾸로 본다. 각 block 의 flag 가:
   • X → spin (아직 안 끝났다)
   • A → 그 sum 을 받아서 누적, 더 앞으로 lookback
   • P → 그 prefix 를 받아서 멈춤
3. publish prefix — 자기 prefix sum 결정. flag = P, prefix publish.
4. local scan + write — 자기 segment 에 prefix 를 더해 출력 write.

강의의 실용적 메시지. “decoupled lookback 을 직접 짜면 두 달 걸린다 — 이미 CUB 안에 있다.” Thrust 의 thrust::inclusive_scan / thrust::exclusive_scan, CUB 의 cub::DeviceScan::ExclusiveSum 같은 라이브러리가 그걸 호출.

// CUB DeviceScan — 한 줄
size_t temp_bytes = 0;
cub::DeviceScan::ExclusiveSum(
    nullptr, temp_bytes, d_in, d_out, n);

void* d_temp;
cudaMalloc(&d_temp, temp_bytes);

cub::DeviceScan::ExclusiveSum(
    d_temp, temp_bytes, d_in, d_out, n);
// 내부적으로 decoupled lookback — ~95% SoL

강의의 마지막 메시지 — 자기 알고리즘을 평가할 때 “이게 SoL 의 몇 % 인가” 의 frame 을 항상 머리에 두기. 그리고 그 % 가 알고리즘 자체의 한계인지, 구현의 한계인지 분리해 보기.