컴퓨트 파이프라인

컴퓨트 파이프라인은 그래픽 카드에 계산을 시키기 전에 실행 준비를 미리 끝내 둡니다. 돌릴 계산 프로그램과 그 프로그램이 데이터를 받을 칸을 한 덩어리로 묶어 둡니다. 실행할 때는 이 덩어리를 골라 몇 번 돌릴지만 정합니다.

상세

이 절은 컴퓨트 파이프라인 안에 무엇이 묶이는지, 왜 실행 전에 미리 만드는지, 한 번의 실행이 어떤 순서로 도는지를 봅니다. 값 천 개를 저마다 두 배로 만드는 계산 하나를 줄곧 예로 씁니다.

GPU(Graphics Processing Unit, 그래픽 처리 장치)는 같은 계산을 동시에 아주 많이 돌리도록 만든 칩입니다. 화면의 점 수백만 개의 색을 한꺼번에 정하려고 만들었습니다. 이 칩을 얹은 판을 흔히 그래픽 카드라고 부릅니다.

GPU 에서 도는 작은 프로그램을 셰이더라고 부릅니다. 그중 그림과 상관없이 메모리의 값을 읽고 계산해 다시 쓰는 셰이더가 컴퓨트 셰이더입니다. 값 천 개를 두 배로 만드는 프로그램이 컴퓨트 셰이더입니다.

컴퓨트 파이프라인은 컴퓨트 셰이더 하나를 GPU 가 곧바로 돌릴 수 있게 준비해 둔 객체입니다. 셰이더 코드만 들고서는 칩이 돌리지 못합니다. 칩의 기계어로 옮겨야 합니다. 데이터를 어디서 받을지도 정해야 합니다. 그 준비를 마친 결과가 컴퓨트 파이프라인입니다.

식당 주방은 문을 열기 전에 조리대마다 칼과 냄비를 꺼내 둡니다. 재료는 주문이 들어올 때마다 올라옵니다. 준비는 하루에 한 번입니다. 주문은 수백 번 들어옵니다.

컴퓨트 파이프라인은 이 조리대 준비에 해당합니다. 계산할 값은 주문의 재료처럼 실행할 때마다 따로 들어옵니다.

파이프라인에 묶이는 세 가지

컴퓨트 파이프라인을 만들 때 넘기는 것은 셋입니다. 셰이더 코드, 그 코드에서 처음 부를 함수의 이름, 데이터를 받을 칸의 배치입니다.

셰이더 코드는 사람이 읽는 셰이더 언어로 적혀 있습니다. 파이프라인을 만드는 순간 GPU 가 알아듣는 기계어로 옮겨집니다. 이렇게 옮기는 일이 컴파일입니다.

한 코드 파일에 함수가 여럿 있을 수 있습니다. 그래서 어느 함수부터 돌릴지를 함께 적습니다. 이 함수 이름을 진입점이라고 부릅니다. 두 배 계산에서는 main 이라는 함수 하나가 진입점입니다.

셰이더가 읽고 쓰는 데이터를 자원이라고 부릅니다. 값을 한 줄로 늘어놓은 메모리 덩어리인 버퍼가 가장 흔한 자원입니다. 두 배 계산은 값 천 개가 든 버퍼 하나를 읽고 거기에 다시 씁니다.

자원 배치는 어떤 종류의 자원이 몇 번 칸으로 들어오는지를 적은 표입니다. 셰이더는 자원을 이름이 아니라 칸 번호로 찾습니다. 두 배 계산의 자원 배치는 「0번 칸에 읽고 쓸 수 있는 버퍼 하나」가 전부입니다.

프로그램이 GPU 를 다룰 때 부르는 함수 모음을 API(Application Programming Interface, 응용 프로그램 인터페이스)라고 합니다. API 문서와 코드는 자원 배치를 파이프라인 레이아웃이라고 부릅니다. 칸 번호는 셰이더 코드에 binding 으로 적혀서 바인딩 번호라고 부릅니다. 이 문서는 둘을 끝까지 자원 배치와 칸이라고 부릅니다.

