AEGLOSLAB 엔비디아 기술 스택

제2부 CUDA 프로그래밍

2.12 컴파일 흐름 — nvcc·PTX·SASS·JIT

nvcc hello.cu -o hello 뒤에서는 호스트 컴파일러와 디바이스 컴파일러가 갈라지고, 커널은 PTX(가상 ISA)를 거쳐 SASS(실제 머신코드)로 내려가며, 어떤 아키텍처용 코드를 바이너리에 넣을지 결정된다. 이 결정을 잘못하면 빌드는 통과해도 다른 GPU에서 런타임에 실패하거나, 매 실행마다 JIT로 다시 컴파일하느라 첫 호출이 느려진다.

2.12.1 이 장에서 배우는 것 / 선수 지식

배우는 것:

  • nvcc가 .cu를 호스트/디바이스 분리 컴파일하는 흐름.
  • PTX(가상 ISA)와 SASS(실제 머신코드) 의 차이, PTX가 JIT로 SASS가 되는 시점.
  • 가상 아키텍처(compute_XX)와 실제 아키텍처(sm_XX), -arch·-code·-gencode로 무엇이 바이너리에 들어가는지.
  • fatbin 다중 임베딩, compute capability 전방·후방 호환, PTX JIT·캐시.
  • cuobjdump·nvdisasm으로 바이너리 안 PTX·SASS 보기.

선수 지식: 0.3(드라이버·툴킷 설치), 2.1(첫 커널·빌드 명령). compute capability 숫자(예: 8.0, 9.0)가 GPU 세대를 식별함을 가정. 구체 버전·번호는 시점 의존(S-5).

2.12.2 왜 두 단계인가 — 가상 ISA와 실제 머신코드

GPU의 실제 명령어 집합(ISA)은 세대마다 바뀐다(Volta·Ampere·Hopper·Blackwell이 서로 다른 머신코드). 컴파일러가 곧장 특정 세대 머신코드만 생성하면 새 GPU마다 모든 바이너리를 다시 빌드해야 한다. 엔비디아는 2단계 컴파일 모델로 푼다.

nvcc 프런트엔드 cicc 등 ptxas 빌드 시점 AOT 또는 드라이버 JIT 런타임 소스 (.cu 디바이스 코드) PTX (Parallel Thread eXecution) 가상 ISA·텍스트, 세대 독립, 전방 호환 '이식 매체' SASS (Streaming ASSembler) 실제 머신코드·바이너리, 특정 sm_XX 고정, GPU가 직접 실행
  • PTX(가상 ISA): 텍스트 중간표현. 특정 칩이 아니라 가상 아키텍처(compute_XX)를 타깃. LLVM IR·JVM 바이트코드와 비슷한 위치 — "아직 기계어가 아닌, 이식 가능한 중간". GPU가 직접 실행 못 함.
  • SASS(실제 머신코드): 특정 실제 아키텍처(sm_XX)의 실제 명령어. GPU가 직접 디코드·실행. 세대 간 비호환. 어셈블러 ptxas가 PTX → SASS 변환.

핵심: PTX는 이식성, SASS는 실행. 빌드 시점에 SASS를 미리 만들어 넣으면(AOT) 그 세대에서 즉시 실행, PTX만 넣으면 처음 실행하는 GPU에서 드라이버가 JIT로 SASS를 만든다(이식성↑, 첫 실행 지연).

2.12.3 호스트/디바이스 분리 컴파일

.cu에는 CPU 호스트 코드(main·cudaMalloc)와 GPU 디바이스 코드(__global__·__device__)가 섞여 있다. nvcc는 이를 갈라 각자 다른 컴파일러로 보낸다.

nvcc 드라이버 조율자 fatbin 임베딩 hello.cu 호스트 경로: 호스트 C++ → 시스템 컴파일러 (gcc/clang/ MSVC) → 호스트 오브젝트 디바이스 경로: __global__/__ device__ → cicc → PTX → ptxas → SASS → fatbin 최종 실행파일 (CPU 기계어 + GPU 디바이스 코드)
  • 디바이스 코드는 cicc(NVVM/LLVM 기반)가 PTX로, ptxas가 SASS로 내리고 fatbinary가 하나의 fatbin에 묶는다.
  • 호스트 코드는 nvcc가 시스템 호스트 컴파일러에 넘긴다. nvcc는 C++ 컴파일러가 아니라 조율자 — 호스트 컴파일러 버전·표준 라이브러리 호환이 빌드 실패의 흔한 원인(-ccbin으로 지정).
  • fatbin은 호스트 오브젝트의 특수 섹션에 임베딩된다. 최종 실행파일 하나에 CPU 기계어 + GPU 디바이스 코드가 함께.

