Scan Algorithm Part 2

Part 1 (L020) 은 한 block 안의 scan 까지 깔았다. 실전에서 scan 을 하는 array 는 거의 항상 그것보다 훨씬 크다 — 10M, 1B element. 한 block 의 thread 한도(1024)를 넘어선다. 이걸 어떻게 짜는가가 part 2 의 본론.

강의가 깐 큰 질문 세 개.

1. multi-block scan 의 구조 — 한 block 안 scan 을 어떻게 “이어 붙이는가”. 같은 reduction tree 가 한 단계 더 위로 올라간다(§ 03–04).
2. 한 단계 더 내려가서 — warp 안에서 scan 을 더 빠르게 — __shfl_up_sync 로 register-only scan(§ 05).
3. 왜 더 빨라지지 않는가 — Izzat 이 NCU 로 보여주는 self-finding: scan 이 memory-bound 처럼 보이지만 실은 __syncthreads() 의 latency 가 dominant. 이건 § 09 의 핵심 지점이고 L024 의 출발점.

part 2 가 part 1 의 코드 위에 올라간다. 잊었으면 못 따라간다 — 5분 복습.

• 입력 — global memory 의 N=1024 element. 이걸 shared memory 로 한 번 load.
• up-sweep — log(BLOCK) step. 각 step 의 thread 절반이 활성. partial sum 이 stride 자리에 누적.
• down-sweep — log(BLOCK) step. exclusive scan 의 경우 root 를 0 으로 set 후 trickle down.
• 출력 — shared memory 에서 global 로 한 번 store.
• cost — 2log(BLOCK) ≈ 20 step (1024 의 경우). 매 step 후 __syncthreads().

강의의 첫 큰 figure. 같은 idea 의 3 단계 launch 패턴. 아이디어 자체는 단순한데 — index 계산과 buffer 관리가 까다롭다.

이 패턴의 weakness — 중간 partial sum 을 HBM 에 적었다 다시 읽는다. 한 array 를 두 번 통과. § 09 에서 measurement 가 이 비용을 정확히 보여준다.

Izzat 이 강의에서 강조한 mental model. multi-block scan 은 사실 recursive 다. block 의 partial sum array 가 또 너무 크면, 그 array 도 같은 multi-block scan 으로 풀면 된다.

각 layer 가 자기 위 layer 의 “unit” 이다. 같은 합산 idea 를 서로 다른 단위 (32, 1024, N) 에 한 번씩 적용. 이게 hierarchical scan 의 정신.

강의의 가장 만족스러운 트릭 중 하나. 32 thread 의 scan 을 — shared memory 도 __syncthreads() 도 안 쓰고 — warp shuffle instruction 만으로 5 step 안에 끝낸다.

// warp-level Hillis-Steele inclusive scan — 32 thread
__device__ float warp_inclusive_scan(float v) {
  for (int stride = 1; stride < 32; stride *= 2) {
    float n = __shfl_up_sync(0xffffffff, v, stride);
    if (threadIdx.x % 32 >= stride) v += n;
  }
  return v;
}

강의에서 Izzat 이 직접 측정한 차이 — warp-level scan 으로 바꾸면 block-level scan 의 시간이 약 30–40% 감소. scan 이 latency-bound 임을 보여주는 첫 강한 신호(§ 09).

block 안 1024 thread 를 32 warp 으로 본다. 각 warp 가 독립적으로 § 05 의 warp scan 을 한다. 그 다음 — 32 개의 “warp 마지막 값” 을 다시 scan (한 warp 으로). 마지막으로 모든 warp 가 자기 prefix 를 받는다.

// block-level scan with warp scans — 약식
__device__ float block_inclusive_scan(float v) {
  int tid  = threadIdx.x;
  int lane = tid & 31;
  int wid  = tid >> 5;
  __shared__ float warp_sums[32];

  v = warp_inclusive_scan(v);                     // A — warp scan
  if (lane == 31) warp_sums[wid] = v;          // B
  __syncthreads();

  if (wid == 0) {                                  // C — warp-of-warp-sums scan
    float w = (lane < 32) ? warp_sums[lane] : 0.0f;
    w = warp_inclusive_scan(w);
    warp_sums[lane] = w;
  }
  __syncthreads();

  if (wid > 0) v += warp_sums[wid - 1];          // D — broadcast
  return v;
}

