제2부 CUDA 프로그래밍

2.8 동기화 I — barrier·atomic

블록 안 수천 스레드가 같은 데이터를 건드릴 때 "누가 언제 읽고 쓰는가"를 다스리지 못하면 결과는 비결정적으로 깨진다. 이 장은 블록 단위 협력의 두 도구 — 블록 barrier(__syncthreads) 와 아토믹(atomic) — 을 다룬다. 워프 내부 통신·협력 그룹은 2.9으로 잇는다.

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

배우는 것:

• 협력의 세 층위(워프·블록·그리드)와 비용을 안다 — 좁을수록 비용이 낮다.
• __syncthreads의 정확한 의미론과 조건분기 안에 넣으면 안 되는 이유(데드락).
• 글로벌·공유 메모리 아토믹의 동작과 충돌(경합) 비용, 히스토그램 같은 경합 패턴 설계(privatization).

선수 지식: 1.2(SIMT·워프), 1.7(공유 메모리·뱅크), 2.3(그리드·블록·워프 설계), 2.6~2.7(공유 메모리·타일링). 워프 프리미티브는 2.9.

2.8.2 왜 동기화가 필요한가 — 협력의 세 층위

빠른 커널은 거의 항상 스레드끼리 데이터를 나눠 협력한다. 협력에는 두 위험이 따른다: 경쟁 상태(race)(A가 쓴 값을 B가 먼저 읽음)와 가시성(visibility)(A의 쓰기가 B 시점에 아직 안 보임). 도구는 협력 범위에 따라 비용이 다르다 — 좁을수록 비용이 낮다.

핵심 요지: 워프 안 통신은 사실상 무비용(레지스터 셔플), 블록 안은 공유 메모리 + barrier 한 번, 블록을 넘는 협력은 비용이 커서 보통 커널 경계로 처리한다. 좋은 커널 설계는 협력을 가능한 한 좁은 범위로 끌어내린다.

2.8.3 블록 barrier — __syncthreads()

__syncthreads()는 블록 안 모든 스레드가 이 지점에 도달할 때까지 대기하는 배리어다. 동시에 메모리 펜스 역할도 한다: 배리어를 통과한 뒤에는 배리어 이전에 블록 내 다른 스레드가 공유·글로벌에 쓴 값이 보인다. 타일링의 전형 골격(2.6~2.7)이 이것이다.

__shared__ float tile[TILE][TILE];
tile[ty][tx] = global[...];   // 1) 각 스레드가 타일에 채우고
__syncthreads();              // 2) 전원이 다 채울 때까지 대기 (+ 가시성)
float v = tile[ty][other_tx]; // 3) 이제 tile 전체를 안전하게 읽음
__syncthreads();              // 4) 다음 타일로 덮어쓰기 전 다시 대기

의미론의 핵심 — 모든 스레드가 같은 배리어에 도달해야 한다. 일부 스레드만 도달하면 나머지는 영원히 기다린다(데드락). 그래서 스레드마다 도달 여부가 달라지는 조건분기 안에 넣으면 안 된다(2.8.6 함정).

// 잘못 — 데드락 위험: n 이상 레인은 이 배리어에 도달 안 함
if (threadIdx.x < n) { ...; __syncthreads(); }

// 올바름 — 배리어는 분기 밖, 전원이 도달
if (threadIdx.x < n) { ... }
__syncthreads();

결과를 받는 변형도 있다: __syncthreads_count(pred)(조건 만족 스레드 수), __syncthreads_and/or(pred)(블록 전체 AND·OR). 동기화와 블록 단위 reduce를 한 번에 한다.

2.8.4 아토믹 연산 — 경합과 비용

여러 스레드가 같은 메모리 위치를 갱신할 때(카운터·히스토그램·누산기), 읽고-더하고-쓰는 과정이 개입하면 갱신이 유실된다. 아토믹은 이 read-modify-write를 분할 불가능하게 만든다.

atomicAdd(&counter, 1);          // 정수·부동소수 합
atomicMax(&maxVal, x);           // 최대
old = atomicCAS(&lock, 0, 1);    // compare-and-swap (락·자유 자료구조의 토대)
atomicExch/Sub/And/Or/Xor/Min(...);

2.8.4.1 충돌 비용 — 직렬화가 본질

아토믹 성능은 얼마나 같은 주소로 몰리느냐(경합) 가 좌우한다. 같은 주소에 동시 갱신이 몰리면 하드웨어가 직렬화한다 — 32 레인이 같은 카운터를 증가시키면 사실상 차례로 처리된다.

