GPU

같은 계산을 아주 많은 데이터에 한꺼번에 거는 장치입니다. 계산을 해내는 자리를 잘게 여러 개 두고 나란히 돌립니다. 혼자서는 일을 시작하지 못합니다. 중앙처리장치가 일감과 데이터를 넘겨줘야 움직입니다.

상세

답안지 수천 장을 같은 채점표로 매긴다고 해 보겠습니다. 한 사람이 처음부터 끝까지 매기는 대신 채점자를 수백 명 두고 각자 한 장씩 같은 표로 매기면 한 번에 끝납니다.

GPU(Graphics Processing Unit, 그래픽 처리 장치)는 같은 계산을 서로 다른 데이터에 한꺼번에 거는 프로세서입니다. 계산을 해내는 자리를 아주 많이 두고, 그 자리 하나하나는 단순합니다. 여러 자리가 같은 명령을 나란히 실행합니다.

이 장치는 혼자 돌지 않습니다. 계산을 맡기는 쪽과 맡는 쪽이 나뉩니다. 벤더에 매이지 않는 계산 API(Application Programming Interface) 명세인 OpenCL 명세가 이 갈라짐에 이름을 붙여 뒀습니다. 플랫폼은 호스트 하나와 거기에 붙은 하나 이상의 디바이스로 이뤄집니다. 디바이스는 하나 이상의 컴퓨트 유닛으로 나뉘고, 컴퓨트 유닛은 다시 하나 이상의 프로세싱 엘리먼트로 나뉩니다. 계산은 프로세싱 엘리먼트 안에서 일어납니다. 프로세싱 엘리먼트는 가상의 스칼라 프로세서라고 정의돼 있습니다. 컴퓨트 유닛은 프로세싱 엘리먼트들과 로컬 메모리로 이뤄지고, 프로세싱 엘리먼트들이 쓸 수 있는 전용 텍스처 필터 유닛을 함께 두기도 합니다.

flowchart TD
    H["호스트"] -->|"커널을 명령으로 제출"| D["디바이스"]
    D --> CU["컴퓨트 유닛"]
    CU --> PE1["프로세싱 엘리먼트"]
    CU --> PE2["프로세싱 엘리먼트"]

프로그램은 호스트 코드와 디바이스 커널 코드 두 벌로 만들어집니다. 호스트 코드는 호스트 프로세서 위에서 돌면서 커널 코드를 명령으로 디바이스에 제출합니다. 디바이스는 그 명령을 자기 프로세싱 엘리먼트 위에서 실행합니다. 일감은 잘게 쪼개져 묶음 단위로 배정됩니다. 묶음 하나인 워크그룹은 하나의 컴퓨트 유닛 위에서 실행되고, 낱개인 워크아이템은 하나 이상의 프로세싱 엘리먼트 위에서 실행될 수 있습니다. 계산을 디바이스의 처리 자리에 어떻게 대응시킬지는 디바이스가 상당한 재량을 가집니다.

한 컴퓨트 유닛 안의 프로세싱 엘리먼트들이 같은 문장 순서를 나란히 실행하면 그 제어 흐름을 수렴됐다고 부릅니다. 프로세싱 엘리먼트마다 흐름이 달라지면 발산했다고 부릅니다. 명세는 하나의 명령 흐름을 여러 프로세싱 엘리먼트에 걸쳐 실행하도록 최적화된 하드웨어가 수렴된 제어 흐름에 잘 맞는다고 적습니다. 커널은 언제나 수렴된 제어 흐름으로 시작하지만, 커널 안의 분기문 때문에 하나의 커널 안에서 수렴과 발산이 둘 다 나타날 수 있습니다.

배경

화면은 점 수백만 개로 이뤄집니다. 그 점 하나하나마다 무슨 색을 칠할지 같은 방식으로 계산해야 합니다. 계산 자리가 하나뿐이면 한 장을 채우는 데 같은 계산을 수백만 번 되풀이해야 합니다. 화면은 1초에도 여러 장 바뀝니다. 점마다 다시 시작되는 이 되풀이가 감당이 안 됐습니다.

바꿔야 할 것은 계산 하나의 속도가 아니라 처리하는 방식이었습니다. 점마다의 계산은 서로 남의 결과를 기다리지 않습니다. 그러면 나눠서 한꺼번에 걸 수 있습니다. 계산 자리를 아주 많이 두고 같은 명령을 서로 다른 데이터에 나란히 거는 전용 장치가 필요했습니다.

그 장치가 맡은 일이 그래픽 처리였고, 계산을 도맡던 중앙처리장치와 짝을 이루는 이름을 받았습니다. 처음에는 정해진 그리기 절차만 수행했습니다. 나중에 그 절차의 일부를 프로그램으로 바꿔 넣을 수 있게 되면서 그리는 일이 아닌 계산까지 같은 장치에 맡기게 됐습니다. 그래픽과 무관한 계산을 이 장치에 넘기기 위한 별도의 계약이 따로 서게 된 것이 그 결과입니다.

예시

Vulkan 의 vkCmdDispatch

그래픽·계산 API 인 Vulkan 은 계산 일감을 디바이스에 밀어 넣는 명령을 이렇게 정의합니다.

void vkCmdDispatch(
    VkCommandBuffer  commandBuffer,
    uint32_t         groupCountX,
    uint32_t         groupCountY,
    uint32_t         groupCountZ);

