GPU 효율 동기화 원시 연산

초록

본 논문은 GPU에서 블록 간 동기화를 위한 장벽, 뮤텍스, 세마포어 구현을 재설계하고, Tesla와 Fermi GPU에서 원자 연산과 일반 메모리 접근의 비용 차이를 정량화한 뒤, 원자 연산 최소화와 스핀·슬립 전략 선택을 위한 설계 원칙을 제시한다.

상세 분석

논문은 먼저 CPU에서 널리 사용되는 동기화 원시 연산(장벽, 뮤텍스, 세마포어)의 구조와 한계를 정리하고, GPU 아키텍처가 갖는 특수성을 강조한다. GPU는 수천 개의 스레드가 워프 단위로 동시에 실행되며, 블록 간 통신은 전역 메모리와 원자 연산에 의존한다. 특히 초기 G80 시리즈는 원자 연산을 지원하지 않아 블록 간 동기화가 사실상 불가능했지만, Tesla(GT200)와 Fermi(GF100)에서는 원자 연산이 도입되었고 성능 차이가 크게 나타난다.

저자들은 12개의 마이크로벤치마크를 설계해 원자 읽기·쓰기와 일반(volatile) 읽기·쓰기의 지연을 측정했다. 결과는 두 GPU 간에 뚜렷한 차이를 보였다. Tesla에서는 원자 연산이 일반 메모리 접근보다 최대 90배 느렸으며, 캐시가 없어 모든 원자 연산이 DRAM을 직접 접근했다. 반면 Fermi는 L2 캐시 덕분에 원자 연산이 최대 3배 정도만 느렸고, 비경합(Non‑contentious) 접근이 경합(Contended) 접근보다 5~10배 빠른 것으로 나타났다. 또한 원자 연산 직후에 일반 메모리 접근을 수행하면 Fermi에서는 메모리 라인이 원자 유닛에 의해 직렬화되어 성능이 급격히 저하되는 현상이 관찰되었다.

이러한 실험 데이터를 바탕으로 저자들은 “GPU 메모리 시스템 추상화”를 제시한다. 핵심 파라미터는 (1) 원자 연산의 절대 비용, (2) 경합 정도에 따른 지연 비율, (3) 원자 연산 후 일반 메모리 접근이 직렬화되는 여부다. 이 추상화는 새로운 동기화 알고리즘 설계 시 원자 연산 횟수를 최소화하고, 가능한 경우 비원자(volatile) 연산으로 대체하도록 가이드한다.

구현 측면에서는 기존의 단순 스핀락을 개선해 두 단계 원자 장벽, 티켓 기반 뮤텍스, 그리고 카운터와 대기 큐를 결합한 세마포어를 제시한다. 각 원시 연산은 블록 단위로 동작하도록 설계돼, 블록 내 스레드들은 __syncthreads() 로 빠르게 협조하고, 블록 간 경쟁은 최소화한다. 특히 세마포어 구현에서는 원자 증가(atomicAdd)와 감소(atomicExch)를 각각 한 번만 사용하고, 이후의 대기/해제는 volatile 플래그를 폴링하거나, 일정 스핀 횟수 후에 “슬립”(스레드가 다른 블록에 양보하도록 커널 종료 후 재시작) 전략을 적용한다.

성능 평가에서는 동일한 워크로드를 Tesla와 Fermi에서 실행해, 새 구현이 기존 스핀락 대비 평균 2.5배~7배 빠름을 보였다. 특히 높은 경합 상황에서 티켓 뮤텍스와 두 단계 장벽이 원자 연산을 한 번만 수행하도록 설계돼, Fermi에서는 캐시 효율 덕분에 거의 선형 확장성을 달성했다. 반면 Tesla에서는 원자 연산 비용이 지배적이므로, 경합이 적은 경우에만 복잡한 알고리즘이 유리했다.

마지막으로 저자들은 “스핀 vs 슬립” 선택 가이드라인을 제시한다. GPU가 원자 연산을 빠르게 처리하고 캐시가 충분히 작동하는 경우(예: 최신 아키텍처)에는 스핀 대기 시간이 짧아 슬립 오버헤드보다 효율적이다. 반대로 원자 연산이 매우 비싼 구형 GPU에서는 일정 스핀 횟수 후 슬립(커널 재시작)으로 전환해 전체 실행 시간을 크게 단축할 수 있다. 이러한 원칙은 향후 GPU 아키텍처가 어떻게 진화하든 적용 가능한 일반적인 설계 지침을 제공한다.