위치 차이도 크다. 공유 메모리 아토믹은 L2를 거치지 않고 SM 안에서 처리돼 글로벌 아토믹보다 일반적으로 훨씬 비용이 낮다(아키텍처 의존). 그래서 히스토그램의 표준 패턴은: 블록마다 공유 메모리에 부분 히스토그램을 아토믹으로 쌓고, 끝에 한 번만 글로벌로 합친다(privatization).

// 경합 완화: privatization(블록별 부분 결과)
__shared__ int hist[NBINS];
for (int i = threadIdx.x; i < NBINS; i += blockDim.x) hist[i] = 0;
__syncthreads();
for (int i = gid; i < N; i += stride)
    atomicAdd(&hist[data[i]], 1);          // 공유 메모리 아토믹 (비용이 낮다)
__syncthreads();
for (int i = threadIdx.x; i < NBINS; i += blockDim.x)
    atomicAdd(&global_hist[i], hist[i]);   // 글로벌은 블록당 NBINS회만

글로벌 아토믹을 N회에서 (블록 수)×NBINS회로 줄였고, 공유 메모리 아토믹은 SM 내부 처리라 더 비용이 낮다.

정직성: atomicAdd의 float·double·FP16·BF16·float2 지원과 공유 메모리 아토믹의 성능 특성은 아키텍처 세대마다 다르다. 구체 가속 수치는 Nsight Compute(2.14)로 직접 측정하라.

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

• [ ] 협력 세 층위(워프·블록·그리드)와 "좁을수록 비용이 낮다"를 설명한다.
• [ ] __syncthreads가 블록 전원이 같은 위치에 도달해야 함을, 조건분기 안에 넣으면 데드락임을 안다.
• [ ] 아토믹 비용이 경합(같은 주소 직렬화) 이 좌우함을 알고, privatization으로 줄인다.

2.8.6 흔한 함정

• 조건분기 안의 __syncthreads. 블록 일부만 도달하면 데드락. 배리어는 모든 스레드가 똑같이 도달하는 위치에. 데이터에 따라 갈리는 if/while 안에 넣지 말 것. compute-sanitizer --tool synccheck로 배리어 발산(barrier divergence)을 검출한다(공유메모리 race는 racecheck, 2.6·2.7).
• 고경합 아토믹을 그대로 둠. 모든 스레드가 한 주소를 두드리면 직렬화로 느려진다. privatization·워프 단위 사전 집계(2.9)로 경합을 줄여라.
• 아토믹 타입 가정. float/double/저정밀 지원은 세대 의존. 지원 여부와 성능을 확인.

2.8.7 연습문제

연습 1 (데드락 만들고 고치기). 조건분기 안에 __syncthreads를 넣어 행이 걸리는 커널을 작성한 뒤 배리어를 분기 밖으로 빼 고쳐라. 무엇이 행을 풀었나?

연습 2 (공유 메모리 히스토그램). unsigned char data[N](0~255)의 256-bin 히스토그램을 작성하라. 블록별 공유 메모리 부분 히스토그램을 atomicAdd로 쌓고 끝에 한 번만 글로벌로 합쳐라.

__global__ void hist256(const unsigned char* data, int N, int* g){
  __shared__ int h[256];
  for(int i=threadIdx.x;i<256;i+=blockDim.x) h[i]=0;
  __syncthreads();
  for(int i=blockIdx.x*blockDim.x+threadIdx.x;i<N;i+=gridDim.x*blockDim.x)
    atomicAdd(&h[data[i]],1);          // 공유 아토믹 (비용이 낮다)
  __syncthreads();
  for(int i=threadIdx.x;i<256;i+=blockDim.x) atomicAdd(&g[i], h[i]); // 글로벌 블록당 256회
}

연습 3 (경합 측정). 모든 스레드가 한 주소에 atomicAdd하는 커널과 흩어진 주소에 하는 커널의 시간을 비교하라. 왜 다른가?

2.8.8 요약 · 다음 장

• 협력은 범위가 좁을수록 비용이 낮다: 워프 → 블록(공유 + __syncthreads) → 그리드(아토믹·펜스).
• __syncthreads는 블록 전원이 같은 위치에 도달해야 한다 — 조건분기 안은 데드락.
• 아토믹 비용은 경합(같은 주소 직렬화) 이 좌우한다. privatization으로 줄인다.

다음 장: 2.9 동기화 II — 워프 shuffle·vote·협력 그룹. 공유 메모리 없이 레지스터끼리 직접 통신하는 워프 프리미티브와, 협력 범위를 1급 객체로 만드는 협력 그룹을 본다.