GPGPU 총정리 ( NVIDIA Tensor Core Programming ) - (5)

CUDA core 같은 경우에는 범용적인 SIMD 연산기 이다. ( 지금까지 해당 파트를 배운 것이고. ) 

반면에 , Tensor Core 는 딱 하나의 연산만 한다. 이 연산이 어떻게 보면 AI 에서 아주 유용하게 nvidia gpu 를 사용할 수 있게 한 꽃이라고 할 수 있다. 

 

 

MMA ( mixed-precision matrix multipy ) 

 ** D = A*B + C 8* 

해당 연산이 Tensor core 에서 수행되는 연산이다. 이 연산이 딥러닝에서 90%이상 차지한다고 보면된다. 

(Attention QK , Linear layer , conv , 등... ) 

 

volta ( v100 ) 에서 처음 도입되고 , 지금은 blackwell 까지 와서 SM 당 4개씩 존재한다고 한다. 

 

cf. 참고를 위해서.

 

 

 

 

연산은 너무 간단해서 깊게 설명할 것 까지는 없으나 , 몇가지 특징이 존재한다. 

 

(1) Mixed precision 

A,B 는 FP16 으로 읽고, 곱셈은 FP16 

accumluation ( +C) 는 FP32 로 합니다. 

 

곱셈은 값의 크기를 크게 바꾸지 않지만, accumulation 은 엄청 많은 항을 더하면서 오차가 쌓이기 때문에 더 높은 정밀도가 필요하다고 한다. 그래서 그냥 +c 부분도 FP16 만으로 누적하면 결과가 망가진다. 

입력은 가볍게 하고 , 누적을 안전하게 하는 방식으로 보면된다. 

 

(2) Warp-level operation

CUDA core는 thread-levl (각 thread 가 자기연산)만 한다. 

Tensor core 에서는 32개 thread (warp)가 협력해서 하나의 행렬 연산(matrix fragment)을 처리한다. 

 

( 그래서 Tensor Core programming에서는 일반 CUDA처럼 “thread 하나가 output element 하나 계산”하는 방식보다, warp 단위로 matrix tile을 처리한다고 이해해야 한다. ) 

 

중요한 내용이라 이해를 위해 예시 (with claude)

시나리오: 16×16 행렬 곱 C = A·B 한 번 계산

방식 1: 기존 CUDA Core 방식

각 thread가 C의 원소 하나를 담당합니다.

 

 

thread (0,0) → C[0][0] 계산 → for k: A[0][k] * B[k][0] 누적
thread (0,1) → C[0][1] 계산 → for k: A[0][k] * B[k][1] 누적
thread (5,7) → C[5][7] 계산 → for k: A[5][k] * B[k][7] 누적
...
thread (15,15) → C[15][15] 계산

총 256개 thread가 각자 독립적으로 자기 원소를 계산해요. 서로 대화 안 합니다. 각 thread 입장에서는 "나는 내 한 칸만 책임진다".

 

방식 2: Tensor Core 방식

32개 thread (1 warp)가 통째로 16×16×16 행렬 곱 하나를 처리합니다.

 

 

warp의 32개 thread가 모두 모여서:
  ├─ thread 0: A의 어떤 조각 + B의 어떤 조각 + C의 어떤 조각을 register에 들고 있음
  ├─ thread 1: A의 다른 조각 + B의 다른 조각 + C의 다른 조각을 register에 들고 있음
  ├─ ...
  └─ thread 31: A의 또 다른 조각 + ...
  
  → wmma::mma_sync() 호출
  → Tensor Core가 32 thread의 register를 한꺼번에 읽어서 
     하드웨어 회로로 행렬곱 처리
  → 결과 D의 조각들을 다시 32 thread의 register에 분배

여기서 thread 하나만 보면 자기가 뭘 계산하는지 모릅니다. "C[5][7]은 내가 책임진다" 같은 1대1 매핑이 없어요. 그냥 "나는 fragment의 일부 데이터를 들고 있고, 32명이 같이 mma_sync를 부르면 어떤 마법이 일어나서 결과가 나온다"인 거예요. 

 

진짜 차이는 뭐냐

(1) 누가 일을 하느냐

                                                                          CUDA Core                                 Tensor Core

 