파이프라인에 안 들어가는 것

자원 배치에는 칸의 종류만 적힙니다. 값 천 개가 든 실제 버퍼는 파이프라인에 들어가지 않습니다. 실제 버퍼는 실행할 때마다 칸에 끼웁니다.

이렇게 갈라 둔 덕분에 파이프라인 하나로 여러 데이터를 처리합니다. 오늘 받은 값 천 개도, 내일 받은 값 천 개도 같은 파이프라인에 버퍼만 바꿔 끼워 돌립니다.

몇 번 돌릴지도 파이프라인 밖에서 정합니다. 값이 천 개면 계산을 천 번쯤, 백만 개면 백만 번쯤 돌립니다. 실행 명령에는 이 횟수 대신 계산 몇 번을 한 묶음으로 친 묶음의 개수를 싣습니다. 묶음은 아래 「한 번의 실행이 도는 순서」에서 풉니다.

flowchart TD
    subgraph P["컴퓨트 파이프라인 · 만들 때 굳는다"]
        S["셰이더 코드 · 기계어로 컴파일해 둠"]
        E["진입점 · main"]
        L["자원 배치 · 0번 칸은 버퍼"]
    end
    subgraph R["실행할 때마다 정한다"]
        B["실제 버퍼 · 값 천 개"]
        N["몇 번 돌릴지 · 묶음 개수"]
    end
    P --> X["실행 명령 한 번"]
    B -->|"0번 칸에 끼운다"| X
    N --> X

위쪽 상자에 든 것은 한 번 만들면 안 바뀝니다. 아래쪽 상자에 든 것은 실행 명령마다 새로 고릅니다. 두 상자가 실행 명령 하나에서 만납니다.

미리 만들어 굳히는 까닭

셰이더 코드를 컴파일하는 일은 짧은 계산이 아닙니다. 코드가 길면 사람이 느낄 만큼 시간이 걸립니다. 이 일을 실행 명령마다 하면 계산보다 준비가 오래 걸립니다.

컴파일은 드라이버가 맡습니다. 드라이버는 운영체제와 GPU 사이에서 프로그램의 명령을 칩이 알아듣는 꼴로 바꿔 주는 소프트웨어입니다.

드라이버는 컴파일 말고 검사도 합니다. 셰이더가 찾는 칸과 넘겨받은 자원 배치가 서로 맞는지 확인해야 합니다. 파이프라인으로 묶어 두면 이 확인도 만들 때 한 번으로 끝납니다.

그래서 파이프라인은 만든 뒤에 못 고칩니다. 고칠 수 있게 두면 드라이버가 실행 때마다 다시 검사하고 다시 컴파일해야 합니다. 셰이더가 하나라도 다르거나 자원 배치가 다르면 파이프라인을 하나 더 만듭니다.

대가는 시작 시간과 개수입니다. 프로그램이 뜰 때 파이프라인을 몰아서 만들므로 첫 계산이 늦어집니다. 셰이더와 자원 배치의 조합이 많으면 파이프라인 수도 그만큼 늘어납니다.

GPU 를 다루는 API 에 따라 이 시작 시간을 줄이는 장치를 둡니다. Vulkan 의 VkPipelineCache 는 한 번 컴파일한 결과를 파일로 남겨 두었다가 다음 실행 때 다시 씁니다. 이런 장치를 파이프라인 캐시라고 부릅니다.

한 번의 실행이 도는 순서

컴퓨트 셰이더는 한 번에 한 개씩 돌지 않습니다. 같은 코드가 수천 번 동시에 돕니다. 그 한 번 한 번을 셰이더 호출이라고 부릅니다. 호출마다 자기 번호를 받아 그 번호의 값 하나를 맡습니다.

호출은 정해진 수만큼 묶여 함께 돕니다. 이 묶음을 작업 그룹이라고 부릅니다. 한 묶음에 호출이 몇 개 드는지는 셰이더 코드에 적혀 있습니다. 그래서 그룹 크기는 파이프라인을 만들 때 함께 굳습니다.

