2018년 국제 컴퓨터 아키텍처 심포지엄(ISCA)에서 존 헤네시와 데이비드 패터슨은 튜링 강연 '컴퓨터 아키텍처의 새로운 황금기'를 발표했다. 두 사람이 튜링상 수상 연구를 하던 1980년대에는 단일 스레드 CPU 성능이 매년 52%씩 늘었지만, 무어의 법칙과 데나드 스케일링이 한계에 부딪힌 2018년에는 그 증가율이 3%로 주저앉았다. 범용 프로세서의 성능 향상이 사실상 멈춘 상황에서, 이들은 특정 영역에 특화된 아키텍처(DSA)가 필요하다고 주장했다. 대표 사례로 든 것이 이미 구글이 상용화한 TPU v1으로, 신경망 추론에서 CPU 대비 29배의 처리량과 80배 나은 에너지 효율을 보였다. 강연은 '앞으로 10년간 새로운 컴퓨터 아키텍처의 캄브리아기 대폭발이 일어날 것'이라는 예측으로 끝맺었다.
그 예측은 현실이 됐다. 오늘날 GPU, TPU, LPU, NPU, DPU, ASIC, 웨이퍼 스케일 엔진, 재구성 가능한 데이터플로, 뉴로모픽, 포토닉, 아날로그까지 수십 종의 아키텍처가 진지하게 개발되고 있다. 다만 실제 대규모 배치에 성공한 것은 아직 소수다. NVIDIA와 AMD의 GPU, TPU와 Trainium 같은 시스톨릭 어레이 가속기, Cerebras의 웨이퍼 스케일 엔진, 그리고 Groq의 LPU 정도다. 현재 NVIDIA가 명확한 선두이고 AMD가 뒤를 잇는데, OpenAI와 Meta 양쪽에서 각각 6GW 규모의 도입 약속을 받았다. TPU는 제미나이 학습에 쓰이며 Anthropic에 최대 100만 개 칩을 공급할 예정이고, Anthropic은 Claude를 100만 개 이상의 Trainium 칩에서 돌리고 있다. Cerebras는 OpenAI의 추론을 맡았고, Groq LPU는 200억 달러 규모의 인수를 통해 NVIDIA로 흡수됐다.
문제의 본질은 데이터 이동이다
AI 연산은 결국 행렬 곱셈이 지배한다. 트랜스포머는 Q/K/V 투영, 어텐션, 출력 투영, FFN으로 이어지는 행렬곱의 연속이며, 그 사이사이에 정규화·활성화·잔차 덧셈 같은 원소별 연산이 끼어든다. 프런티어 모델 하나를 학습하는 데는 10의 25제곱 회에 달하는 곱셈-누산 연산이 필요하다. 중요한 것은 이 행렬곱의 '모양'이 작업에 따라 달라진다는 점이다. 학습과 프리필(프롬프트 입력을 한 번에 통과시키는 단계)은 수천 개의 토큰을 같은 가중치 행렬에 쌓아 올리므로 큰 행렬-행렬 곱(GEMM)이 되고, 연산 강도가 높아 연산 자원에 병목이 걸린다. 반면 디코드는 한 번에 토큰 하나씩만 자기회귀적으로 생성하기 때문에 행렬-벡터 곱(GEMV)으로 전락하고, 토큰 하나를 뽑기 위해 모델의 모든 가중치와 KV 캐시를 통째로 읽어야 한다. 연산 강도가 프리필 대비 수십~수백 배 떨어진다.
추론 시스템은 이 낮아진 강도를 여러 사용자의 디코드 단계를 묶는 연속 배칭, 초안 토큰을 한 번에 검증하는 추측 디코딩, 모델 내부에서 여러 토큰을 예측하는 기법으로 다시 끌어올린다. GEMV를 GEMM으로 승격시켜 행렬곱 유닛의 활용률을 높이는 것이다. 다만 각 사용자가 자기 KV 캐시를 읽어야 하므로, 긴 문맥의 디코드는 가중치 대역폭이 아니라 KV 대역폭에 병목이 걸린다. 결국 모든 아키텍처가 씨름하는 것은 하나의 문제다. 연산 성능은 지수적으로 늘었지만 메모리 대역폭은 그러지 못한, 이른바 '메모리 장벽'이다. 칩을 이해한다는 것은 데이터가 어디에 있고, 어떻게 연산 유닛으로 이동하며, 연산 유닛은 어떻게 생겼고, 칩들이 대규모로 어떻게 통신하는가라는 네 가지 질문으로 압축된다.
NVIDIA GPU: 프로그래밍 모델을 지키며 진화하다
NVIDIA의 철학은 수천 개의 스레드를 CUDA로 노출하고 호스트 CPU가 조율하는 프로그래머블 칩이 병렬 작업에 가장 적합하다는 것이다. 같은 칩이 트랜스포머를 학습하고 추론을 서비스하며 그래픽과 과학 시뮬레이션까지 돌린다. GPU의 핵심은 처리량 지향 코어인 스트리밍 멀티프로세서(SM)로, V100의 80개에서 A100 108개, H100 132개, B200 148개, B300 160개, 차세대 Rubin의 224개까지 늘어왔다. SM 하나는 각자 워프 스케줄러와 레지스터 파일을 가진 네 개의 서브 파티션으로 구성되며, 32개 스레드가 SIMT 방식으로 보조를 맞춰 실행되는 워프 단위로 동작한다. 트랜스포머 블록의 FLOP 중 약 99%가 행렬곱이므로 실질적 연산 성능은 텐서 코어에서 나온다.
텐서 코어의 여섯 세대를 관통하는 흐름은 '행렬곱 명령이 스레드로부터 점점 분리되는' 과정이다. V100 1세대는 워프 32개 스레드가 함께 실행하고 완료까지 대기하는 동기식 16×16×16 FP16 연산이었다. A100은 TF32·BF16·구조적 희소성을, H100은 네이티브 FP8과 함께 128개 스레드가 비동기로 발사하는 wgmma를 도입해, 행렬곱이 백그라운드에서 도는 동안 워프가 다음 타일을 준비할 수 있게 했다. B200은 두 개의 SM에 걸친 256×256×16 연산, 네이티브 FP4, 그리고 누산기 타일을 레지스터 대신 전담하는 256KB 텐서 메모리(TMEM)를 더했다. Blackwell의 tcgen05.mma에 이르면 모든 피연산자가 공유 메모리와 TMEM으로 빠져나가, 스레드 하나가 명령을 발사하고 곧바로 복귀하며 나머지 워프는 다른 일을 할 수 있다. Rubin 6세대는 FP4 처리량을 확장하고 FP6을 더하며, 하드웨어에서 NVFP4 마이크로블록 스케일링을 처리하는 3세대 트랜스포머 엔진과 짝을 이룬다.
이 '분리'가 바로 GPU에서 어텐션 커널이 효율적인 이유다. 행렬곱이 워프를 붙잡지 않기 때문에 워프는 그 사이에 소프트맥스를 돌리거나 마스크를 적용하거나 다음 타일을 미리 불러올 수 있다. FlashAttention-3나 FA4 같은 현대 어텐션 커널은 모두 이 행렬곱과 원소별 연산의 겹침 위에 세워져 있다. 메모리 계층도 같은 방향으로 진화했다. Ampere의 cp.async는 레지스터를 거치지 않는 비동기 복사를, Hopper의 TMA는 스레드 하나가 다차원 타일 서술자만 제출하면 전용 DMA 엔진이 주소 계산을 도맡는 방식을 도입했고, Blackwell은 TMA가 TMEM으로 직접 적재하도록 확장했다. 세대를 거듭할수록 워프가 타일마다 해야 할 일이 하나씩 줄어드는 것, 그것이 NVIDIA가 프로그래밍 모델을 바꾸지 않으면서 메모리 장벽과 싸워온 방식이다.
한국 실무자 관점에서 이 흐름이 주는 시사점은 분명하다. 학습·프리필·디코드는 서로 다른 병목을 가지므로 하드웨어 선택과 배칭 전략을 워크로드별로 나눠 최적화해야 한다는 것, 그리고 NVIDIA의 아키텍처 진화가 프로그래밍 모델의 연속성을 유지한 채 이뤄져 왔기에 CUDA 생태계의 관성이 여전히 강력하다는 점이다. 다만 이 글이 다룬 것은 아직 GPU 중심의 조망이며, TPU·Trainium·Cerebras·Groq가 택한 서로 다른 스케일업·스케일아웃 및 소프트웨어 스택의 세부는 별도의 깊은 검토가 필요한 영역으로 남는다.