일하는 단위	thread 1개	warp 32개 thread
한 번에 처리하는 출력	scalar 1개	행렬 조각 (예: 16×16)
코드 작성 시점	thread 관점	warp 관점

(2) 코드 작성 mental model이 완전히 다름

CUDA Core 코드 짤 때:

 

int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0;
for (int k = 0; k < K; k++) {
    sum += A[row][k] * B[k][col];  // 내 행, 내 열, 내 누적
}
C[row][col] = sum;  // 내 칸에 저장

"나(thread)"가 주어. 내가 어떤 row, 어떤 col을 맡았는지 명확.

 

 

Tensor Core 코드 짤 떄 

wmma::fragment<matrix_a, 16,16,16, half, row_major> a_frag;
wmma::fragment<matrix_b, 16,16,16, half, row_major> b_frag;
wmma::fragment<accumulator, 16,16,16, float> c_frag;

wmma::fill_fragment(c_frag, 0.0f);
wmma::load_matrix_sync(a_frag, A_ptr, 16);   // 32 thread가 같이 부름
wmma::load_matrix_sync(b_frag, B_ptr, 16);   // 32 thread가 같이 부름
wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);  // 32 thread가 같이 부름
wmma::store_matrix_sync(C_ptr, c_frag, 16, mem_row_major);

"나(thread)"가 안 보여요. 주어가 "우리 warp"예요. 각 함수는 32 thread가 동시에 같은 값으로 호출해야 합니다. 한 명이라도 다른 인자로 부르면 undefined behavior.

 

(3) "fragment"가 그 추상화의 핵심

a_frag는 표면적으로는 한 변수처럼 보이지만, 실제로는 32개 thread의 register에 분산 저장된 행렬 조각이에요. 즉:

• thread 0의 a_frag ≠ thread 5의 a_frag — 각자 다른 데이터 일부를 들고 있음
• 그런데 코드상에서는 그냥 a_frag라는 하나의 변수로 보임
• 누가 어떤 원소를 들고 있는지는 하드웨어가 알아서 정함 (opaque layout — 프로그래머가 신경 쓸 필요 없음)

이게 처음에 굉장히 어색합니다. 일반 변수처럼 생겼는데 사실 32명이 나눠 들고 있는 공유 객체. "나는 a_frag[0][0]을 들고 있나?" 물어봐도 의미 없어요. 그건 알 필요가 없게 설계돼 있어요.

그래서 실용적으로 뭐가 바뀌나

p.23 예제 (matmulT) 다시 봐보세요:

 

 

cpp

int warp = (blockDim.x * blockIdx.x + threadIdx.x) / warpSize;
int cx = warp % (Bx/16);
int cy = warp / (Bx/16);

여기서 warp 변수가 핵심이에요. thread index를 32로 나눠서 이 thread가 속한 warp의 ID를 구합니다. 그리고 그 warp ID로 "이 warp가 C의 어느 16×16 타일을 책임지는가"를 정해요.

즉, 일반 CUDA에서 "이 thread가 C의 어느 원소를 책임지는가"를 정했다면, Tensor Core에서는 "이 warp가 C의 어느 타일을 책임지는가"로 한 단계 올라간 거예요. tile은 16×16이니까 한 warp가 256개 원소를 한 번에 처리하는 셈이고, 그게 mma_sync 한 번으로 끝납니다.

한 줄 요약

CUDA Core는 "1 thread = 1 output element", Tensor Core는 "1 warp = 1 output tile". 코드의 주어가 thread에서 warp로 올라가고, 그 warp 안의 32 thread는 fragment를 나눠 들고 협력 호출을 한다.

이 mental model 전환이 익숙해지면 그 다음 페이지들(WMMA API, shared memory 버전 등)이 훨씬 잘 읽힙니다. 더 헷갈리는 부분 있으면 알려주세요. 

 

 

FMA = Fsed Multipy-add 

즉 , a*b+c 를 한번의 연산으로 처리하는 것이다. 

 

여기서 유의할 점은

- 곱셈 한번(a*b) + 덧셈 한번 = floating-point 연산 (FLOP) 2개 

- 하드웨어 입장에서는 회로 하나로 한사이클 처리 = FMA 1개 

 

