병렬 Primitive 고급 단권화

1 warp-level 동기화 모델 ★

2 shfl 4-variant ★

각 intrinsic: (mask, val, delta, width=32). width는 logical warp 크기 32 나누기.

3 shfl_xor butterfly reduce·all-reduce

// warp 32-wide sum, result lane0..31 동일
for(int d=16; d>0; d>>=1)
  v += __shfl_xor_sync(0xffffffff, v, d);

4 vote intrinsic ★

5 match intrinsic sm_70+

• atomic 대신 peer-group aggregation에 사용
• histogram conflict-free 업데이트 패턴의 핵심

6 activemask / sync_warp

• __activemask(): 현재 convergent lane bitmask 조회
• __syncwarp(mask): warp 내부 barrier
• Volta ITS 이후 __syncwarp 필수 — lockstep 가정 금지

7 mask 생성 규칙 ★

8 warp reduce → broadcast

float warp_sum(float v){
  for(int d=16; d>0; d>>=1)
    v += __shfl_down_sync(0xff..ff, v, d);
  return __shfl_sync(0xff..ff, v, 0); // bcast
}

9 PTX 대응 ↗ V03 §9

• shfl.sync.{idx,up,down,bfly}.b32
• vote.sync.{all,any,ballot}.pred
• match.any.sync.b{32,64}

1 계층 객체 ★

헤더: <cooperative_groups.h>, namespace cg = cooperative_groups.

2 기본 패턴

namespace cg = cooperative_groups;
auto block = cg::this_thread_block();
auto warp  = cg::tiled_partition<32>(block);
auto quad  = cg::tiled_partition<4>(warp);

block.sync();        // == __syncthreads
warp.thread_rank();  // 0..31
warp.size();         // 32

3 tiled_partition 계층

4 collective 연산 ★

5 coalesced_group

• divergent branch 내부의 active lane만 모음
• cg::coalesced_threads()로 생성
• atomic 대체 aggregation에 사용

6 grid_group grid-level sync

auto grid = cg::this_grid();
grid.sync();   // 모든 block 대기

persistent kernel · multi-pass reduction을 1-kernel로 구현 가능. ↗ §9

7 cluster_group Hopper 전용

• __cluster_dims__(x,y,z) 또는 launch API
• DSMEM 접근 ↗ V04 §3
• cluster.sync() = cluster 내 block 전체 barrier

8 labeled_partition

histogram privatization, sort-within-warp 등에 활용.

1 7단계 요약표 ★★

PMPP Ch10 / Harris reduction talk 기반. 단계 번호는 관례적 표기.

2 1. Naive (divergent)

for(int s=1; s<BS; s*=2){
  if(tid % (2*s) == 0)
    sd[tid] += sd[tid+s];
  __syncthreads();
}

warp 내 절반 꺼짐 → SIMT divergence.

3 2. Interleaved (no mod)

for(int s=1; s<BS; s*=2){
  int i = 2*s*tid;
  if(i < BS) sd[i] += sd[i+s];
  __syncthreads();
}

divergence ↓ but bank conflict (stride 2,4,8…).

4 3. Sequential addressing ★

for(int s=BS/2; s>0; s>>=1){
  if(tid < s)
    sd[tid] += sd[tid+s];
  __syncthreads();
}

conflict ↓ · divergence ↓. warp 앞쪽이 활성.

5 4. First-add during load

int i = blockIdx.x*(2*BS) + tid;
sd[tid] = g[i] + g[i+BS];   // ½ grid
__syncthreads();

global load 효율 2×, grid ½.

6 5. Unroll last warp

if(tid < 32){
  volatile float* v = sd;
  v[tid]+=v[tid+32]; v[tid]+=v[tid+16];
  v[tid]+=v[tid+8];  v[tid]+=v[tid+4];
  v[tid]+=v[tid+2];  v[tid]+=v[tid+1];
}

warp 내부 __syncthreads 제거. volatile 필수 캐시 우회.

7 6. Warp shuffle ★

float v = sd[tid];
for(int d=16; d>0; d>>=1)
  v += __shfl_down_sync(0xff..ff, v, d);
if(tid==0) out[blockIdx.x] = v;

smem 경유 제거 · register only · Kepler+.

8 7. Multi-block + CUB

• 2단계: block partial → grid reduce
• 또는 atomicAdd로 최종 합산
• 또는 grid.sync()로 1-kernel
• 실전: cub::DeviceReduce::Sum · thrust::reduce

1 Scan 정의 ↗ V01 §12

2 복잡도 비교표 ★★

Hillis-Steele은 Kogge-Stone의 상위 계열 구조 동일.

3 Kogge-Stone regular 구조

for(int o=1; o<BS; o*=2){
  float v=0;
  if(tid>=o) v = sd[tid-o];
  __syncthreads();
  if(tid>=o) sd[tid] += v;
  __syncthreads();
}

