GPU 아키텍처 #6: Ampere - Sparsity 가속과 MIG

시리즈 로드맵

Turing 이후의 과제

Turing은 소비자 GPU에 Tensor Core와 RT Core를 처음 가져왔다. 그러나 데이터센터 시장에서는 Volta V100이 그대로 유지됐다. 2018년부터 2020년 사이 LLM 연구가 급격히 커졌다. GPT-2(2019), GPT-3(2020) 수준의 모델을 학습하려면 이전과 다른 규모의 가속이 필요했다.

동시에 AI 추론 압축 연구가 성숙했다. 모델 가중치의 상당 비율을 0으로 만들어도 정확도 손실이 미미하다는 사실이 실험으로 검증됐다. 하드웨어가 이 희소성(sparsity)을 직접 가속할 수 있다면, 추론 처리량을 두 배로 늘릴 수 있었다.

Ampere(2020)는 세 가지로 이 문제를 공략했다. 3세대 Tensor Core의 2:4 구조적 희소성 가속, TF32와 BF16이라는 새 정밀도 형식, A100에서 하드웨어로 GPU를 분리하는 MIG(Multi-Instance GPU).

다이 스펙: 두 갈래의 Ampere

Ampere는 공정과 타깃 시장에 따라 두 개의 서로 다른 다이로 나뉜다.

GA100은 TSMC 7nm로 공정 세대가 올라갔다. GA102는 수율 확보를 위해 Samsung 8nm를 택했다. Pascal과 Turing이 각각 GP100(고성능)과 GP102/TU102(소비자)로 나뉜 것처럼, Ampere도 두 갈래로 분기했다.

3세대 Tensor Core

TF32: 코드 변경 없는 가속

TF32(TensorFloat-32)는 Ampere에서 처음 도입된 19비트 데이터 형식이다.

FP32 (32비트):  부호 1 | 지수 8 | 가수 23
TF32 (19비트):  부호 1 | 지수 8 | 가수 10   ← FP32와 지수 범위 동일
FP16 (16비트):  부호 1 | 지수 5 | 가수 10
BF16 (16비트):  부호 1 | 지수 8 | 가수 7

TF32는 FP32와 동일한 지수 범위(오버플로/언더플로 없음)를 유지하면서, 가수를 10비트로 줄여 FP16 수준의 정밀도를 갖는다. cuBLAS와 cuDNN은 A100에서 TC GEMM에 TF32를 기본으로 사용한다. FP32로 작성된 기존 학습 코드를 수정하지 않아도 A100에서 자동으로 TF32 가속이 적용된다.

# TF32 비활성화 (정밀도 검증 목적)
NVIDIA_TF32_OVERRIDE=0 python train.py

BF16 Tensor Core

BF16(Brain Float 16)은 Google Brain에서 제안한 형식이다. FP32와 동일한 8비트 지수를 유지하므로 FP16의 지수 오버플로 문제가 없다. 분산 학습에서 gradient의 dynamic range가 넓어도 안전하다. Turing은 BF16을 CUDA Core에서 소프트웨어로만 처리했다. Ampere는 BF16을 Tensor Core에 직접 통합했다.

FP64 Tensor Core (A100 전용)

GA100(A100)은 FP64 Tensor Core를 처음 도입했다. FP64 GEMM을 전용 TC 회로에서 처리해, FP64 CUDA Core 단독 처리 대비 2배 처리량을 얻는다.

과학 시뮬레이션, 기후 모델링 등 FP64 정밀도가 필수인 HPC 워크로드가 주요 대상이다.

2:4 구조적 희소성: Sparsity 가속

3세대 Tensor Core의 가장 독창적인 추가는 2:4 구조적 희소성(2:4 Structured Sparsity)이다. 연속된 4개의 값 중 정확히 2개가 0인 패턴을 하드웨어가 직접 인식하고 건너뛴다.

원래 가중치 행 (8개 값):
  [0.3, 0.0, -0.5, 0.0, 0.8, 0.0, 0.2, 0.0]
      ↑    ↑                   ↑        ↑ (0으로 고정)

2:4 패턴 압축 저장:
  값:    [0.3, -0.5,  0.8,  0.2]        (50% 압축)
  인덱스: [0, 2,      0, 2]             (4-그룹 내 위치, 2비트씩)

추론 시 Tensor Core는 압축 값과 인덱스를 함께 로드해, 0 곱셈을 건너뛰며 동일한 결과를 생성한다. 연산량이 절반으로 줄면서 처리량이 2배가 된다.

