fragment = Tensor Core 가 처리할 타일 조각을 warp 가 레지스터에 담아놓은 형태

fragment = warp 한 개가 16, 16, 16 MMA 연산을 하기 위해 A, B, C 타일을 lane(스레드)별 레지스터에 쪼개 저장한 행렬 조각

즉, fragment 란

• 메모리에 있는 행렬에 대응하는 게 아님
• warp 안에서만 존재하는 레지스터 구조체
• Tensor Core 연산자 (mma_sync) 가 바로 읽어서 쓰기 위한 전용 데이터 구조

각 fragment 는 행렬 타일 하나를 warp 레벨에서 들고 있다.

 

Fragment 데이터 구조가 필요이유

Tensor Core 의 기본 연산이 16, 16, 16

• Warp 32 lanes 는 16, 16 개 C element 를 계산해야 함
• lane 하나가 256 개를 다 들고 있을 순 없음
• 그래서 warp 가 전체 타일을 32 개로 나눔 각 lane 은 8개 element 를 담당

이 8개가 레지스터에 흩어진 모양으로 저장, 그걸 들고 있는 구조체가 fragment.x[] 배열

fragment = warp 레벨에서 16, 16 타일을 lane 들이 나눠 들고 있는 메타 행렬

 

Fragment 종류 3가지

Tensor Core 는 A, B, C operand 가 각각 다르게 생겼기 때문에 fragment 타입도 3개

1) matrix_a fragment

wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a_frag;

• A 타일 16, 16 을 warp 에 분배해서 저장한 구조
• A 는 row-major 또는 col-major 선택 가능
• 내부적으로 lane 별로 A 의 특정 8개 element를 들고 있음

 

2) matrix_b fragmet

wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b_frag;

• A 타일 16, 16 을 warp 에 분배해서 저장한 구조
• B 는 보통 col-major 를 많이 씀

 

3) accumulator fragment (C fragment)

wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;

• C 타일 16, 16 = 256 element 를
• warp 32 lanes 가 8개씩 나누어 들고 있는 구조

 

fragment 가 하는 실제 일

(1) 메모리에서 타일 로드

wmma::load_matrix_sync(a_frag, A_ptr, lda);

• global/shared memory 에 16, 16 로 저장된 행렬을
• warp 의 fragment 형태로 lane register 로 분배

 

(2) Tensor Core 연산 수행

wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

• a_frag, b_frag 에 든 16, 16, 16 데이터를 Tensor Core 에 넣어
• c_frag 의 256 칸 부분합을 한 번에 갱신함

 

(3) 최종 결과를 메모리에 저장

 wmma::store_matrix_sync(C_ptr, c_frag, ldc);

• c_frag 에 들어 있던 lane 별 8개 값들을
• 16, 16 C 타일 형태로 메모리에 재구성해서 저장