GEMM 확장 방향성 - 현재 완전 기본적인 shared memory, tiling 사용 중

3. 그 다음 확장 방향들 (실험 메뉴판)

이제부터는 방향만 던질게. 하나씩 골라서 실험해보면 좋음.

(1) BLOCK_TILE / THREAD_TILE / BK 분리

지금은 TILE x TILE 이 block tile이자 shared tile이자 K tile.

조금 더 일반화하면:

• BM, BN: block tile 크기 (예: 128 x 128)
• BK: K 방향 타일 (예: 8, 16, 32)
• TM, TN: thread tile 크기 (예: 4 x 4)
• blockDim.x, blockDim.y: warp/스레드 조직

이렇게 잡고:

• shared As[BM][BK]
• shared Bs[BK][BN]
• 각 스레드는 (TM x TN) C 서브타일을 담당

=> 너가 지금 만든 테스트 코드에 template 파라미터로 이거 늘려서,
조합 바꾸면서 GFLOPs 비교해 보는 벤치마크로 확장 가능.

---

(2) vectorized load (float4 / float2) + coalescing 체크

지금 로드는 A[row*K + kA], B[kB*N + col]로 스칼라 로드.

실험용 확장:

• reinterpret_cast<const float4*> 로 float4 단위 로드
• shared에도 [row][col] 대신 float4* 형태로 쌓거나, 로드 후 4개 풀어쓰기
• N, K를 4의 배수/8의 배수로 맞춘 케이스와 아닌 케이스 비교

이걸 벤치에서 옵션으로 켜고 끌 수 있게 만들면 “vectorized load 효과”를 체감하기 좋음.

---

(3) double buffering / prefetch 아이디어

현 커널은:

1. 타일 t 로드
2. TILE 만큼 곱
3. 다음 타일 t+1 로드

이걸 ping-pong shared buffer로 바꾸면:

• As0/Bs0, As1/Bs1 두 세트 shared mem
• t를 진행하면서:
  • t 타일 계산 중일 때 t+1 타일을 미리 로드 (지금은 그냥 ld; __syncthreads(); compute 구조라 완전한 overlap은 어렵지만 패턴을 익히는 용도)

나중에 Ampere cp.async 익히면 여기다 그대로 얹으면 됨.

---

(4) perf harness 확장: cublas baseline + parameter sweep

gemm_perf.cu를 아래처럼 키울 수 있음:

1. run_gemm_cublas 추가해서 cublasSgemm 성능 측정
2. 여러 TILE / (TM, TN) 조합을 돌리도록:
   • CaseCfg에 TILE 말고 커널 ID 추가 (예: NAIVE, SHARED_32, SHARED_64, REG_32_2x2, …)
3. M,N,K를 정사각형 + 비정사각형 (예: 1024x768x512, 4096x256x4096 등) 여러 개로 확대
4. 결과를:
   • GFLOPs
   • “우리 커널 / cuBLAS” 비율
   • max diff (검증)

표처럼 찍히게 만들면 나중에 포폴/블로그에 그대로 박아 넣을 수 있음.

---

(5) epilogue 확장: bias + activation까지 달아서 “연산 fusion” 감각 익히기

지금은 순수 GEMM만 함.

• gemm_tiled_* 끝부분에서
  • C[...] = acc; 대신
  • acc += bias[...] (per-N/per-M 중 하나)
  • acc = max(0, acc); 같은 ReLU 붙이기
• 이걸 옵션으로 켜고 끌 수 있게 만들면:
  • “GEMM + bias + act” tuple을 단일 커널로 처리 vs 따로 처리 비교 가능

이건 너가 이미 프레임워크에서 하고 있는 “epilogue fusion” 컨셉이랑 그대로 이어지고,
실제 커널 수준에서 어떤 모습인지 체감하기 좋음.

---

(6) register / occupancy 분석까지 포함

NVCC 옵션: