gpu.cpp — Portable GPU compute using WebGPU

CUDA 가 ML 의 표준이지만, CUDA 가 없는 자리가 많다 — 학생의 MacBook, AR 안경, 로봇 SoC, 웹 브라우저, 일반 사용자의 데스크탑. Austin 의 출발은 그 자리들 위에서 GPU compute 를 어떻게 띄울지의 질문이다.

강의가 답하려는 두 줄 —

1. NVIDIA · AMD · Intel · Apple · Qualcomm 어디든 같은 GPU 코드 한 벌로 도달할 수 있는 spec 이 있는가 — 이 자리에 WebGPU 가 들어온다.
2. 그 spec 을 직접 써보면 보일러플레이트가 살인적인데, 그걸 어떻게 줄일 것인가 — 이 자리에 gpu.cpp 가 들어온다.

강의의 인지적 frame 은 명시적이다. “CUDA 와 같은 자리에 ‘WebGPU + gpu.cpp’ 가 서면 어떤가” — 단, 같은 정점의 perf 를 노리는 게 아니라 접근성과 portability 의 다른 축에 가치를 둔다. 100% 의 perf 의 80% 를 따 내면서 100 배 더 많은 환경에서 돈다는 거래.

그래서 강의 끝에 손에 잡혀야 할 것은 — WebGPU 의 4 추상(adapter, device, buffer, pipeline), WGSL 의 핵심 syntax(@compute, @workgroup_size), gpu.cpp 의 3 함수(createTensor, createKernel, dispatchKernel), 그리고 이 stack 의 한계가 어디서 시작되는지(Tensor Core 미접근, FP16 분기) 다.

WebGPU 라는 이름이 “브라우저 그래픽” 만 떠올리게 하지만, Austin 이 명시적으로 정정한다. WebGPU 는 GPU 자체의 일반 추상 spec이고, 브라우저는 그 한 사용자다. 데스크탑 native 에서도 같은 spec 이 동일하게 돈다.

WebGPU 의 약속을 한 줄로 — “같은 코드가 NVIDIA · AMD · Intel · Apple · Qualcomm 위에서 모두 동작한다”. 단, 각 vendor 의 native API(CUDA, Metal Performance Shaders) 와 같은 정점의 perf 는 아니다. 그 갭의 이유가 § 07 의 본론.

WebGPU 의 셰이더 언어는 WGSL (WebGPU Shading Language) 이다. Rust 에 영향받은 syntax. CUDA · GLSL · HLSL · Metal 의 디자인 결정을 종합한 결과물.

// WGSL — vector add compute shader
@group(0) @binding(0) var<storage, read>       a: array<f32>;
@group(0) @binding(1) var<storage, read>       b: array<f32>;
@group(0) @binding(2) var<storage, read_write> c: array<f32>;

@compute @workgroup_size(256)
fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
    let i = gid.x;
    if (i < arrayLength(&a)) {
        c[i] = a[i] + b[i];
    }
}

WebGPU 를 raw C API 로 직접 부르면 — Austin 의 표현 그대로 — “Vulkan 보다 작지만 여전히 수백 줄 보일러플레이트가 박힌다”. gpu.cpp 는 single-header C++ 라이브러리로 그 보일러플레이트를 3개의 함수로 줄인다.

// gpu.cpp — vector add 전체
#include "gpu.h"
using namespace gpu;

int main(){
  Context ctx       = createContext();      // adapter + device

  Tensor a = createTensor(ctx, {N}, kf32, dataA);
  Tensor b = createTensor(ctx, {N}, kf32, dataB);
  Tensor c = createTensor(ctx, {N}, kf32);

  Kernel k = createKernel(ctx, kAddShader,
                          Bindings{a, b, c},
                          /*workgroups=*/{N/256, 1, 1});

  std::promise<void> p;
  dispatchKernel(ctx, k, p);                 // async
  wait(ctx, p.get_future());

  toCPU(ctx, c, host_out, sizeof(host_out));
}

Austin 이 강의에서 깐 작은 역사 — Vulkan 은 explicit 의 끝까지 갔다. 모든 메모리, 모든 동기화, 모든 pipeline 상태를 사용자가 박는다. 그게 driver 의 부담을 줄였지만 사용자의 부담을 폭발시켰다. WebGPU 는 그 사이의 절충 — explicit 의 핵심은 유지하되 default 를 박는다.

강의의 인사이트 — WebGPU 의 디자인 목표 자체가 “CUDA 의 ergonomics + portability”다. spec 위원회가 “Vulkan 의 explicit 은 너무 멀리 갔다” 는 합의에서 출발했다. 그래서 default value 가 풍부하고 (sampler, alignment, layout 등 자동), 동시에 vendor-extension 을 표준 안에 박지 않는다.

