제2부 CUDA 프로그래밍

2.9 동기화 II — 워프 shuffle·vote·협력 그룹

2.9.1 이 장에서 배우는 것 / 선수 지식

2.8에서 블록 barrier와 아토믹을 봤다. 이 장은 한 단계 더 좁은 협력 — 워프 안 레인끼리 공유 메모리 없이 직접 통신하는 셔플·투표·reduce, 그리고 협력 범위를 1급 객체로 만드는 협력 그룹을 다룬다. 마지막엔 같은 리덕션을 세 방식(셔플·공유메모리·아토믹)으로 구현해 트레이드오프를 본다.

배우는 것:

• 워프 셔플(__shfl_*_sync)·투표(__ballot_sync 등)·reduce로 공유 메모리 없이 워프 내부 협력을 구현한다.
• __syncwarp와 독립 스레드 스케줄링 때문에 _sync 프리미티브가 필요한 이유.
• 협력 그룹으로 워프·타일·블록·그리드 협력을 타입 안전하게 표현한다.
• 메모리 펜스와 그리드 전역 동기화의 한계.

선수 지식: 1.2(SIMT·워프), 2.8(barrier·atomic). 워프가 32 레인 SIMT로 한 명령을 함께 실행한다는 것을 안다고 가정.

2.9.2 __syncwarp와 독립 스레드 스케줄링

Volta 이후 GPU는 독립 스레드 스케줄링(Independent Thread Scheduling) 을 도입하면서, 워프 안 레인들이 항상 같은 명령을 락스텝으로 실행한다는 옛 가정이 깨졌다. 다이버전트 분기 후 자동 재수렴 보장이 약해졌고, 그래서 워프 내부에서 레인끼리 데이터를 교환하는 모든 프리미티브는 명시적 mask를 받는 _sync 버전을 써야 한다(__shfl이 아니라 __shfl_sync).

__syncwarp(mask)는 mask로 지정한 레인들을 재수렴·동기화한다(전체 참여면 0xffffffff).

unsigned full = 0xffffffff;        // 32 레인 모두 참여
val = __shfl_down_sync(full, val, 16);
__syncwarp(full);                  // (필요 시) 재수렴 보장

정직성: _sync 프리미티브 자체가 mask 내 레인 동기화를 포함하므로, 잘 작성한 셔플 리덕션에서는 매 단계 __syncwarp를 따로 부를 필요가 없는 경우가 많다. 그러나 공유 메모리를 워프 내부 통신에 쓸 때(레인 A가 쓰고 레인 B가 읽는)는 __syncwarp가 반드시 필요하다 — 재수렴이 보장되지 않기 때문. 폐기된 volatile 트릭에 의존하지 말 것.

2.9.3 워프 수준 프리미티브 — 레인끼리 직접 말하기

워프 프리미티브는 공유 메모리를 거치지 않고 레지스터끼리 직접 레인 간 데이터를 옮기거나 집계한다. 비용이 거의 없고 배리어도 필요 없어, 리덕션·스캔·정렬·압축의 핵심 빌딩블록이다.

2.9.3.1 셔플 — __shfl_*_sync

__device__ __forceinline__ float warpReduceSum(float val) {
    unsigned mask = 0xffffffff;
    for (int offset = 16; offset > 0; offset >>= 1)
        val += __shfl_xor_sync(mask, val, offset);
    return val;  // 워프의 모든 레인이 총합을 가짐 (xor 버터플라이)
}
// __shfl_down_sync 로 바꾸면 lane 0에만 결과가 모인다

5단계(log₂32)로 끝나며, 공유 메모리 0바이트·배리어 0회다.

2.9.3.2 투표 — __ballot_sync·__any_sync·__all_sync

__ballot_sync는 워프 단위 압축(stream compaction)의 핵심: "내가 출력할 값을 가졌나?"를 투표하면 __popc(1의 개수)로 출력 개수를, 내 비트 아래쪽 __popc로 워프 내 prefix 위치를 한 번에 얻는다.