W = N log N · D = log N · 모든 단계 all-active.

4 Brent-Kung up·down ★

1. Up-sweep: reduction tree 올리기 N 비용
2. Root = identity로 초기화
3. Down-sweep: 좌측 자식으로 부모를, 우측으로 (부모 ⊕ 좌측 기존값)

5 Sklansky divide-and-conquer

[ L ][ R ] → scan(L), scan(R)
                 → R += last(L)

step = log N · 동일 level 모두 독립 → bank conflict 주의.

6 Warp scan (shfl_up) ★

float v = x;
for(int d=1; d<32; d*=2){
  float u = __shfl_up_sync(0xff..ff, v, d);
  if(lane >= d) v += u;
}

inclusive warp scan · 32-wide · 5-step.

7 3-phase hierarchical ★

1. block 내 local scan (warp → block)
2. block 총합 수집 → aux[ ]
3. aux 자체를 scan
4. block i 원소에 aux[i−1] 가산

blk0    blk1    blk2
 ↓        ↓        ↓
local   local   local
  Σ0     Σ1     Σ2
  └─ scan aux ──┘
  +0     +Σ0    +Σ0+Σ1

8 Segmented scan

응용: SpMV reduce, RLE, 계층적 reduction.

1 Histogram 문제

2 privatization 4-level ★

3 shared privatization

__shared__ int sh[K];
// init
for(...) atomicAdd(&sh[b], 1);
__syncthreads();
// flush
if(tid<K) atomicAdd(&hist[tid], sh[tid]);

4 aggregation ★

int b = bin(x);
int peers = __match_any_sync(0xff..ff, b);
int cnt = __popc(peers);
int leader = __ffs(peers)-1;
if(lane == leader)
  atomicAdd(&sh[b], cnt);

contention K → contention warp 수준으로 축소.

5 thread-coarsening

• thread 하나가 C 개 원소 처리
• block 수 ↓ · smem init/flush overhead 분산
• register pressure 와 trade

6 sort-then-reduce

1. 입력을 bin id로 radix sort
2. segmented reduce로 count
3. K 크거나 분포 skewed 시 atomic보다 유리

7 기법 선택 결정표 ★

8 다차원 히스토그램

• bin 공간 = 차원들의 카르테시안 곱
• 선형화 후 동일 pattern
• 2D 이미지 hist: tile 분할 후 shared privat

9 CUB API ↗ §13

• cub::DeviceHistogram::HistogramEven
• cub::DeviceHistogram::HistogramRange
• cub::DeviceHistogram::MultiHistogramEven (channel)

1 문제 정의

2 접근 3분류 ★

3 Bitonic Top-K

1. 각 block이 local bitonic sort
2. 상위 K만 global heap/buffer에 push
3. 두 번째 pass로 K·B → K 축소

4 Radix Top-K ★

1. 최상위 b-bit로 bucket 분류
2. bucket count 누적 → K 들어가는 경계 bucket 결정
3. 경계 bucket만 다시 b-bit 재귀
4. 나머지 bit 확정까지 반복

b = 4 or 8. float은 IEEE monotone reorder 필요 부호·지수 처리.

5 Float bit reorder

6 warp-level top-K

• warp 내 32 lane이 자기 값 보유
• shfl_xor butterfly sort
• K ≤ 32이면 1-warp로 종결

7 CUB · lib 활용

8 적용 맥락

• LLM sampling top-k ↗ V08 §10
• MoE router top-k ↗ V08 §1
• sparse attention indexing

9 결정 트리

K ≤ 32 ──► warp bitonic
K ≤ 256 ─► block bitonic
K 중간 ──► radix select
K 큼 ───► full sort

1 왜 Split-K인가 ★

2 Split-K 아이디어

3 합산 방식 2종 ★

4 Atomic variant 의사코드

// slice s, tile (m,n)
accum = 0;
for(k=s·kseg; k<(s+1)·kseg; k+=BK)
  accum += A[m,k:k+BK] · B[k:k+BK,n];
atomicAdd(&C[m,n], accum);

tile 단위 하나의 atomic. FP32면 hardware atomic 가능.

5 Buffer + Reduce variant

// stage 1: partial
for(s, m, n) C_part[s,m,n] = A_s · B_s;

// stage 2: reduce
C[m,n] = Σ_s C_part[s,m,n];

stage 2는 O(S·M·N) reduce kernel. FP 합 순서 deterministic.

6 S (split 수) 선택

• 목표: #CTA · S ≳ #SM · 2
• S 과대: K 조각당 K_seg가 너무 작아져 MMA latency 은폐 실패
• 실전 S ∈ {2, 4, 8, 16}

7 Atomic vs Buffer 선택표 ★

8 Split-K의 한계

