[Triton] ConSan에서 barrier 다중 도착 시 false positive deadlock 감지 수정

PR 링크: triton-lang/triton#9046 상태: Merged | 변경: +126 / -3

들어가며

Triton의 Concurrency Sanitizer(ConSan)는 GPU 커널의 동시성 버그를 정적으로 감지한다. 하지만 barrier_expect를 무시하고, 각 tma.async_copy_global_to_shared를 독립적인 barrier arrive로 모델링하고 있었다. 이로 인해 두 개 이상의 TMA copy가 같은 barrier를 공유하면, barrier가 과다 도착(overarrival)한 것으로 오판하여 deadlock을 보고했다.

핵심 코드 분석

Before: 각 TMA copy가 barrier arrive로 카운트

if (auto copyOp = dyn_cast<ttng::AsyncTMACopyGlobalToLocalOp>(op)) {
  info.emplace();
  info->trackingKind = MemEffectsOpInfo::TrackingKind::Barrier;
  info->pred = copyOp.getPred();
  info->barriers.push_back({copyOp.getBarrier(), nullptr, 1}); // count=1
  info->operandEffects.emplace_back(MemEffectsOpInfo::Effects::Write,
                                    copyOp.getResult());
}

두 TMA copy가 같은 barrier를 사용하면, ConSan은 2번 arrive한 것으로 모델링하여 overarrival -> deadlock으로 오판한다.

After: barrier_expect를 arrive로, TMA copy는 count=0

if (auto expectOp = dyn_cast<ttng::BarrierExpectOp>(op)) {
  // barrier_expect가 arrive 역할을 함
  info.emplace();
  info->trackingKind = MemEffectsOpInfo::TrackingKind::Barrier;
  info->pred = expectOp.getPred();
  info->barriers.push_back({expectOp.getAlloc(), nullptr, /*count=*/1});
}

if (auto copyOp = dyn_cast<ttng::AsyncTMACopyGlobalToLocalOp>(op)) {
  info.emplace();
  info->trackingKind = MemEffectsOpInfo::TrackingKind::Barrier;
  info->pred = copyOp.getPred();
  // barrier state를 업데이트하지 않음 (see BarrierExpectOp handling)
  info->barriers.push_back({copyOp.getBarrier(), nullptr, /*count=*/0});
  info->operandEffects.emplace_back(MemEffectsOpInfo::Effects::Write,
                                    copyOp.getResult());
}

테스트: 2개 buffer + 1개 barrier

// Two TMA copies contribute to a single expected transaction.
ttng.barrier_expect %bar, 8192, %true
ttng.async_tma_copy_global_to_local %a[...] %a_smem, %bar, %true
ttng.async_tma_copy_global_to_local %b[...] %b_smem, %bar, %true
ttng.wait_barrier %bar, %c0_i32, %true  // 이제 deadlock으로 오탐하지 않음

왜 이게 좋은가

1. False positive 제거: 여러 TMA copy가 같은 barrier를 공유하는 합법적인 패턴에서 deadlock 오탐을 제거한다.
2. 하드웨어 모델에 근접: mbarrier.arrive.expect_tx에 대응하는 barrier_expect를 arrive로 모델링하여 실제 하드웨어 동작에 가까워졌다.
3. 실용적 타협: 전체 byte count 추적 대신 간단한 arrive 모델링으로 대부분의 패턴을 올바르게 처리한다.

정리

이 PR은 ConSan이 barrier_expect를 무시하던 문제를 수정하여, 여러 TMA copy가 같은 barrier를 공유하는 패턴에서 false positive deadlock 감지를 제거했다. barrier_expect를 arrive로, 개별 TMA copy는 count=0으로 모델링한다.

참고 자료

• triton-lang/triton#9046

---

이 글은 AI를 활용하여 PR의 핵심 변경사항을 분석하고 정리한 것입니다. 실제 코드의 맥락은 원본 PR을 참고해 주세요.