기존 작성한 GE_v2 의 regemm 타일드 커널의 1:1 매핑, gemm_tiled_test 작성

기존 ge_v2 의 epilogue, bias, activation, Z-stash, template 정책을 제거한, 알고리즘 코어를 동일하게 구성한 test code 를 통해 커널 최적화를 구현하는 것이 목표

 

구조적 동일함을 항목별 확인을 통해 검증

1. 타일 분할 방식

기존 커널

const int m0 = blockIdx.y * BM_;
const int n0 = blockIdx.x * BN_;
const int tx = threadIdx.x;
const int ty = threadIdx.y;
const int tm0 = m0 + ty * THR_M;
const int tn0 = n0 + tx * THR_N;

테스트 커널

const int m0 = blockIdx.y * BM;
const int n0 = blockIdx.x * BN;
const int tx = threadIdx.x;
const int ty = threadIdx.y;
const int tm0 = m0 + ty * THR_M;
const int tn0 = n0 + tx * THR_N;

동일한 block - tile mapping

동일한 thread - mini-tile mapping

동일한 THR_M / THR_N 방식

 

즉, block 이 1개의 (BM x BN) C-타일을 담당하고 thread 는 그 타일 안에서 THR_M x THR_N 부분을 담당한다.

 

2. shared memory tile 배치 - 동일

기존 커널

__shared__ float As[2][BM_][BK_ + PADK];
__shared__ float Bs[2][BK_][BN_ + PADN];

테스트 커널

__shared__ float As[BM][BK + PADK];
__shared__ float Bs[BK][BN + PADN];

더블 버퍼링은 제거, 타일 구조는 완전히 동일

 

3. 타일 로딩 (load_A_tile / load_B_tile) - 동일

기존 커널

const int tid = ty * TDX + tx;
const int elems = BM_ * BK_;
for (int e = tid; e < elems; e += (TDX * TDY)) {
    r = e / BK_;
    c = e % BK_;
    gm = m0 + r;
    gk = k0 + c;
    As[stage][r][c] = A[gm * p.lda + gk];
}

테스트 커널

const int tid = ty * TDX + tx;
const int elems = BM * BK;
for (int e = tid; e < elems; e += nthreads) {
    r = e / BK;
    c = e % BK;
    gm = m0 + r;
    gk = k0 + c;
    As[r][c] = A[gm * K + gk];
}

인덱싱 / 패턴 / 병렬 분배 방식까지 100% 매핑

• block 내 스레드로 A/B 타일을 나눠서 채우는 방식
• boundary check 후 0-padding
• row-major 기반 global read

 

4. K-loop 구조 - 동일

기존 커널

for (int k0 = 0; k0 < p.K; k0 += BK_) {
    // (필요 시 next tile load)
    for (int kk = 0; kk < BK_; ++kk) {
        load a_vec
        load b_vec
        FMA(acc, a_vec, b_vec)
    }
}

테스트 커널

for (int k0 = 0; k0 < K; k0 += BK) {
    // load tile
    for (int kk = 0; kk < BK; ++kk) {
        load a_vec
        load b_vec
        FMA
    }
}

BK 단위 타일 K-loop 

inside 에서 rank-1 outer-product multiply-accumulate ( micro-kernel )

둘 다 동일한 classical shared memory tiling GEMM 패턴

 

5. 마이크로 커널 ( 레지스터 기반 outer product ) - 동일 

기존 커널 / 테스트 커널 완전 동일

float a_vec[THR_M];
float b_vec[THR_N];

acc[i][j] = fmaf(a_vec[i], b_vec[j], acc[i][j]);

스레드 당 THR_M x THR_N output 을 갖고 각kk 에서 outer-product 방식 FMA 

이건 SGEMM 오픈소스 커널에서 가장 표준적인 micro-kernel 형태

 

6. 모든 K tile 끝나면 C 저장 - 동일

기존 커널

C[m * p.ldd + n] = final_acc;

테스트 커널

C[m * N + n] = acc[i][j];

epilogue 없는 순수 저장 버전

 

차이점으로 

• epilogue 제거
• double-buffering 제거

그 외는 동일

 

7. 실행 결과

==== Tiled GEMM test ====
  M=1024, N=1024, K=1024
  Tile: BM=64, BN=64, BK=16
  Block: (16 x 4), Thread tile: (16 x 4)
[CPU] computing reference...
[Check] max |diff| = 1.831055e-04
[Perf]  avg time = 0.988 ms,  GFLOP/s = 2173.17


==== Tiled GEMM test ====
  M=2048, N=2048, K=2048
  Tile: BM=64, BN=64, BK=16
  Block: (16 x 4), Thread tile: (16 x 4)
[CPU] computing reference...
[Check] max |diff| = 2.441406e-04
[Perf]  avg time = 5.897 ms,  GFLOP/s = 2913.48

현재 성능을 기준으로 커널 최적화 수행 필요!!