# CUDA 커널을 실행하면 내부에서 벌어지는 일 > Clean Markdown view of GeekNews topic #30953. Use the original source for factual precision when an external source URL is present. ## Metadata - GeekNews HTML: [https://news.hada.io/topic?id=30953](https://news.hada.io/topic?id=30953) - GeekNews Markdown: [https://news.hada.io/topic/30953.md](https://news.hada.io/topic/30953.md) - Type: GN+ - Author: [xguru](https://news.hada.io/@xguru) - Published: 2026-06-30T08:58:31+09:00 - Updated: 2026-06-30T08:58:31+09:00 - Original source: [fergusfinn.com](https://fergusfinn.com/blog/what-happens-when-you-run-a-gpu-kernel/) - Points: 2 - Comments: 1 ## Topic Body - 단순한 벡터 덧셈 CUDA 프로그램도 결과 `2.000000`을 얻기까지 **컴파일 파이프라인**, 드라이버 호출, GPU 명령 큐, 워프 스케줄링, 메모리 계층, 완료 세마포어를 거침 - `nvcc`는 호스트 코드와 디바이스 코드를 나눠 `cicc`로 **PTX**, `ptxas`로 **SASS**를 만들고, cubin과 PTX를 fatbin에 묶어 Linux 실행 파일 안에 넣음 - `vadd<<<4096, 256>>>` launch 구문은 호스트 launch stub으로 바뀌며, 인자 `da`, `db`, `dc`, `n`은 CUDA 런타임과 `libcuda.so.1`을 거쳐 드라이버에 전달됨 - GPU 실행은 QMD, pushbuffer, GPFIFO, `GP_PUT`, **doorbell** MMIO 쓰기로 시작되고, RTX 4090의 128개 SM이 4096개 블록과 256개 스레드 구성을 워프 단위로 실행함 - 이 커널은 float 덧셈 1회당 12바이트 전송이 필요한 **낮은 산술 집약도** 때문에 Nsight Compute에서 10.78μs, DRAM 피크의 79.65%, warp issue 5.17%로 메모리 대역폭에 좌우됨 --- ### 예제 커널과 관찰 범위 - 예제 프로그램은 `vadd` CUDA 커널로 두 float 배열을 더해 세 번째 배열에 저장함 - `n = 1 << 20`으로 1,048,576개 float를 처리함 - launch 구성은 `vadd<<<4096, 256>>>(da, db, dc, n)`이며 `4096 * 256 = n`개 스레드를 사용함 - RTX 4090 대상으로 `nvcc -arch=sm_89`로 컴파일해 실행하면 `c[0]=2.000000 c[n-1]=2.000000`이 출력됨 - 이 한 줄의 결과에도 CPU 명령 수천만 개, device file, 약 900개의 `ioctl`, 메모리 매핑된 doorbell 레지스터가 관여함 ### `nvcc`가 실행 파일을 만드는 과정 - `nvcc --keep`를 사용하면 컴파일 파이프라인 산출물을 직접 확인할 수 있음 - `vadd.ptx`: `cicc`가 만든 디바이스 코드의 **PTX** - `vadd.sm_89.cubin`: `ptxas`가 만든 디바이스 코드의 **SASS** - `vadd.fatbin`: cubin과 PTX를 묶은 fatbin - `vadd.cudafe1.stub.c`: 호스트 launch stub과 커널 등록 코드 - `vadd.o`: fatbin이 포함된 최종 호스트 오브젝트 - 호스트 코드는 호스트 컴파일러로 처리되고, 디바이스 커널 `vadd`는 `cicc`와 `ptxas` 단계를 거침 - **PTX**는 가상 ISA로, 타입이 있는 무한한 가상 레지스터를 사용하며 실제 하드웨어 레지스터 수를 직접 반영하지 않음 - 예제 PTX는 `blockIdx.x * blockDim.x + threadIdx.x` 계산, 경계 검사, global load, float add, global store를 포함함 - CUDA 포인터는 기본적으로 generic pointer라서 `cvta.to.global`로 global address로 변환한 뒤 `ld.global`을 사용함 - `mul.wide.s32`는 index를 `sizeof(float)`인 4바이트 단위 오프셋으로 바꾸고 32비트에서 64비트로 확장함 - **SASS**는 아키텍처별 실제 명령어이며, RTX 4090 대상 출력에서는 PTX보다 더 압축된 형태로 나타남 - `S2R`은 `SR_CTAID.X`, `SR_TID.X` 같은 특수 레지스터를 일반 레지스터로 복사함 - PTX의 `mul.wide`와 `add` 조합은 SASS에서 `IMAD.WIDE`로 합쳐짐 - `cvta` 변환은 주소 지정 과정에 흡수됨 - `c[0x0][...]` 피연산자는 driver-managed **constant bank 0**을 가리킴 - 포인터 `a`, `b`, `c`는 `0x160`, `0x168`, `0x170`에 위치함 - `n`은 `0x178`에 위치함 - `blockDim.x` 같은 launch geometry와 ABI 값도 같은 bank에 있음 - cubin은 Linux 실행 파일과 같은 컨테이너 형식인 ELF 파일임 - fatbinary는 cubin과 PTX를 함께 묶음 - 이 RTX 4090에서는 SASS가 실제 실행되지만, PTX는 다른 아키텍처에서 드라이버가 JIT 컴파일할 수 있는 fallback으로 포함됨 - PTX는 verbose plain text라서 `nvcc`가 기본적으로 압축함 ### 호스트 코드가 launch를 준비하는 방식 - 컴파일러 프론트엔드 `cudafe++`는 `main` 이전에 실행되는 숨은 constructor를 삽입함 - 이 constructor는 embedded fatbinary를 CUDA 런타임에 등록함 - 호스트 쪽 함수 포인터 `vadd`와 fatbin 안의 mangled device kernel name을 연결함 - `vadd<<<4096, 256>>>(da, db, dc, n)` 구문은 생성된 **host launch stub**으로 바뀜 - `da`, `db`, `dc`, `n`은 host memory의 argument buffer에 각각 오프셋 `0`, `8`, `16`, `24`로 정렬되어 들어감 - 이 오프셋은 SASS가 constant bank 0에서 읽는 `0x160`, `0x168`, `0x170`, `0x178` 위치와 대응함 - stub은 `__cudaLaunch`를 호출하면서 호스트 쪽 dummy `vadd` 함수 주소를 넘김 - 이 주소는 CPU에서 실행할 함수 주소가 아니라 런타임 등록 테이블을 조회하는 key로 쓰임 - 런타임은 대응되는 device symbol name을 찾은 뒤 closed-source user-mode driver인 `libcuda.so.1`로 넘어감 - 첫 GPU 호출 시 CUDA 런타임은 `libcuda.so.1`을 동적으로 열고 context를 생성함 - `strace`에서는 `/lib/x86_64-linux-gnu/libcuda.so.1`이 열리는 것을 볼 수 있음 - context에는 CPU가 GPU와 통신하는 **channel**이 포함됨 - CUDA 12.2부터 module loading은 기본적으로 lazy임 - 특정 커널이 처음 launch될 때까지 SASS cubin 업로드를 미룸 - `CUDA_MODULE_LOADING`으로 제어 가능함 ### GPU에 작업을 전달하는 명령 큐 - GPU는 CPU처럼 함수 호출을 받아 entry point로 jump하지 않음 - PCIe bus 너머에서 host memory 안의 driver command stream을 읽음 - `cuLaunchKernel`은 완성된 launch command를 이 stream에 넣고 GPU에 알림 - 첫 실행에서는 driver가 커널 SASS를 GPU 메모리로 복사함 - code buffer를 할당하고 SASS를 복사함 - channel에는 host RAM에 있는 두 핵심 구조가 있음 - **pushbuffer**: driver가 GPU command인 method를 쓰는 메모리 영역 - **GPFIFO**: pushbuffer span을 가리키는 pointer ring buffer - GPFIFO entry는 pushbuffer span의 `(base, length)`를 나타내는 두 개의 32비트 word로 구성됨 - GPU와 driver는 두 cursor로 작업 소비와 생산 위치를 추적함 - `GP_GET`: GPU가 어디까지 소비했는지 나타냄 - `GP_PUT`: driver가 어디까지 생산했는지 나타냄 - 둘 다 USERD라는 per-channel 구조에 있음 - 커널 launch 시 driver는 pushbuffer span에 method를 쓰고, GPFIFO entry가 이를 가리키게 한 뒤 `GP_PUT`을 전진시킴 - 현대 GPU에서는 host engine이 cursor를 계속 감시하지 않으므로 **doorbell**이 필요함 - GPU는 process에 작은 register window를 mapping함 - driver는 channel의 work-submit token을 doorbell register에 씀 - host engine은 doorbell을 받은 뒤 `GP_PUT`을 읽고 GPFIFO entry와 pushbuffer span을 DMA로 가져감 ### QMD가 담는 실행 정보 - launch는 `SET_INLINE_QMD_ADDRESS_A/B`와 `LOAD_INLINE_QMD_DATA` method burst로 시작됨 - **QMD(Queue Meta Data)** 는 compute grid의 launch descriptor임 - grid와 block 크기인 `4096`, `256`을 포함함 - thread당 register 수와 shared memory 요구량을 포함함 - 프로그램 시작 주소와 커널 인자를 담은 constant bank 주소를 포함함 - 완료를 알릴 위치도 포함함 - host stub이 패킹한 인자들은 driver가 constant bank로 복사하고, QMD에 그 bank 주소가 기록됨 - QMD는 GPU에 SASS 위치, parallel program 구성 방식, 완료 signal 위치를 알려줌 - `cuLaunchKernel`은 doorbell이 울린 순간 반환함 - 호출은 비동기이므로 CPU는 GPU 작업이 진행되는 동안 계속 실행될 수 있음 ### SM, 워프, 점유율 - host engine은 QMD를 **compute work distributor**에 넘김 - 이 구성 요소는 GPU 전체에 하나 있음 - linear SASS instruction stream을 SM들에 분산해 병렬 프로그램으로 실행하게 함 - 대상 GPU인 GeForce RTX 4090은 **128 SM**을 사용함 - launch는 4096개 block과 block당 256 thread로 구성됨 - 각 SM은 local instruction cache를 가지고, active warp는 program counter를 유지함 - Volta 이후에는 thread별 program counter와 call stack을 갖는 Independent Thread Scheduling 모델이 있음 - issue는 여전히 warp 단위로 이루어짐 - 예제 커널에서는 resource limit이 block residency를 결정함 - block당 `256 threads = 8 warps` - `ptxas`는 thread당 16개 register를 예약함 - register 기준으로는 SM당 16개 block이 가능함 - thread capacity는 SM당 1,536 active threads라서 `1536 / 256 = 6`개 block만 가능함 - 따라서 SM당 최대 6개 block, 즉 48개 warp가 resident 상태가 됨 - SM은 4개 processing block, 즉 sub-partition으로 나뉨 - 48개 resident warp는 4개 sub-partition에 균등 분배됨 - 각 warp scheduler는 full 상태에서 12개 active warp를 관리함 - 매 cycle eligible warp 하나를 골라 32개 lane에 다음 명령을 dispatch함 ### 워프가 eligible 상태가 되는 조건 - GPU는 CPU의 out-of-order 실행처럼 단일 thread에서 동적 의존성을 크게 추출하지 않음 - 많은 resident warp를 두고 stall이 발생하면 다른 warp로 전환해 latency를 숨김 - 컴파일러가 예측 가능한 timing을 schedule하고, hardware scoreboard가 예측하기 어려운 부분을 처리함 - 128비트 SASS instruction에는 `ptxas`가 쓴 **control-code payload**가 들어 있음 - fixed-latency instruction에는 static stall count가 들어감 - yield hint는 scheduler priority를 양보할지 알려줌 - variable-latency operation에는 per-warp physical scoreboard barrier 6개가 사용됨 - 예제 SASS 구간에서 두 `LDG.E`는 같은 scoreboard barrier `B2`를 set함 - `FADD`는 `B2`를 wait-on으로 가짐 - 두 load가 돌아와 barrier가 clear되기 전까지 해당 warp는 ineligible 상태가 됨 - scheduler는 그동안 같은 sub-partition의 다른 warp를 고름 - `FADD`에서 `STG.E`로 넘어가는 구간은 fixed latency로 처리됨 - `FADD`는 `stall=5`를 갖고, `R9` 결과가 준비될 때까지 warp를 몇 cycle park함 - 별도 barrier는 필요하지 않음 - 이 control payload는 `nvdisasm` 기본 출력에서는 숨겨짐 - `cuobjdump -sass`의 raw 128-bit encoding에서 두 번째 64비트 word에 포함됨 - layout은 문서화된 것이 아니라 microbenchmarking으로 재구성된 것임 ### 메모리 접근과 성능 측정 - warp가 `LDG.E`를 실행하면 32개 thread가 각각 주소를 계산함 - 예제는 consecutive float array 접근이라 warp 전체가 `32 * 4 = 128 bytes` 연속 블록을 요청함 - SM load/store unit은 **request coalescing**을 수행함 - 32개의 4바이트 요청을 4개의 32바이트 sector request로 합침 - 연속 접근이 아니었다면 필요한 것보다 더 많은 데이터를 읽을 수 있음 - coalesced request는 먼저 SM local L1 Data Cache를 확인함 - miss가 나면 crossbar interconnect를 거쳐 72MB L2 Cache slice로 감 - L2에서도 miss가 나면 memory controller와 memory bus를 지나 GDDR6X VRAM으로 감 - `STG.E` store도 원칙적으로 반대 방향의 같은 경로를 따름 - Nsight Compute 측정값은 이 커널이 memory-bound임을 보여줌 - `launch__grid_size`: 4,096 - `launch__block_size`: 256 - `launch__registers_per_thread`: 16 - `launch__waves_per_multiprocessor`: 5.33 - `sm__warps_active.avg.pct_of_peak`: 82.77% - `smsp__issue_active.avg.pct_of_peak`: 5.17% - `dram__throughput.avg.pct_of_peak`: 79.65% - `gpu__time_duration.sum`: 10.78μs - 커널은 **산술 집약도**가 매우 낮음 - 두 4바이트 load와 한 4바이트 store, 총 12바이트 전송당 float add 1회를 수행함 - DRAM read 측면에서는 8.4MB를 10.78μs에 읽어 약 780GB/s이며, 피크의 약 4/5 수준임 - 4MB 출력 `c`는 72MB L2에 들어가므로 device-to-host copy가 읽기 전까지 DRAM으로 flush되지 않음 ### 결과가 CPU로 돌아오는 과정 - kernel launch는 doorbell을 울린 순간 CPU로 반환되므로, GPU는 완료 사실을 별도로 알려야 함 - 4096개 block이 모두 retire되면 GPU는 QMD에 담긴 completion semaphore를 post함 - QMD의 fence field는 words 23–24에 있음 - default stream에서 `cudaMemcpy(c, dc, ...)`는 kernel 뒤에 놓임 - GPU copy engine은 semaphore가 올라올 때까지 gated 상태가 됨 - `c`가 아직 72MB L2에 dirty 상태로 있으므로 copy engine read는 DRAM 왕복 없이 L2에서 처리됨 - 데이터는 PCIe를 넘어 host memory로 이동함 - copy가 끝나면 copy engine은 자체 semaphore를 post함 - host의 `cudaMemcpy` 대기가 끝남 - `c`는 다시 일반 host memory가 됨 - `printf`는 `c[0]`와 `c[n-1]`을 RAM에서 읽어 stdout으로 출력함 ### launch 내부를 들여다보는 방법 - open kernel modules를 읽는 것만으로는 `libcuda`가 closed-source라서 일부 동작을 직접 확인하기 어려움 - method write는 syscall을 거치지 않고 이미 mapping된 write-combined buffer에 직접 쓰이므로, pushbuffer를 보려면 memory를 읽어야 함 - `LD_PRELOAD` shim으로 `mmap`을 감싸 `/dev/nvidia*`에서 mapping된 영역을 기록할 수 있음 - test program이 launch 직후 shim의 dump 함수를 호출하면 mapped pushbuffer를 출력할 수 있음 - dump는 `SET_INLINE_QMD_ADDRESS_A`에 해당하는 method burst를 찾음 - pushbuffer method header는 opcode, payload count, subchannel index, register offset을 bit field로 담음 - `0x0318`은 `SET_INLINE_QMD_ADDRESS_A` - `0x0320 + i * 4`는 `LOAD_INLINE_QMD_DATA(i)` - dump에서는 count 66의 increasing-method burst가 보이며, address word 2개와 64개 QMD word, 총 256바이트 QMD가 inline으로 실림 - QMD 안의 word 12는 `0x1000`, word 18은 `0x100`으로 launch의 4096과 256에 해당함 - driver setup은 `ioctl`로 진행됨 - one-kernel program에서 `strace`는 948개의 `ioctl`을 기록함 - 대부분은 one-time setup임 - 주요 file descriptor는 `/dev/nvidiactl`과 `/dev/nvidia-uvm`임 - NVIDIA resource manager ioctl magic byte는 `0x46`, 즉 `'F'`임 - command number `0x2A`는 `NV_ESC_RM_CONTROL`, `0x2B`는 `NV_ESC_RM_ALLOC`로 해석됨 - `nvcc --keep`로 생성되는 `vadd.cudafe1.stub.c`에서는 startup registration 코드도 볼 수 있음 - `__attribute__((__constructor__))`가 붙은 함수가 `main` 전에 실행됨 - `__cudaRegisterBinary`와 `__cudaRegisterEntry`를 통해 host function pointer `vadd`와 device entry point `_Z4vaddPKfS0_Pfi`가 연결됨 ## Comments ### Comment 60772 - Author: neo - Created: 2026-06-30T08:58:32+09:00 - Points: 1 ###### [Hacker News 의견들](https://news.ycombinator.com/item?id=48718863) - 흥미로운 글이었고, **기본 스트림의 세마포어** 설명도 재미있었음 CUDA가 명령 동기화를 암묵적으로 처리해 주고, 병렬 명령은 스트림을 통해 선택적으로 쓰게 해 주는 점이 좋음 처음부터 동기화의 복잡성을 전부 사용자에게 떠넘기는 Vulkan과 대비됨 - 하드웨어 쪽은 일부 **공개 문서**가 있음 메서드 문서나 QMD 형식을 찾으려고 꼭 커널 소스를 읽을 필요는 없음 [https://github.com/NVIDIA/open-gpu-doc/blob/master/classes/c...]() 참고 - 매우 유용했음 특히 **doorbell과 QMD** 부분이 CUDA 실행 문법이 실제로 GPU에 제출되는 내용과 어떻게 이어지는지 보여줘서 가장 도움이 됨 대부분의 설명은 커널, 블록, 워프 근처에서 멈추는데, 이 글은 CPU→드라이버→GPU 경로를 훨씬 따라가기 쉽게 해 줌 - 제어 코드는 글에서 설명한 것보다 조금 더 복잡함 실제로는 제어 워드 안의 비트라기보다 **테이블 조회**에 가까움 - 지금은 커널을 최적화해서 더 빠르게 돌리는 것을 주 업무로 하는 회사들이 있음 그런 회사들이 언젠가 이를 아주 잘하는 **오픈소스 라이브러리**에 밀려날지 궁금함 Nvidia라면 언제든 그런 걸 내놓을 수도 있어 보임 아니면 대형 제공업체들이 추론 속도를 높이는 `moat`로 삼으려고 이 회사들을 인수하면서 더 잘될 수도 있음 - 단기적으로는 **인재 인수형 인수**가 꽤 가능성 있어 보임 다만 kernelbench 같은 관련 벤치마크에서 모델이 발전하는 걸 보면, 더 범용화된 해법들도 결국 나올 수밖에 없다고 봄 문제는 새 하드웨어 세대마다 기존 모델이 본 적 없는 제약이나 기능이 자주 생긴다는 점임 예를 들어 Blackwell의 tcgen05도 한때는 분포 밖 사례였음 모델이 더 잘 일반화하기 시작하면 치명적인 장벽은 아닐 수 있지만, 적어도 지금은 여전히 걸림돌임 [1] [https://kernelbench.com/]() - CUDA를 대규모로 돌리면 **Nvidia 드라이버와 라이브러리 버그**를 처리하는 데 엔지니어 시간이 역겨울 정도로 많이 들어감 Nvidia 라이브러리에 더 의존하는 걸 기대하는 사람은 별로 못 봄 - 아마 그렇지 않을 것 같음 작업 부하의 세부 사항, 즉 정확한 매개변수, 메모리 안의 데이터 표현, 값의 범위 등이 최적화 전략을 크게 갈라놓기 때문임 - HPC 석사를 막 끝냈고 CUDA, MPI+CUDA, OpenCL 수업을 들었는데, 수업 전에 이런 글을 읽었으면 훨씬 도움이 됐을 듯함 특히 **워프가 실행 가능하다는 뜻**을 다룬 부분의 앞뒤가 좋았음 - 먼저, 여러 구석구석을 잘 파고든 좋은 글임 다만 CUDA의 `runtime API`를 거치지 않으면 사용자 공간의 많은 **부두교 같은 부분**이 사라짐 드라이버 API를 쓰고, 커널 소스를 문자열로 받아 NVIDIA의 런타임 컴파일러로 컴파일하면 무슨 일이 벌어지는지 더 잘 볼 수 있음 전부는 아니지만 꽤 많은 부분이 투명해짐 더 “원시적인” 버전은 여기 있음: [https://github.com/NVIDIA/cuda-samples/tree/master/cpp/0_Int...]() 같은 내용을 훨씬 읽기 쉽고, 그래도 완전히 투명한 현대 C++ API 형태로 보려면 이걸 보면 됨: [https://github.com/eyalroz/cuda-api-wrappers/blob/master/exa...]() 내 CUDA API wrappers 헤더 전용 라이브러리의 예제 프로그램임 - 드라이버 API는 CUDA 커널을 **핫 리로드 가능한 셰이더**처럼 다룰 수 있어서 좋음 실행 중에 코드를 바꿔가며 개발할 수 있어서 재미있음 - 베어메탈에서?