gpu.cpp

gpu.cpp is a lightweight library that makes portable GPU compute with C++ simple.

It focuses on general purpose native GPU computation, leveraging the WebGPU specification as a portable low-level GPU interface. This means we can drop in GPU code in C++ projects and have it run on Nvidia, Intel, AMD, and other GPUs. The same C++ code can work on a wide variety of laptops, workstations, mobile devices or virtually any hardware with Vulkan, Metal, or DirectX support.

Technical Objectives: Lightweight, Fast Iteration, and Low Boilerplate

With gpu.cpp we want to enable a high-leverage library for individual developers and researchers to incorporate GPU computation into programs relying on nothing more than a standard C++ compiler as tooling. Our goals are:

• High power-to-weight ratio API: Provide the smallest API surface area that can cover the full range of GPU compute needs.
• Fast compile/run cycles: Ensure projects can build nearly instantaneously, compile/run cycles should be <5 seconds on a modern laptop.
• Minimal dependencies and tooling overhead: A standard clang C++ compiler should be enough, no external library dependencies beyond the WebGPU native implementation.

The implementation aims for a small API surface area with minimum boilerplate. There are a small number of library operations to carry out an broad range of low-level GPU operations. We avoid abstractions that add layers of indirection, making the mapping between the gpu.cpp library to raw WebGPU API clear when it's needed.

In this spirit of fast experimentation, we also want near-instantaneous C++ builds taking no more than a second or two even on modestly capable personal computing devices. With this in mind, we not only keep the API surface area small, but also keep the implementation small and we also provide a prebuilt binary of the Dawn native WebGPU implementation.

The core library implementation in the header-only gpu.hpp source code is around 1000 lines of code. In addition to enabling instantaneous, semi-interactive compilation cycles, the small implementation surface area keeps maintenance burden low and the velocity of improvements high. We also pre-build Google's Dawn WebGPU implementation as a shared library binary. This allows builds to link the shared library with each build and incorporate Google's powerful native WebGPU implementation without paying the cost of re-compiling Dawn during development cycles.

For more advanced users and release deployments, we include cmake examples for building both Dawn with gpu.cpp end-to-end, but this is not required nor recommended for most users to get started.

Quick Start: Building and Running

To build a gpu.cpp project, you will need to have installed on your system:

• clang++ compiler installed with support for C++17.
• python3 and above, to run the script which downloads the Dawn shared library.
• make to build the project.
• Only on Linux systems - Vulkan drivers. If Vulkan is not installed, you can run sudo apt install libvulkan1 mesa-vulkan-drivers vulkan-tools to install them.

The only library dependency of gpu.cpp is a WebGPU implementation. Currently we support the Dawn native backend, but we plan to support other targets and WebGPU implementations (web browsers or other native implementations such as wgpu). Currently we support MacOS, Linux, and Windows (via WSL).

Optionally, Dawn can be built from scratch with gpu.cpp using the cmake build scripts provided - see the -cmake targets in the Makefile. However, this is recommended for advanced users only. Building Dawn dependencies with cmake takes much longer than using the precompiled Dawn shared library.

After cloning the repo, from the top-level gpu.cpp, you should be able to build and run the hello world GELU example by typing:

make

The first time you build and run the project this way, it will download a prebuilt shared library for the Dawn native WebGPU implementation automatically (using the setup.py script). This places the Dawn shared library in the third_party/lib directory. Afterwards you should see libdawn.dylib on MacOS or libdawn.so on Linux. This download only occurs once.

The build process itself should take a few seconds. If the build and executions is successful, you should see the output of the GELU computation:

Hello gpu.cpp!
--------------

  gelu(0.00) = 0.00
  gelu(0.10) = 0.05
  gelu(0.20) = 0.12
  gelu(0.30) = 0.19
  gelu(0.40) = 0.26
  gelu(0.50) = 0.35
  gelu(0.60) = 0.44
  gelu(0.70) = 0.53
  gelu(0.80) = 0.63
  gelu(0.90) = 0.73
  gelu(1.00) = 0.84
  gelu(1.10) = 0.95
  ...

