[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 도구의 도움을 받아 작성되었습니다.
관련 포스트
- [triton] AMD Canonicalize Pointers에서 arith.select의 비대칭 fat pointer 처리 강화
- [Triton] AMD gfx1250 Tensor Descriptor 기반 GEMM 테스트 추가
- [triton] AMD TDM의 Partition-Aware 분할 및 다중 Intrinsic 지원
- [triton] AMD GFX9 Async Copy에서 Shared Memory 순서 버그 수정
- [triton] AMD Async Wait Count에서 Warp Free Variable 및 Register Zero Base 버그 수정
PR Analysis 의 다른글
- 이전글 [triton] Global Sanitizer에 TMA 및 cp.async 연산 부분 지원 추가
- 현재글 : [Triton] AMD RDNA3에서 buffer cache modifier LLVM IR 전파
- 다음글 [ray] Ray Data의 차세대 데이터 소스 API: DataSourceV2 설계 및 최적화 전략
댓글