• 각 CTA에 동일 K_seg → work balance 정적
• M,N이 BM·BN과 정합 안 되면 tile 경계 낭비
• 이 문제를 해결하는 게 Stream-K ↗ §8

9 CUTLASS 제공 ↗ V06

• SplitKSerial (buffer+reduce)
• SplitKParallel (atomic)
• epilogue reduce kernel 자동 생성

1 동기 ★

2 핵심 아이디어

• CTA = SM 수와 동일 (persistent)
• 각 CTA가 MMA iteration의 flat range를 담당
• tile 경계에 걸리면 부분 accumulator 저장
• fixup kernel이 조각을 합산

3 work 분할

iter: 0 1 2 3 4 5 6 7 8 9 ...
CTA0: [ 0 .. 3 ]
CTA1:         [ 4 .. 6 ]
CTA2:                 [ 7 .. 9 ]
                       ↑ tile 경계에서 split

Osama et al. 2023 Stream-K. CUTLASS 3.x에 정식 통합.

4 3 variant

Stream-K는 실제로 DP 블록 + Stream-K 블록을 혼합 운용하기도 한다.

5 Fixup kernel

// per partial slab
for(tile) C[tile] = Σ partials[tile][*];

partial은 on-chip 또는 global workspace에 저장. semaphore/flag로 dependency.

6 Semaphore 패턴

• tile-owner CTA만 최종 write
• 보조 CTA는 partial + flag set
• owner가 flag poll → 합산 → C 기록

7 Split-K vs Stream-K 결정표 ★★

8 핵심 수식

9 CUTLASS 구현 개요 ↗ V06 §13

