Reductions

L002~L005 까지의 모든 예제 (vector add, RGB→grayscale, mean filter, matmul) 는 한 thread 가 자기 출력 element 만 책임지는 embarrassingly parallel 패턴이었다. reduction 은 그 가정이 처음 깨지는 자리. 한 출력 (예: sum) 을 만들기 위해 모든 thread 가 협력해야 한다. block 사이의 통신, atomic, syncthreads, warp shuffle 같은 새 도구들이 모두 이 한 영역에서 등장.

강의가 답하는 세 질문.

1. 병렬 sum 이 왜 trivial 하지 않은가 — 의존성 있는 알고리즘의 일반적 어려움.
2. 같은 sum 이 어떤 점에서 한 단씩 빨라지는가 — naive 부터 7개 kernel 의 진화.
3. PyTorch 의 일반 reduction 은 어떻게 동작하는가 — min, max, norm, mean 이 한 dispatch path 에서.

시퀀셜 sum 은 한 줄 — for i: s += a[i]. 의존성이 있다 (다음 step 이 이전 step 의 결과를 본다). 단순히 thread N 개로 쪼개면 — 각 thread 가 서로의 결과를 기다려야 한다. 답은 tree다.

• n 개 element 의 sum 을 첫 step 에서 n/2 개의 pair 로 모은다 (각 pair 가 1 thread).
• 다음 step 에서 n/4 개, 그 다음 n/8 ... 총 log₂(n) step.
• 각 step 사이에 모든 thread 의 진행이 sync — block 안이면 __syncthreads, block 사이면 kernel boundary 또는 atomic.

이 tree 패턴이 모든 reduction 알고리즘의 base. min, max, norm, dot product, softmax 의 normalize 부분 모두 같은 형태. “연산이 결합 법칙을 만족하면 (associative) tree 로 풀 수 있다” 가 reduction 알고리즘의 충분 조건.

tree reduction 의 step 별 active thread 와 그 지속 거리는 두 방식 — “인접 pair” vs “half-half”. 같은 알고리즘이지만 warp 활용도가 다르다.

이 한 시각 차이가 강의의 핵심 진화. naive (simple_reduce.cu) 가 “인접 pair” 방식이고 — 후속 버전이 모두 sequential addressing 으로.

강의의 첫 reduction. 한 thread 가 인접 pair 를 합치고 sync, 그 결과를 다음 step 의 wider stride 로 다시 합치고 sync. 코드는 짧고 명확하지만 warp divergence 가 심하다.

// simple_reduce.cu — 강의 repo 그대로
__global__ void SimpleSumReductionKernel(float* input, float* output)
{
    unsigned int i = 2 * threadIdx.x;
    for (unsigned int stride = 1; stride <= blockDim.x; stride *= 2) {
        if (threadIdx.x % stride == 0) {                    // ← divergence 의 자리
            input[i] += input[i + stride];
        }
        __syncthreads();
    }
    if (threadIdx.x == 0) *output = input[0];
}

이 첫 형태는 정확하다. 작은 입력 (size = 2048) 에서는 결과가 맞고, 강의의 메시지는 “이게 한참 부족한 출발점이다, 다음 단계들이 같은 일을 빠르게”.

두 번째 형태는 sequential addressing. stride 를 1, 2, 4, ... 로 키우는 대신 N/2, N/4, N/8 ... 로 줄인다. 결과 — active thread 들이 첫 절반에 모이고, 그 다음 첫 1/4 에 모이고, ... 한 warp 이 한 path 를 유지.

// control_divergence_reduce.cu — sequential addressing
__global__ void ConvergedSumReduction(float* input, float* output)
{
    unsigned int t = threadIdx.x;
    for (unsigned int stride = blockDim.x; stride > 0; stride /= 2) {
        if (t < stride) {                          // 첫 stride 개 thread 만 active
            input[t] += input[t + stride];
        }
        __syncthreads();
    }
    if (t == 0) *output = input[0];
}

이 변환의 효과 — warp execution efficiency 가 극적으로 좋아짐. NCU 가 직접 보여준다. 같은 알고리즘, 같은 정확도, 측정 가능한 시간 차이.

이전 두 버전은 모든 step 에서 input[t] += input[t + stride] — global memory 에 read/write. log(n) step × 2 = 2 log(n) HBM access per element. shared memory 로 옮기면 그 모두가 on-chip 으로 사라진다.

