Scan Algorithm

Izzat 의 강의는 PMPP 교과서 7장 의 raw 영상 버전으로 보면 거의 정확하다. 다른 chapter 와 달리 scan 은 알고리즘이 직관과 정반대다 — sequential 로 짜면 5줄, parallel 로 잘 짜려면 한 chapter 가 든다.

강의가 깐 큰 질문 두 개.

1. 왜 scan 이 중요한가 — radix sort, stream compaction, segmented reduction, dynamic memory allocation, sparse matrix 처리 등 거의 모든 GPU 데이터 알고리즘의 빌딩블록. L019 의 CSV 파서가 그 사용 사례 중 하나.
2. obvious 한 sequential 을 어떻게 parallel 로 짜는가 — 두 가지 핵심 알고리즘 (Hillis-Steele, Brent-Kung) 이 답을 준다. 같은 결과, 다른 trade-off.

scan = prefix sum 의 일반화. binary 결합 연산 ⊕ 와 입력 [x0, x1, ..., x_{n-1}] 가 주어지면, 출력은 위치 i 의 누적값. 두 변형이 있다.

일반적으로 ⊕ 는 associative (결합법칙) 면 충분하다 — commutative 일 필요는 없다. min, max, or, and, 행렬 곱 모두 가능. associative 하지 않은 op (예: subtraction) 는 scan 으로 표현할 수 없다.

가장 단순한 parallel scan. step k 에서 thread i 는 자기 값에 2^k 떨어진 element 를 더한다. log(n) step 후 모든 thread 가 자기 자리의 prefix sum 을 들고 있다.

// Hillis-Steele — block 단위 scan (shared memory 위)
__global__ void hillis_steele(const float* in, float* out, int n) {
  __shared__ float buf[2][BLOCK];     // double buffer
  int tid = threadIdx.x;
  buf[0][tid] = in[tid];
  __syncthreads();

  int read = 0, write = 1;
  for (int stride = 1; stride < n; stride *= 2) {
    if (tid >= stride)
      buf[write][tid] = buf[read][tid] + buf[read][tid - stride];
    else
      buf[write][tid] = buf[read][tid];
    __syncthreads();
    read ^= 1; write ^= 1;
  }
  out[tid] = buf[read][tid];
}

두 phase 알고리즘. up-sweep (reduce) — binary tree 를 위로 올라가면서 partial sum 을 모은다. down-sweep — 트리를 다시 내려오면서 누적값을 분배한다.

// Brent-Kung — 약식 (up-sweep)
for (stride = 1; stride < BLOCK; stride *= 2) {
  int idx = (tid + 1) * stride * 2 - 1;
  if (idx < BLOCK) buf[idx] += buf[idx - stride];
  __syncthreads();
}
// down-sweep (exclusive 의 경우 root 를 0 으로 set 후)
for (stride = BLOCK/2; stride > 0; stride /= 2) {
  int idx = (tid + 1) * stride * 2 - 1;
  if (idx + stride < BLOCK) {
    float t = buf[idx];
    buf[idx] += buf[idx - stride];
    buf[idx - stride] = t;
  }
  __syncthreads();
}

PMPP 가 가르치는 정식 frame. work efficiency = 알고리즘이 sequential 과 같은 op 수를 쓰는가. step efficiency = critical path 의 길이가 얼마나 짧은가. 두 축은 trade-off — Hillis-Steele 은 step 이 짧지만 work 가 많다.

강의에서 Izzat 이 직접 보여준 작은 “함정”. naive 한 Hillis-Steele 구현은 buf[tid] = buf[tid] + buf[tid - stride] — 같은 buffer 에 읽고 쓴다. 어떤 thread 가 자기 자리를 update 한 뒤, 다른 thread 가 그 update 된 값을 읽으면 결과가 망가진다.

// 잘못된 in-place — race condition
if (tid >= stride)
  buf[tid] += buf[tid - stride];   // race!
__syncthreads();