희소성 처리 흐름:

학습 (sparsity 적용):
  일반 FP32/FP16 학습 → ASP 라이브러리로 2:4 패턴 pruning → fine-tuning → 정확도 복원

추론 (sparse inference):
  압축 가중치 로드 → Sparse Tensor Core가 dense output 생성 → 2× 처리량

실제로 ResNet-50, BERT, GPT 등 주요 모델에서 2:4 sparsity를 적용해도 정확도 손실이 0.5% 미만으로 보고됐다.

처리량 비교

Ampere SM 구조

GA100과 GA102의 갈림

두 다이는 동일한 Tensor Core 세대를 쓰지만 SM 구성이 다르다.

GA100 Sub-core (A100)                  GA102 Sub-core (RTX 3090)
─────────────────────────────          ──────────────────────────────────
Warp Scheduler                         Warp Scheduler
Dispatch Unit ×2                       Dispatch Unit ×2
16 FP32           (전통)               16 FP32           (전통)
 8 FP64           (HPC용)              16 FP32/INT32     ← 이중 모드 (신규)
16 INT32          (별도)               [FP64 없음]
 1 TC (3세대)                           1 TC (3세대)
 8 LD/ST   4 SFU                        8 LD/ST   4 SFU

SM 합계 (4 Sub-core):                  SM 합계 (4 Sub-core):
  64 FP32 + 32 FP64 + 64 INT32           128 FP32 (+ 선택적 INT32) + 미미한 FP64

GA102의 핵심 변화는 INT32 파이프라인이 FP32도 실행할 수 있게 된 것이다. 클록당 FP32 연산이 Turing TU102의 2배다. 단, FP32와 INT32를 동시에 실행할 수는 없다. 한 사이클에 FP32 전용 64개 + FP32/INT32 듀얼 64개, 총 128개 FP32 연산이 가능하다.

Volta/Turing/Ampere 비교

GA100(A100)은 데이터센터 전용으로 RT Core가 없다. GA102(RTX 30xx)는 2세대 RT Core를 탑재했다. Turing 1세대 대비 BVH 처리량이 향상됐다. GA100의 최대 occupancy는 GV100과 동일하게 64 warps/SM이며, GA102는 48 warps/SM으로 중간값이다.

GA100은 Unified L1+Shared를 192KB까지 늘렸다. 대형 행렬 연산에서 타일을 Shared Memory에 최대한 담을수록 전역 메모리 접근이 줄어들기 때문이다.

MIG: Multi-Instance GPU

MIG(Multi-Instance GPU)는 A100 전용 기능이다. 하나의 물리 GPU를 최대 7개의 독립적인 GPU 인스턴스로 분리한다.

A100 SXM4 (108 SM 전체)
┌─────────────────────────────────────────────────────────┐
│  GPC 0 (14 SM) │ GPC 1 (14 SM) │ ... │ GPC 7 (14 SM)  │
│  L2 슬라이스   │ L2 슬라이스   │     │ L2 슬라이스    │
│  HBM2e 파티션  │ HBM2e 파티션  │     │ HBM2e 파티션   │
└─────────────────────────────────────────────────────────┘

MIG 7개 인스턴스 (각 1/7 A100):
  ┌──────┐ ┌──────┐ ┌──────┐ ┌──────┐ ┌──────┐ ┌──────┐ ┌──────┐
  │ 14SM │ │ 14SM │ │ 14SM │ │ 14SM │ │ 14SM │ │ 14SM │ │ 14SM │
  │  5GB │ │  5GB │ │  5GB │ │  5GB │ │  5GB │ │  5GB │ │  5GB │
  └──────┘ └──────┘ └──────┘ └──────┘ └──────┘ └──────┘ └──────┘
  각 인스턴스: 독립 SM 집합 + 전용 L2 캐시 슬라이스 + HBM 대역폭 파티션

MIG의 격리 수준:

• SM: 각 인스턴스는 자신에게 할당된 SM만 사용한다.
• L2 캐시: L2를 슬라이스로 분할해 인스턴스 간 캐시 간섭 없음.
• 메모리 대역폭: HBM2e 대역폭도 파티션 단위로 보장.
• 에러: 한 인스턴스의 오류가 다른 인스턴스에 영향을 주지 않음.

MIG 인스턴스 크기 옵션(A100 40GB 기준):

