--- title: "쿠다 커널 한 줄이 실제 지피유에서 실행되기까지 벌어지는 일" published: 2026-06-29T13:11:08.000Z canonical: https://jeff.news/article/4498 --- # 쿠다 커널 한 줄이 실제 지피유에서 실행되기까지 벌어지는 일 단순한 벡터 덧셈 쿠다 커널이 컴파일, 런타임 등록, 드라이버 호출, 지피유 명령 큐, 와프 스케줄링, 메모리 계층을 거쳐 실행되는 과정을 깊게 추적한 글이다. 엔비디아 RTX 4090에서 4096개 블록과 256개 스레드로 104만8576개 float를 더하는 예제를 사용하며, 실행 뒤에는 900개가 넘는 ioctl, QMD, GPFIFO, 도어벨 레지스터, 와프 스코어보드 같은 저수준 메커니즘이 동원된다. ## 커널 하나가 생각보다 긴 여행을 함 - 글의 출발점은 진짜 단순한 쿠다(CUDA) 벡터 덧셈 예제임 - `vadd<<<4096, 256>>>`로 커널을 실행함 - 전체 스레드 수는 4096 곱하기 256, 즉 104만8576개 - 각 스레드는 `c[i] = a[i] + b[i]` 한 번만 수행함 - 결과는 당연히 `c[0]=2.000000 c[n-1]=2.000000`임 - 그런데 이 당연한 출력 뒤에는 수천만 개 CPU 명령어, 장치 파일 몇 개, ioctl 약 900개, 메모리 매핑된 도어벨 레지스터가 얽힘 - 글의 재미는 “더하기 한 번”이 아니라 그 더하기가 CPU와 GPU 사이에서 어떻게 실제 작업으로 바뀌는지 끝까지 따라가는 데 있음 > [!IMPORTANT] > 이 예제 커널 자체는 산술이 병목이 아님. float 덧셈 1번을 위해 입력 2개와 출력 1개, 총 12바이트 수준의 데이터를 움직이기 때문에 사실상 메모리 대역폭 벤치마크에 가까움. ## nvcc는 하나의 컴파일러가 아니라 파이프라인임 - `nvcc`는 CUDA 소스를 한 번에 GPU 실행 파일로 바꾸는 마법 버튼이 아니라 여러 컴파일러를 호출하는 드라이버임 - 호스트 코드는 일반 C/C++ 컴파일러로 감 - 디바이스 코드인 `vadd`는 `cicc`를 거쳐 PTX가 됨 - 그 PTX를 `ptxas`가 실제 GPU 아키텍처용 SASS로 변환함 - 마지막엔 cubin과 PTX가 fatbin으로 묶여 실행 파일 안에 들어감 - PTX는 가상 ISA라서 실제 하드웨어 제약을 직접 반영하지 않음 - typed virtual register가 많고, 주소 계산도 꽤 장황함 - 예를 들어 `a[i]` 주소 하나를 만들기 위해 포인터 변환, 인덱스 곱셈, 베이스 주소 더하기 같은 단계가 보임 - CUDA 포인터는 기본적으로 global, shared, local 메모리 중 어디든 가리킬 수 있어서 `cvta.to.global` 같은 변환도 등장함 - SASS로 내려가면 훨씬 하드웨어 냄새가 남 - `S2R`은 `blockIdx.x`, `threadIdx.x` 같은 특수 레지스터 값을 일반 레지스터로 옮김 - `IMAD.WIDE`는 인덱스 곱셈과 주소 더하기를 한 번에 처리함 - `LDG.E`, `FADD`, `STG.E`가 실제 로드, 덧셈, 저장 흐름을 이룸 - 커널 인자도 그냥 스택에 들어가는 게 아니라 constant bank 0에 놓임 - 포인터 `a`, `b`, `c`와 크기 `n`이 고정 오프셋에 배치됨 - 모든 스레드가 같은 인자를 읽기 때문에 constant cache의 브로드캐스트 특성이 잘 맞음 - 글에서는 SASS의 `c[0x0][...]` 오프셋과 호스트 스텁의 인자 패킹 오프셋이 연결되는 걸 보여줌 ## 호스트에서 GPU로 넘어가는 순간 - CUDA 커널 호출 문법은 컴파일러가 호스트 launch stub으로 바꿔 넣음 - `vadd<<<4096, 256>>>(da, db, dc, n)`는 실제로 인자를 버퍼에 정렬해 넣고 `__cudaLaunch`를 호출하는 코드가 됨 - 호스트 쪽 `vadd` 함수 포인터는 CPU에서 실행할 함수가 아니라 디바이스 커널을 찾는 lookup key처럼 쓰임 - 프로그램 시작 전 생성자가 fatbin을 CUDA 런타임에 등록하고, 호스트 함수 포인터와 디바이스 심볼 이름을 매핑해둠 - 첫 GPU 호출 시 CUDA 런타임은 `libcuda.so.1`을 열고 드라이버 쪽으로 넘어감 - 글의 환경에선 `strace`로 `/lib/x86_64-linux-gnu/libcuda.so.1`을 여는 걸 확인함 - 사용자 모드 드라이버 아래에는 커널 모드 드라이버 `nvidia.ko`가 있고, 둘 사이 통신은 장치 파일과 ioctl로 이뤄짐 - CUDA 12.2부터는 모듈 로딩이 기본적으로 lazy라서, 특정 커널의 SASS는 첫 실행 시점에야 GPU 메모리로 올라감 ```mermaid sequenceDiagram participant 호스트코드 as 호스트 코드 participant 쿠다런타임 as CUDA 런타임 participant 드라이버 as NVIDIA 드라이버 participant 명령큐 as Pushbuffer/GPFIFO participant 지피유 as GPU 호스트코드->>쿠다런타임: 커널 스텁이 인자 패킹 쿠다런타임->>드라이버: 디바이스 심볼 조회 후 실행 요청 드라이버->>명령큐: QMD와 메서드 기록 드라이버->>지피유: 도어벨 레지스터 쓰기 지피유->>명령큐: GPFIFO와 pushbuffer 읽기 지피유->>지피유: SM에 블록과 와프 배포 지피유-->>호스트코드: 완료 세마포어와 복사 완료 ``` ## GPU는 함수 호출을 받지 않고 큐를 읽음 - CPU가 GPU의 함수를 직접 호출하는 구조는 아님 - GPU는 PCIe 너머에 있고, 호스트 메모리에 놓인 명령 스트림을 읽음 - 드라이버는 pushbuffer에 GPU 명령을 쓰고, GPFIFO에는 그 pushbuffer 범위를 가리키는 엔트리를 넣음 - `GP_PUT`은 드라이버가 어디까지 작업을 생산했는지, `GP_GET`은 GPU가 어디까지 소비했는지를 나타냄 - 커널 실행 정보는 QMD(Queue Meta Data)에 담김 - 그리드 4096, 블록 256 같은 launch geometry - 스레드당 레지스터 수, 공유 메모리 사용량 - SASS 시작 주소와 커널 인자가 들어 있는 constant bank 주소 - 완료를 알릴 세마포어 위치까지 포함됨 - 최신 GPU는 커서 변경만 보고 자동으로 움직이지 않아서 도어벨을 울려야 함 - 드라이버가 memory-mapped I/O 레지스터에 채널 토큰을 씀 - GPU의 host engine은 그제야 새 작업이 있는 채널을 확인함 - 이후 GPFIFO를 따라 pushbuffer를 DMA로 읽고, QMD를 compute work distributor에 넘김 ## SM 안에서는 와프 스케줄러가 지연을 숨김 - 예제는 RTX 4090 기준으로 설명됨 - 대상 칩은 128개 SM을 가진 AD102 구성 - 블록 하나는 256스레드, 즉 8와프 - `ptxas`가 예약한 레지스터는 스레드당 16개 - SM 하나에 동시에 올라갈 수 있는 블록 수는 리소스 제약으로 결정됨 - RTX 4090의 SM은 최대 1536 active thread, 즉 48와프를 가짐 - 블록당 256스레드이므로 thread capacity 기준으로 SM당 6블록이 한계 - 레지스터만 보면 더 들어갈 여지가 있지만, 여기선 스레드 수가 더 빡센 제약임 - 와프 스케줄링은 CPU의 out-of-order 실행과 철학이 다름 - CPU는 한 스레드 안에서 의존성을 복잡하게 분석해 병렬성을 끌어냄 - GPU는 많은 와프를 상주시켜놓고, 하나가 메모리 로드로 막히면 다른 와프를 실행함 - 그래서 하드웨어 스케줄링을 단순하게 유지하고, 컴파일러가 명령어에 스케줄링 힌트를 많이 박아 넣음 - SASS 명령어에는 실행 제어 정보가 같이 들어 있음 - 고정 지연 명령에는 stall count가 들어감 - 전역 메모리 로드처럼 지연 시간이 예측 어려운 명령에는 scoreboard barrier를 사용함 - 예제에선 두 `LDG.E`가 같은 barrier를 set하고, 그 값을 쓰는 `FADD`가 barrier를 wait함 - 로드가 돌아오기 전까지 해당 와프는 eligible하지 않으니 스케줄러가 다른 와프를 고름 ## 메모리 대역폭이 이 커널의 진짜 병목 - 각 와프의 32개 스레드는 연속된 float 배열 원소를 읽음 - 스레드당 4바이트이므로 와프 전체는 128바이트를 요청함 - 로드/스토어 유닛은 이를 32바이트 sector 요청 4개로 합침 - 접근 패턴이 연속적이라 coalescing이 아주 잘 되는 케이스임 - 데이터는 L1, L2, VRAM 계층을 따라 이동함 - 먼저 SM의 L1 데이터 캐시를 확인함 - miss가 나면 128개 SM과 연결된 72MB L2 캐시를 거침 - 거기서도 miss가 나면 GDDR6X VRAM까지 내려감 - 저장도 기본적으로 반대 경로를 따라감 - Nsight Compute 결과가 이 커널의 성격을 딱 보여줌 - 그리드 크기 4096, 블록 크기 256 - 스레드당 레지스터 16개 - active warp는 피크의 82.77퍼센트 - 실제 issue active는 5.17퍼센트 - DRAM throughput은 피크의 79.65퍼센트 - GPU 실행 시간은 10.78마이크로초 > [!IMPORTANT] > 와프는 많이 떠 있는데 명령 발행률은 5.17퍼센트밖에 안 됨. 연산기가 놀아서가 아니라, 대부분의 시간이 메모리 로드를 기다리는 데 쓰인다는 뜻임. - 출력 버퍼 `c`는 곧바로 VRAM까지 내려가지 않았음 - 글에 따르면 4MB 출력은 72MB L2에 들어가 있었고, 이후 device-to-host 복사가 읽어갈 때 L2에서 바로 PCIe로 넘어감 - Nsight Compute 기준 DRAM read는 8.4MB 수준이고 write는 사실상 없었다고 설명함 - 그래서 10.78마이크로초 동안 읽기 대역폭은 대략 780GB/s, 피크의 약 4/5 수준으로 해석됨 ## 다시 CPU로 돌아오기 - 커널 launch 자체는 비동기라서 도어벨을 울린 뒤 CPU는 바로 돌아옴 - GPU의 모든 블록이 끝나면 QMD에 들어 있던 completion semaphore를 기록함 - 기본 스트림의 `cudaMemcpy(c, dc, ...)`는 커널 뒤에 있으므로 이 완료를 기다림 - 복사 엔진이 GPU 메모리의 결과를 호스트 메모리로 DMA 전송함 - 복사가 끝나면 CPU 입장에선 그냥 일반 메모리에 결과가 들어온 상태가 됨 - `cudaMemcpy`가 반환됨 - `printf`가 `c[0]`과 `c[n-1]`을 읽어 문자열로 포맷함 - stdout에 write syscall을 날리고, 우리는 `2.000000`을 보게 됨 - 글의 결론은 단순하지만 강함. CUDA 커널 호출은 문법상 함수 호출처럼 보이지만 실제로는 컴파일러, 런타임, 드라이버, PCIe, GPU 큐, SM 스케줄러, 캐시 계층이 전부 참여하는 분산 시스템에 가까움 - 성능 튜닝할 때 왜 occupancy, register count, memory coalescing, stream, copy engine 같은 단어가 계속 나오는지 감이 잡힘 - 추상화는 편하지만, 병목은 추상화 아래에서 생긴다는 걸 아주 친절하게 보여주는 글임 --- ## 기술 맥락 - 이 글에서 가장 중요한 선택은 CUDA가 커널 호출을 CPU 함수 호출처럼 처리하지 않는다는 점이에요. GPU는 PCIe 너머에 있는 별도 장치라서, 드라이버가 QMD와 pushbuffer를 만들고 GPU가 그 명령 큐를 읽는 방식으로 일을 시작해요. 왜 이런 구조냐면 CPU와 GPU가 같은 실행 모델이나 스택을 공유하지 않기 때문이에요. - PTX와 SASS를 나눠 쓰는 이유도 실용적이에요. PTX는 앞으로 나올 GPU에서도 다시 컴파일할 수 있는 중간 표현이고, SASS는 지금 타깃 GPU에서 바로 실행되는 기계어예요. 그래서 실행 파일 안에는 RTX 4090용 SASS와 fallback용 PTX가 같이 들어가고, 드라이버는 필요한 경우 PTX를 JIT해서 새 SASS를 만들 수 있어요. - 와프 스케줄링은 GPU가 메모리 지연을 다루는 방식의 핵심이에요. CPU처럼 한 스레드를 똑똑하게 재정렬하기보다, 많은 와프를 올려놓고 막힌 와프를 건너뛰는 전략을 써요. 이게 가능한 이유는 CUDA 코드가 처음부터 수많은 스레드를 병렬로 던지는 모델이기 때문이에요. - 이 예제에서 성능을 결정한 건 덧셈 연산이 아니라 메모리 경로예요. float 덧셈은 너무 싸고, 입력 두 개와 출력 하나를 옮기는 비용이 훨씬 커요. 그래서 Nsight Compute에서 issue active는 낮은데 DRAM throughput은 79.65퍼센트까지 올라가는 이상해 보이는 숫자가 나오는 거예요. - 개발자가 가져갈 실전 교훈은 메모리 접근 패턴을 먼저 보라는 거예요. 여기선 연속된 float를 읽어서 coalescing이 잘 됐지만, 실제 커널에서 stride가 꼬이거나 분기 때문에 와프가 갈라지면 같은 하드웨어에서도 대역폭을 훨씬 못 쓸 수 있어요. CUDA 최적화가 자료구조 설계와 같이 가는 이유가 바로 이 지점이에요. ## 핵심 포인트 - nvcc는 쿠다 코드를 호스트 코드, PTX, SASS, fatbin이 포함된 실행 파일로 만든다 - 커널 실행 시 CUDA 런타임은 fatbin 등록 정보와 호스트 스텁을 이용해 드라이버로 넘어간다 - 드라이버는 QMD를 pushbuffer와 GPFIFO에 넣고 도어벨 MMIO 쓰기로 GPU에 작업을 알린다 - RTX 4090에서는 128개 SM, 블록당 256스레드, 스레드당 16레지스터 조건에서 SM당 최대 6블록이 상주한다 - 예제 커널은 산술보다 메모리 대역폭이 병목이며 Nsight Compute 기준 10.78마이크로초, DRAM 피크의 79.65퍼센트를 사용했다 ## 인사이트 쿠다를 ‘커널 호출하면 병렬로 돈다’ 정도로 이해하던 사람에게는 꽤 좋은 해부도다. 특히 컴파일러가 스케줄링 정보를 명령어에 박아 넣고, 드라이버가 QMD와 도어벨로 GPU를 깨우는 부분은 성능 튜닝할 때 추상화 아래를 보는 감각을 준다.