unsigned active = __activemask();
unsigned vote = __ballot_sync(active, keep);     // keep인 레인들의 비트
int warpCount = __popc(vote);                     // 이 워프가 내보낼 개수
int myRank    = __popc(vote & ((1u<<laneId)-1));  // 내 워프 내 순위

2.9.3.3 워프 reduce (__reduce_*_sync)

CUDA는 셔플 루프 대신 하드웨어 가속 워프 reduce를 제공한다(__reduce_add_sync 등). 정수형에 대해 Ampere 이상(SM 8.0+)에서 단일 명령으로 워프 합·최소·최대를 낸다(시점·아키텍처 의존). 부동소수 reduce·scan은 협력 그룹의 cg::reduce가 이식성이 좋다. 직접 작성한 셔플 루프는 모든 아키텍처에서 동작하므로 학습·이식성 기준값으로 좋다.

2.9.4 협력 그룹(Cooperative Groups)

__syncthreads·__shfl_sync는 "블록 전체"·"워프 32레인"이라는 고정 범위에 묶여 있다. 협력 그룹(CG)은 협력하는 스레드 집합을 명시적 객체로 만들어, 워프보다 작은 타일·블록·그리드까지 같은 인터페이스로 동기화·집계한다.

#include <cooperative_groups.h>
#include <cooperative_groups/reduce.h>
namespace cg = cooperative_groups;
__global__ void kernel() {
    cg::thread_block block = cg::this_thread_block();
    block.sync();                                   // == __syncthreads()
    auto tile = cg::tiled_partition<32>(block);      // 워프 또는 더 작은 타일
    float v = tile.shfl_down(myVal, 16);             // == __shfl_down_sync
    unsigned vote = tile.ballot(pred);               // == __ballot_sync
    float sum = cg::reduce(tile, myVal, cg::plus<float>());  // 타입 안전, 이식성
}

주요 그룹 타입(API는 CUDA 버전에 따라 진화):

coalesced_group은 다이버전트 코드에서 "지금 함께 실행 중인 레인만" 으로 그룹을 만들어 안전하게 집계한다 — if 안에서 워프 일부만 활성일 때 아토믹 경합을 줄이는 정석:

if (predicate) {
    cg::coalesced_group g = cg::coalesced_threads();
    int total = cg::reduce(g, 1, cg::plus<int>());
    if (g.thread_rank() == 0) atomicAdd(counter, total);  // 워프당 아토믹 1회
}

2.9.5 메모리 펜스와 그리드 동기화 한계

2.9.5.1 메모리 펜스 — 가시성 vs 동기화

배리어(__syncthreads)는 "도달 동기화 + 가시성"을 함께 준다. 때로 순서·가시성만 필요하고 대기는 불필요 — 이때 펜스를 쓴다.

전형 용도: 한 블록이 결과를 글로벌에 쓰고 "완료" 플래그를 세울 때, 데이터 쓰기가 플래그 쓰기보다 먼저 보이도록 __threadfence()를 데이터와 플래그 사이에. 없으면 다른 블록이 "완료"는 봤는데 데이터는 아직 못 보는 상태가 생긴다. CUDA의 메모리 모델은 약한 일관성이라 명시적 펜스·아토믹 없이는 블록 간 쓰기 순서를 가정할 수 없다(현대적으로는 cuda::atomic·cuda::memory_order, libcu++).

참고 — cuda::pipeline·cuda::barrier(범용 비동기 배리어, CUDA 11+ libcu++). 이 장의 __syncthreads·펜스는 직접 작성한 동기화다. cp.async(비동기 글로벌→공유메모리 복사, 3.4·4.3에서 CUTLASS 문맥으로 다룸)를 표준 C++ 스타일 API로 감싼 것이 cuda::pipeline이고, 짝을 이루는 cuda::barrier는 "몇 스레드가 도달했나 + 비동기 복사가 끝났나"를 함께 추적하는 배리어다. 복사와 계산을 겹치는(overlap) 손짜기 패턴(2.10의 스트림 오버랩과 같은 목적을 블록 내부에서)을 라이브러리 코드 없이 직접 쓸 때 등장한다 — CUTLASS 같은 라이브러리가 내부적으로 쓰는 것과 같은 계열의 도구를 애플리케이션 코드에서 직접 쓰는 경우다. 이 책은 개념만 짚는다: 직접 작성할 일은 드물고(대개 CUTLASS·프레임워크가 대신 해줌), API는 CUDA 버전 의존이 크다.

