《GPU Mode》
L040
2024 · DEC
High priority
transcript · available
CUDA Docs for Humans
NVIDIA 의 공식 docs 는 “완전하지만 사람을 위해 쓰여 있지 않다”. 같은 사실이 5 군데에 있고, 검색이 어렵고, 학습 경로가 안 보인다. Charles Frye 가 짠 GPU Glossary 는 같은 자료를 사람이 학습하기 좋은 형태로 다시 짠 hypertextual reference. CUDA 의 stack — programming model · runtime · driver · PTX — 을 한 화면에서 navigate 하는 방법, ML 연구자가 어디서 막히는지의 패턴, LLM 을 docs 학습 도구로 어떻게 쓰는지까지.
CUDA stack
GPU Glossary
PTX
runtime API
driver API
SM
PMPP
LLM-aided learning
C
Speaker
Charles Frye
Modal · ML educator · GPU Glossary 저자
§ 01강의가 풀려는 문제· 왜 docs 가 어려운가
“CUDA 가 어려운 게 아니라, CUDA 의 docs 가 어렵다” 의 측정
Charles 의 출발점은 단순하다 — ML 연구자/엔지니어가 CUDA 를 배우려 할 때 가장 큰 마찰이 “어디서부터 어떻게 읽어야 하는지 모른다”는 것. NVIDIA 의 docs 는 자료가 모자라서가 아니라 너무 많고 분산되어 있어서 학습 경로가 안 잡힌다.
강의의 본질은 “같은 사실의 다른 layout”. NVIDIA 의 공식 자료는 reference 로 매우 좋다 — 어떤 함수의 정확한 의미를 찾을 때. 그러나 학습 시퀀스로는 안 좋다 — 무엇이 중요하고 무엇이 가지치기인지 표시가 없다.
강의의 인지적 frame
Charles 의 입장 — “docs 가 학습용이려면 (1) hypertextual (한 단어 클릭 → 정의), (2) task-oriented (그 일을 하려면 무엇을 알아야 하는가), (3) navigable (5 분 만에 어디든 갈 수 있어야)” 의 세 조건을 만족해야 한다. NVIDIA 의 docs 는 (1) 만 부분적, (2) (3) 은 거의 없음. 그 빈 자리를 사람이 메우는 작업이 GPU Glossary.
“좋은 학습 자료는 사실의 정확함이 아니라 사실들 사이의 관계를 보여주는 자료다.”Charles Frye · 강의 도입부
그래서 강의의 자산은 — (1) CUDA stack 의 4 layer (programming model · runtime · driver · PTX) 의 위치 그림, (2) ML 연구자가 가장 자주 빠지는 함정의 list, (3) docs 를 navigate 하는 트릭, (4) LLM 을 학습 도구로 쓰는 패턴.
§ 02official docs 의 한계· 완전 ≠ 학습 가능
“같은 사실이 5 군데 있고, 어디가 정본인지 모른다”
NVIDIA docs 의 가장 큰 문제는 분산성. 같은 개념이 — Programming Guide, Best Practices Guide, PTX ISA, Toolkit Reference, 각 라이브러리(cuBLAS / cuDNN / NCCL) 의 자체 docs 에 — 약간씩 다른 표현으로 나타난다. 어디를 정본으로 봐야 하는지의 휴리스틱이 사람마다 다르다.
official docs 가 잘 못하는 것
분산 · 중복 · 중복
같은 개념(예: occupancy) 이 Programming Guide 에 한 정의, Best Practices Guide 에 다른 강조, NCU 의 metric description 에 또 다른 정의. 어떤 사용자에게 어떤 정의가 맞는지 표시 없음.
새 사용자는 “어디서부터 읽어야 한다” 의 답을 찾기 위해 Reddit / Stack Overflow / blog 를 떠돈다. docs 자체가 그 navigation 을 해주지 않는다.
official docs 가 잘 하는 것
완전성 · 정확성
한 함수의 정확한 signature, 한 instruction 의 정확한 의미. 찾을 자리를 알면 정확한 답이 있다.
그래서 GPU Glossary 의 입장 — “공식 docs 를 대체하지 않는다. 들어가는 입구를 만들어 준다.” Glossary 의 거의 모든 항목이 NVIDIA 의 official 자료로 cross-link.
하나의 docs 가 모든 사용자를 다룰 수 없다
강의에서 Charles 가 짚은 — “NVIDIA 입장에서 docs 는 hardware engineer 와 ML 연구자와 game developer 를 모두 만족시켜야 한다. 그래서 일반화될 수밖에 없다.” 특정 청중을 위한 docs 는 community 가 만든다 — Glossary 가 ML 청중을 위한 한 시도.
§ 03task-oriented 재구성· topic → task
“이 일을 하려면 무엇을 알아야 하는가” 의 형태로 자료를 재정렬
topic-oriented (이 개념의 정의) 와 task-oriented (이 일을 하려면) 가 다른 자료. 학습 자료는 보통 task-oriented 가 더 효율적. Glossary 는 둘 다 하는 형태.
L0 · 진입
"PyTorch 가 GPU 위에서 무엇을 하는가" 같은 큰 질문ML 연구자의 자연스러운 출발점. 이 질문이 stack 의 어디로 매핑되는지 표시.
curriculum entry
L1 · 큰 그림
CUDA stack 의 4 layerprogramming model · runtime · driver · PTX. 같은 단어 “CUDA” 가 4 개 다른 것을 가리킨다.
disambig
L2 · 개념
SM · warp · thread block · shared memory…각 개념의 한 문단 정의 + 다른 개념과의 관계 + official docs 링크.
hypertextual
L3 · 구체
함수 / instruction 단위NVIDIA 의 official reference 로 link out. Glossary 자체는 한 줄 요약만.
link out
L4 · 다음 단계
PMPP · GPU Mode lectures · LLM다음 학습 자료로의 명시적 링크. “이 다음에 무엇을 보면 좋은가” 가 자료의 끝마다 박혀 있다.
curriculum exit
강의에서 Charles 가 보여준 한 화면 — Glossary 의 한 항목 (예: compute capability) 가 다른 항목들과 거미줄처럼 연결. 한 단어 클릭 → 정의 → 관련 개념 → 공식 docs 의 시퀀스가 5 초 안에 완료. 이게 Glossary 의 의도된 사용 방식.
“docs 의 가치는 한 항목의 길이가 아니라 항목들 사이의 link 밀도에 있다.”Charles Frye
§ 04PMPP 와의 관계· 교과서 + reference
책 한 권 + reference 한 본 — 학습의 두 축
PMPP (책)
선형적 학습 경로
Chapter 1 → 2 → 3 의 sequential. 한 번 읽고 나면 mental model 이 생긴다. 그러나 한 항목 빠르게 lookup 하기는 어려움.
GPU Glossary (Charles)
hypertextual reference
선형 X. 어디든 click 으로 진입. 이미 mental model 이 있는 사람이 빠르게 lookup 하거나 학습할 때 best.
official docs (NVIDIA)
완전한 reference
정확한 signature, 정확한 의미. 그러나 학습 entry point 가 없음. ”알아야 할 게 정해진” 사용자를 위한 자료.
Charles 의 추천 시퀀스 — (1) PMPP 1~5 장으로 mental model. (2) Glossary 로 navigate 하면서 PMPP 에서 헷갈렸던 것들을 cross-check. (3) 실제 작업에서 official docs 로 정확한 signature lookup.
strict vs casual 관계
Glossary 는 PMPP 의 경쟁자가 아니라 보완재. 책이 mental model 을 만들고, Glossary 가 그 mental model 위에서 빠른 lookup 을 가능하게 한다. 같은 문제를 다른 형태로 해결.
§ 05자주 빠지는 함정· async · 메모리 모델
ML 연구자가 가장 자주 막히는 5 가지
Glossary 를 짜면서 Charles 가 관찰한 — ML 연구자가 GPU 를 배울 때 가장 자주 헷갈리는 자리들. 흥미롭게도 거의 모두가 “async / 동시성” 의 변종.
async kernel launchCPU 가 launch 한 시점에 GPU 가 끝나지 않는다. time.time() 으로 측정 안 됨 (L001 의 출발점).
stream 의 위치한 GPU 안에서 여러 stream 이 병렬. PyTorch 가 보통 default stream 만 쓰지만, 분산이나 graph capture 에서는 명시적.
unified memory vs explicit copycudaMallocManaged 의 “HBM 처럼 보이는 memory” 가 사실은 page fault 로 옮겨짐. 빠른 코드에선 explicit copy.
register vs shared vs HBM한 thread 가 가지는 자리 (register), 한 block 이 가지는 자리 (shared), 모두가 보는 자리 (HBM). 크기와 latency 가 100~1000× 차이.
PTX vs CUDA C++PTX 는 virtual ISA. CUDA C++ 가 PTX 를 거쳐 SASS 로 컴파일. PTX 한 줄 ≠ SASS 한 줄 (L037).
runtime API vs driver APIruntime 은 cudaXxx (간단), driver 는 cuXxx (상세). 대부분의 ML 사용자는 runtime 만. driver 는 dynamic load 같은 lower-level 시.
CC (compute capability)SM 의 “세대” + “version”. 7.0 = Volta, 8.0 = Ampere, 9.0 = Hopper. 같은 코드가 CC 에 따라 다르게 컴파일.
함정의 공통점
위 7 개의 공통점 — “같은 단어가 여러 의미로 쓰이거나, 보이는 것보다 더 많은 일이 뒤에서 일어난다”. ML 연구자는 추상에 익숙한 청중이라, 추상 아래의 동시성/계층 구조에 약하다. Glossary 의 가장 중요한 자리는 이 함정 7 개를 명시적으로 표시하는 것.
§ 06자료 navigation 트릭· search · cross-link
모르는 단어를 만났을 때 5 분 안에 답에 도달하는 워크플로
- (1) Glossary 에서 단어 검색 — Glossary 가 가장 짧고 hypertextual. 한 단어가 여러 정의를 가질 수 있다는 사실을 첫 화면에서 disambig.
- (2) PMPP 의 관련 chapter — Glossary 가 정의를 주면, PMPP 에서 그 정의의 mental model 이 어떻게 구성되는지.
- (3) NVIDIA docs 의 정확한 signature — 함수/instruction 차원의 정확한 의미. Glossary 의 each entry 에 link out.
- (4) GPU Mode lecture 의 deep dive — 같은 개념의 실제 사용 사례와 함정. 이 archive 의 각 페이지가 그 자리.
- (5) NVIDIA blog / Colfax / SemiAnalysis — 최신 hardware 의 새 기능에 대한 reverse-engineering 류 자료. Hopper, Blackwell 같은 새 architecture 는 docs 가 부족.
Charles 의 specific tip — “같은 단어 (예: thread)” 가 CPU/GPU 에서 다른 의미 임을 항상 의심. CUDA thread 는 CPU thread 와 매우 다른 단위 (32 thread = warp 가 SIMD 처럼 lockstep). 단어를 만나면 “이게 어느 layer 의 thread 인가?” 를 묻는 습관.
“좋은 학습은 정의에 도달하는 것이 아니라 정의들 사이의 차이를 인식하는 것.”학습 노트
§ 07LLM 활용· docs 학습의 새 도구
docs 가 다 들어 있는 LLM 에게 학습 partner 의 자리를 준다
강의의 흥미로운 부분. Charles 가 명시한 한 줄 — “NVIDIA docs 는 거의 다 LLM 의 학습 데이터에 들어 있다. 그래서 LLM 에게 ‘설명해 달라’ 는 게 docs 를 직접 읽는 것보다 빠른 경우가 많다.”
LLM 활용 패턴 4 가지
(a) 한 단어 disambig — “CUDA thread 와 CPU thread 의 차이를 한 문단으로.” (b) 시퀀스 추천 — “Volta 의 tensor core 를 배우려면 어떤 PMPP chapter 를?” (c) 코드 설명 — Triton/CUDA snippet 한 개를 LLM 에게 “여기서 일어나는 일은?” (d) 함정 검증 — “이 코드가 async 적으로 잘못된 부분이 있는가?”
Charles 가 강조한 단점 — LLM 은 최신 hardware 의 새 기능 에 약하다. cutoff 이전의 자료만 보장. Hopper 의 TMA 같은 건 LLM 의 답이 부정확하거나 outdated 일 수 있어서, 직접 official docs 또는 Colfax / Tri Dao 의 자료를 확인.
“LLM 은 reference 를 외우고 있다. 사람이 잘 못하는 일을 잘 한다. 그 자리에 학습 partner 로 둔다.”Charles Frye
§ 08커뮤니티 ratchet· PR 가능한 형태
Glossary 가 GitHub repo 인 이유 — 한 사람이 다 못 채운다
Glossary 의 design 결정 중 하나 — open source. GitHub 에 markdown 으로 적혀 있고, PR 가능. 이유는 단순하다 — 한 사람이 GPU 의 모든 면을 cover 할 수 없다.
Charles 는 ML 연구자 / educator 의 시각으로 시작. SASS / uarch 의 깊은 부분은 비교적 얇다. 그 자리에 Arun (L037) 같은 expert 가 PR 로 채울 수 있는 형태. 커뮤니티가 ratchet 으로 채우는 docs 가 한 사람이 fully 짜는 docs 보다 결국 좋아진다는 가설.
PR 의 좋은 패턴
(a) 새 entry 추가 — 빠진 개념. (b) 기존 entry 의 link 추가 — 다른 entry 와의 관계. (c) 정의의 정정 — 잘못된 표현. (d) 새 hardware 갱신 — Blackwell 같은 세대가 release 되면. 각각이 작은 PR 로 가능하다는 점이 유지보수의 핵심.
강의에서 Charles 가 짧게 — “Modal 이 hosting 하지만 NVIDIA 의 자료가 아니고, 누구나 PR 할 수 있다. 잘못된 부분이 있으면 issue 또는 PR 환영.” 라는 입장.
§ 09다음 자료· 시간 차원의 학습 곡선
Glossary 다음에 무엇을 — Charles 의 추천 경로
- PMPP 6~10 장 — convolution / sparse / scan / reductions. mental model 의 두 번째 layer.
- GPU Mode lectures L001 ~ L010 — profiling, performance checklist, fusion 같은 실용 시퀀스. Glossary 의 개념을 실제 코드 안에서 본다.
- flash attention paper 시리즈 — FA1 / FA2 / FA3. 같은 알고리즘이 hardware 와 같이 진화한 사례. attention 이 GEMM 만큼 효율적이 된 자리 (L036).
- CUTLASS / CuTe docs — kernel 직접 짜는 단계. Hopper 의 새 자원을 어떻게 expose 하는지의 표현. 학습 곡선이 가파름.
- Colfax research blog — Hopper 시대 이후의 새 패턴들 (warp specialization, ping-pong) 의 reverse-engineering 자료.
- cuobjdump + godbolt — 자기가 짠 커널의 SASS 직접 읽기 (L037). docs 의 마지막 layer.
강의에서 Charles 가 직접 — “이 archive (GPU Mode lectures) 도 Glossary 와 짝을 이루는 자료. 같은 개념을 강의 형태로, deeper 로 다룬다.” 자연스러운 다리.
§ 10기억할 메모와 코드· glossary 포인터
다시 열었을 때 5분 안에 손에 잡혀야 할 것
official docs 의 한계
완전하지만 학습 가능하지 않다. 한 사실이 여러 자리. 정본 결정 어려움.
CUDA stack 4 layer
programming model · runtime API · driver API · PTX. 같은 “CUDA” 단어가 4 개 의미.
task-oriented 재구성
topic-oriented (정의) 와 task-oriented (이 일을 하려면) 의 차이. Glossary 는 둘 다.
PMPP + Glossary + lectures
책 (mental model) + reference (lookup) + 강의 (deep dive). 셋이 짝지어 학습.
자주 빠지는 7 함정
async, stream, unified memory, register/shared/HBM, PTX vs C++, runtime vs driver, CC.
함정의 공통점
같은 단어 다른 의미 + 보이는 것보다 많은 일이 뒤에. ML 연구자가 약한 자리.
LLM 학습 partner
docs 가 LLM 학습 데이터에 들어 있다. disambig / sequence / 코드 설명에 활용. 최신 hardware 는 official 확인.
PR 가능한 docs
Glossary 가 GitHub. 누구나 PR. 커뮤니티 ratchet 의 가설.
손에 새기기 — 실습 시퀀스
- Glossary 30 분 둘러보기 — 메인 페이지에서 무작위 entry 5 개. 각각의 link out 따라가며 “5 클릭 안에 어디든 도달” 의 패턴 손에.
- 한 함정 직접 검증 — 위 7 함정 중 자기가 약한 것 하나 골라, Glossary + PMPP + lecture 의 세 자료로 cross-check. mental model 가 맞는지.
- PR issue 만들기 — Glossary 에서 빠진 또는 헷갈리는 entry 발견하면 GitHub issue. 작은 contribution 으로 시작.
- LLM partner 실험 — 모르는 CUDA 개념 한 개 골라 Claude / GPT 에게 “PMPP chapter X 를 한 문단으로” 요청. 그 답을 PMPP 에서 검증.
- 자기 학습 경로 표시 — Glossary 의 어느 entry 가 자기에게 새로웠는지 markdown 한 페이지에. 6 개월 뒤 다시 볼 때의 baseline.
§ 11다른 강의로 이어지는 길· connections
Glossary 의 entry 들이 시리즈 안 어디에 다시 등장하는가
§ 12열린 질문· open questions
다음에 다시 들었을 때 직접 검증해야 할 것들
- Glossary 의 새 hardware 갱신 속도 — Hopper / Blackwell 의 새 기능이 언제 들어오는가. 커뮤니티 PR 가설이 실제로 작동하는지 시간이 지나며 확인.
- LLM 의 GPU 지식 cutoff — 모델별로 PTX/SASS 의 깊이 차이. 어떤 모델이 “Hopper 의 wgmma 를 정확히 이해하는가” 를 prompt 로 검증.
- “PR 가능한 docs” 의 한계 — open source 가설이 실제로 다른 expert 들의 contribution 을 끌어들이는가. 강의 시점 기준 contribution 패턴이 어디로 가는지.
- NVIDIA 의 docs 정책 변화 — Hopper / Blackwell 시대에 NVIDIA 가 docs 를 더 학습 친화적으로 만드는 방향이 보이는가. 그러면 Glossary 의 자리가 어떻게 변하는가.
- 다른 vendor (AMD ROCm, Apple Metal) — 같은 형태의 “Glossary” 가 있는가. CUDA 만 이 자리가 채워져 있고 다른 vendor 는 docs 부족이 더 심한 듯.
검증 메모
Glossary 는 진행형 자료. 이 노트의 “7 함정” list 는 강의 시점의 강조점이고, Glossary 의 실제 entry 와 일대일은 아닐 수 있다. 자기 학습 시점의 Glossary entry list 를 기준으로 다시 확인.