diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 1d780b9c..e6c0fe19 100644 --- a/.github/workflows/ci.yml +++ b/.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 diff --git a/CHANGELOG.md b/CHANGELOG.md index 73a1791b..534512bb 100644 --- a/CHANGELOG.md +++ b/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 diff --git a/Makefile b/Makefile index fe328043..6f43dcb0 100644 --- a/Makefile +++ b/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 diff --git a/build.rs b/build.rs index 0b5f1741..e6f66b36 100644 --- a/build.rs +++ b/build.rs @@ -32,6 +32,7 @@ fn main() { ("WGPUComputePipeline", "WGPUComputePipelineImpl"), ("WGPUDevice", "WGPUDeviceImpl"), ("WGPUInstance", "WGPUInstanceImpl"), + ("WGPUPipelineCache", "WGPUPipelineCacheImpl"), ("WGPUPipelineLayout", "WGPUPipelineLayoutImpl"), ("WGPUQuerySet", "WGPUQuerySetImpl"), ("WGPUQueue", "WGPUQueueImpl"), diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 820d17e8..ed845573 100644 --- a/examples/CMakeLists.txt +++ b/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) diff --git a/examples/pipeline_cache/CMakeLists.txt b/examples/pipeline_cache/CMakeLists.txt new file mode 100644 index 00000000..09010a29 --- /dev/null +++ b/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}) diff --git a/examples/pipeline_cache/main.c b/examples/pipeline_cache/main.c new file mode 100644 index 00000000..f3749cfb --- /dev/null +++ b/examples/pipeline_cache/main.c @@ -0,0 +1,149 @@ +#include "framework.h" +#include "webgpu-headers/webgpu.h" +#include "wgpu.h" +#include +#include +#include + +#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; +} diff --git a/examples/pipeline_cache/shader.wgsl b/examples/pipeline_cache/shader.wgsl new file mode 100644 index 00000000..e80c1384 --- /dev/null +++ b/examples/pipeline_cache/shader.wgsl @@ -0,0 +1,9 @@ +@group(0) +@binding(0) +var values: array; + +@compute +@workgroup_size(1) +fn main(@builtin(global_invocation_id) global_id: vec3) { + values[global_id.x] = values[global_id.x] * 2u; +} diff --git a/ffi/wgpu.h b/ffi/wgpu.h index 4f7d4e0d..f29ec181 100644 --- a/ffi/wgpu.h +++ b/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); diff --git a/src/conv.rs b/src/conv.rs index 4ca15dd0..14b58d73 100644 --- a/src/conv.rs +++ b/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( + _descriptor: T, + extras: Option<&native::WGPUPipelineDescriptorExtras>, +) -> Option { + let pipeline_cache = extras?.pipelineCache.as_ref()?; + Some(pipeline_cache.id) +} + pub fn from_u64_bits>(value: u64) -> Option { if value > u32::MAX.into() { return None; diff --git a/src/lib.rs b/src/lib.rs index 087c5acd..b12536b5 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -3,9 +3,10 @@ use conv::{ from_u64_bits, map_adapter_type, map_backend_type, map_bind_group_entry, map_bind_group_layout_entry, map_device_descriptor, map_instance_backend_flags, - map_instance_descriptor, map_pipeline_layout_descriptor, map_query_set_descriptor, - map_query_set_index, map_sampler_border_color_extras, map_shader_module, - map_shader_runtime_checks, map_surface, map_surface_configuration, CreateSurfaceParams, + map_instance_descriptor, map_pipeline_cache_extras, map_pipeline_layout_descriptor, + map_query_set_descriptor, map_query_set_index, map_sampler_border_color_extras, + map_shader_module, map_shader_runtime_checks, map_surface, map_surface_configuration, + CreateSurfaceParams, }; use parking_lot::Mutex; use smallvec::SmallVec; @@ -201,6 +202,19 @@ pub struct WGPUInstanceImpl { context: Arc, } +pub struct WGPUPipelineCacheImpl { + context: Arc, + id: id::PipelineCacheId, +} +impl Drop for WGPUPipelineCacheImpl { + fn drop(&mut self) { + if !thread::panicking() { + let context = &self.context; + context.pipeline_cache_drop(self.id); + } + } +} + pub struct WGPUPipelineLayoutImpl { context: Arc, id: id::PipelineLayoutId, @@ -2075,8 +2089,10 @@ pub unsafe extern "C" fn wgpuDeviceCreateComputePipeline( // TODO(wgpu.h) zero_initialize_workgroup_memory: false, }, - // TODO(wgpu.h) - cache: None, + cache: follow_chain!( + map_pipeline_cache_extras((descriptor), + WGPUSType_PipelineDescriptorExtras => native::WGPUPipelineDescriptorExtras) + ), }; let (compute_pipeline_id, error) = @@ -2105,6 +2121,46 @@ pub unsafe extern "C" fn wgpuDeviceCreateComputePipeline( })) } +#[no_mangle] +pub unsafe extern "C" fn wgpuDeviceCreatePipelineCache( + device: native::WGPUDevice, + descriptor: Option<&native::WGPUPipelineCacheDescriptor>, +) -> native::WGPUPipelineCache { + let (device_id, context, error_sink) = { + let device = device.as_ref().expect("invalid device"); + (device.id, &device.context, &device.error_sink) + }; + let descriptor = descriptor.expect("invalid descriptor"); + + let desc = wgc::pipeline::PipelineCacheDescriptor { + label: string_view_into_label(descriptor.label), + data: if descriptor.data.is_null() { + None + } else { + Some(Cow::Borrowed(std::slice::from_raw_parts( + descriptor.data, + descriptor.dataSize, + ))) + }, + fallback: descriptor.fallback != 0, + }; + + let (pipeline_cache_id, error) = context.device_create_pipeline_cache(device_id, &desc, None); + if let Some(cause) = error { + handle_error( + error_sink, + cause, + desc.label, + "wgpuDeviceCreatePipelineCache", + ); + } + + Arc::into_raw(Arc::new(WGPUPipelineCacheImpl { + context: context.clone(), + id: pipeline_cache_id, + })) +} + #[no_mangle] pub unsafe extern "C" fn wgpuDeviceCreatePipelineLayout( device: native::WGPUDevice, @@ -2360,8 +2416,10 @@ pub unsafe extern "C" fn wgpuDeviceCreateRenderPipeline( }), // TODO(wgpu.h) multiview_mask: None, - // TODO(wgpu.h) - cache: None, + cache: follow_chain!( + map_pipeline_cache_extras((descriptor), + WGPUSType_PipelineDescriptorExtras => native::WGPUPipelineDescriptorExtras) + ), }; let (render_pipeline_id, error) = context.device_create_render_pipeline(device_id, &desc, None); @@ -2985,6 +3043,41 @@ pub unsafe extern "C" fn wgpuInstanceRelease(instance: native::WGPUInstance) { Arc::decrement_strong_count(instance); } +// PipelineCache methods + +#[no_mangle] +pub unsafe extern "C" fn wgpuPipelineCacheGetData( + pipeline_cache: native::WGPUPipelineCache, + data: *mut u8, + data_size: usize, +) -> usize { + let pipeline_cache = pipeline_cache.as_ref().expect("invalid pipeline cache"); + let context = &pipeline_cache.context; + + let Some(result) = context.pipeline_cache_get_data(pipeline_cache.id) else { + return 0; + }; + let count = result.len(); + + if !data.is_null() { + let count = count.min(data_size); + std::ptr::copy_nonoverlapping(result.as_ptr(), data, count); + } + + count +} + +#[no_mangle] +pub unsafe extern "C" fn wgpuPipelineCacheAddRef(pipeline_cache: native::WGPUPipelineCache) { + assert!(!pipeline_cache.is_null(), "invalid pipeline cache"); + Arc::increment_strong_count(pipeline_cache); +} +#[no_mangle] +pub unsafe extern "C" fn wgpuPipelineCacheRelease(pipeline_cache: native::WGPUPipelineCache) { + assert!(!pipeline_cache.is_null(), "invalid pipeline cache"); + Arc::decrement_strong_count(pipeline_cache); +} + // PipelineLayout methods #[no_mangle]