실행 명령은 작업 그룹을 몇 개 돌릴지만 받습니다. 이 명령을 디스패치라고 부릅니다. 앞에서 말한 「몇 번 돌릴지」가 이 그룹 개수로 넘어갑니다. 전체 호출 수는 그룹 개수에 그룹 크기를 곱한 값입니다.

그룹 크기가 64 이고 값이 천 개면 계산은 이렇게 나옵니다.

Math.ceil(1000 / 64)  // 16 그룹
16 * 64               // 1024 호출
1024 - 1000           // 24 호출이 논다

그룹 개수는 올림으로 구합니다. 내림하면 끝의 값 몇 개를 맡을 호출이 모자랍니다. 대신 남는 호출 24 개가 생깁니다.

호출 번호는 그룹이 바뀌어도 0 으로 돌아가지 않고 이어집니다. 그룹 0 이 0~63 번을, 그룹 1 이 64~127 번을 받습니다. 마지막 그룹 15 는 960~1023 번을 받습니다. 그래서 남는 호출 24 개는 전부 마지막 그룹의 끝에 몰립니다.

flowchart TD
    subgraph D["디스패치 한 번 · 작업 그룹 16 개"]
        G0["작업 그룹 0 · 호출 0~63 → 값 0~63"]
        G1["작업 그룹 1 · 호출 64~127 → 값 64~127"]
        GM["작업 그룹 2~14 · 13 개 접음"]
        subgraph G15["작업 그룹 15"]
            W["호출 960~999 · 값을 맡음"]
            I["호출 1000~1023 · 24 개가 논다"]
        end
        G0 ~~~ G1 ~~~ GM ~~~ W ~~~ I
    end

그룹마다 호출이 64 개인 것은 파이프라인을 만들 때 굳었습니다. 그룹이 16 개인 것은 디스패치가 정했습니다. 남는 호출들은 자기 번호가 천 이상인 것을 보고 아무 일도 하지 않습니다.

명령은 하나씩 바로 GPU 로 가지 않습니다. 커맨드 버퍼라는 목록에 차례로 적혔다가 한 번에 넘어갑니다. 한 번에 넘기면 GPU 에 명령을 건네는 횟수가 줄어듭니다.

넘어간 커맨드 버퍼는 큐에 들어갑니다. 큐는 GPU 가 받은 명령을 들어온 차례대로 꺼내 처리하는 통로입니다. 커맨드 버퍼를 큐에 넣는 것은 프로그램이 합니다.

sequenceDiagram
    participant 프로그램
    participant 드라이버
    participant GPU
    프로그램->>드라이버: 셰이더와 자원 배치로 파이프라인을 만든다
    드라이버-->>프로그램: 컴파일과 검사를 마친 파이프라인
    Note over 프로그램,드라이버: 여기까지는 시작할 때 한 번
    Note over 프로그램: 파이프라인 고르기 · 버퍼 끼우기 · 디스패치를 커맨드 버퍼에 적는다
    프로그램->>드라이버: 커맨드 버퍼를 큐에 넣는다
    드라이버->>GPU: 큐의 명령을 차례로 넘긴다
    Note over GPU: 호출 1024 개가 저마다 값 하나를 두 배로
    GPU-->>프로그램: 끝났다고 알린다

그림의 위 두 줄이 준비입니다. 나머지는 실행 때마다 되풀이됩니다. 값이 바뀌면 버퍼만 바꿔 끼우고 다시 디스패치합니다. 파이프라인을 다시 만드는 줄은 되풀이에 끼지 않습니다.

WebGPU 코드로 본 모습

이 소절은 위 순서를 실제 코드로 적어 봅니다. 브라우저에서 GPU 를 쓰는 API 인 WebGPU로 씁니다. 줄 수가 적어 각 부분이 어디에 대응하는지 잘 보입니다.

먼저 셰이더 코드입니다. WGSL(WebGPU Shading Language, WebGPU 셰이더 언어)로 적습니다. 칸 번호와 그룹 크기와 진입점이 전부 이 코드 안에 들어 있습니다.

