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

9bow · 7월 1, 2026, 9:30오전

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으로 출력함

전체 경로

커널 소스는 cicc로 PTX, ptxas로 SASS가 되고, fatbinary가 PTX fallback과 함께 cubin을 담은 fatbin으로 묶어 linker가 평범한 Linux 실행 파일에 붙임. main 이전에 실행되는 constructor가 그 fatbin을 등록하며 호스트 stub을 mangled device name에 연결함. 첫 launch가 cubin을 GPU로 lazy 업로드함. cuLaunchKernel은 launch 구성으로 QMD를 만들어 pushbuffer에 GPU method로 쓰고 GP_PUT을 전진시킨 뒤 단일 MMIO 쓰기로 doorbell을 울리며, 그 순간 GPU의 host engine이 작업을 가져와 QMD를 compute work distributor에 넘김. distributor가 4096개 block을 128개 SM에 full occupancy로 분산하고, SM당 4개 warp scheduler가 컴파일러가 써둔 stall count를 가진 128비트 명령을 issue하며, coalesced 메모리 경로가 DRAM 피크의 4/5 대역폭으로 입력을 끌어와 100만 개 lane 각각에서 덧셈 1회를 계산함. 그다음 완료 세마포어와 copy engine이 그 결과를 버스 너머 printf가 기다리던 곳으로 다시 가져오고, 우리는 다음을 알게 됨:

c[0]=2.000000 c[n-1]=2.000000

부록: launch 내부를 들여다보는 방법

커널 launch의 각 부분을 관찰하는 데에는 여러 기법이 쓰임. 일부는 open kernel modules를 정독해 얻을 수 있지만, libcuda가 closed-source라 소스만으로는 확인할 수 없는 주장은 아래 진단 훅으로 알아냄.

인터포지션 훅(interposition hook)

driver의 method write는 syscall을 거치지 않고 이미 mapping된 write-combined buffer에 직접 쓰이므로, 보려면 memory를 읽어야 함
LD_PRELOAD shim으로 mmap을 감싸 /dev/nvidia* 파일에서 mapping된 모든 영역을 기록하고, launch 직후 test program이 호출하는 dump 함수로 그 영역을 출력함
shim을 공유 라이브러리로 컴파일(gcc -shared -fPIC -o shim.so shim.c -ldl)한 뒤 LD_PRELOAD=./shim.so ./vadd로 실행하면, driver가 channel용으로 mapping한 write-combined buffer를 훑어 launch의 method burst를 덤프함

pushbuffer 명령 스트림 디코딩

pushbuffer method는 header word 1개와 그 뒤의 data word들로 구성됨. header는 clc46f.h의 NVC46F_DMA_INCR_* 매크로로 정의된 4개 필드를 패킹함
- bits 31:29 opcode: 0x1은 increasing-method write, 0x3은 non-increasing write, 0x4는 immediate-data write
- bits 28:16 count: payload word 수
- bits 15:13 subchannel index: 명령을 특정 백엔드 엔진 컨텍스트로 라우팅
- bits 11:0 method의 register offset을 4로 나눈 값
compute class별 method는 아키텍처마다 정의됨: clc3c0.h(Volta), clc5c0.h(Turing), clc6c0.h/clc7c0.h(Ampere), clc9c0.h(Ada), clcbc0.h(Hopper), clcdc0.h(Blackwell). Ada의 clc9c0.h는 클래스 번호 0xC9C0만 정의하고 Ampere method set을 상속하는 29줄 stub이라, 실제 정의는 Ampere 헤더에서 읽음
- 0x0318 SET_INLINE_QMD_ADDRESS_A: inline-QMD burst를 pushbuffer에 직접 스트리밍하는 경로(LOAD_INLINE_QMD_DATA(i), offset 0x0320 + i * 4)
- 0x02b4 SEND_PCAS_A: out-of-line 경로. VRAM의 다른 곳에 있는 QMD를 가리키는 포인터만 전달
덤프는 inline 경로를 보여줌: SET_INLINE_QMD_ADDRESS_A에서 시작하는 count 66의 increasing-method burst. 66개 word는 address word 2개(0x0318/0x031c)와 LOAD_INLINE_QMD_DATA word 64개(0x0320부터)로, 256바이트 QMD가 inline으로 실림. 그 안에서 word 12는 0x1000, word 18은 0x100으로 vadd<<<4096, 256>>>의 4096과 256에 해당함

디바이스 메모리 읽기와 QMD 레이아웃

