[triton] GSan 테스트에서 nanosleep 대신 Atomic 기반 동기화로 전환

PR 링크: triton-lang/triton#9818 상태: Merged | 변경: +37 / -22

들어가며

GPU Concurrency Sanitizer(GSan)는 커널 간의 데이터 레이스를 감지하는 도구입니다. GSan의 테스트는 의도적으로 데이터 레이스를 발생시켜 감지 여부를 검증하는데, 기존에는 nanosleep(500_000)으로 스레드 실행 순서를 제어했습니다. 이 방식은 GPU 클럭 속도나 부하에 따라 비결정적이어서 테스트가 불안정(flaky)했습니다.

핵심 코드 분석

Before (nanosleep 기반):

@triton.jit
def _raw_kernel(ptr, scratch_ptr):
    pid = tl.program_id(0)
    if pid == 0:
        tl.store(ptr, 1)
    else:
        nanosleep(500_000)  # 500us 대기 -- 불안정
        value = tl.load(ptr)
        tl.store(scratch_ptr, value)

After (atomic polling 기반):

@triton.jit
def atomic_poll(counter_ptr, expected):
    while tl.atomic_add(counter_ptr, 0, sem="relaxed") < expected:
        nanosleep(100)

@triton.jit
def _raw_kernel(ptr, scratch_ptr, counter_ptr):
    pid = tl.program_id(0)
    if pid == 0:
        tl.store(ptr, 1)
        tl.atomic_add(counter_ptr, 1, sem="relaxed")  # 신호
    else:
        atomic_poll(counter_ptr, 1)  # 대기
        value = tl.load(ptr)
        tl.store(scratch_ptr, value)

왜 이게 좋은가

nanosleep(500_000)은 (1) GPU 클럭에 따라 실제 대기 시간이 다르고, (2) 500us가 부족하면 순서가 뒤집혀 의도한 레이스가 발생하지 않고, (3) 과도하면 테스트가 느려집니다. Atomic polling은 이벤트 기반 동기화로, pid 0이 작업을 완료하고 counter를 증가시키면 pid 1이 즉시 진행합니다. sem="relaxed"를 사용하여 메모리 순서 의미론은 완화하면서도 값 전달은 보장합니다. RAW/WAR/WAW 모든 테스트 패턴에 동일하게 적용되어 일관성이 높습니다.

정리

GSan의 RAW/WAR/WAW 및 TMA 관련 테스트 7개에서 nanosleep 기반 동기화를 atomic counter polling으로 교체하여 테스트 안정성을 향상시켰습니다.

참고 자료

• triton-lang/triton#9818

이 분석은 AI가 실제 코드 diff를 기반으로 작성했습니다.