단순한 벡터 덧셈 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

가 연결됨
댓글과 토론