CUDA 커널을 실행하면 내부에서 벌어지는 일

• 단순한 벡터 덧셈 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가 연결됨