// shared_reduce.cu — 강의 repo 그대로
__global__ void SharedMemoryReduction(float* input, float* output)
{
    __shared__ float input_s[BLOCK_DIM];
    unsigned int t = threadIdx.x;
    // 첫 add 도 동시에 — 한 thread 가 두 element 를 한 번에
    input_s[t] = input[t] + input[t + BLOCK_DIM];
    for (unsigned int stride = blockDim.x / 2; stride >= 1; stride /= 2) {
        __syncthreads();
        if (t < stride) {
            input_s[t] += input_s[t + stride];           // ← shared, 빠름
        }
    }
    if (t == 0) *output = input_s[0];
}

이 일화의 메시지를 풀어 보면 — 변환의 효과는 입력 사이즈에 의존한다. 큰 입력 (수천만 element) 에서는 shared 로 옮기는 게 큰 차이를 만들지만, 작은 입력에서는 거의 차이 없을 수 있다. NCU 의 metric 이 좋아 보이는 것과 wall-clock 시간이 정렬되지 않는 자리.

L008 §06 의 thread coarsening 이 reduction 에 가장 자연스럽게 적용된다. 한 thread 가 시작할 때 4 또는 8 element 를 register 에 누적해 sum 을 만든 뒤, 그 sum 들을 shared 에 적고 tree reduction.

// reduce_coarsening.cu — coarsen factor 4 의 골격
__global__ void CoarsenedReduction(float* input, float* output, int n)
{
    __shared__ float input_s[BLOCK_DIM];
    unsigned int t = threadIdx.x;
    unsigned int base = blockIdx.x * blockDim.x * 4 + t;

    float sum = 0;                                    // register 에 누적
    sum += input[base + 0 * blockDim.x];
    sum += input[base + 1 * blockDim.x];
    sum += input[base + 2 * blockDim.x];
    sum += input[base + 3 * blockDim.x];
    input_s[t] = sum;

    for (unsigned int stride = blockDim.x / 2; stride >= 1; stride /= 2) {
        __syncthreads();
        if (t < stride) input_s[t] += input_s[t + stride];
    }
    if (t == 0) atomicAdd(output, input_s[0]);     // block 끝에 1번 atomic
}

강의의 multistream-reduce.cu 와 segment_reduce.cu 가 마지막 두 형태. multi-stream 은 여러 다른 reduction (예: per-row sum 의 row 별) 을 다른 CUDA stream 에서 동시에 띄움 — 한 reduction 의 launch overhead 를 다른 reduction 의 compute 와 overlap.

segment reduce 는 한 input 을 여러 segment 로 나누고 각 segment 의 sum 을 따로 — “per-row sum” 의 한 형태. PyTorch 의 tensor.sum(dim=…) 이 이 형태로 lower 됨.

강의의 마지막 한 페이지는 — “PyTorch 안에 min, max, sum 의 분리된 kernel 이 없다, 한 일반 reduction kernel 이 op 인자로 받음”. 이 사실의 의미가 흥미롭다. 같은 코드가 여러 op 에 reuse 된다는 추상의 가치, 그리고 — 커스텀 reduction 도 같은 path 위에 layered 가능.

강의에서 깐 가장 미묘한 자리. accuracy.py 와 sensitivity.py 가 보여주는 사실 — float 산술은 결합 법칙을 만족하지 않는다. (a + b) + c ≠ a + (b + c) 인 케이스가 있다. 그런데 tree reduction 은 시퀀셜과 다른 순서로 sum — 결과가 다를 수 있다.

accuracy.py 의 핵심 예 — 1000.0 (fp32) 에 0.001 을 1000번 더한다. fp32 의 mantissa 가 1000 근처에서는 0.001 을 표현하지 못함 — 그래서 더해도 1000 그대로. 즉 시퀀셜이면 답이 1000.0 인데 — tree 로 0.001 들을 먼저 모으면 1.0 이 되고 그게 1000.0 에 더해져 1001.0. 시퀀셜과 tree 의 답이 다르다.

# accuracy.py — 강의 repo (요약)
import torch
large_value = torch.tensor([1000.0], dtype=torch.float32)
small_values = torch.full((1000,), 0.001, dtype=torch.float32)

