Implement pipeline caches

.github/workflows/ci.yml

@@ -313,12 +313,14 @@ jobs:
           make example-enumerate_adapters
           make example-texture_arrays
           make example-immediates
+          make example-pipeline_cache
       - name: Run examples debug
         run: |
           make run-example-capture
           make run-example-compute
           make run-example-enumerate_adapters
           make run-example-immediates
+          make run-example-pipeline_cache
       - name: Build examples release
         run: |
           make example-capture-release
@@ -327,9 +329,11 @@ jobs:
           make example-enumerate_adapters-release
           make example-texture_arrays-release
           make example-immediates-release
+          make example-pipeline_cache-release
       - name: Run examples release
         run: |
           make run-example-capture-release
           make run-example-compute-release
           make run-example-enumerate_adapters-release
           make run-example-immediates-release
+          make run-example-pipeline_cache-release

CHANGELOG.md

@@ -73,6 +73,22 @@ The format is loosely based on [Keep a Changelog](https://keepachangelog.com/en/
 - `WGPUNativeFeature_StorageTextureArrayNonUniformIndexing`, `WGPUNativeFeature_Multiview`, `WGPUNativeFeature_ShaderFloat32Atomic`, `WGPUNativeFeature_TextureAtomic`, `WGPUNativeFeature_TextureFormatP010`, `WGPUNativeFeature_PipelineCache`, `WGPUNativeFeature_ShaderInt64AtomicMinMax`, `WGPUNativeFeature_ShaderInt64AtomicAllOps`, `WGPUNativeFeature_TextureInt64Atomic`, `WGPUNativeFeature_ShaderBarycentrics`, `WGPUNativeFeature_SelectiveMultiview`, `WGPUNativeFeature_MultisampleArray`, `WGPUNativeFeature_CooperativeMatrix`, `WGPUNativeFeature_ShaderPerVertex`, `WGPUNativeFeature_ShaderDrawIndex`, `WGPUNativeFeature_AccelerationStructureBindingArray`, `WGPUNativeFeature_MemoryDecorationCoherent`, `WGPUNativeFeature_MemoryDecorationVolatile` @lisyarus
 - `wgpuCommandEncoderClearTexture` @lisyarus
 - `wgpuDeviceCreateShaderModuleTrusted` @lisyarus
+- Pipeline caches, which let the result of shader compilation be reused between runs of a program. Requires `WGPUNativeFeature_PipelineCache`. @rael-g
+  - `WGPUPipelineCache` object, with `wgpuDeviceCreatePipelineCache`, `wgpuPipelineCacheGetData`, `wgpuPipelineCacheAddRef` and `wgpuPipelineCacheRelease`
+  - `WGPUPipelineDescriptorExtras` struct to be chained in `WGPURenderPipelineDescriptor` or `WGPUComputePipelineDescriptor`
+    ```c
+    WGPUPipelineCache pipeline_cache = wgpuDeviceCreatePipelineCache(
+        device, &(const WGPUPipelineCacheDescriptor){
+                    // Data previously returned by wgpuPipelineCacheGetData, if any.
+                    .data = data,
+                    .dataSize = data_size,
+                });
+
+    WGPUPipelineDescriptorExtras extras = {
+        .chain = { .sType = WGPUSType_PipelineDescriptorExtras },
+        .pipelineCache = pipeline_cache,
+    };
+    ```
 - Sampler address modes ClampToBorder and ClampToZero supported: @lisyarus
   - `WGPUNativeAddressMode_ClampToBorder` address mode
   - `WGPUSamplerBorderColor` enum

Makefile

@@ -170,6 +170,18 @@ example-enumerate_adapters-release: examples-release
 run-example-enumerate_adapters-release: example-enumerate_adapters-release
 	cd examples/triangle && "../build/RelWithDebInfo/enumerate_adapters/enumerate_adapters"
 
+example-pipeline_cache: examples-debug
+	cd examples/build/Debug && cmake --build . --target pipeline_cache
+
+run-example-pipeline_cache: example-pipeline_cache
+	cd examples/pipeline_cache && "../build/Debug/pipeline_cache/pipeline_cache"
+
+example-pipeline_cache-release: examples-release
+	cd examples/build/RelWithDebInfo && cmake --build . --target pipeline_cache
+
+run-example-pipeline_cache-release: example-pipeline_cache-release
+	cd examples/pipeline_cache && "../build/RelWithDebInfo/pipeline_cache/pipeline_cache"
+
 example-texture_arrays: examples-debug
 	cd examples/build/Debug && cmake --build . --target texture_arrays
 

build.rs

@@ -32,6 +32,7 @@ fn main() {
         ("WGPUComputePipeline", "WGPUComputePipelineImpl"),
         ("WGPUDevice", "WGPUDeviceImpl"),
         ("WGPUInstance", "WGPUInstanceImpl"),
+        ("WGPUPipelineCache", "WGPUPipelineCacheImpl"),
         ("WGPUPipelineLayout", "WGPUPipelineLayoutImpl"),
         ("WGPUQuerySet", "WGPUQuerySetImpl"),
         ("WGPUQueue", "WGPUQueueImpl"),

examples/CMakeLists.txt

@@ -34,6 +34,7 @@ add_subdirectory(compute)
 add_subdirectory(enumerate_adapters)
 add_subdirectory(immediates)
 add_subdirectory(metal_interop)
+add_subdirectory(pipeline_cache)
 add_subdirectory(texture_arrays)
 add_subdirectory(triangle)
 

examples/pipeline_cache/CMakeLists.txt

@@ -0,0 +1,24 @@
+cmake_minimum_required(VERSION 3.20)
+project(pipeline_cache LANGUAGES C)
+
+add_executable(pipeline_cache main.c)
+
+if (MSVC)
+    add_compile_options(/W4)
+else()
+    add_compile_options(-Wall -Wextra -Wpedantic)
+endif()
+
+include_directories(${CMAKE_SOURCE_DIR}/../ffi)
+include_directories(${CMAKE_SOURCE_DIR}/../ffi/webgpu-headers)
+include_directories(${CMAKE_SOURCE_DIR}/framework)
+
+if (WIN32)
+    set(OS_LIBRARIES d3dcompiler ws2_32 userenv bcrypt ntdll opengl32 Propsys RuntimeObject)
+elseif(UNIX AND NOT APPLE)
+    set(OS_LIBRARIES "-lm -ldl")
+elseif(APPLE)
+    set(OS_LIBRARIES "-framework Foundation -framework CoreFoundation -framework QuartzCore -framework Metal")
+endif()
+
+target_link_libraries(pipeline_cache framework ${WGPU_LIBRARY} ${OS_LIBRARIES})

examples/pipeline_cache/main.c

@@ -0,0 +1,149 @@
+#include "framework.h"
+#include "webgpu-headers/webgpu.h"
+#include "wgpu.h"
+#include <stdbool.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#define LOG_PREFIX "[pipeline_cache]"
+
+static void handle_request_adapter(WGPURequestAdapterStatus status,
+                                   WGPUAdapter adapter, WGPUStringView message,
+                                   void *userdata1, void *userdata2) {
+  UNUSED(status)
+  UNUSED(message)
+  UNUSED(userdata2)
+  *(WGPUAdapter *)userdata1 = adapter;
+}
+static void handle_request_device(WGPURequestDeviceStatus status,
+                                  WGPUDevice device, WGPUStringView message,
+                                  void *userdata1, void *userdata2) {
+  UNUSED(status)
+  UNUSED(message)
+  UNUSED(userdata2)
+  *(WGPUDevice *)userdata1 = device;
+}
+
+// Chains pipeline_cache onto the descriptor, so the cache serves and records
+// the compiled shader code.
+static WGPUComputePipeline create_pipeline(WGPUDevice device,
+                                           WGPUShaderModule shader_module,
+                                           WGPUPipelineCache pipeline_cache) {
+  WGPUPipelineDescriptorExtras pipeline_cache_extras = {
+      .chain =
+          (const WGPUChainedStruct){
+              .sType = (WGPUSType)WGPUSType_PipelineDescriptorExtras,
+          },
+      .pipelineCache = pipeline_cache,
+  };
+
+  return wgpuDeviceCreateComputePipeline(
+      device, &(const WGPUComputePipelineDescriptor){
+                  .nextInChain = (const WGPUChainedStruct *)&pipeline_cache_extras,
+                  .label = {"compute_pipeline", WGPU_STRLEN},
+                  .compute =
+                      (const WGPUComputeState){
+                          .module = shader_module,
+                          .entryPoint = {"main", WGPU_STRLEN},
+                      },
+              });
+}
+
+int main(int argc, char *argv[]) {
+  UNUSED(argc)
+  UNUSED(argv)
+  frmwrk_setup_logging(WGPULogLevel_Warn);
+
+  WGPUInstance instance = wgpuCreateInstance(NULL);
+  assert(instance);
+
+  WGPUAdapter adapter = NULL;
+  wgpuInstanceRequestAdapter(instance, NULL,
+                             (const WGPURequestAdapterCallbackInfo){
+                                 .callback = handle_request_adapter,
+                                 .userdata1 = &adapter
+                             });
+  assert(adapter);
+
+  // Not a failure: there is simply nothing to demonstrate on a backend that
+  // does not implement pipeline caches.
+  if (!wgpuAdapterHasFeature(adapter,
+                             (WGPUFeatureName)WGPUNativeFeature_PipelineCache)) {
+    printf(LOG_PREFIX " pipeline caches are unsupported on this adapter\n");
+    wgpuAdapterRelease(adapter);
+    wgpuInstanceRelease(instance);
+    return EXIT_SUCCESS;
+  }
+
+  WGPUFeatureName required_device_features[1] = {
+      (WGPUFeatureName)WGPUNativeFeature_PipelineCache,
+  };
+
+  WGPUDevice device = NULL;
+  wgpuAdapterRequestDevice(adapter,
+                           &(const WGPUDeviceDescriptor){
+                               .requiredFeatureCount = 1,
+                               .requiredFeatures = required_device_features,
+                           },
+                           (const WGPURequestDeviceCallbackInfo){
+                               .callback = handle_request_device,
+                               .userdata1 = &device
+                           });
+  assert(device);
+
+  WGPUShaderModule shader_module =
+      frmwrk_load_shader_module(device, "shader.wgsl");
+  assert(shader_module);
+
+  // Created without data, so it starts out empty; compiling a pipeline through
+  // it populates it.
+  WGPUPipelineCache pipeline_cache = wgpuDeviceCreatePipelineCache(
+      device, &(const WGPUPipelineCacheDescriptor){
+                  .label = {"pipeline_cache", WGPU_STRLEN},
+              });
+  assert(pipeline_cache);
+
+  WGPUComputePipeline compute_pipeline =
+      create_pipeline(device, shader_module, pipeline_cache);
+  assert(compute_pipeline);
+
+  size_t data_size = wgpuPipelineCacheGetData(pipeline_cache, NULL, 0);
+  printf(LOG_PREFIX " cache data size=%zu\n", data_size);
+
+  uint8_t *data = NULL;
+  if (data_size > 0) {
+    data = malloc(data_size);
+    assert(data);
+    size_t written = wgpuPipelineCacheGetData(pipeline_cache, data, data_size);
+    assert(written == data_size);
+  }
+
+  wgpuComputePipelineRelease(compute_pipeline);
+  wgpuPipelineCacheRelease(pipeline_cache);
+
+  // Stands in for persisting the data to disk and passing it back on a later
+  // run of the program.
+  WGPUPipelineCache reloaded_pipeline_cache = wgpuDeviceCreatePipelineCache(
+      device, &(const WGPUPipelineCacheDescriptor){
+                  .label = {"reloaded_pipeline_cache", WGPU_STRLEN},
+                  .data = data,
+                  .dataSize = data_size,
+                  .fallback = true,
+              });
+  assert(reloaded_pipeline_cache);
+
+  WGPUComputePipeline reloaded_compute_pipeline =
+      create_pipeline(device, shader_module, reloaded_pipeline_cache);
+  assert(reloaded_compute_pipeline);
+
+  printf(LOG_PREFIX " pipeline recreated from cache data\n");
+
+  free(data);
+  wgpuComputePipelineRelease(reloaded_compute_pipeline);
+  wgpuPipelineCacheRelease(reloaded_pipeline_cache);
+  wgpuShaderModuleRelease(shader_module);
+  wgpuDeviceRelease(device);
+  wgpuAdapterRelease(adapter);
+  wgpuInstanceRelease(instance);
+  return EXIT_SUCCESS;
+}

examples/pipeline_cache/shader.wgsl

@@ -0,0 +1,9 @@
+@group(0)
+@binding(0)
+var<storage, read_write> values: array<u32>;
+
+@compute
+@workgroup_size(1)
+fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
+    values[global_id.x] = values[global_id.x] * 2u;
+}

ffi/wgpu.h

@@ -39,6 +39,8 @@ typedef enum WGPUNativeSType
     WGPUSType_PrimitiveStateExtras = 0x0003000A,
     /** Identifies @ref WGPUSamplerDescriptorExtras. */
     WGPUSType_SamplerDescriptorExtras = 0x0003000B,
+    /** Identifies @ref WGPUPipelineDescriptorExtras. */
+    WGPUSType_PipelineDescriptorExtras = 0x0003000C,
     WGPUNativeSType_Force32 = 0x7FFFFFFF
 } WGPUNativeSType;
 
