int8 tensorcore matmul for Turing

강의 도입에서 Erik 이 분명히 한 자리 — "this is not state-of-the-art fast or even close. it's about getting the basics right." 같은 시리즈 두 주 뒤에 H100 cuBLAS-beat 강의가 잡혀 있는 자리에서, 일부러 고른 old GPU(Turing) + 가장 작은 도구셋(CUDA C++ 만, no CUTLASS, no inline asm). 추상이 다 걷힌 자리에서 tensor core 가 어떻게 도는지를 손에 잡히게 만드는 게 목표다.

그래서 이 강의는 한 큰 결과를 보여주는 게 아니라 — 점진적으로 빨라지는 5개의 커널과 그 사이의 NCU diff를 본다. naive (~4 TFLOPs) → column-major → SMEM tile → register tile → ... 의 단계마다 어떤 metric 이 무엇 때문에 어떻게 변하는지를 같은 도구(Nsight Compute) 의 같은 화면에서 본다.

이 강의의 가치 — 같은 도구(NCU)가 같은 코드의 5개 변형에서 어떻게 다른 hint 를 주는지의 reading 자료. NCU 를 처음 본 사람에게 가장 좋은 1시간짜리 worked example 이다.

강의 첫 5분이 집중한 자리. INT8 matmul 이 단순히 “float 을 int 로 바꾼” 게 아니라는 것 — register width 가 안 맞는다. INT8 두 개를 곱하면 결과는 INT16, accumulator 는 INT32. 즉 input 폭과 output 폭이 다르다. 이게 hardware instruction 의 모양을 결정한다.

// CPU (AVX-512 VNNI) — 한 instruction
__m512i acc = _mm512_dpbusd_epi32(acc, a, b);
//        ^^^^ accumulate INT32
//             dot product of 4×INT8

// CUDA non-tensor core (DP4A)
int acc = __dp4a(a4, b4, acc);

// CUDA tensor core (Turing+)
nvcuda::wmma::fragment<...> a_frag, b_frag, c_frag;
mma_sync(c_frag, a_frag, b_frag, c_frag);

Turing 의 INT8 tensor core 는 warp 단위 instruction 이다. 32개 thread 가 같이 호출해야 한다. 한 번의 호출이 16×16×16 matmul 을 처리. 16 = M, 16 = N, 16 = K.

// 강의의 첫 tensor core kernel — 약식
__global__ void tc_matmul(
    const int8_t* A, const int8_t* B,
    int* C, int M, int N, int K)
{
    // 16x16 fragments
    fragment<matrix_a, 16,16,16, int8_t, row_major> aF;
    fragment<matrix_b, 16,16,16, int8_t, row_major> bF;  // ← row major B = strided
    fragment<accumulator, 16,16,16, int>          cF;

    fill_fragment(cF, 0);

    for (int k = 0; k < K; k += 16) {
        load_matrix_sync(aF, &A[i*K + k], K);
        load_matrix_sync(bF, &B[k*N + j], N);
        mma_sync(cF, aF, bF, cF);
    }

    store_matrix_sync(&C[i*N + j], cF, N, mem_row_major);
}

강의 첫 큰 surprise. 위의 코드를 15000×15000 크기 입력으로 돌리면 — 1.5초, 약 4 TFLOPs. Turing 의 INT8 tensor core peak 가 60+ TFLOPs 이므로 일 자릿수 utilization. 이유를 NCU 가 가리킨다.

그리고 NCU 가 같이 보여주는 두번째 화면 — warp stall reasons.

강의의 첫 fix. row-major 두 행렬에서 K 차원으로 sum 을 돌면, B 의 access 는 K 만큼 strided = uncoalesced. column-major 로 두면 K 차원이 contiguous = coalesced. 변경은 fragment 의 layout flag 한 줄과 — B 를 미리 column-major 로 transpose 하는 별도 kernel.

다음 fix. block 안의 여러 warp 가 같은 A 행, 같은 B 열을 반복해서 읽는다. 한 번 SMEM 에 올려두면 — 그 block 안에선 register 로 stage 가능. PMPP 책의 표준 tiling 패턴이 그대로.

