AMD ROCm: MI355X(CDNA 4)용 4웨이브 인터리브 FP8 GEMM 커널 최적화로 웨이브당 레지스터 예산 2배 확보
AMD ROCm 블로그에서 CDNA 4 아키텍처 기반 Instinct MI355X 가속기에 대한 FP8 GEMM(혼합 정밀도 행렬 곱셈) 커널의 새로운 4웨이브 인터리브 방식을 소개합니다. 이 최적화는 8웨이브 핑퐁 방식에서 웨이브당 512개 레지스터 전체 VGPR 예산을 갖는 4웨이브 방식으로 전환하며, XOR 기반 스위즐로 LDS 메모리 충돌을 제거하고 MFMA 명령어와 데이터 로딩의 정밀한 중첩으로 메모리 레이턴시를 숨깁니다.
이 기사는 AI가 1차 출처를 기반으로 생성했습니다.
AMD ROCm 팀이 CDNA 4 아키텍처 기반 AMD Instinct MI355X GPU의 FP8 GEMM(General Matrix Multiply) 커널에 대한 새로운 최적화 전략의 상세 기술 설명을 공개했습니다. 이 최적화는 표준 8웨이브 핑퐁 방식을 뛰어넘어, 웨이브당 레지스터 예산을 대폭 늘리고 하드웨어 수준에서 메모리 충돌을 제거하는 4웨이브 인터리브 설계를 도입합니다.
FP8 GEMM에서 4웨이브 방식이 8웨이브를 능가하는 이유
핵심 차이는 VGPR(Vector General Purpose Registers) 레지스터 할당에 있습니다. 8웨이브 설계에서는 각 웨이브가 레지스터 공간을 나눠 써야 합니다. 4웨이브 방식은 SIMD 유닛 하나당 웨이브 하나를 할당하여 512개 VGPR 전체 예산에 접근할 수 있습니다.
4웨이브 방식의 타일: 요소 기준 128×128(8웨이브의 64×128 대비). 전역 타일: 256×256×128. 네 웨이브 각각은 16×16 MFMA 연산을 사용해 64×64 블록 4개를 처리합니다.
4웨이브 인터리브가 MI355X에서 메모리 레이턴시를 제거하는 방법
4웨이브 방식은 하드웨어 스케줄러에 의존하지 않고 명령어의 정밀한 인터리빙으로 레이턴시를 해결합니다. 구현에는 __builtin_amdgcn_sched_barrier(0)를 사용합니다. 각 interleaved_block은 16개 MFMA 명령어와 8개 LDS-to-register 로드(ds_read_b128)를 중첩합니다.
XOR 기반 LDS 스위즐이 뱅크 충돌을 제거하는 방법
MI355X의 LDS는 64개 뱅크를 가집니다. MFMA 명령어는 컬럼 우선 순서로 접근하여 8배 뱅크 충돌을 일으킵니다. XOR 스위즐:
row_bits = (offset % (16 * 128) >> 7) / 2
mask = row_bits << 4
swizzled_offset = offset ^ mask결과: 12.5% 대신 LDS 전체 대역폭 활용.
하드웨어 플랫폼과 칩렛 인식 최적화
CDNA 4 아키텍처의 AMD Instinct MI355X, 8칩렛(XCD) 토폴로지. ROCm 팀은 모턴 커브 유사 매핑을 사용한 칩렛 인식 그리드 스위즐을 추가했습니다.
자주 묻는 질문
- 4웨이브 방식이 FP8 GEMM에서 8웨이브보다 유리한 이유는 무엇인가요?
- 4웨이브 방식은 각 웨이브가 SIMD 유닛 하나를 독점하여 512개 VGPR 레지스터 전체를 사용할 수 있습니다. 8웨이브 방식은 레지스터를 나눠야 하므로 각 웨이브의 예산이 절반으로 줄어듭니다.
- XOR 기반 LDS 스위즐은 어떤 문제를 해결하나요?
- MI355X의 LDS는 64개 뱅크를 가지며, MFMA 명령어는 컬럼 우선 순서로 접근해 최대 8배 뱅크 충돌이 발생합니다. XOR 스위즐은 오프셋을 재배치하여 충돌을 제거하고 LDS 전체 대역폭을 활용합니다.