[Triton] AMD RDNA3에서 buffer cache modifier LLVM IR 전파
PR 링크: triton-lang/triton#9777 상태: Merged | 변경: +171 / -7
들어가며
Buffer cache modifier(.cg, .cs, .cv, .wt)는 GPU 커널에서 캐시 동작을 제어하는 힌트다. 예를 들어 streaming 데이터에 .cs(non-temporal)를 지정하면 캐시 오염을 방지한다. Triton의 AMD 백엔드에서 CDNA3/4에는 이미 구현되어 있었지만, RDNA3에서는 aux=0으로 항상 기본값이 사용되어 cache modifier가 무시되고 있었다.
핵심 코드 분석
Before
RDNA3 GPU에서 cache modifier 지정 시, CPol 비트가 설정되지 않아 항상 aux=0(기본 캐시 동작)으로 출력되었다.
After (RDNA3 매핑 테이블)
Load .cg → aux=1 (GLC: bypass GL1)
Load .cs → aux=7 (GLC|SLC|DLC: non-temporal, skip MALL)
Load .cv → aux=7 (GLC|SLC|DLC: non-temporal, skip MALL)
Store .cs → aux=7 (GLC|SLC|DLC: non-temporal, skip MALL)
Store .wt → aux=7 (GLC|SLC|DLC: non-temporal, skip MALL)
After (dispatch 방식 변경)
// Before: 개별 GPU 모델(GK_GFX*)로 분기
// After: ISAFamily 단위로 분기하여 자동 커버리지 확보
왜 이게 좋은가
- 성능 잠재력 해방: bandwidth-bound 워크로드에서 non-temporal 힌트는 캐시 오염을 방지하여 유효 대역폭을 높인다.
- ISAFamily 기반 dispatch: 새 칩이 추가되어도 ISAFamily만 맞으면 자동으로 올바른 cache modifier가 적용된다.
- MLIR FileCheck 테스트: gfx1150을 대표 타겟으로 load/store 각 modifier에 대한 correctness 테스트가 포함되었다.
정리
Cache modifier가 올바르게 전파되지 않으면, 커널 작성자가 의도한 캐시 전략이 하드웨어에 도달하지 못한다. 이 PR은 RDNA3 아키텍처의 GLC/SLC/DLC 비트 매핑을 추가하고, GPU 모델 대신 ISAFamily 단위 dispatch로 유지보수성을 높였다.
참고 자료
이 글은 AI 도구의 도움을 받아 작성되었습니다.
관련 포스트
- [sglang] SGLang: AMD GPU 환경에서의 DeepSeek-V4 성능 최적화 분석
- [sglang] AMD GPU 최적화: Triton 커널 퓨전을 통한 Qwen2 MoE 공유 전문가 게이팅 성능 향상
- [triton] Triton AMD StreamK GEMM 커널의 Race Condition 해결: 동기화 로직 최적화 분석
- [triton] [AMD Triton] LLVM InstCombine의 함정을 피하는 법: TDM 텐서 클램핑 최적화
- [triton] Triton AMD 커널 최적화: TDM 로드 파이프라이닝 개선을 통한 성능 향상
PR Analysis 의 다른글
- 이전글 [triton] Global Sanitizer에 TMA 및 cp.async 연산 부분 지원 추가
- 현재글 : [Triton] AMD RDNA3에서 buffer cache modifier LLVM IR 전파
- 다음글 [ray] Ray Data의 차세대 데이터 소스 API: DataSourceV2 설계 및 최적화 전략
댓글