SMEM 으로 block 안에서 reuse 했지만, warp 안의 같은 register fragment 는 여전히 여러 번 다시 load 된다. 한 warp 가 더 큰 출력 tile (예: 32×32 또는 32×64) 를 계산하면 — 같은 A fragment 를 여러 B fragment 와 곱한다 = register reuse.

강의에서 별도 섹션처럼 풀어준 자리. NCU 의 “Warp State Statistics” 가 보여주는 stall 카테고리들의 의미. 매 단계 NCU diff 를 읽을 때 이 어휘가 손에 있어야 한다.

강의에서 Erik 가 한 정리 — "naive 에서 column-major 로 가면 LG throttle 이 가장 많이 줄고, SMEM tile 추가하면 long scoreboard 가 줄고, register tile 추가하면 short scoreboard 까지 줄어든다. math pipeline stall 이 늘어나면 compute bound 에 다가가는 것".

다음 주 강의(L045 H100 cuBLAS 격파)와 이 강의의 코드 시퀀스가 같은 도구로 풀리지 않는 이유 — Hopper 가 새로운 hardware 의 자리를 추가했기 때문.

강의에서 6개월 뒤 다시 돌아왔을 때 가장 빨리 복원해야 하는 사실들과 — 직접 손에 박아야 할 코드 자료들.

손에 새기기 — 실습 시퀀스

1. naive INT8 mma.sync 커널 베이스라인 — wmma fragment 로 16×16×16 호출. 5000² 정도로 작게 시작. cuda.Event 로 ms 측정.
2. NCU 한 번 돌리기 — ncu --set full -o trace ./naive. 메모리 그래프와 warp stall 분포 확인. 특히 LG Throttle / Long Scoreboard 비율.
3. column-major B — 별도 transpose 커널 작성, B 를 column-major 로 변환 후 fragment 의 layout flag 변경. NCU diff 로 traffic 감소 확인.
4. SMEM tile 추가 — block tile size 128×128, K-stage 32. __shared__ int8_t As[128][32] 같은 declare. bank conflict 회피용 padding 한 줄 추가.
5. register tile 키우기 — warp 당 출력 32×32 또는 32×64. fragment 여러 개를 한 warp 가 들고 있게.
6. NCU diff 5단계 모두 비교 — ncu compare 또는 같은 trace 두 개를 GUI 에서 열기. stall 분포 변화의 양상을 보고서로 정리.
7. cuBLAS INT8 baseline 측정 — 같은 크기로 cuBLAS 의 INT8 GEMM. 본 강의의 결과가 cuBLAS 의 몇 % 인지 확인. 완벽 격차가 아니라 격차의 origin 이 학습 목표.

L043 의 5단계 사다리 (column-major → SMEM → register tile → double buffering) 가 시리즈 안에서 어디에 다시 호출되는지 묶어둔다.

학습 노트로 정리하면서 의도적으로 비워둔 자리들 — 강의 안에서 부분적으로만 등장한 주제, 또는 후속 자료가 더 깊게 다룰 주제.

• cuBLAS INT8 GEMM 의 절대 격차 — 본 강의는 의도적으로 SOTA 와 비교하지 않는다. 동일 사이즈에서 본 강의의 끝 커널이 cuBLAS 의 몇 % 인지는 직접 측정 필요.
• SMEM bank conflict 의 정확한 패턴 — INT8 8개가 한 word 인 layout 에서 어느 stride 가 conflict 를 만드는지의 일반화. NCU 의 shared store/load stalls 를 직접 본다.
• double buffering 단계의 자세한 코드 — 강의 시간 부족으로 마지막 변형은 스킵. async copy 가 없는 Turing 에서 어떻게 SMEM stage / compute overlap 을 구현하는지 — 별도 자료로 확인.
• 다양한 행렬 크기에서의 sweep — 본 강의는 15000² 만 본다. 1024², 4096² 등 작은 크기에서 launch overhead / tail effect 가 어떻게 변하는지 미확인.
• occupancy 가 떨어지는 자리의 정확한 register count — Turing 의 64K register file 안에서 한 warp 가 몇 개 이상 fragment 를 들면 occupancy 가 깨지는지 — --ptxas-options=-v 로 확인.
• signed vs unsigned INT8 의 성능 차이 — VNNI / DP4A 모두 mixed-sign instruction. tensor core 도 동일한지, 별도 instruction 인지 확인 필요.