AICF Kernel Engineering Report ( GEMM & BiasAdd )

GEMM Kernel Engineering

1.1 연산 특성 분석

Operation Type : Compute-Bound

핵심 병목

• Global Memory Latency
• Register Pressure
• Tensor Core Utilization

Data Reuse 가 높은 연산, 핵심 전략은 DRAM 접근을 최소화하고, 연산 파이프라인이 놀지 않게 만드는 것

 

1.2 Memory Hierarchy 전략

Step1 - Shared Memory Tiling

Naive GEMM

C[i,j] += A[i,k] * B[k,j]

k-loop 마다 DRAM 접근 발생 - Latency 문제

해결

• Block 단위 32 32 TILE 분해
• Global -> Shared Memory preload
• Thread Block 전체가 재사용

효과

Global Memory Traffic 감소

O(N³) → O(N³ / T)

Arithmetic Intensity 증가 - Roofline 상 Compute 영역 이동

 

Step2 - Shared Memory Bank Conflict 제거

Ampere Shared Memory

• 32 Banks
• 32 threads per warp

32 32 tile 을 그대로 쓰면

• column access 시 동일 bank 충돌

해결

__shared__ half tileA[32][32 + 8];

Padding 적용

결과

• Bank Conflict - 0
• Shared Memory Throughput 안정화
• Warp stall 제거

 

Step 3 - Tensor Core (WMMA) 활성화

CUDA Core 대신 Tensor Core 직접 사용

wmma::load_matrix_sync
wmma::mma_sync
wmma::store_matrix_sync

연산 단위 변화

Thread-level	Warp-level
Scalar FMA	16×16 MMA
CUDA Core	Tensor Core

핵심 설계 포인트

• Shared - Register Fragment
• Fragment Accumulator
• Register spiling

해석

• DRAM 병목 제거
• Tensor Core 활성화
• Compute Utilization 상승

Memory bound -> Compute bound 전환 성공

 

 

BiasAdd Kernel Engineering

2.1 연산 특성 분석

BiasAdd 의 본질은 얼마나 빨리 읽고 쓰느냐,

연산 자체는 중요하지 않다.

 

2.2 Bandwidth Saturation 전략

Step 1 - Coalesced Access

Warp 32 threads - 연속 주소 매핑

idx = blockIdx.x * blockDim.x + threadIdx.x;

목표

• 32개 요청 - 1 ~ 2 DRAM 트랜잭션

 

Step 2 - Vectorized I/O

Scaler 접근의 문제

• Instruction issue 증가
• Memory Controller Overhead 증가

해결

const half2* A_vec = reinterpret_cast<const half2*>(A);
half2 v = A_vec[idx];
half2 r = __hadd2(v, bias);

효과

• Instruction Count 감소
• DRAM Burst 효율 상성
• Issue Slot 낭비 제거

 

Step 3 - Grid-Stride Loop

대규모 텐서 대응

for (int i = idx; i < N; i += stride)

• Launch Overhead 감소
• Occupancy 유지

 

2.3 Alignment 전략

Vectorized 접근 조건

• 4Byte
• 16Byte

정렬 불일치 시 - Fallback Scalar Kernel 분기

이 설계는 Correctness 와 Performance 를 동시에 보장하는 구조