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 조각들 캐시
'GPU-KERNEL' 카테고리의 다른 글
| kernel_hierarchy_test (0) | 2025.11.19 |
|---|---|
| 작성된 gemm_tiled_kernel<float, TILE><<<grid, block>>> 의 의미 (0) | 2025.11.19 |
| GEMM 커널 테스트 코드 작성 및 최적화 (1024^3 크기 실험) (0) | 2025.11.18 |
| ILP ( Instruction - Level Parallelism ) (0) | 2025.11.18 |
| 다음 추가 개념 리스트 (0) | 2025.11.18 |