Fragment layouy visualize test

• 타일 크기 m, n, k  = 16
• warp 1개, 32 lanes 가 16 x 16 타일의 담당
• 커널 동작
  • matrix_a / matrix_b
  • wmma::load_matrix_sync 로 fragment 로드
  • 각 lane 이 frag.x[i] 를 읽어서 row, col 복원
  • accumulator
  • 각 lane 이 기록
  • wmma::store_matrix_sync 로 저장, 
  • 값 복원

 

Accumulator fragment (C) - warp 타일 구조

=== [accumulator]: raw C (value = lane*100 + frag_idx) ===
...
=== [accumulator]: decoded mapping (row,col,lane,frag_idx) ===
( 0, 0): lane  0, frag_idx  0
( 0, 1): lane  0, frag_idx  1
( 0, 2): lane  1, frag_idx  0
( 0, 3): lane  1, frag_idx  1
...
( 8, 0): lane  0, frag_idx  2
( 8, 1): lane  0, frag_idx  3
...
( 8, 8): lane  0, frag_idx  6
( 8, 9): lane  0, frag_idx  7
...

2.1 lane - (row, col) 분담

각 lane 당 8 개 요소 보유, warp 1개가 16, 16 타일을 8, 8 sub-tile 4 개로 쪼개고, 각 lane 은 8 element를 맡는다.

 

B 는 A 와 동일한 warp 타일 구조를 공유하지만, (K, N) 좌푝에서 col-major 로 배치된 걸 16, 16 으로 투영한 

0, 1, 800, 801, 8, 9, 808, 809, ...

0, 1, 8, 9, 800, 801, 808, 809, ...

A, B 의 lane 0 dump 출력, 같은 좌표 집합인데 바뀐 순서