@group(0) @binding(0)   // 바인드 그룹 0 · 칸 0
var<storage, read_write> v: array<f32>;

@compute @workgroup_size(64)   // 그룹 크기
fn main(@builtin(global_invocation_id) id: vec3u) {
  if (id.x < 1000u) {          // 남는 24 개
    v[id.x] = v[id.x] * 2.0;
  }
}

WebGPU 는 칸을 몇 개씩 모아 한 번에 끼웁니다. 이 모음을 바인드 그룹이라고 부릅니다. @group(0) 은 0번 바인드 그룹을, @binding(0) 은 그 안의 0번 칸을 가리킵니다. 이 group 은 호출을 묶는 작업 그룹과 상관이 없습니다.

global_invocation_id 는 호출마다 받는 자기 번호입니다. if 줄이 번호가 천 이상인 호출을 걸러 냅니다. 앞에서 본 남는 호출 24 개가 여기서 아무 일도 안 하고 끝납니다.

다음은 파이프라인을 만드는 코드입니다. 프로그램이 시작할 때 한 번 부릅니다.

const pipeline = device.createComputePipeline({
  layout: "auto",
  compute: { module, entryPoint: "main" },
});

module 은 위 셰이더 코드를 넘겨 만든 셰이더 모듈입니다. entryPoint 가 진입점입니다. layout: "auto" 는 자원 배치를 셰이더 코드의 칸 번호에서 읽어 내라는 뜻입니다.

마지막은 실행할 때마다 부르는 코드입니다. 실제 버퍼는 바인드 그룹에 담겨 칸에 끼워집니다. 코드의 bindGroup 은 0번 칸에 값 천 개가 든 버퍼를 짝지어 둔 바인드 그룹입니다.

const pass = encoder.beginComputePass();
pass.setPipeline(pipeline);
pass.setBindGroup(0, bindGroup);
pass.dispatchWorkgroups(16);  // 1024 호출
pass.end();
device.queue.submit([encoder.finish()]);

encoder 는 명령을 적는 기록장입니다. beginComputePass() 는 컴퓨트 명령을 적을 구간을 엽니다. 이 구간을 컴퓨트 패스라고 부릅니다. pass.end() 가 그 구간을 닫습니다.

setBindGroup(0, bindGroup) 의 0 은 셰이더의 @group(0) 과 같은 바인드 그룹 번호입니다. finish() 가 적은 것을 커맨드 버퍼로 닫습니다. queue.submit 이 그 커맨드 버퍼를 큐에 넣습니다. 값이 바뀌어도 이 여섯 줄만 다시 돌면 됩니다.

그래픽스 파이프라인과의 차이

그래픽스 파이프라인은 도형을 받아 화면에 칠할 픽셀 색을 내놓는 길입니다. 꼭짓점을 다루는 단계와 도형이 덮는 픽셀을 골라내는 단계와 색을 칠하는 단계가 차례로 이어집니다.

컴퓨트 파이프라인에는 그런 이어짐이 없습니다. 컴퓨트 셰이더 한 단계가 전부입니다. 앞 단계에서 받는 것도 뒤 단계로 넘기는 것도 없습니다. 메모리에서 읽고 메모리에 씁니다.

컴퓨트 쪽이 묶는 설정이 훨씬 적습니다. 묶는 설정이 적으니 설정 조합에 따라 파이프라인이 불어나는 일도 덜합니다.

단계가 하나뿐인데도 파이프라인이라 부르는 까닭은 API 가 둘을 같은 종류의 객체로 다루기 때문입니다. 둘 다 「실행 전에 굳혀 두는 설정 객체」입니다. 고르는 방법도 같습니다. 이 이름에서 파이프라인은 단계가 이어진 줄이라는 뜻보다 그 설정 객체라는 뜻에 가깝습니다.

API 마다 부르는 이름

GPU 를 다루는 API 는 여럿입니다. 컴퓨트 파이프라인을 가리키는 객체와 호출의 이름도 API 마다 다릅니다.