@@ -1474,6 +1476,45 @@ typedef struct WGPUSamplerDescriptorExtras {
     WGPUSamplerBorderColor samplerBorderColor;
 } WGPUSamplerDescriptorExtras WGPU_STRUCTURE_ATTRIBUTE;
 
+/**
+ * An opaque handle to a pipeline cache, used to reuse the result of shader
+ * compilation between executions of the program.
+ *
+ * Requires @ref WGPUNativeFeature_PipelineCache.
+ */
+typedef struct WGPUPipelineCacheImpl* WGPUPipelineCache WGPU_OBJECT_ATTRIBUTE;
+
+typedef struct WGPUPipelineCacheDescriptor {
+    WGPUChainedStruct const * nextInChain;
+    WGPUStringView label;
+    /**
+     * The data used to initialise the cache, which must have been returned by a
+     * previous call to @ref wgpuPipelineCacheGetData. May be NULL, in which case
+     * the cache starts out empty.
+     *
+     * Cache data is specific to the device, driver version and hardware it was
+     * produced on, and is treated as invalid anywhere else. See @ref fallback
+     * for how that case is handled.
+     */
+    WGPU_NULLABLE uint8_t const * data;
+    size_t dataSize;
+    /**
+     * Whether to create an empty cache when @ref data is invalid, rather than
+     * failing. Recommended to set to true.
+     */
+    WGPUBool fallback;
+} WGPUPipelineCacheDescriptor WGPU_STRUCTURE_ATTRIBUTE;
+
+/**
+ * Chained onto @ref WGPURenderPipelineDescriptor or
+ * @ref WGPUComputePipelineDescriptor to reuse compiled shader code from a
+ * pipeline cache.
+ */
+typedef struct WGPUPipelineDescriptorExtras {
+    WGPUChainedStruct chain;
+    WGPU_NULLABLE WGPUPipelineCache pipelineCache;
+} WGPUPipelineDescriptorExtras WGPU_STRUCTURE_ATTRIBUTE;
+
 #ifdef __cplusplus
 extern "C"
 {
@@ -1489,6 +1530,27 @@ extern "C"
     WGPUBool wgpuDevicePoll(WGPUDevice device, WGPUBool wait, WGPU_NULLABLE WGPUSubmissionIndex const *submissionIndex);
     WGPUShaderModule wgpuDeviceCreateShaderModuleSpirV(WGPUDevice device, WGPUShaderModuleDescriptorSpirV const *descriptor);
 
+    /**
+     * Creates a pipeline cache, which can be passed to pipeline creation via
+     * @ref WGPUPipelineDescriptorExtras to reuse previously compiled shader code.
+     *
+     * Requires @ref WGPUNativeFeature_PipelineCache.
+     */
+    WGPUPipelineCache wgpuDeviceCreatePipelineCache(WGPUDevice device, WGPUPipelineCacheDescriptor const *descriptor);
+    /**
+     * Returns the data of a pipeline cache, suitable for persisting to disk and
+     * passing back to @ref wgpuDeviceCreatePipelineCache on a later run.
+     *
+     * Following the same convention as @ref wgpuInstanceEnumerateAdapters, this
+     * is called twice: once with `data` set to NULL to obtain the required size,
+     * then again with a buffer of at least that size. `dataSize` is ignored when
+     * `data` is NULL. Returns the size of the data, which is 0 when the cache has
+     * no data to return.
+     */
+    size_t wgpuPipelineCacheGetData(WGPUPipelineCache pipelineCache, WGPU_NULLABLE uint8_t *data, size_t dataSize);
+    void wgpuPipelineCacheAddRef(WGPUPipelineCache pipelineCache);
+    void wgpuPipelineCacheRelease(WGPUPipelineCache pipelineCache);
+
     void wgpuSetLogCallback(WGPULogCallback callback, void *userdata);
 
     void wgpuSetLogLevel(WGPULogLevel level);

src/conv.rs

@@ -2057,6 +2057,20 @@ pub fn map_sampler_border_color_extras(
     }
 }
 
+/// Generic over the pipeline descriptor type, since the same extras struct is
+/// chained onto both render and compute pipeline descriptors.
+///
+/// # Safety
+///
+/// The `pipelineCache` field of `extras`, when set, must be a valid pipeline cache.
+pub unsafe fn map_pipeline_cache_extras<T>(
+    _descriptor: T,
+    extras: Option<&native::WGPUPipelineDescriptorExtras>,
+) -> Option<wgc::id::PipelineCacheId> {
+    let pipeline_cache = extras?.pipelineCache.as_ref()?;
+    Some(pipeline_cache.id)
+}
+
 pub fn from_u64_bits<T: bitflags::Flags<Bits = u32>>(value: u64) -> Option<T> {
     if value > u32::MAX.into() {
         return None;