본문 바로가기

GPU-KERNEL

GEMM: Naive vs Shared-Memory Tiled 성능 비교 (간단 정리)

1. 실험 개요

GPU에서 행렬 곱셈(GEMM, C = A×B)의 두 가지 구현을 비교했다.

  • Naive GEMM
    • Shared memory 사용 없음
    • 스레드당 C(row, col) 하나 계산
    • A/B를 global memory에서 직접 읽음
  • Shared-Memory Tiled GEMM (TILE = 32)
    • A/B를 32×32 타일 단위로 shared memory에 로드
    • 타일 내부에서 데이터 재사용 후 accumulate
    • 스레드당 C(row, col) 하나 계산

문제 크기:
1024³, 2048³, 4096³

2. 측정 결과

성능 요약

Size                                          Naive (GFLOP/s)                    Tiled (GFLOP/s)                       Speedup (time ratio)
1024³ 749.5 869.3 0.86x
2048³ 771.8 899.8 0.86x
4096³ 764.1 908.0 0.84x

정확도: 모든 케이스에서
max |diff| = 0 (FP32 동일 결과)


3. 해석

  • Shared memory 타일링은
    약 15–20% 정도의 성능 향상을 제공했다.
  • 그 이유:
    • Naive도 B축 접근은 coalesced가 잘 나와 기본 캐시 효율이 높다.
    • Tiled는 global memory 트래픽을 줄여주지만,
      shared load + sync 오버헤드가 함께 존재한다.
  • 이 때문에 단순 타일링만으로는 “몇 배” 향상은 나오지 않는다.

4. 결론

“Shared memory만 적용한 1단계 타일링”의 실질적 이득은 약 15–20% 수준이다.
더 큰 성능 향상은 다음 기법들을 함께 적용해야 얻을 수 있다.

  • 스레드당 다중 출력 계산 (레지스터 타일링)
  • Double-buffering / cp.async
  • Tensor Core (WMMA)

 

“C 행렬의 TILE×TILE 부분을 계산하는 블록이고,
그 영역에 필요한 A/B의 일부분을 shared memory에 올려서,
각 스레드가 자신이 맡은 C 원소 하나를 완성시키는 구조”

이걸 그냥 한 줄로 요약하면:

  • Block = C 타일 전체 담당
  • Thread = C 타일 내부 원소 하나 담당
  • Shared tile = 그 C 타일을 계산하는 데 필요한 A/B 조각들 캐시