TL;DR
RTX 5090 등 소비자용 GPU에서 cuBLAS가 비효율적인 커널을 사용하여 성능이 저하되는 현상을 분석하고, 이를 해결한 커스텀 TMA 커널을 제시했다.
배경
RTX 5090에서 FP32 배치 연산 시 cuBLAS가 가용 연산 성능의 약 40%만 사용하는 문제를 발견하고, 이를 해결하기 위해 직접 작성한 커스텀 커널과 벤치마크 결과를 공유했다.
의미 / 영향
이 토론은 최신 하이엔드 소비자용 GPU인 RTX 5090조차 소프트웨어 라이브러리 최적화 미비로 인해 제 성능을 내지 못하고 있음을 시사한다. 개발자들은 표준 라이브러리에만 의존하기보다 TMA와 같은 최신 하드웨어 기능을 직접 활용하는 커스텀 최적화의 중요성을 재확인했다.
커뮤니티 반응
NVIDIA가 소비자용 GPU의 성능을 의도적으로 제한하거나 최적화를 방치하고 있다는 의혹에 대해 활발한 토론이 진행되었으며, 작성자의 커스텀 커널 성능에 대해 긍정적인 반응이 많았다.
주요 논점
cuBLAS가 소비자용 GPU에서 성능을 제대로 내지 못하고 있으며, 이는 데이터로 증명된 사실이다.
단순 버그일 가능성이 있으며, 향후 CUDA 업데이트를 통해 해결될 수 있는지 지켜봐야 한다.
합의점 vs 논쟁점
합의점
- RTX 5090에서 cuBLAS의 현재 성능은 하드웨어 잠재력에 비해 현저히 낮다.
- TMA를 활용한 비동기 데이터 로딩이 최신 GPU 아키텍처에서 성능 향상의 핵심이다.
논쟁점
- 이러한 성능 저하가 NVIDIA의 의도적인 급 나누기(Segmentation) 전략인지 단순한 최적화 누락인지에 대한 의견이 갈린다.
실용적 조언
- RTX 5090 환경에서 대규모 FP32 배치 연산이 필요한 경우 cuBLAS 대신 TMA 기반의 커스텀 커널 사용을 고려해야 한다.
- 성능 병목 현상 분석 시 NCU 프로파일링과 SASS 스케줄링 확인을 통해 실제 하드웨어 활용률을 점검해야 한다.
섹션별 상세
__global__ __launch_bounds__(256) void fused_matmul( const __grid_constant__ CUtensorMap A_tma, const __grid_constant__ CUtensorMap B_tma, float* C) {
extern __shared__ __align__(128) char dsmem[];
// ... (중략)
for (int t = 0; t < K/BK; t++) {
int s = t % 2; // Current buffer
mbarrier_wait(mbarrier[s], phase[s]);
if (tid == 0 && t + 1 < nt) {
tma_load_2d(next_buf_a, &A_tma, next_k, bm, next_mbar);
tma_load_2d(next_buf_b, &B_tma, bn, next_k, next_mbar);
}
// Compute: all 256 threads do FMA from shared memory
// ... (중략)
}
}TMA와 이중 버퍼링을 사용하여 cuBLAS보다 높은 성능을 내는 커스텀 행렬 곱셈 커널 예시
용어 해설
- cuBLAS
- — NVIDIA가 제공하는 GPU 가속 BLAS(Basic Linear Algebra Subprograms) 라이브러리로, 행렬 연산과 같은 선형 대수 계산을 최적화하여 수행하는 핵심 소프트웨어 구성 요소이다.
- TMA
- — Tensor Memory Accelerator의 약자로, NVIDIA Hopper 및 Blackwell 아키텍처에서 도입된 하드웨어 기능이며 전역 메모리와 공유 메모리 간의 데이터 전송을 비동기적으로 처리하여 연산 효율을 높인다.
- SASS
- — NVIDIA GPU에서 실제로 실행되는 저수준 어셈블리 명령어로, 컴파일된 CUDA 코드가 하드웨어 스케줄러에 의해 어떻게 처리되는지 분석할 때 사용되는 최종 기계어 형태이다.
- Double Buffering
- — 데이터를 처리하는 동안 다음 데이터를 미리 로드하기 위해 두 개의 버퍼를 교대로 사용하는 기법으로, 메모리 대기 시간을 연산 시간과 중첩시켜 전체 처리량을 극대화하는 최적화 방식이다.
- FMA
- — Fused Multiply-Add의 약자로, 곱셈과 덧셈을 하나의 명령어로 처리하는 연산 방식이며 부동 소수점 연산의 정밀도를 높이고 실행 속도를 단축시키는 핵심 연산 단위이다.
언급된 도구
NVIDIA GPU용 선형 대수 라이브러리
병렬 컴퓨팅 플랫폼 및 프로그래밍 모델
NVIDIA Nsight Compute 프로파일링 도구
AI 요약 · 북마크 · 개인 피드 설정 — 무료
출처 · 인용 안내
인용 시 "요약 출처: AI Trends (aitrends.kr)"를 표기하고, 사실 확인은 원문 보기 기준으로 진행해 주세요. 자세한 기준은 운영 정책을 참고해 주세요.