# 시퀀셜
res = large_value.clone()
for v in small_values:
    res += v
# res = ~1000.0  (작은 값들이 mantissa 에 쌓이지 못함)

# tree (PyTorch 의 sum)
res2 = large_value + small_values.sum()
# res2 = ~1001.0  (작은 값들이 먼저 모여 1.0 이 됨)

그리고 nondeterminism.py 가 보여주는 더 큰 함정 — 같은 input, 같은 코드, 다른 run 에서 결과가 다르다. 이유는 atomic 의 순서가 비결정. block 끝의 atomic add 가 어떤 순서로 일어나느냐가 GPU 의 scheduling 에 의존. “training 의 reproducibility 가 깨지는 자리”.

7개 reduction 의 진화와 각 단계의 핵심 한 줄.

손에 새기기 — 실습 시퀀스

1. 7개 .cu 모두 빌드 · 측정 — repo clone, 모든 reduction kernel 의 시간을 자기 GPU 에서. 강의의 진화가 자기 측정에서도 같은 방향으로 나오는가.
2. 입력 사이즈 sweep — size = 1024, 2048, 65536, 1M, 100M. 각 형태의 시간 변화. shared 의 효과가 작은 입력에서는 안 보일 수 있다.
3. warp execution efficiency 측정 — NCU 의 해당 카운터를 simple vs convergent 에 대해 본다. divergence 제거의 정량적 효과.
4. coarsen factor sweep — N = 1, 2, 4, 8, 16, 32. 시간과 register 사용량의 관계.
5. accuracy.py 직접 돌리기 — 시퀀셜 sum 과 PyTorch sum 의 차이를 확인. 1000 + 0.001×1000 의 케이스가 자기 환경에서 재현되는가.
6. nondeterminism.py 의 결과 — 같은 코드를 여러 번 돌려 결과가 매번 같은지 (또는 다른지) 확인. torch.use_deterministic_algorithms(True) 의 효과.
7. warp shuffle 추가 도전 — strict tree reduction 의 마지막 32 thread 단계를 __shfl_xor_sync 로 바꿔본다 (강의에서 본격 안 다루지만 NVIDIA whitepaper 의 표준).
8. 한 페이지 plan — 자기 모델의 한 reduction (예: layernorm 의 mean) 을 골라 7개 형태 중 어느 것이 적합한지 — bound 와 segment 수 기반 결정.

tree reduction · segment reduce · online reduction 의 패턴이 시리즈 거의 모든 LLM kernel 강의에서 다른 옷을 입고 등장.

강의에서 흐릿한 자리들과, 자기 GPU 에서 직접 봐야 손에 박힐 사실들.

• warp shuffle 의 본격 형태 — 강의에서 직접 다루지 않음. __shfl_xor_sync, __shfl_down_sync 가 마지막 32 thread reduction 의 표준. shared 와 sync 가 사라짐 → 빠름. NVIDIA whitepaper 또는 CUB 의 BlockReduce 코드 참조.
• shared_reduce 가 안 빨라지는 이유의 정확한 진단 — Mark 의 코드 주석이 그대로 남아 있는 자리. 입력 사이즈 변화로 효과가 보이는가? 아니면 본질적인 다른 이유가 있는가? (확인 필요)
• cooperative groups — block 사이 동기화의 새 메커니즘. cudaLaunchCooperativeKernel. 강의에서 안 다룸. 한 kernel 으로 큰 reduction 을 끝낼 수 있게 함.
• CUB 라이브러리 — production reduction 의 사실상 표준. cub::DeviceReduce::Sum. 강의에서 직접 인용 안 함. PyTorch 의 일반 reduction kernel 도 일부 CUB 사용.
• Hopper 의 distributed shared memory — block 끼리 직접 shared 를 share 가능. 큰 reduction 의 새 옵션.
• online softmax 와의 관계 — 이 강의의 “단순 reduction” 이 attention 의 online softmax 로 가는 길은 — running max 도 같이 추적. L012 FlashAttention 에서.
• fp16/bf16 reduction 의 정확도 — accuracy.py 가 fp32 의 케이스. fp16 에서는 mantissa 가 짧아 같은 함정이 더 일찍 발생. mixed precision training 에서의 reduction 은 fp32 accumulator 가 표준.