gpu.cpp 의 portability 의 한 차원이 다른 vendor(WebGPU 이 줌)이라면, 다른 한 차원은 browser ↔ native다. 같은 코드가 두 form factor 위에서 돈다.

# 같은 코드의 두 빌드 — gpu.cpp 의 CMake
# 1. Native — Dawn 정적 링크
cmake -B build -DGPUCPP_BACKEND=native
cmake --build build
./build/main                    # 데스크탑 GPU 위에서 실행

# 2. Browser — Emscripten
emcmake cmake -B build_web
cmake --build build_web
python -m http.server -d build_web 8000
# 브라우저로 localhost:8000 — 같은 코드 같은 GPU

Austin 이 강의에서 정직하게 깐 갭 — WebGPU 는 vendor-portable spec 이다 보니, NVIDIA mma / Tensor Core, FP8/INT4 같은 vendor specific 명령에 직접 닿지 않는다. GEMM perf 의 절반 이상이 이 명령에서 오는 모던 GPU 에서 결정적인 한계.

강의의 데모 시퀀스에서 Austin 이 직접 보여준 사례들과, gpu.cpp 가 가능하게 하는 새 form factor 들의 시나리오.

강의 끝 Q&A 에서 다룬 다음 단계들. WebGPU 는 정적 spec 이 아니라 — Khronos / W3C 가 활발히 extension 을 추가하는 중.

손에 새기기 — 실습 시퀀스

1. gpu.cpp 클론 + 첫 빌드 — git clone github.com/AnswerDotAI/gpu.cpp && cmake. native binary 가 자기 GPU 위에서 vector add 를 도는 것을 확인.
2. 같은 코드를 브라우저 빌드 — emscripten 으로. localhost:8000 에서 같은 셰이더가 도는 것을 확인. 동시에 두 빌드를 띄워두면 portability 의 의미가 손에 잡힌다.
3. WGSL 의 첫 셰이더 — @compute @workgroup_size(256) 으로 elementwise pow 를 짠다. global_invocation_id 의 의미를 직접 출력해서 본다.
4. gpu.cpp 의 GEMM 단계 — repo 안 examples/matmul 시리즈. naive → shared → register tile 을 한 단계씩 따라간다. 같은 코드의 CUDA 버전이 있다면 PTX 와 WGSL → MSL/SPIR-V transpile 결과를 diff.
5. WebGPU raw 와 gpu.cpp 의 비교 — 같은 vector add 를 raw WGPUDevice API 로 짜본다. 보일러플레이트의 양을 직접 손에 박는다.
6. shader hot-reload 데모 재현 — 한 process 가 도는 동안 WGSL 파일을 watch + reload. animation 의 transition 이 자연스럽게 흐르는지 확인.
7. LLM inference 데모 띄우기 — gpu.cpp 의 examples 에 Llama 1B 데모. 브라우저 탭에서 토큰이 나오는 것을 본 후, profiler 로 어떤 단계가 가장 오래 걸리는지 확인.
8. cooperative_matrix proposal 읽기 — W3C 의 proposal 문서. 같은 추상이 NVIDIA mma · Apple AMX · Intel XMX 를 어떻게 한 입구로 묶으려 하는지 본다.

• cooperative_matrix extension 의 stable 시점 — 강의 시점(2024 mid)에는 proposal. 그 후 어떤 backend(Dawn? wgpu?) 가 먼저 stable 로 받았는지.
• WGSL → MSL / SPIR-V / DXIL 의 lowering 결과 diff — 같은 GEMM 셰이더가 세 backend 에서 만드는 IR 의 차이. tint --dump 로 확인 가능.
• gpu.cpp 의 LLM inference 가 vendor-native 와의 perf 갭 — Llama 1B 같은 모델을 같은 GPU 위 cuBLAS / Metal MPS 로 돌렸을 때와의 토큰당 latency 비교.
• browser 안 LLM inference 의 실용 한계 — 메모리, latency, cold start. wasm + WebGPU 가 도달 가능한 모델 크기의 상한.
• subgroups extension 위 reduction — Triton 의 tl.sum 같은 패턴을 WGSL 로 짤 때 코드 양과 perf.
• AR/로봇 SoC 위 실측 — Snapdragon · Mali · 등 모바일 GPU 에서 같은 gpu.cpp 코드의 perf characteristic. CUDA-grade desktop GPU 와 어떤 차원이 다른지.
• portable performance vs portable correctness — vendor 별 같은 셰이더의 floating point 결정성. ML 학습 reproducibility 와의 충돌 자리.