2.9.5.2 그리드 전역 동기화의 한계

__syncthreads는 블록 안만 동기화한다. "그리드 전체가 한 지점에서 만나기"는 기본적으로 불가능에 가깝다 — 블록들이 SM에 시차를 두고 스케줄되기 때문. 정석은 커널을 끊는 것(단계 1 커널 → 단계 2 커널, 커널 경계가 자연스러운 그리드 배리어).

협력 그룹의 grid_group::sync()는 한 커널 안에서 가능케 하지만 제약이 따른다: cudaLaunchCooperativeKernel로 실행, 모든 블록이 동시 상주해야 함(그리드 크기가 점유율로 묶임), 디바이스·드라이버 지원 확인 필요.

견해: 그리드 동기화는 단계 간 데이터를 재사용하는 특수한 경우(영속 커널·일부 그래프 알고리즘)에 가치가 있다. 일반 다단계 파이프라인은 커널을 끊거나 CUDA Graphs(2.11) 로 런치 오버헤드를 줄이는 편이 단순·이식적이다. 멀티 GPU 전역 동기는 NCCL 집합연산(4.13)이 현실적 도구.

2.9.6 실습 — 같은 리덕션 세 가지 방식

큰 배열 합을 (A) 워프 셔플, (B) 공유 메모리, (C) 글로벌 아토믹으로 구현해 트레이드오프를 본다.

// reduce.cu — nvcc -O3 -arch=sm_80 reduce.cu -o reduce
__device__ __forceinline__ float warpReduceSum(float v) {
    for (int o = 16; o > 0; o >>= 1) v += __shfl_down_sync(0xffffffff, v, o);
    return v;
}
// (A) 워프 셔플 + 블록 합 (공유메모리 최소, 배리어 1회)
__global__ void reduceShuffle(const float* x, float* out, int N) {
    float sum = 0.f;
    for (int i = blockIdx.x*blockDim.x+threadIdx.x; i < N; i += gridDim.x*blockDim.x) sum += x[i];
    sum = warpReduceSum(sum);
    __shared__ float warpSums[32];
    int lane = threadIdx.x & 31, wid = threadIdx.x >> 5;
    if (lane == 0) warpSums[wid] = sum;
    __syncthreads();
    if (wid == 0) {
        sum = (threadIdx.x < (blockDim.x+31)/32) ? warpSums[lane] : 0.f;
        sum = warpReduceSum(sum);
        if (lane == 0) atomicAdd(out, sum);    // 블록당 글로벌 아토믹 1회
    }
}
// (B) 공유 메모리 트리 리덕션 (배리어 log2(blockDim)회)
__global__ void reduceShared(const float* x, float* out, int N) {
    extern __shared__ float s[]; int tid = threadIdx.x; float sum = 0.f;
    for (int i = blockIdx.x*blockDim.x+tid; i < N; i += gridDim.x*blockDim.x) sum += x[i];
    s[tid] = sum; __syncthreads();
    for (int stride = blockDim.x>>1; stride > 0; stride >>= 1) {
        if (tid < stride) s[tid] += s[tid+stride];
        __syncthreads();
    }
    if (tid == 0) atomicAdd(out, s[0]);
}
// (C) 글로벌 아토믹만 — 모든 스레드가 한 주소로 (최악의 경합)
__global__ void reduceAtomic(const float* x, float* out, int N) {
    for (int i = blockIdx.x*blockDim.x+threadIdx.x; i < N; i += gridDim.x*blockDim.x)
        atomicAdd(out, x[i]);
}

예상(개략 — 정확한 시간은 GPU·N 의존): shuffle 가장 빠름, shared 비슷~약간 느림, atomic 크게 느림(16M 갱신이 한 주소에서 직렬화). 세 합은 같아야 한다(float 누적 미세 오차 가능).

