마이크로 커널은 GEMM 가은 큰 커널을 내부에서 더 작고 구조화된 연산 단위로 쪼개어 최적화하는 핵심 기법
CUDA 프로그래밍 전체 " 타일링 - 워프 구조 - 공유 메모리 - 레지스터 블록 " 의 큰 틀은 익숙
마이크로 커널은 그 중에서도 레지스터 단계의 가장 내부 루프를 전문적으로 설계하는 영역
GEMM = 큰 문제
타일링 ( 64, 64 ) = 중간 문제
워프 / 스레드 타일 ( 16, 4 ) = 더 작은 문제
마이크로 커널 = 각 스레드가 레지스터에서 반복적으로 수행하는 극소 단위 연산 패턴
마이크로 커널이 왜 중요한가
GEMM 성능의 70 % 이상은 사실 마이크로 커널 설계에서 정해진다.
이 영역이 결정하는 것은
- 레지스터 배열 방식
- FMA 발행 패턴
- A/B 레지스터 prefetch
- Unroll factor
- Instruction scheduling
- L1/L2/Shared latency hiding
- Warp-level accumulator layout
라이브러리가 최적화 하는 부분이 마이크로 커널
마이크로 커널의 기본 구조
전통적인 GEMM 내부 루프 구조
for k in 0..BK:
load A_frag from shared
load B_frag from shared
C_frag += A_frag * B_frag마이크로 커널은 이걸 더 정교하게 바꾼 버전
예 - 1 warp 내 32 threads 가 유지하는 레지스터 배치
A_reg[4]; // 4개 레지스터 fragment
B_reg[2];
C_reg[8]; // 8개 partial accumulators그리고 k 루프마다.
#pragma unroll
for (int kk = 0; kk < unroll; kk++) {
preload_next_A();
preload_next_B();
fma(C_reg[0], A_reg[0], B_reg[0]);
fma(C_reg[1], A_reg[0], B_reg[1]);
fma(C_reg[2], A_reg[1], B_reg[0]);
...
}여기에는 다음 전략이 숨어 있음
- FMA instruction pipeline 정확히 맞추기
- 레지스터 bank conflict 최소화
- cp.async/미리 가져오기로 latency 숨기기
- Stall 없이 warp schedulaer 가 full throughput 유지
- Warp 내 공유메모리 bank pattern 조정
- Accumulator spilling 없이 유지
즉, 완전히 "레지스터 단위 수학 머신" 을 설계하는 것과 같은 느낌
넓게 보면 이런 의미
딥러닝에서 GEMM 성능을 끝가지 잡는 건 다음 두 개
- 타일링 구조
- 마이크로 커널
마이크로 커널은 커널 엔진의 심장부!!??, 여기서 최적화가 곧 전체 GEMM 성능을 결정한다.
지금까지 작성해온 kernel 과의 차이
현재 GEMM 구조
- 공유 메모리 타일링 되어 있음
- Warp tile 도 존재
- 레지스터 accumulator 도유지
- Pipeline 의 적용
하지만 내부의 레지스터 계산 단위는 정형화된 micro-kernel 형태가 아님, ( C_reg tile 의 레이아웃, A/B unroll factor, scheduling 최적화는 미정 )
마이크로 커널을 이해하면 cuBLAS FP32 의 구조를 볼 수 있음
왜 이제서야 안거야 마이크로 커널을??!?!?
일반 CUDA 개발자들이 공부하지 않는 계층??!!
- NVIDIA 세미나
- BLIS/BLAS 논문
- CUTLASS Gemm core
- HPC 연구실
- cuBlAS 내부 구조 분석
이런 영역에서만 등장하는 용어,