클라우드 서비스 제공자가 A100 하나를 여러 사용자에게 제공하거나, 소형 추론 작업을 다수 병렬 실행할 때 활용된다.

NVLink 3.0과 메모리

NVLink 3.0

DGX A100 시스템은 8개의 A100을 NVLink 3.0 + NVSwitch 2.0으로 연결한다. GPU당 NVLink 총 대역폭은 600 GB/s, NVSwitch 2.0은 스위치당 7.2 Tb/s 비전이선 대역폭을 제공한다.

HBM2e와 GDDR6X

A100 SXM4 80GB의 2 TB/s는 V100 대비 2.2배다. GDDR6X는 PAM4 시그널링으로 동일 클록에 두 배의 데이터를 전송한다.

L2 캐시 파티셔닝

A100의 40MB L2는 V100 대비 6.7배다. 단순히 크기만 늘어난 것이 아니다. A100 다이는 SM 배열을 기준으로 두 개의 물리적으로 분리된 20MB L2 파티션으로 나뉜다. 각 파티션은 인접한 절반의 SM 집합에 귀속된다. 자신의 파티션에 있는 데이터는 낮은 지연으로 접근하고, 반대편 파티션 데이터는 내부 크로스바를 경유해 접근 비용이 높아진다.

이 구조가 MIG와 연결된다. MIG 인스턴스는 연속된 SM 블록과 해당 SM에 귀속된 L2 파티션 슬라이스를 함께 할당받는다. 인스턴스 간 L2 간섭이 하드웨어 수준에서 차단된다.

Ampere는 L2 캐시 상주 여부를 명시적으로 제어하는 API도 함께 도입했다.

// 영구 L2 영역 크기 설정 (예: 20MB 고정)
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, 20ULL * 1024 * 1024);

// 특정 버퍼를 영구적으로 L2에 유지 (커널 간에도 유지됨)
cudaStreamAttrValue attr = {};
attr.accessPolicyWindow.base_ptr  = (void*)weight_ptr;
attr.accessPolicyWindow.num_bytes = weight_size;
attr.accessPolicyWindow.hitRatio  = 1.0f;
attr.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;  // L2에 고정
attr.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;    // 미스 시 bypass
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);

cudaAccessPropertyPersisting으로 표시된 데이터는 커널 경계를 넘어 L2에 잔류한다. 반복적으로 접근하는 가중치 테이블이나 임베딩 테이블에 활용하면 HBM 왕복을 줄일 수 있다. cudaAccessPropertyStreaming은 반대로 해당 접근이 L2를 오염시키지 않도록 bypass 처리한다.

cp.async: 비동기 메모리 복사

Ampere는 cp.async PTX 명령어를 새로 도입했다. 전역 메모리에서 공유 메모리로 데이터를 복사할 때, 레지스터를 거치지 않고 직접 복사한다.

Turing 이전 (레지스터 경유):           Ampere cp.async (직접 복사):
Global Mem → 레지스터 → Shared Mem     Global Mem ──→ Shared Mem
(레지스터 소비, blocking)              (레지스터 없음, non-blocking)

cp.async는 비동기(non-blocking)다. 복사를 요청하고 즉시 다음 명령을 실행할 수 있다. 파이프라인을 명시적으로 구성해 현재 타일 계산과 다음 타일 데이터 적재를 겹친다.

#include <cuda/pipeline>

__global__ void tiled_matmul(float *A, float *B, float *C, int N) {
    __shared__ float smA[2][TILE][TILE];
    __shared__ float smB[2][TILE][TILE];

    cuda::pipeline<cuda::thread_scope_thread> pipe = cuda::make_pipeline();
    int stage = 0;

    // 첫 타일 프리패치
    cuda::memcpy_async(smA[stage], A + ..., sizeof(smA[0]), pipe);
    cuda::memcpy_async(smB[stage], B + ..., sizeof(smB[0]), pipe);
    pipe.producer_commit();

    for (int tile = 0; tile < N / TILE; tile++) {
        int next = 1 - stage;
        // 다음 타일 비동기 적재 (현재 타일 계산 중 진행)
        cuda::memcpy_async(smA[next], A + ..., sizeof(smA[0]), pipe);
        cuda::memcpy_async(smB[next], B + ..., sizeof(smB[0]), pipe);
        pipe.producer_commit();

        pipe.consumer_wait();  // 현재 타일 적재 완료 대기
        __syncthreads();

        // 현재 타일로 TC GEMM 실행
        wmma::mma_sync(c_frag, a_frag, b_frag, c_frag);

        pipe.consumer_release();
        stage = next;
    }
}

