GPU 연산 최적화 - 내가 잘못생각하고 있었음, gpu-cpu 간 이동이 발생하는 경우에 대해 사실은 모두 gpu 내에서 실행

딥러닝 프레임워크의 성능은 대부분 행렬 연산(GEMM), 합성곱(Conv2D), 그리고 그 뒤에 붙는 **에필로그 연산(bias, activation, residual add 등)**에서 결정됩니다. GPU 최적화는 크게 커널 내부 최적화와 커널 간 최적화(fusion) 두 가지 층위에서 이뤄집니다.

 

1. 커널 내부 최적화 기법

1.1 타일링(Tiling)

• 문제: Naïve 구현은 각 스레드가 같은 데이터를 반복적으로 global memory에서 읽음 → 대역폭 병목.
• 해결: 데이터를 타일 단위로 shared memory에 로드하고 여러 스레드가 재사용.
• 효과: DRAM 접근량 ↓, 산술강도(FLOPs/byte) ↑.

1.2 메모리 접근 최적화

• Coalesced Access: 스레드들이 연속된 주소를 읽도록 레이아웃 정렬.
• Vectorized Load/Store: float4, __half2 단위 접근으로 메모리 트랜잭션 효율 개선.
• Shared Memory Padding: bank conflict 방지.

1.3 레지스터/스레드 활용

• Register Blocking: 각 스레드가 작은 sub-tile을 레지스터에 보관 → 반복적 사용.
• Occupancy Tuning: block 크기와 레지스터 사용량을 균형 있게 잡아 SM 활용 극대화.
• Loop Unrolling: K dimension 전개로 ILP(Instruction-Level Parallelism) 확보.

1.4 파이프라이닝 & 비동기 전송

• Double Buffering: 다음 타일을 가져오는 동안 현재 타일 연산 진행.
• cp.async (Ampere+): global→shared 복사를 비동기로 수행, load/compute overlap.

1.5 Tensor Core 활용

• WMMA API (mma.sync, ldmatrix): FP16/BF16/TF32를 사용해 행렬 곱을 텐서코어에서 처리.
• 효과: 같은 연산량 대비 수십 배 더 높은 throughput.

 

2. 커널 간 최적화 (Fusion)

2.1 에필로그(Epilogue) 융합

• 문제: GEMM → bias add → activation을 각각 별도 커널로 실행하면 global mem에 쓰고 다시 읽어야 함.
• 해결: GEMM 내에서 누산(accumulator)을 레지스터 상태에서 곧바로 acc = act(acc + bias) 처리 후 저장.
• 효과:
  • 메모리 트래픽 절반 이하로 감소.
  • 커널 런치 횟수 감소.
  • 레지스터/SM 지역성 활용 극대화.

2.2 전/후처리 융합

• Conv2D + bias + BN + ReLU 같은 연산을 하나의 커널에서 처리.
• Dropout, Residual Add도 epilogue 단계에서 결합 가능.

2.3 장단점

• 장점: 메모리 접근 감소, 속도 향상.
• 단점: 레지스터 압박으로 occupancy가 떨어질 수 있음. 필요에 따라 fuse granularity를 조절해야 함.

 

3. 성능 분석 방법

3.1 Nsight Compute / Systems

• 커널별 FLOPs, DRAM Throughput, L2 Hit Rate, Shared Mem Throughput 측정.
• Warp Stall 이유(Memory Dep., Sync, Pipe Busy 등) 확인.

3.2 Roofline Model

• **Arithmetic Intensity (FLOPs/Byte)**를 기준으로, 연산이 메모리 바운드인지 compute 바운드인지 확인.
• 타일링/에필로그 융합이 산술 강도를 얼마나 끌어올렸는지 체크.

 

4. 최적화 체크리스트

• Global memory 접근은 Coalesced 되어 있는가?
• Shared memory에 타일링해서 재사용하고 있는가?
• 레지스터 blocking 및 occupancy 균형이 맞는가?
• cp.async 같은 비동기 로드로 load/compute overlap 했는가?
• Tensor Core를 활용했는가 (FP16/BF16/TF32)?
• Bias/Activation 등 후처리를 epilogue에서 융합했는가?
• Nsight로 DRAM bytes, FLOPs, occupancy를 측정했는가?

---

5. 결론

GPU 연산 최적화는 크게 두 층위에서 이뤄집니다:

1. 커널 내부: 메모리 계층(shared, register) 최적 활용, 타일링, 파이프라이닝, 텐서코어 사용.
2. 커널 간: 불필요한 메모리 왕복을 없애기 위한 epilogue 융합과 연산 fusion.

즉, “같은 GPU” 위에서도 구현 방식 차이만으로 성능은 수 배 이상 달라질 수 있습니다.