이걸 딱 잡고 왼쪽 아래 그림을 살펴보면, 

(우리는 C = A*B+C , 즉 다시말해 C+=A*B 의 형태를 보기로 약속한다. )

1. 입력 (A,B) - FP 16 으로 저장

2. 곱셈 : FP16 x FP16 , 결과는 full precision 으로 받는다 ( 오버플로 이런거 방지 ) 

3. 누적 : 결과를 FP32 에 누적한다 ( 이렇게하면 오차가 상대적으로 적어지고 ) 

4. 출력 : FP32 

 

 

(volta 기준)

이번에는 오른쪽 아래 그림을 살펴보자 : 4x4x4 MMA 

D(4×4) = A(4×4) × B(4×4) + C(4×4)

 

volta 에서 tensor core 1개가 한클럭에 처리하는 최소 단위이다. 4x4 행렬 3개를 받아서 (a,b,c) -> 4x4 행렬을 하나 뱉는다. 

 

FMA 개수

- 출력 D (4x4) = 16개

- D[I][J] 가 A 의 row 한줄, B 의 col 한줄 + C 하나 이므로 , 총 4번의 FMA 를 해야한다. (4X4 니까그냥..) 

- 총 FMA = 16개 원소 * 4 ( 바로 위에서 구한게 4번의 FMA ) = 64 FMA 

즉,  64 FP16/FP32 FMA operations per clock per Tensor Core 

(FLOPS 로 계산하면 64*2 = 128 FLOPS. )

 

SM 단위로 살펴보자. 

Volta 에서 하나의 SM 에는 tensor core 가 8개 존재한다. 

 

따라서..

TC 1개당 64 fma/clock

*8개 (SM당 )

--> 512 FMA / CLOCK / SM 

(FLOPS = 1024 ) 

이다. 

 

 

---

 

 

WMMA = warp matrix mltiply-accumulate 

이름 그대로 warp 가 협력해서 하는 MMA 를 호출하는 cuda c++ api 이다. 

(위에서 설명했듯이 , Tensor core 는 thread 하나가 혼자하는 연산이 아니라 , warp 단위로 협력해서 수행하는ㅇ ㅕㄴ산. ) 

 

 

코드를 하나씩 살펴보자. 

#include <mma.h>
using namespace nvcuda;