중간 산출물: --keep(.ptx·.cubin·.fatbin 남김), --verbose(하위 도구 호출 출력). 단계 끊기: -ptx·-cubin·-fatbin.

2.12.4 가상 vs 실제 아키텍처 — -arch·-code

이 장의 핵심이자 혼동이 잦은 지점. nvcc에는 두 종류의 아키텍처 이름이 있다.

종류접두사예의미결과
가상compute_compute_80PTX가 타깃하는 기능 집합PTX
실제sm_sm_80실제 GPU 세대SASS
  • -arch=<가상>: PTX를 어떤 가상 아키텍처로(= 어떤 기능까지). PTX 생성 입력.
  • -code=<실제 또는 가상>: ptxas가 어떤 실제 아키텍처용 SASS를, 그리고 PTX 자체를 바이너리에 넣을지.
PTX 생성 ptxas ptxas SASS 안 만듦 -arch=compute_80 PTX(compute_80) -code=sm_80 SASS(sm_80) → fatbin (그 GPU 즉시 실행) -code=sm_90 SASS(sm_90) → fatbin -code=compute_80 PTX 자체를 fatbin에 (JIT용 씨앗)

규칙: -code에 sm_XX면 그 세대 SASS AOT(JIT 없이 즉시 실행). compute_XX면 PTX가 fatbin에(미래 GPU JIT 씨앗, 전방 호환). -arch의 가상 번호는 -code의 실제 번호보다 높을 수 없다.

축약: -arch=sm_80 하나 = -arch=compute_80 -code=sm_80(그 세대 SASS만, PTX 없음). -arch=compute_80 하나 = PTX만. 즉 -arch=sm_XX는 "그 칩 전용 바이너리", -arch=compute_XX는 "PTX만, 어디서든 JIT".

2.12.5 fatbin과 -gencode — 한 바이너리에 여러 아키텍처

실무 바이너리는 여러 세대 GPU에서 모두 실행되어야 한다. fatbin은 여러 아키텍처 코드(SASS 여럿 + PTX)를 한 묶음에 담고, 런타임에 현재 GPU에 맞는 코드를 고른다. -gencode는 -arch·-code 짝을 여러 번 지정해 누적한다.

# 권장 패턴: 지원할 각 세대 SASS + 미래용 PTX 하나
#   (sm_80=Ampere SASS, sm_90=Hopper SASS, 마지막 compute_90=미래 GPU용 PTX)
nvcc kernel.cu -o app \
  -gencode arch=compute_80,code=sm_80 \
  -gencode arch=compute_90,code=sm_90 \
  -gencode arch=compute_90,code=compute_90 # 최상위 PTX → 미래 GPU JIT 씨앗

(번호는 시점 의존, 확인 필요. 패턴 — 여러 SASS + 마지막 PTX 하나 — 이 요점.)

아니오 예 예 아니오 런타임: 현재 GPU capability = X.Y fatbin에 sm_XY SASS가 있는가? 그 SASS 로드·실행 (가장 빠름, JIT 없음) fatbin에 X.Y 이하 PTX가 있는가? 드라이버 JIT로 sm_XY SASS 생성 (첫 실행 지연) 실행 실패 (no kernel image is available)

-arch=native는 빌드 머신 GPU capability를 자동 감지해 그 세대만 빌드 — 로컬엔 편하나 배포엔 위험(빌드/배포 머신 GPU가 다르면 실패). 대규모 배포는 지원 세대를 명시적으로 -gencode로 나열. 트레이드오프: 세대를 많이 넣을수록 fatbin·빌드 시간↑.

2.12.6 compute capability 호환 — 전방·후방

후방 호환 (SASS) — 같은 Major 안 전방 호환 (미래 GPU) — PTX JIT 경유 같은 Major 8.x 더 높은 minor 드라이버 JIT Major가 바뀌면 (8.x→ 9.x) SASS 못 씀 sm_80 SASS sm_86 등에서 실행 compute_80 PTX 미래 아키텍처(sm_ 90, sm_100…)에서 SASS 생성 PTX를 넣어둔 덕분 · 첫 실행에 JIT 비용
  • SASS 후방 호환은 같은 Major 안에서만 일반적으로 보장. Major가 바뀌면 미리 만든 SASS는 못 쓴다.
  • 전방 호환(미래 GPU)은 오직 PTX를 통해서만. fatbin에 PTX 씨앗을 넣어두면 출시 당시 없던 새 GPU에서도 JIT로 실행된다 — 2.12.5 패턴이 마지막에 최상위 PTX 하나를 두는 이유.
  • 단 JIT는 드라이버가 그 PTX 버전을 이해할 만큼 새로워야 한다. 너무 오래된 드라이버는 새 툴킷 PTX를 JIT 못 함 → 드라이버 업데이트.