QMD 구조는 cla0c0qmd.h에 32비트 경계를 넘는 multi-word(MW) 비트 필드로 정의됨. 여러 주소형 필드를 담지만 종류가 다름
- PROGRAM_OFFSET MW(287:256)(Word 8): channel code base 기준 32비트 entry-point offset(64비트 포인터 아님)
- CONSTANT_BUFFER_ADDR_LOWER(i)/ADDR_UPPER(i): 인자를 담은 Constant Bank 0가 Word 29-30에 위치
- RELEASE0_ADDRESS_LOWER/UPPER MW(767:736)(Word 23-24): fence/세마포어용
- CIRCULAR_QUEUE_ADDR_LOWER/UPPER MW(319:288)(Word 9-10)
이 주소들은 CPU가 직접 못 읽음(plain load는 fault, cudaMemcpy/cuMemcpyDtoH도 거부). 그래서 작은 peek 커널로 GPU를 통해 512바이트를 읽어냄

__global__ void peek(const unsigned char* src, unsigned char* dst) {
    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
         i < 512;
         i += blockDim.x * gridDim.x) {
        dst[i] = src[i];
    }
}

QMD의 각 주소 필드를 가리켜 보면 정확히 하나가 512바이트 SASS 전체를 반환함. 매치는 Word 48에서 나옴(qmd[48] -> 0x74167b272300 512 / 512 bytes match)
driver의 program 필드는 Word 8(PROGRAM_OFFSET)인데 SASS가 Word 48에서 매치되는 이유: Word 8은 32비트 offset만 담고, word 48/49는 하드웨어 소유 필드인 HW_ONLY_INNER_GET(MW(1566:1536))/HW_ONLY_INNER_PUT(MW(1598:1568))임. launch 후 덤프에서 이 word들은 완전한 64비트 GPU 가상 주소를 담고, word 48 값을 역참조하면 커널 SASS가 나옴. scheduler가 launch 시 program offset을 이 scheduler 소유 필드로 resolve하는 것으로 읽힘

driver의 ioctl 디코딩

libcuda는 메모리와 GPU 객체를 driver device file에 ioctl을 걸어 평범하게 설정함. one-kernel program에 strace를 걸면 948개가 기록되며(대부분 one-time setup이고, 꾸준한 launch loop는 훨씬 적게 발생), 거의 전부 /dev/nvidiactl과 /dev/nvidia-uvm 두 file descriptor에 집중됨
magic byte 0x46은 'F'로 NVIDIA resource manager의 ioctl magic임. command number는 open kernel modules의 nv_escape.h로 디코딩됨: 0x2A는 NV_ESC_RM_CONTROL, 0x2B는 NV_ESC_RM_ALLOC

SASS 제어 워드 디코딩

eligibility 부분의 stall count, barrier, yield 비트는 ptxas가 각 명령의 두 번째 64비트 word 상단에 패킹하는 21비트 control field에서 옴. cuobjdump -sass가 mnemonic 옆에 출력함

 20    17 16       11 10   8 7    5 4 3    0
┌────────┬───────────┬──────┬──────┬─┬──────┐
│ reuse  │ wait mask │ read │write │Y│stall │
│  (4)   │    (6)    │ barr │ barr │ │ (4)  │
└────────┴───────────┴──────┴──────┴─┴──────┘

두 3비트 인덱스는 명령이 set하는 scoreboard barrier를, 6비트 mask는 wait할 barrier를, Y는 yield 비트를, stall은 static cycle count를 나타냄. 레이아웃은 문서화된 것이 아니라 microbenchmarking으로 재구성된 것임. 가장 명확한 공개 재구성은 Citadel의 Volta 아키텍처 분석 논문과 maxas control-code 노트임

NVCC 호스트 등록 콜백

nvcc --keep로 생성되는 vadd.cudafe1.stub.c에서 컴파일러가 GPU 코드를 startup에 등록하는 코드를 직접 볼 수 있음
__attribute__((__constructor__))가 붙은 __sti____cudaRegisterAll이 main 전에 실행되어 device binary를 CUDA 런타임에 등록함. 실행되면 __cudaRegisterEntry가 호스트 함수 포인터 vadd를 mangled device entry point _Z4vaddPKfS0_Pfi에 매핑해, cudaLaunchKernel이 launch 시 조회하는 해시 테이블을 구성함

// vadd.cudafe1.stub.c
static void __sti____cudaRegisterAll(void) __attribute__((__constructor__));

static void __nv_cudaEntityRegisterCallback(void **__T4) {
    __cudaRegisterEntry(__T4, (void(*)(const float*, const float*, float*, int))vadd,
                        _Z4vaddPKfS0_Pfi, -1);
}

static void __sti____cudaRegisterAll(void) {
    __cudaRegisterBinary(__nv_cudaEntityRegisterCallback);
}

Hacker News 의견

