쿠다 커널 한 줄이 실제 지피유에서 실행되기까지 벌어지는 일

커널 하나가 생각보다 긴 여행을 함

• 글의 출발점은 진짜 단순한 쿠다(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 사이에서 어떻게 실제 작업으로 바뀌는지 끝까지 따라가는 데 있음

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 메모리로 올라감

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마이크로초

• 출력 버퍼 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 최적화가 자료구조 설계와 같이 가는 이유가 바로 이 지점이에요.