2.12.7 PTX JIT와 JIT 캐시

PTX만 있고 맞는 SASS가 없으면 드라이버가 런타임에 PTX→SASS를 컴파일(Just-In-Time)한다. 그 PTX를 처음 쓰는 시점에 일어나 콜드스타트 지연(첫 커널 런치·모듈 로드가 느림)으로 나타난다. 매 실행마다 재컴파일은 낭비이므로 드라이버는 JIT 결과를 디스크에 캐시(compute cache)한다 — 같은 PTX·드라이버·GPU면 다음부터 캐시된 SASS 재사용.

첫 실행 다음 무효화 드라이버 JIT 느림 캐시 적중 PTX SASS 캐시 저장 PTX 저장된 SASS 로드 (빠름) 드라이버 버전· GPU·PTX 변경 다시 JIT

환경 변수(이름·동작 확인 필요): CUDA_CACHE_DISABLE=1(캐시 끄기, 진단용), CUDA_CACHE_PATH(위치), CUDA_FORCE_PTX_JIT=1(SASS가 있어도 JIT 강제, 전방 호환 테스트).

함정: 컨테이너·CI에서 캐시 디렉터리가 매번 새로 생기면 모든 실행이 콜드스타트("JIT 폭탄"). 배포에서 JIT 지연이 문제면 해당 GPU용 SASS를 -gencode로 미리 넣어 JIT 경로를 피하는 것이 근본 해법.

2.12.8 워크스루 — 여러 -gencode로 빌드하고 들여다보기

nvcc --list-gpu-arch     # 이 툴킷이 아는 가상 아키텍처(compute_XX) 목록
nvcc --list-gpu-code     # 실제 아키텍처(sm_XX) 목록
# 여러 세대 SASS + PTX (번호는 본인 환경에 맞게)
nvcc vadd.cu -o vadd \
  -gencode arch=compute_80,code=sm_80 \
  -gencode arch=compute_90,code=sm_90 \
  -gencode arch=compute_90,code=compute_90
# fatbin 안 구성 보기
cuobjdump vadd          # 어떤 sm_XX SASS, 어떤 PTX가 있는지
cuobjdump -ptx vadd     # 임베딩된 PTX 텍스트
cuobjdump -sass vadd    # 임베딩된 SASS 디스어셈블
# PTX/cubin 따로 만들어 보기
nvcc -ptx vadd.cu -arch=compute_80 -o vadd.ptx
nvcc -cubin vadd.cu -arch=sm_80 -o vadd.cubin
nvdisasm -c vadd.cubin            # SASS + 기본 블록

관찰(개략): cuobjdump vadd에 sm_80·sm_90 ELF + compute_90 PTX 항목. -ptx는 ld.global.f32·add.f32 같은 PTX(가상 무한 레지스터), -sass는 실제 명령(LDG·FADD·STG 류, 세대별 mnemonic·한정된 실제 레지스터). cuobjdump는 fatbin/실행파일을 까서 무엇이 들었나를, nvdisasm은 단일 cubin SASS가 어떻게 생겼나를 본다.

2.12.9 중간 점검 (체크포인트)

  • [ ] -arch=sm_80과 -arch=compute_80이 fatbin에 각각 무엇을 넣는지(SASS만 vs PTX만) 안다.
  • [ ] -gencode 권장 패턴이 왜 마지막에 PTX 하나를 두는지(전방 호환) 설명한다.
  • [ ] no kernel image is available 오류가 fatbin에 맞는 SASS도 JIT 가능한 PTX도 없다는 뜻임을 안다.
  • [ ] cuobjdump -ptx·-sass·nvdisasm을 구분해 쓴다.

증상: 첫 실행만 느리고 이후 빠름 → PTX JIT 콜드스타트 + 캐시(그 GPU용 SASS를 미리 넣었나 확인). 컨테이너에서 매 실행 느림 → JIT 캐시 디렉터리가 휘발성(CUDA_CACHE_PATH를 영속 볼륨, 또는 SASS AOT). no kernel image → -gencode 누락. PTX JIT compilation failed → 드라이버가 너무 오래됨(업데이트).