Computed 10000 values of GELU(x)

If you need to clean up the build artifacts, you can run:

make clean

Hello World Tutorial: A GELU Kernel

As a real-world example for how to use gpu.cpp, let's start with a practical-but-simple example of a GPU kernel from neural networks.

GELU is a non-linear embarassingly parallel operation often used in modern large language model transformer-based architectures.

It takes as input a vector of floats and applies the GELU function to each element of the vector. The function is nonlinear, attenuating values below zero to near zero, approximating the y = x identity function for large positive values. For values close to zero, GELU smoothly interpolates between the identity function and the zero function.

The GELU code below will illustrate the three main aspects of setting up a GPU computation with gpu.cpp:

1. The code that runs on the GPU (in WebGPU Shading Language, or WGSL), implementing the compute operation.

2. The code that runs on the CPU (in C++) that sets up the GPU computation by allocating and preparing resources. For high performance, this code should be run ahead-of-time from the hot paths of the application.

3. The code that runs on the CPU (in C++) that dispatches the GPU computation and retrieves the results. The key concern of hot-path dispatch code is to eliminate or minimize any unnecessary resource allocation or data movement (offloading such concerns to step 2). A secondary consideration is that GPU dispatches are asynchronous. We work with standard C++ asynchronous primitives to manage the asynchronous aspect of kernel dispatch.

Here's a GELU kernel implemented (based on the CUDA implementation in llm.c) as on-device WebGPU WGSL code and invoked from the host using gpu.cpp library functions and types. It can be compiled using a standard C++ compiler (we recommend Clang):

#include <array>
#include <cstdio>
#include <future>

#include "gpu.hpp"

using namespace gpu; // createContext, createTensor, createKernel,
                     // dispatchKernel, wait, toCPU Bindings,
                     // Tensor, Kernel, Context, Shape, kf32

static const char *kGelu = R"(
const GELU_SCALING_FACTOR: f32 = 0.7978845608028654; // sqrt(2.0 / PI)
@group(0) @binding(0) var<storage, read_write> inp: array<{{precision}}>;
@group(0) @binding(1) var<storage, read_write> out: array<{{precision}}>;
@compute @workgroup_size({{workgroupSize}})
fn main(
    @builtin(global_invocation_id) GlobalInvocationID: vec3<u32>) {
    let i: u32 = GlobalInvocationID.x;
    if (i < arrayLength(&inp)) {
        let x: f32 = inp[i];
        out[i] = select(0.5 * x * (1.0 + tanh(GELU_SCALING_FACTOR
                 * (x + .044715 * x * x * x))), x, x > 10.0);
    }
}
)";

int main(int argc, char **argv) {
  Context ctx = createContext();
  static constexpr size_t N = 10000;
  std::array<float, N> inputArr, outputArr;
  for (int i = 0; i < N; ++i) {
    inputArr[i] = static_cast<float>(i) / 10.0; // dummy input data
  }
  Tensor input = createTensor(ctx, Shape{N}, kf32, inputArr.data());
  Tensor output = createTensor(ctx, Shape{N}, kf32);
  std::promise<void> promise;
  std::future<void> future = promise.get_future();
  Kernel op = createKernel(ctx, {kGelu, /* 1-D workgroup size */ 256, kf32},
                           Bindings{input, output},
                           /* number of workgroups */ {cdiv(N, 256), 1, 1});
  dispatchKernel(ctx, op, promise);
  wait(ctx, future);
  toCPU(ctx, output, outputArr.data(), sizeof(outputArr));
  for (int i = 0; i < 16; ++i) {
    printf("  gelu(%.2f) = %.2f\n", inputArr[i], outputArr[i]);
  }
  return 0;
}