commandBuffer 는 이 명령을 기록해 둘 커맨드 버퍼입니다. groupCountX · groupCountY · groupCountZ 는 각 차원으로 디스패치할 로컬 워크그룹의 개수입니다. 이 명령이 실행되면 groupCountX × groupCountY × groupCountZ 개의 로컬 워크그룹으로 이뤄진 글로벌 워크그룹이 조립됩니다. 무엇을 계산할지가 아니라 묶음을 몇 개 만들지를 숫자로 적어 넘기는 자리라는 것이 시그니처에 그대로 드러납니다.

CUDA 의 커널 실행

범용 계산 쪽에서는 언어 확장이 같은 일을 받습니다. CUDA C++ 는 커널이라 부르는 함수를 정의할 수 있게 해 줍니다. 커널은 호출되면 일반 C++ 함수처럼 한 번만 실행되는 것이 아니라, 서로 다른 N개의 스레드가 N번 나란히 실행합니다.

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    ...
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);
    ...
}

<<<1, N>>> 이 스레드를 몇 개 띄울지를 정합니다. 내장 변수 threadIdx 는 스레드마다 값이 다르므로 i 도 스레드마다 다릅니다. VecAdd() 를 실행하는 N개의 스레드는 각각 한 쌍의 덧셈만 수행합니다. 크기 N인 벡터 두 개를 더해 결과를 벡터 C 에 담는 일이 이렇게 끝납니다.

리눅스 DRM 의 장치 노드

운영체제는 GPU 를 장치 노드로 내보냅니다. 리눅스 커널의 DRM(Direct Rendering Manager) 코어는 여러 인터페이스를 응용 프로그램에 내주는데, 보통은 libdrm 의 대응 래퍼 함수를 통해 씁니다. 어느 장치를 여느냐에 따라 사용자 공간이 할 수 있는 조작이 달라집니다. 주로 ioctl 이 갈립니다.

card<num>
controlD<num>
renderD<num>

프라이머리 노드는 언제나 만들어지고 이름은 card<num> 입니다. 지금은 쓰이지 않는 컨트롤 노드 controlD<num> 도 함께 만들어집니다. 인증 단계를 거치지 않고 GPU 접근만 내주려고 나중에 렌더 노드가 도입됐습니다. 렌더 노드는 렌더 클라이언트만 상대하므로 모드 설정이나 특권 ioctl 을 여기서는 실행할 수 없습니다. 이름은 renderD<num> 이고 디바이스마다 하나씩 만들어집니다.

PyTorch 가 텐서를 올리는 자리

프레임워크는 이 장치를 인자 하나로 지목합니다. PyTorch 는 텐서를 만들 때 device 를 붙여 어디에 자리 잡을지 정합니다.

x_cpu = torch.empty(2)
x_gpu = torch.empty(2, device=cuda)

같은 torch.empty() 호출인데 아래 줄만 GPU 쪽에 만들어집니다. 계산할 데이터를 이쪽으로 옮기는 일이 별도의 동작이라는 것이 이 한 줄 차이에 드러납니다.

경계

메인보드에 붙어 나오는 내장 그래픽도 GPU 인가. 맞습니다. Vulkan 명세는 물리 디바이스의 종류를 열거하면서 VK_PHYSICAL_DEVICE_TYPE_INTEGRATED_GPU 를 호스트에 내장돼 있거나 호스트와 밀접하게 결합된 디바이스라고 정의합니다. 인터링크를 통해 호스트에 붙는 별도 프로세서는 VK_PHYSICAL_DEVICE_TYPE_DISCRETE_GPU 로 따로 두고, 가상화 환경의 가상 노드는 VK_PHYSICAL_DEVICE_TYPE_VIRTUAL_GPU 로 둡니다. 호스트에 붙는 방식이 갈릴 뿐 셋 다 GPU 라는 이름을 달고 있습니다. 명세는 각 정의에 보통이라는 한정어를 달아 뒀고, 이 종류 값이 정보 제공 목적일 뿐 시스템 동작에 직접 영향을 주지는 않는다고 못 박습니다.

관련 항목

GPU를 이루는 실행 단위

컴퓨트 유닛 · 프로세싱 엘리먼트 · 워크그룹 · 워크아이템 · 워프 · 스레드 블록 · 스레드 · 커널

GPU가 그림을 그리며 거치는 처리 단계

셰이더 · 파이프라인 · 래스터화 · 텍스처 · 프레임버퍼

GPU가 갖는 실행 성질

SIMD(Single Instruction Multiple Data) · SIMT(Single Instruction Multiple Threads) · 병렬 처리 · 처리량 · 지연

GPU 메모리에 얹히거나 오가는 하드웨어와 데이터

VRAM(Video RAM) · 메모리 대역폭 · 호스트-디바이스 전송 · PCIe(Peripheral Component Interconnect Express) · 텐서

운영체제가 GPU를 불러 쓰는 통로

운영체제 · 디바이스 드라이버 · DRM · libdrm · 노드 · 커맨드 버퍼 · 인증 · 물리 디바이스 · ioctl

GPU에 계산을 넘기는 API·언어·프레임워크

API · CUDA · OpenCL · Vulkan · OpenGL · WebGPU · C++ · GLSL(OpenGL Shading Language) · HLSL(High-Level Shading Language) · WGSL(WebGPU Shading Language) · PyTorch

GPU에 일감을 넘기는 주체

호스트 · CPU(Central Processing Unit)

GPU를 대신할 수 있는 다른 가속기

NPU(Neural Processing Unit) · TPU(Tensor Processing Unit) · FPGA(Field-Programmable Gate Array) · ASIC(Application-Specific Integrated Circuit)

GPU가 쓰이는 분야

그래픽스 · 렌더링 · 게임 개발 · 머신러닝 · 학습 · 추론 · 배치 처리