CUDA 커널 한 줄이 실행되기까지: GPU 내부에서 실제로 벌어지는 모든 과정

단순한 결과 뒤에 숨겨진 복잡한 여정

이 글은 CUDA 커널을 실행했을 때 GPU 내부에서 실제로 어떤 일이 벌어지는지를 단계별로 정리한 기술 블로그입니다.
겉보기에는 단순한 벡터 덧셈 CUDA 프로그램이 2.000000이라는 결과를 출력하는 과정처럼 보이지만, 그 뒤에서는 컴파일 파이프라인, CUDA 런타임과 드라이버, GPU 명령 큐, 워프 스케줄링, 메모리 계층, 완료 신호 처리까지 매우 복잡한 흐름이 작동합니다.
본 글에서는 예제 vadd 커널을 기준으로, CUDA 커널이 컴파일되는 순간부터 GPU에서 실행되고 결과가 다시 CPU로 돌아오기까지의 전체 내부 동작을 기술적으로 풀어 설명합니다.

반응형

---

예제 CUDA 커널과 관찰 범위

예제로 사용된 CUDA 커널은 두 개의 float 배열을 더해 결과를 세 번째 배열에 저장하는 가장 기본적인 벡터 덧셈 커널입니다.

• 처리 데이터 수:
  n = 1 << 20 → 1,048,576개의 float
• 커널 실행 구성:
  vadd<<<4096, 256>>>(da, db, dc, n)
  → 총 4096 × 256 = 1,048,576개의 스레드
• 대상 GPU: GeForce RTX 4090
• 컴파일 옵션:
  nvcc -arch=sm_89

실행 결과는 단순합니다.

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

하지만 이 한 줄의 출력 뒤에는 수천만 개의 CPU 명령, 약 900개의 ioctl 호출, GPU와 통신하기 위한 메모리 매핑된 doorbell 레지스터까지 관여합니다.

---

nvcc가 실행 파일을 만드는 전체 컴파일 파이프라인

CUDA 컴파일러 nvcc는 하나의 소스 코드를 호스트 코드와 디바이스 코드로 분리해 처리합니다.

nvcc --keep 옵션을 사용하면 다음과 같은 중간 산출물을 확인할 수 있습니다.

• vadd.ptx
  → cicc가 생성한 디바이스 코드의 PTX
• vadd.sm_89.cubin
  → ptxas가 PTX를 실제 GPU 명령어(SASS)로 변환한 결과
• vadd.fatbin
  → cubin과 PTX를 함께 묶은 바이너리
• vadd.cudafe1.stub.c
  → 호스트 쪽 launch stub 코드
• vadd.o
  → 최종 호스트 오브젝트 파일

PTX와 SASS의 차이

• PTX
  • 가상 ISA
  • 타입이 있는 무한한 가상 레지스터 사용
  • 실제 하드웨어 제약을 직접 반영하지 않음
• SASS
  • GPU 아키텍처에 종속적인 실제 명령어
  • PTX보다 훨씬 압축된 형태

RTX 4090에서는 SASS가 직접 실행되며, PTX는 다른 아키텍처에서 JIT 컴파일을 위한 fallback 용도로 포함됩니다.

---

호스트 코드가 커널 launch를 준비하는 방식

vadd<<<4096, 256>>>(da, db, dc, n) 구문은 컴파일 과정에서 host launch stub 함수 호출로 변환됩니다.

launch stub의 핵심 역할

• 커널 인자(da, db, dc, n)를 argument buffer에 정렬해 패킹
• CUDA 런타임 함수 __cudaLaunch 호출
• 커널을 식별하기 위한 호스트 함수 포인터 전달

이 함수 포인터는 실제로 CPU에서 실행될 함수 주소가 아니라, 런타임 내부 등록 테이블을 조회하기 위한 key 역할을 합니다.

첫 GPU 호출 시 발생하는 일

• CUDA 런타임이 동적으로 libcuda.so.1을 로드
• GPU context 생성
• CPU와 GPU 사이의 통신 채널 초기화

CUDA 12.2부터는 lazy module loading이 기본값이어서, 실제 커널이 처음 실행될 때까지 SASS 업로드가 지연됩니다.

---

GPU에 작업을 전달하는 명령 큐 구조

GPU는 CPU처럼 함수 호출로 직접 진입하지 않습니다.
대신 호스트 메모리에 있는 드라이버 커맨드 스트림을 PCIe를 통해 읽습니다.

핵심 구성 요소

• pushbuffer
  • 드라이버가 GPU 명령(method)을 기록하는 메모리 영역
• GPFIFO
  • pushbuffer의 위치와 길이를 가리키는 링 버퍼

GPU와 드라이버는 다음 두 커서를 사용해 작업 상태를 추적합니다.

• GP_PUT: 드라이버가 어디까지 명령을 썼는지
• GP_GET: GPU가 어디까지 명령을 소비했는지