해결은 double buffering — read 와 write 를 다른 buffer 로 분리, 매 step 마다 swap. buf[2][BLOCK] 으로 두 개를 두고 read, write index 를 toggle.

강의의 미묘한 디테일. Volta 이전 GPU 에서는 같은 warp 의 32 thread 가 lockstep 으로 도는 게 자연스러운 가정이었다 — warp-internal 동기화가 implicit. Volta 부터 independent thread scheduling 이 들어가서 — warp 안 thread 들이 더 이상 lockstep 으로 도는 게 보장되지 않는다.

그리고 강의의 또 한 자리 — 같은 warp 안 thread 들 사이에서도 ordering 보장이 약해서 “thread A 가 buf[5] 에 쓰고 thread B 가 그 직후 buf[5] 를 읽는다” 같은 패턴은 __syncthreads() 또는 __syncwarp() 없이는 안전하지 않다. scan 처럼 stride 마다 의존성이 바뀌는 알고리즘에서 매 step 후 반드시 sync.

// 안전한 패턴 (Volta+)
for (stride = 1; stride < BLOCK; stride *= 2) {
  float v = (tid >= stride) ? buf[read][tid - stride] : 0.0f;
  __syncthreads();
  buf[write][tid] = buf[read][tid] + v;
  __syncthreads();
  read ^= 1; write ^= 1;
}

강의 끝부분에서 짧게 짚는 자리지만, scan 이 어디에 쓰이는지 이해하지 않으면 “왜 이렇게 깊이 들어가는가” 의 동기가 약해진다. 대표 응용 5가지.

손에 새기기 — 실습 시퀀스

1. Hillis-Steele 직접 짜기 — n=1024 의 single-block scan. shared memory + double buffer. CPU sequential 결과와 비교.
2. Brent-Kung up/down sweep — index 계산 (idx = (tid+1) * 2*stride − 1) 을 종이에 그려보고 코드로. exclusive 와 inclusive 두 버전.
3. work-vs-step 측정 — 두 알고리즘의 add 호출 수를 카운트해서 sequential 대비 비율 확인. n=1024 면 Hillis-Steele 약 10×, Brent-Kung 약 1×.
4. volatile/__syncwarp 실험 — Hillis-Steele 의 inner loop 에서 sync 없이 돌려보고 결과가 nondeterministic 한지 확인.
5. warp shuffle 버전 — __shfl_up_sync 로 32-thread Hillis-Steele. shared memory 안 쓰고 register-only.
6. compaction 만들기 — random array 의 “positive only” compaction 을 mask + exclusive scan + scatter 로 직접.
7. NCU 비교 — 두 알고리즘의 instruction count, achieved occupancy, scoreboard stall. work-efficient 가 항상 빠른지 직접 검증.

• shared memory bank conflict — Brent-Kung 의 index 패턴이 bank conflict 를 만드는 자리가 있다. PMPP 가 padding trick (CONFLICT_FREE_OFFSET) 으로 해결. 강의에서는 짧게 언급.
• warp shuffle scan — 강의에서는 일반 shared memory 버전 위주. __shfl_up_sync 의 정식 코드는 L021 의 영역.
• inclusive ↔ exclusive 변환의 비용 — exclusive 를 inclusive 로 바꿀 때 한 step 의 추가 work. 강의의 BLOCK 안에서는 무시할 수 있지만 multi-block 으로 가면 의미 있음.
• thread coarsening — 한 thread 가 여러 element 처리해서 work 를 키우는 대신 BLOCK 의 thread 수를 줄이는 트릭. L021 의 본격 주제.
• strict CCCL 비교 — CUB 의 BlockScan 이 정확히 어떤 알고리즘 (RAKING/WARP_SCANS/MATRIX) 을 default 로 쓰는지 직접 source 확인.
• numerical stability — float scan 은 더하는 순서에 따라 round-off 누적 차이가 있다. associative 가 numerical 로는 정확히 성립하지 않음. 강의에서 언급 안 됨 — 후속 자료로 보강.