• TileScheduler 추상
• StreamKScheduler dispatch policy
• fixup은 별도 kernel or fused epilogue
• workspace = (#SM · BM · BN) partial buffer

1 정의 ★

2 왜 쓰는가

• uneven work (tile별 iter 수 변동)
• launch overhead 누적 회피
• grid-level sync 필요 시 (cooperative)
• Stream-K의 기반 ↗ §8

3 기본 스켈레톤

__global__ void persist(..., int* ctr){
  while(1){
    int id = (tid==0) ?
      atomicAdd(ctr, 1) : 0;
    id = __shfl_sync(0xff..ff, id, 0);
    if(id >= N_WORK) break;
    process_work(id);
  }
}

4 work scheduling 3 패턴

5 grid size 선택 ★

cooperative launch 시: grid ≤ GPU가 동시에 launch 가능한 상한.

6 grid.sync() 활용

• multi-pass reduction을 1-kernel로
• partial → sync → final
• launch overhead × pass 회피

7 사용 상황 표

8 주의 사항

• debug 어려움: grid-level race
• atomic counter는 cudaMemsetAsync로 매 invocation 초기화
• early-exit CTA가 있으면 cooperative sync 막힘
• register pressure ↑ (loop body 길어짐)

9 활용 예

• FA forward persistent ↗ V07
• Stream-K GEMM
• graph BFS frontier expand
• inference decoder loop

1 패턴 개요 ★

• Hopper WGMMA warp specialization ↗ V04 §9
• CUTLASS cooperative / pingpong
• single-GPU에선 shared ring 중심

2 Ring buffer 자료구조

3 stage 수 선택

stage↑ → smem ↑ → occupancy ↓ trade.

4 Sync mechanism 3종

5 Lock-free ring 의사코드

// producer (warp P)
while((head+1)%N == tail) ;
buf[head] = load_tile();
__threadfence_block();
head = (head+1)%N;

// consumer (warp C)
while(head == tail) ;
t = buf[tail];
compute(t);
tail = (tail+1)%N;

head/tail은 shared, 각각 전담 warp만 write. memory fence로 ordering.

6 mbarrier phase 기반 Hopper

• 각 stage에 full-bar · empty-bar
• producer: arrive(full), wait(empty)
• consumer: wait(full), arrive(empty)
• phase bit 토글로 2-stage 라운드트립

7 warp specialization

warp 0-3   : producer (TMA load)
warp 4-7   : consumer A (WGMMA)
warp 8-11  : consumer B (WGMMA)  pingpong

Hopper setmaxnreg로 producer ↓ reg / consumer ↑ reg 재분배.

8 성능 이득 요인

• compute ↔ load 완전 중첩
• register file 효율 재분배
• instruction issue port 분산
• global load latency 은폐

9 단점·함정

• deadlock: arrive/wait 짝 안 맞으면 영구 정지
• stage 수 잘못 잡으면 stall
• smem 사용량 stage·tile 크기 곱으로 팽창
• debug 비용 ↑ 코드보다 설계

1 GPU 정렬 3-family ★

2 Bitonic Sort 개념

stage k (k=1..log N):
  substage j (j=k..1):
    compare-swap (i, i ^ (1<<(j-1)))

3 Bitonic 코드

for(k=2; k<=N; k*=2)
  for(j=k/2; j>0; j>>=1){
    p = tid ^ j;
    if(p > tid){
      up = ((tid & k) == 0);
      if((a[tid]>a[p])==up) swap(a[tid],a[p]);
    }
    __syncthreads();
  }

data-independent → warp divergence ↓. power-of-2 크기.

4 Merge Path ★

↗ V01 §13 PMPP Ch12 co-rank 함수 정의.

5 Radix Sort 개요 ★

1. b-bit씩 digit 추출 (LSB → MSB)
2. 각 digit에 대해 counting sort (stable)
3. counting sort 내부 = histogram + scan + scatter
4. b=4 (16 bucket) 또는 b=8 (256) 일반적

6 Radix 한 iteration

input → digit(input) → local hist
      → scan hist → global hist
      → scan global → offsets
      → scatter to output

Scan · Histogram 재사용 ↗ §4, §5.

7 Float 정렬 bit trick

radix sort를 float·double에 적용하는 표준 전처리.

8 Merge vs Radix 선택

9 Segmented sort

• 여러 작은 정렬을 동시에
• 각 segment 독립 key·value
• cub::DeviceSegmentedRadixSort
• MoE · grouped GEMM 전처리 ↗ V08

1 연산 3종 ★

A: sparse · B,x,Y: dense · M: sparsity mask.

2 Storage 포맷 ↗ V01 §14

3 SpMV 전략 3종 ★

4 Vector SpMV 의사코드

// warp = 32 thread / row
r = blockIdx.x;
float s = 0;
for(j=row_ptr[r]+lane; j<row_ptr[r+1]; j+=32)
  s += val[j] * x[col_idx[j]];
// warp reduce
for(d=16; d>0; d>>=1)
  s += __shfl_down_sync(0xff..ff, s, d);
if(lane==0) y[r] = s;

5 Merge-based SpMV

Merrill & Garland 2016. skew 심한 그래프에 강함.

6 SpMM 전략

• 출력 열을 여러 묶음으로 tiling
• 각 tile에 대해 A의 row 순회
• B의 해당 row를 dense로 load
• accumulate

row-by-row × B-column-tile = 일반적 패턴. cuSPARSE 기본 경로.

7 SDDMM 특성

• nnz(M) 만큼만 inner product
• Graph attention, GNN에 핵심

8 2:4 structured sparsity

9 Library 맵

• cuSPARSE: CSR/COO SpMV/SpMM
• cuSPARSELt: 2:4 structured
• CUB: primitive block
• DGL·PyG: graph frontend

1 CUB 계층 ★

CUB = "CUDA UnBound". NVIDIA/cub on GitHub.

2 Device API 핵심

3 CUB 2-call pattern ★

void* d_temp = nullptr;
size_t temp_bytes = 0;
cub::DeviceReduce::Sum(
  d_temp, temp_bytes, d_in, d_out, N);
cudaMalloc(&d_temp, temp_bytes);
cub::DeviceReduce::Sum(
  d_temp, temp_bytes, d_in, d_out, N);

1회차: 필요한 temp 크기 조회. 2회차: 실제 실행.

4 Block primitive

5 Warp primitive

• WarpReduce<T>
• WarpScan<T>
• shfl 기반 내부 구현

6 Thrust 대응 맵 ★

7 선택 기준

• Thrust: STL-like, prototyping 빠름, iterator 추상
• CUB: lower level, custom kernel 내부에서 block/warp collective
• Thrust는 내부적으로 CUB 호출

8 기타 library

1 Primitive complexity ★

2 Warp intrinsic 표

3 Reduction 단계 1-liner

1. naive mod
2. interleaved index
3. sequential addr
4. first-add load
5. unroll last warp
6. shfl_down warp
7. multi-block/CUB

4 Scan 4종 (W / D)

5 GEMM 분할 선택 ★

6 결정 트리 (primitive)

Task ?
├ scalar out ──► Reduction
├ per-elem out ──► Scan
├ bin count ───► Histogram
├ sorted perm ─► Radix sort
├ k largest ───► Top-K (radix)
├ y = A·x ────► SpMV (merge)
└ mask sel ───► Compact (scan)

7 구현 단계 원칙

1. 라이브러리(CUB/Thrust) 우선
2. 안 맞으면 block primitive + custom kernel
3. hot path만 warp shuffle로 최적화
4. 결정론 · 수치 오차 검증

8 출처

• PMPP 4e Ch 10–14
• CUDA CPG §7 warp-level
• Cooperative Groups guide (CUDA 11+)
• CUB docs · Thrust docs
• Stream-K: Osama et al. 2023
• Merge-based SpMV: Merrill & Garland 2016
• Harris: Optimizing Parallel Reduction (NVIDIA)