1. 왜 첫 강의가 profiler인가

Mark가 잡은 독자는 "CUDA tutorial hell"에 지친 PyTorch 사용자다. CUDA 전문가가 어딘가 어두운 방에서 마법처럼 kernel을 짠다는 식의 이야기를 걷어내고, 필요한 순간에 필요한 만큼만 내려가는 흐름을 보여준다.

그래서 출발점은 torch.square 같은 작은 연산이다. 이미 PyTorch가 잘하는 일을 일부러 다시 구현해 본다. 목적은 PyTorch를 이기는 것이 아니라, custom kernel이 Python 프로그램 안으로 들어오고 profiler에서 어떻게 보이는지 손으로 확인하는 것이다.

흥미로운 점은 첫 예제가 빠르지 않아도 된다는 데 있다. 실제로 custom Triton 코드가 PyTorch보다 느릴 수 있다. 이 실패가 오히려 좋은 출발점이다. 느리다는 사실을 알았으면, 이제 "왜 느린가"를 물을 수 있다.

2. PyTorch 안으로 CUDA를 들여오기

강의의 실전 축은 PyTorch extension이다. CUDA는 보통 C/C++ 쪽에서 작성되지만, PyTorch는 load_inline 같은 경로를 통해 Python 코드에서 C++/CUDA 코드를 컴파일하고 불러올 수 있게 해 준다.

여기서 중요한 감각은 host와 device의 경계다. Python은 kernel을 호출하고 tensor pointer를 넘긴다. 실제 병렬 계산은 GPU 위에서 실행되는 __global__ 함수가 한다. CUDA 코드 자체는 낯설어 보여도, 입출력 tensor를 받고 index를 계산해 한 element씩 처리한다는 구조는 단순하다.

load_inline은 학습용으로 특히 좋다. 컴파일 캐시나 배포 문제까지 한 번에 해결하려 들지 않고, "내 kernel이 PyTorch 프로그램에서 호출된다"는 첫 연결을 빨리 만든다.

3. profiler는 한 가지가 아니다

강의는 profiler를 계단처럼 다룬다. 먼저 autograd profiler로 큰 덩어리의 시간을 본다. 그 다음 PyTorch profiler로 Chrome trace를 열어 CPU activity와 CUDA activity가 어떻게 이어지는지 본다. trace의 화살표, 즉 flow event는 Python 쪽 호출이 어떤 CUDA launch로 연결되는지 보여준다.

이 단계에서 이미 많은 것을 알 수 있다. 어떤 kernel이 자주 호출되는지, CPU가 GPU에게 일을 던지는 시간이 보이는지, 실제 GPU kernel 시간이 지배적인지 확인할 수 있다. 성능 문제를 "GPU가 느림"이라고 뭉뚱그리지 않는 연습이다.

마지막에는 Nsight Compute, 즉 ncu로 내려간다. 여기서는 kernel grid, block size, achieved occupancy, memory coalescing, shared memory 관련 힌트가 나온다. 다만 ncu는 환경을 탄다. 클라우드 드라이버나 권한 문제로 profiling이 안 될 수 있으니, 도구가 실패하는 경우도 개발 경험의 일부로 받아들여야 한다.

4. Triton은 쉬운 길이지만 자동 승리는 아니다

Triton은 Python 안에서 GPU kernel을 쓰게 해 주는 DSL이다. CUDA C++보다 훨씬 짧게 쓸 수 있고, block 단위 프로그래밍 모델이 PyTorch 사용자에게 친절하게 다가온다.

하지만 Triton을 쓴다고 저절로 빨라지지는 않는다. block size, num_warps, memory access pattern이 맞지 않으면 PyTorch의 기존 kernel보다 느릴 수 있다. 강의에서 square 예제는 바로 이 점을 보여준다. 성능은 언어의 이름이 아니라 실행된 kernel의 모양에서 나온다.

Triton이 생성한 PTX를 들여다볼 수 있다는 점도 좋다. 고수의 영역처럼 보이는 compiler 산출물을 작게 열어 보면, global memory load/store나 register 사용 같은 단서가 보인다. 처음에는 다 읽지 못해도 괜찮다. "볼 수 있다"는 것 자체가 큰 진전이다.

5. 성능 힌트는 질문으로 읽자

profiler가 "occupancy가 낮다", "memory coalescing을 보라", "block size가 이상하다" 같은 신호를 줄 때, 그것을 바로 처방으로 받아들이면 위험하다. occupancy를 올렸다고 항상 빨라지는 것은 아니고, shared memory를 쓴다고 늘 이기는 것도 아니다.

그래도 좋은 질문 목록은 생긴다. launch overhead가 큰가? 같은 일을 너무 많은 kernel로 쪼갰는가? memory access가 연속적인가? PyTorch baseline은 어떤 kernel을 호출하는가? Triton이 알아서 해 주는 부분과 CUDA로 직접 관리해야 하는 부분은 어디인가?

6. 이 강의에서 가져갈 태도

첫 강의의 결론은 "CUDA는 쉽다"가 아니다. 더 정확히는 "CUDA로 내려가는 첫 계단은 생각보다 낮다"에 가깝다. PyTorch에서 extension을 붙이고, Triton kernel을 작성하고, profiler로 보는 일은 충분히 손에 잡힌다.

반대로 최적화는 여전히 정직하다. 내 kernel이 왜 느린지 모르면 더 낮은 층으로 내려가야 한다. profiler, PTX, Nsight Compute는 겁주는 도구가 아니라 질문을 정교하게 만들어 주는 도구다.

Source Links

• YouTube
• Slides PDF
• Slides Google
• Code
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_001.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_001.md

1. PMPP식으로 CUDA를 보기

이 강의는 PMPP(Programming Massively Parallel Processors) 앞부분을 CUDA Mode식으로 복습한다. 병렬화는 말로는 쉽다. 하지만 실제로 높은 성능을 내려면 sequential loop를 그냥 여러 조각으로 나누는 것보다 더 많은 생각이 필요하다.

예를 들어 prefix sum처럼 처음 보기에는 순차적인 문제가 있다. 반대로 vector addition처럼 거의 embarrassingly parallel한 문제도 있다. GPU 공부의 초반에는 이 차이를 구분하는 눈이 중요하다. 모든 문제가 같은 모양으로 병렬화되지는 않는다.

Andreas는 memory bandwidth도 계속 강조한다. LLM inference처럼 token을 하나씩 만들 때 batch가 작으면 GPU 계산 유닛을 충분히 못 쓸 수 있다. 즉 병렬성이 있다고 자동으로 빠른 것은 아니다. 데이터 이동과 실행 모양이 맞아야 한다.

2. host, device, 그리고 귀찮지만 중요한 copy

CUDA의 기본 프로그램은 host CPU와 device GPU 사이를 오간다. GPU memory를 할당하고, CPU memory의 데이터를 GPU로 복사하고, kernel을 launch하고, 결과를 다시 가져온다. 처음에는 이 순서가 장황해 보이지만 성능을 이해하려면 꼭 지나가야 한다.

강의의 vector addition 예제는 이 흐름을 가장 단순한 형태로 보여준다. 한 thread가 한 output element를 맡는다. d_는 device pointer, h_는 host pointer처럼 이름을 나눠 두면 머릿속에서도 데이터가 어디 사는지 덜 헷갈린다.

PyTorch를 쓰면 이 중 많은 부분이 감춰진다. 하지만 감춰졌다는 말은 없어진다는 뜻이 아니다. tensor가 CUDA device에 있으면 PyTorch는 그에 맞는 CUDA kernel을 호출하거나 생성한다. 나중에 느린 코드를 만났을 때 이 경계를 기억해야 한다.

3. grid, block, thread

CUDA kernel은 많은 thread로 실행된다. 이 thread들은 grid, block, thread라는 계층으로 조직된다. 각 thread는 blockIdx, threadIdx, blockDim 같은 built-in 값을 이용해 자신이 맡을 global index를 계산한다.

처음 익힐 공식은 단순하다. 1차원에서는 대개 idx = blockIdx.x * blockDim.x + threadIdx.x 꼴이다. 전체 원소 수가 block size로 딱 나누어떨어지지 않으므로, kernel 안에는 if (idx < n) 같은 guard가 필요하다.

block 안의 thread들은 shared memory를 공유할 수 있다. block끼리는 독립적이고 실행 순서를 가정하면 안 된다. 이 제약은 귀찮은 규칙이 아니라 CUDA 알고리즘 설계의 뼈대다.

4. 세 예제의 공통 구조

vector addition은 한 thread가 a[i] + b[i]를 계산한다. RGB to grayscale도 비슷하다. 한 pixel을 맡은 thread가 RGB 값을 읽고 grayscale 값을 쓴다. data가 2D 이미지라면 index 계산이 조금 더 입체적일 뿐이다.

mean filter는 한 걸음 더 간다. 한 output pixel을 계산하려면 주변 pixel 여러 개를 읽어야 한다. 그래서 "한 thread가 한 output element"라는 전략은 그대로지만, memory access 패턴은 더 복잡해진다. 여기서 이후 강의의 shared memory와 tiling 이야기가 자연스럽게 기다리고 있다.

matrix multiplication도 언급된다. 단순한 전략은 한 output matrix element를 한 thread가 맡고, row와 column의 dot product를 계산하는 것이다. 아직 빠른 matmul은 아니지만, GPU kernel을 생각하는 기본 문법으로는 아주 좋은 출발점이다.

5. compilation도 실행 모델의 일부다

CUDA code는 보통 nvcc를 거쳐 PTX 같은 중간 표현으로 내려간다. 이때 __global__ 함수는 host에서 launch할 수 있는 kernel이라는 뜻을 갖는다. triple angle bracket 문법은 "이 함수를 GPU에서 grid/block 모양으로 실행하라"는 CUDA C++ 확장이다.

강의 자료에는 순수 CUDA C 예제와 PyTorch extension 예제가 함께 있다. 같은 kernel 아이디어가 standalone CUDA 프로그램으로도, Python에서 호출되는 extension으로도 나타난다. 둘을 나란히 보면 PyTorch가 편의를 주는 층과 CUDA가 실제 일을 하는 층이 분리되어 보인다.

6. 초반에 조심할 것

thread를 많이 launch하는 것을 CPU thread 생성 비용과 같은 식으로 상상하면 안 된다. GPU에서는 아주 많은 thread를 띄우는 것이 자연스럽다. 대신 block 크기, memory access, guard 조건, shared memory 사용 가능 범위를 신경 써야 한다.

이 강의는 아직 최적화 끝판왕을 만들지 않는다. 목표는 CUDA 프로그램의 해부학을 익히는 것이다. data가 어디 있고, 누가 계산하고, 어떻게 launch되는지 알면 이후 profiler와 memory hierarchy 이야기가 훨씬 선명해진다.

Source Links

• YouTube
• Slides
• Slides Google
• Code
• Supplemental code
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_002.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_002.md

1. 낯선 CUDA를 익숙한 Python으로 당기기

강의는 이미지 tensor에서 시작한다. channel, height, width 모양의 rank-3 tensor를 보고, PyTorch와 matplotlib가 축 순서를 다르게 기대한다는 아주 현실적인 이야기부터 한다. CUDA 입문이지만 첫 장면은 여전히 Python 데이터 다루기다.

이 선택이 좋다. GPU kernel은 결국 data layout 위에서 돈다. 이미지의 한 pixel을 grayscale로 바꾸는 코드를 Python으로 쓰고, 그 다음 같은 일을 CUDA가 하도록 옮기면 "병렬화"가 추상 단어가 아니라 index 계산 문제로 내려온다.

Jeremy의 핵심 습관은 먼저 Python으로 작은 kernel runner를 흉내 내는 것이다. 진짜 GPU thread는 아니지만, block과 thread를 loop로 돌며 각 thread가 어떤 일을 하는지 눈으로 볼 수 있다.

2. range가 아니라 blocks와 threads

CUDA는 Python의 range(n)처럼 하나씩 값을 넘겨주는 모델이 아니다. 전체 일을 block들로 나누고, 각 block 안에 thread들이 있다. thread는 blockIdx, threadIdx, blockDim으로 자기 위치를 계산한다.

처음에는 이 구조가 불필요하게 이상해 보인다. 하지만 이유가 있다. 같은 block 안의 thread들은 shared memory를 공유하고, 필요하면 block 단위로 동기화할 수 있다. 나중에 빠른 kernel을 만들 때 이 구조가 바로 성능의 지렛대가 된다.

강의에서는 256 threads per block 같은 관용적인 값을 놓고, 필요한 block 수를 ceiling division으로 계산한다. 마지막 block이 남는 원소보다 더 많은 thread를 가질 수 있으므로 guard 조건이 꼭 필요하다.

3. grayscale kernel의 교훈

grayscale 예제는 GPU 입문에 알맞다. 각 output pixel은 독립적이고, 입력의 RGB 세 값을 읽어 하나의 값을 쓴다. 한 thread가 한 output pixel을 맡으면 된다.

이때 중요한 것은 계산식보다 indexing이다. channel-first tensor를 어떻게 읽을지, row와 column을 어떻게 만들지, output 범위를 벗어난 thread는 어떻게 빠져나올지 정해야 한다. CUDA의 많은 초보 버그는 수학이 아니라 이 좌표 변환에서 나온다.

그리고 여기서 이미 debugging 전략이 나온다. 처음부터 CUDA에서만 작성하면 확인하기 어렵다. Python version을 먼저 만들고, 작은 입력에서 결과를 비교하면 훨씬 덜 헤맨다.

4. PyTorch extension으로 연결하기

CUDA kernel을 Python에서 쓰려면 binding이 필요하다. PyTorch는 C++/CUDA extension을 통해 이 경로를 제공한다. 강의에서는 Python에서 작성한 kernel 구조를 CUDA C로 옮기고, PyTorch가 compile/load한 함수를 호출한다.

CUDA의 triple angle bracket 문법은 처음 보면 이상하다. 하지만 읽는 법은 단순하다. 어떤 kernel을, 몇 개의 block과 몇 개의 thread로, 어떤 argument와 함께 실행할 것인지 적는다.

이 방식은 교육용 장난만은 아니다. Jeremy는 가능한 한 Python으로 먼저 생각하고 CUDA로 옮기는 습관을 실제 작업 방식으로도 선호한다고 말한다. 복잡한 kernel일수록 debug 가능한 작은 형태가 필요하기 때문이다.

5. matmul로 넘어가는 길

강의 후반은 matrix multiplication으로 향한다. 가장 단순한 matmul kernel은 output matrix의 한 cell을 한 CUDA thread가 맡고, dot product를 계산한다. 빠르지는 않지만 구조를 보기에는 좋다.

여기서 2D grid와 2D block이 등장한다. row와 column은 각각 block index와 thread index의 x/y 좌표에서 나온다. grayscale에서 배운 1D 감각을 2D로 확장하는 셈이다.

빠른 matmul은 다음 강의에서 shared memory와 tiling으로 이어진다. 지금은 "정확한 naive CUDA matmul"을 만드는 것만으로 충분하다. 정확함 없이 최적화로 달려가면, 나중에 profiler가 아니라 점괘를 읽게 된다.

6. 혼자 해볼 작은 과제

먼저 Python으로 grayscale 함수를 쓰고, 같은 입력에 대해 PyTorch 연산 결과와 비교해 보자. 그 다음 CUDA extension version을 만들고 결과가 byte 단위로 얼마나 가까운지 확인한다.

그 다음에는 일부러 block size를 바꿔 보자. 128, 256, 512 threads per block에서 launch 모양은 어떻게 달라지는가? 성능보다 먼저 "몇 개의 thread가 실제로 유효한 일을 했는가"를 계산해 보는 것이 좋다.

Source Links

• YouTube
• Code
• Colab
• Notebook
• Supplemental code
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_003.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_003.md

1. block은 SM 위에 올라간다

CUDA kernel을 launch할 때 우리는 block 안의 thread 수와 전체 block 수를 정한다. 실행 시 각 thread block은 어떤 Streaming Multiprocessor(SM)에 배정된다. 프로그래머가 "이 block은 저 SM으로"라고 직접 지정하지 않는다.

한 SM에는 여러 block이 동시에 올라갈 수 있다. 다만 register, shared memory, 최대 thread 수 같은 resource limit이 허락해야 한다. 그래서 block size는 성능에 영향을 준다. 너무 작으면 일을 충분히 못 채우고, 너무 크면 동시에 올라갈 block 수가 줄어들 수 있다.

강의에서 언급되는 RTX 3090 예처럼 한 SM이 수용할 수 있는 thread 수가 정해져 있을 때, 256이나 512 같은 block size가 깔끔하게 맞는 경우가 있다. 숫자를 외우기보다 "내 launch shape가 hardware limit과 어떻게 맞물리는가"를 보는 습관이 중요하다.

2. warp divergence는 조용히 시간을 먹는다

NVIDIA GPU에서 warp는 32개 thread가 lockstep으로 실행되는 단위다. 같은 warp 안의 thread들이 서로 다른 branch를 타면, GPU는 한쪽 branch를 실행하는 동안 다른 thread들을 비활성화하고, 다시 반대쪽을 실행한다.

이것이 warp divergence다. if가 항상 나쁜 것은 아니지만, warp 안에서 조건이 갈라지면 일부 lane이 놀게 된다. CPU의 branch와 같은 감각으로 보면 놓치기 쉽다.

특히 thread index에 따라 서로 다른 경로를 타는 코드는 조심해야 한다. 빠른 kernel은 종종 "모든 thread가 같은 일을 하되, 다른 data를 만진다"는 단순함을 유지하려고 애쓴다.

3. occupancy는 만능 점수가 아니다

occupancy는 SM에 얼마나 많은 warp/block이 올라와 있는지를 보는 지표다. latency를 숨기려면 기다리는 warp가 충분히 있어야 하므로 occupancy는 중요하다.

하지만 occupancy가 높다고 항상 빠르지는 않다. register를 줄여 occupancy를 올렸더니 instruction이 늘거나 local memory spill이 생기면 오히려 느려질 수 있다. Nsight Compute의 occupancy 계산은 좋은 힌트지만, 최종 판정은 benchmark와 profile이 한다.

강의는 register와 local memory 이야기도 건드린다. 변수가 많고 compiler가 register에 다 담지 못하면 느린 local memory로 spill될 수 있다. "변수 하나 더"가 공짜가 아닐 때가 있다.

4. memory hierarchy와 data locality

많은 GPU kernel은 memory access에 묶인다. global memory에서 값을 읽고 계산한 뒤 다시 쓰는 패턴은 단순하지만, data reuse가 없으면 bandwidth 한계에 빨리 닿는다.

shared memory는 block 안의 thread들이 공유하는 빠른 on-chip memory다. 적절히 쓰면 global memory에서 여러 번 읽을 값을 한 번 가져와 block 안에서 재사용할 수 있다. 하지만 shared memory 자체도 제한된 resource라서, 쓰면 occupancy와 trade-off가 생긴다.

이 지점에서 tiling이 등장한다. 큰 문제를 shared memory에 들어갈 작은 tile로 나누고, tile 안에서 계산을 많이 한다. matmul, convolution, attention이 모두 이 발상 위에 서 있다.

5. arithmetic intensity로 감 잡기

Thomas는 flop per byte, 즉 memory transfer 한 byte당 얼마나 많은 floating-point operation을 하는지 보는 관점을 소개한다. 이 값이 낮으면 memory-bound일 가능성이 크고, 높으면 compute-bound에 가까워진다.

예를 들어 elementwise 연산은 대개 값을 읽고 조금 계산하고 바로 쓴다. 이런 연산은 kernel launch overhead와 memory bandwidth에 민감하다. 여러 elementwise 연산을 fusion하면 같은 data를 여러 번 global memory에 왕복시키지 않아도 된다.

반대로 matmul은 같은 값을 여러 번 재사용할 수 있어 arithmetic intensity를 높일 여지가 크다. 그래서 tiling과 shared memory가 특히 중요해진다.

6. 실전 읽기 포인트

profile을 볼 때 먼저 kernel 시간이 긴지, launch overhead가 많은지, memory throughput이 높은지, occupancy가 낮은지 나눠 보자. 하나의 숫자만 보지 않는 것이 좋다.

그리고 64-bit indexing 같은 사소해 보이는 선택도 성능에 영향을 줄 수 있다. 강의에서 PyTorch kernel을 옮기다 index type 때문에 느려졌던 경험담은 좋은 경고다. GPU에서는 "정확하지만 불필요하게 큰 타입"도 비용이 될 수 있다.

Source Links

• YouTube
• Slides
• Code
• Notebook
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_004.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_004.md

1. 지난 입문에서 한 발 더

이 강의는 Jeremy의 CUDA for Python Programmers 두 번째 편이다. 앞 강의에서 grayscale과 naive matmul을 통해 blocks와 threads를 배웠다면, 여기서는 성능을 위해 memory access를 바꾸기 시작한다.

matrix multiplication은 좋은 무대다. 계산식은 단순하지만 빠르게 만들기는 어렵다. output cell 하나를 만들려면 한 row와 한 column의 dot product가 필요하고, naive 구현은 같은 input 값을 여러 thread가 global memory에서 반복해서 읽는다.

PyTorch baseline과 비교해 정확성을 확인하는 습관은 그대로 유지된다. 먼저 pure Python의 작은 sample로 생각하고, CUDA kernel로 옮기고, 결과가 PyTorch와 가까운지 본다.

2. shared memory의 범위

shared memory는 모든 thread가 다 보는 전역 공간이 아니다. 특정 block 안의 thread들이 공유하는 on-chip memory다. CUDA 프로그래밍 관점에서는 "한 block이 협업할 수 있는 책상" 정도로 생각하면 좋다.

이 제한이 곧 설계의 모양을 만든다. block 하나가 output tile 하나를 맡고, 그 tile 계산에 필요한 input tile들을 shared memory에 올린다. 같은 block의 thread들은 그 shared memory 값을 재사용한다.

물론 shared memory는 작다. 모든 matrix를 올릴 수 없으니 tile로 쪼갠다. 큰 문제를 작은 조각으로 나누고, 조각마다 load, synchronize, compute를 반복한다.

3. Python으로 CUDA의 모양을 흉내 내기

Jeremy는 곧장 CUDA C로 뛰지 않는다. Python으로 dim3 같은 작은 구조를 만들고, block과 thread loop를 직접 돌려 CUDA launch를 흉내 낸다. 처음 보면 우회로 같지만, 실제로는 debugging 비용을 줄이는 지름길이다.

shared memory도 Python tensor로 흉내 낸다. block마다 필요한 크기의 임시 공간을 만들고, thread들이 그 공간을 채우고 읽는다고 생각한다. 이렇게 하면 CUDA의 dynamic shared memory를 쓰기 전에 알고리즘의 순서를 확인할 수 있다.

중요한 변화는 kernel runner의 구조다. 이제 각 thread가 자기 output만 바로 계산하는 것이 아니라, "shared memory 채우기"와 "dot product 계산"이라는 두 단계에 참여한다. 단계 사이에는 동기화가 필요하다.

4. dynamic shared memory

CUDA에서는 launch할 때 dynamic shared memory 크기를 지정할 수 있다. kernel 안에서는 하나의 contiguous block으로 보이고, 프로그래머가 그 안을 A tile과 B tile 영역으로 나눠 쓴다.

강의의 matmul에서는 tile width에 따라 필요한 shared memory가 정해진다. 대략 A tile과 B tile 두 장이 필요하므로 tile width의 제곱에 두 배를 곱한 크기를 생각하게 된다. dtype 크기까지 고려해야 실제 byte 수가 나온다.

이 방식은 유연하지만 책임도 따른다. shared memory index를 잘못 계산하면 서로 다른 값이 겹쳐 쓰이고, 동기화가 빠지면 어떤 thread는 아직 채워지지 않은 값을 읽을 수 있다.

5. 빠름은 memory reuse에서 나온다

naive matmul에서는 같은 input 값이 여러 output cell 계산에 반복해서 쓰인다. shared memory tiling은 그 반복 읽기를 줄인다. global memory에서 한 번 가져온 tile을 block 안에서 여러 번 쓴다.

이때 계산량 자체는 크게 달라지지 않는다. 달라지는 것은 data movement다. GPU 성능 최적화에서 "계산을 줄이지 않았는데 빨라지는" 많은 사례가 바로 여기서 나온다.

Triton이나 CUTLASS가 하는 일도 큰 방향은 비슷하다. tile을 잘 고르고, memory movement를 줄이고, hardware가 좋아하는 모양으로 계산을 배치한다. 이 강의는 그 자동화된 도구들이 숨기는 기본 동작을 손으로 만져 보는 시간이다.

6. 혼자 확인할 것

작은 matrix에서 naive CUDA matmul과 shared memory matmul의 결과를 비교하자. 그 다음 tile width를 바꿔 보며 shared memory 사용량과 실행 시간이 어떻게 달라지는지 본다.

성능이 기대만큼 안 나오더라도 괜찮다. 이 강의의 목표는 최종 matmul library를 만드는 것이 아니라, shared memory를 쓰는 kernel의 구조를 이해하는 것이다. 빠른 GEMM은 이 구조 위에 훨씬 많은 디테일을 더한 물건이다.

Source Links

• YouTube
• Code
• Notebook
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_005.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_005.md

1. optimizer는 학습의 뒷정리가 아니다

training loop를 볼 때 forward와 backward가 주인공처럼 보인다. 하지만 optimizer step도 매 iteration마다 실행된다. Adam류 optimizer라면 parameter, gradient, moment buffer들을 읽고 쓰며 여러 elementwise 연산을 수행한다.

Jane은 PyTorch optimizer의 여러 구현 경로를 비교한다. 단순 for-loop 방식, foreach 방식, fused 방식은 모두 같은 수학을 하더라도 GPU 실행 모양이 다르다. 그리고 GPU에서는 이 실행 모양이 시간을 만든다.

작은 tensor가 많으면 특히 launch overhead가 도드라진다. 각 parameter마다 add, mul, lerp, div 같은 연산이 따로 kernel로 나가면, 계산보다 "일을 시키는 비용"이 커질 수 있다.

2. for-loop에서 foreach로

가장 직관적인 optimizer 구현은 parameter를 하나씩 돌며 update한다. 이해하기 쉽지만, parameter가 많으면 같은 종류의 CUDA kernel을 아주 많이 launch하게 된다.

foreach 구현은 여러 tensor에 대해 같은 연산을 batch처럼 처리한다. Python-level loop를 줄이고, dispatcher 아래에서 묶인 연산을 실행한다. 여기서 이미 launch 수가 줄어드는 효과가 있다.

하지만 foreach도 만능은 아니다. 여러 tensor pointer와 shape/stride 정보를 넘겨야 하고, 실제 kernel이 어떻게 batching을 처리하는지에 따라 overhead가 남는다. 그래도 "하나씩 하지 말고 묶어라"는 방향은 분명하다.

3. fused optimizer의 의미

fusion은 단순히 코드 줄 수를 줄이는 일이 아니다. GPU에서 fusion은 중간 결과를 global memory에 쓰고 다시 읽는 일을 줄이고, 여러 작은 kernel launch를 하나로 합치는 일이다.

Adam update를 생각해 보자. gradient를 읽고, first moment와 second moment를 갱신하고, bias correction과 learning rate를 적용하고, parameter를 갱신한다. 이것을 여러 kernel로 쪼개면 buffer들이 반복해서 global memory를 오간다.

fused kernel은 가능한 한 한 번 읽어 계산하고 한 번 쓴다. 그래서 계산량이 아주 큰 작업보다 memory-bound elementwise update에서 특히 효과가 크다.

4. CUDA kernel의 signature를 읽기

강의는 simplified CUDA kernel signature를 통해 tensor pointer가 어떻게 넘어가는지도 설명한다. PyTorch tensor의 data_ptr는 device memory 주소를 가리키는 큰 정수처럼 보일 수 있다. host에서 그 주소를 직접 dereference하는 것이 아니라, CUDA kernel 안에서 GPU가 사용한다.

stride도 중요하다. tensor가 contiguous라고 가정하면 편하지만 PyTorch tensor는 view, transpose 등으로 다양한 stride를 가질 수 있다. production optimizer는 이런 일반성을 어느 정도 감당해야 한다.

여기서 성능과 일반성의 tension이 생긴다. 더 많은 shape와 stride를 지원하려면 kernel이 복잡해지고, 특화된 fast path를 만들면 API 표면이 늘어난다. PyTorch 내부 구현이 어려운 이유다.

5. launch overhead를 숫자로 의심하기

강의의 메시지는 아주 실용적이다. CUDA kernel launch는 공짜가 아니다. 연산량이 작거나 tensor가 많이 쪼개져 있으면, 계산보다 launch overhead가 더 눈에 띌 수 있다.

그래서 optimizer 최적화는 "더 빠른 곱셈"보다 "몇 번 launch하는가", "같은 data를 몇 번 memory에서 읽는가"로 시작하는 편이 좋다. profiler trace에서 optimizer step 구간을 열어 kernel이 빗방울처럼 흩어져 있다면 fusion 후보일 수 있다.

unified memory 같은 주제도 언급되지만, PyTorch에서 바로 편하게 쓸 수 있는 해법은 아니다. 일반 사용자는 먼저 foreach/fused 옵션과 tensor layout, parameter grouping을 확인하는 쪽이 현실적이다.

6. 실험 과제

작은 모델에서 AdamW를 for-loop, foreach, fused 옵션으로 돌려 보고 profiler trace를 비교해 보자. wall-clock만 보지 말고 optimizer step 안의 kernel launch 개수와 각 kernel 시간을 세어 보면 좋다.

parameter를 일부러 많은 작은 tensor로 쪼갠 모델과, 비교적 큰 tensor 몇 개로 된 모델을 비교해 보자. 같은 optimizer라도 launch overhead가 어떤 경우에 더 크게 보이는지 금방 감이 온다.

Source Links

• YouTube
• Slides
• Code / speaker GitHub
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_006.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_006.md

1. quantization은 왜 GPU kernel 문제가 되는가

GPT Fast나 SAM Fast 같은 사례에서 quantization은 중요한 속도/메모리 도구로 등장한다. 하지만 model weight를 4-bit로 저장한다고 끝나는 일이 아니다. 실제 inference에서 그 압축된 값을 어떻게 읽고, 풀고, 곱하고, 누적할지가 GPU kernel의 문제로 바뀐다.

Charles는 CUDA와 Triton의 경계에서 quantization을 다룬 경험을 공유한다. 예전에는 GPU에서 쓸 만한 quantized kernel이 부족했고, 필요한 모델과 format이 기존 라이브러리와 딱 맞지 않는 경우가 많았다.

여기서 "직접 kernel을 쓴다"는 요구가 나온다. quantization algorithm은 논문 이름으로 끝나지 않고, memory layout과 dequantization 위치와 Tensor Core 사용 가능성으로 내려온다.

2. dynamic, static, weight-only

dynamic quantization은 activation을 실행 중에 관찰해 scale을 정하고 낮은 bit 표현으로 바꾸는 방식이다. compute-bound 상황에서는 도움이 될 수 있고, PyTorch와 torch.compile만으로도 꽤 많은 부분을 처리할 수 있다.

weight-only quantization은 activation은 그대로 두고 weight만 낮은 bit로 저장한다. 이 경우 선택지가 갈린다. weight를 먼저 dequantize해서 일반 matmul을 할 수도 있고, low-bit weight를 읽으며 mixed dtype 곱셈을 시도할 수도 있다.

그래서 "int4"라는 말은 모호하다. W4A16인지, accumulation은 FP32인지, scale은 per-channel인지 per-group인지, zero point가 있는지 없는지까지 말해야 kernel의 일을 알 수 있다.

3. Triton의 장점과 마법성

Triton은 quantized kernel 실험에 매력적이다. CUDA C++보다 짧고, Python 생태계와 가깝고, block 단위로 memory access와 program id를 표현할 수 있다.

강의에서도 Triton kernel이 손으로 쓴 CUDA/CUTLASS 계열 kernel에 꽤 가까운 성능을 내는 사례가 언급된다. 특히 빠르게 실험하고 PyTorch로 붙이는 데 강하다.

하지만 이 편리함은 때로 마법처럼 느껴진다. 왜 빨라졌는지, 왜 느려졌는지 알기 위해서는 결국 generated code, memory access, block size, register pressure를 봐야 한다. Triton은 생각을 없애 주는 도구가 아니라 반복 속도를 높여 주는 도구다.

4. block size와 layout이 발목을 잡는다

low-bit weight는 packed format으로 저장된다. int4 두 개가 byte 하나에 들어갈 수 있고, scale과 metadata도 같이 움직인다. 이때 thread들이 연속적인 memory를 잘 읽는지, unpack 비용이 어디서 발생하는지, block size가 Tensor Core나 cache에 맞는지가 중요하다.

강의에는 block size가 성능을 제한할 수 있다는 관찰이 나온다. 같은 quantization 방식이어도 tile shape와 grouping이 다르면 memory coalescing과 reuse가 달라진다.

또한 quantization은 layer마다 민감도가 다르다. 모든 layer를 같은 bit로 밀어붙이면 정확도나 성능에서 손해를 볼 수 있다. kernel 개발과 model 품질 실험이 서로 떨어져 있지 않다.

5. 정확도와 성능을 같이 봐야 한다

quantization은 성능 최적화이면서 numerical approximation이다. 빠른 kernel이 나와도 model 품질이 무너지면 쓸 수 없다. 반대로 정확도는 괜찮지만 dequantization overhead 때문에 느리면 역시 곤란하다.

그래서 실험은 최소 세 값을 같이 봐야 한다. memory 사용량, latency/throughput, 정확도 지표다. 여기에 kernel 관점의 profiler 지표, 예를 들어 memory bandwidth와 occupancy를 붙이면 원인이 더 잘 보인다.

6. 따라 해볼 읽기 순서

먼저 PyTorch로 dynamic quantization 또는 weight-only quantization의 reference를 만든다. 그 다음 같은 연산을 Triton으로 구현하고, dequantize를 kernel 밖에서 할 때와 안에서 할 때를 비교한다.

마지막으로 generated PTX나 profiler를 보며 packed load, unpack, multiply-accumulate 중 어디가 병목인지 찾는다. "4-bit니까 4배 빠르겠지"라는 생각은 잠시 접어 두는 편이 안전하다.

Source Links

• YouTube
• Slides
• Code / speaker GitHub
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_007.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_007.md

1. 체크리스트가 필요한 이유

Mark는 이 강의를 "CUDA performance checklist"로 잡는다. 빠른 kernel을 만들 때 매번 천재적인 아이디어가 필요한 것은 아니다. 반복해서 확인해야 할 기본 질문들이 있다.

첫 질문은 memory다. GPU에는 큰 DRAM/HBM이 있고, 작지만 빠른 SRAM/shared memory가 있다. 빠르게 만들고 싶다면 가능한 한 느린 memory 왕복을 줄이고, 빠른 memory에서 재사용해야 한다.

다만 SRAM이 빠르니 무한히 넣으면 된다는 이야기는 아니다. SRAM은 면적과 열, 용량의 제약이 있다. 그래서 shared memory는 귀한 자원이고, 잘못 쓰면 occupancy를 깎아먹는다.

2. coalescing부터 의심하자

global memory access가 coalesced인지 보는 것은 기본 중의 기본이다. 인접한 thread들이 인접한 memory를 읽으면 hardware가 효율적으로 transaction을 묶을 수 있다. 반대로 흩어진 주소를 읽으면 bandwidth를 제대로 쓰기 어렵다.

강의 자료의 coalesce.cu는 이 차이를 직접 실험하기 위한 예제다. Nsight Compute에서 memory throughput, L1 cache throughput, replay 관련 지표를 보면 access pattern의 흔적이 보인다.

coalescing은 코드 한 줄의 문제가 아니라 data layout과 thread mapping의 문제다. thread index가 증가할 때 memory address도 자연스럽게 증가하는지 먼저 확인해 보자.

3. occupancy는 heuristic이다

occupancy를 높이면 latency hiding에 도움이 된다. 하지만 강의는 occupancy를 최종 목표로 보지 않는다. Nsight Compute가 theoretical occupancy와 achieved occupancy를 보여 주더라도, kernel 시간이 줄었는지 함께 봐야 한다.

block size를 바꾸면 occupancy가 변한다. CUDA occupancy calculator나 runtime helper가 추천값을 줄 수 있다. 하지만 추천값은 시작점이지 정답지가 아니다.

tile quantization이나 wave quantization처럼, block과 problem size가 딱 맞지 않아 일부 block이 덜 찬 채 실행되는 문제도 있다. 이런 현상은 "thread를 많이 띄웠는데 왜 느리지"라는 질문의 답이 된다.

4. shared memory와 tiling

DRAM을 덜 쓰려면 reuse가 필요하다. shared memory tiling은 global memory에서 가져온 값을 block 안에서 여러 번 쓰게 해 준다. matmul, convolution, softmax, normalization 계열에서 자주 나오는 패턴이다.

하지만 shared memory로 옮기는 것 자체에도 비용이 있다. load, store, synchronization이 필요하고 bank conflict 같은 세부 문제도 생긴다. reuse가 충분하지 않으면 오히려 복잡하기만 할 수 있다.

그래서 checklist의 태도는 "shared memory를 쓰자"가 아니라 "global memory 왕복을 줄일 reuse가 있는가"다. reuse가 있다면 shared memory, register tiling, fusion 중 무엇이 맞는지 고른다.

5. coarsening, privatization, divergence

thread coarsening은 한 thread가 여러 element를 처리하게 하는 기법이다. launch와 indexing overhead를 줄이고 memory access를 더 잘 묶을 수 있지만, register pressure를 늘릴 수도 있다.

privatization은 여러 thread가 같은 global location을 다투지 않도록 중간 결과를 private/local 공간에 모았다가 나중에 합치는 방식이다. reduction이나 histogram류에서 자주 보인다.

divergence는 warp 안의 thread들이 서로 다른 경로를 타는 문제다. branch 자체보다 "같은 warp 안에서 조건이 갈리는가"가 핵심이다. 체크리스트는 이런 문제를 하나씩 profile로 확인하게 해 준다.

6. 좋은 실험의 모양

강의 자료는 coalesce.cu, coarsening.cu, divergence.cu, occupancy.cu, tiling.cu처럼 한 주제에 하나씩 작은 kernel을 둔다. 이 구성이 좋다. 한 번에 모든 최적화를 넣으면 무엇이 효과였는지 모른다.

각 kernel을 compile하고 ncu로 profile한 뒤, 한 가지 가설만 바꿔 다시 측정해 보자. 성능 공부에서 가장 비싼 버릇은 "이게 빠를 것 같아서"를 확인 없이 쌓는 것이다.

Source Links

• YouTube
• Slides
• Code
• Supplemental code
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_008.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_008.md

1. reduction은 어디에나 있다

reduction은 vector의 max를 구하거나 합을 구하는 단순한 예제로 시작한다. 하지만 실제로는 훨씬 자주 나온다. normalization은 평균과 분산을 줄이고, softmax는 max와 sum을 구하며, attention도 tile별 reduction을 품고 있다.

PyTorch에서는 torch.max, torch.sum 같은 API를 부르면 내부에서 CUDA reduction kernel이 실행되거나 생성된다. 사용자는 한 줄을 쓰지만, 아래에서는 많은 thread가 협력해 값을 접는다.

이 강의의 목표는 유명한 reduction algorithm을 naive version에서 시작해 점점 빠르게 만드는 것이다. GPU 성능 공부에 아주 좋은 형태다. 단순하고, 병목이 분명하고, 개선 단계가 눈에 보인다.

2. serial reduction에서 parallel tree로

serial reduction은 한 변수를 들고 왼쪽부터 오른쪽까지 훑는다. GPU에서는 이 방식이 아깝다. 많은 값을 동시에 비교하거나 더할 수 있기 때문이다.

parallel reduction tree는 값을 짝지어 줄인다. 첫 단계에서 여러 thread가 각자 두 값을 합치고, 다음 단계에서 절반의 thread가 다시 합친다. 이렇게 반복하면 log 단계로 하나의 결과에 가까워진다.

하지만 이 과정에는 divergence가 있다. 단계가 진행될수록 active thread 수가 줄어든다. warp 단위 실행에서는 놀고 있는 lane이 생긴다. 좋은 kernel은 이 비효율을 줄이기 위해 indexing과 memory layout을 조정한다.

3. global memory에서 shared memory로

naive reduction은 중간 결과를 global memory에 쓰고 다시 읽을 수 있다. 이는 느리다. block 안에서 처리할 수 있는 부분은 shared memory에 올려 두고 줄이는 편이 낫다.

강의의 shared memory version은 첫 load 이후 reduction을 shared memory에서 진행한다. Nsight Compute로 보면 cache hit나 memory throughput 지표가 달라질 수 있다. 다만 shared memory에 전체 input이 들어간다고 가정하면 큰 입력에서 틀린 결과가 나온다.

이 지점이 중요하다. 작은 예제에서 맞는 kernel이 큰 문제에서 틀릴 수 있다. reduction은 block-local 결과와 global result를 나누어 생각해야 한다.

4. correctness가 은근히 어렵다

floating point reduction은 순서에 따라 결과가 조금 달라질 수 있다. 병렬 reduction은 serial loop와 덧셈 순서가 다르므로 bitwise로 완전히 같지 않을 수 있다. 이것은 bug일 수도 있고 floating point의 정상적인 성질일 수도 있다.

강의 자료에 accuracy.py, nondeterminism.py, sensitivity.py가 있는 이유가 여기에 있다. reduction은 성능뿐 아니라 numerical behavior도 함께 봐야 한다.

max reduction처럼 비교 중심인 경우와 sum reduction처럼 누적 오차가 생기는 경우를 구분하자. softmax에서는 안정성을 위해 max를 먼저 빼는 stabilized softmax가 필요하다.

5. PyTorch의 일반 reduction을 상상하기

PyTorch에는 min kernel, max kernel이 각각 완전히 따로 있다기보다, 다양한 reduction을 처리하는 일반화된 machinery가 있다. dtype, dimension, stride, contiguous 여부를 감당해야 하므로 단순 예제보다 훨씬 복잡하다.

우리가 작은 CUDA 파일로 배우는 것은 그 내부의 뼈대다. thread가 일부 값을 맡고, block 안에서 줄이고, block 결과를 다시 합친다. 실제 framework는 여기에 broadcasting, shape, stride, dispatch, generated kernel 같은 층을 붙인다.

Triton으로도 reduction을 작성할 수 있다. block size와 mask, tl.max, tl.sum 같은 primitive가 나오지만, 원리는 같다. 어디서 줄이고 어디서 memory를 오가는지가 핵심이다.

6. 실험 과제

simple_reduce.cu, shared_reduce.cu, reduce_coarsening.cu를 차례로 읽고 같은 입력에서 결과와 profile을 비교해 보자. 입력 크기를 shared memory에 들어가는 크기와 훨씬 큰 크기로 나누어 실행하면 가정이 어디서 깨지는지 보인다.

그 다음 softmax를 reduction 관점에서 다시 보자. max reduction, exp, sum reduction, normalize가 어떤 순서로 이어지는지 적어 보면 attention kernel이 덜 낯설어진다.

Source Links

• YouTube
• Slides
• Code
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_009.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_009.md

1. CUDA ninja 혼자서는 제품이 안 된다

Oscar의 출발점은 경험적이다. CUDA kernel을 좋아하고 잘 쓰는 사람은 팀에 필요하다. 하지만 real application에서는 그 사람이 모든 호출 지점을 직접 붙잡고 있을 수 없다.

그래서 목표는 "모두를 CUDA ninja로 만들기"가 아니라, CUDA를 모르는 사용자도 좋은 성능을 얻을 수 있는 라이브러리를 제공하는 것이다. abstraction은 여기서 성능의 적이 아니라 배포의 조건이 된다.

물론 abstraction은 공짜가 아니다. 너무 높은 abstraction은 성능과 제어를 잃고, 너무 낮은 abstraction은 사용자에게 CUDA 세부사항을 떠넘긴다. production library는 이 중간 지점을 찾는 일이다.

2. host side가 먼저 보인다

CUDA 최적화라고 하면 kernel 내부만 떠올리기 쉽다. 하지만 Oscar는 실제 application에서는 host side code를 먼저 보아야 한다고 말한다. CPU thread가 GPU에 일을 어떻게 enqueue하는지, 어디서 synchronize되는지, memory allocation이 어디서 block되는지가 중요하다.

stream을 쓰면 copy와 compute를 겹칠 수 있다. 하지만 stream을 쓴다고 자동으로 overlap이 생기지는 않는다. dependency와 synchronization 지점이 잘못되면 CPU도 GPU도 기다린다.

production 환경에서는 allocation도 신경 써야 한다. 반복 루프 안에서 매번 GPU memory를 할당하고 해제하면 latency가 튄다. memory manager나 buffer pool이 필요한 이유다.

3. vertical fusion과 horizontal fusion

강의는 fusion을 kernel 내부 수식 합치기보다 더 넓게 다룬다. vertical fusion은 producer와 consumer 사이의 latency를 줄이는 방향이다. 한 단계의 output을 다음 단계가 바로 소비하게 만들어 memory 왕복이나 기다림을 줄인다.

horizontal 또는 concurrent fusion은 독립적인 작업들을 동시에 진행해 자원을 더 잘 채우는 쪽에 가깝다. 여러 stream, 여러 copy, 여러 kernel을 어떻게 겹칠지 보는 문제다.

이 구분은 유용하다. 어떤 병목은 memory-bound kernel 자체에 있고, 어떤 병목은 host scheduling이나 stream dependency에 있다. 둘을 같은 "fusion"이라는 말로 뭉개면 처방이 흐려진다.

4. producer-consumer와 delay buffer

강의 후반의 중요한 그림은 producer-consumer model이다. producer가 데이터를 만들고 consumer가 사용한다. 목표는 consumer가 producer를 기다리지 않게 하고, producer도 빈 buffer를 기다리지 않게 하는 것이다.

이를 위해 여러 buffer를 돌려 쓰는 방식이 나온다. 어떤 pointer는 몇 iteration 전에 synchronize되어 이제 안전하게 재사용할 수 있다. Oscar는 이런 delay buffer 개념으로 stream 작업의 pipeline을 만든다.

이 방식은 단순한 double buffering보다 일반적이다. real-time 또는 soft real-time application에서는 latency spike를 줄이고 일정한 처리 흐름을 만드는 것이 중요하다. kernel time만 빠른 것과 제품이 부드럽게 도는 것은 다르다.

5. Python binding은 사용성을 넓힌다

강의는 C++/CUDA 중심이지만 Python 사용자도 염두에 둔다. 딥러닝 쪽 사용자의 상당수는 Python에서 작업하므로, 라이브러리가 Python에서 접근 가능해야 더 넓게 쓰인다.

다만 Python binding은 얇은 포장이어야 한다. performance-critical path는 C++/CUDA 쪽에서 관리하고, Python은 configuration과 high-level 호출을 맡는 편이 좋다.

여기서 API 설계가 중요해진다. 사용자가 CUDA stream과 buffer lifetime을 전부 알아야 한다면 라이브러리라고 부르기 어렵다. 하지만 필요할 때 advanced user가 제어할 수 있는 escape hatch는 남겨야 한다.

6. production-ready의 뜻

이 강의에서 production-ready는 "벤치마크가 빠르다"보다 넓다. 설치와 빌드가 가능하고, 사용자 API가 안정적이고, memory와 stream lifetime이 안전하며, profiling 가능한 구조를 갖춘 상태에 가깝다.

CUDA kernel 개발자는 종종 device code에 애정을 많이 준다. 그 마음은 이해된다. 하지만 제품에서는 host code, packaging, logging, error handling, synchronization policy가 같은 무게로 중요해진다.

Source Links

• YouTube
• Slides folder
• Code / speaker GitHub
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_010.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_010.md

1. 0을 곱하지 않는다는 말의 함정

수학적으로 0 곱하기 무엇은 즉시 0이다. 그렇다면 sparse tensor는 당연히 빨라야 할 것 같다. 하지만 dense matmul 방식 그대로 모든 위치를 훑으며 0을 곱하면 거의 이득이 없다.

Jesse는 이 차이를 분명히 한다. sparsity는 "0이 많다"가 아니라 "0을 건너뛸 수 있는 representation과 kernel이 있다"까지 포함해야 성능 기술이 된다.

따라서 sparse representation이 필요하다. nonzero 값만 저장하거나, block 단위 mask를 저장하거나, hardware가 이해하는 semi-structured pattern을 사용한다. 어떤 format을 고르느냐가 kernel의 모양을 결정한다.

2. unstructured, block, semi-structured

unstructured sparsity는 위치가 자유롭다. 가장 flexible하지만 GPU가 좋아하는 규칙적인 memory access와는 거리가 멀다. metadata가 많고 load pattern이 흩어질 수 있다.

block sparsity는 작은 block 단위로 0인지 아닌지를 정한다. 예를 들어 32x32나 4x4 같은 block을 생각할 수 있다. flexibility는 줄지만, block 단위로 건너뛰고 dense-like 연산을 할 수 있어 GPU에 더 친절하다.

semi-structured sparsity는 hardware나 library가 지원하는 특정 pattern을 따른다. 제약이 강하지만, 지원되는 path에서는 성능을 얻기 쉽다. "얼마나 자유로운 sparsity인가"와 "얼마나 빠르게 실행 가능한가"는 trade-off다.

3. pruning은 offline, sparse matmul은 online

sparsity는 보통 두 단계로 나눠 생각하면 편하다. 먼저 offline에서 어떤 weight를 0으로 만들지 결정하고 sparse format으로 변환한다. inference 때는 compressed representation과 sparse kernel을 사용한다.

이 구분은 latency를 이해하는 데 중요하다. pruning 자체가 inference path에 들어오면 느리다. inference에서는 이미 준비된 sparse weight를 빠르게 읽고 계산해야 한다.

하지만 model을 sparse하게 만드는 방법도 정확도에 영향을 준다. 작은 값부터 제거할지, retraining을 할지, SparseGPT 같은 방법을 쓸지에 따라 품질이 달라진다.

4. quantization과 같이 쓰면 더 복잡하다

강의는 sparsity와 quantization을 함께 쓰는 어려움도 다룬다. 둘 다 memory와 compute를 줄이려는 기술이지만, composition은 단순 덧셈이 아니다.

예를 들어 sparse matmul을 cuSPARSELt 같은 library call로 처리하고, 앞뒤에 quantize/dequantize kernel을 붙이면 operator fusion이 깨질 수 있다. memory 왕복이 늘어나면 sparsity에서 얻은 이득을 다시 잃는다.

반대로 sparse와 quantization을 하나의 fused kernel로 묶으려면 kernel 작성 난도가 올라간다. metadata, scale, packed value, sparse index를 모두 다뤄야 한다. 이 지점이 active research와 engineering의 경계다.

5. library 지원의 현실

GPU sparse kernel은 dense kernel만큼 보편적이지 않다. dense GEMM은 cuBLAS, CUTLASS, Triton 등 강한 도구가 많지만, sparse는 format과 pattern에 따라 지원 범위가 갈린다.

강의에서는 sparse kernel이 없어서 직접 구현하거나, library를 쓰되 fusion이 안 되어 성능이 기대보다 낮아지는 상황이 나온다. 이것이 sparsity가 논문 그림만큼 쉽게 제품 성능으로 이어지지 않는 이유다.

또한 sparse format은 model architecture와도 맞아야 한다. LLM과 Vision Transformer에서 어떤 layer를 sparse하게 만들지, activation sparsity를 볼지 weight sparsity를 볼지에 따라 적용점이 달라진다.

6. 실험 과제

작은 matrix를 dense, unstructured sparse, block sparse 세 format으로 만들어 보자. nonzero 비율이 같아도 memory access와 metadata 크기가 어떻게 다른지 계산해 보는 것이 먼저다.

그 다음 sparse matmul 앞뒤에 dequantize나 activation function을 붙였을 때 kernel launch와 memory traffic이 어떻게 늘어나는지 profiler로 확인해 보자. sparsity의 적은 종종 "0"이 아니라 "그 0을 설명하는 부가 정보"다.

Source Links

• YouTube
• Slides
• Code / speaker GitHub
• Related folder
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_011.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_011.md

1. attention을 memory 문제로 보기

attention은 두 번의 matmul과 softmax로 이루어진다. 평범하게 구현하면 QK^T score matrix를 만들고, softmax를 적용하고, 다시 V와 곱한다. sequence length가 길어지면 score matrix가 커져 memory가 빠르게 부담된다.

FlashAttention의 핵심은 이 중간 matrix를 크게 materialize하지 않는 것이다. tile을 읽고, 필요한 softmax 통계를 갱신하고, output accumulator를 업데이트한다. HBM 왕복을 줄이는 것이 목표다.

Thomas는 이를 matmul tiling의 연장선에서 설명한다. shared memory/SRAM으로 tile을 옮기고, 그 안에서 계산을 많이 한 뒤, 필요한 결과만 밖으로 내보낸다.

2. softmax가 문제를 어렵게 만든다

QK^T만 있었다면 tiled matmul로 생각하면 된다. 하지만 softmax는 row 전체의 max와 sum을 필요로 한다. tile 하나만 보고는 전체 softmax를 바로 알 수 없다.

그래서 online softmax가 필요하다. 새 tile을 볼 때마다 현재까지의 max와 normalization sum을 갱신하고, 기존 output accumulator를 새 max 기준으로 rescale한다. 이 trick 덕분에 score row 전체를 한 번에 들고 있지 않아도 된다.

stabilized softmax는 float32에서도 중요하다. 큰 logit을 그대로 exponentiate하면 overflow가 날 수 있다. max를 빼는 안정화는 attention kernel에서 성능만큼이나 correctness의 문제다.

3. block과 head의 mapping

CUDA kernel에서는 grid의 block들이 각자 attention 작업의 일부를 맡는다. 강의에서는 head dimension과 sequence tile을 어떻게 block에 배치할지 이야기한다. 한 block 안의 cooperating thread들이 shared memory를 써야 하므로, block mapping은 매우 중요하다.

head dimension이 크거나 tile size가 커지면 shared memory 사용량이 늘어난다. shared memory가 부족하면 tile을 줄이거나 register와 memory 사용을 다시 조정해야 한다.

이런 선택은 단순히 코드 취향이 아니다. block size, shared memory, register 수는 occupancy와 spill에 직접 영향을 준다. FlashAttention 구현이 어려운 이유는 알고리즘보다 이 resource juggling에 있다.

4. register spilling과 profiler

Thomas는 register spilling도 강조한다. thread가 너무 많은 local 상태를 들고 있으면 compiler가 register에 다 넣지 못하고 local memory로 spill할 수 있다. 그러면 빠르려고 만든 kernel이 느린 memory를 다시 두드리게 된다.

Nsight Compute 같은 profiler는 register 사용량, spill, occupancy 관련 신호를 준다. FlashAttention처럼 복잡한 kernel에서는 코드만 보고 성능을 확신하기 어렵다.

강의 자료에는 CUDA Python, inline compilation, Numba 같은 경로도 나온다. 어떤 도구를 쓰든 핵심은 같다. generated kernel이 실제로 shared memory와 register를 어떻게 쓰는지 확인해야 한다.

5. compiler fusion과 hand-written kernel

자동 compiler fusion은 elementwise나 memory-bound 연산에서 큰 도움을 준다. 하지만 FlashAttention처럼 softmax 통계와 matmul tiling이 얽힌 알고리즘을 compiler가 일반 PyTorch graph에서 완전히 회복하기는 어렵다.

그래서 Triton tutorial이나 hand-written CUDA 구현이 여전히 의미를 갖는다. 알고리즘 자체가 memory traffic을 줄이도록 재배치되어 있기 때문이다. 단순히 여러 op를 붙이는 fusion과는 층위가 다르다.

FlashAttention은 "attention을 빠르게 하는 kernel"이면서 동시에 GPU programming 교과서다. tiling, shared memory, online reduction, numerical stability, register pressure가 한데 모여 있다.

6. 읽고 구현할 때의 순서

먼저 작은 attention을 PyTorch로 구현하고 intermediate score matrix의 shape와 memory size를 계산해 보자. 그 다음 FlashAttention pseudo-code에서 어떤 값만 유지하면 되는지 표시한다.

작은 tile size로 CUDA 또는 Triton version을 만들어 correctness를 맞추고, 마지막에 profiler로 shared memory 사용량과 register spill을 본다. 처음부터 fastest FlashAttention을 목표로 잡으면 마음이 좀 다친다. 작게, 정확하게, 그 다음 빠르게.

Source Links

• YouTube
• Code
• CUDA file
• Notebook
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_012.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_012.md

1. 긴 문맥은 왜 아픈가

long context 모델은 매력적이다. 긴 문서, audio, video, multimodal 기록을 더 많이 넣을 수 있고, world model 같은 방향에서도 큰 context는 중요한 재료가 된다.

하지만 attention은 sequence length에 대해 memory가 빠르게 커진다. 강의에서는 Ring Attention paper의 예를 빌려, batch size 1이어도 매우 긴 token 수에서는 memory가 엄청나게 필요하다는 점을 강조한다.

여기서 먼저 솔직해야 한다. Ring Attention은 attention의 quadratic scaling 자체를 없애지 않는다. 필요한 계산은 여전히 많이 한다. 다만 memory와 communication을 조직해서 더 긴 sequence를 여러 device에서 처리하려 한다.

2. FlashAttention에서 이어지는 blockwise softmax

Ring Attention을 이해하려면 FlashAttention의 online softmax를 먼저 잡아야 한다. attention row를 block으로 나누어 보고, 각 block의 max와 normalization 정보를 합쳐도 전체 softmax와 같은 결과를 만들 수 있다.

Andreas는 softmax를 blockwise로 계산하는 법을 길게 설명한다. 한 block에서 얻은 local max와 sum, output을 다음 block의 정보와 결합한다. 새 max가 더 크면 이전 accumulator를 rescale해야 한다.

이 결합 법칙이 중요하다. block을 어떤 순서로 보더라도 필요한 통계를 갱신할 수 있어야 sequence를 나누고 돌릴 수 있다. Ring Attention은 이 성질을 distributed setting으로 가져간다.

3. distributed parallelism의 위치

training에는 data parallelism, tensor/model parallelism, pipeline parallelism 등 여러 형태가 있다. Ring Attention은 sequence dimension을 나누는 sequence parallelism 쪽으로 생각할 수 있다.

각 device가 sequence의 일부를 맡고, attention에 필요한 K/V block을 ring 형태로 주고받는다. 이렇게 하면 모든 device가 전체 K/V를 한 번에 들고 있지 않아도 된다.

다만 communication이 사라지는 것은 아니다. block을 넘기는 비용과 계산을 어떻게 overlap할지가 성능의 핵심이 된다. 긴 context에서는 memory를 줄이는 대신 communication schedule이 새 병목이 될 수 있다.

4. inference와 training을 구분하자

강의에서 중요한 경고는 Ring Attention이 주로 training 쪽 기술로 보인다는 점이다. token-by-token inference에서는 KV cache와 decoding pattern이 다르고, FlashAttention 자체도 long context inference에 그대로 최적화되어 있지 않을 수 있다.

상용 long-context 모델들이 내부적으로 무엇을 쓰는지는 공개되어 있지 않은 경우가 많다. Ring Attention과 비슷한 아이디어가 들어갈 수는 있지만, 확정적으로 말하기 어렵다.

따라서 독자는 "long context = Ring Attention"으로 외우기보다, 어떤 workload에서 memory가 문제인지, compute가 문제인지, communication이 문제인지 나누어 보아야 한다.

5. reduction step의 비용

blockwise softmax는 여러 block의 통계를 합치는 reduction을 필요로 한다. Ring Attention에서도 각 device가 본 K/V block들에 대해 online state를 갱신하고, 최종 output을 얻는다.

이 reduction step은 attention matmul에 비해 작아 보일 수 있지만, 긴 sequence와 여러 device에서는 무시하기 어렵다. profiler로 communication, matmul, reduction이 각각 어디서 시간을 쓰는지 봐야 한다.

Andreas는 shared memory 질문에도 선을 긋는다. sequence parallelism은 device 간 분할 전략이고, 각 device 내부의 CUDA kernel은 여전히 shared memory를 쓸 수 있다. 층위를 섞지 않는 것이 중요하다.

6. 읽기와 실험 포인트

먼저 howto_log_sum_exp.ipynb처럼 log-sum-exp와 blockwise softmax 결합을 작은 숫자로 따라가 보자. 이 부분이 잡히면 Ring Attention의 수식은 훨씬 덜 무섭다.

그 다음 2개 또는 4개 device가 있다고 가정하고 K/V block이 ring으로 이동하는 표를 그려 보자. 각 step에서 어떤 device가 어떤 block을 보고, online max/sum/output을 어떻게 갱신하는지 적는다.

Source Links

• YouTube
• Slides
• Slides Google
• Notebook
• Code
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_013.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_013.md

1. 왜 Triton을 쓰는가

Umer Adil의 강의는 “Triton이 멋지다”에서 시작하지 않는다. 시작점은 더 실용적이다. PyTorch 연산 조합만으로는 느리고, CUDA C++로 내려가기는 부담스러울 때, 중간에 놓인 선택지가 Triton이다. Python 안에서 kernel을 쓰되, GPU memory access와 tile shape는 직접 생각한다.

그래서 Triton은 좋은 첫 커널을 빨리 쓰게 해준다. vector add 같은 커널은 금방 나온다. 하지만 “빠른” 커널은 별개의 문제다. block size, num_warps, num_stages, mask 처리, load/store 패턴을 놓치면 PyTorch보다 느린 커스텀 커널이 된다. GPU 프로그래밍에서 이것보다 슬픈 일도 많지만, 꽤 대표적인 슬픔이다.

CUDA를 대체한다기보다, CUDA로 내려가기 전에 머무를 수 있는 넓은 작업대에 가깝다. 아주 미세한 shared memory layout, bank conflict, 특정 instruction scheduling까지 붙잡아야 하면 CUDA가 필요할 수 있다. 반대로 LLM의 pointwise fusion, 작은 matmul 변형, custom activation처럼 tile 단위로 표현이 잘 되는 일은 Triton이 꽤 자연스럽다.

2. program_id가 blockIdx처럼 보일 때 조심하자

CUDA를 배운 사람은 program_id를 보면 곧장 blockIdx를 떠올린다. 완전히 틀린 비유는 아니다. 다만 Triton에서 한 program은 thread block이라기보다 출력 tile을 계산하는 논리적 일감이다. 그 안에서 tl.arange(0, BLOCK_SIZE)가 tile 안의 element index vector를 만든다.

vector add를 보면 구조가 선명하다. pid = tl.program_id(0)로 내가 몇 번째 tile인지 알고, offsets = pid * BLOCK_SIZE + tl.arange(...)로 이 tile이 읽을 위치를 만든다. tl.load(x + offsets, mask=offsets < n)로 경계 밖을 막고, 계산한 뒤 tl.store한다. 코드는 NumPy처럼 보이지만 실제로는 tile 전체를 대상으로 하는 SIMD식 표현이다.

matmul도 같은 생각의 2D 확장이다. pid_m, pid_n이 output matrix의 tile 좌표가 되고, M/N/K 방향의 offset vector를 만든 뒤 K축을 따라 누적한다. 처음부터 Tensor Core를 떠올리기보다, “내 program은 C의 어느 직사각형을 책임지는가?”를 먼저 묻는 편이 좋다.

3. 디버깅은 interpret 모드부터

Triton의 친절한 부분 하나는 TRITON_INTERPRET=1이다. kernel을 GPU에서 바로 돌리지 않고 CPU 쪽 interpreter로 실행하면, kernel 안에서 breakpoint()와 print를 쓸 수 있다. 처음 tl.load의 모양이 머릿속과 다른지 확인하기에 아주 좋다.

물론 이것은 성능 측정용이 아니다. interpreter는 느리고, 모든 GPU 동작을 똑같이 재현하지도 않는다. 그래도 학습 단계에서는 큰 선물이다. CUDA kernel 안에서 printf를 뿌리며 어둠 속을 걷는 것보다, 특정 program_id에서 멈춰 offset과 mask를 보는 편이 훨씬 덜 고독하다.

강의 notebook의 triton_util.py도 이런 용도다. 특정 program에서만 breakpoint를 걸거나 print하게 해준다. tile 단위 사고가 흔들릴 때는 전체 실행을 보지 말고, program 하나를 손에 들고 보는 것이 빠르다.

4. benchmarking과 autotune은 “예쁜 코드”를 깨운다

Triton에는 benchmark helper와 autotune이 있다. 여기서 중요한 점은 autotune이 마법사가 아니라는 것이다. 후보 configuration을 사람이 정해주고, Triton이 그중 좋은 것을 고른다. 후보가 나쁘면 결과도 나쁘다.

가장 자주 만지는 값은 BLOCK_SIZE, num_warps, num_stages다. BLOCK_SIZE는 tile 크기와 memory transaction 모양을 바꾸고, num_warps는 한 program을 수행하는 병렬성을 바꾸며, num_stages는 pipelining 깊이에 영향을 준다. 작은 matrix에서는 PyTorch가 더 빠를 수 있고, 큰 matrix에서는 Triton kernel이 따라잡거나 앞설 수 있다. 그래서 size sweep 없이 “빠르다”라고 말하면 보통 너무 이르다.

실전에서는 baseline을 꼭 둔다. PyTorch eager, torch.compile, 기존 Triton 예제, 가능하면 cuBLAS나 CUTLASS까지 놓고 본다. 커스텀 kernel의 목표는 자존심이 아니라 병목 제거다.

5. Triton의 rough edge도 기술의 일부다

Umer는 Triton이 아직 거친 부분이 있다고 솔직히 말한다. compiler error가 친절하지 않을 때가 있고, 어떤 코드는 기대와 다르게 compile되거나 성능이 흔들린다. 이건 Triton을 피하라는 뜻이 아니라, 디버깅과 측정을 더 작은 단위로 하라는 뜻이다.

PTX를 dump해서 보는 습관도 여기서 나온다. 내가 쓴 tile 연산이 실제로 어떤 load/store와 instruction으로 내려갔는지 확인하면, “컴파일러가 알아서 해주겠지”의 범위가 조금씩 보인다. Triton은 추상화지만, 추상화 아래에 HBM, register, warp, instruction이 사라지는 것은 아니다.

6. 손으로 해볼 것

먼저 vector add를 직접 쓰고, 일부러 mask를 빼서 마지막 tile에서 무슨 일이 나는지 본다. 그 다음 interpreter로 특정 program_id의 offsets를 출력한다. 마지막으로 block size를 64, 128, 256, 1024로 바꿔 작은 N과 큰 N에서 PyTorch와 비교한다.

matmul은 처음부터 “최고 성능”을 목표로 하지 말자. naive Triton matmul을 쓰고, BLOCK_M/N/K, num_warps, num_stages만 sweep해도 꽤 많은 것을 배운다. 성능 숫자보다 중요한 질문은 이것이다. 이 kernel은 memory-bound인가, compute-bound인가, 아니면 launch overhead에 묶였는가?

Source Links

• YouTube
• Lecture 014 repo folder
• Notebook
• triton_util.py
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_014.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_014.md

1. 왜 굳이 CUTLASS를 배우나

Eric Auld는 CUTLASS를 처음부터 만만한 도구로 소개하지 않는다. 오히려 “배울 만한 보상이 있지만, 진입 장벽도 있다”는 톤이다. cuBLAS, cuDNN 같은 host-side library는 아주 좋다. 정해진 연산을 호출하면 library가 Tensor Core 사용 여부, kernel 선택, 여러 최적화를 알아서 처리한다.

문제는 모델이 늘 정해진 연산만 하지 않는다는 데 있다. FlashAttention처럼 fusion이 필요하거나, epilogue에 특이한 연산이 붙거나, layout이 표준적이지 않거나, matmul과 그 주변을 한 kernel 안에서 만지고 싶을 때가 온다. 이때 CUTLASS는 vendor급 kernel을 완전히 새로 쓰지 않고 조립할 수 있는 부품 상자가 된다.

다만 이 부품 상자는 C++ template로 되어 있고, 표기법도 낯설다. 강의 초반의 “underscore convention” 이야기가 그래서 중요하다. A_mk 같은 이름은 그냥 예쁜 변수명이 아니라 tensor의 mode와 index를 표시하는 힌트다. CUTLASS 코드는 변수명부터 좌표계다.

2. CuTe의 중심: tensor는 pointer와 layout이다

CuTe에서 tensor는 크게 두 가지를 가진다. 실제 data를 가리키는 engine, 그리고 좌표를 offset으로 바꾸는 layout이다. layout은 shape와 stride로 이루어진다. “2차원 tensor”라고 말해도, 그것이 row-major인지 column-major인지, 혹은 더 복잡한 nested layout인지는 layout이 정한다.

이 관점은 처음에는 과하게 형식적이다. 하지만 GPU kernel에서는 이 형식성이 힘이 된다. coalesced load, shared memory tile, MMA operand layout은 모두 “좌표를 어떤 순서로 펴는가”에 달려 있다. CUTLASS는 이걸 주먹구구 index arithmetic 대신 composable layout 연산으로 다루려 한다.

3. tiling은 큰 tensor를 작은 말로 바꾸는 일

GEMM을 빠르게 하려면 큰 matrix를 tile로 잘라야 한다. 이 말 자체는 익숙하다. CUTLASS/CuTe가 흥미로운 지점은 tiling을 단순히 “M을 128씩 자른다”가 아니라 layout의 변환으로 본다는 점이다. 큰 tensor의 좌표를 outer tile 좌표와 inner tile 좌표로 나눈다.

이 분해가 되면 같은 언어로 여러 층의 tile을 표현할 수 있다. thread block tile, warp tile, MMA instruction tile이 모두 같은 문제의 다른 해상도다. Tensor Core는 특정 shape의 MMA를 수행하고, kernel은 그 instruction tile을 더 큰 output tile 속에 배치한다.

4. Tensor Core와 library의 거리

Tensor Core를 직접 쓰려면 operand layout, register 배치, shared memory swizzle, instruction shape를 알아야 한다. CUTLASS는 이 작업을 모두 숨기지는 않는다. 대신 숨길 것과 드러낼 것을 나눈다. 사용자는 tile shape, element type, layout, epilogue 같은 고수준 선택을 하고, library는 architecture-specific machinery를 제공한다.

이 강의는 FlashAttention을 예로 드는 이유가 분명하다. FlashAttention은 단순히 matmul 하나가 아니라 QK, softmax, PV, online normalization이 얽힌 fused algorithm이다. 이런 곳에서는 “좋은 GEMM 호출”만으로 끝나지 않는다. 연산을 어떤 tile 단위로 streaming하고, 어떤 중간값을 register/shared memory에 둘지 결정해야 한다.

5. 읽는 법: 수식보다 좌표를 따라가자

CUTLASS 코드를 읽을 때 가장 좋은 습관은 template 이름을 한 번에 이해하려 하지 않는 것이다. 먼저 tensor의 shape, layout, tiling을 찾는다. 그 다음 어떤 coordinate가 어떤 layout을 거쳐 memory offset이 되는지 따라간다. 수식처럼 읽되, 실제 pointer까지 내려간다.

또 하나는 “static한 것과 dynamic한 것”을 구분하는 것이다. 큰 문제 크기는 runtime에 올 수 있지만, 작은 tile shape와 layout은 compile time에 박혀 있을 때가 많다. GPU compiler가 강하게 최적화하려면 이런 compile-time 정보가 매우 중요하다. CUTLASS의 긴 type은 이 정보를 compiler에게 먹이는 방식이기도 하다.

6. 손으로 해볼 것

작은 2D array를 하나 정하고 shape (M, N), stride (N, 1)과 (1, M)을 직접 써보자. 같은 coordinate (i, j)가 row-major와 column-major에서 offset이 어떻게 달라지는지 계산한다. 그 다음 4x4 matrix를 2x2 tile로 자르고, outer coordinate와 inner coordinate를 나눠 써본다.

코드를 본다면 CUTLASS 전체를 바로 빌드하려 하지 말고 CuTe layout 예제를 먼저 읽는 편이 좋다. make_shape, make_stride, make_layout, make_tensor가 어떤 좌표계를 만드는지 종이에 그리면, template 괄호가 조금 덜 위협적으로 보인다.

Source Links

• YouTube
• Eric Auld GitHub
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_015.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_015.md

1. 슬라이드보다 trace를 본 강의

이 강의는 준비된 슬라이드 강의라기보다 실제 모델을 열고 profiling view를 오가며 보는 세션이다. Taylor의 배경도 그쪽이다. Lightning AI와 PyTorch 쪽에서 profiling, benchmarking, model optimization을 다뤄온 사람답게, “이론적으로 병목일 것 같은 곳”보다 trace에서 보이는 증거를 먼저 잡는다.

GPU 성능 문제를 볼 때 흔한 실수는 kernel 이름만 보는 것이다. gemm, embedding, pointwise 같은 이름이 보이면 뭔가 안 듯하지만, 실제 질문은 더 구체적이다. 이 kernel은 어느 Python line에서 왔나? 앞뒤에 host gap이 있나? batch size를 올리면 같은 문제가 유지되나? memory peak 때문에 GPU가 굶고 있나?

2. PyTorch profiler가 보여주는 것

PyTorch profiler는 Chrome trace 형태로 CPU runtime, CUDA runtime, GPU kernel lane을 함께 보여준다. kernel을 클릭하면 stack trace를 통해 어느 Python 호출에서 launch되었는지 추적할 수 있다. 이 연결이 아주 중요하다. GPU lane만 보면 “무언가 느리다”이고, stack trace까지 보면 “내 코드의 이 부분이 느리다”가 된다.

Timeline에서 먼저 볼 것은 gap이다. GPU가 쉬고 있는데 CPU가 kernel을 준비하고 있다면 host overhead나 synchronization을 의심한다. 여러 kernel이 겹쳐 보인다면 CUDA stream과 scheduler가 어떻게 일을 배치하는지도 본다. 정상적인 CUDA 프로그램은 host가 launch하고, GPU는 나중에 실행한다. 이 비동기성을 모르면 trace를 잘못 읽기 쉽다.

3. 빠른 kernel이 빠른 모델을 보장하지 않는다

강의 중 한 예시는 embedding kernel이다. kernel만 보면 아주 빠르거나, baseline보다 두 배 빨라 보일 수 있다. 그런데 모델에 넣으면 전체 속도가 기대만큼 나오지 않는다. 이유는 간단할 수 있다. kernel 자체가 병목이 아니거나, batch size가 너무 작아 GPU를 채우지 못하거나, 앞뒤 memory movement가 더 크다.

Taylor가 반복해서 보는 축은 batch size다. batch size 1에서는 많은 inference workload가 memory-bound처럼 보이고, GPU SM을 충분히 쓰지 못한다. batch를 키우면 arithmetic intensity가 달라지고, kernel launch overhead도 상대적으로 작아진다. 반대로 memory peak가 커져서 더 이상 batch를 못 올리는 경우도 있다.

4. memory profiler는 “왜 batch를 못 키우나”를 말해준다

성능 문제는 종종 시간보다 memory에서 시작한다. PyTorch의 memory profiler는 allocation timeline과 peak memory를 보여준다. 어떤 첫 batch가 유난히 크거나, activation이 오래 살아남거나, model object가 scope 밖에서도 참조되어 memory를 잡고 있으면 batch size를 못 올린다.

이런 문제는 kernel 최적화로 해결되지 않는다. reference를 끊거나, activation lifetime을 줄이거나, gradient accumulation과 micro-batch를 조정해야 한다. profiling의 좋은 점은 이런 결정을 감으로 하지 않게 만든다는 것이다. “느린 것 같다” 대신 “여기서 peak가 생겨 batch 4가 안 된다”라고 말할 수 있다.

5. flame graph와 stack trace는 친구다

trace가 넓은 지도라면 flame graph는 호출 비용의 단면이다. PyTorch profiler의 stack 정보를 flame graph로 바꾸면 어떤 Python path가 많은 CUDA work를 만들어내는지 보기 쉽다. 특히 framework code와 user code가 섞인 모델에서는 이것이 유용하다.

다만 flame graph도 전부는 아니다. CPU 시간이 크다고 GPU 병목이 아니고, GPU kernel 시간이 크다고 host가 문제 없는 것도 아니다. profiling view를 여러 개 열고 같은 사건을 다른 각도에서 보는 습관이 필요하다.

6. 손으로 해볼 것

작은 Transformer block이나 embedding-heavy 모델을 하나 잡고 PyTorch profiler를 켠다. 먼저 batch size 1과 batch size 4를 비교한다. GPU idle gap, kernel count, 가장 긴 kernel, memory peak를 각각 기록한다.

그 다음 의도적으로 torch.cuda.synchronize()를 한 줄 넣어본다. trace에서 host와 device timeline이 어떻게 달라지는지 확인하면 CUDA의 비동기 실행이 몸에 들어온다. 마지막으로 memory profiler를 켜고 첫 batch와 이후 batch의 allocation 차이를 본다.

Source Links

• YouTube
• Taylor Robbie LinkedIn
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_016.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_016.md

1. collective는 여러 GPU의 약속이다

NCCL은 여러 GPU 사이의 point-to-point와 collective communication을 다루는 NVIDIA library다. collective는 참여한 rank들이 같은 종류의 통신에 함께 들어오는 약속이다. all-gather는 각 rank의 조각을 모두에게 모아주고, broadcast는 한 rank의 값을 모두에게 보낸다. scatter는 반대로 rank별로 다른 조각을 나눠준다.

딥러닝에서 가장 중요한 예는 all-reduce다. 각 GPU가 자기 mini-batch로 gradient를 계산한 뒤, 모든 rank의 gradient를 reduce하고 결과를 다시 모든 rank가 갖는다. 보통 평균 gradient를 얻기 위해 sum 후 rank 수로 나눈다. 이렇게 해야 각 GPU의 model parameter가 같은 방향으로 update된다.

2. DDP는 all-reduce를 숨기지만 없애지는 않는다

PyTorch DistributedDataParallel을 쓰면 Python 코드에는 all-reduce가 직접 보이지 않는다. forward, loss, backward, optimizer step만 있는 것처럼 보인다. 하지만 profiler를 보면 backward 중간중간 NCCL all-reduce kernel이 나타난다.

이 숨김은 좋은 추상화지만, 성능을 볼 때는 투명해져야 한다. DDP는 autograd hook을 이용해 gradient가 준비되는 순간 bucket 단위 all-reduce를 시작한다. 모든 backward 계산이 끝난 뒤 한 번에 통신하는 것보다, 앞쪽 layer의 gradient를 먼저 보내는 편이 compute와 communication을 겹칠 수 있다.

3. rank, communicator, stream

NCCL API를 보면 communicator가 등장한다. communicator는 어떤 rank들이 함께 통신하는지, 각 rank가 누구인지, 어떤 장치를 쓰는지에 대한 문맥이다. 분산 코드를 읽을 때 rank는 “프로세스 번호” 정도로 시작해도 되지만, 실제로는 device, process group, communicator가 함께 움직인다.

stream도 중요하다. communication이 compute와 겹치려면 all-reduce가 compute stream을 막지 않아야 한다. 물론 dependency는 지켜야 한다. gradient가 계산되기 전에 all-reduce를 시작할 수는 없다. 그래서 이벤트와 stream dependency가 들어간다. 성능은 결국 “겹칠 수 있는 것은 겹치고, 기다려야 하는 것은 정확히 기다리는” 쪽으로 간다.

4. toy DDP로 보는 평균 gradient

Dan Johnson은 toy model로 DDP의 작동을 설명한다. rank 0과 rank 1이 서로 다른 input을 가지고 forward를 하고, 각자 local gradient를 계산한다. 그 다음 all-reduce를 통해 두 gradient의 평균을 각 rank가 받는다. 이후 optimizer step을 하면 두 model replica는 계속 같은 parameter를 유지한다.

이 예제의 장점은 DDP를 신비한 framework magic에서 끌어내린다는 것이다. 결국 일어나는 일은 “각 rank가 계산한 숫자를 모아 같은 숫자로 맞춘다”다. 큰 모델과 많은 GPU에서도 원리는 같다. 다만 bucket, overlap, topology, algorithm 선택이 성능을 좌우한다.

5. profiler에서 NCCL을 찾는 법

Nsight Systems나 PyTorch profiler를 열면 compute kernel 사이에 NCCL 관련 kernel과 memcpy/communication activity가 보인다. 여기서 볼 것은 all-reduce가 backward와 겹치는지, 긴 communication tail이 남는지, 특정 bucket이 너무 늦게 준비되는지다.

통신은 공짜가 아니다. 빠른 GEMM을 아무리 잘 써도 all-reduce가 step 끝에 길게 남으면 throughput이 막힌다. 반대로 통신이 잘 overlap되면 같은 network 비용도 덜 아프게 느껴진다. 분산 학습 성능은 kernel과 network의 합이 아니라, timeline 위 배치의 문제다.

6. 손으로 해볼 것

2 GPU가 있다면 작은 linear model을 DDP로 돌리고 profiler trace를 저장한다. 코드에는 all-reduce가 없는데 trace에는 NCCL call이 나타나는 지점을 찾는다. 그 다음 bucket size를 바꾸거나 model layer 수를 바꿔 overlap 모양이 어떻게 달라지는지 본다.

GPU가 하나뿐이라면 개념 실험만 해도 좋다. 두 rank의 gradient vector를 손으로 두 개 만들고 all-reduce sum/average 결과를 계산해본다. DDP는 이 작은 산술을 거대한 model parameter 전체에 반복해서 적용하는 장치다.

Source Links

• YouTube
• Lecture 017 repo folder
• Slides
• ddp_example.py
• ddp_simple.py
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_017.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_017.md

1. fusion이 필요한 순간

Kapil Sharma의 강의는 추천 모델과 DLRM류 workload에서 시작한다. embedding lookup, dense MLP, feature interaction, activation이 섞인 모델은 작은 연산이 많고 memory bandwidth에 민감하다. 이런 모델에서 add, mul, relu, sigmoid 같은 pointwise op를 따로 kernel로 실행하면 HBM을 여러 번 왕복한다.

예를 들어 tensor에 1을 더하는 kernel을 세 번 띄우는 대신 한 kernel에서 3을 더하면 launch도 한 번이고 memory도 덜 건드린다. 아주 단순한 예지만 fusion의 핵심이 잘 보인다. 계산이 복잡해서 빠른 것이 아니라, 불필요한 중간 저장을 없애서 빠르다.

2. torch.compile과 CUDA graph의 위치

강의에서는 torch.compile이 생성한 Triton pointwise kernel과 CUDA graph 이야기도 나온다. torch.compile은 graph를 보고 여러 pointwise op를 하나로 묶을 수 있다. CUDA graph는 반복되는 launch sequence를 capture해서 launch overhead를 줄이는 쪽에 가깝다. 둘은 같은 문제가 아니다.

Fusion은 kernel 수와 memory traffic을 줄인다. CUDA graph는 launch 경로의 CPU overhead를 줄인다. 둘 다 trace에서 “작은 kernel이 우르르 나오는” 상황을 덜 아프게 만들지만, 해결하는 층위가 다르다. 그래서 profiler에서 먼저 봐야 한다. 나는 launch overhead에 묶였는가, HBM traffic에 묶였는가?

3. generated Triton을 읽는 즐거움과 위험

Kapil은 compiler가 만든 Triton kernel을 열어본다. Triton POI처럼 pointwise operation을 나타내는 kernel이 보이고, 그 안에 add/mul/activation이 합쳐져 있다. generated code는 사람이 처음부터 쓴 코드보다 지저분할 수 있지만, compiler가 어떤 연산을 묶었는지 확인하기에는 좋다.

여기서 중요한 태도는 “compiler가 해줬으니 끝”이 아니다. compiler가 fusion하지 못하는 graph break가 있을 수 있고, dynamic shape 때문에 CUDA graph capture가 어려울 수 있으며, padding은 overhead를 만든다. 자동화는 출발점이지 면죄부가 아니다.

4. 손으로 쓰는 fusion: CUDA와 Triton

강의 자료에는 pointwise_add_relu_fused.cu, Triton fused add/mul/relu 예제, LoRA on MLP fused kernel 같은 코드가 연결되어 있다. 손으로 fused kernel을 쓰면 compiler가 놓친 패턴을 직접 최적화할 수 있다. 특히 LoRA처럼 base matmul 주변에 작은 rank update와 activation이 붙는 경우, 중간 tensor를 줄이는 효과가 크다.

하지만 손으로 쓰는 순간 책임도 온다. dtype, broadcasting, stride, alignment, boundary mask, backward pass까지 직접 생각해야 한다. inference-only path라면 부담이 덜하지만, training kernel이면 gradient까지 확인해야 한다. fusion이 성능만큼 correctness를 요구한다는 사실을 잊으면 안 된다.

5. mega-kernel의 유혹

Q&A에서 “그럼 matmul과 activation까지 다 하나로 묶으면 되나?”라는 질문이 자연스럽게 나온다. 답은 조심스럽다. 이론적으로는 가능하지만, 한 kernel이 너무 많은 일을 하면 occupancy, register pressure, shared memory 사용량, tuning space가 커진다. 하나 바꾸면 전체를 다시 튜닝해야 한다.

Persistent kernel 같은 더 큰 아이디어도 있지만, 아무 모델에나 붙이는 만능 도구는 아니다. 보통은 pointwise chain, epilogue, memory-bound 작은 연산부터 묶는 것이 안전하다. matmul 자체는 이미 cuBLAS, CUTLASS, Triton matmul이 매우 세게 최적화되어 있으니, 그 주변의 불필요한 왕복을 먼저 본다.

6. 손으로 해볼 것

PyTorch에서 y = sigmoid((x + a) * b)를 세 줄로 작성하고 profiler로 kernel 수를 본다. 그 다음 torch.compile을 켜고 kernel 수가 줄어드는지 확인한다. 마지막으로 같은 연산을 작은 Triton kernel로 직접 써서 memory read/write 횟수를 생각해본다.

응용으로는 작은 LoRA layer를 잡아보자. base output, low-rank update, bias, activation을 어디까지 합칠 수 있는지 종이에 먼저 그린다. fusion 후보를 고르는 기준은 단순하다. 중간 tensor가 크고, 연산은 가벼우며, shape가 안정적인 곳부터다.

Source Links

• YouTube
• Lecture 018 repo folder
• FusedKernels.ipynb
• Triton fused add/mul/activation
• Pointwise add ReLU CUDA
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_018.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_018.md

1. pandas 같은 API, CUDA 같은 뒤편

Devavret Makkar는 cuDF maintainer였고, 강의 시점에는 Voltron Data에서 GPU-native query engine을 만들고 있다. 그래서 강의의 관점이 흥미롭다. ML kernel만 보던 사람에게 “dataframe도 kernel들의 묶음”이라는 사실을 보여준다.

cuDF는 pandas와 비슷한 Python API를 제공하지만, backend는 CUDA다. CSV나 Parquet을 읽고, filter, group-by, join 같은 연산을 GPU에서 수행한다. TCS 같은 query engine은 SQL query를 execution plan으로 바꾸고, 그 plan을 GPU에서 실행한다. 사용자는 dataframe을 보지만, 안쪽에는 columnar memory와 hash table, scan, atomic이 있다.

2. 1 Billion Row Challenge가 좋은 장난감인 이유

강의의 toy workload는 1 Billion Row Challenge다. city와 temperature가 들어 있는 큰 text file에서 city별 min/mean/max 같은 aggregate를 구하는 문제다. CPU에서도 parsing과 aggregation이 만만치 않은데, GPU로 옮기면 병렬성이 생기는 동시에 새로운 제약도 생긴다.

GPU kernel 안에서는 마음대로 global allocation을 늘릴 수 없다. 가능한 city set이나 hash table 크기를 미리 잡아야 한다. file을 chunk로 나누고, 각 thread/block이 line boundary를 찾아 parsing하고, city를 찾아 aggregate slot에 atomic으로 더한다. 잘 짠 tailored kernel은 한 번의 pass로 많은 일을 할 수 있다.

3. column은 data buffer만이 아니다

cuDF의 column은 단순한 pointer 하나가 아니다. data buffer, null mask, type 정보, 문자열이면 offsets와 chars buffer까지 가진다. table은 이런 column들의 묶음이고, view는 소유권 없이 buffer를 가리키는 창에 가깝다.

이 구조를 이해해야 group-by가 보인다. key가 하나의 integer column이면 hash가 쉽다. 하지만 string column이나 여러 key column이면 row index를 받아 각 column의 값을 읽고, hash를 combine하고, equality comparator도 여러 column을 비교해야 한다. dataframe API 한 줄 뒤에는 꽤 복잡한 device-side accessor가 있다.

4. hash group-by와 GPU의 불편함

Hash table을 GPU에서 만들 때는 CPU처럼 “필요하면 resize”가 쉽지 않다. 보통 충분히 큰 table을 미리 만들고, sentinel value와 atomic compare-and-swap 같은 방식으로 slot을 차지한다. collision이 나면 probing한다. key가 이미 있으면 aggregate를 update하고, 없으면 새 key를 넣는다.

이 방식은 병렬성이 좋지만 memory와 contention에 민감하다. key distribution이 나쁘면 atomic이 몰리고, hash table이 너무 작으면 probing이 길어진다. 반대로 너무 크게 잡으면 memory를 낭비한다. GPU data processing은 kernel만 빠르게 하는 것이 아니라 data structure를 GPU답게 만드는 일이다.

5. query engine으로 커지면 문제가 바뀐다

단일 GPU에 file이 올라가고 한두 연산을 하는 수준을 넘으면 query engine이 필요해진다. SQL은 filter, join, group-by, sort를 조합한다. engine은 query를 plan으로 바꾸고, 어떤 operator를 어떤 순서로 실행할지, intermediate를 어디에 둘지 결정한다.

여기서 중요한 비용은 data movement다. 같은 dataset을 query마다 다시 읽으면 GPU가 빠른 이득이 줄어든다. 가능한 한 data를 GPU에 머물게 하고, intermediate table의 lifetime을 짧게 관리해야 한다. ML에서 activation memory를 보는 눈이 data analytics에서는 column buffer와 shuffle memory를 보는 눈으로 바뀐다.

6. 손으로 해볼 것

작은 CSV를 하나 만들고 pandas와 cuDF의 API가 어떻게 대응되는지 살펴본다. GPU가 없다면 개념 실험으로 충분하다. city, temperature column을 두고 group-by mean을 만들려면 어떤 buffer, hash key, aggregate value가 필요한지 적어보자.

더 깊게 보려면 string key group-by를 생각한다. key 자체를 hash table에 복사할 것인가, row index를 저장하고 comparator가 column view를 통해 비교하게 할 것인가? 이 질문 하나에 GPU memory layout의 냄새가 꽤 진하게 난다.

Source Links

• YouTube
• Devavret GitHub
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_019.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_019.md

1. inclusive와 exclusive scan

Scan은 associative operator를 입력 배열 앞에서부터 누적한 결과다. addition을 예로 들면 inclusive scan은 x[0] + ... + x[i]를 y[i]에 둔다. exclusive scan은 x[i]를 빼고 x[0] + ... + x[i-1]를 둔다. 첫 값은 identity, 덧셈이면 0이다.

이 연산은 parallel computing의 작은 스위스칼 같은 존재다. stream compaction에서 각 thread가 결과를 어디에 써야 하는지 정하고, radix sort와 histogram, parsing, graph algorithm에서도 자주 나온다. 특히 exclusive scan은 “내 앞에 몇 개가 있었나”를 알려주기 때문에 위치 계산에 강하다.

2. 큰 scan은 segment로 나눈다

전체 배열을 한 번에 모든 thread가 scan할 수는 없다. thread block 사이 synchronization이 비싸고 제한적이기 때문이다. 그래서 먼저 input을 segment로 나누고, 각 block이 자기 segment를 scan한다. 각 block은 자기 segment의 partial sum도 남긴다.

그 다음 partial sum 배열을 다시 scan한다. 마지막으로 각 block은 앞 block들의 prefix sum을 자기 segment 결과에 더한다. 이 구조는 여러 kernel launch가 필요할 수 있다. Lecture 20은 이 전체 구조 중 첫 단추, 즉 single block scan에 집중한다.

3. single block scan의 첫 구현

single block 안에서는 shared memory가 핵심이다. global memory에서 input을 읽어 shared memory buffer에 넣고, 여러 step의 scan 연산을 shared memory 안에서 수행한 뒤 결과를 global memory에 쓴다. 매 step마다 thread들이 서로의 값을 읽으므로 __syncthreads() 위치가 중요하다.

가장 직접적인 병렬 scan은 stride를 1, 2, 4, ...로 키우며 각 thread가 i - stride 값을 더하는 방식이다. 이 방식은 이해하기 쉽지만 각 step마다 같은 buffer를 읽고 쓰면 hazard가 생긴다. 한 thread가 방금 쓴 값을 다른 thread가 같은 step에서 읽어버리면 원래 의도한 이전 단계 값이 아니다.

4. double buffering과 boundary check

이 문제를 피하는 쉬운 방법은 double buffering이다. 한 buffer에서 읽고 다른 buffer에 쓰며 step마다 역할을 바꾼다. shared memory를 두 배로 쓰지만, 같은 step 안에서 읽기와 쓰기가 섞이는 위험을 줄인다. 강의에서는 이런 bug를 일부러 찾아보는 흐름도 나온다.

또 하나의 습관은 모든 memory access에 boundary check를 붙이는 것이다. stride가 커질수록 threadIdx - stride가 음수가 될 수 있고, 마지막 block에서는 global index가 n을 넘을 수 있다. GPU kernel에서 out-of-bounds는 조용히 틀린 값을 만들 수 있다. 조용한 버그가 제일 시끄럽게 시간을 잡아먹는다.

5. work efficiency를 묻기

Sequential scan은 O(n) work를 한다. 단순 병렬 scan은 step마다 많은 thread가 더하므로 O(n log n) work가 될 수 있다. parallel time은 줄었지만 총 work는 늘어난다. 이때 “work efficient한가?”라는 질문이 나온다.

이 강의에서는 최적 scan의 모든 답을 내기보다, 왜 더 좋은 알고리즘이 필요한지 감각을 만든다. shared memory를 쓰면 global memory traffic은 줄지만, synchronization과 extra work가 남는다. 다음 강의는 이 비용을 더 줄이는 방향으로 간다.

6. 손으로 해볼 것

입력 [3, 6, 7, 4, 8, 2, 1, 9]의 inclusive/exclusive scan을 손으로 계산한다. 그 다음 stride 1, 2, 4 step을 표로 그려 각 thread가 어느 값을 읽는지 표시한다. 같은 buffer를 제자리에서 update하면 어느 step에서 값이 오염되는지도 찾아보자.

CUDA로 구현한다면 block size를 작게 잡고 shared memory buffer를 출력해보는 것이 좋다. 성능보다 correctness가 먼저다. scan은 작은 index 실수 하나가 전체 뒤쪽 결과를 다 밀어버린다.

Source Links

• YouTube
• Slides
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_020.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_020.md

1. 지난 시간의 비용을 다시 본다

Lecture 20의 scan은 shared memory를 이용해 single block scan을 만들었다. 하지만 단순한 Kogge-Stone식 접근은 O(n log n) work를 할 수 있고, step마다 synchronization이 필요하다. parallelism을 얻는 대신 work와 sync 비용을 냈다.

GPU에서 sync는 공짜가 아니다. 특히 scan은 latency-bound 성격이 강해서 __syncthreads()가 많이 들어가면 성능이 쉽게 막힌다. 그래서 part 2는 같은 연산을 더 적은 work와 더 적은 sync로 하는 방법을 찾는다.

2. Brent-Kung: 덜 일하는 scan

Brent-Kung scan은 reduction tree와 down-sweep에 가까운 구조로 work를 줄인다. 모든 step에서 모든 thread가 바쁘지는 않다. 어떤 단계에서는 active thread가 줄어 idle warp가 생긴다. 대신 전체 work는 더 효율적이다.

여기서 trade-off가 보인다. Kogge-Stone은 parallel step이 많고 work도 많지만 구조가 단순하다. Brent-Kung은 work efficient하지만 일부 단계에서 병렬성이 덜 찬다. 어떤 쪽이 빠른지는 hardware, block size, sync cost에 따라 달라진다. 그래서 알고리즘 이름만 외우면 부족하고, profiler로 봐야 한다.

3. thread coarsening과 register tiling

Thread coarsening은 thread 하나가 element 하나만 처리하지 않고 여러 element를 맡게 하는 방식이다. 예를 들어 한 thread가 4개나 8개의 consecutive element를 register에 읽어 sequential scan을 한다. 그 thread는 자기 local partial sum을 block-level scan에 제공하고, 나중에 앞 thread들의 prefix를 더한다.

이 방식은 좋은 이유가 있다. thread 내부 sequential scan은 work efficient하고 sync가 필요 없다. shared memory 접근도 줄어든다. 대신 register 사용량이 늘고 occupancy가 떨어질 수 있다. occupancy는 높을수록 무조건 좋은 값이 아니다. memory latency를 숨길 만큼 충분하면 되고, 더 많은 register로 work를 줄이는 편이 나을 수 있다.

4. warp-level primitive의 힘

Warp 내부 thread들은 더 싸게 협력할 수 있다. shuffle instruction을 이용하면 shared memory와 block-wide sync 없이 값 교환이 가능하다. block 전체를 바로 scan하기보다, 각 warp가 자기 scan을 하고, warp별 partial sum을 모아 scan한 뒤 다시 각 warp에 더하는 구조를 만들 수 있다.

이렇게 하면 block-wide __syncthreads() 수가 줄어든다. 강의에서도 scan은 latency-bound이고 sync가 큰 bottleneck이라고 짚는다. warp-level scan은 이 병목을 직접 겨냥한다. GPU에서 “같은 block 안”이라는 말은 충분히 구체적이지 않다. warp 안인지, warp 사이인지가 다르다.

5. partial sums를 다시 global로 보내는 문제

여러 block으로 나뉜 scan에서는 각 block의 partial sum을 처리해야 한다. 가장 단순한 방법은 local scan kernel, partial sums scan kernel, add kernel을 따로 launch하는 것이다. 이해하기 쉽지만 global memory에 partial sums를 쓰고 다시 읽는 왕복이 생긴다.

강의 후반에는 single-pass scan 아이디어도 언급된다. block들이 순서대로 partial sum 상태를 보고 앞 block의 prefix를 이어받는 방식이다. 여기에는 block 간 synchronization을 직접 할 수 없다는 CUDA의 제약, memory ordering, status flag 같은 문제가 따라온다. Lecture 24의 decoupled look-back으로 이어지는 다리다.

6. 손으로 해볼 것

8개 element scan을 Kogge-Stone과 Brent-Kung 표로 각각 그려보자. 각 step에서 몇 개 thread가 active한지 세고, 총 add 수를 비교한다. 그 다음 한 thread가 2개 element를 맡는 coarsening 버전을 그려 local scan과 block scan이 어떻게 나뉘는지 본다.

CUDA 실습에서는 warp shuffle 기반 scan을 작은 block에서 구현해보는 것이 좋다. 같은 결과를 shared memory 버전과 비교하고, Nsight Compute에서 shared memory transaction과 synchronization 비용이 어떻게 달라지는지 확인한다.

Source Links

• YouTube
• Slides
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_021.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_021.md

1. 왜 autoregressive decoding이 답답한가

LLM decoding은 기본적으로 token을 하나씩 만든다. 다음 token을 만들려면 이전 token이 필요하다. batch가 작고 latency가 중요한 serving에서는 큰 model forward 한 번을 돌려 token 하나만 얻는 일이 많다. 이때 workload는 memory-bound가 되기 쉽다. 큰 weight를 HBM에서 읽는데, batch가 작아 compute를 충분히 채우지 못한다.

Speculative decoding은 이 병목을 다른 방향에서 건드린다. 작은 draft model이나 n-gram proposer가 여러 후보 token을 먼저 제안한다. 큰 target model은 이 후보들을 한 번에 score하고, acceptance rule에 따라 앞에서부터 몇 개를 받아들인다. 운이 좋으면 target model forward 한 번으로 token 여러 개를 얻는다.

2. vLLM의 원칙과 tension

Cade Daniel은 vLLM을 ease of use, performance, hardware support가 만나는 프로젝트로 설명한다. PagedAttention, tensor parallelism, multi-LoRA, chunked prefill, guided decoding, quantization 같은 기능들이 같은 serving engine 안에 들어간다. 이 넓이가 spec decoding을 어렵게 만든다.

vLLM은 high-throughput batch inference로 유명하지만, spec decoding은 low-latency small-batch에서도 중요하다. 두 목표는 항상 같은 방향이 아니다. throughput을 위한 scheduling과 batching 정책이 latency 최적화와 충돌할 수 있다. 그래서 “알고리즘은 간단한데 framework에 넣기는 어렵다”가 된다.

3. acceptance rate가 성능의 심장이다

Spec decoding의 이득은 accepted tokens per step에서 나온다. 후보를 많이 내도 target이 대부분 거절하면 proposal overhead만 생긴다. 반대로 acceptance rate가 높으면 target forward 횟수가 줄어 time per token이 내려간다.

여기서 draft model 선택이 중요하다. 작은 Transformer draft model은 proposal quality가 좋을 수 있지만 자체 비용이 있다. n-gram proposer는 CPU-only로 단순하게 돌릴 수 있고 특정 반복 패턴에서는 잘 맞는다. Medusa나 EAGLE 같은 방법은 target model 주변에 multi-token prediction head를 붙이는 방향이다. 좋은 proposer는 싸고, target distribution과 충분히 닮아야 한다.

4. lossless rejection sampling과 실용적 변형

Rejection sampling은 target model의 분포를 망치지 않으면서 draft proposal을 받아들이는 방법이다. “그냥 target top-1과 같으면 accept” 같은 greedy rule보다 엄밀한 장점이 있다. 다만 구현에서는 proposal probability와 target probability를 함께 다루어야 하고, token별 acceptance를 prefix 순서로 처리해야 한다.

강의는 vLLM 안에서 proposer worker, scorer, verifier가 어떻게 나뉘는지도 짚는다. spec request와 non-spec request가 같은 batch 안에 섞일 수 있고, accepted token 수가 request마다 다르다. 즉, 알고리즘의 수식보다 scheduler와 data structure가 더 복잡해지는 순간이 온다.

5. memory-bound일 때 특히 빛난다

Spec decoding은 모든 상황에서 같은 이득을 주지 않는다. 큰 batch로 이미 compute-bound에 가까우면 target forward 한 번에 여러 token을 score하는 이득이 줄 수 있다. 반대로 small batch, live inference, latency-sensitive setting에서는 weight loading 비용이 커서 token 여러 개를 한 번에 검증하는 효과가 크다.

성능을 볼 때는 target model time, draft/proposal time, verification overhead, accepted tokens per step을 분리해서 봐야 한다. speedup 하나만 보면 원인을 잃는다. acceptance rate가 낮은지, proposer가 비싼지, scheduler overhead가 큰지에 따라 고칠 곳이 다르다.

6. 손으로 해볼 것

작은 vocabulary 예제를 만들고 draft distribution과 target distribution을 직접 적어본다. draft가 token 세 개를 제안했을 때, 앞에서부터 accept/reject가 어떻게 진행되는지 손으로 따라가면 prefix acceptance의 의미가 보인다.

vLLM 코드를 읽는다면 proposer, scorer, verifier가 분리된 파일을 찾고 data가 어떤 형태로 오가는지 먼저 본다. kernel 최적화보다 먼저 request scheduling과 KV cache 상태가 어떻게 변하는지 보는 편이 좋다.

Source Links

• YouTube
• Slides
• Cade Daniel
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_022.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_022.md

1. Tensor Core는 matrix multiply의 가속기다

Vijay Thakkar와 Pradeep Ramani의 강의는 NVIDIA 내부에서 CUTLASS와 CuTe를 다루는 사람들의 시각을 보여준다. Tensor Core는 matrix multiply-add를 아주 빠르게 수행하는 hardware unit이다. ML workload에서 GEMM과 convolution이 너무 중요하니, GPU architecture도 그쪽으로 깊이 진화했다.

Volta에서는 warp-level Tensor Core instruction이 등장했고, Ampere를 거쳐 Hopper에서는 warp-group MMA가 중요해졌다. Hopper의 WGMMA는 네 개 warp가 협력해 더 큰 matrix operation을 수행하고, operand는 register뿐 아니라 shared memory descriptor를 통해 공급된다. instruction이 커질수록 프로그래머가 맞춰야 할 주변 조건도 커진다.

2. tile hierarchy가 성능의 지도다

큰 GEMM은 output matrix의 tile로 나뉜다. thread block이 큰 output tile을 맡고, warp 또는 warp-group이 그 안의 작은 tile을 맡고, MMA instruction이 더 작은 instruction tile을 계산한다. 이 계층을 잘 맞추어야 data reuse가 생긴다.

문제는 data layout이다. Tensor Core가 원하는 operand 모양과 global memory의 row/column-major layout은 그대로 맞지 않는다. shared memory에 올릴 때 swizzle하고, vectorized load를 만들고, bank conflict를 피해야 한다. CUTLASS와 CuTe는 이 복잡한 일을 tile programming model로 감싸준다.

3. CUTLASS는 “instruction wrapper”가 아니다

CUTLASS는 단순히 mma()를 부르기 쉽게 하는 얇은 wrapper가 아니다. tile scheduler, collective copy, collective MMA, epilogue, pipeline stage, shared memory layout까지 kernel 전체를 구성하는 template library다. 좋은 Tensor Core kernel은 copy와 math가 겹치고, shared memory buffer가 잘 관리되고, output tile 사이 load balancing도 맞아야 한다.

그래서 CUTLASS 코드가 복잡하다. 하지만 복잡함에는 이유가 있다. architecture마다 instruction shape, memory movement, synchronization primitive가 바뀐다. library는 그 차이를 감추는 동시에 고성능을 위한 선택지를 노출해야 한다. 쉬운 API와 peak performance 사이에는 늘 긴 줄다리기가 있다.

4. Hopper에서 바뀐 것들

Hopper는 cluster, distributed shared memory, Tensor Memory Accelerator(TMA), warp-group MMA 같은 기능을 강조한다. TMA는 global memory에서 shared memory로 tensor-shaped block을 비동기적으로 옮길 수 있고, bounds와 swizzle 같은 정보를 descriptor로 가진다. 단순 copy loop가 아니라 hardware-assisted data movement에 가깝다.

WGMMA도 중요하다. Ampere식 warp-level MMA에서는 operand가 주로 register에서 오지만, Hopper WGMMA는 shared memory operand와 descriptor가 핵심이다. TMA와 WGMMA가 모두 async 성격을 가지므로 barrier와 pipeline 관리가 성능에 직접 연결된다.

5. 왜 직접 쓰기 어려운가

Tensor Core를 잘 쓰려면 어느 dtype을 쓸지, tile shape는 어떤지, stage count는 얼마인지, shared memory layout은 어떻게 할지, copy instruction은 무엇인지, epilogue는 어떤지 모두 맞춰야 한다. 하나만 틀려도 compile은 되지만 성능이 안 나올 수 있다.

이 강의의 메시지는 “모두가 PTX를 외워야 한다”가 아니다. 오히려 CUTLASS/CuTe 같은 library가 왜 필요한지를 보여준다. peak에 가까운 kernel은 hardware detail과 programming abstraction 사이의 계약이다. 계약서를 안 읽고 서명하면 언젠가 성능 청구서가 온다.

6. 손으로 해볼 것

작은 GEMM을 대상으로 tile hierarchy를 그려보자. 전체 C matrix, CTA tile, warp tile, MMA tile을 서로 다른 색으로 표시한다. 각 단계에서 A와 B tile이 어디에 저장되는지, 재사용은 어디서 일어나는지 써본다.

그 다음 CUTLASS GEMM 예제에서 tile shape, stage count, element type을 찾아본다. 코드를 전부 이해하려 하지 말고, “이 type parameter가 memory movement를 바꾸는가, MMA shape를 바꾸는가, epilogue를 바꾸는가?”만 분류해도 큰 진전이다.

Source Links

• YouTube
• Slides
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_023.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_023.md

1. scan은 왜 다시 등장하나

Jake Hemstad와 Georgii Evtushenko는 scan을 병렬 컴퓨팅의 강력한 building block으로 본다. 순차적으로 보이는 문제를 병렬로 바꿔주는 힘이 있기 때문이다. Mamba 같은 modern architecture에서도 scan 계열 연산이 중요해지며, ML 사람에게도 더 이상 먼 알고리즘이 아니다.

Lecture 20과 21이 scan을 직접 구현하는 법을 다뤘다면, Lecture 24는 “정말 빠르게 만들려면 무엇을 기준으로 삼아야 하는가”를 묻는다. 답은 speed-of-light analysis다. benchmark 숫자 하나보다, hardware 한계와 비교한 위치를 보자는 이야기다.

2. speedup은 종종 착시다

어떤 구현이 baseline보다 200배 빠르다고 하자. 멋져 보인다. 그런데 peak memory bandwidth의 2%밖에 쓰지 못한다면 이야기가 달라진다. baseline이 너무 나빴을 뿐, 새 구현도 아직 hardware를 거의 못 쓰고 있을 수 있다.

Scan은 arithmetic intensity가 낮다. prefix sum은 element를 읽고 결과를 쓰는 memory-bound operation으로 보는 것이 자연스럽다. 최소한 input을 한 번 읽고 output을 한 번 써야 하므로 대략 2N memory operation이 lower bound가 된다. 따라서 speed-of-light는 peak bandwidth에서 이 필수 traffic을 처리하는 시간으로 잡을 수 있다.

3. hierarchy scan의 한계

전통적인 hierarchical scan은 block scan을 하고, partial sums를 저장하고, 다시 scan하고, 결과를 더한다. 이해하기 쉽지만 extra global memory traffic이 생긴다. memory operation이 2N보다 많아지므로, 아무리 잘해도 speed-of-light의 일정 비율을 넘기 어렵다.

강의에서는 이런 식으로 알고리즘을 숫자로 본다. “이 구현은 3N traffic이 필요하니 2N lower bound 기준으로 최대 66% 근처가 한계” 같은 사고다. 정확한 숫자는 조건에 따라 달라지지만, 이 계산 방식이 중요하다. 알고리즘의 구조가 성능 상한을 미리 정한다.

4. decoupled look-back의 아이디어

빠른 device-wide scan에서 중요한 알고리즘이 decoupled look-back이다. 각 block이 자기 tile scan을 하고, 앞 block의 prefix 상태를 보며 전체 prefix를 이어받는다. partial sums를 별도 kernel로 다시 scan하지 않고, 한 pass에 가깝게 처리하려는 아이디어다.

하지만 여기에는 memory ordering 문제가 있다. block이 “내 prefix 값이 준비됐다”는 flag를 쓰고, 다음 block이 그 값을 읽는다. flag와 value가 잘못된 순서로 보이면 틀린 결과가 나온다. 그래서 status flag, memory barrier, packed state 같은 세부가 중요하다. 빠른 scan은 알고리즘과 memory model이 만나는 지점이다.

5. CUB/CCCL이 왜 강한가

NVIDIA의 CUB/CCCL scan은 단순히 좋은 코드 한 파일이 아니다. block scan, warp scan, decoupled look-back, architecture별 tuning, 많은 parameter search가 쌓인 결과다. 강의에서 “scan이 speed of light에 가까운가”를 계속 묻는 이유도 여기에 있다. library가 좋은지 판단하려면 baseline 대비 speedup보다 bandwidth 효율을 봐야 한다.

또한 speed-of-light analysis는 성능 regression을 잡는 데도 유용하다. GPU가 바뀌면 peak bandwidth와 practical bandwidth가 달라진다. CI에서 여러 GPU의 bandwidth 대비 효율을 보면, 절대 시간만 보는 것보다 건강한 지표가 된다.

6. 손으로 해볼 것

먼저 scan의 최소 traffic을 계산한다. float32 N개를 scan하면 input read N개와 output write N개, 총 8N bytes가 lower bound다. 가지고 있는 GPU의 measured copy bandwidth를 기준으로 최소 시간을 추정해본다.

그 다음 직접 구현한 scan이나 library scan의 처리량을 GB/s로 바꾼다. “몇 ms”보다 “speed-of-light의 몇 퍼센트인가”를 적어보자. 이 한 줄이 performance engineering의 표정을 꽤 바꾼다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_024.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_024.md

1. CUDA Mode에서 AMD를 보는 이유

Haocong Wang의 강의는 ROCm과 Composable Kernel을 소개한다. CUDA Mode라는 이름이 붙어 있어도, 목표는 NVIDIA 하나를 외우는 것이 아니라 parallel programming의 구조를 배우는 것이다. AMD GPU에서도 Tensor Core에 해당하는 matrix unit, HBM, shared/local memory, kernel fusion 문제가 똑같이 등장한다.

ROCm은 AMD의 open compute platform이고, CK는 그 위에서 ML tensor operator를 만들기 위한 library다. GEMM, convolution, attention 같은 연산을 고성능으로 구현하려면 hardware-specific detail과 reusable abstraction이 함께 필요하다. 이 균형은 CUTLASS를 볼 때와 매우 닮았다.

2. CK의 역사: convolution에서 tensor operator로

초기 CK는 MIOpen의 convolution kernel과 깊게 연결되어 있었다. convolution이 중요했고, 기존 monolithic kernel 방식으로는 fusion과 다양한 shape 대응이 어려웠다. 시간이 지나며 AI Template 같은 framework와 함께, convolution, GEMM, attention을 더 일반적인 tensor operator로 표현하는 방향이 선명해졌다.

강의에서 중요한 문장은 convolution과 GEMM을 unified하게 본다는 흐름이다. NHWC tensor를 GEMM의 M/K 차원으로 mapping하고, tensor view를 transform해 implicit GEMM convolution을 만든다. 즉, 연산의 이름보다 coordinate transform이 더 근본적인 언어가 된다.

3. tensor view는 좌표 변환 장치다

CK의 첫 번째 기둥은 tensor view다. naive tensor view는 memory에 놓인 tensor를 shape/stride로 바라본다. transform tensor view는 dimension을 merge하거나 split하고, 어떤 축을 고정하거나, 여러 logical dimension을 새 coordinate로 바꾼다.

이 장치는 단순한 편의 기능이 아니다. convolution을 GEMM처럼 실행하려면 input tensor의 window와 channel, batch와 spatial coordinate를 GEMM M/N/K로 바꿔야 한다. 손으로 index 식을 흩뿌리면 kernel이 금방 읽기 어려워진다. tensor view는 이 mapping을 library가 이해할 수 있는 형태로 만든다.

4. tile programming과 static distributed tensor

두 번째 기둥은 tile programming이다. GPU kernel은 결국 HBM에서 tile을 읽고, register나 LDS(local data share, NVIDIA의 shared memory와 비슷한 역할)로 옮기고, 다시 coalesced하게 store한다. CK는 이 과정을 tile-level operator로 표현한다.

Static distributed tensor라는 표현도 나온다. 이는 tile 안의 data가 thread들에 어떻게 분산되어 있는지를 compile-time 정보로 다루는 쪽에 가깝다. load tile, store tile, shuffle tile 같은 연산이 reusable primitive가 된다. 좋은 kernel framework는 “빠른 kernel 하나”가 아니라 “빠른 kernel을 계속 만들 수 있는 말”을 제공한다.

5. FlashAttention과 customization

CK는 ROCm FlashAttention에도 연결된다. FlashAttention은 QK, softmax, PV를 하나의 streaming/fused kernel로 묶어 HBM traffic을 줄이는 대표 사례다. AMD에서도 같은 algorithmic idea를 hardware에 맞게 구현해야 한다. Triton 논문/구현의 아이디어가 CK/C++ kernel로 옮겨지는 장면이다.

Customization도 CK의 중요한 장점으로 소개된다. 기존 GEMM kernel에 quantization, epilogue, additional op를 붙여 새로운 fused kernel을 만들 수 있다. 기능만 되게 하는 것은 비교적 쉽지만, 최적 성능을 내려면 tile shape, memory movement, hardware unit 사용을 다시 맞춰야 한다. 이 점은 모든 vendor에서 공평하게 어렵다.

6. 손으로 해볼 것

작은 NHWC convolution을 하나 잡고, 이것이 implicit GEMM의 M/N/K로 어떻게 바뀌는지 써보자. batch와 output spatial position이 M으로, filter/channel 쪽이 K로 들어가는 식의 mapping을 직접 그리면 CK의 tensor view가 왜 필요한지 보인다.

ROCm 환경이 있다면 CK example을 바로 빌드해도 좋지만, 없어도 slide의 transform tensor view 그림을 따라가는 것만으로 충분하다. NVIDIA의 CUTLASS/CuTe와 AMD의 CK가 같은 문제를 다른 문법으로 푸는 모습을 비교해보자.

Source Links

• YouTube
• Lecture 025 repo folder
• Slides PDF
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_025.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_025.md

1. CUDA Mode가 SYCL Mode로 잠깐 바뀐 날

Patric Zhao의 강의는 Intel GPU와 SYCL programming을 소개한다. 강의 초반의 메시지는 분명하다. 인터넷에는 NVIDIA GPU를 low-level로 다루는 자료가 많지만, non-NVIDIA GPU programming 자료는 상대적으로 적다. 그래서 이 강의는 portability와 vendor diversity 측면에서 가치가 있다.

SYCL은 Khronos 계열의 C++ 기반 heterogeneous programming model이다. CPU, GPU, FPGA 같은 device를 대상으로 할 수 있고, Intel oneAPI/DPC++ 생태계에서 많이 쓰인다. CUDA가 NVIDIA GPU에 강하게 묶여 있다면, SYCL은 더 넓은 device abstraction을 목표로 한다. 물론 abstraction이 넓다고 performance detail이 사라지지는 않는다.

2. device, queue, command

SYCL 프로그램은 먼저 device를 고른다. CPU selector, GPU selector처럼 어떤 device에서 실행할지 선택하고, 그 device에 대응하는 queue를 만든다. queue는 device로 보낼 command의 통로다. kernel launch와 memory copy도 queue에 submit된다.

이 구조는 CUDA stream과 닮은 부분이 있지만 용어와 추상화가 다르다. queue 하나 안에서는 dependency가 순서대로 추적될 수 있고, 여러 queue를 쓰면 developer가 event dependency를 더 명시적으로 관리해야 한다. Patric은 SYCL graph도 언급한다. 여러 memory copy와 kernel 사이 dependency를 graph로 잡아 runtime이 실행하게 하는 방향이다.

3. memory model: buffer와 USM

Heterogeneous programming의 첫 문제는 memory다. CPU와 discrete GPU는 보통 독립 memory space를 가진다. data를 GPU에서 쓰려면 옮겨야 하고, 다시 CPU에서 읽으려면 되돌려야 한다. SYCL은 buffer/accessor 방식과 USM(Unified Shared Memory) 방식으로 이 문제를 다룬다.

Buffer 방식은 runtime에게 data access pattern을 더 많이 알려준다. 어떤 kernel이 어떤 buffer를 읽고 쓰는지 알 수 있으므로 dependency를 자동으로 잡기 쉽다. USM은 pointer 기반이라 CUDA에 익숙한 사람에게 편하지만, 명시적 copy와 lifetime 관리 책임이 더 직접적으로 온다. 편한 방식과 runtime이 최적화하기 쉬운 방식이 항상 같지는 않다.

4. 첫 kernel: parallel_for

SYCL kernel은 보통 queue.submit 안의 command group에서 parallel_for로 작성한다. range는 전체 work item 수를 정하고, item은 global index 같은 정보를 제공한다. 더 세밀하게는 nd_range를 써서 global range와 local range를 함께 지정한다. CUDA의 grid/block과 비슷한 감각을 만들 수 있다.

여기서도 질문은 같다. 어떤 부분을 parallel하게 만들 것인가? 모든 코드를 GPU에 올리는 것은 의미가 없다. computation-intensive하고 data parallel한 hot spot을 골라야 한다. CPU에서 잘하는 control-heavy 작업과 GPU에서 잘하는 throughput 작업을 나누는 눈이 필요하다.

5. Intel GPU architecture와 local memory

강의 후반은 Intel GPU architecture와 memory mode, shared local memory를 다룬다. NVIDIA의 shared memory처럼, Intel 쪽에서도 local memory를 잘 쓰면 global memory traffic을 줄이고 data reuse를 늘릴 수 있다. FlashAttention 같은 fused kernel에서도 결국 HBM에서 한 번 읽은 tile을 local/shared memory에 두고 재사용하는 사고가 반복된다.

Patric은 memory transfer time과 compute time을 분리해서 보는 관점도 보여준다. 어떤 경우에는 kernel compute보다 host-device transfer가 지배적이다. 이때는 dtype을 줄이거나, transfer와 compute를 overlap하거나, data를 device에 오래 머물게 하는 전략이 필요하다. vendor가 달라도 병목의 얼굴은 비슷하다.

6. 손으로 해볼 것

SYCL hello-world vector add를 읽고 CUDA vector add와 대응시켜보자. queue는 무엇에 해당하는가, parallel_for의 index는 CUDA의 무엇과 비슷한가, buffer/accessor가 memory dependency를 어떻게 표현하는가를 표로 적어보면 좋다.

그 다음 같은 vector add를 buffer 방식과 USM 방식으로 비교한다. 코드가 어느 쪽이 더 직접적인지, dependency 표현은 어느 쪽이 더 명시적인지 본다. SYCL을 배울 때는 syntax보다 execution model을 먼저 잡는 편이 덜 헷갈린다.

Source Links

• YouTube
• Slides
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_026.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_026.md

1. 왜 local GPU인가

Austin Huang은 강의를 H100 cluster가 아니라 laptop, phone, workstation에서 시작한다. MacBook Pro 한 대는 작아 보이지만, 매년 팔리는 수를 생각하면 세상에 깔린 local GPU compute는 무시하기 어렵다. Pixel 같은 phone도 high-end 쪽부터 점점 더 많은 GPU 연산 능력을 갖고, AR, robotics, real-time multimodal model처럼 latency와 privacy가 중요한 use case가 생긴다.

문제는 이 compute가 균질하지 않다는 점이다. NVIDIA CUDA만 생각하면 편하지만, 개인 장치 세계에는 Apple GPU, AMD, Intel, mobile GPU가 섞여 있다. 그래서 보통은 PyTorch, game engine, inference runtime 같은 큰 framework를 통해 간접적으로 GPU를 쓴다. 흔한 작업에는 괜찮지만, 낯선 algorithm을 작게 실험하려면 “runtime에 PR을 내야 하나?” 같은 이상한 벽을 만난다.

2. WebGPU와 gpu.cpp의 자리

WebGPU는 browser 전용 장난감이 아니라 GPU resource를 요청하고, buffer를 만들고, shader를 dispatch하는 API specification이다. native 쪽에서는 Dawn 같은 구현을 통해 Vulkan, DirectX, Metal backend로 내려갈 수 있다. Austin의 설명대로 “Vulkan이나 DirectX나 Metal을 돌릴 수 있으면 WebGPU도 기대해볼 수 있다”는 식의 mental model이 유용하다.

하지만 raw WebGPU도 작지 않다. 강의에서는 작은 activation kernel 하나를 launch하는 데도 WebGPU direct API는 수백 줄 setup이 필요하다고 보여준다. Vulkan은 더 길다. gpu.cpp는 여기서 compute workload에 필요한 부분만 잡는다. device code, resource allocation, hot-path dispatch. 이 세 가지만 작게 감싸면, 거대한 engine을 만들지 않고도 C++ project 안에서 GPU kernel을 날릴 수 있다.

3. CUDA와 닮았지만 CUDA는 아니다

WGSL(WebGPU Shading Language)은 gpu.cpp가 만든 문법이 아니다. WebGPU standard의 일부다. CUDA kernel처럼 thread id를 보고 index를 계산하며, workgroup_size는 CUDA block size와 비슷한 감각으로 읽을 수 있다. 그래서 CUDA를 아는 사람은 squint하면 대충 읽힌다. 다만 memory model, binding, syntax는 WebGPU 쪽의 규칙을 따른다.

Host side도 CUDA 예제와 일부 닮아 있다. context를 만들고, input/output buffer를 만들고, kernel과 buffer를 bind하고, workgroup 수를 지정해 dispatch한다. 끝나면 wait하고 host로 copy한다. 이 유사성이 중요하다. 완전히 새로운 사고를 강요하기보다, CUDA에서 익힌 “device code와 host orchestration을 나눈다”는 감각을 portable API 위로 옮긴다.

4. 작은 library가 중요한 이유

gpu.cpp의 core implementation은 강의에서 약 천 줄 정도로 소개된다. 이 크기는 단지 귀엽다는 뜻이 아니다. compile iteration이 짧고, project에 넣었을 때 빌드가 과하게 무거워지지 않는다는 뜻이다. Dawn 자체를 매번 빌드하면 몇 분이 걸릴 수 있지만, shared library로 연결하면 사용자가 컴파일하는 것은 자기 C++ snippet에 가깝다.

Austin이 보여준 demo도 이 철학과 잘 맞는다. simulation loop에서 kernel을 반복 dispatch하거나, Shadertoy처럼 x, y, time의 함수로 animation을 만들고, 실행 중에 다른 compute kernel로 바꿔 끼운다. GPU buffer state가 유지되기 때문에 kernel이 바뀌어도 animation transition이 자연스럽게 이어진다. “GPU compute를 app 내부의 live material로 다룬다”는 느낌이 있다.

5. matmul 최적화가 portability의 시험지다

강의 후반의 matmul 이야기는 gpu.cpp가 단순 wrapper인지 아닌지를 보는 시험이다. 유명한 CUDA matmul worklog처럼 naive kernel에서 시작해 shared memory tiling 같은 최적화를 더해 본다. 목표는 WebGPU/WGSL에서도 CUDA에서 배운 병목 감각, 즉 memory reuse, coalescing, tile shape 같은 것들을 표현할 수 있는지 확인하는 것이다.

물론 portable API가 vendor library를 항상 이긴다는 뜻은 아니다. Tensor Core 같은 특정 hardware feature를 얼마나 직접 쓸 수 있는지, WebGPU standard가 어떤 기능을 언제 받아들이는지는 별개의 문제다. 다만 local GPU 연구의 첫 장애물이 “800줄 setup”이면 아무도 재미있게 실험하지 못한다. gpu.cpp는 그 첫 장애물을 낮춘다.

6. 손으로 해볼 것

gpu.cpp repository의 vector add나 작은 activation example을 읽고, 같은 코드를 CUDA vector add와 나란히 놓아보자. workgroup_size, global id, buffer binding, dispatch count가 각각 CUDA의 어떤 개념과 닮았는지 표로 쓰면 빠르게 감이 온다.

그 다음 작은 matmul을 naive하게 작성하고, tile을 도입해 읽기/쓰기 횟수를 줄여 본다. 최고 성능보다 중요한 것은 같은 algorithm을 CUDA와 WebGPU에서 어떻게 다르게 표현해야 하는지 보는 것이다. portability는 추상 명사가 아니라, index 계산과 buffer binding에서 매번 드러난다.

Source Links

• YouTube
• Slides
• Code: gpu.cpp
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_027.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_027.md

1. 병목은 GPU utilization 숫자 뒤에 숨어 있다

Byron Hsu는 LLM training의 두 병목에서 출발한다. 하나는 memory다. 큰 model에서는 weight, gradient, optimizer state만으로도 GPU memory를 크게 먹고, activation은 FSDP나 ZeRO로도 자동으로 사라지지 않는다. 그래서 forward 중 activation이 쌓이고 backward 중 gradient가 쌓이는 profile을 보면, 어디서 memory spike가 나는지 직접 보인다.

다른 하나는 time이다. GPU utilization이 100%라고 해서 유용한 FLOPs가 충분히 나오는 것은 아니다. GPU가 “바쁘다”와 “잘 일한다”는 다르다. PyTorch profiler timeline을 보면 elementwise op가 줄줄이 launch되고, small kernel들이 많아 launch overhead와 memory traffic이 커지는 경우가 있다. Liger Kernel은 이 작은 낭비들을 LLM training operator 단위에서 줄인다.

2. 왜 Triton인가

CUDA로 production kernel을 쓰면 control은 많지만 개발 비용도 크다. Triton은 Python 안에서 kernel을 쓰고 PTX까지 내려가는 길을 제공한다. Byron은 Triton이 PyTorch만큼 쉽지는 않다고 선을 긋는다. 여전히 stride, physical layout, block size, numerical precision을 이해해야 한다. 그래도 CUDA보다 빠르게 kernel idea를 구현하고 팀 안에서 리뷰하기 쉽다.

Liger Kernel이 겨냥하는 지점은 바로 이 중간층이다. 연구용 prototype kernel이 아니라, training loop에 넣어도 되는 RMSNorm, fused linear cross entropy 같은 operator를 Triton으로 제공한다. 즉 “한 번 빠른 kernel을 짠다”보다 “kernel을 test하고, benchmark하고, convergence까지 확인하는 습관을 library에 넣는다”가 더 중요하다.

3. RMSNorm은 작은 예제지만 좋은 예제다

RMSNorm은 LLM에서 자주 등장하고, reduction과 elementwise scaling이 함께 있는 좋은 입문 kernel이다. 단순히 x / rms(x)를 계산한다고 끝이 아니다. 어느 dtype으로 accumulate할지, bf16 input에서 어디까지 fp32로 올릴지, backward가 reference와 맞는지 확인해야 한다.

Triton benchmark decorator를 이용하면 여러 size에서 kernel 시간을 비교하기 쉽다. 그러나 benchmark만으로는 충분하지 않다. PyTorch profiler에서 실제 training graph 안의 op들이 어떻게 줄었는지 봐야 한다. Byron이 강조하는 관점은 “성능 표 하나”가 아니라 “profile timeline에서 operation이 정말 줄었는가”다.

4. fused linear cross entropy의 memory 절약

Liger Kernel의 대표 예시는 fused linear cross entropy다. 일반적인 LM head에서는 hidden state에 linear layer를 적용해 거대한 logits를 만들고, 그 logits를 cross entropy에 넘긴다. Vocabulary가 크면 logits tensor 자체가 큰 memory spike가 된다. Training에서는 이 spike가 batch size와 sequence length 선택을 막는다.

Fusion의 아이디어는 linear와 cross entropy 사이의 큰 intermediate를 global memory에 오래 남기지 않는 것이다. 필요한 조각을 계산하고 loss/reduction에 바로 써서 memory를 줄인다. 이때 numerical detail은 더 까다로워진다. softmax normalization, target index, backward gradient가 reference와 맞아야 하고, memory를 줄인 대신 loss가 조금씩 어긋나면 production에서는 실패다.

5. contiguity는 사소하지 않다

강의에서 특히 실전적인 부분은 contiguity다. PyTorch tensor는 logical shape과 physical storage가 다를 수 있다. transpose를 했다고 data가 실제 memory에서 재배치되는 것은 아니며, stride만 바뀔 수 있다. Triton kernel은 physical view로 load/store를 하기 때문에 이 차이를 놓치면 값은 그럴듯하게 틀린다. 제일 얄미운 종류의 bug다.

그래서 custom kernel을 넣을 때는 input contract를 분명히 해야 한다. contiguous만 받는지, arbitrary stride를 받을지, 내부에서 .contiguous()를 부를지 결정해야 한다. .contiguous()는 correctness를 쉽게 만들지만, 큰 copy 때문에 성능 이득을 지워버릴 수 있다. 빠른 kernel이 느린 copy를 몰래 데리고 오면 꽤 서운하다.

6. 손으로 해볼 것

RMSNorm을 작은 Triton kernel로 구현하고 PyTorch reference와 forward/backward를 비교해보자. dtype을 bf16, fp32로 바꾸며 tolerance를 기록하면 precision 문제가 눈에 들어온다.

그 다음 profiler를 켜고 작은 training step에서 fused operator가 launch 수와 peak memory를 어떻게 바꾸는지 본다. 마지막으로 일부러 non-contiguous input을 넣어 kernel이 실패하는지, 조용히 틀리는지, 명확히 error를 내는지 확인해보자. Production kernel의 절반은 빠른 코드고, 나머지 절반은 이런 불편한 질문들이다.

Source Links

• YouTube
• Slides
• Lecture 28 README links
• Colab 1
• Colab 2
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_028.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_028.md

1. 표면은 Python, 안쪽은 compiler

Kapil Sharma의 강의는 Triton을 쓰는 사람에게 “내 kernel이 어디로 가는가?”를 보여준다. Triton은 researcher와 engineer가 Python style DSL로 GPU kernel을 빠르게 작성하게 해준다. 그래서 PyTorch code에서 @triton.jit를 붙이면 뭔가 알아서 빠른 kernel이 생기는 듯 보인다.

하지만 그 안에는 compiler pipeline이 있다. CUDA도 host code와 device code가 분리되어 nvcc, host compiler, PTX, cubin을 거친다. Triton도 마찬가지로 여러 IR lowering 단계를 거친다. 차이는 CUDA C++ source가 아니라 Triton DSL에서 시작하고, MLIR ecosystem을 적극적으로 쓴다는 점이다.

2. CUDA compilation과 Triton compilation

CUDA 쪽에서는 nvcc toolchain이 host C++ code를 일반 compiler로 보내고, device code는 PTX 또는 cubin 쪽으로 내린다. PTX는 NVIDIA의 virtual ISA이고, 최종 실행은 architecture-specific SASS로 이어진다. Compiler Explorer 같은 도구로 vector add kernel의 PTX/SASS를 볼 수 있다.

Triton은 Python DSL을 common Triton IR로 바꾸고, target에 맞게 Triton GPU IR, NVIDIA GPU 관련 IR, LLVM IR, PTX, cubin으로 내려간다. 이 흐름은 “Triton이 portable할 수 있는 이유”와 “portable하지만 backend마다 다른 최적화가 필요한 이유”를 동시에 설명한다. Common IR이 있다고 해서 모든 hardware detail이 사라지는 것은 아니다.

3. Triton cache는 공부 도구다

강의에서 가장 실용적인 장면은 cache directory를 직접 여는 부분이다. Triton은 기본적으로 home 아래 .triton/cache에 compile artifact를 둔다. Cache를 지우고 vector add kernel을 compile하면 TTIR, TTGIR, PTX, cubin, launcher 관련 shared library 같은 파일이 생긴다.

이 파일들을 보면 program_id, pointer arithmetic, masked load/store가 어떤 IR로 표현되는지 보인다. 성능 문제가 생겼을 때 “Triton이 이상해요”에서 멈추지 않고, 어느 단계에서 layout conversion이 생겼는지, load/store가 예상과 다른지, PTX가 어떤 모양인지 확인할 수 있다. 좋은 debugging은 추측의 양을 줄이는 일이다.

4. MLIR와 dialect라는 언어

MLIR의 ML은 machine learning이 아니라 Multi-Level Intermediate Representation이다. Compiler를 만들기 위한 toolkit에 가깝다. Triton은 MLIR 위에서 Triton 전용 dialect, Triton GPU dialect, NVIDIA GPU 쪽 dialect를 정의하고, 여러 pass를 통해 IR을 바꾼다.

Kapil은 constant propagation, dead code elimination, inlining, canonicalization, common subexpression elimination 같은 pass도 언급한다. 일부는 MLIR 생태계의 standard pass이고, Triton은 Python binding으로 이를 사용한다. 여기서 얻을 교훈은 단순하다. Triton kernel을 잘 쓰려면 syntax만 보지 말고, compiler가 볼 수 있는 형태로 정보를 주어야 한다.

5. target이 늘어날수록 질문도 늘어난다

강의 중 질문은 Metal, HIP, Triton CPU 같은 target으로 이어진다. 어느 IR layer에 다른 DSL을 꽂을 수 있는가? Triton IR을 common layer로 삼을 수 있는가? CPU backend에서는 x86 assembly로 내려갈 수 있는가? 이 질문들은 모두 같은 방향을 본다. GPU programming은 이제 단일 vendor API만의 문제가 아니라, compiler interface의 문제가 되었다.

그렇다고 “한 IR이면 모든 것이 해결된다”는 식으로 읽으면 곤란하다. IR이 높을수록 만들기 쉽지만 hardware-specific scheduling 정보는 덜 들어 있다. IR이 낮을수록 hardware에 가까우나 사람이 쓰기 어렵다. Triton의 가치는 그 사이에서 꽤 좋은 타협점을 제공한다는 데 있다.

6. 손으로 해볼 것

Triton vector add 예제를 하나 만들고, cache directory를 비운 뒤 실행해보자. 생성된 TTIR, TTGIR, PTX를 열어 program_id, mask, load, store가 어디에 나타나는지 표시한다. 이 작업은 한 번만 해도 Triton을 보는 눈이 달라진다.

다음에는 block size를 바꾸고 artifact가 어떻게 바뀌는지 비교한다. 시간이 있으면 TRITON_CACHE_DIR를 별도 위치로 지정해 실험별 artifact를 보존하자. Kernel tuning notebook 옆에 IR diff를 남기면, 나중에 자기 자신에게 꽤 친절한 사람이 된다.

Source Links

• YouTube
• Lecture 029 repo folder
• Slides PDF
• vector_add.py
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_029.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_029.md

1. Training에는 저장할 것이 너무 많다

Thien Tran은 post-training quantization을 짧게 복습한 뒤 바로 training으로 넘어간다. Inference에서는 보통 weight-only quantization이나 activation quantization을 생각한다. Weight를 int4/int8로 저장하고 matmul 안에서 dequantize하거나, activation scale을 dynamic하게 계산해 더 낮은 precision Tensor Core를 쓰는 식이다.

Training은 더 복잡하다. Forward pass에서는 backward를 위해 activation을 저장하고, backward에서는 gradient가 생기며, optimizer step에는 Adam 같은 optimizer state가 필요하다. Adam은 parameter에 비례해 두 개의 state를 갖기 때문에, model parameter보다 optimizer state가 더 큰 부담이 되기도 한다. Memory profile을 보면 activation이 쌓이고 사라지는 모양, gradient가 생겼다가 optimizer step 뒤 줄어드는 모양이 보인다.

2. 목표는 memory와 speed, 둘 다다

Quantized training의 첫 목표는 memory footprint를 줄이는 것이다. Activation, model parameters, optimizer state가 모두 후보가 된다. 특히 optimizer state quantization은 이해하기 쉽다. Adam state가 parameter의 두 배 크기라면, 여기를 8bit나 4bit로 줄이는 것만으로도 큰 memory 절약이 가능하다.

두 번째 목표는 speed다. 낮은 precision Tensor Core, 예를 들어 int8이나 FP8 경로를 사용할 수 있으면 throughput이 오른다. 하지만 speed는 공짜가 아니다. Scale 계산, quantize/dequantize, rounding, reduction이 추가된다. 이 overhead가 줄어든 memory traffic보다 커지면, bit를 낮췄는데 wall time은 별로 좋아지지 않는 상황이 생긴다.

3. Low-bit optimizer kernel의 핵심

Optimizer kernel을 생각해보자. 일반적으로 parameter, gradient, optimizer state를 읽고, Adam이나 SGD update를 계산한 뒤 parameter와 state를 다시 쓴다. Low-bit optimizer에서는 저장된 state를 dequantize하고 update를 계산한 뒤 다시 quantize해야 한다.

강의에서 중요한 포인트는 dequantized optimizer state를 global memory에 쓰지 않는 것이다. 낮은 bit로 저장한 state를 읽어 kernel 안에서 fp32 같은 계산 형태로 풀고, update를 한 뒤 다시 낮은 bit로 압축해 저장한다. 중간의 큰 fp32 state가 global memory에 나타나지 않아야 memory 절약이 살아 있다. 가능하면 shared memory 안에서 scale 계산과 quantization을 처리한다.

4. Group size와 scale은 사소한 knob이 아니다

Quantization에는 scale이 필요하다. Tensor-wise scaling은 단순하지만 range가 큰 tensor에서 error가 커질 수 있다. Row-wise나 block-wise scaling은 품질을 개선할 수 있지만, scale을 계산하기 위한 reduction과 metadata가 필요하다. Group size가 커지면 scale 수는 줄지만 quantization error가 커질 수 있고, group size가 작아지면 scale 계산과 저장 비용이 올라간다.

Optimizer state에서는 이 tradeoff가 더 예민하다. Update가 반복되기 때문에 작은 numerical bias가 누적될 수 있다. Stochastic rounding 같은 기법이 등장하는 이유도 여기에 있다. 이 강의는 algorithm paper를 깊게 파기보다, 구현자가 kernel 안에서 어떤 state를 어느 precision으로 들고 있어야 하는지 보여주는 데 초점을 둔다.

5. Tensor Core를 쓰려면 data가 그 모양이어야 한다

낮은 precision은 성능상 매력적이지만, hardware가 좋아하는 layout과 dtype으로 들어가야 한다. int8, FP8 Tensor Core를 쓰려면 matmul input이 올바른 packing, alignment, scale 처리와 맞아야 한다. “dtype을 바꿨으니 빠르겠지”는 GPU 앞에서 자주 깨지는 순진한 기대다.

Triton이나 CUTLASS를 사용할 때도 마찬가지다. Quantized path는 memory bandwidth를 줄이는 대신 dequantization과 scale load가 생긴다. Compute-bound인지 memory-bound인지 먼저 봐야 한다. Weight를 더 작게 만들어도 kernel이 scale 처리에 묶이면 이득은 줄어든다.

6. 손으로 해볼 것

작은 Adam optimizer를 PyTorch로 작성하고, optimizer state를 fp32에서 int8 block-wise 저장으로 바꿔보자. 처음에는 빠른 kernel보다 state size, scale 계산, dequantize-update-quantize 순서를 정확히 적는 것이 좋다.

그 다음 tiny model training을 몇 step만 돌려 loss curve와 parameter difference를 비교한다. Quantized training은 “한 step의 max error”만 봐서는 부족하다. 반복되는 update에서 error가 어떤 방향으로 쌓이는지 봐야 한다. 아주 작은 실험이어도 이 감각은 꽤 오래 간다.

Source Links

• YouTube
• Lecture 030 repo folder
• Slides PDF
• Notebook
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_030.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_030.md

1. Apple GPU는 그냥 옆에 있는 hardware가 아니다

Nikita Shulga는 자신의 laptop을 예로 든다. CPU peak와 GPU peak를 napkin math로 비교하면, laptop 안의 Apple GPU도 충분히 무시하기 어려운 compute를 갖고 있다. 문제는 CUDA capable GPU가 아니라는 점이다. NVIDIA 쪽 지식만으로는 이 compute를 직접 다루기 어렵다.

PyTorch에는 MPS backend가 있다. MPS는 Metal Performance Shaders이고, 결국 professionally written Metal kernel들의 묶음에 가깝다. 다만 모든 operator가 항상 준비되어 있지는 않다. 새 op가 필요하거나 backend coverage가 비어 있으면, 직접 Metal kernel을 알아야 한다. 이 강의는 그 첫 발판이다.

2. Metal의 큰 단어들

Metal shader는 GPU에서 실행되는 함수다. CUDA kernel과 역할이 비슷하다. Host side에서는 device, command queue, command buffer, encoder, pipeline state 같은 객체가 등장한다. 이름이 많아 겁이 나지만, 실행 흐름은 익숙하다. GPU가 실행할 code를 준비하고, buffer를 묶고, thread grid를 정하고, command buffer에 넣어 실행한다.

Nikita는 CUDA를 알면 Metal의 상당 부분을 이미 안다고 말한다. 맞는 말이다. Parallel index, memory hierarchy, threadgroup, synchronization 같은 질문은 GPU 언어마다 반복된다. 다른 점은 API의 ceremony와 Apple GPU/MPS/PyTorch가 만나는 접점이다.

3. PyTorch에서 Metal kernel을 넣는 길

강의에는 PyTorch extension 관점도 나온다. Metal shader를 PyTorch operator처럼 쓰려면 ATen/MPS utility와 연결해야 한다. Nikita는 MPS operations utility 쪽의 Metal shader library class를 사용해 kernel source를 JIT compile하고 invocation framework를 쓰는 방향을 언급한다.

여기서 중요한 것은 “Metal을 따로 배우기”가 아니라 “PyTorch op의 missing piece를 Apple GPU에서 채우기”다. Input tensor가 MPS device에 있고, Metal buffer로 접근되며, command buffer에 dispatch되어 결과 tensor가 채워진다. CUDA extension을 쓸 때와 마찬가지로 framework boundary가 성능과 correctness의 큰 부분을 차지한다.

4. GEMM을 쓰면 진짜 공부가 시작된다

Vector add는 API를 익히기 좋지만, GPU를 이해하게 만드는 것은 GEMM이다. Lecture material에는 gemm_perf_studies.mm가 있다. GEMM은 global memory load, threadgroup memory reuse, register accumulation, tile shape가 모두 드러난다. CUDA에서 shared memory tiling을 배웠다면, Metal에서는 threadgroup memory로 비슷한 사고를 해볼 수 있다.

다만 Apple GPU의 execution unit, memory behavior, available matrix instructions는 NVIDIA와 같지 않다. CUDA의 block/thread 이름을 그대로 번역하면 절반만 맞다. 좋은 연습은 “CUDA의 무엇과 닮았나”와 “여기서는 무엇이 다르게 드러나나”를 분리해서 보는 것이다.

5. MPS와 custom Metal 사이

가능하면 MPS가 제공하는 빠른 kernel을 쓰는 것이 좋다. 이미 잘 최적화되어 있고, PyTorch와의 integration도 되어 있다. 하지만 MPS가 제공하지 않는 operator나 새 실험을 해야 한다면 custom Metal kernel이 필요하다. 이때 목표는 vendor library를 무조건 이기는 것이 아니라, 없는 path를 정확하게 채우는 것이다.

Custom kernel은 lifecycle과 synchronization도 신경 써야 한다. Command buffer를 언제 commit하고 wait할지, tensor memory가 언제 살아 있는지, CPU와 GPU가 같은 buffer를 언제 볼 수 있는지 명확해야 한다. GPU programming에서 “계산식은 맞는데 실행 순서가 틀림”은 꽤 흔한 고전이다.

6. 손으로 해볼 것

Lecture 031의 gemm_perf_studies.mm를 열고, host side setup과 shader code를 색으로 나누어보자. Device, pipeline, buffer, dispatch가 어디서 만들어지는지 표시하면 Metal API의 뼈대가 보인다.

그 다음 작은 elementwise op를 PyTorch MPS tensor에 연결하는 흐름을 따라가 보자. CUDA extension을 만든 경험이 있다면, CUDA stream, kernel launch, tensor pointer에 해당하는 Metal/MPS 쪽 단어를 짝지어 적어보면 좋다.

Source Links

• YouTube
• Lecture 031 repo folder
• Slides PDF
• gemm_perf_studies.mm
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_031.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_031.md

1. 최적화 library가 마주친 진짜 일

Daniel Han은 Unsloth가 처음에는 fine-tuning을 빠르게 하고 memory를 줄이는 optimization library였다고 설명한다. 그런데 실제 사용자가 붙으면서 일이 커졌다. Gemma의 BOS token, RoPE, LayerNorm, activation function 같은 model-specific bug를 만나고, model analysis와 compatibility가 중요한 일이 되었다.

이 점이 강의의 매력이다. GPU kernel 강의처럼 시작하지만, production training library의 생활감이 강하다. Kernel이 빨라도 tokenizer가 틀리면 model quality가 무너진다. Cross entropy가 memory를 덜 써도 gradient accumulation denominator가 틀리면 loss가 다른 의미가 된다.

2. RMSNorm과 dtype discipline

Unsloth의 초기 kernel 중 하나는 RMSNorm이다. RMSNorm은 작아 보이지만 dtype discipline이 필요하다. 어떤 값을 fp32로 upcast할지, 언제 bf16/fp16으로 내릴지, backward에서 gradient가 reference와 맞는지 확인해야 한다. Daniel은 kernel code에 upcast/downcast 관련 주석이 많다는 점을 보여준다. 이 주석들은 장식이 아니라 correctness의 흔적이다.

Triton은 빠르게 kernel을 만들기 좋지만, compiler가 모든 numerical choice를 대신 책임지지는 않는다. 특히 LLM fine-tuning에서는 작은 오차가 반복 update를 거쳐 크게 보일 수 있다. “동작한다”와 “수렴한다” 사이에는 꽤 넓은 강이 있다.

3. Cross entropy는 memory 문제이면서 수학 문제다

LM training에서 cross entropy는 vocabulary dimension 때문에 큰 memory를 쓴다. LM head에서 logits를 만들고 loss를 계산하는 과정은 memory spike의 주범이 될 수 있다. Unsloth는 cross entropy 쪽에서도 kernel을 손보고, 필요한 upcast를 kernel 내부로 옮겨 memory usage를 줄이는 방향을 다룬다.

하지만 더 중요한 사건은 gradient accumulation의 denominator 문제다. 여러 microbatch를 accumulate할 때 loss를 어떤 denominator로 normalize하는지 틀리면, 겉보기에는 training이 돌아가도 gradient scale이 달라진다. Daniel은 이 문제를 강하게 강조한다. GPU 최적화는 수식을 바꾸지 않는다는 약속 위에서만 의미가 있다.

4. FlashAttention과 dependency의 현실

Unsloth는 FlashAttention도 활용한다. Attention은 long context fine-tuning에서 memory와 speed를 좌우한다. 다만 external dependency는 늘 현실적인 문제를 데려온다. 특정 hardware, 특정 dtype, 특정 build 환경에서 잘 맞아야 하고, package user가 설치에서 막히면 library 경험이 나빠진다.

그래서 Unsloth의 engineering은 “직접 Triton으로 다 쓰기”와 “검증된 external kernel을 잘 쓰기” 사이의 균형이다. 모든 kernel을 직접 쓰는 것은 control을 주지만 maintenance cost도 준다. 좋은 library는 여기서 너무 용감하지도, 너무 의존적이지도 않아야 한다.

5. Quantization은 싸게 돌리기 위한 현실적 도구다

강의에는 quantized model과 LoRA/fine-tuning 이야기도 나온다. 70B 같은 큰 model은 그대로 올리기 어렵고, quantization은 VRAM을 줄여 더 싼 장비에서 돌릴 수 있게 한다. Weight-only quantization, activation quantization, KV cache quantization은 각각 memory-bound와 compute-bound 상황에서 다른 의미를 갖는다.

Unsloth가 다루는 사용자 경험은 여기서 중요해진다. Pre-quantized model을 제공하고, LoRA training path를 맞추며, accuracy를 잃지 않는 범위를 찾아야 한다. Kernel이 빠른 것만으로는 부족하다. 사용자가 “그냥 fine-tune했는데 잘 됐다”고 느끼려면, 수많은 작은 default가 맞아야 한다.

6. 손으로 해볼 것

작은 language model에서 gradient accumulation을 직접 구현하고, denominator를 token count 기준과 batch count 기준으로 바꿔 loss scale을 비교해보자. 틀린 normalize가 얼마나 조용히 training을 바꾸는지 체감할 수 있다.

그 다음 RMSNorm reference와 Triton version을 만들고, forward뿐 아니라 backward gradient까지 비교한다. 마지막으로 nontrivial tokenizer setting, 예를 들어 BOS/EOS handling을 바꿨을 때 학습 데이터가 어떻게 달라지는지 확인하자. GPU 성능 작업은 생각보다 자주 tokenizer 옆자리에 앉아 있다.

Source Links

• YouTube
• Slides
• Daniel Han
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_032.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_032.md

1. Quantization의 목적이 바뀌었다

Wang Lei는 mixed precision computing의 배경에서 시작한다. 예전에는 lower precision이 hardware low precision instruction, 즉 Tensor Core 같은 unit을 더 잘 쓰기 위한 도구로 많이 이해되었다. 최근 LLM에서는 model이 너무 커져서 GPU memory에 올리는 것 자체가 문제가 되었다. Weight-only quantization은 memory를 줄이는 동시에, memory-bound inference에서 speed도 줄 수 있다.

하지만 낮은 bit weight를 쓰면 계산이 자동으로 쉬워지는 것은 아니다. Weight는 int4, int2처럼 packed되어 있고, activation은 보통 fp16/bf16일 수 있다. GEMM 전에 dequantize를 해야 하며, scale과 zero point도 고려해야 한다. 이 모든 것을 대충 처리하면 memory를 아낀 만큼 layout 변환과 unpacking에서 시간을 잃는다.

2. BitBLAS와 Ladder의 관점

BitBLAS는 kernel library이고, Ladder는 end-to-end compiler로 소개된다. 목표는 quantized model을 실제 hardware instruction에 맞는 fast kernel로 이어주는 것이다. TVM이나 Triton 같은 system이 compute와 schedule을 분리하는 것처럼, BitBLAS도 “무엇을 계산하는가”와 “어떤 layout과 tile로 실행하는가”를 나눠 생각한다.

강의에서 흥미로운 지점은 “Triton-like language”인 Tile Language(TL) 언급이다. Python DSL로 kernel을 쓰는 편의는 가져오되, low-bit GEMM에 필요한 layout와 schedule abstraction을 더 직접적으로 다루려는 시도다. Low-bit workload에서는 일반 matmul보다 layout이 더 앞에 나온다.

3. Layout은 성능의 절반이다

NVIDIA Tensor Core는 peak throughput을 내려면 data가 특정 모양으로 들어오기를 기대한다. BitBLAS 강의에서 반복되는 주제는 data layout이다. Layout이 맞지 않으면 memory access가 나빠지거나, Tensor Core를 충분히 먹이지 못하거나, kernel 사이에 변환이 추가된다.

Layout transform을 별도 kernel로 자주 수행하면 memory traffic이 다시 커진다. 그래서 layout을 operator graph 위로 전파하는 생각이 나온다. Elementwise op나 memory-intensive 작은 op에서는 layout을 유지하고, compute-intensive op끼리는 fusion하거나 shared memory 안에서 data를 재사용한다. 좋은 compiler는 “이 layout을 지금 바꿔야 하나?”를 계속 묻는다.

4. Dequantization은 어디에서 할 것인가

Weight-only quantization에서는 matmul 중 weight를 dequantize해야 한다. 이때 global memory에서 unpack한 뒤 fp16 weight를 써버리면 memory 절약의 의미가 약해진다. 더 나은 방향은 packed weight를 읽고, shared memory나 register 가까운 곳에서 필요한 tile만 풀어 MMA에 넣는 것이다.

하지만 이것도 쉽지는 않다. Packed bit를 꺼내는 instruction, scale load, zero point 처리, tile shape가 모두 얽힌다. 작은 shape에서는 overhead가 크게 보이고, 큰 compute-bound GEMM에서는 dequantization이 상대적으로 묻힐 수 있다. 그래서 workload shape별 benchmark가 필요하다. Low-bit kernel은 “한 가지 빠른 kernel”로 끝나기 어렵다.

5. Correctness와 performance를 같이 잡기

BitBLAS 같은 system은 correctness도 까다롭다. Quantized path에서는 reference fp16/fp32 output과 bit-exact로 맞지 않는 경우가 자연스럽다. 그래서 tolerance와 error metric을 정하고, layer별 performance와 quality impact를 같이 봐야 한다.

강의의 compiler 관점은 여기서 도움이 된다. Compute definition은 reference semantics를 잡고, schedule은 performance를 바꾼다. 이 둘을 분리하면, optimization을 바꿀 때 “수학을 바꿨는지, 실행 방법만 바꿨는지”를 더 분명히 추적할 수 있다.

6. 손으로 해볼 것

작은 int4 weight-only linear layer를 만들어보자. Weight를 pack하고, scale을 group별로 저장한 뒤, matmul 안에서 dequantize하는 naive implementation을 작성한다. 그 다음 dequantized weight를 미리 풀어놓는 baseline과 memory traffic을 비교한다.

가능하면 BitBLAS tutorial이나 lecture repo를 열어 layout transform이 어디서 일어나는지 찾아보자. Kernel code보다 먼저 tensor가 어떤 physical layout으로 흐르는지 그려보면, low-bit GEMM이 왜 compiler problem인지 선명해진다.

Source Links

• YouTube
• Lecture 033 repo folder
• Slides PDF
• Wang Lei GitHub
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_033.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_033.md

1. Quantization은 어디에 쓰는가

Hicham Badri는 quantization을 넓게 잡고 시작한다. Weight quantization은 큰 model을 더 적은 VRAM에 올리고, memory-bound inference를 빠르게 할 수 있다. Activation quantization은 많은 user를 serve하는 compute-bound 상황에서 Tensor Core path를 더 잘 쓰게 할 수 있다. KV cache quantization은 long prompt나 long context에서 memory를 줄이는 데 중요하다.

즉 “quantization”이라는 단어 하나로 여러 병목을 덮어버리면 안 된다. Weight가 병목인지, activation bandwidth가 병목인지, KV cache capacity가 병목인지에 따라 kernel 모양이 달라진다. 같은 int4라도 inference prefill, decode, fine-tuning에서 의미가 다르다.

2. Linear quantization부터 붙잡자

강의는 linear quantization과 nonlinear quantization을 구분한다. Linear quantization은 scale과 zero point 또는 symmetric scale을 이용해 float value를 integer range로 mapping한다. Nonlinear quantization이나 vector quantization도 있지만, kernel 구현을 설명할 때는 linear quantization이 가장 손에 잡힌다.

핵심은 group이다. Tensor 전체에 scale 하나를 쓰면 단순하지만 정확도가 떨어질 수 있다. Group-wise quantization은 더 정확하지만, group별 scale을 읽고 적용해야 한다. Low-bit kernel에서 scale load는 생각보다 크게 보인다. Weight data는 4배 줄였는데 scale과 unpack이 발목을 잡는 장면이 흔하다.

3. Dequantize를 matmul과 붙이기

가장 naive한 방법은 quantized weight를 미리 dequantize해 fp16 weight를 만들고, 일반 matmul을 부르는 것이다. 쉽지만 memory 절약이 사라진다. 더 좋은 방법은 matmul kernel 안에서 packed weight를 읽고, 필요한 tile만 dequantize해 dot product에 넣는 것이다.

여기서 Triton이 매력적이다. Python 가까운 syntax로 custom matmul variant를 빠르게 실험할 수 있고, PyTorch와 integration도 좋다. Hicham은 Triton을 맹목적으로 배우기보다 CUDA concepts, warps, memory hierarchy를 같이 이해해야 한다고 말한다. Triton이 쉽다는 말은 GPU가 쉬워졌다는 말이 아니다. 단지 실험의 마찰이 줄었다는 말에 가깝다.

4. Packing, repacking, contiguity

Low-bit kernel은 data representation이 까다롭다. int4 weight 두 개를 byte 하나에 pack할 수 있고, int2는 더 복잡해진다. Hardware가 좋아하는 tile layout과 model checkpoint가 저장한 layout이 다르면 repack이 필요하다. Repack 비용은 한 번만 내면 괜찮을 수 있지만, hot path에 들어오면 곧바로 성능을 잡아먹는다.

강의에서 나오는 contiguity 문제도 중요하다. Transpose된 tensor는 logical shape이 맞아도 physical memory가 contiguous하지 않을 수 있다. .contiguous()를 호출하면 correctness는 쉬워지지만 copy 비용이 생긴다. Low-bit path에서는 packing까지 얽히므로, tensor stride와 packing layout을 kernel contract에 명확히 적어야 한다.

5. Triton의 장점과 한계

Triton은 빠른 iteration과 readability가 강점이다. TVM 기반 solution은 compiler build와 customization이 무겁고, CUDA는 개발 속도가 느릴 수 있다. Triton은 아이디어가 괜찮은지 빨리 볼 수 있다. PyTorch integration도 좋고, torch.compile ecosystem과 가까운 것도 장점이다.

하지만 performance ceiling이 항상 최고라는 뜻은 아니다. 특히 Tensor Core를 매우 세밀하게 먹이는 GEMM, TMA, architecture-specific path에서는 CUDA/CUTLASS 계열이 더 강할 수 있다. Low-bit Triton kernel은 “충분히 빠르고, 고치기 쉽고, 배포하기 쉬운” 지점을 노린다. 이 균형이 실제 library에서는 자주 이긴다.

6. 손으로 해볼 것

작은 linear layer에서 int4 weight-only path를 만들어보자. 먼저 dequantize 후 PyTorch matmul을 부르고, 다음에는 Triton kernel 안에서 dequantize하는 version을 만든다. 두 version의 memory allocation과 runtime을 비교하면 fusion의 의미가 보인다.

그 다음 일부러 transpose된 non-contiguous input을 넣어보자. Kernel이 stride를 제대로 처리하는지, .contiguous()를 호출해 성능이 얼마나 바뀌는지 확인한다. Low-bit kernel debugging은 bit packing보다 stride에서 먼저 미끄러지는 경우가 많다.

Source Links

• YouTube
• Slides
• Lecture 34 README
• Hicham Badri GitHub
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_034.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_034.md

1. SGLang은 front end와 backend를 함께 본다

SGLang은 programming interface와 inference backend를 함께 제공하는 LLM serving system이다. Yineng Zhang은 SGLang이 가볍고 customizable한 open-source inference engine으로 쓰인다고 소개한다. 강의의 초점은 세 가지다. CPU overhead hiding, FlashInfer integration, 그리고 TurboMind GEMM optimization 계획이다.

Serving system에서는 모델 kernel만 빠르면 끝나지 않는다. Request가 계속 들어오고, 각 request는 prompt 길이와 generation 상태가 다르다. Scheduler는 continuous batching으로 이들을 섞어 GPU에 보낸다. 이 과정의 CPU logic이 blocking이면 GPU는 빈칸을 만든다.

2. CPU overhead는 idle GPU로 나타난다

강의는 다른 engine들의 CPU overhead 분석을 인용하며, 일부 경우 CPU scheduling 때문에 큰 시간이 낭비될 수 있다고 말한다. Scheduler는 input message를 받고, 이전 step 결과를 처리하고, stop condition을 확인하고, prefix matching과 request reorder를 수행하고, 다음 batch를 위한 memory를 할당한다. 이 일은 대부분 CPU work다.

Blocking scheduler에서는 CPU가 다음 batch를 준비하는 동안 GPU가 기다린다. 이상적인 구조는 GPU model worker가 현재 batch를 실행하는 동안 CPU scheduler가 다음 batch를 준비하는 것이다. 이렇게 pipeline을 만들면 CPU scheduling overhead를 숨길 수 있다. Serving에서 overlap은 optional optimization이 아니라 생존 기술에 가깝다.

3. Prefix cache는 scheduling 문제다

SGLang은 prefix caching을 early feature로 강조한다. LLM serving에서는 여러 request가 같은 system prompt나 shared prefix를 가질 수 있다. 이 prefix의 KV cache를 재사용하면 prefill work를 줄일 수 있다. 하지만 cache를 잘 쓰려면 request matching, eviction, memory allocation, batch construction이 함께 움직여야 한다.

강의에서는 radix tree 같은 prefix matching과 request reorder 이야기가 나온다. Cache hit을 많이 얻으려면 비슷한 prefix를 가진 request를 잘 묶어야 하지만, latency와 fairness도 고려해야 한다. Serving scheduler는 언제나 “GPU 효율”과 “사용자 지연” 사이에서 줄타기를 한다.

4. FlashInfer와 kernel integration

FlashInfer는 LLM inference를 위한 high-performance GPU kernel library로 소개된다. Attention, sampling, normalization, activation kernel 같은 serving hot path를 제공한다. SGLang은 FlashInfer team과 가까이 협업하며 새로운 optimization을 integration하고 benchmark한다고 설명한다.

여기서 중요한 점은 kernel library를 붙이는 일도 system work라는 것이다. 좋은 kernel을 호출하려면 input layout, batch shape, KV cache format, sampling path가 맞아야 한다. Kernel 하나가 빠른데 scheduler가 그 shape을 만들지 못하면 이득은 제한된다. Inference engine은 kernel과 batch planner가 함께 최적화되는 물건이다.

5. GEMM과 SM utilization

강의 후반에는 TurboMind GEMM optimization과 Stream-K에서 영감을 받은 방식도 언급된다. 목표는 workload를 나누어 SM idle time을 줄이는 것이다. Decode phase의 shape은 batch와 sequence 상태에 따라 작고 불규칙해질 수 있다. 이때 일반 GEMM이 모든 SM을 잘 채우지 못하면 tail effect가 생긴다.

Serving에서 GEMM은 training의 큰 dense GEMM과 다르게 보일 때가 많다. Request마다 길이가 다르고 batch가 계속 변한다. 그래서 kernel 성능표 하나보다 workload distribution을 보는 것이 중요하다. SGLang 강의는 “engine-level benchmark”가 왜 필요한지 잘 보여준다.

6. 손으로 해볼 것

작은 toy serving loop를 만들어보자. CPU가 request를 받고 batch를 구성한 뒤 GPU 작업을 sleep이나 dummy kernel로 흉내낸다. Blocking scheduler와 next-batch preparation을 overlap한 scheduler를 비교하면 idle gap이 눈에 보인다.

다음에는 prefix가 같은 request들을 일부러 섞어 넣고, cache hit을 높이기 위해 reorder하면 latency가 어떻게 바뀌는지 기록해보자. Serving optimization은 kernel보다 queue를 먼저 그려야 할 때가 많다.

Source Links

• YouTube
• Lecture 035 repo folder
• Slides PDF
• Speaker profile
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_035.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_035.md

1. Attention의 병목은 큰 score matrix다

Jay Shah는 attention을 O = softmax(QK^T)V에서 차근차근 시작한다. Query와 Key를 곱하면 sequence length에 대해 큰 score matrix가 생긴다. Naive implementation은 이 score matrix를 HBM에 쓰고, softmax를 위해 다시 읽고, V와 곱는다. Long context에서는 이 HBM traffic이 핵심 병목이 된다.

FlashAttention의 기본 아이디어는 score matrix를 HBM에 만들지 않는 것이다. Q tile을 잡고 K/V tile을 block by block으로 읽으며, online softmax state를 유지하고 output을 누적한다. Backward에서도 forward attention matrix를 저장하지 않고 recomputation을 활용한다. 수학은 같게 두고 memory hierarchy를 정면으로 이용한다.

2. Hopper에서 새로 보는 두 도구: WGMMA와 TMA

FlashAttention 3는 Hopper, 특히 H100의 feature를 적극적으로 쓴다. WGMMA는 warp group 단위의 asynchronous MMA primitive다. 네 개 warp가 함께 Tensor Core 연산을 수행하는 식으로 이해하면 된다. TMA(Tensor Memory Accelerator)는 single thread가 bulk tensor movement를 issue할 수 있는 asynchronous load mechanism이다.

이 둘은 kernel 구조를 바꾼다. Ampere식 cp.async multi-stage pipeline도 좋지만, Hopper에서는 producer가 TMA로 global memory에서 shared memory로 tile을 가져오고, consumer warp group이 WGMMA로 연산하는 식의 역할 분리가 더 강력해진다. Hardware가 asynchronous하게 일할 수 있으면, software schedule도 그에 맞게 짜야 한다.

3. Warp specialization과 ping-pong scheduling

강의에서 중요한 phrase는 producer와 consumer다. 일부 warp group은 memory load를 맡고, 다른 warp group은 MMA와 softmax computation을 맡는다. 이 inter-warp-group overlap이 warp specialization이다. Producer가 다음 tile을 가져오는 동안 consumer는 현재 tile을 계산한다.

Ping-pong scheduling은 두 consumer group이 번갈아 일하도록 구성한다. 한쪽이 WGMMA를 수행하는 동안 다른 쪽은 softmax나 준비 작업을 하는 식으로 gap을 줄인다. 이상적인 그림처럼 완벽히 겹치지는 않지만, barrier와 pipeline을 이용해 상당히 근접하게 만들 수 있다. 여기서 CUTLASS/CuTe는 tile layout과 pipeline object를 표현하는 언어가 된다.

4. Register와 shared memory는 schedule의 일부다

FlashAttention kernel은 register pressure가 크다. Q/K/V tile, softmax state, accumulator를 들고 있어야 하기 때문이다. Hopper에서는 producer와 consumer가 다른 register budget을 갖도록 조정하는 방식도 중요하다. Consumer에게 더 많은 register를 주어 spilling을 피하고, producer는 TMA issue 중심으로 가볍게 둔다.

Shared memory layout도 단순 buffer가 아니다. TMA로 load한 뒤 consumer가 쓰기 좋은 형태로 놓아야 하고, 필요하면 internal transpose도 한다. Memory movement와 compute schedule이 따로 놀면 WGMMA를 빨리 만들어도 전체 kernel은 느려진다. GPU kernel에서 layout은 항상 실행 계획의 일부다.

5. Persistent kernel과 work scheduling

FlashAttention 3는 persistent kernel 관점도 다룬다. Logical work tile 수만큼 CTA를 무작정 띄우는 대신, GPU를 채울 정도의 physical CTA를 띄워두고 work tile을 계속 할당한다. 이렇게 하면 불규칙한 work나 tail effect를 줄이고 SM utilization을 안정시키는 데 도움이 된다.

다만 persistent scheduling은 complexity를 데려온다. Work queue, tile assignment, synchronization이 필요하고, 모든 shape에서 같은 이득을 주지는 않는다. 이 강의는 FlashAttention 3를 “기발한 trick 하나”가 아니라 Hopper hardware와 algorithm이 맞물린 총체적인 schedule로 보게 한다.

6. 손으로 해볼 것

작은 attention을 naive하게 구현하고 score matrix의 크기와 HBM read/write 양을 계산해보자. 그 다음 tile 단위 online softmax를 Python으로 구현해 score matrix를 만들지 않고도 같은 output을 얻는지 확인한다.

CUTLASS code를 바로 다 읽기 어렵다면, FlashAttention 3 slide에서 producer/consumer 역할과 TMA/WGMMA가 등장하는 부분만 먼저 따라가자. 어떤 warp가 data를 옮기고, 어떤 warp가 compute를 하는지 색으로 칠하면 schedule이 덜 무섭다.

Source Links

• YouTube
• Lecture 036 repo folder
• Slides PDF
• Colfax Research blog
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_036.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_036.md

1. 왜 SASS를 보나

Arun Demeure는 강의 초반에 중요한 경고를 한다. H100 같은 GPU에서 완전한 custom SASS를 open-source assembler만으로 편하게 쓰기는 어렵다. 그렇다면 왜 SASS를 공부할까? 답은 더 좋은 CUDA와 Triton을 쓰기 위해서다.

High-level language는 빨리 시작하게 해준다. 하지만 kernel이 정말 한계에 가까워질수록 vector load, register allocation, uniform register, instruction scheduling 같은 detail이 성능을 좌우한다. SASS를 읽으면 compiler가 내 의도를 어떻게 바꾸었는지 볼 수 있다. GPU 최적화에서 “내 코드는 이렇게 생겼는데요”는 종종 변명거리가 되지 못한다.

2. PTX는 실행 instruction이 아니다

PTX는 virtual ISA다. 실제 GPU가 실행하는 것은 SASS다. Arun은 아주 단순한 integer addition 예시로 PTX와 SASS가 1:1이 아님을 보여준다. PTX에서는 두 입력 addition처럼 보이는 것이 SASS에서는 IADD3처럼 세 입력 instruction으로 합쳐질 수 있다. 사용하지 않던 operand slot이 다른 계산과 합쳐지는 식이다.

Bitwise operation도 마찬가지다. 여러 logical operation이 LOP3 같은 instruction으로 합쳐질 수 있다. 이 변환은 compiler가 한다. 그래서 source line과 SASS instruction의 mapping은 자주 애매하다. 한 instruction이 여러 source operation에서 왔을 수 있고, 한 source line이 여러 instruction으로 흩어질 수 있다.

3. Architecture마다 사소하게, 그래서 크게 다르다

Volta, Turing, Ampere, Hopper는 비슷해 보이지만 instruction과 resource model이 계속 바뀐다. 예를 들어 uniform register는 per-warp 값처럼 동작하며, compiler가 이를 사용할 수 있게 code를 써야 할 때가 있다. CUDA나 PTX에서 직접 모든 것을 통제할 수 있는 것은 아니다.

또한 일부 constant memory access는 instruction 안에 섞여 보이지만 cache miss가 나면 큰 latency를 만들 수 있다. 이런 detail은 흔히 rare case처럼 보이지만, performance counter나 SASS를 읽을 때 설명력이 생긴다. Low-level 공부의 가치는 모든 것을 외우는 데 있지 않다. 이상한 profiler 결과를 만났을 때 물어볼 질문이 늘어나는 데 있다.

4. SM은 하나의 큰 덩어리가 아니다

Arun은 SM 내부 partition 관점도 강조한다. Warp scheduler, register allocation, pipeline은 per-SM 하나로만 생각하면 너무 거칠다. 예를 들어 sub-core/partition 단위로 warp가 배치되고, register allocation threshold에 따라 occupancy가 계단식으로 바뀔 수 있다.

Pipeline 이름도 오해를 부른다. Integer pipe가 integer만 하는 것이 아니고, fp32 pipe도 어떤 integer multiply를 처리할 수 있다. Newer architecture에서는 fp32-only pipe가 추가되는 식으로 변화가 있다. 결국 bottleneck은 “FP32 FLOPs가 부족함”처럼 간단히 끝나지 않는다. Decode/fetch/schedule rate, pipeline mix, memory instruction이 함께 본문이다.

5. Occupancy만으로 latency를 숨길 수 없다

Little's law 관점도 나온다. Throughput을 유지하려면 latency를 덮을 만큼 충분한 parallel work가 필요하다. 여기서 parallelism은 thread-level occupancy만 뜻하지 않는다. Instruction-level parallelism, 동시에 outstanding load가 얼마나 있는지, warp들이 program의 어느 위치에 흩어져 있는지도 중요하다.

Profiler에서 warp stall을 볼 때도 조심해야 한다. “warp가 stall했다”는 말은 원인 하나를 말해주지 않는다. Memory dependency인지, scheduler pressure인지, pipeline contention인지, register spill인지 더 파고들어야 한다. SASS와 microarchitecture 지식은 이 진단을 덜 흐리게 만든다.

6. 손으로 해볼 것

작은 CUDA kernel을 Compiler Explorer나 cuobjdump/nvdisasm으로 보고 PTX와 SASS를 비교해보자. Add 두 개, bitwise op 몇 개처럼 아주 단순한 code가 오히려 좋다. Compiler가 instruction을 합치는 장면을 직접 보면 PTX를 보는 태도가 달라진다.

그 다음 같은 kernel을 조금 바꾸어 register count와 occupancy가 어떻게 변하는지 Nsight Compute로 본다. SASS를 읽는 목적은 멋진 assembly 암송이 아니다. Source를 어떻게 바꾸면 compiler와 hardware가 덜 힘들어하는지 배우는 것이다.

Source Links

• YouTube
• Lecture 037 repo folder
• Slides PDF
• Speaker GitHub
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_037.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_037.md

1. 왜 ARM CPU kernel인가

Scott Roy는 PyTorch/ExecuTorch 관점에서 low-bit ARM operator를 소개한다. ARM CPU는 mobile device와 Apple Mac에서 이미 널리 쓰이고, NVIDIA의 Grace Hopper, Grace Blackwell 같은 server-side system에서도 중요해지고 있다. Local LLM을 생각하면 GPU만큼 CPU path도 현실적인 target이다.

강의의 operator는 크게 두 종류다. 하나는 dynamic quantized activation과 low-bit weight를 갖는 linear operator이고, 다른 하나는 low-bit weight embedding operator다. Low-bit는 1bit부터 8bit까지를 다룬다. 이 operator들은 eager, torch.compile, AOT, ExecuTorch surface에서 쓰일 수 있도록 설계된다.

2. Affine quantization을 작게 잡기

Quantization은 float weight를 integer와 scale/zero point로 표현하는 일이다. 예를 들어 fp32 weight를 3bit integer 범위로 줄이면 memory footprint가 크게 감소한다. 하지만 계산할 때는 이 integer를 다시 float 의미로 해석해야 한다. 그래서 unpack과 dequantization이 kernel의 일부가 된다.

ARM CPU에서는 이 작업을 vector instruction과 cache behavior에 맞춰야 한다. GPU kernel처럼 수천 thread를 던지는 문제가 아니라, CPU threadpool과 SIMD vector lane을 잘 쓰는 문제다. 같은 “low-bit”라도 병목의 얼굴이 다르다.

3. Operator tiling과 microkernel tiling

강의에서 중요한 구분은 operator level tiling과 kernel 안쪽 tiling이다. Linear operator는 matrix를 output tile로 나누고, 각 tile을 threadpool이 병렬로 처리한다. 각 tile을 실제로 계산하는 lower-level kernel은 single-threaded지만 vectorized되어 ARM SIMD instruction을 사용한다.

Microkernel은 더 안쪽에서 register-level tiling을 한다. Scott은 regular C++ struct와 function으로 kernel config를 표현한다고 설명한다. CUDA처럼 별도 언어를 쓰는 것이 아니라, C++ 안에서 tile size, zero handling, bit-width별 unpack routine을 조합한다. GPU programming과 닮은 개념을 CPU vectorization 문법으로 푸는 셈이다.

4. Tile size는 완벽한 공식보다 workload에서 온다

Q&A에서는 cache size를 보고 tile size를 어떻게 정하느냐는 질문이 나온다. 답은 아주 이론적인 auto-tuning이 아니라, GEMV workload와 vectorization에 맞춘 practical choice에 가깝다. M dimension tile은 사실상 1에 가깝고, weight 쪽 tile은 vector lane과 unpack 효율을 고려한다.

강의에서는 ARM CPU용 GEMV kernel이 tile size와 zero 여부 등에 따라 여러 variant를 갖고, bit-width별 unpack routine도 따로 있다고 설명한다. Low-bit CPU kernel은 하나의 generic loop보다 여러 작은 specialist를 갖는 쪽이 유리하다. Code size와 performance 사이의 익숙한 tradeoff다.

5. PyTorch surface와 local inference

이 operator들은 torchchat과 ExecuTorch에서 사용할 수 있다고 소개된다. torchchat은 desktop에서 LLM을 locally run하는 PyTorch solution이고, ExecuTorch는 mobile/edge deployment 쪽과 연결된다. 즉 이 작업은 microbenchmark용 장난감이 아니라, 실제 user device에서 model을 돌리기 위한 path다.

Performance slide에서는 Llama 3.1 8B를 M1 Mac Pro에서 benchmark하는 장면이 언급된다. 정확한 숫자는 원자료에서 확인해야 하지만, 읽을 때 포인트는 분명하다. GPU가 없어도, 또는 GPU path가 충분치 않아도, low-bit CPU kernel은 local model experience를 좌우할 수 있다.

6. 손으로 해볼 것

작은 matrix-vector multiply를 잡고 weight를 4bit로 pack해보자. 가장 먼저 naive unpack 후 fp32/fp16 GEMV를 만들고, 그 다음 tile별로 unpack하면서 dot product를 하는 version을 만든다. Unpack을 hot loop 어디에 두는지가 곧 성능 질문이 된다.

Apple Silicon machine이 있다면 NEON/Accelerate baseline과 비교해보자. 없다면 C++ code만 읽어도 좋다. GPU Mode 강의 안에서 이 CPU lecture가 주는 교훈은 넓다. 좋은 kernel thinking은 hardware가 바뀌어도 tile, layout, vectorization이라는 비슷한 단어로 돌아온다.

Source Links

• YouTube
• Lecture 038 repo folder
• Slides PDF
• Speaker GitHub
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_038.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_038.md

1. 큰 training을 작게 읽기

2024년 마지막 GPU Mode 강의는 TorchTitan codebase walk-through에 가깝다. Mark Saroufim은 TorchTitan을 100개 이상 GPU scale의 training을 쉽게 시도할 수 있는 repository로 소개한다. 핵심 code는 대략 5천 줄의 Python이고, config 등을 포함해도 매우 작다. 그래서 한두 시간 안에 repo를 책처럼 읽는 것이 가능하다는 점이 강의의 분위기다.

Tianyu Liu는 performance 문서보다 더 큰 scale의 3D parallelism experiment도 paper에 있다고 보충한다. 정확한 최신 숫자는 repo/paper에서 확인해야 하지만, 이 강의에서 중요한 것은 scale record가 아니다. PyTorch distributed primitive를 어떻게 composable하게 엮는지 보는 것이다.

2. Copy-paste하라는 reference architecture

TorchTitan의 철학은 “너무 숨기지 않는다”에 가깝다. Mark는 강의 중 code가 적고 copy-paste해서 자기 project에 맞게 고치기를 원한다고 말한다. Megatron이나 DeepSpeed처럼 자체 core function을 깊게 갖는 방식과 다르게, TorchTitan은 PyTorch에 있는 기능을 드러내어 조합하는 쪽을 택한다.

이 접근은 교육적으로 좋다. Single GPU 작업에 익숙한 사람도 train.py에서 config, data loader, model build, parallelize, optimizer, checkpointing 흐름을 따라갈 수 있다. Distributed training을 처음 볼 때 제일 무서운 것은 수식보다 repo가 너무 커서 입구를 못 찾는 일이다. TorchTitan은 입구를 작게 만든다.

3. Data loader도 distributed object다

강의에서 조용히 중요한 부분은 data loading이다. Data parallel rank마다 다른 data를 받아야 하지만, tensor parallel이나 pipeline parallel처럼 같은 sample을 함께 처리하는 non-DP rank들은 같은 data를 봐야 한다. 그래서 DP degree와 DP rank가 data loader에 들어간다.

또 하나는 checkpointable data loader다. 거대한 token dataset을 한 번만 훑는 상황에서는 mid-epoch resumption이 중요하다. Training을 중간에 멈췄다가 checkpoint에서 재개할 때 dataset 처음으로 돌아가면 안 된다. Distributed training의 correctness는 model weight뿐 아니라 data iterator state에도 걸려 있다.

4. Model은 평범하게 만들고 parallelism을 얹기

TorchTitan은 model definition 자체를 최대한 distributed-aware하게 만들지 않으려 한다. 먼저 Llama 3.1 같은 transformer model을 평범하게 build하고, 그 위에 parallelism strategy를 적용한다. Tianyu는 이것이 Megatron류와 다른 점이라고 설명한다. Model을 침범하지 않고 parallelism을 조합하려는 방향이다.

물론 완전히 무침습적일 수는 없다. Tensor parallelism은 QKV projection, MLP, output head 같은 layer에 sharding과 communication hook을 넣어야 한다. Loss parallel은 vocabulary dimension이 큰 cross entropy를 sharded computation으로 처리한다. FSDP는 parameter/gradient/optimizer state shard를 관리한다. 중요한 것은 이 조합의 순서와 boundary를 code에서 볼 수 있다는 점이다.

5. Communication overlap과 activation checkpointing

TorchTitan은 async tensor parallel 같은 기능도 다룬다. 기본 tensor parallel communication은 blocking일 수 있다. Matrix를 여러 조각으로 나누고 일부 multiplication을 먼저 수행하면 communication과 computation을 동시에 진행할 수 있다. PyTorch의 symmetric memory 같은 기능이 이 방향과 연결된다.

Activation checkpointing도 중요한 memory-speed tradeoff다. 모든 activation을 저장하면 memory가 커지고, 모두 recompute하면 time이 커진다. TorchTitan은 full checkpointing과 selective checkpointing mode를 제공한다. Operator-based selective checkpointing은 offline benchmark와 경험에 기반한 list를 사용한다고 설명된다. 여기에도 정답 하나는 없다. 저장할 것과 다시 계산할 것을 workload에 맞게 고르는 일이다.

6. Distributed checkpoint와 손으로 해볼 것

Large training에서 checkpointing은 torch.save 한 번으로 끝나지 않는다. TorchTitan은 PyTorch Distributed Checkpoint(DCP)를 사용한다. 여러 GPU가 model shard를 나누어 저장하고, async mode로 CPU-to-disk stage를 overlap할 수 있다. GPU-to-CPU staging과 disk write를 어떻게 겹칠지에 따라 training interruption이 달라진다.

강의에서는 zero-overhead checkpointing 논의도 언급된다. Optimizer step 직후 저장할 state를 CPU로 옮기면서 다음 forward를 진행하고, backward가 시작되기 전에 copy가 끝나도록 맞추는 식이다. 아직 fault tolerance까지 모두 해결한 것은 아니지만, checkpoint가 training loop의 성능 feature라는 점은 분명하다.

TorchTitan repo에서 train.py를 열고, config parsing부터 model build, parallelize call, checkpoint call까지 한 화면짜리 call graph를 그려보자. Distributed training을 배울 때는 처음부터 NCCL 내부를 파기보다 “어느 시점에 어떤 parallelism이 적용되는가”를 잡는 편이 낫다.

작은 2-GPU 환경이 있다면 data loader가 DP rank별로 다른 sample을 주는지 확인한다. GPU가 없다면 code를 읽으며 TP, FSDP, activation checkpointing, compile이 적용되는 순서를 적어보자. TorchTitan은 실행하지 못해도 꽤 좋은 교과서다.

Source Links

• YouTube
• TorchTitan repository
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_039.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_039.md

1. 문서가 없었던 것이 아니라 길이 끊겨 있었다

Charles Frye는 Modal에서 만든 GPU Glossary를 보여주며 강의를 시작한다. nvcc compiler flag를 읽다가 compute capability가 나오고, compute capability를 따라가면 SM architecture가 나오고, 다시 compiler와 PTX로 이어진다. 각각의 정보는 NVIDIA 문서, 오래된 white paper, textbook, blog post에 있다. 문제는 초심자가 그 길을 스스로 찾아야 한다는 점이다.

그래서 이 강의의 분위기는 "CUDA 문서는 나쁘다"라기보다 "좋은 자료가 서로 충분히 연결되어 있지 않다"에 가깝다. Charles는 PTX ISA 문서, CUDA C++ Programming Guide, Best Practices Guide 같은 공식 자료를 높게 평가한다. 다만 이런 자료는 종종 PDF나 오래된 문서 형태라서, 웹에서 기대하는 식의 발견성과 상호 링크가 부족하다.

여기서 좋은 학습법도 나온다. 새로운 software system이나 programming model을 만나면 trace를 만들어서 클릭해 보라는 것이다. PyTorch profiler trace를 따라가며 ATen, host launch와 device execution의 비동기성, CUDA runtime과 driver의 경계를 배웠다는 이야기가 나온다. 멋진 kernel을 쓰기 전에, 먼저 내가 밟고 있는 stack을 보자는 말이다.

2. CUDA는 하나가 아니라 여러 층이다

강의에서 가장 중요한 정리는 "CUDA가 무엇인가?"라는 질문을 한 번에 답하지 않는 것이다. CUDA는 Python/PyTorch에서 보이는 software platform이기도 하고, C/C++나 Fortran에 들어가는 programming model이기도 하며, compute unified device architecture라는 hardware design 철학이기도 하다.

PyTorch 같은 application은 CUDA runtime API, CUDA driver API, libcuda, cudart 같은 층을 통해 GPU에 일을 보낸다. 이때 Python 사용자는 대부분 이 경계를 직접 보지 않는다. 하지만 문제가 생기면 에러 메시지는 아주 낮은 층에서 올라온다. 그래서 "PyTorch에서 CUDA 쓴다"는 말은 편하지만, debugging할 때는 너무 뭉툭한 말이 된다.

programming model 쪽에서는 kernel, grid, thread block, thread, shared memory 같은 개념이 나온다. 중요한 점은 이것들이 먼저 추상 개념이라는 것이다. thread block은 programming model의 단위이고, cooperative thread array는 PTX 쪽 구현 언어에서 보이는 표현이다. 둘을 구분하면 문서를 읽을 때 헷갈림이 줄어든다.

3. shared memory는 말 그대로만 들으면 위험하다

CUDA의 shared memory는 thread block 안의 thread들이 공유하는 programmer-managed memory다. 그런데 같은 말을 hardware 관점에서 하면 SM 안의 L1 data cache와 scratchpad 성격이 섞인 공간을 가리키게 된다. 강의 중간 질문에서도 이 지점이 바로 교정된다. 단순한 cache라고 부르면 hardware-managed cache를 떠올리기 쉽지만, CUDA shared memory는 programmer가 배치와 동기화를 관리한다.

이 차이는 성능과도 연결된다. programmer-managed memory는 귀찮지만, 그만큼 context switch나 memory traffic의 영향을 직접 제어할 수 있다. GPU kernel을 최적화할 때 shared memory를 "빠른 배열" 정도로만 이해하면 곧 막힌다. 누가 언제 읽고 쓰는지, block 안에서 어떤 synchronization이 가능한지, block 밖과는 어떤 보장이 없는지를 함께 봐야 한다.

Charles가 강조하는 CUDA programming model의 제약도 여기서 이해된다. block 사이에는 마음대로 synchronization할 수 없고, block들은 임의 순서로 실행될 수 있어야 한다. 처음에는 답답하지만, 바로 이 제약 덕분에 같은 kernel이 SM 수가 다른 GPU에서도 비교적 자연스럽게 scale한다.

4. PTX는 단순한 중간 산물이 아니다

Charles의 두 번째 강한 주장은 CUDA stack에서 PTX가 아주 중요하다는 것이다. PTX는 compiler가 내보내는 intermediate representation이면서 virtual ISA다. 실행 시점에는 실제 GPU의 SM에 맞는 SASS/cubin 쪽으로 JIT 또는 runtime compile된다.

이 층이 있기 때문에 forward compatibility가 가능해진다. 하나의 binary나 package가 새 GPU에서 다시 compile되어 실행될 수 있고, 때로는 새 hardware의 이점을 어느 정도 얻는다. 최고 성능을 위해 SASS를 직접 겨냥하는 경우도 있지만, 배포성과 호환성을 생각하면 PTX라는 완충층이 큰 의미를 가진다.

CPU 쪽 x86와 microcode 이야기도 잠깐 나온다. 요지는 NVIDIA가 비교적 늦게 온 덕분에 "사람과 compiler가 겨냥하는 instruction set"과 "실제 hardware가 실행하는 binary instruction" 사이를 더 명시적으로 나눌 수 있었다는 관찰이다. 문서 지도를 만들 때 PTX를 빼놓으면 CUDA의 허리가 사라진다.

5. 좋은 문서는 실행 가능한 문서에 가까워진다

강의 후반은 GPU Glossary의 다음 단계 이야기다. LLM이 읽기 좋은 llms.txt 형식, retrieval이나 graph RAG, 그림과 텍스트를 함께 다루는 방법이 언급된다. 단순히 사이트 오른쪽 아래에 chat widget을 붙이는 방식은 충분하지 않다는 반성도 있다.

더 흥미로운 방향은 문서 안에서 작은 code snippet을 실제로 실행해 보는 것이다. Charles는 Modal 환경에서 libcuda를 지워 보거나 linking 방식을 바꿔 보며 어떤 층이 깨지는지 확인했다고 말한다. 이런 식의 실험은 용어를 정의하는 것보다 강하다. stack의 경계가 손에 잡히기 때문이다.

앞으로 보강하고 싶은 내용으로 synchronization hierarchy, thread block cluster, performance debugging 용어, occupancy, bank conflict, GPU fleet management, NVML, nvidia-smi, multi-GPU execution, NCCL, NVLink, PCIe 등이 나온다. 즉 이 문서는 CUDA kernel 문법만이 아니라 GPU를 실제로 운영하는 사람의 지도까지 넓히려 한다.

6. 손으로 해볼 것

첫 과제는 아주 작게 시작하면 된다. PyTorch 프로그램 하나를 profiler로 찍고, trace에서 CPU op, CUDA runtime call, kernel execution이 어떻게 분리되어 보이는지 확인해 보자. "Python에서 한 줄"이 GPU에서는 어떤 launch와 device work로 보이는지 보는 것만으로도 stack 감각이 달라진다.

두 번째는 CUDA 문서 읽기 지도 만들기다. compute capability, SM, shared memory, PTX, cubin, SASS 여섯 단어를 골라 각 단어가 어느 문서에서 정의되는지, 어떤 단어로 다시 링크되어야 하는지 적어 보자. 모르는 단어를 줄이는 것보다, 모르는 단어 사이의 길을 만드는 것이 이 강의의 태도다.

Source Links

• YouTube
• Slides
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_040.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_040.md

1. 왜 또 attention engine인가

강의는 vLLM의 paged attention에서 출발한다. autoregressive inference serving에서는 요청마다 길이가 다르고, 최대 길이에 맞춰 cache를 잡으면 memory fragmentation이 커진다. paged attention은 KV cache를 작은 page 단위로 관리하고, attention kernel이 page table을 직접 따라가도록 만든다. 먼저 contiguous cache로 gather한 뒤 FlashAttention을 돌리는 방식보다 memory 효율이 좋다.

하지만 paged attention은 끝이 아니다. RadixAttention류의 prefix caching은 공유 prefix를 tree 구조로 관리하고, speculative decoding은 draft model이 만든 후보 tree를 검증한다. 긴 context에서는 Quest 같은 방식처럼 page별 reduced key로 중요도를 먼저 보고 top-k page만 본다. KV cache는 점점 더 "연속된 tensor"에서 멀어진다.

Zihao의 문제의식은 여기 있다. serving framework가 새 cache policy를 낼 때마다 attention kernel을 새로 쓰면 유지보수가 어려워진다. FlashInfer는 여러 cache management 전략을 kernel이 이해할 수 있는 공통 형식으로 바꾸려 한다.

2. block sparse로 KV cache를 설명한다

FlashInfer는 KV cache와 attention mask를 block sparse representation으로 본다. sparse matrix의 원소가 scalar 하나가 아니라 작은 block이라는 점이 중요하다. GPU Tensor Core는 일정한 MMA shape를 좋아하므로, 완전히 자유로운 sparse pattern보다 block 단위 구조가 hardware에 더 잘 맞는다.

paged attention도 이 관점으로 볼 수 있다. 전체 KV address space를 column으로 보고, 각 request가 접근하는 page를 nonzero block으로 두면 page table이 block sparse matrix가 된다. radix tree나 speculative decoding tree도 attention에서 허용되는 위치를 sparse structure로 표현할 수 있다.

다만 16x16이나 128x128 block을 고정하면 sparse ratio가 높은 tree attention에서 낭비가 커진다. 강의에서는 vector sparse에 가까운, column 수를 더 유연하게 잡는 설계가 나온다. 16x16 block 하나에 nonzero가 하나만 있어도 256개 element를 활성화하는 대신, 더 좁은 vector 단위로 낭비를 줄이는 식이다. 그래도 Tensor Core를 쓰기 위해 global memory에서 흩어진 row/column을 shared memory로 모은 뒤, on-chip에서는 dense하게 먹이는 구조를 유지한다.

3. dense처럼 계산하려면 먼저 잘 모아야 한다

block sparse 구조라고 해서 Tensor Core를 포기하는 것은 아니다. attention의 head dimension은 보통 contiguous하고 크기도 충분하다. 그래서 global memory에서는 page index를 따라가며 흩어진 KV를 읽더라도, shared memory에 올린 뒤에는 dense tile처럼 Tensor Core에 먹일 수 있다.

Hopper의 TMA는 contiguous tensor copy에는 강력하지만, page index를 거치는 nonlinear memory access에는 바로 맞지 않는다. 그래서 FlashInfer는 sparse KV loading에서는 Ampere식 asynchronous copy인 ldgsts로 fallback한다. 성능이 완전히 무너지지는 않지만, pointer arithmetic과 register pressure가 늘어난다. 강의의 평가에서도 FlashAttention 3 template에서 sparse와 dense의 gap이 더 커지는 이유가 여기에 있다.

이 지점이 재미있다. "최신 hardware feature를 쓰면 빠르다"가 아니라, data layout이 feature의 전제와 맞아야 한다. TMA가 좋은 도구라도 indirect index array를 따라가는 access pattern에는 다른 전략이 필요하다.

4. compiler는 단순하고, runtime은 현실적이다

FlashInfer는 attention variant를 바꾸기 위해 compiler와 runtime layer를 둔다. compile time에는 CUDA/CUTLASS template과 Jinja string generation을 써서 kernel을 만든다. Zihao는 ML compiler 경험이 있지만 여기서는 단순함 때문에 Jinja template을 선택했다고 말한다. type checking과 debugging은 약하지만, 사용자는 생성된 CUDA code를 볼 수 있고 torch.load류의 흐름으로 빠르게 붙일 수 있다.

customization은 FlexAttention에서 영향을 받았다. score transform, logits transform, logits mask 같은 functor를 넣어 sigmoid attention, sliding window, paged attention, custom mask, RoPE fusion 같은 변형을 표현한다. 단, 모든 attention을 다 포괄하지는 않는다. QK를 만든 뒤 elementwise transform과 scan/online softmax류를 거쳐 V와 곱하는 계열이 주 영역이다. linear attention처럼 계산 구조가 완전히 다른 것은 밖에 있다.

runtime 쪽은 serving의 변동성을 맡는다. prefill은 query length와 KV length가 크고, decode는 query length가 1에 가깝다. append나 incremental prefill은 그 중간이다. FlashInfer는 query tile size 16, 32, 64, 128 같은 여러 kernel을 준비하고, workload 특성에 따라 고른다.

5. variable length는 scheduler 문제다

serving batch 안에는 긴 request와 짧은 request가 섞인다. 그대로 CTA에 나누면 어떤 SM은 일이 끝나고 놀고, 어떤 SM은 긴 request에 붙잡힌다. FlashInfer의 runtime scheduler는 attention work를 tile로 쪼갠 뒤, 큰 작업부터 가장 덜 바쁜 CTA에 배치하는 식으로 load balancing을 한다.

split-K와 비슷한 느낌이 있지만, atomic reduction을 피하고 deterministic output을 유지하려는 선택이 들어간다. floating point accumulation order가 바뀌면 결과가 조금 달라질 수 있고, serving system에서는 이 작은 차이가 재현성 문제로 이어질 수 있다. 그래서 partial output을 만들고, 두 번째 contraction kernel에서 정해진 순서로 merge한다.

API도 serving framework와 맞춘다. init, plan, run 단계가 있고, plan은 generation step마다 scheduling 정보를 만든다. 같은 KV cache layout이 여러 layer에서 재사용되기 때문에 planning cost는 amortize될 수 있다. CUDA Graph와 torch.compile과의 호환성도 언급된다.

6. 성능 숫자를 읽는 태도

강의에는 여러 benchmark가 나온다. page size 1의 paged attention에서 FA2 template은 sparse KV와 contiguous KV의 gap이 10% 안쪽으로 보이고, FA3 template은 TMA를 못 쓰는 영향으로 gap이 더 커진다. variable length에서는 constant distribution보다 uniform이나 Zipf-like distribution에서 load balancing 효과가 더 잘 드러난다.

end-to-end에서는 SGLang과 결합해 Llama 3.1 8B/70B, 1 H100/4 H100 설정에서 비교한 결과가 언급된다. transcript 기준으로 "up to 2x latency reduction" 같은 표현이 나오지만, 정확한 표와 조건은 원본 slide 확인이 필요하다. 좋은 독자는 숫자만 외우지 않는다. 어떤 distribution을 썼는지, attention이 병목인지, plan overhead가 amortize되는지, TMA를 쓸 수 있는 access pattern인지 함께 본다.

마지막으로 로드맵도 중요하다. NanoFlow식 intra-device parallelism, 제한된 SM budget에서 kernel 실행, CUDA green context, fully reproducible serving kernel, FlexAttention alignment, DP acceleration 같은 주제가 나온다. FlashInfer는 attention kernel 하나가 아니라 serving system과 같이 움직이는 CUDA kernel library가 되려 한다.

7. 손으로 해볼 것

작은 batch를 만들고 request별 KV length를 [128, 2048, 64, 1024]처럼 섞어 보자. attention work를 동일 크기 tile로 나눈 뒤 greedy하게 CTA에 배치하면 각 CTA의 총 work가 어떻게 달라지는지 손으로 계산해 보면 runtime scheduler의 필요가 보인다.

두 번째로 paged KV cache를 block sparse matrix로 그려 보자. row는 request, column은 page address로 두고, 각 request가 접근하는 page만 색칠한다. 이 그림 위에 shared prefix가 있는 request들을 추가하면 왜 "하나의 block size"가 항상 좋은 답이 아닌지 보인다.

Source Links

• YouTube
• Slides
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_041.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_041.md

1. 예전 ML framework의 목표와 지금의 목표

Adam Paszke는 PyTorch 시절의 목표를 generality와 usability로 설명한다. 다양한 연구 아이디어를 빠르게 돌리는 것이 중요했고, 성능을 어느 정도 잃어도 연구 속도가 이겼다. 하지만 Transformer가 너무 성공하면서 상황이 바뀌었다. 많은 workload가 matmul, attention, normalization, pointwise op 주변으로 모였고, 몇 퍼센트 성능 차이가 큰 비용 차이가 된다.

그래서 kernel language의 목적도 달라진다. 처음부터 PyTorch나 JAX 대신 쓰는 생산성 도구가 아니다. 기본 framework로 충분히 빠르면 그대로 쓰면 된다. kernel language는 성능이 아쉬울 때 꺼내는 도구이고, 그때는 peak performance가 첫 번째 목표가 된다.

Mosaic GPU와 Palace는 이 위치에 있다. Python front end를 쓰지만, 편함만을 위해 low-level hardware detail을 숨기지는 않는다. 오히려 TMA, WGMMA, register allocation hint, block cluster처럼 성능에 영향을 주는 knob을 드러낸다. 대신 반복적이고 실수하기 쉬운 pipeline wiring은 helper로 줄인다.

2. Python DSL의 여러 맛

강의 초반에는 C++ template, string-based metaprogramming, Triton식 AST parsing, JAX/Palace식 tracing의 tradeoff가 나온다. C++ template은 강력하지만 error message와 debugging이 고통스러울 수 있다. string template은 단순하고 실용적이지만 scope, name collision, type checking 문제가 있다.

Triton은 Python AST를 읽고 constexpr 같은 staged programming을 제공한다. Palace는 JAX처럼 tracing한다. Python if는 trace 시점에 한 branch만 남고, Python for는 unroll되는 식이다. 일반 JAX에서는 이 제약이 불편할 때가 있지만, kernel에서는 오히려 자연스럽다. kernel programming에서는 unroll과 specialization이 성능 전략이기 때문이다.

Adam이 원하는 것은 "manual kernel을 잘 쓰는 playground"다. compiler가 완벽한 결정을 해 주기를 기다리지 않고, 사람이 tile size, pipeline, memory space, warp specialization 같은 중요한 선택을 직접 한다. 하지만 선택과 무관한 잡일은 자동화한다.

3. Palace의 기본 감각: reference와 block spec

Palace kernel의 인자는 일반 JAX array가 아니라 shaped reference다. x_ref[...]처럼 읽으면 그 시점의 값을 register 쪽 immutable array로 가져오고, z_ref[...] = ...처럼 쓰면 output reference에 저장한다. C의 pointer처럼 생겼지만, shape가 있는 reference라서 NumPy식 slicing 감각을 유지한다.

parallelism은 grid로 만든다. 하나의 program invocation이 output의 어느 tile을 맡는지 정하고, program_id에 해당하는 grid coordinate로 slice를 잡는다. Triton의 block pointer와 비슷한 문제를 Palace는 BlockSpec으로 더 명시적으로 표현한다.

BlockSpec은 두 가지를 말한다. 첫째, input/output을 어떤 tile shape로 자를 것인가. 둘째, grid coordinate가 그 tile coordinate로 어떻게 mapping되는가. 이 access pattern을 kernel body 밖에서 선언하면 runtime/compiler가 memory transfer와 compute를 pipeline할 수 있다. 여기서 중요한 말은 "optimization이 우연히 발견된다"가 아니라 "programmer가 access pattern을 약속했기 때문에 pipeline이 가능하다"다.

4. modern GPU는 TPU를 닮아간다

Adam은 GPU와 TPU의 수렴을 흥미롭게 설명한다. classic GPU programming model은 많은 work를 flight에 띄워 memory latency를 숨기는 데 강했다. 하지만 modern GPU에서 WGMMA 같은 큰 Tensor Core instruction을 쓰면 한 block이 shared memory와 register를 많이 먹는다. 한 SM에 여러 block을 올려 latency를 숨기는 방식이 약해진다.

그래서 Hopper 이후의 GPU는 더 계획적인 pipeline을 요구한다. TMA로 HBM-to-SMEM transfer를 명시적으로 관리하고, WGMMA로 warp group 단위 matmul을 돌리며, block cluster와 TMA multicast로 block 사이 shared memory에 data를 뿌린다. 이 모습은 TPU의 static scheduling 감각과 가까워진다.

Mosaic GPU는 이 흐름을 전제로 한다. "portable한 하나의 kernel"보다 "hardware generation의 feature를 제대로 쓰는 kernel"을 우선한다. portability는 좋지만, 목표가 최고 성능이면 backend-specific API를 쓰는 것이 자연스럽다.

5. TMA, WGMMA, warp specialization을 어떻게 숨기지 않을까

Mosaic GPU는 SMEM allocation을 추상화하지 않는다. run_scoped로 shared memory와 barrier를 잡고, GMEM-to-SMEM copy를 명시할 수 있다. TMA descriptor 생성과 host/device descriptor 관리 같은 지루한 부분은 helper가 맡는다. 즉 control은 남기고 잡음은 줄인다.

emit_pipeline은 더 높은 수준의 helper다. programmer가 "이 compute를 여러 step으로 돌리고, 각 step에서 어떤 slice를 읽고 쓸지"를 BlockSpec으로 말하면, helper가 async copy와 wait, double buffering에 가까운 pipeline 구조를 만들어 준다. TMA swizzle이나 tiling transform도 block spec에 붙일 수 있다. Tensor Core가 좋아하는 layout으로 copy 도중 data를 재배치하는 식이다.

WGMMA도 같은 흐름이다. Hopper의 warp group MMA는 shared memory에서 operand를 읽고 accumulator에 누적한다. Mosaic GPU는 이를 Python DSL에서 호출할 수 있게 한다. block cluster와 TMA multicast도 직접 assembly를 엮는 대신 "이 copy를 cluster의 특정 dimension에 collective하게 하라"는 형태로 표현하려 한다.

6. FlashAttention 3와 profiler가 보여준 것

강의에서 가장 구체적인 demo는 FlashAttention 3 kernel이다. Adam은 Palace/Mosaic GPU 구현이 Hopper에서 70%대 Tensor Core utilization을 얻었고, kernel도 100-150줄 정도의 Python으로 표현된다고 설명한다. 정확한 숫자는 slide와 benchmark 조건을 함께 확인해야 하지만, 메시지는 분명하다. 낮은 층을 쓰더라도 boilerplate를 줄이면 kernel이 읽을 수 있는 크기로 남는다.

또 하나 인상적인 부분은 work group level profiler다. JAX named_scope를 Palace 안에서도 쓰게 하고, kernel 내부 event를 SMEM에 기록한 뒤 GMEM으로 복사해 Chrome tracing이나 Perfetto에서 본다. 각 process를 SM처럼, 각 line을 warp group처럼 해석해 QK, softmax, PV, barrier wait가 시간축에서 어떻게 겹치는지 본다.

FlashAttention 3에서는 한 warp group이 softmax를 계산하는 동안 다른 warp group이 matmul을 하도록 barrier로 stagger한다. softmax는 ALU를 쓰고 matmul은 Tensor Core를 쓰므로, 두 hardware unit을 동시에 바쁘게 만들려는 전략이다. head dimension 64와 256에서 softmax와 matmul의 상대 비용이 달라지는 모습도 trace에서 보인다.

7. 손으로 해볼 것

Palace를 바로 설치하지 않아도, 먼저 matmul의 BlockSpec을 종이에 써 보자. C의 (pid_m, pid_n) tile을 계산하려면 A는 (pid_m, 0) tile, B는 (0, pid_n) tile, C는 (pid_m, pid_n) tile을 본다. 이 mapping을 lambda로 적는 연습만 해도 "pointer arithmetic 대신 access pattern 선언"이라는 감각이 잡힌다.

그 다음 FlashAttention을 hardware resource 관점으로 나눠 보자. QK와 PV는 Tensor Core, softmax는 ALU, KV loading은 memory pipeline, barrier는 synchronization이다. trace를 상상하며 어느 resource가 비는지 적어 보면 왜 warp specialization이 단순한 마법 단어가 아닌지 보인다.

Source Links

• YouTube
• Lecture 042 repo folder listed in local notes
• Slides link listed in local notes
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_042.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_042.md

1. 왜 하필 int8이고 왜 Turing인가

Erik은 처음부터 목표를 분명히 한다. 이 코드는 최고 성능용 production GEMM이 아니라 교육용이다. CUTLASS 같은 거대한 header library도 쓰지 않고, inline assembly도 직접 쓰지 않는다. CUDA toolkit의 WMMA 계열 API로 Tensor Core의 기본 사용법을 익힌다.

int8 matmul은 좋은 학습 재료다. 8-bit integer 두 개를 곱하면 결과는 더 넓은 accumulator가 필요하고, CPU에서는 AVX512 VNNI 같은 전용 dot product instruction이 등장한다. CUDA에도 Tensor Core를 쓰지 않는다면 dp4a 같은 네 개짜리 8-bit dot product instruction이 있다. Tensor Core로 가면 이 생각을 tile 단위로 확장한다.

또 int8은 결과 검증이 상대적으로 깔끔하다. floating point matmul은 accumulation order가 바뀌면 roundoff가 달라지고 tolerance를 정해야 한다. integer matmul은 overflow 조건을 조심해야 하지만, 같은 수학을 하면 exact result를 기대할 수 있다. 교육용 benchmark와 자동 채점에 잘 맞는다.

2. test가 좋아야 kernel을 고칠 수 있다

강의에는 온라인 제출 시스템 이야기가 나온다. 사용자가 code를 제출하면 correctness test, compute sanitizer, benchmark를 돌리고 점수를 준다. 이때 단순히 "틀렸습니다"라고 말하면 학습이 잘 되지 않는다. 어느 block boundary에서 처음 틀렸는지, 어떤 structured input에서 깨지는지 알려줘야 한다.

검증 방법도 여러 가지다. 작은 matrix는 눈으로 보기 쉽지만 실제 GPU block과 Tensor Core fragment 경로를 충분히 건드리지 못한다. identity matrix나 block-constant structured matrix는 indexing bug를 찾는 데 좋다. 큰 matrix의 전체 reference matmul은 비싸므로 Freivalds algorithm 같은 probabilistic check도 쓸 수 있다. C = A B라면 임의 vector x에 대해 C x와 A (B x)를 비교하는 방식이다.

이 이야기는 kernel 최적화 전에 아주 중요하다. 빠른데 틀린 kernel은 그냥 빠른 쓰레기다. 조금 무뚝뚝하지만, GPU는 그런 농담을 자주 친다.

3. scalar matmul에서 fragment matmul로 시선을 올리기

첫 구현은 naive scalar matmul을 16x16 fragment 단위로 바꾼 것에 가깝다. A는 n x k, B는 k x n int8이고, C는 n x n int32 accumulator다. scalar 관점에서는 C[i,j] 하나가 K 방향 inner product다. Tensor Core 관점에서는 16x16 tile fragment를 곱하고 accumulator fragment에 더한다.

CUDA WMMA API에는 fragment type, load_matrix_sync, mma_sync, store_matrix_sync, fill_fragment 같은 기본 함수가 있다. mma_sync가 fragment multiply-add를 수행한다. 여기서부터 중요한 경고가 따라온다. fragment 내부 storage layout은 unspecified이고 architecture마다 바뀔 수 있다. 즉 fragment 안의 element mapping을 마음대로 가정하면 API contract 밖으로 나간다.

첫 버전은 돌아가지만 느리다. Tensor Core를 호출했는데도 spec sheet의 speed-of-light와 멀다. 이유는 늘 그렇듯 "where's my data?"다. B가 row-major로 저장되어 K 방향 access가 strided가 되고, global memory load가 비효율적이다.

4. 첫 번째 교정: B를 column-major로 둔다

A와 B가 모두 row-major이면 matmul에서 B를 K 방향으로 읽을 때 stride가 커진다. Tensor Core 자체는 빠르지만, data가 제때 못 오면 warp는 기다린다. Nsight Compute의 memory chart에서는 L2와 L1 사이, device memory와 L2 사이 traffic이 크게 보이고, warp stall에서는 LG throttle과 long scoreboard가 눈에 띈다.

B를 column-major 형식으로 바꾸면 access pattern이 좋아진다. 별도 transpose kernel이 필요하지만, 큰 matmul은 O(n^3)이고 transpose는 O(n^2)라 충분히 큰 matrix에서는 overhead가 작다. transcript 기준으로 transpose overhead가 몇 ms 수준이고 kernel 시간이 크게 줄어드는 예가 나온다.

여기서 profiler 읽기의 작은 함정도 나온다. 최적화 뒤 kernel이 빨라졌는데 normalized warp stall 비율이 오히려 커질 수 있다. 분모인 "useful instruction을 낸 cycle"도 줄었기 때문이다. counter 하나만 보고 판단하면 안 된다. timing, instruction count, memory traffic, stall reason을 같이 봐야 한다.

5. register reuse: warp가 더 많은 output을 맡는다

다음 최적화는 register reuse다. 한 warp가 output fragment 하나만 계산하면 input fragment를 재사용할 기회가 적다. Erik은 예시로 warp 하나가 3x3 output fragments를 계산하도록 하며, A 쪽 3개, B 쪽 3개 input fragment를 load해 9개 accumulator fragment에 누적한다.

이 생각은 scalar tiled matmul과 같다. 단지 scalar element가 아니라 16x16 fragment를 단위로 볼 뿐이다. PMPP에서 thread coarsening이라고 부르는 감각이 여기서는 warp coarsening에 가깝다. data reuse가 늘고 global load와 index calculation이 줄어든다.

물론 공짜는 없다. accumulator fragment와 input fragment를 더 많이 들고 있어야 하므로 register usage가 늘어난다. 아직 spill이 나지 않더라도 occupancy가 떨어진다. 하지만 줄여야 할 latency 자체가 크게 줄었으므로 net win이 될 수 있다. 어느 정도까지 coarsening할지는 benchmark로 찾아야 한다.

6. shared memory는 빠르지만 공짜는 아니다

그 다음은 shared memory tiling이다. 여러 warp가 같은 input tile을 재사용할 수 있게 global memory에서 shared memory로 옮겨 두고, load_matrix_sync가 shared memory에서 fragment를 읽게 한다. Turing에는 Hopper의 TMA가 없으므로 thread들이 직접 vectorized load/store로 copy한다. 16-byte alignment를 compiler에 알려주기 위해 int4로 reinterpret하는 작은 기술도 나온다.

shared memory를 쓰면 device memory traffic은 크게 줄지만, 새로운 병목이 생긴다. __syncthreads() 때문에 barrier stall이 생기고, register spill이 local memory/L1/L2 traffic으로 보일 수 있다. 또 shared memory load 자체가 MIO throttle, short scoreboard 같은 stall reason으로 등장한다. shared memory는 global memory보다 빠르지만 무한히 빠른 공간은 아니다.

강의 후반의 재미있는 최적화는 fragment layout을 직접 해석하지 않고 우회하는 방법이다. load_matrix_sync로 memory layout에서 register layout으로 한 번 변환한 뒤, fragment 내부 storage를 blind copy하듯 store/load direct helper로 다룬다. element mapping을 알 필요는 없고, store와 load가 서로 inverse라는 점만 쓴다. 이로 shared memory access와 address calculation을 줄이며 추가 speedup을 얻는다.

7. 손으로 해볼 것

먼저 16x16 fragment 하나를 scalar element 하나처럼 생각하고, block matrix multiplication 식을 써 보자. C_ij += A_ik * B_kj에서 각 문자가 scalar가 아니라 fragment라고 보면 WMMA 코드가 훨씬 덜 낯설다.

그 다음 Nsight Compute를 켜고 같은 kernel의 두 version을 비교해 보자. B row-major version과 B column-major version을 만들고, wall time만 보지 말고 memory traffic, instruction count, LG throttle, long scoreboard를 같이 적는다. counter가 왜 좋아졌는지, 혹은 왜 직관과 다르게 보이는지 설명하는 연습이 이 강의의 진짜 과제다.

Source Links

• YouTube
• GPU Mode lectures repo
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_043.md
• Local transcript: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_043.md

1. 지금 확인된 것과 확인되지 않은 것

사용 가능한 local notes에 따르면 이 강의의 제목은 NVIDIA Profiling이고, transcript fetch는 실패했다. speaker, slides, code, related repo folder도 확인되지 않았다. 그러므로 아래 내용은 강의 내용의 재구성이 아니라, 이 강의를 다시 볼 때 붙잡을 reading guide다.

특히 어떤 profiler를 썼는지, Nsight Systems인지 Nsight Compute인지, NVTX range나 CUDA event를 어떻게 다뤘는지, 어떤 benchmark를 보여줬는지는 확인 필요다. 숫자나 tool option을 지어내지 않는 편이 낫다. profiling 강의에서 틀린 counter 이름 하나는 작은 오타가 아니라 잘못된 debugging 습관이 된다.

2. profiling은 시간을 쪼개는 일이다

NVIDIA Profiling을 읽을 때 첫 질문은 "전체가 느리다"가 아니라 "어느 층에서 시간이 쓰이는가"여야 한다. Python overhead인지, kernel launch overhead인지, H2D/D2H copy인지, kernel 내부 memory stall인지, Tensor Core utilization 문제인지 분리해야 한다.

보통 trace 수준 도구는 timeline을 보여주고, kernel 수준 도구는 counter를 보여준다. 이 구분을 머릿속에 두면 좋다. timeline은 "언제 무엇이 실행됐나"를 묻고, counter는 "그 kernel 안에서 왜 기다렸나"를 묻는다.

3. counter는 답이 아니라 다음 질문이다

occupancy가 낮다고 항상 occupancy를 올려야 하는 것은 아니다. register를 더 쓰면서 memory traffic을 줄이면 occupancy는 내려가도 kernel은 빨라질 수 있다. long scoreboard가 높다고 무조건 prefetch를 넣는 것도 아니다. 실제로는 uncoalesced access, cache miss, dependency chain, insufficient independent work가 섞여 있을 수 있다.

좋은 profiling 습관은 counter를 처방전처럼 읽지 않는 것이다. counter는 다음 실험을 고르게 해 준다. memory throughput이 낮고 warp가 memory를 기다린다면 layout을 바꿔 보고, shared memory stall이 크다면 bank conflict나 load pattern을 의심한다. 하지만 실험은 하나씩 해야 한다. 두 개를 동시에 바꾸면 profiler가 답을 해도 사람이 못 알아듣는다.

4. 작은 kernel 하나로 도구를 익히자

이 강의를 다시 볼 때는 vector add보다 조금 더 구조가 있는 kernel을 곁에 두면 좋다. 예를 들어 row-wise reduction, transpose, tiled matmul 같은 작은 kernel을 준비하고, naive version과 optimized version을 만든다. 그러면 profiler counter의 변화가 더 잘 보인다.

특히 transpose는 memory coalescing을 보기 좋고, reduction은 synchronization과 occupancy를 보기 좋다. matmul은 Tensor Core, shared memory, register pressure, arithmetic intensity를 한꺼번에 다루지만 처음 도구를 익히기에는 조금 클 수 있다.

5. 확인할 것

원본 영상을 볼 때는 세 가지를 확인하자. 첫째, 강의자가 어떤 tool을 중심으로 설명하는가. 둘째, "항상 먼저 보는 화면"이 무엇인가. 셋째, counter를 보고 실제 code change로 이어지는 예제가 있는가.

자료가 보강되면 이 partial 원고는 실제 예제 중심으로 다시 써야 한다. 지금 버전은 독자를 속이지 않는 placeholder에 가깝다. 빈칸을 솔직하게 남기는 것도 reader의 품질이다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_044.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_044.md

1. 지금 말할 수 있는 범위

local notes가 확인해 주는 것은 제목, speaker handle, YouTube link, transcript failure뿐이다. slides, code, paper, repo folder는 없다. 그러므로 이 글은 강의의 실제 최적화 기법을 요약하지 않는다. H100, cuBLAS, GEMM specialization을 읽을 때의 틀만 제공한다.

특히 "outperforming cuBLAS"라는 문장은 언제나 조건부로 읽어야 한다. 모든 shape, dtype, batch, transpose flag, epilogue, alignment, warmup, workspace, stream setting에서 cuBLAS를 이긴다는 뜻일 가능성은 낮다. 보통은 특정 problem family에서 library의 일반성을 줄이고 더 공격적인 선택을 했다는 뜻이다.

2. cuBLAS는 약한 상대가 아니다

cuBLAS는 단순한 baseline이 아니다. NVIDIA가 architecture별로 오래 다듬은 dense linear algebra library다. 그래서 cuBLAS보다 빠른 kernel이 나오려면 대개 더 좁은 문제를 잡아야 한다. 예를 들어 fixed shape, fixed dtype, fixed layout, no transpose, 특정 epilogue fusion, 특정 batch 구조처럼 제약이 들어간다.

H100에서는 WGMMA, TMA, warp specialization, shared memory swizzle, register allocation 같은 주제가 중요할 수 있다. 하지만 이 강의가 실제로 무엇을 썼는지는 확인 필요다. reader에서는 "아마 썼을 것이다"라고 쓰지 않는다. 대신 원본을 볼 때 확인할 checklist로 남긴다.

3. H100 kernel에서 물어볼 질문

H100 kernel을 읽을 때는 data movement가 먼저다. TMA로 global memory에서 shared memory로 tile을 옮기는가, swizzle로 bank conflict를 피하는가, double buffering이나 multi-stage pipeline으로 copy와 WGMMA를 겹치는가를 본다.

다음은 warp specialization이다. producer warp group과 consumer warp group을 나누는지, register를 setmaxnreg류의 hint로 재분배하는지, barrier가 어디에 놓이는지 확인해야 한다. 이런 부분은 작은 code change가 큰 성능 차이를 낼 수 있고, 반대로 debugging이 매우 어려운 구간이기도 하다.

4. benchmark를 믿기 전에 보는 것

timing은 반드시 synchronization을 포함해 올바르게 재야 한다. CUDA event를 쓰더라도 warmup, stream, graph capture 여부, input allocation 위치, clock 상태가 영향을 준다. cuBLAS와 custom kernel이 같은 data layout을 받는지, custom kernel만 유리한 pre-processing을 빼먹지 않았는지도 봐야 한다.

그리고 speedup은 평균 하나보다 분포가 낫다. 특정 K에서만 빠른지, 작은 M/N에서 launch overhead가 지배하는지, 큰 matrix에서 memory bandwidth가 충분한지 봐야 한다. "cuBLAS보다 빠름"은 headline이고, engineer에게 필요한 것은 조건표다.

5. 손으로 해볼 것

원본을 보기 전 준비 운동으로, 하나의 fixed GEMM shape를 정해 보자. 예를 들어 M=N=K=4096 같은 큰 square GEMM과, LLM inference에서 나올 법한 skinny GEMM을 하나씩 잡는다. 각 shape에서 cuBLAS가 얼마나 일반적인 요구를 처리해야 하는지, custom kernel이 무엇을 가정할 수 있는지 적어 보자.

강의를 다시 볼 때는 모든 성능 claim 옆에 "조건"을 붙인다. dtype, shape range, GPU, clock, baseline, fusion 여부, pre/post-processing 포함 여부를 채우면 좋은 reader 원고로 확장할 수 있다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_045.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_045.md

1. local GEMM과 distributed GEMM은 병목의 모양이 다르다

한 GPU 안의 GEMM에서는 Tensor Core utilization, shared memory tiling, register reuse, TMA pipeline 같은 질문이 먼저 온다. Distributed GEMM에서는 여기에 통신이 붙는다. A와 B를 어떻게 나누고, partial result를 어디서 합치며, 어떤 collective를 쓰고, compute와 communication을 어떻게 겹칠지가 성능을 좌우한다.

local notes에는 title과 speaker, slide link만 있고 transcript는 실패했다. 따라서 강의가 row/column parallelism, 2D SUMMA, ring, all-gather, reduce-scatter 중 무엇을 다뤘는지는 확인 필요다. 이 글에서는 distributed GEMM을 읽을 때 필요한 기본 틀만 잡는다.

2. partition이 algorithm을 정한다

GEMM C = A B에서 A를 row-wise로 나누면 각 GPU는 C의 row shard를 계산할 수 있지만 B 전체가 필요할 수 있다. B를 column-wise로 나누면 C의 column shard를 얻지만 A 전체가 필요하다. 2D partition은 memory와 communication을 더 균형 있게 만들 수 있지만 scheduling과 collective가 복잡해진다.

어떤 partition이 좋은지는 cluster topology와 matrix shape에 달려 있다. NVLink, PCIe, InfiniBand, node boundary가 모두 다르다. 같은 FLOP 수라도 data movement가 다르면 distributed GEMM의 시간이 달라진다. 그래서 distributed GEMM은 local roofline만으로 설명되지 않는다.

3. overlap이 없으면 통신이 그대로 보인다

여러 GPU에서 matmul을 하려면 어차피 data를 주고받아야 한다. 중요한 것은 그 시간이 compute 뒤에 줄 서느냐, compute와 겹치느냐이다. tile 단위로 쪼개고, 한 tile을 계산하는 동안 다음 tile을 all-gather하거나 reduce-scatter하면 통신 비용을 일부 숨길 수 있다.

다만 overlap은 공짜가 아니다. stream, event, buffer lifetime, collective ordering이 모두 맞아야 한다. 너무 작은 tile은 launch와 collective overhead를 키우고, 너무 큰 tile은 overlap 기회를 줄인다. H100 같은 빠른 GPU에서는 compute가 짧아져 통신을 숨기기 더 어려워질 수도 있다.

4. correctness도 분산된다

distributed GEMM에서는 성능만 아니라 correctness debugging도 어려워진다. 한 rank의 shard shape가 틀리거나, collective group이 어긋나거나, accumulation order가 달라지면 결과가 조금씩 혹은 크게 달라진다. mixed precision이나 low precision accumulation이 끼면 재현성 문제도 커진다.

강의를 볼 때는 benchmark보다 먼저 data layout diagram을 찾자. 각 GPU가 어떤 A/B/C shard를 들고 시작하고 끝나는지, 중간 buffer가 어떤 shape인지 그릴 수 있어야 한다. 그림이 안 그려지면 performance number도 오래 못 간다.

5. 손으로 해볼 것

GPU 4개를 2x2 grid로 두고 C = A B를 계산한다고 상상해 보자. A와 B를 각각 2D block으로 나눈 뒤, 각 GPU가 어떤 block을 갖고 있고 어떤 block을 받아야 C_ij를 계산할 수 있는지 써 보자. 그 다음 각 step에서 필요한 all-gather 또는 send/recv를 표시한다.

원본 강의를 다시 볼 때는 slide의 algorithm 이름을 확인하고, 이 손그림과 비교하면 좋다. 특히 "communication volume", "overlap", "topology aware" 같은 단어가 나오면 실제로 어떤 buffer와 collective를 말하는지 붙잡아야 한다.

Source Links

• YouTube
• Slides
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_046.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_046.md

1. 지금 자료의 한계

local notes에는 제목, YouTube link, transcript failure만 있다. speaker, slides, code, repo folder는 확인되지 않았다. 따라서 이 글은 KernelBot이 실제로 어떤 Discord command, backend queue, GPU pool, scoring system을 썼는지 설명하지 않는다.

대신 GPU kernel benchmark bot을 읽을 때 필요한 원칙을 정리한다. Lecture 43의 int8 matmul 강의에도 온라인 제출 시스템 이야기가 나왔기 때문에, 이 강의는 그 community infrastructure 쪽을 더 직접적으로 다뤘을 가능성이 있다. 단, 그것도 추정이므로 원본 확인 전까지는 가능성으로만 둔다.

2. benchmark bot은 채점기이자 선생님이다

좋은 benchmark bot은 빠른 코드만 줄 세우지 않는다. correctness, sanitizer, timeout, reproducibility, warmup, variance를 함께 관리해야 한다. GPU kernel은 undefined behavior가 있어도 어떤 input에서는 빠르게 "맞아 보일" 수 있다. compute sanitizer와 structured test가 필요한 이유다.

또 feedback이 너무 빈약하면 학습이 멈춘다. "wrong answer"만 주면 사용자는 어디서 indexing이 깨졌는지 모른다. 반대로 정답 code를 너무 많이 보여주면 학습 과제가 사라진다. benchmark bot은 competitive programming judge와 teaching assistant 사이의 애매한 위치에 있다.

3. 공정한 GPU benchmark는 생각보다 까다롭다

GPU benchmark는 같은 code라도 clock, thermal state, driver, stream, graph capture, input allocation, random seed에 영향을 받는다. Discord bot처럼 shared infrastructure를 쓰면 queue 상태와 이전 job의 영향도 고려해야 한다. 공정함은 단순히 같은 GPU에서 돌린다는 뜻이 아니다.

따라서 bot은 benchmark policy를 명시해야 한다. 몇 번 warmup하는지, median을 쓰는지 min을 쓰는지, correctness failure와 timeout을 어떻게 처리하는지, cuBLAS/PyTorch baseline을 어떤 조건으로 부르는지 정해야 한다. leaderboard가 재미있으려면 먼저 신뢰할 수 있어야 한다.

4. 보안과 격리도 kernel 문제다

사용자가 제출한 CUDA/C++ code를 서버에서 실행한다면 sandboxing이 필요하다. compile time abuse, file system access, long-running kernel, GPU hang, excessive memory allocation 같은 문제가 생길 수 있다. 이 강의가 실제로 다뤘는지는 확인되지 않았지만, benchmark bot 설계에서는 빼놓기 어렵다.

학습용 시스템이라도 infrastructure는 production 냄새가 난다. timeout, process isolation, GPU reset policy, log redaction, artifact cleanup이 있어야 한다. kernel 하나 재는 일도 여럿이 같이 쓰면 작은 운영체제가 된다.

5. 손으로 해볼 것

작은 benchmark challenge를 직접 설계해 보자. 예를 들어 row-wise layer norm kernel을 과제로 정하고, input shape 세 개, correctness tolerance, baseline, timeout, scoring formula를 적는다. 그 다음 사용자가 일부러 틀릴 법한 case를 세 개 만든다. non-contiguous input, odd dimension, large hidden size 같은 식이다.

원본 강의를 볼 때는 KernelBot이 이 문제들을 어떻게 풀었는지 확인하면 된다. Discord UI보다 더 중요한 것은 test와 benchmark의 신뢰성이다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_047.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_047.md

1. 지금 손에 있는 것

확실한 항목은 강의 제목, speaker 이름, YouTube URL, 그리고 public transcript/caption 수집이 실패했다는 사실이다. source bundle은 slides, code, papers, related repo folder를 찾지 못했다.

이 정도 자료로는 "강의가 어떤 playbook을 제시했다"라고 말할 수 없다. training playbook인지, serving playbook인지, cluster 운영 노하우인지도 아직 빈칸이다. reader 원고는 그 빈칸을 덮기보다 표시해야 한다.

2. playbook은 정답 목록이 아니라 절차일 수 있다

큰 GPU 시스템에서 문제는 한 층에만 머물지 않는다. kernel이 빨라도 data loader가 늦을 수 있고, GPU memory가 충분해 보여도 checkpoint I/O가 막을 수 있으며, single-node 성능이 좋아도 multi-node communication이 전체 속도를 잡을 수 있다.

그래서 playbook이라는 단어는 조심해서 읽을 가치가 있다. 좋은 playbook은 "항상 이 설정을 써라"가 아니라, 병목을 찾고, 조건을 적고, 작은 변경을 시험하고, 되돌릴 수 있게 만드는 순서에 가깝다. 원본을 볼 때는 speaker가 절차를 말하는지, 특정 stack을 말하는지 분리해서 적자.

3. 숫자는 단위와 조건으로만 읽힌다

ultra-scale 이야기에서 성능 숫자는 특히 쉽게 멋있어 보인다. 하지만 tokens/sec, samples/sec, latency percentile, GPU utilization, MFU, cost per token은 서로 다른 질문에 답한다.

원본 영상에 그래프가 나오면 값보다 먼저 분모를 적자. GPU 종류, GPU 수, interconnect, model, precision, sequence length, global batch size, warmup 여부가 빠지면 비교가 흔들린다. scale-out 결과라면 single-node baseline도 같이 필요하다.

4. 자료가 보강되면 채울 자리

강의가 training 중심이면 ZeRO류 sharding, activation checkpointing, pipeline bubble, all-reduce overlap 같은 항목이 들어올 수 있다. inference 중심이면 prefill/decode, batching, KV cache, SLO, admission control 쪽을 봐야 한다.

하지만 지금은 어느 쪽도 강의 내용으로 귀속하지 않는다. 제목과 관련 링크만으로는 범위를 고를 수 없다. 이 노트의 역할은 나중에 원본을 보며 "무엇이 실제로 언급되었고, 무엇은 독자의 배경지식이었나"를 가르는 것이다.

5. 작은 독서 준비

원본을 열기 전에 자신이 아는 training 또는 serving path를 한 장으로 그려 보자. 각 edge 옆에는 측정 가능한 지표를 붙인다. 예를 들어 all-reduce는 bandwidth와 overlap, attention은 TFLOP/s와 HBM traffic, queue는 p95/p99 latency처럼 둔다.

강의를 본 뒤에는 speaker가 실제로 다룬 층만 표시한다. 표시되지 않은 층은 "중요하지 않다"가 아니라 "이 강의에서는 아직 근거가 없다"로 남겨 둔다.

Source Links

• YouTube
• Ultra-Scale Playbook link recorded in transcript stub
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_048.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_048.md

1. 자료 한계부터 적는다

local archive에는 speaker, YouTube link, transcript failure만 있다. slides, code, repo folder는 없다. 따라서 int4, int8, quantization scheme, Apple GPU architecture, MPSGraph, Metal Shading Language 중 어떤 세부 내용을 다뤘는지는 확인 필요다.

다만 GPU Mode 맥락에서 Low Bit Metal Kernels는 분명한 위치가 있다. CUDA가 아닌 Apple GPU에서 low-bit inference kernel을 어떻게 구현하고 측정할 것인가다. CUDA 용어를 그대로 가져가면 일부는 도움이 되지만, 일부는 오해를 만든다.

2. low bit kernel의 공통 질문

low bit kernel은 storage를 줄여 memory bandwidth를 아끼려는 경우가 많다. int4나 int8 weight를 더 작은 memory footprint로 읽고, compute 직전에 dequantize하거나 scale을 적용한다. 이때 병목은 단순히 arithmetic이 아니라 unpack, scale load, layout, vectorization, memory coalescing이 된다.

CUDA든 Metal이든 질문은 비슷하다. bit-packed data를 어떤 단위로 읽을 것인가, scale/zero point를 어디에 둘 것인가, dequantize를 register에서 할 것인가, matmul이나 dot product와 어떻게 fuse할 것인가. 하지만 instruction set과 memory hierarchy가 다르므로 답은 달라질 수 있다.

3. Metal에서는 용어를 다시 확인해야 한다

CUDA의 block, warp, shared memory, Tensor Core 같은 단어를 Apple Metal에 그대로 대응시키면 위험하다. Metal에는 threadgroup, SIMD-group, threadgroup memory 같은 별도 용어와 실행 모델이 있다. Apple GPU의 matrix acceleration path도 NVIDIA Tensor Core와 같은 API/ISA로 노출되지 않을 수 있다.

그래서 이 강의를 볼 때는 CUDA 번역표를 만들기보다 Metal의 원래 vocabulary를 먼저 적는 편이 좋다. threadgroup size, memory alignment, vector type, simdgroup operation, buffer layout, pipeline state 같은 항목을 확인해야 한다.

4. benchmark는 device class를 분리해서 봐야 한다

Apple Silicon은 laptop, desktop, mobile 계열이 섞여 있다. memory bandwidth, thermal envelope, GPU core count가 다르다. "Metal kernel이 빠르다"는 말도 어떤 chip과 OS, power state, tensor shape에서 빠른지 확인해야 한다.

또 PyTorch/MPS나 Core ML 같은 상위 stack과 비교할 때 data transfer와 conversion cost를 포함했는지 봐야 한다. low-bit kernel은 unpack/dequant overhead가 작아 보여도 end-to-end에서는 format conversion이 더 클 수 있다.

5. 손으로 해볼 것

간단한 int4 weight-only matrix-vector multiplication을 설계해 보자. weight 2개를 byte 하나에 pack하고, group size 32마다 scale 하나를 둔다고 가정한다. 한 thread가 몇 byte를 읽고, 어디서 unpack하며, scale을 몇 번 재사용하는지 적어 보면 low-bit kernel의 병목이 보인다.

원본 강의를 볼 때는 speaker가 이 일반 흐름 중 어디를 Metal-specific하게 바꿨는지 확인하자. 그 지점이 rich 원고의 중심이 된다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_049.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_049.md

1. 확인된 것은 제목과 speaker뿐이다

local notes에 따르면 speaker는 Umar Jamil이고, transcript와 slide/code는 없다. 따라서 이 글은 Umar가 실제로 어떤 code를 보여줬는지, 어떤 순서로 배웠는지, 어떤 mistake를 강조했는지 재구성하지 않는다.

대신 "learning journey"라는 제목에 맞춰 독자가 따라갈 수 있는 안전한 학습 지도만 남긴다. 이 강의가 보강되면 실제 episode, code snippet, benchmark 결과로 바꿔야 한다.

2. CUDA는 실행 모델을 배우는 층이다

CUDA를 배울 때 처음 붙잡아야 하는 것은 syntax가 아니라 execution model이다. host가 kernel을 launch하고, grid와 block이 thread를 만들고, 각 thread가 자기 index를 계산해 data를 처리한다. 그 위에 register, shared memory, global memory, synchronization, occupancy가 붙는다.

FlashAttention을 이해하려면 이 기본기가 필요하다. attention은 단순히 수식이 아니라 tile을 어떻게 나누고, Q/K/V를 어떤 memory hierarchy에 두고, softmax를 online으로 유지하며, HBM traffic을 줄이는 실행 문제이기 때문이다.

3. Triton은 tile 감각을 빠르게 준다

Triton은 CUDA보다 Python에 가까운 표면을 제공하지만, 핵심은 여전히 tile이다. program_id가 어떤 output tile을 맡고, tl.arange가 tile 안 offset을 만들고, tl.load/tl.store가 mask와 함께 memory access를 표현한다.

학습 여정에서 Triton은 좋은 중간층이 된다. CUDA의 모든 thread-level detail로 바로 내려가기 전에, "내 program이 어떤 tile을 계산하는가"를 훈련할 수 있다. FlashAttention 같은 algorithm을 볼 때도 tile 단위 사고가 먼저 잡혀야 한다.

4. FlashAttention은 algorithm과 kernel의 접점이다

FlashAttention은 attention 수식을 바꾸는 것보다 memory traffic을 바꾸는 algorithm이다. QK 전체 matrix를 HBM에 만들고 softmax한 뒤 PV를 계산하는 대신, tile을 streaming하면서 online softmax를 유지한다. 덕분에 큰 intermediate를 쓰지 않고 HBM read/write를 줄인다.

이 아이디어는 CUDA와 Triton 양쪽에서 구현될 수 있다. 중요한 것은 "수학적으로 같은 attention"과 "hardware가 보는 data movement"를 동시에 보는 것이다. 학습자는 여기서 수식, tile, memory, profiler가 한 줄로 이어지는 경험을 한다.

5. 좋은 학습 순서

먼저 PyTorch attention을 쓰고 profiler로 병목을 본다. 다음으로 naive Triton matmul이나 softmax를 구현해 tile과 mask를 익힌다. 그 다음 CUDA shared memory tiled matmul을 보며 thread block과 synchronization을 익힌다. 마지막으로 FlashAttention의 online softmax와 fused kernel을 읽는다.

이 순서는 빠른 성능을 보장하는 순서가 아니라 덜 길을 잃는 순서다. CUDA를 먼저 완전히 끝내려 하면 너무 넓고, FlashAttention 논문만 읽으면 kernel의 몸이 안 보인다. 중간에 작은 benchmark를 계속 두는 것이 좋다.

6. 손으로 해볼 것

softmax(QK^T)V를 작은 행렬로 직접 계산한 뒤, Q block 하나와 K/V block 여러 개를 순서대로 보며 online softmax의 running max와 running sum이 어떻게 갱신되는지 써 보자. 이 손계산은 FlashAttention의 심장이다.

원본 강의를 볼 때는 Umar가 어떤 "막혔던 지점"을 이야기하는지 기록하자. learning journey 강의의 가치는 완성된 지식보다, 어디서 헷갈리고 어떻게 넘어갔는지에 있다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_050.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_050.md

1. 확인된 자료

local notes에는 speaker Jake Cannell, YouTube link, transcript failure만 있다. slides, code, paper, repo folder는 없다. 따라서 어떤 GPU model을 다뤘는지, gaming GPU와 data center GPU를 어떻게 비교했는지, benchmark 결과가 무엇인지는 확인 필요다.

이 source-limited 원고는 consumer GPU performance를 읽을 때의 조심스러운 틀이다. 특히 4090, 3090, laptop GPU, Apple GPU 같은 이름을 제목만 보고 끌어오지 않는다.

2. 같은 FLOP 숫자라도 같은 GPU가 아니다

consumer GPU는 peak FLOPs가 높아 보여도 data center GPU와 제약이 다르다. ECC, HBM 대 GDDR, NVLink 유무, PCIe bandwidth, power limit, cooling, sustained clock, MIG 같은 partitioning feature가 다를 수 있다. kernel이 순간적으로 빠른 것과 장시간 안정적으로 serving하는 것은 다른 문제다.

또 data center GPU에는 H100처럼 TMA, WGMMA, cluster-level feature가 강조되는 세대가 있고, consumer line은 feature exposure나 performance ratio가 다를 수 있다. 정확한 비교는 architecture와 compute capability, driver/toolchain 조건을 확인해야 한다.

3. consumer GPU에서 병목은 더 실용적이다

개인 workstation에서는 VRAM 용량이 먼저 한계가 될 수 있다. 큰 model을 올리기 위해 quantization, CPU offload, paging, smaller batch를 선택하면 kernel peak보다 memory footprint와 transfer가 중요해진다. fan curve나 thermal throttling도 긴 benchmark에서는 무시하기 어렵다.

또 multi-GPU를 하더라도 NVLink가 없거나 bandwidth가 제한될 수 있다. 이 경우 tensor parallelism이나 distributed inference에서 communication cost가 data center node와 다르게 보인다. "내 4090 두 장"과 "H100 NVLink node"는 숫자만 GPU 두 개로 같을 뿐, 시스템은 다르다.

4. benchmark는 생활 조건을 포함한다

consumer GPU benchmark는 실사용 조건을 함께 적어야 한다. case open/closed, power limit, driver version, background display load, OS, batch size, quantization format이 결과에 영향을 줄 수 있다. data center benchmark보다 더 "집안 사정"을 탄다.

그렇다고 consumer GPU가 덜 중요하다는 뜻은 아니다. 오히려 많은 개발자가 이 환경에서 kernel을 배우고 prototype을 만든다. 다만 reader는 data center GPU의 optimization recipe가 그대로 옮겨진다고 가정하지 않아야 한다.

5. 손으로 해볼 것

자신의 GPU 또는 흔한 consumer GPU 하나를 골라 spec sheet를 적어 보자. VRAM 용량, memory bandwidth, PCIe generation/lane, Tensor Core 지원 dtype, compute capability, power limit을 적고, H100 같은 data center GPU와 차이를 표로 만든다.

원본 강의를 볼 때는 speaker가 어떤 차이를 가장 중요하게 봤는지 표시하자. 이 source-limited 원고는 그 표를 중심으로 다시 써야 한다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_051.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_051.md

1. 자료 한계

local notes에는 speaker Tanishq Kumar, YouTube link, transcript failure만 있다. slides, code, paper link는 없다. 제목은 low precision의 scaling laws를 말하지만, 어떤 model size, dataset, precision format, metric을 다뤘는지는 확인되지 않았다.

그러므로 이 원고는 강의 내용의 요약이 아니다. 원본을 다시 볼 때 어떤 축을 기록해야 하는지 알려주는 source-limited 원고다. scaling law는 숫자와 조건이 생명이라, 근거 없는 매끈한 설명을 넣으면 오히려 해롭다.

2. low precision은 단순한 압축이 아니다

precision을 낮추면 memory footprint와 bandwidth가 줄고, hardware가 지원하면 throughput도 올라간다. FP16, BF16, FP8, INT8, INT4 같은 format은 training과 inference에서 각각 다른 tradeoff를 가진다. 하지만 낮은 bit가 항상 이득은 아니다. scaling factor, outlier, accumulation dtype, optimizer state, activation distribution이 함께 움직인다.

특히 training에서는 stability가 중요하다. inference quantization에서는 calibration과 quality drop을 본다면, training low precision에서는 gradient, update, loss scaling, convergence가 들어온다. 제목의 scaling laws가 어느 쪽을 말하는지는 확인 필요다.

3. scaling law라면 반드시 조건표가 필요하다

scaling law는 "모델이 커지면 어떻게 된다"는 이야기처럼 들리지만, 실제로는 실험 조건의 묶음이다. parameter count, token count, dataset, architecture, optimizer, learning rate schedule, precision policy, loss metric을 함께 봐야 한다.

low precision scaling law라면 여기에 quantization granularity, scale format, clipping, stochastic rounding, accumulation precision, hardware target도 붙는다. 어느 하나가 빠지면 같은 결론을 재현하기 어렵다.

4. kernel reader에게 중요한 연결점

GPU Mode 독자에게 low precision은 model quality만의 주제가 아니다. dtype이 바뀌면 kernel도 바뀐다. int8 Tensor Core matmul, FP8 GEMM, low-bit dequant fusion, KV cache quantization, bandwidth-bound attention 같은 문제가 바로 연결된다.

하지만 kernel throughput만 보고 precision policy를 정할 수는 없다. 더 빠른 dtype이 quality를 크게 떨어뜨리면 end-to-end cost가 오히려 나빠질 수 있다. 반대로 아주 작은 quality drop으로 memory capacity가 열리면 serving economics가 크게 좋아질 수 있다. scaling law는 이 균형을 정량화하려는 시도다.

5. 손으로 해볼 것

작은 LLM inference scenario를 하나 만들고, FP16 weight와 INT8 weight-only quantization의 memory footprint를 비교해 보자. weight memory, KV cache memory, activation memory를 나눠 적으면 "precision을 낮췄는데도 어디가 여전히 큰가"가 보인다.

원본 강의를 볼 때는 모든 graph에 대해 x축, y축, model family, precision format, baseline, metric을 적자. 특히 "law"라고 부르는 식이 있다면 fit range와 residual을 확인해야 한다. 예쁜 직선은 늘 의심하고, 그 의심은 꽤 건강하다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_052.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_052.md

1. Q&A는 모르는 곳을 드러내는 형식이다

torch.compile Q&A라는 제목은 강의가 정리된 튜토리얼보다 질문과 답변에 가까웠을 가능성을 말해 준다. 하지만 transcript가 없으므로 어떤 질문이 나왔는지, Richard Zou가 어떤 사례를 들었는지는 확인되지 않는다.

그래서 이 reader 원고는 torch.compile의 공식 동작을 자세히 설명하는 글이 아니다. 원본을 다시 볼 때 적어야 할 질문의 틀이다. 어떤 model에서 compile을 켰는가, compile time과 runtime을 어떻게 나눴는가, graph break는 어디서 났는가, speedup은 어떤 shape에서 나온 것인가. 이런 질문을 챙겨야 강의를 다시 볼 때 손에 잡힌다.

2. torch.compile은 Python을 곧바로 CUDA로 바꾸지 않는다

초보자가 자주 기대하는 그림은 단순하다. PyTorch code를 넣으면 compiler가 알아서 최적 CUDA kernel을 내보낸다. 실제 시스템은 그보다 복잡하다. Python-level program을 추적하고, tensor shape와 dtype에 대한 가정을 세우고, graph를 만들고, backend가 실행 가능한 형태로 낮춘다.

이 과정에서 성능 이득은 여러 층에서 올 수 있다. Python overhead가 줄 수도 있고, operator가 fuse될 수도 있고, backend가 더 나은 kernel을 고를 수도 있다. 반대로 dynamic control flow, unsupported op, shape 변화, side effect가 있으면 graph break나 recompilation 비용이 생길 수 있다.

3. speedup 숫자보다 조건이 먼저다

torch.compile 성능을 읽을 때는 "몇 배 빨라졌다"보다 조건표가 중요하다. model, batch size, sequence length, dtype, device, PyTorch version, backend, warmup 횟수, compile time 포함 여부가 없으면 숫자는 금방 미끄러진다.

특히 Q&A 강의라면 질문자가 겪은 문제의 context가 답의 절반이다. training인지 inference인지, single GPU인지 multi-GPU인지, custom op가 있는지, dynamic shape가 많은지에 따라 좋은 답이 달라진다. 좋은 reader는 답변만 모으지 않고 질문의 전제도 같이 적는다.

4. graph break는 실패가 아니라 진단 신호다

compile이 잘 안 될 때 graph break를 단순한 에러로만 보면 답답하다. 하지만 graph break는 compiler가 어디까지 프로그램을 이해했는지 알려 주는 신호다. Python object, data-dependent branch, unsupported operation, mutation 같은 부분이 드러난다.

원본 강의를 볼 때는 "이런 코드를 쓰면 깨진다"보다 "왜 깨졌는가"를 적는 편이 더 좋다. 그래야 같은 문제를 다른 model에서 만났을 때, 코드를 어떻게 바꾸거나 compile boundary를 어디에 둘지 판단할 수 있다.

5. 손으로 해볼 것

작은 PyTorch module 하나를 골라 eager와 torch.compile을 비교해 보자. 먼저 warmup을 분리하고, compile time을 뺀 steady-state latency를 재고, shape를 하나만 쓸 때와 여러 shape를 섞을 때를 나눠 본다.

그 다음 graph break 로그를 켜고 어떤 부분이 compile되지 않는지 확인한다. 이 작은 실험은 Q&A를 읽는 눈을 만들어 준다. 답을 외우는 것보다, 내 코드가 compiler에게 어떻게 보이는지 보는 쪽이 훨씬 오래 간다.

Source Links

• YouTube
• Slides / notes document
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_053.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_053.md

1. 제목만으로 algorithm을 만들지 않는다

확인된 정보는 제목, YouTube link, transcript failure뿐이다. PPO, DQN, policy gradient, vectorized environment, JAX, PyTorch 중 무엇을 썼는지는 notes에서 확인되지 않는다. speaker도 local notes에는 Missing으로 남아 있다.

그래서 이 원고는 LeanRL의 내부 구조를 설명하지 않는다. "Small RL Models at the Speed of Light"라는 제목을 만났을 때, 강의를 보며 어떤 장부를 열어야 하는지 정리한다. 작은 model을 빠르게 학습시킬 때 병목이 어디로 이동하는가, 그리고 GPU를 쓰는 일이 항상 답인가를 묻는 글이다.

2. 작은 RL은 GPU보다 loop가 먼저 보일 때가 많다

RL training은 supervised learning처럼 큰 batch tensor 하나를 계속 밀어 넣는 모양이 아닐 수 있다. environment를 step하고, observation을 모으고, action을 샘플링하고, reward와 done을 처리한 뒤 update한다. model이 작으면 neural network forward보다 이 control loop가 더 비싸질 수 있다.

따라서 "빠른 RL"은 kernel optimization만의 문제가 아니다. environment vectorization, rollout buffer layout, CPU-GPU transfer, random number generation, logging, Python dispatch를 같이 봐야 한다. 작은 모델일수록 사소한 overhead가 주인공이 된다.

3. "speed of light"는 상한선을 묻는 말로 읽자

제목의 표현은 과감하지만, reader에서는 과장 대신 상한선을 묻는 말로 읽는 편이 좋다. 작은 RL workload에서 이론적으로 얼마나 빨라질 수 있는가. environment가 초당 몇 step을 낼 수 있는가. learner는 그 data를 얼마나 빨리 소비하는가. 둘 중 어느 쪽이 기다리는가.

이렇게 보면 benchmark도 더 정직해진다. frames per second, samples per second, episode throughput, wall-clock-to-reward 같은 지표는 서로 다르다. 강의를 다시 볼 때 어떤 지표를 speed라고 부르는지 먼저 적어야 한다.

4. LeanRL을 볼 때 확인할 것

자료가 보강되면 가장 먼저 LeanRL이 무엇을 단순화했는지 확인해야 한다. codebase 구조인지, algorithm 구현인지, vectorized rollout인지, compiler/JIT 경로인지, 아니면 benchmark methodology인지에 따라 원고의 중심이 달라진다.

또 작은 RL 모델이라고 해도 "작다"의 의미가 여러 가지다. parameter 수가 작은가, environment가 가벼운가, observation이 작은가, horizon이 짧은가, batch가 작은가. 강의의 실제 예제가 무엇인지 알기 전까지는 결론을 비워 두는 편이 안전하다.

5. 손으로 해볼 것

가벼운 toy environment를 하나 잡고, random policy로 environment step throughput만 먼저 재 보자. 그 다음 작은 MLP policy forward를 붙이고, 마지막으로 learner update를 붙인다. 각 단계가 초당 몇 sample을 처리하는지 따로 적는다.

이 실험은 LeanRL 원본을 볼 때 좋은 기준점이 된다. speaker가 어디를 줄였는지, 그리고 그 줄임이 algorithmic change인지 systems change인지 구분할 수 있다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_054.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_054.md

1. unified라는 말은 늘 양쪽으로 당긴다

unified device accelerator language는 매력적인 목표다. CPU, GPU, 다른 accelerator를 하나의 언어 모델로 표현하면 code reuse와 portability가 좋아질 수 있다. 하지만 accelerator 성능은 대개 device-specific detail에서 나온다.

transcript가 없어 Modular가 어떤 해법을 제시했는지는 확인되지 않는다. 이 reader는 따라서 Modular의 구체 syntax나 compiler path를 설명하지 않는다. 대신 unified language를 읽을 때 놓치기 쉬운 tradeoff를 정리한다.

2. portability는 공짜가 아니다

accelerator language가 portable하려면 memory hierarchy, parallelism, synchronization, vectorization, layout 같은 차이를 표현할 방법이 필요하다. 너무 추상화하면 성능을 잃고, 너무 세부적이면 unified라는 장점이 약해진다.

CUDA는 NVIDIA GPU에 깊이 붙어 있고, Triton은 tile-level tensor program을 Python스럽게 표현한다. Modular의 언어가 이 사이 어디에 놓이는지는 원본 강의에서 확인해야 한다. 어떤 층까지 사용자가 직접 말하고, 어떤 층부터 compiler가 맡는지가 핵심이다.

3. 언어 설계는 compiler 약속이다

device language는 문법만의 문제가 아니다. 사용자가 어떤 정보를 compiler에게 주고, compiler가 어떤 최적화를 약속하는지의 계약이다. shape가 static인지 dynamic인지, aliasing을 어떻게 다루는지, memory placement를 어떻게 표시하는지에 따라 가능한 최적화가 달라진다.

unified language라면 더 많은 계약이 필요하다. 한 target에서는 자연스러운 layout이 다른 target에서는 느릴 수 있다. compiler가 target별 rewrite를 해도 semantic은 유지되어야 한다. 이 지점에서 type system, IR, scheduling abstraction이 중요해진다.

4. 성능 주장은 baseline을 봐야 한다

이런 강의에서 benchmark가 나온다면 반드시 baseline을 확인해야 한다. hand-written CUDA와 비교하는가, vendor library와 비교하는가, PyTorch eager와 비교하는가, Triton과 비교하는가. 비교 대상에 따라 결론이 완전히 달라진다.

또 "unified"의 성공은 한 device에서 빠른 것만으로는 부족하다. 여러 target에서 같은 source가 얼마나 적게 바뀌는지, target-specific tuning hook은 얼마나 필요한지, debug와 profiling이 가능한지도 함께 봐야 한다.

5. 손으로 해볼 것

간단한 reduction kernel을 상상해 보자. CUDA에서는 thread block과 shared memory를 어떻게 쓸지, Triton에서는 block size와 mask를 어떻게 둘지, high-level tensor API에서는 어떤 op로 표현할지 각각 적어 본다.

그 다음 unified language가 이 세 표현 중 어디에 가까워야 할지 생각해 보자. 성능을 얻으려면 어떤 정보가 반드시 드러나야 하는지도 같이 적는다. 원본 강의를 볼 때 이 목록이 좋은 체크리스트가 된다.

Source Links

• YouTube
• Slides
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_055.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_055.md

1. kernel benchmark는 작은 거짓말을 잘 만든다

kernel benchmark는 짧고 숫자가 선명해서 믿기 쉽다. 하지만 warmup을 빼먹거나, synchronization을 안 하거나, tensor shape를 편한 것만 고르거나, memory allocation을 같이 재면 숫자는 쉽게 흔들린다. GPU는 비동기 실행이라 wall-clock timing 하나로는 특히 조심해야 한다.

local archive에는 이 강의의 transcript가 없으므로 구체 사례는 만들지 않는다. 대신 benchmarking tale을 읽을 때 필요한 질문을 남긴다. 무엇을 재는가, 어디서 synchronize하는가, baseline은 무엇인가, cache와 allocator 상태는 어떻게 통제하는가.

2. 올바른 측정은 kernel 밖에서 시작한다

좋은 benchmark는 kernel code만 보지 않는다. input tensor 생성, host-to-device transfer, allocation, stream, synchronization, warmup, repeat count를 모두 통제한다. kernel launch overhead를 포함할지 뺄지도 명확해야 한다.

PyTorch와 비교할 때도 같은 연산인지 확인해야 한다. broadcasting, dtype promotion, layout conversion, contiguous call이 숨어 있으면 baseline이 다른 일을 하고 있을 수 있다. "내 kernel이 빠르다" 전에 "둘이 같은 일을 한다"가 먼저다.

3. 평균 하나로는 부족하다

kernel timing은 평균만 보면 놓치는 것이 많다. p50, p95, variance, outlier를 같이 봐야 한다. 특히 짧은 kernel은 launch overhead와 scheduler noise가 커서 작은 차이를 과장하기 쉽다.

또 benchmark는 shape 하나로 끝나지 않는다. block size가 잘 맞는 shape에서는 빨라도, ragged shape나 non-contiguous input에서는 느릴 수 있다. production workload에 가까운 shape distribution을 따로 모아 두면 benchmark가 훨씬 덜 장식적이 된다.

4. profiler가 숫자에 이유를 붙인다

time만으로는 병목을 알 수 없다. memory bandwidth가 한계인지, occupancy가 낮은지, register pressure가 큰지, instruction mix가 나쁜지, uncoalesced access가 있는지는 profiler counter를 봐야 한다.

Nsight Compute나 PyTorch profiler 같은 도구는 benchmark의 뒷면을 보여 준다. 단, profiler도 overhead와 sampling/collection 조건이 있으므로 timing 숫자와 profiler run을 구분해서 해석하는 것이 좋다.

5. 손으로 해볼 것

elementwise add kernel을 하나 직접 만들고 PyTorch baseline과 비교해 보자. 첫 번째 버전은 일부러 synchronization을 빼고 재고, 두 번째 버전은 CUDA event나 명시적 sync를 넣어 잰다. 숫자가 어떻게 달라지는지 보는 것만으로도 benchmarking 감각이 확 달라진다.

그 다음 shape를 크게, 작게, non-contiguous하게 바꿔 본다. 빠른 kernel을 만드는 일만큼이나, 빠르다고 말할 수 있는 조건을 쓰는 일이 중요하다는 걸 몸으로 알게 된다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_056.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_056.md

1. CuTe는 "예쁜 tensor wrapper"보다 더 낮은 층에 있다

CuTe는 CUTLASS 생태계에서 layout과 tensor abstraction을 다루는 중요한 구성요소로 알려져 있다. 하지만 이 노트는 강의의 실제 예제나 slide 순서를 확인하지 못했으므로, 특정 API 사용법을 단정하지 않는다.

GPU Mode 독자에게 CuTe가 중요한 이유는 matmul 같은 kernel이 결국 layout 문제이기 때문이다. data가 논리적으로는 matrix여도, hardware가 보는 것은 address, stride, tile, thread mapping이다. CuTe는 이 관계를 더 수학적이고 조합 가능한 방식으로 표현하려는 도구로 읽을 수 있다.

2. layout은 모양이 아니라 주소 함수다

초보자는 tensor shape를 먼저 본다. 하지만 kernel 작성자는 shape와 함께 layout을 봐야 한다. 같은 (M, N) matrix라도 row-major, column-major, tiled layout, swizzled layout은 memory access와 coalescing이 다르다.

CuTe를 읽을 때 좋은 관점은 layout을 "index를 address로 바꾸는 함수"로 보는 것이다. tile을 나누고, thread가 맡을 element를 정하고, shared memory에 옮길 때도 결국 이 함수들이 합성된다. 이름은 귀엽지만, 내용은 꽤 날카롭다.

3. tiling은 performance와 correctness를 같이 건드린다

GEMM kernel에서 tile 크기는 Tensor Core utilization, shared memory 사용량, register pressure, occupancy, memory coalescing에 모두 영향을 준다. 좋은 tile은 계산을 많이 재사용하게 하지만, 너무 크면 resource를 많이 먹는다.

CuTe 같은 abstraction은 tile과 layout을 명시적으로 다루게 해 준다. 이것은 편의 기능이 아니라 안전장치이기도 하다. 복잡한 kernel에서 좌표 변환을 손으로 흩뿌리면 작은 off-by-one과 stride 착각이 성능과 correctness를 동시에 깨뜨린다.

4. slide를 볼 때는 예제 하나를 끝까지 따라가자

원본 slide가 보강된 자료로 확인되므로, rich 원고로 바꿀 때는 예제 하나를 처음부터 끝까지 따라가는 것이 좋다. 예를 들어 작은 matrix tile이 global memory에서 shared memory로, 다시 MMA operand로, 마지막 output tile로 가는 경로를 그리면 CuTe의 의미가 훨씬 선명해진다.

API 이름을 많이 나열하는 것보다 좌표가 어떻게 바뀌는지 보여 주는 편이 reader에 더 좋다. CuTe는 결국 "이 좌표가 어디에 놓이는가"를 놓치지 않게 하는 언어이기 때문이다.

5. 손으로 해볼 것

작은 4 x 4 matrix를 2 x 2 tile 네 개로 나눠 보자. 각 element (i, j)가 row-major memory offset i * 4 + j로 가는 표를 만들고, tile 좌표와 tile 내부 좌표로 다시 표현해 본다.

그 다음 shared memory tile에 transpose해서 넣는다고 가정하고 offset 표를 하나 더 만든다. 이 손표가 CuTe layout algebra를 읽는 가장 소박한 입구다. 기계가 하는 일을 사람이 한 번 손으로 해 보는 것, 묘하게 강력하다.

Source Links

• YouTube
• GPU Mode lecture folder
• Slides PDF
• Slides mirror
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_057.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_057.md

1. disaggregated라는 말의 범위를 먼저 묻자

Disaggregated LLM Inference는 꽤 넓은 제목이다. prefill과 decode 분리일 수도 있고, compute와 memory 자원 분리일 수도 있으며, serving cluster에서 역할별 worker를 나누는 이야기일 수도 있다. local notes에는 세부 자료가 없으므로 어느 쪽이라고 단정하지 않는다.

다만 LLM inference에서 disaggregation이 자주 등장하는 이유는 있다. prompt를 처리하는 prefill은 큰 matrix 연산에 가깝고, token을 하나씩 생성하는 decode는 KV cache를 계속 읽는 latency-sensitive 작업에 가깝다. 두 단계는 같은 model을 쓰지만 병목과 좋은 batch 모양이 다르다. 강의가 이 구도를 어떻게 썼는지는 원본에서 다시 확인해야 한다.

2. prefill과 decode는 다른 생물이다

prefill은 입력 prompt 전체를 한 번에 처리한다. sequence length가 길면 attention과 GEMM 작업량이 크고, GPU를 꽤 잘 채울 수 있다. throughput-oriented batching이 잘 맞는 경우가 많다.

decode는 매 step 새 token 하나를 만들며 반복된다. batch가 커져도 per-token latency와 KV cache memory traffic이 중요해진다. request마다 남은 길이가 다르고, early stopping도 있다. 그래서 scheduling과 cache placement가 성능을 크게 흔든다.

3. 분리하면 새 비용도 생긴다

disaggregation은 공짜가 아니다. prefill과 decode를 다른 worker나 device로 나누면 KV cache 이동, queueing, admission control, fault handling, backpressure 같은 비용이 생긴다. 특히 KV cache는 작지 않다. 잘못 옮기면 계산을 줄이려다 network와 memory traffic을 키울 수 있다.

또 분리된 system은 debugging이 어려워진다. 하나의 request가 여러 worker를 지나가므로 trace와 metric이 필요하다. p95 latency가 나빠졌을 때 prefill queue인지, cache transfer인지, decode batcher인지 찾을 수 있어야 한다.

4. benchmark는 workload mix가 생명이다

LLM serving benchmark는 prompt length, output length, arrival rate, batching policy, model size, quantization, GPU type이 함께 있어야 읽힌다. disaggregated inference에서는 여기에 worker ratio와 network topology까지 붙는다.

원본 강의를 볼 때는 어떤 traffic assumption을 썼는지 먼저 적자. 짧은 prompt가 많은 chat workload와 긴 문서 summarization workload는 전혀 다른 system을 요구할 수 있다. 같은 architecture가 한쪽에서는 멋지고 다른 쪽에서는 과할 수 있다.

5. 손으로 해볼 것

가상의 serving workload를 하나 만든다. prompt length 평균, output length 평균, request arrival rate를 정하고, prefill time과 decode token time을 아주 단순한 식으로 둔다. 그 다음 prefill worker와 decode worker 수를 바꾸며 어디에 queue가 쌓이는지 계산해 보자.

숫자는 거칠어도 괜찮다. 이 작은 queueing model은 disaggregation이 왜 system design 문제인지 보여 준다. kernel 하나가 빨라도 queue가 막히면 사용자는 느리다고 느낀다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_058.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_058.md

1. 이름을 설명으로 바꾸지 말자

FastVideo가 논문 이름인지, library인지, 내부 시스템인지, 단순한 강의 제목인지는 source note만으로 판별되지 않는다. 그러니 diffusion, transformer, VAE, scheduler, attention trick 같은 익숙한 말을 강의 내용처럼 붙이면 안 된다.

지금 할 수 있는 일은 작다. 비디오 모델을 읽을 때 필요한 비용 축을 미리 세워 두고, 원본에서 확인된 항목만 나중에 채운다. source-thin draft는 화려한 요약보다 좋은 빈 표에 가깝다.

2. video workload는 차원이 하나 더 늘어난다

이미지 모델은 height와 width를 다룬다. video는 여기에 frame 축이 붙는다. 그래서 tensor shape 하나도 (batch, frames, channels, height, width)처럼 읽어야 한다.

frame 수가 늘면 activation memory, attention 대상, decode 또는 sampling 반복의 총량이 같이 움직일 수 있다. 단, 이 강의가 정확히 어떤 video model을 대상으로 했는지는 아직 알 수 없다. 아래 내용은 강의 요약이 아니라 원본을 읽기 위한 계산 감각이다.

3. fast의 의미를 먼저 고른다

"빠르다"는 말은 하나가 아니다. training throughput이 빨라진 것일 수 있고, inference latency가 줄어든 것일 수 있다. 같은 품질에 더 적은 step을 쓰는 algorithmic change일 수도 있고, 같은 연산을 더 잘 배치한 systems work일 수도 있다.

원본을 볼 때 첫 표시점은 여기다. speaker가 무엇을 빠르게 만들었는가. wall-clock인지, GPU time인지, cost인지, memory footprint인지, 품질 대비 속도인지가 정해져야 나머지 주장을 읽을 수 있다.

4. 품질을 빼면 video 성능은 반쪽이다

video generation 계열이라면 속도와 품질을 같이 봐야 한다. temporal consistency, motion artifact, prompt alignment, flicker 같은 문제가 성능 숫자 뒤에 숨어 있을 수 있다.

강의가 다른 종류의 video processing을 다뤘다면 품질 지표도 달라진다. encoding, retrieval, understanding, generation은 모두 다른 metric을 쓴다. 따라서 이 노트는 특정 metric을 못 박지 않고, "비교가 같은 출력 품질을 유지하는가"라는 질문만 남긴다.

5. 손계산으로 준비하기

작은 표를 만들자. 행은 frames, resolution, batch, dtype, step count로 두고, 열은 memory, compute, latency, quality risk로 둔다. 각 칸에는 수식이 아니라 "늘어남", "줄어듦", "측정해야 함" 정도만 적어도 충분하다.

원본을 볼 때 FastVideo가 이 표의 어느 행을 건드리는지 표시한다. 아무 것도 표시할 수 없다면, 이 노트를 더 길게 쓰지 말고 caption이나 slide를 다시 구하는 것이 먼저다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_059.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_059.md

1. linear attention은 하나의 algorithm 이름이 아니다

linear attention은 넓은 가족 이름이다. kernel feature map을 쓰는 방식, recurrent state를 쓰는 방식, chunking을 쓰는 방식, causal/non-causal 형태가 모두 다를 수 있다. notes에는 구체 variant가 없으므로 어떤 수식도 원본 강의의 내용처럼 쓰면 안 된다.

그래도 공통 질문은 있다. quadratic attention이 sequence length에 대해 부담스러우니, K/V 전체와 Q의 모든 pair를 직접 만들지 않는 길을 찾는다. 하지만 complexity를 낮춘 뒤에도 GPU에서 빠르게 실행되는지는 별개의 문제다.

2. Big-O가 줄어도 kernel은 삐질 수 있다

quadratic attention은 큰 intermediate와 HBM traffic 때문에 비싸다. linear attention은 이론적으로 더 좋은 scaling을 약속할 수 있다. 그러나 GPU는 큰 dense matmul을 아주 잘한다. 반대로 작은 scan, recurrent update, irregular memory access는 hardware를 덜 채울 수 있다.

따라서 optimization의 질문은 "O(n)이니까 빠르다"가 아니다. 어느 sequence length부터 이득인가, batch와 head dimension이 어느 정도인가, state를 register/shared/global memory 어디에 둘 수 있는가, backward pass는 어떻게 되는가를 봐야 한다.

3. causal 구조는 병렬화를 어렵게 한다

sequence를 앞에서 뒤로 누적하는 causal linear attention은 recurrent한 냄새가 난다. 각 token이 이전 state에 의존하면, 단순히 모든 position을 독립적으로 계산하기 어렵다. parallel scan이나 chunking이 필요할 수 있다.

이 부분이 GPU 최적화의 재미있는 지점이다. algorithm은 순차적으로 보이지만, chunk 안에서는 병렬화하고 chunk 사이에는 prefix를 전달하는 식으로 구조를 바꿀 수 있다. 원본 강의가 어떤 전략을 택했는지는 확인 필요다.

4. benchmark는 FlashAttention과 다르게 읽어야 한다

linear attention을 standard attention 또는 FlashAttention과 비교할 때는 quality와 speed를 함께 봐야 한다. 같은 model architecture인지, retraining이 필요한지, perplexity나 downstream metric이 어떻게 바뀌는지 없으면 성능 비교가 반쪽이다.

또 짧은 sequence에서는 quadratic attention이 이미 충분히 빠를 수 있다. linear attention의 이득은 긴 sequence, memory pressure, streaming decode 같은 조건에서 더 잘 드러날 수 있다. 그래서 shape sweep이 중요하다.

5. 손으로 해볼 것

길이 8짜리 toy sequence를 만들고, 일반 attention이 만드는 8 x 8 score matrix의 element 수를 세어 보자. 그 다음 linear attention variant가 있다고 가정하고 per-token state 크기를 d x d 또는 d로 놓았을 때 필요한 저장량을 비교해 본다.

정확한 variant는 원본 강의에서 확인해야 한다. 하지만 이 작은 비교만으로도 "quadratic matrix를 없앤 뒤 무엇을 새로 들고 다니는가"라는 질문이 생긴다. 그 질문이 linear attention 최적화의 입구다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_060.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_060.md

1. 이름에서 출발하되 설명하지 않는다

D-Matrix Corsair가 hardware인지, software stack인지, compiler/runtime인지, benchmark 대상인지 이 자료만으로는 알 수 없다. 제목은 단서지만 증거는 아니다.

그래서 이 원고는 Corsair의 구조나 성능을 설명하지 않는다. GPU Mode reader에 필요한 것은 그럴듯한 배경 설명이 아니라, 원본을 볼 때 놓치면 안 되는 질문의 순서다.

2. 먼저 범위를 적는다

첫 번째 질문은 "무엇을 위한 시스템인가"다. 대상 workload가 정해지지 않으면 성능 주장도 의미가 흐려진다. dense linear algebra인지, sparse pattern인지, attention류 연산인지, 더 넓은 end-to-end system path인지는 원본에서 직접 확인해야 한다.

두 번째 질문은 "사용자가 만지는 인터페이스가 무엇인가"다. kernel을 직접 쓰는가, library call을 쓰는가, graph compiler에 맡기는가. 이 경계가 정해져야 optimization 이야기가 어느 층의 이야기인지 보인다.

3. kernel이라는 단어를 과하게 믿지 않는다

notes에서 잡힌 kernel 신호는 도움이 되지만 충분하지 않다. GPU에서 kernel은 launch되는 device code를 뜻하는 경우가 많다. 다른 stack에서는 kernel이 더 큰 graph primitive나 library 내부 단위를 가리킬 수도 있다.

원본을 볼 때는 "kernel"이 누구의 단어인지 적자. speaker가 직접 말한 용어인지, 자동 추출 결과인지, reader가 가져온 배경지식인지 분리해야 한다. 이 구분이 없으면 source-thin draft가 순식간에 제품 소개문처럼 변한다.

4. 성능 주장은 조건표 없이는 보류한다

throughput, latency, efficiency 같은 단어가 나오면 곧바로 숫자를 믿지 말고 조건표를 채운다. 비교 대상, input shape, dtype, batch, sequence 또는 tensor shape, warmup, end-to-end 포함 범위가 필요하다.

특히 낯선 시스템은 boundary가 중요하다. host preprocessing, data transfer, compilation time, queueing이 포함되는지 빠지는지에 따라 같은 숫자도 다른 이야기가 된다. 이 강의의 실제 boundary는 아직 이 파일 안에 없다.

5. 원본을 볼 때 쓸 빈 표

다음 열을 가진 표를 준비한다: identity, target workload, programming interface, memory model, execution unit, benchmark baseline, unsupported or unclear items. 강의를 보며 speaker가 말한 것만 채운다.

빈칸은 실패가 아니다. source가 얇은 상태에서 빈칸을 보존하는 것이, 출처 없는 추론으로 독자를 안심시키는 것보다 훨씬 낫다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_061.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_061.md

1. scheduling language는 "어떻게 실행할까"를 말한다

컴파일러를 처음 배울 때는 보통 program semantics에 집중한다. 무엇을 계산하는가. 하지만 고성능 kernel에서는 어떻게 계산하는가가 거의 같은 무게를 가진다. loop를 어떻게 나누고, memory를 어디에 두고, parallelism을 어떻게 배치하는지가 성능을 만든다.

Exo 계열의 핵심 관심사는 이런 scheduling 선택을 언어로 다루는 데 있다고 볼 수 있다. 단, local notes에는 구체 API나 예제가 없으므로 Exo 2의 실제 설계를 단정하지 않는다.

2. schedule은 optimization log이기도 하다

고성능 코드를 손으로 고치다 보면 algorithm과 optimization이 뒤섞인다. loop tiling, unrolling, vectorization, memory staging이 원래 계산식 사이에 스며들어, 나중에는 무엇이 의미이고 무엇이 성능 trick인지 구분하기 어렵다.

scheduling language는 이 둘을 떼어 놓으려는 방향이다. 먼저 계산을 쓰고, 그 계산을 target에 맞게 바꾸는 transformation을 별도로 기록한다. 그러면 최적화 과정을 읽고, 재적용하고, 검증하기 쉬워진다.

3. "growing"은 기능 추가보다 언어 확장의 문제다

제목의 Growing a scheduling language는 단순히 feature를 더한다는 말보다 넓게 읽힌다. schedule language가 실제 workload를 만나면 새로운 transformation, 새로운 memory pattern, 새로운 hardware constraint를 표현해야 한다. 언어가 자라야 하는 이유다.

하지만 언어가 커질수록 위험도 있다. primitive가 많아지면 배우기 어렵고, transformation끼리 상호작용이 복잡해진다. 좋은 scheduling language는 표현력과 reasoning 가능성 사이에서 균형을 잡아야 한다.

4. GPU 독자에게 중요한 연결점

CUDA나 Triton을 직접 쓰는 사람에게 scheduling language는 남의 일이 아니다. block tiling, shared memory staging, warp-level mapping, vectorized load는 모두 schedule 선택이다. 다만 CUDA에서는 이 선택이 code 구조 안에 박혀 있고, Exo 같은 접근에서는 transformation으로 드러날 수 있다.

원본 강의를 볼 때는 어떤 target을 예로 들었는지 확인해야 한다. CPU SIMD인지, embedded accelerator인지, GPU인지에 따라 scheduling primitive의 맛이 달라진다. GPU Mode reader라면 GPU와 연결되는 지점을 따로 표시해 두면 좋다.

5. 손으로 해볼 것

간단한 matrix multiply loop를 세 겹으로 써 보자. 그 다음 algorithm은 그대로 두고, i, j, k loop를 tile한다고 말로 적어 본다. tile 크기를 바꾸면 memory reuse와 parallelism이 어떻게 달라지는지 짧게 설명한다.

이 연습의 목적은 빠른 matmul을 직접 완성하는 것이 아니다. 계산과 schedule을 분리해서 말하는 습관을 만드는 것이다. 그 습관이 Exo 2 같은 강의를 훨씬 읽기 쉽게 만든다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_062.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_062.md

1. search-based라는 말은 후보 공간을 전제로 한다

compiler가 search를 한다는 것은 아무 곳이나 뒤진다는 뜻이 아니다. 먼저 탐색할 공간이 있어야 한다. tile size, loop order, fusion boundary, memory placement, vector width, parallel mapping 같은 선택지를 어떤 형태로 표현할지 정해야 한다.

local notes에는 구체 compiler 이름이나 예제가 없다. 따라서 이 원고는 AutoTVM, Ansor, Triton autotune 같은 특정 시스템의 설명이 아니다. 원본을 볼 때 어떤 search space를 만들고 어떤 기준으로 후보를 고르는지 확인하기 위한 가이드다.

2. rule-based와 search-based는 적이 아니다

고성능 compiler에는 사람이 만든 rule이 필요하다. 말도 안 되는 schedule을 줄이고, hardware 제약을 지키고, correctness를 보장해야 한다. 하지만 rule만으로 모든 shape와 hardware, operator variant를 덮기는 어렵다.

search-based 접근은 이 빈틈을 메운다. 사람이 좋은 후보 공간을 설계하고, compiler가 그 안에서 측정하거나 예측한다. 중요한 건 search가 magic이 아니라는 점이다. 공간을 잘못 만들면 아무리 열심히 찾아도 좋은 답이 없다.

3. cost model은 편하지만 거짓말할 수 있다

모든 후보를 실제 GPU에서 benchmark하면 시간이 오래 걸린다. 그래서 cost model이나 learned model이 후보의 성능을 예측할 수 있다. 하지만 GPU 성능은 cache, occupancy, memory coalescing, compiler lowering, clock state 같은 요인에 민감하다.

좋은 search-based compiler는 예측과 실제 측정을 어떻게 섞는지가 중요하다. cost model을 신뢰하되 blind faith로 두지 않고, calibration과 validation을 계속 해야 한다. 원본 강의에 benchmark가 있다면 prediction error도 함께 보고 싶다.

4. 탐색 비용도 시스템 비용이다

autotuning은 결과 kernel이 빠르면 멋져 보이지만, tuning time도 실제 비용이다. shape가 자주 바뀌거나 model deployment가 잦으면 긴 search는 부담이 된다. cache 가능한가, offline tuning인가, online adaptation인가를 봐야 한다.

또 search result는 hardware와 software version에 묶일 수 있다. driver, compiler, GPU architecture가 바뀌면 최적 schedule이 달라질 수 있다. "한 번 찾은 답"이 영원한 답은 아니다.

5. 손으로 해볼 것

Triton의 작은 matmul이나 reduction kernel을 하나 골라 block size 후보 세 개만 정해 보자. 각 후보에 대해 예상되는 memory traffic, parallelism, register pressure를 말로 적고, 실제 benchmark 결과와 비교한다.

틀려도 괜찮다. search-based compiler를 이해하는 가장 좋은 방법은 내 머릿속 cost model이 얼마나 자주 빗나가는지 보는 것이다. 컴파일러도 결국 그 빗나감을 줄이려는 장치다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_063.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_063.md

1. GPU가 두 개면 자동으로 두 배가 되지 않는다

single GPU에서 빠른 code가 multi-GPU에서 그대로 빠른 것은 아니다. data를 나누고, 결과를 모으고, GPU 사이 communication을 맞추고, host와 device의 역할을 정해야 한다. 이 과정에서 synchronization 하나가 전체 throughput을 잡아먹을 수 있다.

local archive에는 구체 자료가 없으므로 CUDA-aware MPI, NCCL, peer-to-peer, unified memory, multi-process 중 무엇을 다뤘는지는 확인 필요다. 이 노트는 multi-GPU 강의를 읽을 때의 기본 질문을 정리한다.

2. 먼저 분할 방식을 정해야 한다

multi-GPU programming의 첫 질문은 "무엇을 나눌 것인가"다. data parallelism은 batch를 나누고 gradient나 결과를 합친다. tensor parallelism은 model의 matrix 연산 자체를 나눈다. pipeline parallelism은 layer나 stage를 나눈다. inference에서는 prefill/decode 역할을 나눌 수도 있다.

분할 방식에 따라 communication pattern이 달라진다. all-reduce, all-gather, reduce-scatter, point-to-point send/recv가 어디에 놓이는지 봐야 한다. computation이 빨라도 communication이 critical path에 있으면 scaling은 금방 꺾인다.

3. topology는 코드 밖에 있지만 성능 안에 있다

GPU 사이 연결은 중요하다. 같은 machine 안에서도 NVLink, PCIe, switch topology에 따라 bandwidth와 latency가 다르다. node를 넘어가면 network가 더 큰 변수가 된다. multi-GPU code는 hardware topology를 모르면 반쯤 눈을 감고 쓰는 셈이다.

원본 강의를 볼 때는 실험 machine 구성을 꼭 적어야 한다. GPU model, GPU 수, interconnect, driver/library version, process layout이 benchmark 해석에 필요하다. "8 GPU"라는 말만으로는 아무것도 충분하지 않다.

4. overlap은 좋지만 복잡하다

communication을 compute와 겹치면 scaling을 개선할 수 있다. 하지만 overlap은 stream, event, dependency, buffer lifetime을 제대로 다뤄야 한다. 잘못하면 race condition이나 숨은 synchronization이 생긴다.

또 overlap이 항상 가능한 것은 아니다. 다음 계산이 communication 결과를 바로 필요로 하면 기다릴 수밖에 없다. 따라서 algorithm 구조에서 overlap 가능한 구간을 찾는 일이 먼저다. profiler trace는 여기서 매우 유용하다.

5. 손으로 해볼 것

두 GPU가 있다고 가정하고 큰 matrix multiply 결과를 row-wise로 나눠 계산하는 toy plan을 그려 보자. input이 어디에 복사되어야 하는지, 각 GPU가 어떤 output row를 만들지, 마지막 결과를 어디서 합칠지 적는다.

그 다음 tensor parallel matmul처럼 K dimension을 나누는 경우를 생각해 보자. 이번에는 partial sum을 어떻게 합칠지 필요하다. 두 그림의 communication이 다르다는 점이 multi-GPU programming의 첫 문턱이다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_064.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_064.md

1. neighborhood attention은 locality를 attention에 넣는다

일반 self-attention은 모든 token이 모든 token을 볼 수 있다. vision에서 image patch나 feature map token이 많아지면 이 global attention은 비용이 커진다. neighborhood attention은 각 위치가 주변 window 안의 token만 보도록 제한하는 방향으로 이해할 수 있다.

다만 원본 강의가 어떤 version, paper, implementation, benchmark를 다뤘는지는 transcript 없이 확인되지 않는다. 따라서 여기서는 구체 수식이나 성능 수치를 만들지 않는다. 핵심은 local window가 algorithm과 kernel 양쪽에 어떤 질문을 던지는지다.

2. 비용은 줄지만 정보 흐름도 바뀐다

neighborhood attention은 attention 대상 수를 줄여 compute와 memory를 줄일 수 있다. 특히 2D feature map에서는 주변 patch만 보는 구조가 convolution의 locality와 닮아 있다. GPU 관점에서는 attention score matrix 전체를 만들지 않아도 되는 가능성이 생긴다.

하지만 global interaction이 줄어드는 tradeoff가 있다. 멀리 떨어진 token 사이 정보가 바로 섞이지 않으면 layer를 여러 번 거치거나 window를 이동시키는 설계가 필요할 수 있다. speed만 보고 model behavior를 잊으면 안 된다.

3. kernel에서는 gather pattern이 중요하다

local window는 이론적으로 비용을 줄이지만, memory access가 쉬워지는 것은 아니다. 각 query 위치마다 주변 key/value를 모아야 하므로 gather pattern, boundary handling, layout, vectorization이 성능을 좌우한다.

vision tensor는 보통 batch, height, width, channel 축을 가진다. layout이 NHWC인지 NCHW인지, head dimension이 어떻게 놓이는지에 따라 coalescing이 달라진다. neighborhood attention kernel은 "적은 계산"을 "좋은 memory access"로 바꾸는 일이 관건이다.

4. window size는 model과 hardware의 공동 선택이다

window가 커지면 더 넓은 context를 보지만 비용이 늘어난다. window가 작으면 빠르지만 표현력이 제한될 수 있다. dilation, shifted window, hierarchical structure 같은 선택지도 있을 수 있으나, 원본 강의가 무엇을 다뤘는지는 확인 필요다.

benchmark를 볼 때는 resolution, window size, head dimension, batch size, dtype, layout을 함께 적어야 한다. neighborhood attention은 shape에 민감한 operation이다. 한 resolution에서 빠른 구현이 다른 resolution에서도 자동으로 빠르다고 볼 수 없다.

5. 손으로 해볼 것

8 x 8 patch grid에서 각 위치가 3 x 3 neighborhood만 본다고 가정하자. 중앙 patch 하나와 모서리 patch 하나가 각각 몇 개의 key를 보는지 세어 본다. 그 다음 global attention의 64개 key와 비교한다.

이 작은 계산은 neighborhood attention의 장점을 바로 보여 준다. 동시에 boundary mask가 왜 귀찮은지도 드러난다. 예쁜 아이디어가 kernel에서는 늘 약간의 장부 정리를 요구한다.

Source Links

• YouTube
• Slides
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_065.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_065.md

1. 지금 확인된 자료

local notes에는 speaker가 Lanxiang Hu로 기록되어 있고, paper link는 GameArena: Evaluating LLM Reasoning through Live Computer Games다. transcript, slides, code, GPU Mode repo folder는 확인되지 않았다. 따라서 영상에서 실제로 어떤 demo, interface, 분석 chart를 설명했는지는 확인 필요다.

paper abstract 기준으로 GameArena는 static benchmark의 data contamination과 saturation 문제를 피하려는 dynamic benchmark다. 세 개의 game을 통해 deductive reasoning, inductive reasoning 같은 능력을 보려 하고, 사람과 LLM이 함께 참여하는 interaction log를 남긴다. paper abstract는 2000개가 넘는 game session, 100명 user study, 다섯 개 state-of-the-art LLM 평가를 언급하지만, 이 숫자의 조건과 세부 결과는 원문에서 다시 확인해야 한다.

2. benchmark를 "문제집"이 아니라 "환경"으로 보기

일반 benchmark는 입력과 정답이 고정되어 있다. 그래서 한 번 널리 알려지면 model이 실제로 reasoning을 잘하는지, benchmark distribution을 외웠는지 구분하기 어려워진다. Game Arena가 흥미로운 지점은 평가를 live environment로 옮긴다는 데 있다. model은 한 번 답하고 끝나는 것이 아니라, game state를 보고 다음 행동을 고른다.

이렇게 되면 평가 단위도 바뀐다. 최종 승패만 보는 것이 아니라, 중간 결정, 정보 갱신, 상대 행동에 대한 반응을 함께 봐야 한다. reasoning을 점수 하나로 납작하게 만드는 대신, 과정의 흔적을 남기려는 시도라고 읽으면 좋다.

3. systems 관점에서 중요한 점

LLM evaluation도 결국 system design이다. 어떤 game을 고를지, 사람이 얼마나 개입할지, model에게 어떤 정보를 줄지, latency가 interaction quality에 영향을 주는지 모두 결과를 바꾼다. 특히 live game은 benchmark server, UI, logging, replay, privacy 같은 운영 문제가 따라온다.

GPU 성능 자체를 다루는 강의는 아니지만, inference serving과 연결되는 질문이 있다. model이 게임 중간에 여러 번 호출된다면 latency와 cost가 user engagement에 영향을 준다. reasoning benchmark가 커질수록 "좋은 prompt"만이 아니라 안정적인 serving pipeline도 평가의 일부가 된다.

4. 조심해서 읽을 부분

game benchmark는 재미있지만, 재미와 reasoning이 항상 같은 방향으로 움직이지는 않는다. 어떤 game은 전략보다 language fluency나 interface 적응력이 더 중요할 수 있다. 사람과 함께 하는 환경이라면 participant pool, instruction wording, session length도 결과에 섞인다.

그래서 원본을 볼 때는 세 가지를 적어 두자. 첫째, 세 game이 각각 어떤 reasoning 능력을 겨냥하는가. 둘째, LLM과 사람에게 주어진 정보가 같은가. 셋째, 분석이 win rate만 보는지, step-by-step decision도 보는지. 여기서 대충 넘어가면 benchmark가 예쁜 포스터가 되고, 읽는 사람은 아무것도 재현하지 못한다.

5. 혼자 해볼 작은 실험

작은 text game 하나를 골라 두 model에게 같은 state를 주고 다음 action을 고르게 해 보자. 중요한 것은 model이 이겼는지가 아니라, state 변화에 따라 이전 가정을 버리는지 기록하는 것이다. session log를 table로 만들면 reasoning benchmark가 왜 static QA보다 귀찮고, 동시에 왜 더 흥미로운지 바로 보인다.

이 강의는 자료가 보강되면 paper의 game design과 lecture demo 중심으로 다시 써야 한다. 지금 버전은 Game Arena를 과장하지 않고 읽기 위한 지도다.

Source Links

• YouTube
• Paper
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_066.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_066.md

1. 지금 확인된 자료

local notes에는 speaker가 Jeff Hammond로 기록되어 있고, slides는 Google Drive, code는 ParRes Kernels/tree/main/Cxx11로 연결되어 있다. transcript가 없으므로 lecture가 어떤 kernel 예제를 live로 다뤘는지, NCCL과 NVSHMEM을 어떤 순서로 비교했는지는 알 수 없다.

그래도 제목만으로도 읽는 축은 꽤 분명하다. NCCL은 all-reduce, broadcast, all-gather 같은 collective communication을 다루는 대표 library다. NVSHMEM은 GPU thread가 remote memory access와 synchronization을 더 직접적으로 표현할 수 있게 하는 PGAS 계열 model로 읽을 수 있다.

2. collective는 "다 같이 약속한 동작"이다

NCCL을 쓸 때 programmer는 보통 "이 tensor를 모든 GPU에서 reduce하자"처럼 group operation을 요청한다. library는 topology, protocol, chunking, scheduling을 안쪽에서 처리한다. 사용자는 높은 수준의 collective semantics를 얻고, 대신 세부 통신 순서는 library에게 맡긴다.

이 abstraction은 training에서 특히 강하다. data parallel all-reduce, tensor parallel all-gather, pipeline boundary communication처럼 pattern이 반복될 때 collective API는 읽기 쉽고 실수도 줄인다. 하지만 모든 communication이 예쁜 collective로 떨어지지는 않는다. irregular exchange, fine-grained handoff, computation과 communication의 촘촘한 overlap을 직접 표현하고 싶을 때 다른 model이 필요해진다.

3. NVSHMEM을 볼 때의 질문

NVSHMEM을 "NCCL보다 low-level"이라고만 기억하면 조금 아쉽다. 더 좋은 질문은 communication을 kernel 내부의 control flow와 얼마나 가까이 붙일 수 있느냐다. GPU가 remote data를 가져오거나 밀어 넣고, 필요한 순간에 signal을 기다리는 식의 표현은 특정 알고리즘에서 CPU scheduling 부담을 줄일 수 있다.

물론 그 자유에는 비용이 있다. memory ordering, synchronization, progress, topology를 더 많이 이해해야 한다. collective가 숨겨 주던 세부 사항이 다시 programmer 앞으로 나온다. 이 강의의 code link를 볼 때도 "API가 신기하다"보다 "어떤 pattern에서 이 노출이 이득인가"를 먼저 물어보면 좋다.

4. 성능에서 조심할 점

multi-GPU 성능은 bandwidth 숫자 하나로 끝나지 않는다. 작은 message는 latency와 synchronization overhead가 지배할 수 있고, 큰 message는 link bandwidth와 topology가 중요해진다. computation과 communication이 겹치는지, stream synchronize가 숨어 있는지도 봐야 한다.

특히 NCCL과 NVSHMEM을 비교한다면 같은 일을 하고 있는지부터 확인해야 한다. collective와 one-sided operation은 API semantics가 다르기 때문에, benchmark가 공정하려면 data movement 양, synchronization 조건, correctness check를 맞춰야 한다.

5. 손으로 해볼 것

두 GPU가 있다면 작은 all-reduce benchmark와 point-to-point style exchange benchmark를 나눠 돌려 보자. message size를 바꾸며 latency 구간과 bandwidth 구간이 어디서 갈라지는지 적어 둔다. GPU가 없다면 ParRes code를 읽으며 communication call이 어디에서 synchronization을 요구하는지만 표시해도 좋다.

원본 slides를 볼 때는 NCCL과 NVSHMEM의 결론을 "A가 B보다 빠르다"로 적지 말자. 어느 topology, 어느 message size, 어느 synchronization model에서 그런지까지 같이 적어야 다음 사람이 다시 실험할 수 있다.

Source Links

• YouTube
• Slides
• Code
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_067.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_067.md

1. 지형도라는 제목의 함정

Landscape라는 말은 넓은 survey를 기대하게 만든다. 하지만 실제 강의가 어떤 기술을 포함했는지는 아직 모른다. NCCL, MPI, NVSHMEM, UCX, GPUDirect, runtime system 중 무엇이 중심이었는지 source bundle은 말해 주지 않는다.

그래서 이 글은 "이 기술들이 강의에 나왔다"가 아니라 "이런 분류축으로 들으면 좋다"에 머문다. 통신 주제는 이름이 비슷하고 층이 겹쳐서, 출처 없는 나열이 가장 쉽게 생긴다.

2. GPU-centric을 두 단계로 나눠 듣기

첫 단계는 data path다. GPU buffer가 host memory를 거치지 않고 다른 GPU나 network 쪽과 더 직접 연결되는지 묻는다. 이 질문은 bandwidth, copy overhead, topology와 이어진다.

두 번째는 control path다. 통신을 CPU가 시작하고 기다리는지, GPU execution 가까이에서 시작하거나 진행을 표현할 수 있는지 묻는다. 두 질문은 관련 있지만 같지 않다. data path가 GPU-aware여도 control은 host API 중심일 수 있다.

3. communication 병목은 한 문장이 아니다

"통신이 느리다"는 말은 너무 크다. 작은 message의 latency가 문제일 수 있고, 큰 tensor의 bandwidth가 문제일 수 있다. synchronization 때문에 GPU가 쉬는 것일 수도 있고, topology와 collective algorithm이 맞지 않는 것일 수도 있다.

GPU-centric communication을 읽을 때는 그래서 symptom을 쪼개야 한다. profiler timeline에서 gap이 보이는가, network counter가 찼는가, collective가 serialized되었는가, kernel과 transfer가 겹쳤는가. 강의가 어떤 증상을 사례로 삼는지 원본에서 확인한다.

4. ML systems와 이어지는 자리

대규모 training이나 distributed model system에서는 all-reduce, all-gather, reduce-scatter, send/recv가 model path 안으로 들어온다. 통신은 더 이상 training loop 바깥의 부록이 아니라 layer latency와 memory pressure를 바꾸는 요인이다.

다만 이 강의가 특정 model family를 중심으로 삼았다는 근거는 현재 없다. GPU Mode 독자는 자신의 배경지식으로 연결해 보되, 강의 본문에 있었던 예제와 독자의 응용을 따로 표시해야 한다.

5. 읽으면서 만들 표

원본 영상을 볼 수 있다면 등장한 API나 runtime 이름을 행으로 둔다. 열은 data path, control path, operation type, synchronization model, topology assumption, example workload로 둔다.

표를 채우다 보면 "GPU-centric"이 단일 기능명이 아니라 여러 선택의 묶음이라는 점이 드러난다. 이 강의가 그 중 어디까지를 landscape로 잡았는지가 나중에 rich 원고의 중심이 된다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_068.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_068.md

1. 지금 확인된 자료

local notes에는 speaker가 Roberto Castro and Andrei Panferov로 기록되어 있고, paper는 Quartet: Native FP4 Training Can Be Optimal for Large Language Models다. code link는 IST-DASLab/Quartet과 isT-DASLab/qutlass다. transcript fetch는 실패했다.

paper abstract에 따르면 Quartet은 Llama-type model을 대상으로 low-precision scaling law를 살피고, FP4 precision에서 accuracy와 computation의 trade-off를 다룬다. "state-of-the-art"나 "optimal" 같은 표현은 논문 안의 조건을 확인해야 한다. reader에서는 결론을 먼저 믿기보다 실험 설정을 먼저 읽어야 한다.

2. 4-bit training은 weight-only quantization과 다르다

많은 사람이 4-bit라고 들으면 inference용 weight quantization을 떠올린다. 하지만 training에서 4-bit를 쓴다는 것은 forward, backward, gradient, accumulation, optimizer update 중 어디를 낮은 precision으로 처리하는지 따져야 한다는 뜻이다. weight만 줄이는 문제보다 훨씬 까다롭다.

FP4는 표현할 수 있는 값의 범위와 간격이 매우 거칠다. scale 선택, outlier 처리, accumulation dtype, stochastic rounding 여부가 곧 학습 안정성으로 이어진다. Quartet을 읽을 때도 "몇 bit인가"보다 "어느 tensor가 어떤 format으로 계산되는가"를 먼저 봐야 한다.

3. scaling law가 있다면 조건이 본문이다

paper abstract는 low-precision scaling law를 언급한다. 이런 주장은 매력적이지만, scaling law는 식보다 조건표가 중요하다. model family, parameter size, token count, optimizer, learning rate schedule, dataset, baseline precision, evaluation metric이 함께 있어야 의미가 생긴다.

특히 4-bit training에서는 hardware generation도 조건이다. Blackwell의 FP4 support를 전제로 한 결과라면, 이전 GPU에서 같은 code를 흉내 내도 같은 cost model이 나오지 않을 수 있다. code link의 qutlass는 kernel 쪽을 볼 때 중요한 단서가 될 가능성이 있다.

4. GPU Mode 독자에게 중요한 연결점

Quartet은 algorithm paper인 동시에 kernel paper처럼 읽어야 한다. FP4 training이 이론적으로 가능해도, kernel이 conversion과 scaling overhead에 먹히면 실제 throughput 이득은 줄어든다. 반대로 hardware path에 잘 올라가면 memory bandwidth, Tensor Core throughput, energy efficiency가 함께 움직인다.

여기서 좋은 질문은 "FP16 대비 몇 배 빠른가" 하나가 아니다. 어떤 layer에서 이득이 크고, 어느 부분이 여전히 higher precision fallback을 요구하는지 봐야 한다. training 전체 step time에서 FP4 kernel이 차지하는 비중도 확인해야 한다.

5. 혼자 해볼 것

간단한 linear layer를 놓고 FP16, FP8, FP4 가상 policy를 표로 비교해 보자. weight, activation, gradient, accumulation, optimizer state를 행으로 두고 dtype을 채우면 "4-bit training"이라는 말이 얼마나 많은 설계 선택을 숨기는지 보인다.

원본 강의나 paper를 다시 볼 때는 모든 benchmark 옆에 hardware, model size, token count, metric을 적자. low precision은 숫자가 멋질수록 조건을 더 끈질기게 봐야 한다. 이건 의심이 아니라 예의에 가깝다.

Source Links

• YouTube
• Paper
• Code: Quartet
• Code: qutlass
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_069.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_069.md

1. 지금 확인된 자료

local notes에는 speaker가 mike64t로 기록되어 있고, Google Slides link가 있다. code와 paper link는 없다. transcript가 없으므로 PCCL이 정확히 어떤 project 이름인지, 어떤 failure model을 가정하는지, 어떤 benchmark를 보여줬는지는 확인 필요다.

제목이 말하는 두 단어는 분명하다. collective와 fault tolerance다. all-reduce 같은 collective는 모든 rank가 정해진 순서로 참여해야 한다. 그런데 rank 하나가 느려지거나 죽거나 network path가 흔들리면, 전체 training job이 같이 멈출 수 있다.

2. collective는 실패에 약한 약속이다

collective는 "모두가 들어와야 끝나는" 연산이다. 그래서 성능이 좋고 semantics가 깔끔한 대신, partial failure에는 민감하다. 한 rank가 collective call에 도달하지 못하면 나머지는 기다린다. timeout과 retry를 붙일 수는 있지만, 그 순간 correctness와 progress의 정의가 어려워진다.

fault tolerant collective는 여기서 질문을 바꾼다. 실패를 드문 accident로 볼 것인가, 아니면 large-scale job에서 반드시 만나는 상태로 볼 것인가. 후자라면 collective library는 단순히 빠른 ring을 만드는 것을 넘어 membership, recovery, consistency까지 생각해야 한다.

3. LLM training에서 왜 중요한가

큰 training job은 GPU 수가 늘수록 failure surface도 넓어진다. 하루짜리 실험이라면 재시작으로 버틸 수 있지만, 긴 pretraining run에서는 한 번의 failure가 큰 비용이 된다. communication layer가 failure를 더 잘 다루면 checkpoint interval, wasted compute, operator intervention이 모두 달라진다.

하지만 fault tolerance는 공짜가 아니다. logging, redundancy, recovery protocol, extra synchronization이 latency와 bandwidth를 잡아먹을 수 있다. PCCL을 읽을 때는 "얼마나 견디는가"와 "정상 상태에서 얼마를 더 내는가"를 같이 봐야 한다.

4. 확인해야 할 세부 질문

첫째, failure model이다. process crash만 다루는지, GPU reset, network partition, silent data corruption까지 보는지 확인해야 한다. 둘째, recovery granularity다. collective 하나를 재시도하는지, communicator를 재구성하는지, training step을 rollback하는지에 따라 system design이 달라진다.

셋째, integration layer다. PCCL이 NCCL replacement인지 wrapper인지, PyTorch distributed와 어떻게 연결되는지, application code가 얼마나 바뀌는지 봐야 한다. fault tolerance는 library 안에서만 멋져도 실제 stack에 못 들어가면 오래 살기 어렵다.

5. 혼자 해볼 것

작은 all-reduce pseudo-code를 그리고 한 rank가 중간에 멈추는 경우를 손으로 추적해 보자. 나머지 rank가 어디에서 block되는지, timeout을 넣으면 tensor state가 어떻게 되는지 적어 보면 fault tolerant collective가 왜 어려운지 금방 드러난다.

원본 slides를 볼 때는 PCCL의 API 이름보다 failure scenario 그림을 먼저 찾자. 어떤 failure를 견디는지 알면, 성능 숫자도 비로소 읽힌다.

Source Links

• YouTube
• Slides
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_070.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_070.md

1. 지금 확인된 자료

local notes에는 speaker가 Sewon Min으로 기록되어 있고, GitHub lecture folder와 ScaleML_FlexOlmo.pdf가 있다. transcript와 code는 없다. slide 첫 부분은 standard MoE가 보통 training/inference efficiency를 위해 쓰이지만, 이 talk에서는 서로 다른 party가 소유한 distributed dataset에 학습을 가능하게 하는 방향을 다룬다고 소개한다.

slide에는 centralized data access가 현실에서는 binary가 아니며, data가 later available이 되거나 sharing restriction을 가질 수 있다는 문제의식이 보인다. 따라서 이 노트는 FlexOlmo를 "더 빠른 MoE"보다 "data 사용 조건을 model architecture에 반영하려는 시도"로 읽는다.

2. data access는 현실에서 지저분하다

언어모델 학습 설명은 종종 모든 data가 한 bucket에 들어 있다고 가정한다. 현실에서는 그렇지 않다. public data와 proprietary data가 나뉘고, 어떤 data는 공유할 수 없고, 어떤 data는 나중에 생긴다. 중앙 학습을 기본값으로 두면 이런 조건은 모두 예외 처리가 된다.

FlexOlmo가 흥미로운 이유는 이 예외를 architecture 문제로 끌고 들어온다는 점이다. data owner가 raw data를 직접 공유하지 않고도 contribution을 만들 수 있다면, open model 생태계의 협업 방식도 달라질 수 있다.

3. MoE가 여기서 하는 일

MoE는 보통 parameter를 많이 두되 token마다 일부 expert만 쓰는 efficiency 구조로 소개된다. FlexOlmo slide는 MoE를 data ownership 문제에도 적용하려 한다. 각 dataset 또는 party의 contribution을 expert처럼 다루면, data를 합치지 않고도 model을 구성하거나 나중에 특정 contribution을 빼는 상상을 할 수 있다.

물론 "expert를 합치면 된다"는 말은 쉽고, 실제 training은 어렵다. routing, expert specialization, shared representation, merging 후 interference, opt-out 후 성능 변화가 모두 문제가 된다. slide에는 news data expert를 opt-out하는 결과가 언급되지만, 그 조건과 metric은 PDF 본문에서 다시 확인해야 한다.

4. GPU/시스템 독자에게 중요한 연결점

MoE는 algorithm만의 문제가 아니다. expert가 늘어나면 routing, all-to-all communication, memory placement, batching이 시스템 병목이 된다. FlexOlmo가 data flexibility를 얻기 위해 MoE를 쓴다면, 그 flexibility가 training과 serving cost를 어떻게 바꾸는지도 함께 봐야 한다.

특히 distributed data owner가 각자 expert를 학습하는 구조라면 checkpoint format, compatibility, tokenizer, optimizer state, evaluation protocol도 중요해진다. 좋은 architecture idea가 실제 open model workflow로 내려오려면 이런 지루한 부분까지 맞아야 한다.

5. 혼자 해볼 것

작은 toy corpus 세 개를 domain별로 나누고, "이 domain을 나중에 opt-out해야 한다"는 조건을 붙여 보자. dense model 하나를 학습하는 경우와 expert별로 분리하는 경우를 diagram으로 비교하면 FlexOlmo의 문제의식이 선명해진다.

원본 slides를 볼 때는 FlexOlmo가 해결하려는 policy problem과 model problem을 분리해서 적자. data governance와 GPU efficiency를 한 문장에 섞으면 멋져 보이지만, debugging은 훨씬 어려워진다.

Source Links

• YouTube
• Slides folder
• Slides PDF
• ScaleML Bootcamp
• Speaker site
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_071.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_071.md

1. 지금 확인된 자료

local notes에는 speaker가 Guangxuan Xiao로 기록되어 있고, GitHub folder와 StreamingLLM.pdf가 있다. transcript와 code link는 notes에는 없지만, slide text에는 StreamingLLM code/blog link가 등장한다. 이 reader에서는 lecture notes의 source와 slide 내용만 사용한다.

slide는 streaming application, multi-round dialogue, decoding 단계의 KV cache memory 문제를 출발점으로 삼는다. 긴 context에서 이전 token의 key/value를 모두 저장하면 memory가 sequence length에 따라 커진다. 그래서 window attention처럼 최근 token만 남기는 접근이 자연스럽지만, 초기 token을 버리면 model이 무너질 수 있다고 설명한다.

2. KV cache가 long-context의 몸값이다

decoder-only Transformer는 새 token을 만들 때 현재 query와 과거 token의 key/value를 사용한다. 이미 계산한 K/V를 cache하면 재계산은 피하지만, 길이가 늘수록 memory가 쌓인다. serving 관점에서는 이 memory가 batch size와 sequence length를 직접 누른다.

window attention은 cache를 최근 token으로 제한해 memory를 줄인다. 문제는 model이 긴 문맥에서 특정 초기 token에 계속 attention을 주는 현상이다. slide는 이를 attention sink라고 부르고, 관련 없는 것처럼 보여도 attention을 끌어당기는 token이 있다고 설명한다.

3. StreamingLLM의 조심스러운 약속

slide는 StreamingLLM이 attention sink token의 KV와 sliding window KV를 함께 보존한다고 설명한다. 이렇게 하면 모든 과거 token을 들고 있지 않아도 stream을 오래 처리할 수 있다. 또 slide는 recomputation baseline 대비 speedup claim을 제시하지만, 정확한 조건은 PDF의 figure와 실험 설정을 확인해야 한다.

중요한 단서도 slide에 직접 있다. StreamingLLM이 infinite context를 주는 것은 아니다. cache에서 evict된 token은 attend할 수 없다. 즉 "무한히 긴 텍스트를 완전히 기억한다"가 아니라, streaming deployment에서 안정적으로 굴러가는 cache policy에 가깝다.

4. GPU Mode 독자에게 중요한 연결점

KV cache는 model architecture와 GPU memory system이 만나는 자리다. cache layout, quantization, paging, batch scheduling, prefill/decode 분리, attention kernel 모두 여기로 모인다. long-context modeling을 논문 아이디어로만 읽으면 serving 병목을 놓치기 쉽다.

attention sink를 보존하는 policy는 kernel에도 영향을 준다. attention이 보는 key/value index가 연속 window만이 아니라 sink prefix와 recent window의 조합이 되기 때문이다. 실제 구현에서는 gather pattern, memory coalescing, cache compaction을 확인해야 한다.

5. 혼자 해볼 것

toy decoder setting을 놓고 KV cache memory를 계산해 보자. layer 수, KV head 수, head dim, dtype byte, batch size, sequence length를 곱하면 long-context가 왜 곧 memory 문제인지 금방 보인다.

원본 slide를 볼 때는 "attention sink를 몇 개 보존하는가", "position id를 어떻게 처리하는가", "어떤 model family에서 평가했는가", "evicted token을 다시 볼 수 없는 한계를 어떻게 설명하는가"를 따로 적어 두자.

Source Links

• YouTube
• Slides folder
• Slides PDF
• ScaleML Bootcamp
• Speaker site
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_072.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_072.md

1. 제목은 넓고 근거는 좁다

Quantization in Large Models라는 제목 하나에는 많은 방법이 들어갈 수 있다. PTQ, QAT, weight-only, activation quantization, KV cache quantization, FP8 training, INT4 serving이 모두 후보처럼 보인다.

하지만 후보는 강의 내용이 아니다. 지금 draft는 method catalog를 만들지 않고, 원본을 볼 때 어떤 축으로 받아 적어야 하는지 정리한다. ScaleML bootcamp 링크는 맥락 단서지만, 강의의 세부 주장까지 대신해 주지는 않는다.

2. 먼저 tensor를 고른다

quantization을 읽을 때 "몇 bit인가"보다 먼저 "무엇을 낮추는가"를 본다. weight만 낮추는지, activation도 낮추는지, KV cache까지 낮추는지에 따라 memory, compute, accuracy risk가 달라진다.

training까지 다루면 이야기는 더 조심스러워진다. gradient, optimizer state, loss scaling, numerical stability가 들어온다. 원본에서 어느 범위를 다뤘는지 표시하지 않으면, large model quantization이라는 말이 너무 쉽게 모든 것을 덮는다.

3. granularity는 정확도와 비용을 동시에 바꾼다

scale을 tensor 전체에 하나만 두면 단순하지만 오차가 커질 수 있다. channel-wise나 group-wise scale은 더 섬세하지만 metadata와 dequant 비용을 만든다. large model에서는 이 overhead도 무시하기 어렵다.

그래서 quantization scheme은 model-level claim과 kernel-level claim을 같이 가진다. accuracy가 유지되는지, memory traffic이 줄었는지, hardware path가 빠른 dtype을 실제로 쓰는지, scale handling이 병목이 되지 않는지 함께 봐야 한다.

4. benchmark를 읽는 순서

baseline을 먼저 적는다. BF16/FP16과 비교하는지, FP8과 비교하는지, 이미 optimized된 low-bit baseline과 비교하는지에 따라 의미가 달라진다.

그다음 metric을 분리한다. perplexity나 downstream score는 품질을 말하고, latency/throughput/memory footprint는 system behavior를 말한다. 둘 중 하나만 있으면 quantization의 절반만 본 것이다.

5. 작게 재현해 볼 질문

transformer block 하나를 골라 tensor 목록을 만든다. embedding, QKV projection, attention score, value, output projection, MLP, KV cache 옆에 dtype과 scale granularity를 적어 본다.

그 표를 보면 "모델을 INT4로 만들었다" 같은 한 문장이 얼마나 많은 선택을 생략하는지 보인다. 원본을 볼 때 speaker가 실제로 고른 선택지만 이 표에 표시하면 된다.

Source Links

• YouTube
• ScaleML Bootcamp
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_073.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_073.md

1. 지금 확인된 자료

local notes에는 speaker가 Songlin Yang으로 기록되어 있고, GitHub folder와 path_talk.pdf가 있다. slide 제목은 Modern Position Encoding in Transformers: RoPE/Yarn and PaTH다. transcript와 code는 없다.

slide 초반은 positional information의 필요성에서 시작한다. causal mask가 없는 self-attention은 input token을 unordered set처럼 다룰 수 있으므로, word order를 구분하려면 position 정보를 넣어야 한다. 이어 absolute positional encoding, RoPE, YaRN, PaTH가 등장한다.

2. RoPE의 좋은 점과 긴 문맥의 불편함

RoPE는 query와 key에 회전을 적용해 attention logit이 상대 위치에 의존하도록 만든다. slide는 rotation matrix의 composition과 inverse 같은 성질을 짚고, RoPE attention score가 absolute position이 아니라 i - j에 연결된다고 설명한다.

이 구조는 elegant하다. 하지만 slide는 RoPE가 pre-training context보다 긴 sequence에서 perplexity가 빠르게 나빠질 수 있다고 설명한다. YaRN이나 position interpolation 같은 방법은 이런 long-context extrapolation 문제를 다루는 쪽으로 읽을 수 있다.

3. PaTH가 던지는 질문

slide에서 PaTH는 RoPE와 비교되며 data-dependent generalized Householder matrix라는 표현이 나온다. RoPE가 data-independent rotation이라면, PaTH는 position 처리에 더 표현력 있는 transformation을 넣으려는 방향으로 보인다. slide에는 PaTH가 특정 이론적 문제를 풀 수 있다는 주장도 나오지만, 이 부분은 원문 증명을 확인해야 한다.

중요한 것은 PaTH를 "RoPE보다 세다"로만 읽지 않는 것이다. 더 강한 표현력은 학습 안정성, kernel cost, checkpoint conversion, long-context generalization을 함께 건드린다. slide에는 RoPE 기반 pretrained checkpoint를 PaTH로 바꾸는 cheap conversion 이야기도 보이지만, 실제 품질과 비용은 세부 실험을 봐야 한다.

4. GPU/커널 관점의 연결

position encoding은 model math처럼 보이지만 kernel에도 영향을 준다. RoPE는 attention 앞에서 Q/K에 적용되는 비교적 규칙적인 elementwise transform이다. fused attention kernel 안에 넣을 수 있고, layout과 vectorization이 중요하다.

PaTH처럼 더 복잡한 transformation이 들어가면 compute overhead와 memory access pattern을 확인해야 한다. long-context 성능을 얻기 위해 attention kernel이 느려지면 serving에서는 trade-off가 생긴다. architecture idea와 kernel path를 같이 읽는 습관이 필요하다.

5. 혼자 해볼 것

작은 2D rotation matrix로 RoPE의 핵심 성질을 손으로 확인해 보자. R(a)^T R(b) = R(b-a)가 왜 상대 위치로 이어지는지 직접 써 보면, slide의 수식이 훨씬 덜 낯설어진다.

원본 slide를 볼 때는 RoPE, YaRN, PaTH 각각에 대해 "무엇을 바꾸는가", "long-context에서 무엇을 개선한다고 주장하는가", "kernel 비용은 어떻게 변하는가"를 표로 정리하자.

Source Links

• YouTube
• Slides folder
• Slides PDF
• ScaleML Bootcamp
• Speaker site
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_074.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_074.md

1. 지금 확인된 자료

local notes에는 speakers가 William Brandon과 Simran Arora로 기록되어 있고, GPU Programming Fundamentals.pdf와 ThunderKittens.pdf가 있다. transcript와 code link는 없다. slide에는 CUDA program의 block/thread launch와 hardware-first perspective가 먼저 나오고, ThunderKittens slide에는 "simple abstractions for hardware-efficient AI kernel programming"이라는 방향이 보인다.

따라서 이 원고는 강의 전체를 완성된 tutorial로 재현하지 않는다. 대신 두 축을 잡는다. 하나는 GPU가 thread, memory hierarchy, bandwidth/latency로 움직인다는 기초 감각이고, 다른 하나는 AI kernel을 tile 단위로 표현해 hardware 선호와 가까워지려는 시도다.

2. hardware-first로 kernel을 보기

CUDA program은 grid, block, thread로 보이지만, 성능은 그 아래 hardware에서 결정된다. warp가 어떻게 실행되는지, register와 shared memory가 얼마나 가까운지, global memory access가 얼마나 비싼지 알아야 같은 code도 다르게 보인다.

초보자는 종종 "thread를 많이 쓰면 빠르다"에서 멈춘다. 실제로는 memory coalescing, occupancy, register pressure, synchronization, arithmetic intensity가 같이 움직인다. fundamentals part는 이 지도를 깔아 주는 역할로 읽으면 좋다.

3. ThunderKittens의 읽기 포인트

ThunderKittens slide는 AI hardware와 tile-based programming을 큰 축으로 둔다. 여기서 tile은 단순한 loop blocking이 아니라, matmul과 attention 같은 AI kernel에서 data movement와 compute를 함께 조직하는 단위다. 좋은 tile abstraction은 programmer가 hardware 세부를 모두 손으로 쓰지 않아도 효율적인 path를 타게 도와준다.

하지만 abstraction은 늘 긴장 속에 있다. 너무 높으면 hardware가 원하는 layout과 synchronization을 표현하기 어렵고, 너무 낮으면 CUDA를 직접 쓰는 것과 다르지 않다. ThunderKittens를 볼 때는 API가 귀여운지보다, 어떤 hardware primitive를 안전하게 감싸는지 보는 편이 낫다.

4. 성능에서 조심할 점

AI kernel은 대부분 "연산을 많이 한다"와 "데이터를 잘 움직인다" 사이의 균형이다. matmul은 tile이 잘 맞으면 compute-bound에 가까워질 수 있지만, attention은 sequence length, mask, KV cache layout에 따라 memory-bound 성격이 강해질 수 있다.

ThunderKittens 같은 framework를 평가할 때도 generated code 또는 backend path를 봐야 한다. 같은 high-level tile expression이 GPU generation마다 다르게 최적화될 수 있다. slide가 언급하는 hardware generation 변화의 문제의식도 여기와 맞닿아 있다.

5. 혼자 해볼 것

작은 tiled matmul을 손으로 그려 보자. A와 B tile이 shared memory에 들어오고, thread들이 register accumulator를 갱신한 뒤, C tile을 저장하는 흐름을 표시한다. 이 그림이 익숙해지면 ThunderKittens 같은 tile abstraction이 왜 필요한지 더 잘 보인다.

원본 slides를 볼 때는 fundamentals와 ThunderKittens를 따로 보지 말자. "이 abstraction은 방금 배운 hardware 사실 중 무엇을 숨기고, 무엇을 드러내는가"를 계속 물으면 강의가 더 오래 남는다.

Source Links

• YouTube
• Slides folder
• GPU Programming Fundamentals PDF
• ThunderKittens PDF
• ScaleML Bootcamp
• Speaker site
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_075.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_075.md

1. 지금 확인된 자료

local notes에는 speaker가 Mark Saroufim으로 기록되어 있고 YouTube link만 있다. transcript fetch는 실패했다. slides, code, paper, repo folder는 확인되지 않았다.

따라서 BackendBench가 실제 project인지, benchmark suite인지, PyTorch backend 비교인지, Triton/CUDA kernel correctness harness인지는 아직 단정할 수 없다. 이 글은 title이 가리키는 correctness problem을 중심으로 읽기 지도를 만든다.

2. LLM kernel correctness는 생각보다 미끄럽다

kernel이 빠르다고 바로 쓸 수는 없다. attention, RMSNorm, rotary embedding, quantized matmul, fused optimizer처럼 LLM kernel은 dtype, shape, mask, stride, broadcasting, numerical tolerance가 얽힌다. 조금 틀린 kernel은 unit test 몇 개를 통과하고도 production에서 조용히 model quality를 망칠 수 있다.

특히 low precision과 fusion이 들어가면 "정답"도 단순하지 않다. FP32 reference와 bitwise match를 요구할 수 없는 경우가 많고, tolerance를 너무 넓게 잡으면 버그를 놓친다. correctness benchmark는 이 균형을 정해야 한다.

3. backend가 많아질수록 test가 언어가 된다

CUDA, Triton, vendor library, custom DSL, compiler backend가 같은 operation을 구현할 수 있다. 이때 benchmark가 latency table만 있으면 서로 대화하기 어렵다. 같은 input domain에서 같은 semantics를 만족하는지 먼저 맞춰야 한다.

BackendBench라는 이름이 의미 있으려면, backend 간 비교 가능한 contract를 제공해야 한다. input shape, dtype, layout, tolerance, determinism, supported hardware를 명시해야 한다. 그래야 새 backend가 들어와도 "빠르다" 전에 "같은 일을 한다"를 말할 수 있다.

4. 성능과 correctness를 같이 읽기

correctness test는 성능과 분리되지 않는다. 어떤 backend는 특정 shape에서만 빠르고, 특정 dtype에서만 정확할 수 있다. 또 deterministic mode를 켜면 느려질 수도 있다. production 선택은 항상 correctness envelope와 performance envelope를 함께 본다.

LLM serving에서는 rare shape도 중요하다. batch size 1 decode path, 긴 sequence, padding이 많은 request, speculative decoding의 odd shape가 kernel을 다른 경로로 보낼 수 있다. test suite가 예쁜 square matmul만 보면 실제 문제를 놓친다.

5. 혼자 해볼 것

Triton으로 작은 RMSNorm이나 matmul wrapper를 만들고 PyTorch reference와 비교하는 test를 짜 보자. contiguous tensor만 넣지 말고 stride가 다른 view, 작은 shape, 큰 shape, 극단값을 넣어 본다. performance benchmark보다 먼저 correctness grid를 만드는 습관을 들이면 kernel 개발이 훨씬 덜 불안해진다.

원본 강의를 볼 때는 BackendBench가 어떤 operation과 backend를 포함하는지, tolerance를 어떻게 정하는지, failure를 어떻게 report하는지 확인하자. 좋은 correctness tool은 실패했을 때 사람이 다음 행동을 알 수 있게 해 준다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_076.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_076.md

1. 지금 확인된 자료

local notes에는 speaker가 Tri Dao로 기록되어 있고 YouTube link만 있다. transcript fetch는 실패했다. slides와 code는 확인되지 않았다. 따라서 이 강의가 Triton, ThunderKittens, CuTe, TVM, MLIR 계열 중 무엇을 중심으로 다뤘는지는 알 수 없다.

그러므로 이 reader는 특정 DSL 소개가 아니다. GPU kernel DSL을 평가할 때의 질문 목록이다. Tri Dao의 다른 작업을 떠올려도 여기서는 title 이상의 내용을 강의에 귀속시키지 않는다.

2. DSL은 편의 문법이 아니라 성능 contract다

GPU kernel DSL은 programmer에게 더 높은 수준의 표현을 준다. 하지만 좋은 DSL은 단지 CUDA 문법을 숨기는 데서 끝나지 않는다. 어떤 parallelism이 생기는지, memory hierarchy를 어떻게 쓰는지, compile-time constant와 runtime value가 어떻게 나뉘는지 표현해야 한다.

Triton을 예로 들면 program id, block size, mask, vectorized load/store가 언어의 핵심 감각이다. CUDA보다 짧게 쓸 수 있지만, tiling과 memory access pattern을 모르면 여전히 느린 kernel이 나온다. DSL은 마법 지팡이가 아니라 더 짧은 손잡이다.

3. abstraction의 좋은 긴장

너무 low-level이면 programmer가 모든 bank conflict와 pipeline을 직접 챙겨야 한다. 너무 high-level이면 compiler가 중요한 결정을 추측해야 한다. GPU DSL의 어려움은 이 사이의 적절한 노출을 찾는 것이다.

좋은 DSL은 자주 쓰는 pattern을 쉽게 만들되, 성능이 안 나올 때 내려갈 통로를 남긴다. tile shape, num warps, stages, memory layout, vectorization을 조정할 수 있어야 한다. 반대로 너무 많은 knob이 있으면 사용자는 autotuner 없이는 길을 잃는다.

4. LLM kernel에서 DSL이 빛나는 자리

LLM workload는 비슷한 pattern이 반복된다. matmul, attention, normalization, activation, quantize/dequantize, KV cache update가 계속 나온다. 이 반복성은 DSL과 compiler가 재사용 가능한 pattern을 만들기 좋은 환경이다.

하지만 LLM kernel은 shape가 고정되어 있지 않다. prefill과 decode가 다르고, batch size와 sequence length가 흔들리고, dtype과 hardware generation도 바뀐다. DSL이 실제로 유용하려면 빠른 happy path뿐 아니라 shape 다양성도 견뎌야 한다.

5. 혼자 해볼 것

같은 vector add나 layernorm을 CUDA와 Triton으로 각각 써 보고, code 길이보다 드러난 개념이 어떻게 다른지 적어 보자. CUDA는 thread/block과 memory를 직접 드러내고, Triton은 block tensor와 mask를 더 앞에 둔다. 이 차이를 느끼면 DSL 논의가 훨씬 구체적이 된다.

원본 강의를 볼 때는 어떤 DSL을 좋다고 말하는지보다, 어떤 실패 사례를 보여주는지 보자. DSL의 철학은 잘 될 때보다 안 될 때 더 잘 드러난다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_077.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_077.md

1. 지금 확인된 자료

local notes에는 speakers가 Muhammad Awad, Muhammad Osama, Brandon Potter로 기록되어 있고, GitHub lecture folder와 Iris PDF가 있다. slide에는 Github.com/ROCm/iris가 보이며, RDMA support는 work in progress라고 적혀 있다. transcript와 code link는 notes에는 없다.

slide의 첫 메시지는 분명하다. multi-GPU programming을 Triton에서 first-class citizen으로 만들고 싶다는 것이다. familiar PyTorch/Triton-like API, host/device abstraction, SHMEM-like remote memory access가 핵심 단어로 보인다.

2. 기존 방식의 불편함

slide는 contemporary approach를 CPU가 communication을 시작하고, host-device synchronization과 bulk-synchronous phase가 생기는 방식으로 그린다. GPU kernel 사이에 RCCL 같은 communication kernel이 들어가고, stream synchronization 때문에 remote communication이 일찍 시작되지 못할 수 있다는 문제의식을 보여 준다.

Iris가 겨냥하는 자리는 이 control path다. communication을 host가 큰 덩어리로 지시하는 대신, GPU 쪽 computation 흐름 가까이에서 remote memory operation과 synchronization을 표현하려는 방향이다.

3. API를 볼 때의 질문

slide는 Iris가 Triton, PyTorch, Pythonic API를 강조한다고 말한다. 좋은 API라면 programmer가 device rank, symmetric heap, remote pointer, signal/wait 같은 개념을 너무 낯설지 않게 다룰 수 있어야 한다. 동시에 communication ordering을 숨기기만 하면 correctness bug가 생길 수 있다.

따라서 Iris를 읽을 때는 예제 code의 짧음보다 semantics를 먼저 보자. remote write가 언제 보이는지, synchronization primitive가 어떤 guarantee를 주는지, deadlock을 어떻게 피하는지, memory allocation이 모든 GPU에서 어떻게 맞춰지는지 확인해야 한다.

4. 성능에서 조심할 점

fine-grained communication은 잘 맞으면 computation overlap을 만든다. 하지만 너무 잘게 쪼개면 latency와 synchronization overhead가 커진다. collective library가 bulk transfer에서 잘하는 일을 Iris-style RMA가 항상 이기는 것은 아니다.

벤치마크를 볼 때는 message size, GPU topology, PCIe/NVLink/Infinity Fabric 경로, scale-up과 scale-out의 차이를 구분해야 한다. slide가 RDMA를 WIP로 표시했다면 node 밖 통신에 대한 결론은 특히 조심해서 읽어야 한다.

5. 혼자 해볼 것

두 GPU에서 block GEMM을 한다고 상상하고, 한 GPU의 tile이 다른 GPU의 partial result를 언제 필요로 하는지 diagram으로 그려 보자. conventional collective phase와 Iris-style fine-grained remote access를 나란히 그리면 어떤 경우에 overlap이 생길 수 있는지 보인다.

원본 slide를 볼 때는 Iris가 Triton에 추가하는 primitive를 표로 만들자. allocation, remote access, synchronization, locking, example kernel을 분리해 적으면 "multi-GPU programming made easier"가 실제로 무엇을 쉽게 만드는지 확인할 수 있다.

Source Links

• YouTube
• Slides folder
• Slides PDF
• GPU Mode lectures index
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_078.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_078.md

1. 지금 확인된 자료

local notes에는 Mengdi Wu와 Xinhao Cheng이 speaker로 기록되어 있다. 자료 링크로는 GPU Mode YouTube, lecture_079/Mirage-2025-09-13.pdf, 그리고 mirage-project/mirage repo가 있다. transcript는 공개 caption 경로에서 가져오지 못했다.

따라서 이 reader는 Mirage의 내부 구현을 단정하지 않는다. "MPK", "mega-kernel", "Compiling LLMs"라는 제목이 가리키는 설계 질문을 읽기 지도처럼 정리한다. 성능 배수, 지원 model, compiler pass 이름 같은 것은 slide나 repo에서 직접 확인해야 한다.

2. 왜 mega-kernel인가

LLM inference나 training path는 많은 operator로 이뤄진다. operator마다 kernel을 따로 launch하고 intermediate tensor를 memory에 남기면, 계산 자체보다 launch overhead와 memory 왕복이 눈에 띌 수 있다. fusion은 이 비용을 줄이려는 오래된 전략이다.

mega-kernel이라는 말은 그 fusion의 범위를 더 크게 상상하게 만든다. 단순히 bias + activation을 붙이는 정도가 아니라, LLM의 여러 단계가 한 kernel 안에서 이어질 수 있는지 묻는 것이다. 이 질문은 멋있지만 위험하다. kernel이 커질수록 register pressure, shared memory 사용, occupancy, scheduling이 모두 어려워진다.

3. compiler가 해야 하는 일

"컴파일해서 mega-kernel을 만든다"는 말은 사람이 손으로 긴 CUDA 파일을 쓰는 것과 다르다. compiler는 graph 수준의 의존성, tile 수준의 data reuse, thread/block 수준의 mapping을 함께 봐야 한다. 하나만 잘해도 전체가 빠르지 않을 수 있다.

이 지점에서 Mirage를 볼 때의 질문은 단순하다. search space를 어떻게 제한하는가? 후보 kernel이 맞는지 어떻게 검증하는가? dynamic shape나 여러 batch size가 들어오면 어떻게 대응하는가? notes에는 이 답이 없다. 하지만 repo와 slide를 읽을 때는 이 세 질문을 들고 들어가면 길을 덜 잃는다.

4. 성능보다 먼저 correctness

큰 fusion은 debugging을 어렵게 만든다. 작은 kernel 여러 개라면 중간 tensor를 찍어 볼 수 있지만, mega-kernel은 중간 상태가 kernel 안에 숨어 있을 수 있다. 그래서 reference implementation과의 비교, dtype별 tolerance, edge shape test가 더 중요해진다.

LLM 쪽에서는 "거의 맞는" 출력도 문제다. low precision, softmax, normalization, quantization이 섞이면 오차가 어디서 커지는지 찾기 어렵다. Mirage repo를 직접 만진다면 첫 실험은 speedup table이 아니라 작은 model 또는 작은 tensor에서 correctness envelope를 만드는 편이 낫다.

5. 혼자 해볼 것

작은 MLP block을 PyTorch eager로 만들고, torch.compile 또는 Triton kernel fusion이 어디까지 합쳐 주는지 관찰해 보자. profiler timeline에서 kernel launch 개수와 memory traffic을 먼저 본다. 그다음 "이 경계를 하나 더 없애려면 어떤 정보가 필요할까?"라고 적어 본다.

Mirage를 읽을 때는 repo의 hello world, 지원하는 model/operator, generated code를 찾자. slide에 benchmark가 있다면 input shape, hardware, baseline, compile time을 같이 적는다. mega-kernel은 빠른 숫자 하나보다 "어떤 조건에서 가능한가"가 더 중요하다.

Source Links

• YouTube
• Slides folder
• Slides PDF
• Code · mirage-project/mirage
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_079.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_079.md

1. 지금 확인된 자료

local notes에는 YouTube 링크와 Modal blog 링크가 있다. GPU Mode lecture repo folder와 slide는 확인되지 않았다. transcript fetch도 실패했다. 따라서 이 글은 "FA4가 정확히 무엇을 한다"가 아니라 "FA4 설명을 따라갈 준비"에 가깝다.

FlashAttention이라는 이름은 이미 많은 배경을 요구한다. attention score matrix를 통째로 materialize하지 않고, tile 단위로 softmax state와 output을 갱신하는 감각이 먼저 있어야 한다. FA4라는 숫자는 새 버전의 세부 구현을 가리키겠지만, notes만으로는 그 세부를 안전하게 옮길 수 없다.

2. attention은 왜 memory 문제가 되는가

scaled dot-product attention은 말로는 간단하다. Q와 K를 곱하고, softmax를 하고, V를 곱한다. 하지만 sequence length가 길어지면 score matrix가 커지고, 이를 HBM에 쓰고 다시 읽는 비용이 커진다.

FlashAttention 계열의 기본 감각은 이 큰 중간 결과를 피하는 것이다. block 단위로 score를 계산하고, online softmax에 필요한 running max와 normalization state를 유지하면서 output을 누적한다. 이 설명은 FlashAttention family에 대한 일반 배경이며, FA4의 구체 구현 claim은 아니다.

3. FA4에서 확인할 질문

FA4 설명을 읽을 때는 "새로운 알고리즘인가, 새로운 kernel engineering인가, 둘 다인가"를 먼저 나누자. attention의 수학은 그대로인데 hardware mapping, pipeline, layout, instruction 선택이 바뀌었을 수도 있다. 반대로 masking, grouped-query attention, variable length 같은 workload 요구가 중심일 수도 있다.

Modal 글을 볼 때는 benchmark 숫자만 따라가지 말고 baseline과 조건을 확인한다. GPU generation, dtype, head dimension, sequence length, causal 여부, batch shape가 바뀌면 attention kernel의 승패가 달라진다. 빠른 kernel은 항상 특정 조건과 함께 읽어야 한다.

4. CuTe와 kernel 표현

notes는 CuTe도 핵심 term으로 잡고 있다. CuTe는 tiled GPU kernel을 표현할 때 layout과 tensor view를 정교하게 다루는 도구다. FlashAttention 같은 kernel은 "무엇을 계산하는가"보다 "어떤 tile을 언제 어디에 놓는가"가 성능을 크게 좌우하므로, CuTe류 abstraction과 자연스럽게 만난다.

다만 이 강의가 CuTe를 얼마나 깊게 다뤘는지는 알 수 없다. 원본을 볼 때는 FA4 kernel이 layout, MMA shape, memory movement를 어떤 언어로 설명하는지 주목하자. 이름은 어렵지만 질문은 친숙하다. data를 덜 움직이고, 움직일 때는 hardware가 좋아하는 모양으로 움직이는가?

5. 혼자 해볼 것

작은 attention을 PyTorch로 구현하고, score matrix를 실제로 저장하는 버전과 block 단위로 누적하는 toy 버전을 비교해 보자. 성능을 production kernel과 비교하려는 실험이 아니라, memory footprint가 어떻게 달라지는지 보는 실험이다.

그다음 Modal article을 읽으며 FA4가 어느 층에서 새로워지는지 표시한다. algorithm, layout, pipeline, hardware instruction, benchmark methodology로 줄을 나누면 "빠르다"가 아니라 "왜 빠른가"를 더 차분하게 볼 수 있다.

Source Links

• YouTube
• Code/article · Modal reverse-engineer FlashAttention 4
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_080.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_080.md

1. 세 단어만 확실하다

제목에서 확실히 보이는 단어는 purely functional, data-parallel, array programming이다. 이 조합은 고수준 array expression과 compiler lowering을 떠올리게 한다.

하지만 어떤 system을 다뤘는지는 아직 모른다. Futhark 같은 특정 언어를 말했을 수도 있고, 더 넓은 array programming 관점을 말했을 수도 있다. source가 없는 이름은 예로만 다루고 강의 내용으로 쓰지 않는다.

2. functional array가 GPU에 던지는 약속

GPU programming은 흔히 index, memory, mutation을 직접 다루는 일처럼 보인다. functional array style은 계산을 map, reduce, scan, stencil 같은 transformation으로 표현하려 한다.

그 약속은 compiler가 더 큰 식을 보고 fusion, parallelization, memory planning을 할 수 있다는 데 있다. 하지만 약속은 구현이 아니다. 실제 성능은 expression을 어떤 schedule과 tile, memory movement로 낮추는지에 달려 있다.

3. 순수함은 비용을 없애지 않는다

side effect가 적으면 의미는 깨끗해진다. race condition을 줄이고 compiler reasoning을 돕는 장점도 있다. 하지만 GPU 비용은 사라지지 않는다. memory traffic, synchronization, occupancy, register pressure는 여전히 계산서에 남는다.

좋은 functional array system은 사용자가 성능을 전혀 보지 않아도 되는 마법이 아니다. 오히려 깨끗한 의미를 유지하면서 필요한 곳에서는 schedule이나 layout에 개입할 수 있어야 한다.

4. 원본에서 표시할 포인트

강의를 볼 때는 먼저 primitive set을 적는다. map/reduce/scan만 다루는지, scatter/gather와 irregular indexing을 어떻게 표현하는지, in-place update가 필요한 알고리즘을 어떤 식으로 바꾸는지 본다.

다음은 inspection이다. generated kernel이나 IR을 볼 수 있는가. benchmark가 있다면 handwritten CUDA, Triton, vendor library, 또는 다른 DSL 중 무엇과 비교했는가. 이 비교 대상이 빠지면 "high-performance"라는 말이 떠 있게 된다.

5. 작은 실험감

y = reduce(x * x + 1) 같은 계산을 떠올려 보자. eager execution에서는 곱, 더하기, reduce가 각각 중간 tensor를 만들 수 있다. fused expression에서는 중간 값을 memory에 쓰지 않고 지나갈 가능성이 있다.

이 손실험의 목적은 특정 framework를 평가하는 것이 아니다. 고수준 식이 실제 memory traffic으로 바뀌는 순간을 보는 것이다. 원본 강의가 어떤 예제를 쓰는지는 영상을 보고 나서만 채운다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_081.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_081.md

1. 지금 확인된 자료

local notes에는 Jason Ansel, Oguz Ulgen, Will Feng이 speaker로 기록되어 있고 YouTube 링크만 있다. transcript는 실패했고, slides/code/paper/repo folder는 missing이다.

그래서 이 장은 Helion 소개를 흉내 내지 않는다. "high-level DSL for ML kernels"라는 제목을 읽는 데 필요한 질문표를 만든다. DSL은 코드가 짧아지는 도구이기도 하지만, 더 깊게는 성능 지식을 어떤 형태로 표현할지 정하는 언어다.

2. high-level의 약속

ML kernel을 직접 쓰는 일은 반복적인 괴로움이 있다. shape, stride, dtype, mask, vectorization, tile size, memory hierarchy를 모두 챙겨야 한다. high-level DSL은 이 중 일부를 compiler와 runtime에 맡기고, programmer가 operation의 의도를 더 크게 쓰게 하려 한다.

하지만 GPU에서 high-level은 공짜가 아니다. abstraction이 높아질수록 compiler가 더 많은 결정을 대신 해야 한다. 좋은 결정이면 생산성과 성능을 함께 얻는다. 나쁜 결정이면 사용자는 예쁜 코드로 느린 kernel을 얻게 된다. 예쁘게 망하면 디버깅이 더 서럽다.

3. DSL의 진짜 시험대

ML kernel DSL은 elementwise 예제에서는 좋아 보이기 쉽다. 진짜 시험은 attention, normalization, quantization, matmul epilogue, ragged shape처럼 실제 workload의 지저분함을 만날 때다. 여기서 DSL이 너무 좁으면 production에 못 들어간다.

또 하나의 시험은 fallback과 inspection이다. kernel이 느리거나 틀렸을 때 generated code, IR, profiler trace를 볼 수 있어야 한다. high-level DSL이 black box가 되면 성능 엔지니어는 답답해진다. 좋은 도구는 높은 문과 낮은 지하실을 둘 다 갖고 있다.

4. Helion을 볼 때의 질문

원본 강의를 볼 때는 Helion이 어떤 사용자를 겨냥하는지 확인하자. CUDA 전문가의 반복 작업을 줄이는 도구인지, PyTorch 사용자에게 kernel authoring을 열어 주는 도구인지, compiler 연구 플랫폼인지에 따라 좋은 설계가 달라진다.

그다음 실제 예제를 본다. DSL syntax가 어떤 정보를 요구하는가? autotuning을 제공하는가? PyTorch tensor와 어떻게 연결되는가? correctness test는 어떻게 쓰는가? notes에는 답이 없지만, 이 질문들은 강의를 보는 동안 곁에 둘 만하다.

5. 혼자 해볼 것

작은 layernorm이나 fused bias-GELU를 하나 골라 CUDA, Triton, 가능한 high-level DSL 스타일 pseudo-code로 나란히 써 보자. 코드 길이보다 "성능에 필요한 정보를 어디에 적는가"를 비교한다.

Helion 자료를 볼 때는 benchmark 표가 나오면 바로 믿지 말고, shape와 baseline을 같이 적는다. DSL은 언어라서, 최고의 예문보다 어색한 문장을 어떻게 처리하는지가 더 중요하다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_082.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_082.md

1. 지금 확인된 자료

local notes에는 YouTube 링크만 있고 speaker는 missing이다. transcript fetch는 실패했다. slides와 code도 없다. 그러므로 이 글은 특정 formal system이나 논문을 설명하지 않는다.

대신 kernel derivation이라는 표현 자체를 붙잡는다. derivation은 "결과 코드를 갑자기 보여주는 것"이 아니라, naive specification에서 optimized kernel까지 가는 길을 단계별 변환으로 설명한다는 뜻으로 읽을 수 있다.

2. 왜 formalized가 필요한가

GPU kernel 최적화는 자주 이야기식으로 전해진다. tile을 잡고, memory를 줄이고, warp를 맞추고, occupancy를 본다. 숙련자에게는 자연스럽지만, 배우는 사람에게는 어느 순간 마술처럼 보인다.

formalized derivation은 이 마술을 줄이려는 시도다. 먼저 수학적 specification을 둔다. 그다음 loop transformation, tiling, fusion, layout change, parallel mapping을 하나씩 적용한다. 각 단계가 의미를 보존한다고 말할 수 있으면, 최종 kernel은 "빠르지만 수상한 코드"가 아니라 "유도된 코드"가 된다.

3. FlashAttention을 예로 떠올리기

notes에는 FlashAttention이 관련 term으로 감지되어 있다. 강의에서 실제로 어떤 예제를 썼는지는 확인할 수 없지만, FlashAttention은 derivation 관점으로 읽기 좋은 대상이다. naive attention에서 시작해 score matrix materialization을 없애고, online softmax를 도입하고, tile 단위 accumulation으로 바꾸는 흐름이 있기 때문이다.

다만 이것을 강의 내용으로 단정하면 안 된다. 원본을 볼 때 FlashAttention이 실제 예제인지, 아니면 notes의 keyword 감지에만 잡힌 것인지 확인하자. 확인되기 전까지는 "가능한 읽기 포인트"로만 둔다.

4. formal과 practical 사이

formal derivation은 correctness를 다루기 좋지만, GPU 성능은 hardware detail을 피할 수 없다. 같은 의미를 보존하는 변환이라도 register pressure가 너무 커지거나 memory coalescing이 깨지면 느리다. formal system이 useful하려면 cost model이나 schedule 선택과 만나는 지점이 있어야 한다.

또 numerics도 조심해야 한다. floating point에서는 연산 순서를 바꾸면 bitwise 결과가 달라질 수 있다. 따라서 "의미 보존"이 정확히 무엇을 뜻하는지, tolerance와 dtype policy가 어떻게 정의되는지 확인해야 한다.

5. 혼자 해볼 것

matrix transpose나 softmax처럼 작은 kernel 하나를 골라보자. 먼저 naive specification을 쓴 뒤, memory coalescing을 위해 loop order를 바꾸고, tile을 넣고, shared memory를 쓰는 과정을 한 줄씩 기록한다. 각 단계마다 왜 같은 결과라고 믿는지 적는다.

원본 강의를 볼 때는 proof나 formal notation이 어렵더라도 겁먹지 말자. 핵심은 빠른 코드를 천재의 손맛으로만 남기지 않고, 재현 가능한 변환의 줄로 만드는 것이다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_083.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_083.md

1. 지금 확인된 자료

local notes에는 speaker가 Paulius Micikevicius로 기록되어 있고 YouTube와 lecture_084/numerics-and-ai.pdf 링크가 있다. code와 transcript는 없다.

Paulius Micikevicius라는 이름 때문에 mixed precision이나 AI numerics를 떠올릴 수 있지만, 이 노트에서는 notes가 확인해 준 범위를 넘지 않는다. 다만 numerics라는 주제 자체는 GPU Mode 독자에게 아주 중요하다. 빠른 kernel이 틀린 숫자를 만들면, 빠르게 틀릴 뿐이다.

2. AI에서 숫자는 계산 결과만이 아니다

AI workload는 floating point format의 영향을 크게 받는다. FP32, TF32, FP16, BF16, FP8 같은 선택은 memory bandwidth, Tensor Core throughput, overflow/underflow, training stability와 얽힌다. inference에서도 quantization과 accumulation precision이 output quality에 영향을 준다.

numerics를 성능 튜닝 뒤에 붙이는 체크박스로 보면 위험하다. kernel을 설계할 때부터 어떤 값이 큰 dynamic range를 갖는지, 어디서 accumulation이 일어나는지, loss scaling이나 normalization이 필요한지 생각해야 한다.

3. 낮은 precision의 약속과 비용

낮은 precision은 memory traffic을 줄이고 Tensor Core를 더 잘 쓰게 해 준다. 그래서 AI system에서 매력적이다. 하지만 낮은 precision은 표현 가능한 숫자의 범위와 간격을 바꾼다. 작은 차이가 사라지거나, 큰 값이 overflow될 수 있다.

중요한 것은 format 이름을 외우는 게 아니다. 어떤 tensor를 낮은 precision으로 저장해도 되는지, 어떤 연산은 더 높은 precision으로 accumulate해야 하는지 구분하는 감각이다. 특히 softmax, reduction, optimizer state처럼 누적과 scale이 얽힌 곳은 조심해야 한다.

4. kernel 개발자가 볼 포인트

numerics는 framework 사용자만의 문제가 아니다. custom CUDA/Triton kernel을 쓰면 rounding, accumulation order, fast math option, atomic operation의 nondeterminism까지 직접 마주한다. reference와 비교할 때도 bitwise match만 고집할지, 상대/절대 오차를 어떻게 둘지 정해야 한다.

slide를 볼 때는 format별 장단점만 보지 말고, 실제 workflow를 찾자. training에서 어떤 test를 통과해야 하는가? inference에서 accuracy regression을 어떻게 잡는가? benchmark는 latency와 quality를 같이 보여 주는가? 숫자의 세계는 조용하지만, bug가 나면 아주 시끄럽다.

5. 혼자 해볼 것

작은 matmul과 softmax를 FP32, FP16, BF16으로 실행해 보고 FP32 reference와 비교해 보자. random input만 쓰지 말고 큰 값, 작은 값, 거의 같은 값이 섞인 input을 넣는다. error가 shape와 value range에 따라 어떻게 달라지는지 보는 것이 목표다.

원본 slide를 읽을 때는 "어떤 format이 빠르다"보다 "어떤 상황에서 안전한가"를 표시하자. 좋은 numerics 공부는 hardware 표와 model 품질 사이에 다리를 놓는다.

Source Links

• YouTube
• Slides folder
• Slides PDF
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_084.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_084.md

1. 지금 확인된 자료

local notes에는 speaker가 Jack Hopkins로 기록되어 있고 YouTube 링크만 있다. transcript fetch는 실패했다. slides, code, paper, repo folder는 missing이다.

GPU Mode의 다른 강의보다 kernel 중심성은 낮아 보인다. 그래도 learning environment는 ML systems 관점에서 중요하다. RL은 model만으로 굴러가지 않는다. 환경 실행, state 수집, action 적용, reward 계산, logging, reproducibility가 모두 system이다.

2. 왜 Factorio인가

Factorio는 단순한 grid world가 아니다. 자원, 제작, 운송, 계획, 장기 목표가 얽힌다. 이런 환경은 agent가 짧은 반응만 잘해서는 해결하기 어렵고, planning과 tool use, long-horizon credit assignment를 요구할 수 있다.

하지만 이 문단은 Factorio 일반 특성에 대한 설명이지, 강의가 실제로 무엇을 보여줬다는 claim은 아니다. 원본을 볼 때는 Factorio Learning Environment가 어떤 task abstraction을 제공하는지, observation과 action space가 어떻게 정의되는지 확인해야 한다.

3. learning environment의 어려움

좋은 environment는 재미있는 task만 있으면 되는 것이 아니다. 빠르게 많이 돌릴 수 있어야 하고, deterministic replay가 가능해야 하며, failure를 분석할 수 있어야 한다. agent가 실패했을 때 "멍청해서"인지 "환경 wrapper가 이상해서"인지 구분할 수 있어야 한다.

복잡한 환경일수록 reward 설계도 조심스럽다. reward가 너무 촘촘하면 agent가 편법을 찾고, 너무 sparse하면 학습 신호가 약하다. benchmark로 쓰려면 task 난이도와 평가 metric도 명확해야 한다.

4. GPU Mode 독자가 볼 포인트

GPU kernel 강의는 아니지만, systems 독자에게는 throughput이 여전히 중요하다. RL/OpenEnv 계열 workload에서는 environment step이 병목인지, model inference가 병목인지, rollout storage가 병목인지 나눠 봐야 한다.

원본 강의를 볼 때는 환경이 local simulation인지, remote service인지, parallel rollout을 어떻게 지원하는지 확인하자. 또 language agent를 염두에 둔 환경인지, classic RL policy를 위한 환경인지에 따라 API 모양이 달라질 수 있다.

5. 혼자 해볼 것

간단한 toy environment를 하나 만들고, random policy로 초당 step 수를 재 보자. 그다음 observation serialization, reward 계산, logging을 하나씩 켜며 어디서 느려지는지 본다. RL 시스템은 종종 model보다 주변 배관에서 시간이 샌다.

원본을 볼 때는 demo의 멋짐보다 evaluation protocol을 보자. seed, task set, success metric, replay/debug tool이 있으면 환경은 연구 도구에 가까워진다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_085.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_085.md

1. 지금 확인된 자료

local notes에는 speaker가 Vicki Wang으로 기록되어 있고 YouTube, lecture folder, cute_dsl_introduce.pdf가 있다. code link는 notes에 없다.

CuTe는 CUTLASS 계열에서 layout과 tiled tensor를 표현하는 중요한 abstraction으로 알려져 있다. 하지만 이 강의가 어떤 버전의 CuTe DSL, 어떤 syntax, 어떤 exercise를 다뤘는지는 transcript 없이 확정할 수 없다. 여기서는 입문자가 붙잡을 mental model을 정리한다.

2. thread보다 tile을 먼저 본다

CUDA를 처음 배울 때는 block, thread, index 계산이 앞에 나온다. CuTe류 abstraction은 관점을 조금 바꾼다. data가 어떤 layout을 갖고, 그 layout을 어떤 tile로 나누며, tile이 hardware operation에 어떻게 mapping되는지 먼저 생각하게 만든다.

이 관점은 matmul 같은 kernel에서 특히 유용하다. 성능은 단순히 thread 수가 아니라 data reuse, memory coalescing, MMA shape, shared memory staging으로 결정된다. CuTe DSL은 이런 구조를 코드에서 더 명시적으로 다루게 하려는 도구로 읽을 수 있다.

3. layout은 장식이 아니다

GPU kernel에서 layout은 성능 그 자체에 가깝다. 같은 값이라도 memory에 놓인 순서와 접근 패턴이 다르면 coalescing, bank conflict, cache behavior가 달라진다. CuTe를 배울 때 layout notation이 낯설어도 피하면 안 된다.

좋은 연습은 작은 2D tensor를 손으로 그리는 것이다. logical coordinate가 어떤 stride를 거쳐 address가 되는지, tile을 자르면 어떤 thread가 어떤 element를 읽는지 표시해 보자. 처음에는 느리지만, 이 그림이 없으면 나중에 profiler 숫자가 외계어처럼 보인다.

4. 입문에서 조심할 점

CuTe DSL은 CUDA보다 high-level처럼 보일 수 있지만, low-level detail을 없애지는 않는다. 오히려 detail을 다른 문법으로 드러낸다. 그래서 "문법을 외우기"보다 "layout, tile, copy, mma가 각각 무엇을 책임지는지"를 분리해서 보는 편이 좋다.

원본 slide를 볼 때는 예제 하나를 끝까지 따라가자. input/output shape, tile shape, shared memory staging, final store까지 끊지 말고 이어야 한다. 중간에 모르는 symbol이 나오면 그 symbol이 performance decision인지 단순 type annotation인지 구분한다.

5. 혼자 해볼 것

작은 GEMM을 골라 naive PyTorch, Triton, CuTe DSL pseudo-structure로 비교해 보자. 실제 구현 전이라도 tile shape와 memory movement를 그림으로 그리면 충분히 얻는 게 있다.

가능하다면 slide의 첫 예제를 그대로 실행 가능한 환경에서 돌려 보고, generated code나 profiler trace를 확인하자. CuTe는 눈으로만 보면 추상적이고, 한 번 깨져 보면 갑자기 현실적이다.

Source Links

• YouTube
• Slides folder
• Slides PDF
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_086.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_086.md

1. 지금 확인된 자료

local notes에는 speaker가 Prajwal Singhania로 기록되어 있고 YouTube 링크만 있다. transcript fetch는 실패했다. slides와 code는 missing이다. notes는 kernel과 NVSHMEM을 핵심 term으로 잡았다.

NVSHMEM은 GPU 중심의 partitioned global address space 통신 library로 설명할 수 있다. 다만 이 강의가 어떤 primitive나 kernel 예제를 다뤘는지는 알 수 없다. 그러니 여기서는 API 사용법보다, NVSHMEM 이야기를 들을 때 어떤 성능과 correctness 질문을 적어야 하는지에 집중한다.

2. 왜 communication kernel인가

multi-GPU program에서는 계산 kernel만 빠르다고 끝나지 않는다. GPU 사이에 activation, gradient, KV cache, partial result가 오가야 한다. 이때 CPU가 큰 collective call을 순서대로 넣는 방식은 단순하지만, fine-grained overlap을 놓칠 수 있다.

communication을 kernel 가까이로 가져오면 GPU가 계산 흐름 안에서 remote memory operation이나 synchronization을 더 직접 표현할 수 있다. 잘 맞으면 latency를 줄이고 overlap을 늘릴 수 있다. 잘못 쓰면 ordering bug와 deadlock이 기다린다. 문은 낮아져도 바닥은 깊다.

3. latency를 볼 때의 함정

low latency communication은 작은 message에서 특히 중요하다. 하지만 latency만 보면 bandwidth를 놓치고, bandwidth만 보면 synchronization cost를 놓친다. message size, GPU topology, intra-node/inter-node 여부, stream interaction을 같이 봐야 한다.

NVSHMEM 같은 도구는 programmer에게 더 많은 통제권을 준다. 그만큼 memory ordering, visibility, barrier semantics를 정확히 알아야 한다. "작동했다"와 "항상 맞다"는 분산 GPU 코드에서 꽤 먼 사이일 수 있다.

4. LLM workload와 연결하기

LLM training에서는 tensor parallelism, pipeline parallelism, optimizer state sharding 같은 곳에서 communication이 중요하다. inference에서도 multi-GPU serving, KV cache placement, expert parallelism이 communication path를 만든다.

이 강의를 볼 때는 어떤 workload를 target으로 삼는지 확인하자. collective 대체인지, point-to-point RMA인지, custom kernel 안의 fine-grained synchronization인지에 따라 평가 기준이 다르다. NCCL과 비교한다면 어떤 조건에서 비교했는지도 꼭 봐야 한다.

5. 혼자 해볼 것

두 GPU가 있다고 가정하고, 한 GPU의 tile 결과를 다른 GPU가 언제 필요로 하는지 timeline을 그려 보자. CPU가 collective을 launch하는 방식과 device-side communication이 가능한 방식을 나란히 그리면 overlap의 가능성과 위험이 같이 보인다.

실제 실험을 한다면 작은 ping-pong latency test부터 시작하자. message size를 바꾸고, synchronization 위치를 바꾸며 측정한다. distributed kernel은 큰 model보다 작은 실험에서 먼저 진실을 말해 준다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_087.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_087.md

1. Tiny가 무엇을 줄였는지 아직 모른다

TinyTPU라는 이름은 작고 교육적인 accelerator를 떠올리게 한다. 하지만 작다는 것이 hardware size인지, ISA인지, compiler stack인지, tutorial scope인지는 원본을 봐야 한다.

이 노트는 특정 구현을 설명하지 않는다. 대신 작은 accelerator 강의를 들을 때 어떤 질문을 던지면 구조가 보이는지 적는다. source가 보강되면 이 틀 위에 실제 설계와 예제를 올리면 된다.

2. 작은 accelerator가 좋은 교재가 되는 이유

대형 accelerator는 많은 층이 동시에 움직인다. systolic array, on-chip buffer, DMA, instruction scheduling, compiler tiling, host runtime이 한꺼번에 나온다.

작은 모형은 그 층을 손으로 따라갈 수 있게 만든다. 값이 array 안에서 어떻게 이동하는지, buffer가 왜 모자라는지, compiler가 왜 tile을 잘라야 하는지가 더 잘 보인다. GPU 독자에게도 이 dataflow 감각은 유용하다.

3. GPU와 비교할 때 조심할 점

GPU는 general-purpose parallel processor에 가깝고, TPU-like design은 matrix operation을 위한 dataflow를 앞에 두는 경우가 많다. 둘 다 병렬 계산을 하지만 programmer가 보는 손잡이가 다르다.

따라서 "GPU보다 빠른가"를 첫 질문으로 두면 놓치는 것이 많다. array 크기, supported dtype, buffer capacity, instruction granularity가 어떤 programming model을 강제하는지 먼저 봐야 한다.

4. compiler와 simulator의 자리

작은 accelerator라도 compiler가 없으면 사용하기 어렵다. tensor program을 tile로 나누고, memory에 맞게 배치하고, 실행 순서를 정해야 한다.

원본을 볼 때는 TinyTPU가 어떤 frontend를 받는지, IR이나 instruction stream을 보여 주는지, simulator가 cycle count나 memory traffic을 어떻게 드러내는지 확인한다. 이 정보가 있어야 "작다"가 단순한 장난감인지, 학습 가능한 system인지 판단할 수 있다.

5. 종이에 먼저 그려 보기

4x4 systolic array를 그리고 4x4 matrix multiply가 몇 cycle에 어떤 값들을 움직이는지 손으로 따라가 보자. weight-stationary나 output-stationary 같은 이름은 그 다음에 붙여도 늦지 않다.

원본에서는 speaker가 실제로 어떤 dataflow를 썼는지만 표시한다. 이 강의가 hardware, simulator, compiler 중 어디에 무게를 두는지는 아직 비워 둔다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_088.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_088.md

1. 지금 확인된 자료

local notes에는 Mehdi Amini와 Jared Roesch가 speaker로 기록되어 있고 YouTube 링크만 있다. transcript fetch는 실패했다. slides, code, paper, repo folder는 missing이다.

따라서 cuTile의 API, NVIDIA 내부/외부 프로젝트 여부, CuTe나 MLIR과의 관계는 단정하지 않는다. 제목의 "(from friends at NVIDIA)"도 배경 힌트일 뿐, 기술 내용을 대신하지 않는다.

2. tile은 GPU 사고의 기본 단위

GPU 성능은 큰 tensor를 작은 작업 조각으로 나누는 방식에 크게 좌우된다. tile은 data reuse를 만들고, memory transaction을 정렬하고, parallel work를 hardware에 맞추는 단위가 된다. 그래서 많은 DSL과 library가 tile abstraction을 중심에 둔다.

하지만 tile은 단어 하나로 끝나지 않는다. logical tile, memory tile, thread tile, MMA tile이 서로 다를 수 있다. 좋은 abstraction은 이 차이를 숨길 때와 드러낼 때를 잘 고른다.

3. abstraction의 장점과 빚

cuTile 같은 도구가 유용하려면 반복적인 index 계산과 layout 실수를 줄여야 한다. programmer가 "이 tile을 이 방식으로 옮기고 계산한다"를 더 직접 표현할 수 있으면, kernel code는 읽기 쉬워질 수 있다.

그 대신 abstraction은 빚을 만든다. generated code가 느릴 때 이유를 추적할 수 있어야 하고, corner case에서 어떤 memory access가 생기는지 볼 수 있어야 한다. tile abstraction이 profiler와 이어지지 않으면, 예쁜 문법 뒤에서 시간이 새어 나간다.

4. NVIDIA toolchain 안에서 볼 질문

원본 강의를 볼 때는 cuTile이 어느 층의 도구인지 확인하자. CUDA C++ library인지, Python DSL인지, MLIR dialect인지, CUTLASS/CuTe와 이어지는 component인지에 따라 해석이 달라진다.

또 target workload를 보자. GEMM, convolution, attention, stencil, collective communication 중 어디를 겨냥하는가? tile abstraction은 범용일수록 어려워지고, 특화될수록 강해진다. 어느 쪽을 택했는지가 중요하다.

5. 혼자 해볼 것

간단한 tiled matmul을 손으로 설계해 보자. block tile, warp tile, MMA tile을 따로 적고, 각 단계에서 어떤 data가 global memory에서 shared memory로, 다시 register로 이동하는지 표시한다.

cuTile 자료를 볼 때는 API 이름보다 예제가 드러내는 decision을 보자. tile 크기는 누가 정하는가? layout 변환은 명시적인가? hardware generation이 바뀌면 code도 바뀌는가? 이 질문들이 abstraction의 실제 무게를 알려 준다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_089.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_089.md

1. 강의 요약이 아니라 준비 노트다

이 파일은 원본 강의의 내용을 재현하지 못한다. caption이 비어 있고, 발표 자료도 없다. speaker가 어떤 incident, project, tool, career advice를 이야기했는지는 영상 확인이 필요하다.

그럼에도 GPU Mode reader에 둘 이유는 있다. Part VII의 주제처럼, 시스템 공부는 kernel syntax만으로 끝나지 않는다. 오래 가는 ML engineering은 문제를 작게 만들고, 숫자를 읽고, 바뀐 환경에서 다시 세우는 능력에 가깝다.

2. resilient skill은 API 암기가 아니다

ML stack은 계속 움직인다. PyTorch, CUDA, compiler, driver, distributed runtime, model architecture가 모두 버전을 바꾼다. 특정 API 하나만 붙잡으면 금방 낡는다.

덜 낡는 것은 습관이다. 작은 reproduction을 만드는 습관, baseline을 남기는 습관, profiler와 log를 같이 보는 습관, release note를 읽는 습관, 고친 뒤 regression을 막는 습관이다. 이 목록은 강의 인용이 아니라 제목을 읽기 위한 준비 프레임이다.

3. debugging은 기술 지식의 압축이다

좋은 debugging은 단순히 에러를 없애는 일이 아니다. 어떤 층에서 문제가 생겼는지 가설을 세우고, 그 가설을 싸게 검증하는 일이다.

GPU/ML systems에서는 이 층이 많다. numerical issue인지, memory pressure인지, kernel launch overhead인지, communication wait인지, data pipeline인지 분리해야 한다. resilient skill은 많은 도구를 아는 것보다 이 분리 능력에서 나온다.

4. 임시 해결도 기록이 있어야 산다

ML engineering에서 workaround는 피하기 어렵다. deadline은 있고, GPU time은 비싸며, dependency는 기다려 주지 않는다. 문제는 임시 해결 자체가 아니라, 왜 들어갔는지 나중에 아무도 모르는 상태다.

원본 강의가 이 주제를 다뤘는지는 아직 알 수 없다. 하지만 reader가 가져갈 질문은 분명하다. fix를 넣을 때 어떤 metric, comment, test, issue가 같이 남아야 팀이 다음 변경을 견딜 수 있는가.

5. 원본을 볼 때 남길 기록

영상을 볼 때 speaker가 드는 실제 사례를 중심으로 적는다. tool 이름보다 사고 순서가 더 중요하다. 증상은 무엇이었고, 어떤 관측으로 범위를 좁혔고, 마지막에 어떤 guardrail을 남겼는지 본다.

개인 연습으로는 최근 ML bug 하나를 postmortem처럼 써 보는 것이 좋다. 증상, 재현 방법, 측정값, 틀린 가설, 실제 원인, 방지 test를 한 페이지에 담으면 이 강의의 제목이 훨씬 덜 추상적으로 느껴진다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_090.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_090.md

1. 지금 확인된 자료

local notes에는 YouTube 링크와 code links가 있다. speaker는 missing이다. code link로는 meta-pytorch/OpenEnv의 GPU Mode tutorial, unslothai/unsloth, huggingface/trl, 그리고 Colab notebook로 보이는 링크가 기록되어 있다. transcript fetch는 실패했다.

그래서 이 장은 RL algorithm lecture로 쓰지 않는다. PPO, GRPO, DPO 같은 구체 알고리즘을 강의가 다뤘다고 말하지 않는다. notes에서 확인되는 범위 안에서 RL/OpenEnv 실험을 시스템으로 읽는 법을 정리한다.

2. agent 실험은 환경이 반이다

RL이나 agent training은 model code만으로 성립하지 않는다. environment가 있고, observation과 action의 contract가 있고, reward 또는 evaluation signal이 있다. 이 contract가 흐리면 좋은 model과 나쁜 실험을 구분하기 어렵다.

OpenEnv라는 이름은 환경을 열어 두고 반복 가능한 실험 단위로 만들려는 방향을 떠올리게 한다. 원본을 볼 때는 environment API, task definition, reset/step/evaluate 흐름을 먼저 확인하자. agent가 멋져 보여도 환경 contract가 흐리면 결과는 오래 못 간다.

3. training stack을 볼 때의 질문

notes에는 Unsloth와 Hugging Face TRL 링크가 있다. 이것은 fine-tuning과 RL training workflow가 연결될 가능성을 보여 주지만, 세부는 확인 필요다. 어떤 model size, 어떤 optimizer, 어떤 reward source를 썼는지는 notes에 없다.

GPU Mode 관점에서는 throughput과 reproducibility를 같이 본다. rollout generation이 병목인지, training step이 병목인지, environment step이 병목인지 나눠야 한다. 또한 checkpoint, seed, dataset/environment version을 남겨야 같은 실험을 다시 할 수 있다.

4. agent demo와 benchmark 사이

agent demo는 설득력이 강하다. 화면에서 뭔가 해내면 좋아 보인다. 하지만 benchmark로 쓰려면 더 차가운 장치가 필요하다. task set, success metric, failure category, random seed, evaluation budget을 정해야 한다.

OpenEnv 자료를 볼 때는 "무엇을 쉽게 실행하게 해 주는가"와 "무엇을 공정하게 비교하게 해 주는가"를 나눠 보자. 좋은 tutorial은 시작을 쉽게 만들고, 좋은 benchmark는 주장을 조심스럽게 만든다.

5. 혼자 해볼 것

작은 text-based environment를 하나 정하고, random agent와 simple heuristic agent를 같은 metric으로 비교해 보자. 그다음 training framework를 붙이기 전에 logging schema를 먼저 만든다. episode id, seed, prompt, action, reward, final outcome을 남기는 것만으로도 실험의 품질이 달라진다.

원본 Colab이나 tutorial을 열 때는 실행 성공보다 재현성을 먼저 본다. dependency version, GPU requirement, expected runtime, saved artifacts가 적혀 있는지 확인하자. agent 실험은 재미있을수록 장부가 필요하다.

Source Links

• YouTube
• Code · OpenEnv GPU Mode tutorial
• Code · Unsloth
• Code · Hugging Face TRL
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_091.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_091.md

1. 지금 손에 있는 것

local notes에는 speaker가 Loubna Ben Allal로 기록되어 있고 YouTube link만 있다. transcript fetch는 실패했고, slides/code/paper/repo folder는 없다. 그래서 이 강의가 Hugging Face식 practical training guide인지, small model pretraining recipe인지, fine-tuning 운영 playbook인지는 아직 단정할 수 없다.

다만 "smol training playbook"이라는 제목은 좋은 학습 태도를 알려 준다. 거대한 run 하나를 신앙처럼 기다리기보다, 작고 빠른 실험으로 가설을 부수고 고치는 방식이다. GPU 시스템 공부에서도 이 감각은 꽤 중요하다. 큰 cluster에서 틀리면 비싸고, 작은 run에서 틀리면 배운다.

2. 작은 훈련은 축소판이 아니라 실험 장치다

작은 모델은 production 모델의 미니어처가 아니다. capacity, data mixture, optimizer stability, communication overhead가 모두 다르게 나타난다. 그래서 작은 실험을 "그대로 성능 예측기"로 읽으면 위험하다.

대신 작은 훈련은 decision filter로 쓰는 편이 좋다. tokenization이 깨졌는가, loss curve가 정상적으로 내려가는가, dataloader가 병목인가, checkpoint/resume이 되는가, eval이 training objective와 싸우는가 같은 질문은 작은 run으로도 많이 걸러진다.

3. GPU 관점에서 먼저 볼 신호

training playbook은 model quality 이야기처럼 보이지만, 시스템 쪽에서는 먼저 throughput과 utilization을 본다. step time이 들쭉날쭉하면 학습 곡선도 해석하기 어렵다. dataloader, host-to-device copy, activation checkpointing, optimizer step, logging이 어디서 시간을 쓰는지 분리해야 한다.

작은 run에서는 GPU가 너무 작게 먹이를 받아 underutilized될 수도 있다. 이때 "작은 모델이 느리다"가 아니라 batch size, sequence length, kernel launch overhead, CPU preprocessing 중 무엇이 dominant인지 보아야 한다. 작은 실험의 미덕은 바로 여기 있다. 원인을 바꿔 가며 볼 수 있다.

4. playbook에 들어가야 할 최소 기록

좋은 playbook은 명령어 모음이 아니라 재현 가능한 관찰 기록이다. model size, dataset slice, tokenizer, sequence length, global batch, gradient accumulation, precision, optimizer, learning-rate schedule, seed, hardware, wall-clock, tokens/sec를 남겨야 한다.

특히 작은 training에서는 "잘 된 run"보다 "왜 실패했는지"가 더 귀하다. divergence, NaN, data corruption, eval mismatch, checkpoint incompatibility를 짧게라도 적어 두면 다음 사람이 같은 구덩이에 다시 들어가지 않는다.

5. 과장하지 않는 scaling 읽기

작은 run의 loss curve가 예쁘다고 큰 run도 성공한다고 말할 수는 없다. 반대로 작은 run에서 absolute quality가 낮다고 아이디어를 버릴 필요도 없다. 작은 실험은 주로 방향과 버그를 보는 도구다.

원본을 볼 때는 강의가 어떤 기준으로 "smol"을 정의하는지 확인하자. parameter count인지, dataset size인지, GPU budget인지, 혹은 iteration speed인지에 따라 playbook의 의미가 달라진다. 같은 단어가 귀여워 보여도 budget spreadsheet 앞에서는 꽤 냉정해진다.

6. 혼자 해볼 것

작은 language model fine-tuning run을 하나 정하고, 100 step짜리 smoke test를 만든다. 목표는 좋은 모델을 얻는 것이 아니라 loss, tokens/sec, memory, checkpoint/resume, eval script가 모두 살아 있는지 확인하는 것이다.

그다음 한 가지만 바꾼다. batch size, precision, sequence length, dataloader workers 중 하나를 바꾸고 결과를 표로 남겨 보자. 작은 playbook은 멋진 선언보다 이런 표 한 장에서 시작된다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_092.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_092.md

1. 자료의 바닥선을 먼저 긋기

local notes에는 speaker가 Jeff Ma로 기록되어 있고 YouTube link만 있다. transcript, slides, code, paper, repo folder는 모두 확인되지 않았다. 따라서 Cornserve가 오픈소스 serving framework인지, 연구 prototype인지, product architecture인지 아직 확정할 수 없다.

이럴 때 독자는 제목의 형용사를 곧이곧대로 믿기보다 질문으로 바꾸는 편이 좋다. easy는 누구에게 쉬운가. fast는 어떤 workload에서 빠른가. scalable은 GPU 수, request rate, model count, modality count 중 무엇에 대한 말인가.

2. 멀티모달 serving은 한 모델 serving보다 울퉁불퉁하다

text-only LLM serving은 prefill과 decode의 모양만으로도 충분히 복잡하다. 멀티모달은 여기에 image/audio/video encoder, preprocessing, feature projection, variable-size input이 들어온다. request마다 tensor shape와 latency profile이 크게 달라진다.

그래서 scheduler는 단순히 token 수만 보지 못한다. 이미지 encoder가 먼저 막히는지, text decoder가 막히는지, CPU preprocessing이 발목을 잡는지, KV cache가 memory를 먹는지 분리해야 한다. "멀티모달"이라는 단어 하나가 queue를 여러 줄로 늘려 버린다.

3. easy는 API 문제가 아니라 실패 모드 문제다

serving system이 쉽다는 말은 hello world가 짧다는 뜻만은 아니다. batch가 섞이고, input이 크고, model이 여러 개고, GPU memory가 빠듯할 때 실패를 이해할 수 있어야 쉽다.

멀티모달에서는 특히 admission control이 중요하다. 큰 이미지 몇 장이 들어와 encoder batch를 밀어 버리거나, 긴 text prompt가 decoder queue를 오래 붙잡을 수 있다. framework가 이런 request를 어떻게 제한하고 관찰하게 해 주는지가 "easy"의 실제 내용이다.

4. fast와 scalable을 따로 읽기

fast는 한 request의 latency일 수도 있고, 전체 tokens/sec 또는 requests/sec일 수도 있다. scalable은 GPU를 더 붙였을 때 throughput이 늘어나는지, 여러 model을 얹었을 때 운영 복잡도가 감당되는지, modality가 늘어도 pipeline이 무너지지 않는지의 문제다.

원본을 볼 때는 benchmark table보다 workload definition을 먼저 확인하자. batch size, image resolution, output length, concurrency, GPU type, warmup, cache state가 빠진 숫자는 친절해 보이지만 실전에서는 길을 잃기 쉽다.

5. 구현자가 조심할 점

멀티모달 serving은 memory layout과 data movement가 숨어서 성능을 갉아먹는다. CPU에서 이미지를 만지고, GPU encoder로 옮기고, 다시 LLM 입력으로 붙이는 사이에 copy와 synchronization이 생긴다. 작은 copy도 request rate가 높으면 줄줄 샌다.

또한 model parallelism과 pipeline parallelism이 항상 답은 아니다. encoder와 decoder의 병목 위치가 다르면 GPU 배치도 달라져야 한다. Cornserve가 이 배치를 자동화하는지, 아니면 사용자가 policy를 적는지 확인하면 강의의 실전성이 보인다.

6. 혼자 해볼 것

간단한 vision-language pipeline을 하나 정하고 latency를 구간별로 재 보자. image load/resize, vision encoder, projection, LLM prefill, decode를 따로 기록한다. 처음부터 최적화하지 말고 어디서 시간이 사라지는지부터 본다.

다음에는 서로 다른 image resolution과 output length를 섞어 queue를 만들어 본다. 평균 latency가 아니라 p95 latency가 어떻게 움직이는지 보면 멀티모달 serving이 왜 scheduler 문제인지 몸으로 이해된다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_093.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_093.md

1. 확인된 내용과 조심할 선

local notes에는 speaker가 Tianqi Chen으로 기록되어 있고 YouTube link만 있다. transcript fetch는 실패했고 slides/code/paper/repo folder는 없다. 그래서 이 강의가 TVM의 새 FFI package 소개인지, runtime object model 설명인지, Python/C++ extension 설계 강의인지는 아직 확인해야 한다.

다만 TVM과 FFI의 조합은 충분히 중요한 독해 포인트를 준다. ML compiler는 Python 사용자, C++ runtime, generated kernel, device memory, external library 사이를 계속 오간다. 그 사이의 경계가 흐리면 성능 버그도, lifetime 버그도, packaging 버그도 같이 온다.

2. FFI는 "함수 호출"보다 넓다

Foreign Function Interface라고 하면 다른 언어의 함수를 부르는 얇은 wrapper를 떠올리기 쉽다. compiler/runtime stack에서는 조금 더 복잡하다. object ownership, type conversion, error propagation, dynamic dispatch, serialization, ABI stability가 모두 얽힌다.

TVM 계열 시스템에서는 Python에서 schedule을 만들고, C++ 쪽 IR/runtime object를 만지고, 최종적으로 device code를 실행하는 흐름이 자연스럽다. 이때 FFI가 튼튼해야 사용자는 Python의 생산성을 얻고 runtime은 C++의 제어력을 유지한다.

3. 성능 경계로서의 FFI

FFI overhead는 hot path에 들어가면 금방 문제가 된다. kernel 하나가 milliseconds라면 Python-to-C++ 호출 비용은 작아 보일 수 있지만, 작은 operator를 많이 launch하거나 tuning loop에서 수만 번 호출하면 이야기가 달라진다.

그래서 좋은 FFI 설계는 "어디까지 crossing을 허용할 것인가"를 정한다. fine-grained call을 편하게 만들 것인지, 큰 operation 단위로 묶을 것인지, callback을 어느 방향으로 허용할 것인지가 전체 시스템의 모양을 바꾼다.

4. 타입과 lifetime이 조용한 난제다

compiler IR object는 보통 graph처럼 서로를 참조한다. Python object가 사라졌는데 C++ object가 살아 있거나, 반대로 runtime buffer가 device에서 사용 중인데 host wrapper가 해제되면 문제는 매우 피곤해진다.

또한 dynamic type system과 static type system 사이에는 번역 비용이 있다. Tensor shape, dtype, device, symbolic expression, module handle 같은 값이 어느 쪽에서 검증되는지 알아야 debugging이 가능하다. 에러 메시지가 FFI boundary에서 증발하면 하루가 길어진다.

5. 원본에서 확인할 질문

원본을 볼 때는 tvm-ffi가 기존 TVM runtime FFI와 어떤 관계인지부터 확인하자. 새 package인지, refactor인지, Python extension authoring을 위한 별도 layer인지에 따라 독해가 달라진다.

또한 zero-copy, object registry, packed function, exception mapping, build system, wheel packaging, thread safety 같은 항목이 언급되는지 살펴보면 좋다. 이런 주제들은 화려하지 않지만 compiler를 실제로 쓰게 만드는 볼트와 너트다.

6. 혼자 해볼 것

작은 C++ function을 Python에서 호출하는 binding을 하나 만들어 보고, scalar call과 batched call의 overhead를 비교해 보자. 꼭 TVM이 아니어도 된다. pybind11, ctypes, cffi 중 하나로 FFI boundary의 감각을 잡는 것이 목적이다.

그다음 device buffer handle을 직접 넘긴다고 상상하며 ownership table을 적어 보자. 누가 할당하고, 누가 해제하며, 에러가 나면 누가 정리하는가. FFI 설계는 이 표가 지저분해질 때 본색을 드러낸다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_094.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_094.md

1. 자료 상태

local notes에는 speaker가 Allen Wang and Colin Taylor로 기록되어 있고 YouTube link만 있다. transcript fetch는 실패했고 slides/code/paper/repo folder는 없다. Monarch가 어떤 project의 어떤 version인지, PyTorch distributed와 어떤 관계인지도 이 자료만으로는 말할 수 없다.

그래서 여기서는 single controller programming이라는 문구를 중심으로 읽는다. distributed training/serving에서는 "누가 전체를 보고 있는가"가 API의 모양을 크게 바꾼다.

2. single controller가 던지는 약속

전통적인 SPMD 스타일에서는 각 rank가 같은 program을 실행하고, rank id에 따라 다른 일을 한다. 이 방식은 강력하지만 debugging할 때 머리가 쉽게 둘로, 넷으로, 여덟으로 갈라진다. log도 rank마다 나오고, failure도 rank마다 다르게 보인다.

single controller programming은 사용자가 하나의 driver 관점에서 여러 worker/device를 조작하게 만들겠다는 약속처럼 읽힌다. tensor나 actor가 remote에 있어도, 코드는 더 중앙집중적인 control flow로 표현될 수 있다.

3. 쉬워지는 것과 어려워지는 것

쉬워지는 것은 orchestration이다. 여러 machine에 model shard를 놓고, pipeline stage를 만들고, evaluation job을 섞는 일을 한 파일에서 표현할 수 있다면 생산성이 좋아진다. 특히 연구 code는 control flow가 자주 바뀌기 때문에 이런 장점이 크다.

어려워지는 것은 illusion management다. 코드가 하나처럼 보인다고 실행도 하나처럼 단순해지는 것은 아니다. remote call latency, data movement, failure propagation, partial retry, backpressure는 여전히 존재한다. API가 이것들을 너무 숨기면 예쁜 code가 느린 system이 된다.

4. GPU 시스템에서 중요한 경계

분산 ML의 병목은 compute만이 아니다. activation transfer, parameter shard movement, optimizer state, checkpoint IO, collective communication이 섞인다. single controller가 이런 비용을 어떻게 드러내는지가 중요하다.

또한 scheduling 결정이 controller에 몰리면 편하지만, 너무 세밀한 작업까지 controller가 지시하면 overhead가 커질 수 있다. 반대로 worker에 많은 자율성을 주면 programming model이 다시 복잡해진다. 좋은 설계는 어느 층에서 결정을 고정하고 어느 층에서 지연할지 조심스럽게 고른다.

5. 원본을 볼 때의 체크리스트

Monarch가 tensor abstraction인지, actor/runtime인지, distributed execution engine인지 먼저 확인하자. API 예제가 있다면 remote tensor가 local tensor처럼 보이는지, placement annotation이 명시적인지, async/future가 어떻게 표현되는지 보자.

failure case도 중요하다. 한 worker가 죽었을 때 controller는 어떤 exception을 보는가. collective가 hang 되면 어떻게 끊는가. debug trace가 rank별로 흩어지는가, 아니면 controller view로 재구성되는가. 이런 부분이 lecture의 진짜 실전성을 가른다.

6. 혼자 해볼 것

PyTorch distributed나 Ray 같은 도구로 아주 작은 2-worker job을 만들고, SPMD 스타일과 controller 스타일 pseudo-code를 나란히 써 보자. 같은 all-reduce나 remote matmul을 표현해도 어떤 정보가 보이고 숨는지 비교한다.

그다음 deliberately failure를 넣어 본다. 한 worker에서 exception을 내고, controller가 어떤 stack trace를 보는지 기록하자. 분산 programming model은 정상 경로보다 실패 경로에서 성격이 드러난다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_095.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_095.md

1. 알려진 것과 모르는 것

speaker는 local notes에서 Missing으로 되어 있다. transcript fetch는 실패했고 slides/code/repo folder도 없다. TLX가 특정 library 이름인지, Triton extension인지, compiler IR인지, research language인지는 원본 확인 전까지 확정할 수 없다.

다만 Triton, kernel, warp, fusion, scheduling이라는 단서가 함께 잡힌다. 이 조합은 보통 "타일 단위로 쉽게 쓰되, 더 낮은 실행 제어를 어떻게 줄 것인가"라는 질문으로 이어진다.

2. Triton이 준 좋은 기본값

Triton은 Python 안에서 tile-based GPU kernel을 쓰게 해 준다. program id로 tile을 나누고, block pointer와 vectorized operation으로 memory access를 표현하며, compiler가 많은 lowering을 맡는다. 이 생산성은 실제로 큰 힘이다.

하지만 더 복잡한 kernel로 갈수록 기본값만으로 부족해진다. warp specialization, async copy, shared memory swizzle, inter-warp communication, fusion boundary 같은 결정은 compiler가 늘 맞히기 어렵다. TLX가 이 영역을 다룬다면, 독자는 "무엇을 더 노출하는가"를 봐야 한다.

3. fusion은 공짜 점심이 아니다

fusion은 kernel launch overhead와 global memory traffic을 줄일 수 있다. 그러나 fusion이 커지면 register pressure, occupancy, scheduling complexity가 올라간다. 어떤 op를 붙이면 빨라지고, 어떤 op를 붙이면 느려지는지 profiler 없이 맞히기는 어렵다.

TLX가 fusion을 강조한다면, 단순히 "여러 op를 하나로"가 아니라 intermediate를 어디에 두는지, layout을 어떻게 유지하는지, reduction과 elementwise를 어떻게 섞는지 확인해야 한다. fusion은 요리라기보다 짐 싸기에 가깝다. 너무 많이 넣으면 가방이 안 닫힌다.

4. warp와 scheduling을 드러낸다는 뜻

warp-level control은 성능을 열어 주지만, 동시에 correctness 부담을 키운다. 누가 load하고, 누가 compute하고, 어느 barrier에서 만나며, shared memory를 언제 덮어써도 되는지 명확해야 한다.

좋은 DSL은 low-level primitive를 그냥 노출하는 데서 끝나지 않는다. 사용자가 잘못된 schedule을 만들었을 때 compiler가 잡아 줄 수 있는지, layout conversion과 synchronization을 추적할 수 있는지, generated code를 살펴볼 수 있는지가 중요하다.

5. 원본에서 확인할 질문

TLX가 Triton syntax를 확장하는지, 별도 IR을 제안하는지, 혹은 특정 backend를 위한 compiler pass인지 확인하자. 예제 kernel이 있다면 matmul, attention, reduction, scan 중 무엇인지도 중요하다.

benchmark가 있다면 비교 기준을 조심해서 보자. PyTorch, Triton, CUDA, vendor library 중 무엇과 비교했는지, shape와 dtype은 무엇인지, tuning effort가 같았는지 확인해야 한다. kernel language 비교는 code length만으로 끝나지 않는다.

6. 혼자 해볼 것

Triton으로 작은 fused operation을 하나 작성해 보자. 예를 들어 y = gelu(x @ w + b)를 matmul과 epilogue로 나누어 생각하고, fusion이 memory traffic을 줄이는 대신 register pressure를 늘릴 수 있음을 기록한다.

그다음 num_warps, block size, num_stages를 바꾸며 benchmark를 남긴다. TLX 원본을 볼 때 이 표가 있으면 "이 언어가 무엇을 자동화하고 무엇을 사용자에게 맡기는지" 더 잘 보인다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_096.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_096.md

1. 지금 확인된 자료

local notes에는 speaker가 William Hu로 기록되어 있고 YouTube link만 있다. transcript fetch는 실패했고 slides/code/paper/repo folder는 없다. HipKittens라는 이름의 정확한 repository, target hardware, supported operation은 이 자료만으로 알 수 없다.

그래서 이 글은 title에서 조심스럽게 읽을 수 있는 portability 문제를 중심으로 둔다. HIP는 AMD GPU programming stack과 관련된 이름으로 자주 쓰이고, Kittens라는 이름은 tile/register 중심 kernel abstraction을 떠올리게 한다. 하지만 여기서 구체 API를 지어내면 안 된다.

2. portability는 syntax 변환이 아니다

CUDA kernel을 HIP로 바꾸는 일은 함수 이름을 바꾸는 것으로 끝나지 않는다. warp와 wavefront 크기, memory hierarchy, vectorized load/store, matrix instruction, shared memory bank, compiler behavior가 다르다.

high-level abstraction이 있다면 이 차이를 감춰 주려 할 수 있다. 그러나 성능 kernel에서는 차이를 완전히 숨기는 순간 최적화 손잡이도 같이 사라질 수 있다. portability의 묘미는 "공통 개념"과 "hardware-specific escape hatch"의 균형이다.

3. Kittens류 abstraction을 읽는 포인트

tile abstraction은 kernel writer에게 큰 도움을 준다. register tile, shared memory tile, warp group 역할을 명시하면 code가 수학식과 hardware 사이 어딘가에서 읽힌다. 좋은 abstraction은 반복되는 boilerplate를 줄이고, layout 실수를 줄인다.

하지만 abstraction이 target hardware를 너무 NVIDIA식으로 가정하면 이식이 어렵다. 반대로 모든 hardware를 일반화하려다 보면 가장 중요한 fast path가 흐릿해진다. HipKittens가 흥미롭다면 아마 이 긴장 위에 있을 가능성이 높다.

4. correctness를 먼저 잡아야 한다

porting 작업에서는 "컴파일된다"와 "맞다"와 "빠르다"가 완전히 다른 단계다. dtype별 rounding, boundary mask, non-contiguous stride, alignment, odd shape에서 버그가 잘 난다.

특히 attention이나 matmul kernel은 작은 shape에서 우연히 맞고 큰 shape에서 틀릴 수 있다. reference implementation과 tolerance policy를 먼저 정해야 한다. 이 부분을 건너뛰면 benchmark 표가 모래 위에 선다.

5. 원본에서 확인할 질문

HipKittens가 어떤 upstream idea를 가져왔는지, target GPU가 무엇인지, 어떤 operation을 보여 주는지 확인하자. Flash attention, GEMM, convolution, reduction 중 무엇을 다루는지에 따라 읽는 렌즈가 달라진다.

또한 performance claim이 있다면 같은 hardware에서 vendor library와 비교했는지, CUDA version과 비교했는지, 아니면 portability 자체가 목표인지 분리해 보자. 이식성 강의에서 "느리다/빠르다"는 목표에 따라 다른 뜻을 가진다.

6. 혼자 해볼 것

아주 작은 vector add나 tiled matmul을 CUDA와 HIP pseudo-code로 나란히 써 보자. API 이름보다 execution model 차이를 적는 것이 중요하다. thread grouping, shared memory 이름, synchronization, compile toolchain을 표로 만든다.

그다음 실제 porting checklist를 만든다. build, correctness, profiler, shape sweep, dtype sweep, baseline 비교. portability는 한 번에 성공하는 마법이 아니라 체크박스를 줄여 가는 일에 가깝다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_097.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_097.md

1. 확인된 자료

local notes에는 speaker가 Yusheng Zheng으로 기록되어 있고 YouTube link만 있다. transcript fetch는 실패했고 slides/code/paper/repo folder는 없다. 따라서 특정 observability platform이나 library를 소개하는 강의인지, 일반 원칙을 다루는 강의인지는 아직 모른다.

그래도 GPU observability라는 주제는 Part VII의 시스템 실전과 잘 맞는다. kernel 하나를 최적화하는 것과, 여러 사용자가 공유하는 GPU fleet에서 문제를 알아차리는 것은 다른 기술이다.

2. profiling과 observability의 차이

profiling은 보통 개발자가 짧은 시간 동안 깊게 들여다보는 일이다. Nsight Compute로 kernel counter를 보고, Nsight Systems로 timeline을 본다. 이건 현미경에 가깝다.

observability는 운영 중인 시스템이 계속 보내는 신호로 문제를 찾는 일이다. GPU utilization, memory usage, ECC error, throttling, queue length, request latency, tokens/sec, kernel failure, driver reset 같은 신호가 시간축 위에 쌓인다. 현미경만 들고 datacenter를 돌아다닐 수는 없다.

3. 좋은 metric은 행동을 만든다

metric이 많다고 좋은 것은 아니다. GPU utilization 95%라는 숫자만 보고는 compute가 잘 쓰이는지, memory copy로 바쁜지, kernel이 spin하고 있는지 알기 어렵다. metric은 다음 행동을 가리켜야 한다.

예를 들어 tokens/sec가 떨어졌고 GPU memory는 꽉 찼으며 batch size가 줄었다면 KV cache pressure를 의심할 수 있다. GPU utilization이 낮고 CPU preprocessing queue가 길면 input pipeline을 본다. observability는 숫자 수집보다 원인 가설을 빨리 세우는 구조다.

4. production GPU의 흔한 함정

단일 benchmark에서는 보이지 않던 문제가 production에서 나온다. multi-tenant 환경에서는 noisy neighbor가 생기고, 긴 request가 batch를 붙잡고, thermal throttling이나 power cap이 latency tail을 흔든다.

또한 driver/runtime level error는 application log와 떨어져 있을 수 있다. CUDA error가 늦게 보고되거나, async execution 때문에 원인 위치가 흐려지는 일도 많다. observability system은 host log, runtime event, device metric의 시간을 맞춰야 한다.

5. 원본에서 확인할 질문

강의가 어떤 metric taxonomy를 제안하는지 보자. hardware metric 중심인지, ML serving metric 중심인지, Kubernetes/DCGM 같은 운영 stack 중심인지에 따라 실전 적용이 달라진다.

case study가 있다면 "증상에서 원인까지"의 경로를 눈여겨보자. 좋은 observability 강의는 dashboard screenshot보다 추론 과정을 보여 준다. 어떤 signal을 보고 어떤 가능성을 버렸는지가 더 중요하다.

6. 혼자 해볼 것

작은 inference server나 training script를 하나 띄우고 GPU metric과 application metric을 동시에 기록해 보자. nvidia-smi dmon, DCGM, PyTorch profiler, server latency log 중 가능한 도구를 조합한다.

그다음 일부러 bottleneck을 만든다. CPU preprocessing sleep을 넣거나 batch size를 줄이거나 memory를 압박해 본다. metric이 그 변화를 설명해 주는지 확인하면, observability가 장식인지 도구인지 금방 드러난다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_098.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_098.md

1. 무엇을 알고 무엇을 모르는가

local notes에는 speaker가 Matt Beton으로 기록되어 있고 YouTube link만 있다. transcript fetch는 실패했고 slides/code/paper/repo folder는 없다. 그래서 이 강의가 federated learning, edge inference, volunteer compute, browser/device cluster 중 무엇을 중심으로 하는지 아직 알 수 없다.

consumer device라는 말은 GPU Mode 맥락에서 흥미롭다. data center에서는 hardware, network, driver, scheduler를 통제할 수 있다. 소비자 기기에서는 그 통제가 약해진다. 바로 그 불편함이 주제의 핵심일 가능성이 높다.

2. consumer device는 작고 많고 제멋대로다

소비자 기기는 성능이 균일하지 않다. laptop GPU, desktop GPU, phone NPU, integrated GPU, browser runtime이 섞일 수 있다. 전원 상태, 발열, background app, network가 계속 바뀐다.

distributed ML은 보통 동기화 비용과 straggler 문제를 안고 있다. consumer device에서는 이 문제가 더 커진다. 빠른 기기 몇 대가 느린 기기 하나를 기다리게 되면 전체 효율이 무너진다. 그래서 "분산"이라는 말은 늘 scheduling과 fault tolerance를 동반한다.

3. 어떤 workload가 맞을까

모든 ML workload가 consumer device 분산에 맞지는 않는다. 큰 synchronous training은 network와 straggler에 약하다. 반면 embarrassingly parallel한 evaluation, data preprocessing, hyperparameter sweep 일부, small fine-tuning, federated averaging류 작업은 더 가능성이 있다.

inference도 경우에 따라 나뉜다. model을 여러 device에 shard하면 communication이 커지고, request latency가 불안해질 수 있다. 대신 각 device가 독립 request를 처리하는 식이면 scheduling이 단순해진다.

4. 보안과 신뢰는 성능 문제이기도 하다

consumer device를 쓰면 data privacy, model theft, malicious result, sandboxing 문제가 따라온다. 이건 별도 정책 문제가 아니라 system cost로 돌아온다. 암호화, verification, sandbox overhead가 latency와 throughput에 영향을 준다.

또한 device가 언제든 사라질 수 있다고 가정해야 한다. checkpointing, idempotent task, retry, partial aggregation 같은 설계가 필요하다. datacenter에서 당연하던 "노드가 한동안 살아 있다"는 믿음이 약해진다.

5. 원본에서 확인할 질문

강의가 training을 말하는지 inference를 말하는지 먼저 확인하자. 그다음 target device, runtime, network model, privacy assumption을 본다. 이 네 가지가 없으면 performance claim을 해석하기 어렵다.

benchmark가 있다면 평균 throughput보다 tail latency와 failure rate를 보자. consumer device system은 최고 성능보다 흔들림을 견디는 능력이 중요할 수 있다.

6. 혼자 해볼 것

두 대의 서로 다른 machine이나 하나의 machine 안의 서로 다른 process를 consumer device라고 가정하고 작은 distributed job을 만들어 보자. 일부 worker에 sleep을 넣어 straggler를 만들고, synchronous aggregation이 얼마나 흔들리는지 본다.

그다음 작업 단위를 더 작게 쪼개거나 timeout/retry를 넣어 본다. distributed ML은 algorithm만큼 운영 조건을 먹고 산다는 사실이 손에 잡힌다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_099.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_099.md

1. 100번째 강의의 자료 상태

local notes에는 speakers가 Kimbo Chen, Cam Quilici, Bryan Shan으로 기록되어 있고 YouTube link만 있다. transcript fetch는 실패했고 slides/code/paper/repo folder는 없다. InferenceX의 실제 dashboard, harness, supported model/server 목록은 원본 확인이 필요하다.

그래도 continuous benchmarking이라는 단어는 충분히 구체적이다. 이것은 "한 번 재서 블로그에 올리기"가 아니라, code와 dependency가 바뀔 때마다 inference 성능을 추적하는 문제다.

2. inference benchmark는 쉽게 썩는다

inference stack은 빠르게 변한다. model runtime, CUDA driver, PyTorch, vLLM, TensorRT-LLM, tokenizer, kernel, quantization path, scheduler가 모두 움직인다. 어제의 benchmark가 오늘의 package 조합에서는 다른 의미가 될 수 있다.

그래서 benchmark는 result table보다 environment record가 먼저다. GPU type, driver, CUDA, library version, model revision, prompt/output distribution, concurrency, warmup, cache policy를 남겨야 한다. 숫자는 혼자 서 있지 못한다.

3. 무엇을 재는지가 곧 철학이다

LLM inference는 latency와 throughput이 서로 당긴다. batch를 키우면 throughput은 좋아질 수 있지만 per-request latency는 나빠진다. streaming decode에서는 first token latency와 inter-token latency도 따로 봐야 한다.

continuous benchmark라면 workload suite가 중요하다. batch size 1, 높은 concurrency, 긴 prompt, 긴 output, short chat, quantized model, multi-GPU serving을 나누어야 한다. 평균 하나로는 운영자가 결정을 내리기 어렵다.

4. OSS benchmark의 어려움

open-source stack은 장점도 크고 변동성도 크다. 최신 commit이 빠르지만 불안정할 수 있고, 특정 GPU에서만 좋은 path가 켜질 수 있다. benchmark harness는 failure를 성능 숫자처럼 정직하게 기록해야 한다.

또한 비교 대상의 tuning effort를 맞추어야 한다. 한 runtime은 default config, 다른 runtime은 hand-tuned config로 비교하면 표는 예쁘지만 공정하지 않다. continuous system은 config drift를 막는 장치가 필요하다.

5. 원본에서 확인할 질문

InferenceX가 어떤 runtime과 model을 포함하는지 확인하자. 그리고 benchmark가 synthetic prompt인지 real trace인지, output length를 어떻게 통제하는지, correctness나 response quality를 같이 보는지 살펴보자.

regression detection 방식도 중요하다. 단순 threshold인지, 통계적 noise를 고려하는지, hardware noise와 software regression을 어떻게 구분하는지 봐야 한다. benchmark는 알람이 너무 많아도, 너무 조용해도 쓸모가 줄어든다.

6. 혼자 해볼 것

작은 inference script를 하나 정하고 매일 같은 config로 latency와 tokens/sec를 JSON으로 저장해 보자. package version과 GPU 이름도 같이 남긴다. 며칠만 지나도 benchmark가 data engineering 문제라는 사실이 보인다.

그다음 의도적으로 batch size나 output length를 바꿔 regression처럼 보이게 만들어 보자. benchmark harness가 config 변화와 성능 변화를 구분해 주는지 확인하면 좋은 continuous benchmark의 조건이 드러난다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_100.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_100.md

1. 자료 상태와 읽기 목표

local notes에는 speaker가 Kapil Sharma로 기록되어 있고 YouTube link만 있다. transcript fetch는 실패했고 slides/code/paper/repo folder는 없다. 따라서 강의의 구체 예제, CUTLASS version, target GPU, benchmark 숫자는 확인해야 한다.

그래도 "Learning CUTLASS the hard way"라는 제목은 방향이 분명하다. CUTLASS는 빠른 GEMM을 부르는 wrapper가 아니라, production GEMM kernel을 구성하는 부품 상자에 가깝다. 어려운 길은 괴롭히려는 길이 아니라 내부 구조를 보려는 길이다.

2. CUTLASS를 배우기 어려운 이유

CUTLASS는 template parameter가 많고 이름도 길다. 처음 보면 type이 code보다 더 큰 것처럼 느껴진다. 하지만 그 type들은 대개 tile shape, data type, layout, operator class, architecture, epilogue, pipeline policy를 표현한다.

어려움은 "문법"보다 "mapping"에 있다. M/N/K tile이 threadblock, warp, instruction tile로 어떻게 쪼개지는지, shared memory가 어떤 layout으로 쓰이는지, Tensor Core instruction이 어떤 shape를 먹는지 연결해야 한다.

3. GEMM 하나에도 많은 계약이 있다

GEMM은 C = A B라고 쓰면 간단하지만 kernel에서는 계약이 많다. row-major/column-major, leading dimension, alignment, dtype accumulation, beta/alpha scaling, epilogue fusion, split-K 여부가 모두 영향을 준다.

CUTLASS를 배울 때는 처음부터 모든 feature를 켜기보다 가장 작은 GEMM에서 출발하는 편이 좋다. 같은 shape를 cuBLAS와 비교하고, correctness를 확인하고, profiler로 memory와 tensor core 사용을 본다. 그다음 epilogue나 layout을 하나씩 바꾼다.

4. 성능을 읽는 눈

CUTLASS kernel이 느릴 때는 template을 더 만지기 전에 병목을 분리해야 한다. Tensor Core를 쓰는지, memory coalescing이 되는지, shared memory bank conflict가 있는지, occupancy가 너무 낮은지, problem shape가 tile과 맞지 않는지 본다.

특히 edge shape는 중요하다. 예쁜 multiple-of-128 shape에서는 빠른 kernel도, 실제 LLM workload의 odd dimension이나 small batch에서는 다른 path로 갈 수 있다. "hard way"는 이런 지저분한 모서리를 보는 습관을 포함한다.

5. 원본에서 확인할 질문

강의가 CUTLASS 2.x 중심인지, CUTLASS 3.x/CuTe 중심인지 확인하자. 둘은 읽는 법이 다르다. CuTe가 들어오면 layout algebra와 tensor abstraction을 같이 이해해야 한다.

예제 code가 있다면 template parameter를 한 줄씩 풀어 보자. operation class, architecture tag, tile shape, stage count, epilogue operator가 각각 무엇을 정하는지 표로 옮기면 막연함이 줄어든다.

6. 혼자 해볼 것

CUTLASS example 중 가장 단순한 GEMM을 빌드하고, 같은 input을 cuBLAS와 비교한다. 그다음 M/N/K를 조금씩 바꾸며 performance가 어디서 떨어지는지 기록한다.

추가로 epilogue에 bias나 activation이 붙는 경우를 찾아보자. LLM kernel에서는 GEMM 뒤의 작은 연산을 어떻게 붙이느냐가 memory traffic을 크게 바꾼다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_101.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_101.md

1. 이 강의는 특히 빈칸이 크다

local notes에는 speakers가 Andrei Panferov and Erik Schultheis로 기록되어 있고 YouTube link만 있다. transcript fetch는 실패했고 slides/code/paper/repo folder는 없다. quartet v2가 library, compiler, kernel collection, serving system, research project 중 무엇인지는 확인되지 않았다.

이런 경우 좋은 reader 원고는 내용을 지어내지 않는 데서 시작한다. 모르는 것을 모른다고 표시해 두면 나중에 원본을 볼 때 채워 넣을 자리가 생긴다. 빈칸도 잘 관리하면 자료다.

2. v2 발표에서 먼저 볼 것

어떤 시스템의 v2 발표라면 보통 v1의 pain point가 있다. 성능이 부족했는지, API가 어려웠는지, correctness 문제가 있었는지, target hardware가 바뀌었는지, 운영 경험이 쌓였는지 확인해야 한다.

v2의 핵심은 새 기능 목록보다 trade-off 변화다. 더 빠른 대신 portability를 잃었는가. 더 쉬운 API 대신 low-level control을 숨겼는가. 더 큰 scale을 위해 debugging이 어려워졌는가. 시스템은 늘 어딘가에서 비용을 낸다.

3. 이름보다 interface를 보자

프로젝트 이름만으로는 거의 아무것도 알 수 없다. 원본에서 가장 먼저 찾아야 할 것은 interface다. 사용자는 어떤 code를 쓰는가. input과 output은 무엇인가. 기존 stack과 어디에서 만나는가.

GPU 시스템에서는 interface가 성능 경계가 되기도 한다. tensor layout을 노출하는지, scheduling을 숨기는지, kernel fusion을 자동으로 하는지, distributed placement를 사용자가 적는지에 따라 가능한 최적화가 달라진다.

4. 증거를 나누어 읽기

system talk의 evidence는 여러 층이다. microbenchmark는 특정 kernel이나 primitive의 상한을 보여 줄 수 있다. end-to-end benchmark는 실제 workload 감각을 준다. production anecdote는 failure mode를 알려 준다.

quartet v2 원본을 볼 때는 이 증거들을 섞지 말자. microbenchmark가 좋아도 end-to-end가 자동으로 좋아지는 것은 아니다. 반대로 end-to-end 개선이 있어도 어떤 component가 기여했는지 따로 확인해야 한다.

5. 나중에 채워야 할 질문

quartet v2가 어떤 domain의 system인지 먼저 적어야 한다. kernel DSL인지, inference engine인지, distributed runtime인지, model architecture tool인지에 따라 이 노트의 상당 부분이 달라진다.

그다음 v2에서 바뀐 핵심 mechanism을 하나만 고르자. scheduler, layout, compiler lowering, caching, communication protocol, API ergonomics 중 무엇이 중심인지 확인하면 독해가 선명해진다.

6. 혼자 해볼 것

원본을 보기 전에 "v2 발표 체크리스트"를 만들어 보자. 기존 문제, 새 설계, 증거, trade-off, migration cost, open questions 여섯 칸이면 충분하다.

영상을 본 뒤에는 각 칸을 한 문장으로 채운다. 한 칸이 비면 그 발표가 부족한 것일 수도 있고, 우리가 놓친 것일 수도 있다. 어느 쪽이든 다음 질문이 생긴다.

Source Links

• YouTube
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_102.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_102.md

1. 이 강의는 슬라이드가 있다

local notes에는 speaker가 Jack Carlisle and Jay Shah로 기록되어 있고, GPU Mode repo의 lecture_103/layout_algebra.pdf가 연결되어 있다. transcript fetch는 실패했고 code link는 없다. 따라서 말로 설명된 예시나 Q&A는 빠져 있지만, slide outline과 수식은 어느 정도 따라갈 수 있다.

슬라이드는 CuTe layouts, CuTe layout algebra, categories, categories of layouts, composition algorithm 순서로 구성된다. GEMM의 A/B/C tensor와 attention의 Q/K/V shape를 출발점으로 삼고, tensor를 iterator와 layout L = S : D로 표현한다.

2. layout은 shape와 stride 이상의 함수다

CuTe에서 layout은 shape S와 stride D의 쌍처럼 보인다. 예를 들어 (4, 6) : (6, 1)은 익숙한 2D mapping을 만든다. 하지만 강의는 여기서 한 발 더 간다. nested tuple shape와 stride를 통해 tile, warp, instruction layout까지 표현한다.

중요한 관점은 coordinate function이다. 입력 좌표가 있고, stride와 곱해져 integer offset으로 간다. 이렇게 보면 layout은 메모리 그림이 아니라 함수다. 함수로 보면 합성, 보완, coalesce 같은 연산을 말할 수 있다.

3. nested tuple은 장식이 아니다

슬라이드는 mode, entry, rank, length, size 같은 용어를 정의한다. (4, (5, 5), (2, (3, 3))) 같은 nested tuple은 보기에는 복잡하지만, GPU kernel에서는 자연스럽다. threadblock, warp, lane, register tile처럼 계층이 있기 때문이다.

refinement도 중요한 아이디어다. 어떤 shape의 entry를 같은 size의 nested tuple로 바꾸면 더 세밀한 구조를 표현할 수 있다. 이는 "같은 총 원소 수를 다른 계층으로 보는 법"이다. kernel layout을 이해할 때 꽤 실용적인 언어다.

4. layout algebra의 기본 연산

슬라이드는 interleaved layout이 logical division A ⊘ B와 logical product A ⊗ B 같은 algebraic operation에서 자연스럽게 생긴다고 설명한다. 이 연산들은 composition, complement, coalesce, concatenation 같은 더 기본적인 연산 위에 정의된다.

composition은 특히 조심스럽다. B ◦ A는 shape가 A의 shape를 refine하고, layout function이 ΦB ◦ ΦA가 되도록 하는 layout이다. 그런데 임의의 두 layout을 합성한다고 항상 예쁜 layout으로 떨어지는 것은 아니다. 그래서 tractable layout이라는 제한이 등장한다.

5. tractable layout과 category 언어

슬라이드는 tractable layout을 column-major layout에서 mode permutation, mode removal, stride 0 mode insertion, reparenthesizing으로 얻을 수 있는 layout로 설명한다. compact layout, strided layout, broadcast layout이 예로 나온다.

category 이야기는 멋을 위한 장식이 아니라 합성 가능성을 다루기 위한 언어로 읽는 편이 좋다. object, morphism, composition rule을 가진 category를 소개하고, FinSet, Vect, positive integer divisibility 같은 예를 든다. layout을 morphism처럼 다루면 어떤 합성이 가능한지 더 엄밀하게 말할 수 있다.

6. kernel 작성자에게 왜 중요한가

layout algebra는 수학 강의처럼 보이지만, 실제로는 data movement를 덜 틀리게 만들기 위한 도구다. MMA accumulator layout, wgmma accumulator layout처럼 hardware instruction이 요구하는 배치가 있고, programmer는 register/shared/global memory 사이를 오가며 layout을 바꿔야 한다.

layout을 함수와 합성으로 이해하면 "이 transpose가 그냥 view인지, 실제 shuffle이 필요한지", "coalesce해도 의미가 같은지", "broadcast stride 0이 어디서 생기는지"를 더 차분하게 볼 수 있다. GPU kernel에서 주소 계산은 사소한 housekeeping이 아니라 성능과 correctness의 중심부다.

Source Links

• YouTube
• GPU Mode lecture folder
• Slides PDF
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_103.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_103.md

1. 슬라이드가 말해 주는 큰 흐름

local notes에는 speakers가 Peter Bell, Mario Lezcano, Keren Zhou로 기록되어 있고, GPU Mode repo의 lecture_104/gluon.pdf가 연결되어 있다. transcript fetch는 실패했고 code link는 없다. 그래도 slides에는 agenda, code fragment, Blackwell matmul case study, profiling/sanitizer/tooling 내용이 꽤 분명히 들어 있다.

agenda는 Gluon introduction, Triton과의 비교, language design, Linear Layouts deep dive, optimized Blackwell matmul, developer tools로 이어진다. 이 강의는 추상 이론만이 아니라 "새 hardware에서 kernel expert에게 어떤 손잡이를 줄 것인가"를 묻는다.

2. 왜 새 언어가 필요하다고 말하는가

슬라이드는 Triton을 high-level tile-based GPU programming으로 소개한다. effort 대비 성능이 좋고 compiler가 많은 복잡성을 처리한다는 장점도 분명히 인정한다.

하지만 Blackwell 쪽 예시는 block-level programming model과 warp specialization 사이의 mismatch를 강조한다. 어떤 warp는 load를 맡고, 어떤 warp는 MMA를 기다리고, 어떤 warp는 epilogue를 수행하는 식의 schedule에서는 compiler가 항상 최선의 결정을 찾기 어렵다. Gluon의 목표는 더 낮은 제어와 생산성 사이에 새 자리를 만드는 것이다.

3. Linear Layouts는 layout class를 하나의 언어로 묶는다

슬라이드는 기존에 register, shared memory, tensor memory layout을 위해 여러 class가 있었다고 설명한다. Linear Layouts는 data가 hardware에 어떻게 놓이는지를 함수로 표현하려는 시도다.

정의는 꽤 수학적이다. Linear Layout은 F2 위의 linear map, 즉 XOR와 AND를 쓰는 binary matrix로 설명된다. 이 관점은 단순히 예쁜 수학이 아니라 generic lowering, optimal layout conversion, swizzling, shuffle sequence 도출을 가능하게 하는 기반으로 제시된다.

4. Gluon의 낮은 손잡이들

Gluon slide에는 gl.warp_specialize의 fork-join model, default partition과 worker partition, static shared memory allocator, hardware intrinsic 노출이 나온다. 예시에는 mbarrier, tma.async_load, shared memory allocation 같은 Blackwell/Hopper 계열 async programming 요소가 보인다.

이런 기능은 kernel expert에게 강력하지만 공짜는 아니다. hardware 이해가 필요하고, portability도 제한된다. 슬라이드의 비교표도 Gluon을 "Tiles with Layouts", "Warp Partitions", "SOL performance", "not portable" 쪽으로 놓는다.

5. Blackwell matmul case study

slides는 optimized Blackwell matmul에서 tcgen05.mma, tensor memory, 2-CTA mode, TMA multicast, Cluster Launch Control을 언급한다. 특히 2-CTA mode는 CTA들이 B tile을 공유해 더 큰 MMA를 쓰고 shared memory replication을 줄이는 그림으로 설명된다.

TMA multicast는 여러 CTA가 같은 B tile을 global memory에서 중복 load하지 않도록 하는 방향으로 제시된다. Cluster Launch Control은 persistent kernel에서 work imbalance를 줄이기 위한 dynamic scheduling 맥락으로 설명된다. 정확한 성능 숫자는 slide context를 더 확인해야 하지만, B200 benchmark 예시는 Gluon matmul이 cuBLAS와 비슷한 수준의 TFLOPS를 보였다는 취지로 읽힌다.

6. 도구가 없으면 낮은 제어는 위험하다

Gluon은 낮은 수준을 노출하는 대신 tool 이야기도 같이 한다. Proton profiler는 Triton/Gluon kernel 최적화를 위한 낮은 overhead profiling으로 소개되고, host/device instrumentation과 PC sampling 같은 mode가 나온다.

sanitizer도 중요하다. concurrency sanitizer는 mbarrier를 기다리기 전 TMA load를 쓰는 식의 race를 잡고, invalid instruction sanitizer는 TMA alignment check를 넣는다. floating point sanitizer와 global memory sanitizer도 언급된다. 낮게 내려간 언어일수록 이런 난간이 없으면 생산성이 금방 무너진다.

Source Links

• YouTube
• GPU Mode lecture folder
• Slides PDF
• Gluon tutorials
• Gluon examples
• Linear layout visualizer
• Local notes: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/notes/lecture_104.md
• Local transcript stub: /Users/xavier/dev/gpumodelyrics/gpumode_lecture_archive/transcripts/lecture_104.md