강의 번호
L037
PTX 가 NVIDIA 의 “virtual ISA” 라면, SASS 는 진짜 칩 위에서 도는 머신코드. NVIDIA 가 공식 문서를 안 내는 이 layer 를 어떻게 읽고, 그게 왜 마지막 30% 의 성능을 결정하는지 — Arun Demeure 가 cuobjdump 화면 위에서 한 줄씩 따라가며 깐다. SM 의 issue slot 과 functional unit, ILP 의 의미, warp scheduler 의 세 가지 stall 원인, 그리고 “Triton 도 결국 SASS 까지 내려간다” 의 진짜 의미.
대부분의 GPU 프로그래밍은 PTX 까지만 본다 — 그게 NVIDIA 가 공식 문서를 내주는 마지막 layer 니까. 그런데 진짜 칩 위에서는 SASS 가 돈다. PTX 와 SASS 사이에는 NVIDIA 의 후처리 컴파일러 가 끼어 있고, 그게 instruction reordering, register 재할당, scoreboard 결정을 다 한다.
Arun 의 핵심 입장 — “Triton 으로 짠 좋은 코드도, 직접 쓴 CUDA 도, 결국 같은 SASS 로 컴파일된다. 차이를 결정하는 건 그 SASS 의 모양.” Triton 이 cuBLAS 의 80% 까지 따라잡았는데도 마지막 20% 가 안 풀리는 이유는 PTX 단계에서는 안 보이고, SASS 를 까봐야 보인다.
“PTX 만 보고 성능을 추론하지 마라” — 이게 강의의 가장 큰 메시지. PTX 의 instruction 한 개가 SASS 에서 한 개로 사상되지 않는다. PTX 가 깔끔해 보여도 SASS 에서 spill, redundant load, scheduling stall 이 들어 있을 수 있다. SASS 가 진실.
그래서 강의 끝에 손에 잡혀야 할 자산은 — (1) SASS 한 줄을 읽는 법, (2) SM 의 issue slot 과 functional unit 의 매핑이라는 mental model, (3) warp scheduler 가 한 cycle 마다 무슨 결정을 하는가, (4) cuobjdump / godbolt 로 자기 커널을 까보는 워크플로.
위에서 가장 흥미로운 건 L4. SASS 의 매 instruction 옆에 “이 결과는 4 cycle 뒤에 ready 됨” 이나 “이 register 는 cache 에서 읽어라” 같은 control bits 가 붙어 있다. 이게 hardware scheduler 가 stall 결정을 하는 근거. NVIDIA 의 후처리 컴파일러가 이 control bits 를 결정 — 그래서 같은 PTX 에서 다른 SASS 가 나올 수 있고, scheduling 이 다른 만큼 latency 가 다르다.
강의에서 Arun 이 짚은 한 줄 — “PTX 의 instruction 수와 SASS 의 instruction 수가 다르다. PTX 한 줄이 SASS 여러 줄로 펴지기도 하고, PTX 여러 줄이 SASS 하나로 합쳐지기도 한다.” 그래서 PTX 만 보고 “register 사용량이 적다” 같은 결론을 내리면 SASS 에서 다르게 나올 수 있다.
이 표는 단순한 hardware 이력이 아니다 — 같은 PTX 가 어느 sm_xx 으로 컴파일되느냐에 따라 SASS 가 통째로 달라진다는 사실의 근거. 예를 들어 sm_80 에서 cp.async 가 한 instruction 이지만, sm_70 에서는 일반 load 두 개로 쪼개져서 register 트래픽이 두 배가 된다.
Triton 은 같은 tl.dot 코드를 어느 sm 으로 컴파일하느냐에 따라 다른 SASS instruction 으로 매핑한다 — sm_80 이면 mma.sync, sm_90 이면 wgmma. 그래서 같은 Triton 코드의 SASS 가 GPU 마다 다른 모양이 된다는 사실이 한 번 손에 잡혀야 한다.
강의에서 Arun 이 든 구체적 예 — H100 에서 thread block cluster 가 도입되며 distributed shared memory 가 가능. 같은 PTX 의 ld.shared 가 cluster scope 에서는 다른 SM 의 SMEM 까지 본다. 이런 사실은 PTX manual 에는 한 줄로 적히지만, SASS 에서는 새 instruction (LDGSTS 의 변종) 으로 펴진다.
SM 안의 warp scheduler 는 매 cycle ready warp pool 을 본다. 그 pool 이 비어 있으면 그 cycle 의 functional unit 이 비워진다. 이게 occupancy 가 중요한 이유의 본질.
강의에서 Arun 이 명시한 한 줄 — “warp 가 64 개 거주해도, 그중 ready 인 게 매 cycle 충분하지 않으면 occupancy 만 높고 utilization 은 낮다.” 이 차이가 NCU 의 “Achieved Active Warps Per Scheduler” 와 “Warp Cycles Per Issued Instruction” 두 metric 의 차이.
(a) Long Scoreboard — 메모리 latency 로 register 가 ready 가 안 됨. (b) Short Scoreboard — TC / MUFU 결과 대기. (c) Dispatch Stall — issue slot 이 비어서 그 cycle 못 issue. (a) 가 가장 흔하고, 보통 더 많은 warp 또는 더 작은 footprint 로 푼다.
SM 한 개에는 여러 functional unit 이 있고, 그 unit 들이 동시에 돈다. FP32 FMA 가 도는 cycle 에 LSU 가 다른 warp 의 load 를 처리할 수 있고, tensor core 가 또 다른 warp 의 MMA 를 진행 중일 수 있다 — 이게 GPU 가 “bandwidth bound” 와 “compute bound” 를 동시에 다룰 수 있는 이유.
강의에서 Arun 이 강조한 사실 — tensor core 가 매우 빠르긴 한데, 그게 “FFMA 가 비어 있다” 의 의미가 아니다. softmax 같은 elementwise 연산은 FMA 와 MUFU pipe 위에서 도는데, 그게 tensor core 와 동시에 돌아야 attention 이 효율적으로 도는 (§ FA3 의 ping-pong).
강의의 실용 부분. SASS 를 읽으려면 — (1) cuobjdump --dump-sass 로 binary 에서 추출, 또는 (2) godbolt.org 의 NVCC 모드로 소스 → SASS 한 화면. 후자가 학습용으로 압도적으로 빠르다.
# Triton 으로 짠 fp32 squared sum 의 SASS 일부 (sm_80, A100) # godbolt 에서 그대로 본 모양. 색은 임의 — instruction 영역만. /*0050*/ LDG.E.SYS R6, [R2.64+0x0] ; // global load · long lat /*0058*/ LDG.E.SYS R7, [R2.64+0x4] ; // global load · long lat /*0060*/ FMUL R8, R6, R6 ; // x * x /*0068*/ FMUL R9, R7, R7 ; // y * y /*0070*/ FFMA R10, R8, R9, R10 ; // acc += x*x + y*y /*0078*/ @!P0 BRA `(.L_x_0) ; // loop branch /*0080*/ STG.E.SYS [R12.64], R10 ; // store result
읽을 때 손에 잡히는 사실들 — LDG.E.SYS 가 global load (E = 64-bit address). FFMA 가 fused multiply-add (한 cycle throughput). @!P0 BRA 가 predicated branch. /*0050*/ 가 instruction 의 byte offset (8 byte 씩 증가하니 instruction 당 8 byte 인 sm_70+ 패턴).
R0…Rn 일반 register, P0…P6 predicate. SR special register (threadIdx 등). S2R special-to-register (idx load). UR uniform register (sm_75+). BAR.SYNC barrier. NOP 진짜 비어 있는 cycle.
godbolt 의 NVCC compiler 옵션 — -arch=sm_80 -O3 --ptxas-options=-v — 가 SASS 를 같이 보여준다. Triton 도 같은 trick — TRITON_DEBUG=1 또는 ~/.triton/cache/ 의 디렉토리 안에 *.cubin + *.ptx 가 같이 떨어진다.
진짜 hardware 가 보는 SASS 는 instruction + control bits. cuobjdump 는 기본으로 control bits 를 안 보여주는데, --dump-sass-noindex 류 옵션 또는 별도 도구로 까보면 — 매 instruction 옆에 “이 결과는 N cycle 뒤 ready”, “register cache 에서 읽어라”, “scoreboard B0 까지 기다려라” 같은 hint 가 박혀 있다.
구조적으로 hardware 는 — (1) long-latency instruction (LDG, MMA) 을 issue 할 때 어떤 scoreboard slot (B0..B5) 에 결과를 기록할지 결정, (2) 그 결과를 읽는 instruction 은 그 slot 이 ready 될 때까지 wait. 이 wait 가 명시적으로 제어 bit 로 박혀 있어 hardware scheduler 가 단순화된다.
NVIDIA 의 후처리 컴파일러가 이 control bits 를 영리하게 채워 넣는다. 같은 PTX 를 다른 sm_xx 으로 컴파일하면 SASS 의 control bits 가 달라진다. 그게 latency-bound 한 코드에서 결정적인 차이를 만든다 — 명시적 wait 가 길게 박힌 SASS 와, 영리하게 짧은 wait 만 박힌 SASS 사이.
강의에서 Arun 이 구체적 사례로 든 — backward 가 forward 보다 느려 보일 때, PTX 는 비슷한데 SASS 의 scoreboard wait 분포가 다르다. PTX 만 보고는 안 보이는 차이.
강의의 가장 비실용적이지만 가장 흥미로운 부분. ILP (Instruction-Level Parallelism) — 한 thread 안에서도 독립적인 instruction 들이 동시에 issue 될 수 있다. 그래서 thread 수가 많지 않아도 throughput 이 산다.
예시 — A100 의 FFMA pipe 의 latency 가 4 cycle, throughput 이 1/cycle. 만약 한 thread 가 매 4 cycle 마다 dependent 한 FFMA 한 개만 발행하면 throughput 의 25% 만 사용. 같은 thread 가 4개 독립 FFMA 를 매 cycle 하나씩 발행하면 — 같은 thread 만으로도 4배 이득.
강의에서 Arun 이 짚은 한 줄 — “40 cycle latency / 10 cycle 마다 한 개 throughput 인 unit 은 적어도 40 개 thread/warp 가 동시에 inflight 해야 latency 를 숨긴다.” 이게 occupancy + ILP 의 합. ILP 를 충분히 짜면 occupancy 를 낮춰도 됨.
sm_75+ 의 uniform register (UR) — warp 안 32 thread 가 같은 값을 가지는 register 를 별도 file 로 둔다. compiler 가 이걸 잘 활용하면 일반 register 의 압박이 줄어든다. PTX 에서는 직접 noted 안 되고 SASS 에서만 보임. “같은 일을 누가 하면 좋을까” 의 결정 단위가 GPU 에서 점점 다양해진다.
SM 의 진화를 한 줄로 요약하면 — 점점 더 많은 일이 비동기로 issue 가능해진다. Volta 의 tensor core 도, Ampere 의 cp.async 도, Hopper 의 TMA / WGMMA / cluster 도 모두 같은 방향. 한 thread 가 “큰 일을 시작하고 다른 일을 한다” 가 본질.
cuobjdump --dump-sass binary. godbolt 의 NVCC 모드도 같은 일을 학습용으로 빠르게.cuobjdump --dump-sass binary.cubin · CUDA toolkit 동봉__global__ void k(float*x){x[0]=x[0]*x[0];}. NVCC sm_80 으로 컴파일해서 SASS 한 화면. LDG → FMUL → STG 패턴 직접 본다.~/.triton/cache/ 에 cubin. cuobjdump --dump-sass 로 까서 자기 Triton 코드의 진짜 모양 확인.이 노트의 모든 SASS instruction 형태는 강의 화면을 재구성한 예시. LDG.E.SYS, FFMA 같은 mnemonic 은 sm_80 기준이고, sm_90 은 다른 변종이 있다. 자기 GPU 에서 직접 컴파일해서 본 SASS 가 진실.