강의의 마지막 큰 트릭. 한 block 의 thread 수가 1024 로 cap 되어 있으니, “각 thread 가 K 개 element 를 들고 있게” 하면 block 하나가 1024·K 개 element 의 scan 을 처리.

각 thread 의 일.

1. K 개 element 를 register 에 load — vectorized load (float4) 면 효율적.
2. thread 안에서 sequential scan — K 가 작으니 (4, 8) 빠르다.
3. 각 thread 의 “마지막 값” (= thread 안 K 개의 합) 만 들고 block-level scan 에 들어간다 — § 06.
4. 받은 prefix 를 자기 thread 의 K 개 element 에 다시 더해서 store.

thread coarsening 의 trade-off. K 개 element 를 register 에 들고 있으려면 — 한 thread 의 register footprint 가 K 만큼 커진다. 그게 너무 크면 SM 위에 띄울 수 있는 block 수가 줄어든다 (occupancy 감소).

강의의 가장 통찰적인 자리. Izzat 이 NCU 로 측정해 보여주는 사실 — naive 한 Brent-Kung block scan 은 DRAM throughput 이 peak 의 30–40% 정도밖에 안 나온다. 흔한 첫 직관 — “더 vectorize 해야지, 더 coalesce 해야지” — 가 맞지 않는다.

실제 dominant cost 는 __syncthreads(). 매 step 마다 모든 thread 가 barrier 를 친다. 1024 thread block 에서 10 step → 10 번의 sync. sync 자체가 latency 를 만든다 — instruction 자체보다 “모두가 도착하기를 기다리는” 시간.

손에 새기기 — 실습 시퀀스

1. 3 launch multi-block scan — 1M element 위에서 직접 짜기. block scan + partial scan + add. CPU 결과와 비교.
2. warp_inclusive_scan — __shfl_up_sync 만 써서 32-thread scan. shared memory 안 쓰는 버전.
3. block scan with warp scans — § 06 의 구조. __syncthreads() 가 정말 2 번만 들어가는지 코드로 검증.
4. thread coarsening sweep — K ∈ {1, 2, 4, 8, 16} 로 측정. 어디서 best 인지, occupancy 가 어디서 떨어지기 시작하는지.
5. NCU latency 분석 — “Long scoreboard stalls” 가 얼마인지. __syncthreads 의 stall % 가 차지하는 비중 확인.
6. CUB vs 직접 구현 — cub::DeviceScan::ExclusiveSum 과 자기 구현의 시간 비교. 차이가 어디서 오는지 hint 만 봐도 풍부함.
7. 1B element scaling — 1B element 까지 키워보고 시간이 N 에 정확히 linear 한지. 만약 N log N 이면 어딘가 알고리즘이 나쁜 것.

• NCU bandwidth % 의 실제 숫자 — § 09 의 35/55/75% 는 강의 narrative 를 재구성한 개념적 값. 자기 GPU 에서 직접 측정해야 정확한 분포를 본다 (확인 필요).
• K 의 sweet spot — 강의에서 정확한 best K 를 명시 안 함. arch 별 (Volta/Ampere/Hopper) 로 다를 수 있다. 환경마다 sweep 필요.
• occupancy 와 latency hiding 의 정량 모델 — “register 늘면 occupancy 줄지만 빠를 수 있다” 의 정확한 break-even point 는 강의에서 정량화 안 됨.
• warp 안 partial sum 위치 트릭 — __shfl_up_sync 의 mask 와 lane 검사 패턴이 미묘. 강의에서 다 보여주지는 않음. CUB source 직접 까기.
• HBM 트래픽 분석 — 3 launch 의 partial sum 이 HBM 을 두 번 통과한다. 그 cost 가 sync 의 cost 와 비교해 얼마인지 강의에서 정량 안 됨.
• decoupled lookback 의 메커니즘 — § 09 끝에서 언급되는 “1 launch 짜리 scan”. L024 에서 본격 다뤄짐.