doorbell의 역할

현대 GPU에서는 host engine이 커서를 지속적으로 감시하지 않습니다.
대신 드라이버가 doorbell MMIO 레지스터에 값을 쓰면, GPU가 즉시 새로운 작업이 있음을 인지합니다.

---

QMD: 커널 실행에 필요한 모든 정보 묶음

커널 launch는 QMD(Queue Meta Data)를 GPU에 전달하는 것으로 시작됩니다.

QMD에는 다음 정보가 포함됩니다.

• grid 크기와 block 크기 (4096, 256)
• thread당 레지스터 수
• shared memory 요구량
• 커널 SASS 시작 주소
• 커널 인자가 저장된 constant bank 주소
• 완료 시 신호를 보낼 semaphore 위치

cuLaunchKernel은 doorbell을 울리는 순간 반환되며, GPU 작업은 비동기로 진행됩니다.

---

SM, 워프, 그리고 점유율

RTX 4090에는 총 128개의 SM이 있습니다.

리소스 계산 결과

• block당 스레드 수: 256 → 8 warps
• thread당 레지스터 수: 16
• SM당 최대 active threads: 1,536

이 조건으로 인해:

• SM당 최대 6개 block
• 총 48개 warp가 resident 상태

각 SM은 4개의 sub-partition으로 나뉘며, 각 warp scheduler는 최대 12개의 active warp를 관리합니다.

---

워프 스케줄링과 stall 처리 방식

GPU는 CPU처럼 단일 스레드에서 복잡한 out-of-order 실행을 하지 않습니다.
대신 여러 warp를 동시에 유지하면서 stall이 발생하면 다른 warp로 전환해 latency를 숨깁니다.

• 고정 latency 명령어: static stall count 사용
• 가변 latency 명령어(load 등): scoreboard barrier 사용

예제 커널에서는 두 번의 global load가 같은 barrier를 사용하며, load가 완료될 때까지 해당 warp는 ineligible 상태가 됩니다.

---

메모리 접근 경로와 성능 특성

메모리 접근 흐름

1. warp가 LDG.E 실행
2. 32개 thread가 연속된 128바이트 요청 생성
3. SM에서 coalescing 수행
4. L1 → L2 → GDDR6X VRAM 순으로 접근

Nsight Compute 측정 결과

• 실행 시간: 10.78μs
• DRAM 대역폭 사용률: 79.65%
• warp issue 비율: 5.17%

이 커널은 산술 집약도가 매우 낮은 전형적인 memory-bound 커널입니다.

• float add 1회당
  → load 8바이트 + store 4바이트 = 총 12바이트 전송

---

GPU 결과가 CPU로 돌아오는 과정

커널 실행은 비동기이기 때문에, GPU는 완료 사실을 별도로 알려야 합니다.

• 모든 block이 종료되면 GPU가 QMD에 지정된 semaphore를 post
• 이후 실행되는 cudaMemcpy는 이 semaphore를 기다림
• 결과 데이터는 L2 캐시에 존재하므로 DRAM 왕복 없이 PCIe를 통해 host memory로 복사
• CPU가 printf로 결과 출력

---

내부 동작을 직접 관찰하는 방법

CUDA 드라이버는 closed-source이기 때문에 모든 내부를 직접 확인할 수는 없습니다.
하지만 다음 방법으로 일부 동작을 관찰할 수 있습니다.

• nvcc --keep로 생성된 stub 코드 분석
• strace로 ioctl 호출 추적
• LD_PRELOAD shim으로 /dev/nvidia* mmap 영역 기록
• pushbuffer에서 QMD method burst 덤프

이를 통해 실제로 GPU에 전달되는 QMD 구조와 launch 파라미터를 확인할 수 있습니다.

---

728x90

이 글에서 살펴본 것처럼, CUDA 커널 한 줄의 실행은 단순한 함수 호출이 아닙니다.
컴파일 단계부터 런타임 등록, 드라이버 명령 큐, QMD 생성, 워프 스케줄링, 메모리 계층을 거쳐 결과가 다시 CPU로 돌아오기까지 수많은 단계가 연결되어 있습니다.

이 구조를 이해하면 다음과 같은 관점을 얻을 수 있습니다.

• 커널 성능 병목이 어디에서 발생하는지 더 명확히 볼 수 있음
• launch 오버헤드와 메모리 바운드 특성을 구분해 최적화 가능
• Nsight Compute 수치를 단순 지표가 아닌 내부 동작과 연결해 해석 가능

CUDA 성능 최적화는 단순히 커널 코드를 고치는 문제가 아니라, GPU 전체 실행 모델을 이해하는 문제라는 점을 이 예제가 잘 보여줍니다.

300x250

https://fergusfinn.com/blog/what-happens-when-you-run-a-gpu-kernel/

What happens when you run a CUDA kernel