Optimization Principles 와 Tensor Core Optimization 의 분리

Optimization Principles 의 내용을 정리하던 중

일반적인 내용들이 고전적 GEMM 최적화 흐름 기준,

Tensor Core 의 타일링 구조같은 경우 그 위에 다시 한 번 더 강하게 구조화된 별도의 계층으로

이 둘이 완전히 같은 것이 아님,

 

Tensor Core 는 기존 tiling 위에 warp-matrix tile 을 덧씌운 형태로 연결

 

전통적 GEMM tiling 구조

Global → Shared → Register → Thread-Tile → Warp-Tile

• thread 는 register 에 일부 Ctile 을 가짐
• warp 는 전체 Ctile 을 여러 thread 로 분배하고 협력
• shared memory 는 A/B tile 을 캐싱
  • SIMT 기반 FMA 최적화 구조

이것은 전통 커널 최적화 패턴

 

Tensor Core Tiling 은 Warp-MMA Tile 이라는 완전히 다른 계층 생성

Tensor Core 에서의 tiling 구조는 강제됨

Warp
 └─ MMA instruction tile (16×16×16 FP16 MMA)
     ├─ fragment A (8 registers/thread)
     ├─ fragment B (8 registers/thread)
     └─ fragment C (8 registers/thread)

warp 단위로 tile 크기가 고정

• 16, 16, 16
• thread 가 마음대로 tile 크기를 선택 불가능
• warp 전체가 하나의 MMA operation 을 수행

데이터 배치가 하드웨어에 의해 강제됨

• fragment 구조 
• 명령어가 요구하는 형태로 shared - register 로드
• swizzling 필요
• 레지스터 배치조차도 warp lane 에 따라 고정됨

shared memory tiling 도 Tensor Core aligment 기준에 맞춰 바뀜

• 공유 메모리 tile 패턴이 기존 FP32 패턴과 다름
• 128B aligment 강제
• cp.async + multi-stage pipleing 이 사실상 필수

register tile 은 C fragment 형태로 고정됨

고전적 register tiling 처럼

float acc[TM][TN];

와 같은 형태가 아닌

wmma::fragment<matrix_c, 16, 16, 16, half> c_frag;

이런 식으로 고정 크기 + 고정 분배 형태가 된다.

 

두 방식의 비교 방법

전통적 FP32/FP16 FMA 기반 Tiling

• tile 크기를 개발자가 직접 선택
• thread 에게 TM/TN 만큼의 C sub tile 배정
• warp tiling 도 임의 디자인 가능
• 레지스터 pressure / occupancy 균형을 자유롭게 조절
• shared memory layout 도 개발자가 결정

 

Tensor Core MMA Tiling

• tile 크기 고정
• fragment 구조 조겅
• warp-level tiling 강제
• 각 lane 이 어떤 matrix 조각을 담당할지도 결정되어 있음
• shared/register 배치 방법이 행렬 재배열 기반

 

최적화 원리는 Tensor Core 보다 훨씬 더 근본적인 개념

공유 메모리 타일링, 레지스터 타일링, coalescing, bank conflict, ILP, warp primitives ... 이런 것들은 Tensor Core 이전에도 존재, 

지금도 CUTLASS, FlashAttension, cuBLAS 내부에서도 그대로 사용

즉, 최적화 원리 Optimization Principles 는 Tensor Core 에 종속되지 않음

Tensor Core 는 이 원리 위에 특수한 제약과 기능을 추가한 구조

 

Tensor Core 최적화는 전통적 tiling/pipeline 의 연장선, 그 구조는 다름

Tensor Core 는 다음을 강제

• warp-level MMA tile 고정
• A/B fragment 의 lane mapping 강제
• shared memory swizzling 필수
• cp.async 기반 멀티 스테이지 pipeline
• warp 단위 협업 필수
• register 배치가 하드웨어가 정한 구조여야 함
• tile size 선택 불가능

위 요소들은 기본 최적화 원리와는 계층이 다르다.