CUTLASS 3.x와 cuBLAS 내부 커널은 이 패턴을 광범위하게 사용한다. Shared Memory I/O와 Tensor Core 연산이 겹쳐 SM이 idle 상태로 대기하는 시간이 줄어든다.

스케줄링 변화

Ampere의 스레드 스케줄링은 Volta/Turing과 동일하다. 독립 스레드 스케줄링(per-thread PC)을 그대로 유지한다. 새로운 요소는 비동기 파이프라인이다.

cp.async의 완료 여부는 커밋(commit)과 대기(wait) 명시적 API로 관리한다. 스케줄러는 이를 일반 메모리 의존성처럼 처리한다: 복사가 완료되기 전 해당 Shared Memory를 읽으면 정의되지 않은 동작이다. pipe.consumer_wait()는 완료 보장을 제공한다.

CUDA 11과 소프트웨어

CUDA 11 (2020)

CUDA 11.0은 Ampere GA100 지원으로 출시됐다.

• cp.async / cuda::pipeline API
• TF32, BF16 WMMA fragment 타입
• cudaMemcpyAsync의 내부 경로가 cp.async를 활용
• cooperative_groups::memcpy_async() (Cooperative Groups 기반 복사)
• MIG 관리: nvidia-smi mig 명령어 세트

DLSS 2.0

DLSS 2.0은 Ampere와 함께 출시됐다. 1.0의 게임별 모델에서, 단일 범용 모델로 전환했다. 이전 프레임의 temporal 정보와 모션 벡터를 활용해 업스케일한다. Turing에서도 DLSS 2.0을 사용할 수 있다(Tensor Core가 있으면 동작).

A100과 대규모 언어 모델

A100 SXM4 80GB는 GPT-3(175B 파라미터)를 FP16으로 단일 GPU에 올릴 수 없는 한계를 보여줬다(필요 VRAM: 약 350 GB). 텐서 병렬(Tensor Parallelism)과 파이프라인 병렬(Pipeline Parallelism)을 조합해 DGX A100 한 노드(8×A100) 또는 다중 노드에 분산하는 것이 표준 방법이 됐다. NVLink 3.0의 600 GB/s가 이 통신을 뒷받침한다.

인터커넥트 및 외부 채널

PCIe 호스트 인터페이스

Ampere는 PCIe 4.0 x16을 NVIDIA GPU 최초로 채택했다. 이론 단방향 대역폭은 32 GB/s(PCIe 3.0의 2배)다.

NVENC / NVDEC

GA102(RTX 30xx)는 AV1 하드웨어 인코드를 최초로 지원하는 세대다. A100은 데이터센터 전용으로 NVENC, NVDEC, 디스플레이가 없다.

디스플레이 출력 (소비자 GPU 레퍼런스)

HDMI 2.1은 4K 120Hz, 8K 60Hz를 지원한다.

정리

Ampere는 두 가지 방향을 동시에 열었다. GA100은 FP64 TC, MIG, NVLink 3.0, 40MB L2, 2 TB/s HBM2e로 대규모 AI 학습과 HPC의 교차점을 겨냥했다. GA102는 128 FP32/SM과 2:4 sparsity로 소비자 게임·AI 추론 처리량을 Turing 대비 크게 끌어올렸다. 2:4 구조적 희소성은 이후 Hopper, Blackwell에서도 계속 이어지는 TC 가속의 핵심 기법이 됐다.

다음 포스트에서는 GPU 메모리 시스템을 다룬다: 레지스터 파일, Shared Memory, L1/L2 캐시, DRAM 계층의 구조와 최적화 기법.

References

• NVIDIA. NVIDIA A100 GPU Architecture Whitepaper, 2020. PDF
• NVIDIA. NVIDIA Ampere Architecture In-Depth. NVIDIA Developer Blog, 2020. Link
• NVIDIA. Exploiting NVIDIA Ampere Structured Sparsity with cuSPARSELt. NVIDIA Developer Blog, 2021.
• NVIDIA. Programming Tensor Cores in CUDA 9. NVIDIA Developer Blog, 2017.
• NVIDIA. New Features in CUDA 11.1. NVIDIA Developer Blog, 2020.
• NVIDIA. Multi-Instance GPU User Guide. NVIDIA Documentation.
• NVIDIA. CUDA C++ Programming Guide: Asynchronous Data Copies. NVIDIA Documentation.