Vulkan 은 그래픽스 파이프라인과 컴퓨트 파이프라인을 같은 VkPipeline 형으로 다룹니다. 파이프라인을 고를 때 어느 쪽인지를 함께 밝힙니다.

OpenGL에는 따로 선 파이프라인 객체가 없습니다. 컴퓨트 셰이더를 붙여 만든 프로그램 객체가 그 역할을 맡습니다. 실행은 glDispatchCompute 로 시킵니다.

쓸 때와 안 쓸 때

백엔드 개발자가 컴퓨트 파이프라인을 손으로 만들 일은 드뭅니다. GPU 로 기계 학습을 돌리는 라이브러리가 그 일을 안쪽에서 대신합니다. 브라우저에서 모델을 돌리는 TensorFlow.js 의 WebGPU 백엔드가 그렇습니다. 이런 라이브러리는 연산마다 파이프라인을 만들어 둡니다. 연산을 부를 때마다 버퍼를 끼워 디스패치합니다.

직접 쓸 만한 때는 같은 계산을 값 수만 개 넘게 되풀이할 때입니다. 사진 보정, 행렬 곱, 입자 수만 개의 위치 갱신이 그렇습니다. 한 파이프라인을 여러 번 돌릴수록 만드는 시간이 묻힙니다.

안 맞는 때도 있습니다. 한 번 돌고 끝나는 작은 계산은 파이프라인을 만드는 시간과 데이터를 GPU 로 옮기는 시간이 계산 시간을 넘습니다. 요청마다 셰이더가 달라지는 일도 안 맞습니다. 파이프라인이 끝없이 늘어납니다. 늘어날 때마다 만드는 시간을 치릅니다.

결과를 화면에 그려야 하면 그래픽스 파이프라인이 맡습니다. 둘을 이어 쓰기도 합니다. 컴퓨트 파이프라인으로 계산해 버퍼에 써 두면 그래픽스 파이프라인이 그 버퍼를 읽어 그립니다.

관련 항목

컴퓨트 파이프라인을 이루는 구성 요소

컴퓨트 셰이더 · 셰이더 모듈 · 진입점 · 파이프라인 레이아웃 · 바인드 그룹 레이아웃 · 디스크립터 세트 레이아웃 · 루트 시그니처 · 특수화 상수 · 작업 그룹

컴퓨트 파이프라인을 실행시키는 명령과 통로

디스패치 · 간접 디스패치 · 커맨드 버퍼 · 커맨드 인코더 · 컴퓨트 패스 · 큐 · 배리어 · 펜스

컴퓨트 파이프라인이 읽고 쓰는 자원

자원 · 버퍼 · 스토리지 버퍼 · 유니폼 버퍼 · 텍스처 · 스토리지 텍스처 · 바인드 그룹 · 디스크립터 세트 · 공유 메모리

컴퓨트 파이프라인을 만들 때 거치는 처리 단계

컴파일 · 셰이더 컴파일 · 기계어 · SPIR-V · 중간 표현 · 파이프라인 캐시 · 디바이스 드라이버

컴퓨트 파이프라인과 같은 종류로 묶이는 파이프라인 객체

그래픽스 파이프라인 · 렌더 파이프라인 · 레이 트레이싱 파이프라인 · 파이프라인 상태 객체

컴퓨트 파이프라인을 제공하는 그래픽스 API

Vulkan · Direct3D · Direct3D 12 · Metal · WebGPU · OpenGL

컴퓨트 셰이더를 적는 셰이더 언어

WGSL · GLSL · HLSL · MSL

컴퓨트 파이프라인 대신 GPU 계산을 맡기는 플랫폼

CUDA · OpenCL · SYCL · GPGPU

컴퓨트 파이프라인으로 푸는 계산 종류

기계 학습 · 행렬 곱 · 이미지 필터 · 입자 시뮬레이션 · 리덕션 · 프리픽스 합

컴퓨트 파이프라인이 속하는 상위 분류

파이프라인 · 셰이더 · GPU · 그래픽스 · API