견해: 단순 리덕션은 (A) 셔플(또는 CUB DeviceReduce/BlockReduce)을 기본값으로. 실전에서는 직접 작성하기보다 CUB·Thrust(2.15)의 검증된 리덕션이 가장 빠르고 안전하다 — 위 코드는 동작 원리를 익히는 학습용.

2.9.7 중간 점검 (체크포인트)

• [ ] __shfl_down_sync로 32레인 워프 합을 5단계에 접는 루프를 외워 쓴다.
• [ ] "왜 셔플이 공유 메모리보다 비용이 낮은가"를 한 문장으로(공유·배리어를 거치지 않고 레지스터끼리 직접).
• [ ] Volta+ 워프 프리미티브는 _sync + 정확한 mask여야 함을 안다.
• [ ] 협력 그룹의 coalesced_threads가 다이버전스에서 활성 레인만 안전하게 묶음을 안다.

2.9.8 흔한 함정

• 다이버전트 코드에서 셔플 mask 오용. if 안에서 워프 일부만 활성인데 0xffffffff를 주면 비활성 레인에서 읽으려 해 미정의 동작. __activemask()/coalesced_threads()로 정확한 마스크를.
• _sync 없는 옛 프리미티브. __shfl·__ballot(mask 없는 형)은 Volta+에서 안전하지 않다(deprecated). 반드시 _sync.
• 공유 메모리 워프 통신에 __syncwarp 누락. 레인 A가 쓰고 B가 읽을 때 __syncwarp 없으면 옛 값을 읽을 수 있다. volatile 트릭 금지. compute-sanitizer --tool racecheck로 이런 워프 내 공유메모리 race도 검출된다(2.6·2.7).
• 펜스 없이 블록 간 플래그 통신. 데이터 쓰기와 플래그 쓰기 사이 __threadfence() 없으면 다른 블록이 플래그는 보고 데이터는 못 본다.
• 그리드 동기화를 일반 런치로 시도. grid.sync()는 cudaLaunchCooperativeKernel + 동시 상주 제약 없이는 동작 안 함.

2.9.9 연습문제

연습 1 (셔플 최대). warpReduceSum을 본떠 warpReduceMax(float)를 작성하라. __shfl_xor_sync로 32레인 모두가 최대를 갖게.

__device__ float warpReduceMax(float v){
  for(int o=16;o>0;o>>=1) v = fmaxf(v, __shfl_xor_sync(0xffffffff,v,o));
  return v;   // xor 버터플라이 → 모든 레인이 동일 최대
}

연습 2 (coalesced 집계). 술어 keep가 참인 원소만 카운트하되, 워프 안 활성 레인을 coalesced_threads()로 모아 워프당 글로벌 아토믹 1회로 합쳐라.

if (keep) {
  auto g = cg::coalesced_threads();
  int total = cg::reduce(g, 1, cg::plus<int>());
  if (g.thread_rank()==0) atomicAdd(counter, total);
}

연습 3 (펜스 추론). 한 블록이 result[bid]를 쓴 뒤 ready[bid]=1로 표시하고, 다른 블록이 ready를 폴링해 result를 읽는다. 어디에 어떤 펜스가 필요하며 없으면 무슨 일이?

2.9.10 요약 · 다음 장

• 워프 셔플·투표는 공유 메모리·배리어 없이 레지스터끼리 직접 통신 — 리덕션·스캔·압축의 핵심. Volta+에선 _sync + 정확한 mask 필수.
• 협력 그룹은 협력 범위를 1급 객체로 만들어 워프·타일·블록·그리드를 같은 인터페이스로, coalesced_threads로 다이버전스를 안전하게.
• 같은 리덕션도 셔플 ≪ 공유 ≪ 단일주소 아토믹 순으로 비용이 낮다. 그리드 전역 동기화는 제약이 많아 커널 경계나 CUDA Graphs로 대체. 실무는 CUB·Thrust(2.15)를 먼저.

다음 장: 2.10 스트림·이벤트·비동기 오버랩. 단일 커널 내부 협력에서 한 단계 올라가, 커널·복사·CPU 작업을 시간축에서 겹치는 동시성을 다룬다.