2.12.10 흔한 함정

  • 타깃 아키텍처 누락 → 런타임 실패. 배포 GPU 세대를 -gencode에 안 넣으면 맞는 SASS가 없다. PTX 씨앗도 없으면 no kernel image로 즉사, PTX가 있으면 매번 JIT. 항상 지원 세대 SASS + 최상위 PTX 하나.
  • -arch=native를 배포 빌드에. 빌드 머신 GPU만 타깃되어 다른 배포 머신에서 실패. 릴리스엔 명시적 -gencode.
  • 가상/실제 혼동. -code=compute_80(PTX만)을 SASS인 줄 알고 쓰면 어디서든 첫 실행 JIT. sm_80(SASS) ≠ compute_80(PTX).
  • 드라이버-툴킷 버전 불일치. 새 툴킷 PTX를 오래된 드라이버가 JIT 못 함. PTX JIT 실패. 버전·정책 확인 필요.
  • 호스트 컴파일러 비호환. nvcc가 부르는 gcc/MSVC가 지원 범위 밖이면 빌드 실패. -ccbin으로 지정.
  • fatbin 비대. 모든 세대를 다 넣으면 실행파일이 커지고 빌드가 느려진다. 실제 지원 세대만.

2.12.11 연습문제

연습 1 (fatbin 내용). vadd.cu를 sm_80·sm_90 SASS + compute_90 PTX로 빌드한 뒤 cuobjdump로 항목을 나열하라. SASS 몇 개, PTX 몇 개?

정답(개략)

SASS ELF 두 개(sm_80, sm_90) + PTX 하나(compute_90). cuobjdump vadd의 fatbin 섹션에서 각 항목 arch를 확인. 출력 형식은 툴킷 버전 의존.

연습 2 (no kernel image 재현). 본인 GPU와 다른 Major 세대 SASS 하나만(PTX 없이) 빌드해 실행하면 무슨 오류가? 왜고, 어떻게 고치나?

정답

no kernel image is available for execution on the device. fatbin에 현재 GPU(다른 Major)용 SASS가 없고 JIT할 PTX 씨앗도 없기 때문. 고침: -gencode에 현재 세대 SASS 추가하거나 최소 compute_XX PTX를 넣어 JIT 경로를 연다. (같은 Major의 더 높은 minor였다면 SASS 후방 호환으로 실행됐을 수도.)

연습 3 (JIT 캐시 효과). PTX-only 바이너리를 (a) 캐시 켠 채 두 번, (b) CUDA_CACHE_DISABLE=1로 두 번 실행해 첫 실행 시간을 비교하라. 무엇이 차이를?

정답

(a)는 첫 실행만 JIT(느림), 둘째부터 캐시 적중(빠름). (b)는 두 번 다 JIT — 캐시를 안 쓰므로. 차이는 드라이버 compute cache의 SASS 재사용 여부. 시간은 환경 의존.

연습 4 (배포 빌드 설계, 견해). A100급·Hopper·미래 GPU를 모두 지원하는 배포 바이너리의 -gencode 옵션을 써라. 무엇을 넣고 왜인가?

정답(예시)

각 지원 세대 SASS(arch=compute_80,code=sm_80·arch=compute_90,code=sm_90)로 JIT 없이 즉시 실행되게 하고, 마지막에 최상위 PTX 하나(arch=compute_90,code=compute_90)로 미래 GPU 전방 호환(JIT) 확보. native·PTX-only는 배포 부적합. 번호는 시점 의존이니 실제 세대로 교체. 트레이드오프: 세대 많을수록 바이너리·빌드시간↑, JIT 지연↓.

2.12.12 요약 · 다음 장

  • nvcc는 .cu를 호스트(시스템 컴파일러)와 디바이스(cicc→PTX→ptxas→SASS)로 분리 컴파일하고 디바이스 코드를 fatbin으로 묶어 실행파일에 임베딩한다.
  • PTX는 가상 ISA(이식·전방 호환), SASS는 특정 sm_XX 실제 머신코드(실행). -arch=sm_XX는 SASS AOT, -arch=compute_XX는 PTX. -gencode로 여러 세대 SASS + 최상위 PTX 하나가 정석.
  • 런타임에 맞는 SASS가 있으면 즉시 실행, 없고 PTX가 있으면 JIT(첫 실행 지연·캐시로 재사용), 둘 다 없으면 no kernel image.
  • 함정의 뿌리는 타깃 아키텍처 누락과 드라이버-툴킷 버전 불일치. 이 SASS 세대 종속은 5.2(TensorRT 빌드)에서 엔진을 빌드한 GPU 세대에 고정하는 점과 같은 원리.

다음 장: 2.13 프로파일링 I — Nsight Systems. 지금까지 만든 커널의 실제 실행을 시스템 타임라인으로 측정 — CPU/GPU 동시성, 오버랩, 런치 버블을 눈으로 본다.