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

 (fergusfinn.com)
17P by GN⁺ | ★ favorite | 댓글 1개
  • 단순한 벡터 덧셈 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가 연결됨

함께 보면 좋은 글 β

GN⁺   [-]
Hacker News 의견들

  • 흥미로운 글이었고, 기본 스트림의 세마포어 설명도 재미있었음
    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 커널을 핫 리로드 가능한 셰이더처럼 다룰 수 있어서 좋음
      실행 중에 코드를 바꿔가며 개발할 수 있어서 재미있음

  • 베어메탈에서?

답변달기