🔢 mathematics

Accurate Models of NVIDIA Tensor Cores

이 논문은 V100, A100, H100, B200 등 다양한 NVIDIA 데이터센터 GPU 의 행렬 곱셈 연산에서 발생하는 IEEE 754 비준수 및 비재현성 문제를 해결하기 위해, 8/16/19 비트 부동소수점 포맷을 지원하는 정확한 소프트웨어 에뮬레이션 모델을 제안합니다.

원저자: Faizan A. Khattak, Mantas Mikaitis

게시일 2026-04-07

📖 3 분 읽기🧠 심층 분석

CC BY 4.0

원저자: Faizan A. Khattak, Mantas Mikaitis

원본 논문은 CC BY 4.0 (http://creativecommons.org/licenses/by/4.0/) 라이선스로 제공됩니다. ✨ 이것은 아래 논문에 대한 AI 생성 설명입니다. 저자가 작성하거나 승인한 것이 아닙니다. 기술적 정확성을 위해서는 원본 논문을 참조하세요. 전체 면책 조항 읽기

1. 배경: 왜 이 연구가 필요한가요? (비밀스러운 요리사들)

현대 인공지능 (AI) 은 엄청난 양의 데이터를 처리해야 합니다. 이때 GPU 는 마치 초고속 주방처럼 작동하며, 수천 개의 재료를 한 번에 섞는 '행렬 곱셈' 작업을 하드웨어로 빠르게 처리합니다. 이를 '텐서 코어 (Tensor Core)'라고 부릅니다.

하지만 문제는 이 '초고속 주방'들이 규칙을 제각기 다르게 따른다는 점입니다.

IEEE 754 표준: 일반적인 컴퓨터 계산은 엄격한 국제 표준 (예: 소수점 반올림 규칙) 을 따릅니다.
NVIDIA 의 텐서 코어: 속도를 위해 이 표준을 일부 무시합니다. 예를 들어, 아주 작은 수를 어떻게 다룰지, 중간 계산 결과를 어떻게 반올림할지, 어떤 칩 (V100, A100, H100 등) 에 따라 요리사의 손맛이 다릅니다.

이 때문에 같은 계산을 다른 세대 GPU 에서 하면 결과값이 미세하게 달라질 수 있습니다. 과학자들은 "어? 이 결과가 왜 달라졌지?"라고 혼란을 겪고, AI 모델이 예측하지 못한 오류를 일으킬 수도 있습니다.

2. 문제: 왜 기존 방법으로는 안 됐나요? (눈대중으로 요리하기)

기존 연구자들은 "어떤 입력값을 넣으면 어떤 결과가 나오는지"를 테스트하며 GPU 의 규칙을 추측했습니다. 하지만 이는 눈대중에 가깝습니다.

테스트할 수 있는 경우의 수가 너무 많아서 (수억 가지), 모든 경우를 다 확인하기 어렵습니다.
그래서 만든 모델이 실제 GPU 의 '손맛'과 100% 일치하지 않아, 과학적 실험에서 재현성이 떨어졌습니다.

3. 해결책: 완벽한 요리 레시피 책 만들기 (이 논문의 핵심)

저자들은 **실제 GPU 와 100% 똑같이 작동하는 소프트웨어 모델 (MATLAB 툴박스)**을 만들었습니다. 이를 위해 두 가지 단계를 거쳤습니다.

1 단계: 정밀한 테스트로 '손맛' 파악하기 (GNFT)

먼저, **특수한 테스트 재료 (테스트 벡터)**를 만들어 GPU 에 넣었습니다.

예: "아주 작은 수를 넣었을 때, 이 주방은 버리나요? 아니면 아주 작은 수로 만들어서 섞나요?"
이를 통해 각 GPU 모델 (V100, A100, H100, B200 등) 이 가진 **고유한 특징 (반올림 방식, 작은 수 처리법 등)**을 찾아냈습니다.

2 단계: 무작위 시뮬레이션으로 검증하기 (ISSM)

이제 찾은 특징을 바탕으로 소프트웨어 모델을 만들었습니다. 하지만 이게 맞는지 확인하기 위해 **수천만 개의 무작위 재료 (입력값)**를 만들어 실제 GPU 와 소프트웨어 모델에 동시에 요리하게 했습니다.

결과가 하나도 안 맞으면? → 모델의 '손맛'을 수정합니다.
결과가 완벽하게 일치하면? → 이 모델은 이제 실제 GPU 와 똑같습니다.

이 과정을 반복해서, 비트 단위 (Bit-level) 로 100% 일치하는 모델을 완성했습니다.

4. 주요 발견: 각 세대 GPU 의 특징 (요리사들의 성향)

논문을 통해 밝혀진 흥미로운 사실들입니다:

V100 (구형): 아주 작은 수를 다룰 때, 중간에 버리는 방식이 특이했습니다.
A100/H100 (신형): 작은 수를 다룰 때, '보조 비트 (Extra bits)'를 더 많이 사용하여 정밀도를 높였습니다. 하지만 아주 작은 수 (비정규수) 가 너무 작으면 아예 0 으로 만들어버리는 '한계점'이 있습니다.
B200 (최신): 가장 정밀하지만, 특정 입력 방식 (fp8) 을 쓸 때는 재료를 섞는 순서가 다른 GPU 와 달랐습니다 (교차 섞기 방식).

이처럼 같은 '텐서 코어'라는 이름이라도, 세대마다 요리하는 방식이 미세하게 다릅니다.

5. 이 모델이 주는 혜택 (왜 중요한가요?)

이제 연구자들은 실제 비싼 GPU 를 사지 않아도 MATLAB 에서 이 모델을 돌려볼 수 있습니다.

실험 전 검증: "이 알고리즘을 이 GPU 에서 돌리면 결과가 어떨까?"를 미리 시뮬레이션해 볼 수 있습니다.
정밀한 분석: "왜 이 결과가 나왔지?"라고 의문을 품었을 때, 모델이 그 이유를 정확히 설명해 줍니다.
고정밀 계산: 낮은 정밀도의 GPU 를 이용해 마치 고정밀 계산처럼 결과를 내는 '다중 단어 (Multi-word)' 알고리즘을 개발할 때, 이 모델이 필수적입니다.

6. 결론: 요약

이 논문은 NVIDIA GPU 의 '숨겨진 계산 규칙'을 완벽하게 해독하여, 누구나 쉽게 사용할 수 있는 '가상 GPU'를 만들었다는 것입니다.

마치 실제 주방의 모든 요리사 (GPU) 의 손맛을 분석하여, 그들과 똑같이 요리를 할 수 있는 '완벽한 레시피 책 (소프트웨어 모델)'을 출판한 것과 같습니다. 이제 과학자들은 이 책을 보고, 실제 장비 없이도 AI 와 과학 계산의 정확성을 높일 수 있게 되었습니다.

한 줄 요약:

"NVIDIA GPU 가 속도를 위해 숨겨둔 계산 규칙을 모두 찾아내어, 실제 하드웨어와 100% 똑같이 작동하는 소프트웨어 모델을 만들어 누구나 쉽게 실험할 수 있게 했습니다."

논문 제목: NVIDIA Tensor Cores 의 정밀 모델링 (Accurate Models of NVIDIA Tensor Cores)

1. 문제 제기 (Problem)

비표준적 연산 행동: AI 및 과학적 컴퓨팅에서 행렬 곱셈을 가속화하기 위해 NVIDIA GPU 의 Tensor Core(및 AMD, Intel 의 유사한 단위) 가 널리 사용되고 있습니다. 그러나 이러한 하드웨어는 IEEE 754 부동소수점 표준을 완전히 준수하지 않습니다.
재현성 부족: 벤더와 GPU 아키텍처 세대 (V100, A100, H100, B200 등) 에 따라 반올림 (rounding) 방식, 정규화 지점, 누적기 (accumulator) 폭, 추가 캐리 비트 등 수치적 특성이 다릅니다. 이로 인해 동일한 코드가 다른 하드웨어에서 실행될 때 비트 단위 (bit-level) 로 다른 결과가 도출되어 재현성이 보장되지 않습니다.
검증의 어려움: 기존 연구들은 테스트 벡터를 통해 하드웨어의 특성을 추정하려 했으나, 제한된 하드웨어 접근성과 불완전한 테스트로 인해 다양한 플랫폼에서의 신뢰성이 검증되지 않았습니다. 또한, 기존 모델들은 실제 하드웨어의 미세한 수치적 특성을 완벽하게 반영하지 못해 정밀한 오차 분석이 어려웠습니다.

2. 방법론 (Methodology)

저자들은 MATLAB 기반의 소프트웨어 모델을 개발하여 Tensor Core 의 내부 동작을 비트 단위 정확도로 시뮬레이션하는 새로운 접근법을 제시했습니다. 이 과정은 두 단계로 이루어집니다.

일반화된 수치 특성 테스트 (GNFT, Generalised Numerical Feature Testing):
- 하드웨어의 특정 수치적 특성 (예: 서브노멀 입력 지원 여부, 추가 정렬 비트 존재 여부 등) 을 트리거하는精心하게 설계된 테스트 벡터를 생성합니다.
- 이를 통해 초기 모델의 수치적 특성 (정규화 지점, FMA 블록 크기, 반올림 모드 등) 을 추정합니다.
입력 공간 탐색 방법 (ISSM, Input Space Search Method):
- GNFT 로 추정된 초기 모델과 실제 GPU 하드웨어 간의 오차를 찾기 위해 대규모 무작위 입력 벡터 (약 $10^7$ 개) 를 생성하여 비교합니다.
- 불일치가 발견되면, 해당 사례를 분석하여 모델의 파라미터 (예: 정렬 비트 수, 누적기 폭, 반올림 로직) 를 수정하는 반복적 정제 (Iterative Refinement) 과정을 거칩니다.
- 이 과정은 모델이 하드웨어와 비트 단위 (bit-identical) 로 일치할 때까지 반복됩니다.

3. 주요 기여 (Key Contributions)

다양한 GPU 아키텍처의 정밀 모델 개발:
- V100, A100, A2, A30, H100, H200, L40S, RTX 1000 (Ada), B200 등 9 가지 GPU 모델에 대한 31 개의 정밀 모델을 MATLAB 에서 구현했습니다.
- 8 비트 (fp8), 16 비트 (fp16, bf16), 19 비트 (tf19) 등 다양한 혼합 정밀도 (mixed-precision) 입력 포맷을 지원합니다.
새로운 수치적 특성의 발견 및 문서화:
- 정렬 비트 (Alignment Bits): 이전 연구보다 정밀하게 추가 정렬 비트 ( $n_{eab}$ ) 의 수를 규명했습니다 (예: H100/B200 의 fp16/bf16 입력 시 2 비트).
- 지수 제한 (Exponent Limitation): 서브노멀 영역에서 정렬 시 사용되는 최대 지수 값이 하드웨어에 따라 제한됨을 발견했습니다 (예: A100/H100 에서 $-133 $또는$ -126$으로 고정).
- 비정규화 누적 (Denormalized Accumulation): 중간 곱셈 결과들이 정규화되지 않은 채로 누적되는 동작을 규명했습니다.
- 인터리빙 패턴 (Interleaved Pattern): H100/H200/B200 에서 fp8 입력을 처리할 때, 입력 벡터가 두 개의 Tensor Core 호출에 교차하여 분배되는 인터리빙 패턴을 발견했습니다.
오픈 소스 MATLAB 툴박스 (MATLAB Tensor Core v0.4.1) 제공:
- GitHub 를 통해 공개된 이 툴박스는 실제 하드웨어 없이도 Tensor Core 의 행동을 시뮬레이션할 수 있게 합니다.
- 사용자 정의 파라미터 (FMA 크기, 정렬 비트, 반올림 모드 등) 를 통해 커스텀 Tensor Core 모델을 생성할 수 있는 기능을 제공합니다.

4. 결과 (Results)

비트 단위 정확도 달성: 무작위 테스트 벡터 ( $10^7$ 개 이상) 를 통한 검증에서 개발된 MATLAB 모델은 실제 GPU 하드웨어의 출력과 비트 단위 (bit-level) 로 완전히 일치함을 확인했습니다.
모델 비교 및 분석:
- V100: 4 개의 FMA 블록 크기 ( $N_{FMA}=4$ ), 23 비트 분수 정렬.
- A100: 8 개의 FMA 블록, 1 비트 추가 정렬 ( $n_{eab}=1$ ), tf19 입력 시 4 개 FMA.
- H100/H200/B200: 16 개 (fp16/bf16) 또는 8 개 (tf19) FMA 블록, 2 비트 추가 정렬 ( $n_{eab}=2$ ). fp8 입력 시 인터리빙 처리 및 32 개 FMA 블록 (wgmma 명령어 사용 시) 동작 규명.
- L40S/Ada RTX 1000: A100 과 유사하지만 fp8 입력 시 13 비트 분수 정렬을 사용.
응용 사례:
- 테스트 벡터 인증: 기존 연구에서 제안된 테스트 벡터들이 특정 수치적 특성 (예: $n_{eab}$ ) 을 고려하지 않아 잘못된 결론을 내릴 수 있음을 모델 시뮬레이션을 통해 증명했습니다.
- 고정밀도 GEMM 에뮬레이션: 저정밀도 Tensor Core 를 사용하여 고정밀도 행렬 곱셈을 에뮬레이션하는 'Multi-word' 알고리즘의 오차 분석을 수행했습니다. 모델에 따라 오차 크기가 크게 달라질 수 있음을 보여주었습니다 (예: B200 모델의 반올림 방식에 따른 오차 차이).

5. 의의 및 중요성 (Significance)

과학적 컴퓨팅의 재현성 확보: AI 와 과학적 계산 분야에서 하드웨어 종속적인 수치적 오차를 이해하고 예측할 수 있는 기반을 마련했습니다.
혼합 정밀도 알고리즘 개발 지원: 개발자들이 실제 하드웨어를 보유하지 않더라도 정밀한 수치 시뮬레이션을 통해 알고리즘을 검증하고 최적화할 수 있게 되었습니다.
표준화 기여: IEEE 754 및 Open Compute Project (OCP) 와 같은 표준화 기구들이 저정밀도 연산의 표준을 제정하는 데 필요한 구체적인 데이터와 통찰력을 제공합니다.
미래 연구의 기반: 제공된 MATLAB 툴박스는 향후 Tensor Core 아키텍처의 변화에 대응하고, 새로운 혼합 정밀도 알고리즘을 연구하는 데 필수적인 도구로 활용될 것입니다.

이 논문은 단순히 하드웨어의 동작을 모방하는 것을 넘어, 왜 그와 같은 동작이 발생하는지 (정렬 비트, 정규화 지점, 인터리빙 등) 를 체계적으로 규명하고 이를 소프트웨어 모델로 구현함으로써, 신뢰할 수 있는 수치 컴퓨팅 환경을 조성하는 데 중요한 기여를 했습니다.