__global__ void wmma_ker(half *a, half *b, float *c) {

 

__global__ 이므로 GPU 에서 도는 커널이다. 입력 A,B 는 fp16 포인터 , 출력 c 는 fp32 포인터이다. ( mixed precision ) 

 

 

fragment 선언

wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::col_major> a_frag;
wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::row_major> b_frag;
wmma::fragment<wmma::accumulator, 16, 16, 16, float> c_frag;

fragment는 "warp 32개 thread가 나눠 들고 있는 행렬 조각"을 담는 컨테이너이다.

( 행렬 전체가 아니라 , Tensor core 연산에 필요한 작은 행렬 조각 ) 

일반 변수처럼 보이지만 실제로는 32 thread의 register에 흩어져 저장되고, 어떤 스레드가 뭘 가지고 있는지는 알 수 없다 ( 컴파일이 알아서 정함 ! ) 

 

인자 정리

matrix_a  : fragment 의 역할 

M,N,K ( 2,3,4 번째 인자) : 16 ,16 ,16 

half : element type ( fp16 )

col_major : layout  

 

다시 좀 풀어서 설명을 해보자면, 

a_frag는 변수 하나처럼 생겼지만, 실제로는 warp 안의 32개 thread가 16×16 행렬을 256조각으로 나눠서 각자 register에 들고 있는 상태이다. 

( 그러면 한 thread 당 = 256 / 32 = 8개 원소를 가지고 있게 되겠다 - 다만 , 각 thread 별로 접근은 불가능 ! ) 

 

M,N,K 란? 

A: M × K (M행, K열)

B: K × N (K행, N열)

로 정의된다 

다만 이 숫자는 이후에 말하는 표에서만 골라서 연산할 수가 있다.

 

col_major / row_major 

: c++ 배열이 기본적으로 row_major 이기는 하나 , col_major 인 경우도 있다. 따라서 tensor 에 넘겨줄때 해당 행렬이 어떤 방식으로 저장되어 있는지 알려주어야 한다. 

 

 

**특이점*

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

 

C/D(accumulator) 에는 lay_out 이 없다. 

C/D 의 layout 은 선언 시점이 아니라 , load&stor 시점에 정해진다고 한다. 

 

 

누적 변수 0으로 초기화

wmma::fill_fragment(c_frag, 0.0f);

c += a*b 를 하기 이전에 c 를 먼저 0으로 설정하는 모습.

여기서도 유의해야할 점은 32thread 가 모두 같이 호출된다는 점 ! 

 

 

메모리에서 fragment로 데이터 로드 

wmma::load_matrix_sync(a_frag, a, 16);
wmma::load_matrix_sync(b_frag, b, 16);

앞에서  a_frag 를 선언하기는 했지만 , 아직 데이터는 없는 상태이다. 

원본 데이터는 global memoryt a 배열에 저장되어 있는데 , 이걸 fragement 로 옮기는 것이 load_matrix_sync 함수이다. 

 

global memory a 배열에서는 256 개 원소가 쭉 있을 것이고 이걸 loabl_matrix_sync 로 가지고 오면 

-> warp 32 thread register 에 각 thread 마다 8개씩 넣는다 ! ( = a_frag ) 

 

인자에서 가장 헷갈리는 부분이 마지막에 있는 "16" 부분이다. 

이게 leading dimension (= stride ) 를 말하는 부분이다. 

(메모리는 1차원인데 , 행렬은 2차원 이므로 . 어디서 한행이 끝나고 다음행 시작인지를 알려주기 위함 ) 

 

지금 우리 예시에서는 16x16 을 가정하고 있지만 , 1024x1024 같이 큰 행렬이 있다면 이를 16x16 타일로 쪼개서 처리해야한다. 

이떄는 stride 가 1024 가 되고 ( 다음 행으로 옮기는데 1024 개를 건너 띄워야 하니까.. ) 이 경우에는 stride=1024 값을 주면되는 식이다. 

// 조각만 따로 저장된 경우
wmma::load_matrix_sync(a_frag, a, 16);    // stride = 조각 너비

// 큰 행렬(너비 1024)의 일부를 읽는 경우
wmma::load_matrix_sync(a_frag, &A[tile_pos], 1024);  // stride = 원본 너비

 

cf.

함수뒤에 익숙한 _sync 가 존재하는데 여기서는 warp 의 32thread 가 모두 여기 올 때 까지 기다린 다는 의미이다 

( 어찌보면 당연하다. 여기서는 warp 단위로 실행한다고 했으니. ) 

fragment 는 32 thread 가 다 같이 모여서 256 개 원소를 협력해서 읽어야 한다 ( 누구라도 빠지면 데이터가 불완전 ) ! 

(일반적인 CUDA 에서 thread 별로 데이터를 관리하는 것과 아예 다르다 ! ) 

 

 

wmma::load_matrix_sync(a_frag, a, 16); 

에서 32thread 가 동시에 부르면..

 

• 32 thread가 barrier에서 모임 (_sync)
• 하드웨어가 256개 원소(16×16)를 32 thread에게 분배 (한 명당 평균 8개)
• 각 thread가 자기 몫을 global memory에서 읽어 자기 register에 적재
• 이제 a_frag가 데이터로 채워짐 → mma_sync 호출 준비 완료

이런 루트로 내부에서 일어난다고 보면 된다. 

 

 

 

---

 

 

 

( 전반적인 설명은 필기본으로 대체 ) 

 

 

 

 

< NVFP4 ( 4비트 추론 ) >

NVFP4 는 숫자 하나를 4비트로 표현한다 ( 표현가능한 값 16 개 ) 

 

NVFP4 (E2M1): [S][E×2][M×1] = 4 bits 

표현범위는 대략 -6~6

실제 값: 0, 0.5, 1, 1.5, 2, 3, 4, 6 (그리고 음수 버전)

 

값을 8개 밖에 사용하지 못하는데 어떻게 이걸 추론에 사용할까 ? 

(weight 같은 경우는 실수값으로 매우 다양한데 , 이를 뭉개버리면 정보 손실이 크다. ) 

 

NVFP4는 크게 두가지 해결책을 사용해서 이를 해결한다 . 

 

(1) high-precision scale encoding 

"값 자체" <=> "배율(sclae) " 을 따로 저장한다 

 

실제 값 == (4bit index) * (배율 scale) 

 

 

(2) two-level micro-block scaling

해당 배율을 어떻게 저장할지에 대한 방법이다. 

 

1. 16개씩 묶어서 micro-block마다 배율 (E4M3 FP8)

텐서 전체를 16개 값씩 잘게 쪼갠다. 이 16개 묶음을 micro-block이라 하고, 각 묶음마다 하나의 배율을 공유한다. 

그래서 각 블록마다 FP8 배율 1개를 공유하는 형식이다. 

 

2. 텐서 전체에 배율을 하나더 (FP32 ) 

그리고 그 위에 텐서 전체에 적용되는 글로벌 배율을 FP32 로 하나 더 얹는다. 

 

복원값 = 4비트 index * 블록 FP8 배율 * 텐서 32배율

 

이 두가지 배율 덕분에 4비트라는 엄청난 압축에도 정확도를 비슷하게 가져갈 수 있다. 

NVFP4 = 4비트 거친 값(E2M1) + 16개 블록마다 FP8 배율 + 텐서당 FP32 배율. 

 

이를 주로 추론 가속용에 사용한다 

최근들어 LLM 같은 경우에 inference 의 비용을 줄이는게 더 중요하다 ( 학습은 한번 비싼 값으로 하는건 ok , 하지만 사용자들이 사용하는 추론 비용은 줄여야한다 ) 라는게 트랜드이고 정설인 것으로 아는데 , 엔비디아가 이에 맞추서 이런 하드웨어를 미리 만들어둔게 아닐까 싶다 

(허용되는 조합 - 참고용 ) 

 

 

---

 

Matrix Multiplication Using Tensor Cores 

이제 배운 것들의 총 정리본이 되겠다. 

이파트는 (아마도) 시험에 가장 잘 나오기도 쉽고 무엇보다 차원계산을 좀 머리속에 담아둔 채로, 특히나 지금까지는 thread level 의 컨셉을 가지고 했다면 , 여기서는 warp 단위로 생각해야하는게 가장 좀 까다롭다. 

하지만 한번 하고나면 또 근본적으로 많이 다를게 없기 때문에 괜찮다. 

한편 살펴보도록 하겠다 !. 

 

 

 

Matrix Multiplication without Using Shared Memory 

__global__ void matmulT(float* __restrict C, const half* __restrict A, const half* __restrict B,
                        int Ay, int Ax, int Bx) {
    int warp = (blockDim.x * blockIdx.x + threadIdx.x) / warpSize; // grid 내 warp 번호

    int cx = warp % (Bx / 16);   // 이 warp가 맡은 C 타일의 (x=열, 
    int cy = warp / (Bx / 16);   //                          y=행) 좌표

    int Atile_pos = cy * 16 * Ax; // A 첫 타일 시작 위치 (행 방향)
    int Btile_pos = cx * 16;      // B 첫 타일 시작 위치 (열 방향)

    // fragment 선언 (전부 16x16x16)
    wmma::fragment<wmma::matrix_a, 16,16,16, half, wmma::row_major> a_frag;
    wmma::fragment<wmma::matrix_b, 16,16,16, half, wmma::row_major> b_frag;
    wmma::fragment<wmma::accumulator, 16,16,16, float> c_frag;
    wmma::fill_fragment(c_frag, 0.0f);   // C 타일을 0으로 초기화

    for (int k = 0; k < Ax / 16; k++) {                 // 타일 단위 누적 루프
        wmma::load_matrix_sync(a_frag, &A[Atile_pos], Ax);  // A 타일 로드
        wmma::load_matrix_sync(b_frag, &B[Btile_pos], Bx);  // B 타일 로드
        wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);     // C += A*B
        Atile_pos += 16;        // A에서 오른쪽으로 한 타일 이동
        Btile_pos += 16 * Bx;   // B에서 아래로 한 타일 이동
    }
    wmma::store_matrix_sync(&C[(cy * Bx + cx) * 16], c_frag, Bx, wmma::mem_row_major);
}