Here we see the GPU code is quoted in a domain specific language called WGSL (WebGPU Shading Language). In a larger project, you might store this code in a separate file to be loaded at runtime (see examples/shadertui for a demonstration of live WGSL code re-loading).

The CPU code in main() sets up the host coordination for the GPU computation. We can think of the use of gpu.cpp library as a collection of GPU nouns and verbs.

The "nouns" are GPU resources modeled by the type definitions of the library and the "verbs" actions on GPU resources, modeled by the functions of the library. The ahead-of-time resource acquisition functions are prefaced with create*, such as:

• createContext() - constructs a reference to the GPU device context (Context).
• createTensor() - acquires a contiguous buffer on the GPU (Tensor).
• createKernel() - constructs a handle to resources for the GPU computation (Kernel), taking the shader code as input and the tensor resources to bind.

These resource acquisition functions are tied to resource types for interacting with the GPU:

• Context - a handle to the state of resources for interacting with the GPU device.
• Tensor - a buffer of data on the GPU.
• KernelCode - the code for a WGSL program that can be dispatched to the GPU. This is a thin wrapper around a WGSL string and also includes the workgroup size the code is designed to run with.
• Kernel - a GPU program that can be dispatched to the GPU. This accepts a KernelCode and a list of Tensor resources to bind for the dispatch computation. This takes an argument Bindings that is a list of Tensor instances and should map the bindings declared at the top of the WGSL code. In this example there's two bindings corresponding to the input buffer on the GPU and the ouptut buffer on the GPU.

In this example, the GELU computation is performed only once and the program immediately exits so preparing resources and dispatch are side-by-side. Other examples in the examples/ directory illustrate how resource acquisition is prepared ahead of time and dispatch occurs in the hot path like a render, model inference, or simulation loop.

Besides the create* resource acquisition functions, there are a few more "verbs" in the gpu.cpp library for handling dispatching execution to the GPU and data movement:

• dispatchKernel() - dispatches a Kernel to the GPU for computation. This is an asynchronous operation that returns immediately.
• wait() - blocks until the GPU computation is complete. This is a standard C++ future/promise pattern.
• toCPU() - moves data from the GPU to the CPU. This is a synchronous operation that blocks until the data is copied.
• toGPU() - moves data from the CPU to the GPU. This is a synchronous operation that blocks until the data is copied. In this particular example, toGPU() is not used because there's only one data movement from CPU to GPU in the program and that happens when the createTensor() function is called.

This example is available in examples/hello_world/run.cpp.

Other Examples: Matrix Multiplication, Physics Sim, and SDF Rendering

You can explore the example projects in examples/ which illustrate how to use gpu.cpp as a library.

After you have run make in the top-level directory which retrieves the prebuilt Dawn shared library, you can run each example by navigating to its directory and running make from the example's directory.

An example of tiled matrix multiplication is in examples/matmul. This implements a WebGPU version of the first few kernels of Simon Boehm's How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklog post. It currently runs at ~ 2.5+ TFLOPs on a Macbook Pro M1 Max laptop, which has a theoretical peak of 10.4 TFLOPs. Contributions to optimize this further are welcome.

A parallel physics simulation of an ensemble of double pendulums simulated in parallel with different initial conditions on the GPU is shown in examples/physics.

We also show some examples of signed distance function computations, rendered in the terminal as ascii. A 3D SDF of spheres is shown in examples/render and a shadertoy-like live-reloading example is in examples/shadertui.

Interestingly, given a starting example, LLMs such as Claude 3.5 Sonnet can be quite capable at writing low-level WGSL code for you - the other shaders in the shadertui example are written by the LLM.

Who is gpu.cpp for?

gpu.cpp is aimed at enabling projects requiring portable on-device GPU computation with minimal implementation complexity and friction. Some example use cases are:

• Development of GPU algorithms to be run on personal computing devices
• Direct standalone implementations of neural network models
• Physics sim