흥미로운 글이었고, 기본 스트림의 세마포어 설명도 재미있었음. CUDA가 명령 동기화를 암묵적으로 처리해 주고, 병렬 명령은 스트림을 통해 선택적으로 쓰게 해 주는 점이 좋음. 처음부터 동기화의 복잡성을 전부 사용자에게 떠넘기는 Vulkan과 대비됨
하드웨어 쪽은 일부 공개 문서가 있음. 메서드 문서나 QMD 형식을 찾으려고 꼭 커널 소스를 읽을 필요는 없음 (GitHub - NVIDIA/open-gpu-doc: Documentation of NVIDIA chip/hardware interfaces · GitHub 참고)
매우 유용했음. 특히 doorbell과 QMD 부분이 CUDA 실행 문법이 실제로 GPU에 제출되는 내용과 어떻게 이어지는지 보여줘서 가장 도움이 됨. 대부분의 설명은 커널, 블록, 워프 근처에서 멈추는데, 이 글은 CPU에서 드라이버, GPU로 이어지는 경로를 훨씬 따라가기 쉽게 해 줌
제어 코드는 글에서 설명한 것보다 조금 더 복잡함. 실제로는 제어 워드 안의 비트라기보다 테이블 조회에 가까움
지금은 커널을 최적화해서 더 빠르게 돌리는 것을 주 업무로 하는 회사들이 있음. 그런 회사들이 언젠가 이를 아주 잘하는 오픈소스 라이브러리에 밀려날지 궁금함. Nvidia라면 언제든 그런 걸 내놓을 수도 있어 보이고, 아니면 대형 제공업체들이 추론 속도를 높이는 moat로 삼으려고 이 회사들을 인수하면서 더 잘될 수도 있음. 단기적으로는 인재 인수형 인수가 꽤 가능성 있어 보임
- 다만 kernelbench 같은 관련 벤치마크에서 모델이 발전하는 걸 보면, 더 범용화된 해법들도 결국 나올 수밖에 없다고 봄. 문제는 새 하드웨어 세대마다 기존 모델이 본 적 없는 제약이나 기능이 자주 생긴다는 점임. 예를 들어 Blackwell의 tcgen05도 한때는 분포 밖 사례였음. 모델이 더 잘 일반화하기 시작하면 치명적인 장벽은 아닐 수 있지만, 적어도 지금은 여전히 걸림돌임 (https://kernelbench.com/)
- CUDA를 대규모로 돌리면 Nvidia 드라이버와 라이브러리 버그를 처리하는 데 엔지니어 시간이 역겨울 정도로 많이 들어감. Nvidia 라이브러리에 더 의존하는 걸 기대하는 사람은 별로 못 봄. 작업 부하의 세부 사항, 즉 정확한 매개변수, 메모리 안의 데이터 표현, 값의 범위 등이 최적화 전략을 크게 갈라놓기 때문임
HPC 석사를 막 끝냈고 CUDA, MPI+CUDA, OpenCL 수업을 들었는데, 수업 전에 이런 글을 읽었으면 훨씬 도움이 됐을 듯함. 특히 워프가 실행 가능하다는 뜻을 다룬 부분의 앞뒤가 좋았음
먼저, 여러 구석구석을 잘 파고든 좋은 글임. 다만 CUDA의 runtime API를 거치지 않으면 사용자 공간의 많은 부두교 같은 부분이 사라짐. 드라이버 API를 쓰고, 커널 소스를 문자열로 받아 NVIDIA의 런타임 컴파일러로 컴파일하면 무슨 일이 벌어지는지 더 잘 볼 수 있음. 더 "원시적인" 버전은 NVIDIA cuda-samples(GitHub - NVIDIA/cuda-samples: Samples for CUDA Developers which demonstrates features in CUDA Toolkit · GitHub)에, 같은 내용을 훨씬 읽기 쉽고 투명한 현대 C++ API로 보려면 cuda-api-wrappers(GitHub - eyalroz/cuda-api-wrappers: Thin, unified, C++-flavored wrappers for the CUDA APIs · GitHub)에 있음
- 드라이버 API는 CUDA 커널을 핫 리로드 가능한 셰이더처럼 다룰 수 있어서 좋음. 실행 중에 코드를 바꿔가며 개발할 수 있어서 재미있음
  - 베어메탈에서?

원문

출처 / GeekNews

더 읽어보기

[GN] 함께 보면 좋은 글β

알려드립니다

이 글은 국내외 IT 소식들을 공유하는 GeekNews의 운영자이신 xguru님께 허락을 받아 GeekNews에 게제된 AI 관련된 소식을 공유한 것입니다.

출처의 GeekNews 링크를 방문하시면 이 글과 관련한 추가적인 의견들을 보시거나 공유하실 수 있습니다!

아래쪽에 좋아요를 눌러주시면 새로운 소식을 정리하고 공유하는데 힘이 됩니다~