diff --git a/Makefile b/Makefile index 6b8127d5..d6002a99 100644 --- a/Makefile +++ b/Makefile @@ -39,6 +39,8 @@ endif .PHONY: check test doc clear \ lib-native lib-native-release \ example-capture example-compute example-triangle \ + example-push_constants example-push_constants-release \ + run-example-push_constants run-example-push_constants-release \ example-capture-release example-compute-release example-triangle-release \ run-example-capture run-example-compute run-example-triangle \ run-example-capture-release run-example-compute-release run-example-triangle-release @@ -118,6 +120,18 @@ examples-debug: lib-native examples-release: lib-native-release cd examples && $(MKDIR_CMD) "build/RelWithDebInfo" && cd build/RelWithDebInfo && cmake -GNinja -DCMAKE_BUILD_TYPE=RelWithDebInfo -DCMAKE_EXPORT_COMPILE_COMMANDS=1 ../.. +example-push_constants: examples-debug + cd examples/build/Debug && cmake --build . --target push_constants + +run-example-push_constants: example-push_constants + cd examples/push_constants && "../build/Debug/push_constants/push_constants" + +example-push_constants-release: examples-release + cd examples/build/RelWithDebInfo && cmake --build . --target push_constants + +run-example-push_constants-release: example-push_constants-release + cd examples/push_constants && "../build/RelWithDebInfo/push_constants/push_constants" + example-capture: examples-debug cd examples/build/Debug && cmake --build . --target capture diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 23b87380..d11858fa 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -32,6 +32,7 @@ add_subdirectory(framework) add_subdirectory(capture) add_subdirectory(compute) add_subdirectory(enumerate_adapters) +add_subdirectory(push_constants) add_subdirectory(texture_arrays) add_subdirectory(triangle) diff --git a/examples/push_constants/CMakeLists.txt b/examples/push_constants/CMakeLists.txt new file mode 100644 index 00000000..2c00ba2e --- /dev/null +++ b/examples/push_constants/CMakeLists.txt @@ -0,0 +1,24 @@ +cmake_minimum_required(VERSION 3.20) +project(push_constants LANGUAGES C) + +add_executable(push_constants 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) +elseif(UNIX AND NOT APPLE) + set(OS_LIBRARIES "-lm -ldl") +elseif(APPLE) + set(OS_LIBRARIES "-framework CoreFoundation -framework QuartzCore -framework Metal") +endif() + +target_link_libraries(push_constants framework ${WGPU_LIBRARY} ${OS_LIBRARIES}) diff --git a/examples/push_constants/main.c b/examples/push_constants/main.c new file mode 100644 index 00000000..6018eaa8 --- /dev/null +++ b/examples/push_constants/main.c @@ -0,0 +1,253 @@ +#include +#include +#include + +#include "framework.h" +#include "webgpu-headers/webgpu.h" + +#define LOG_PREFIX "[push_constants]" + +static void handle_request_adapter(WGPURequestAdapterStatus status, + WGPUAdapter adapter, char const *message, + void *userdata) { + UNUSED(status) + UNUSED(message) + *(WGPUAdapter *)userdata = adapter; +} +static void handle_request_device(WGPURequestDeviceStatus status, + WGPUDevice device, char const *message, + void *userdata) { + UNUSED(status) + UNUSED(message) + *(WGPUDevice *)userdata = device; +} +static void handle_buffer_map(WGPUBufferMapAsyncStatus status, void *userdata) { + UNUSED(userdata) + printf(LOG_PREFIX " buffer_map status=%#.8x\n", status); +} + +int main(int argc, char *argv[]) { + UNUSED(argc) + UNUSED(argv) + frmwrk_setup_logging(WGPULogLevel_Warn); + + uint32_t numbers[] = {0, 0, 0, 0}; + uint32_t numbers_size = sizeof(numbers); + uint32_t numbers_length = numbers_size / sizeof(uint32_t); + + WGPUInstance instance = wgpuCreateInstance(NULL); + assert(instance); + + WGPUAdapter adapter = NULL; + wgpuInstanceRequestAdapter(instance, NULL, handle_request_adapter, + (void *)&adapter); + assert(adapter); + + WGPUSupportedLimitsExtras supported_limits_extras = { + .chain = + { + .sType = WGPUSType_SupportedLimitsExtras, + }, + .limits = + { + .maxPushConstantSize = 0, + }, + }; + WGPUSupportedLimits supported_limits = { + .nextInChain = &supported_limits_extras.chain, + }; + wgpuAdapterGetLimits(adapter, &supported_limits); + + WGPURequiredLimitsExtras required_limits_extras = { + .chain = + { + .sType = WGPUSType_RequiredLimitsExtras, + }, + .limits = supported_limits_extras.limits, + }; + WGPURequiredLimits required_limits = { + .nextInChain = &required_limits_extras.chain, + .limits = supported_limits.limits, + }; + + WGPUFeatureName requiredFeatures[] = { + WGPUNativeFeature_PushConstants, + }; + WGPUDeviceDescriptor device_desc = { + .label = "compute_device", + .requiredFeatures = requiredFeatures, + .requiredFeatureCount = 1, + .requiredLimits = &required_limits, + }; + + WGPUDevice device = NULL; + wgpuAdapterRequestDevice(adapter, &device_desc, handle_request_device, + (void *)&device); + assert(device); + + WGPUQueue queue = wgpuDeviceGetQueue(device); + assert(queue); + + WGPUShaderModule shader_module = + frmwrk_load_shader_module(device, "shader.wgsl"); + assert(shader_module); + + WGPUBuffer storage_buffer = wgpuDeviceCreateBuffer( + device, &(const WGPUBufferDescriptor){ + .label = "storage_buffer", + .usage = WGPUBufferUsage_Storage | WGPUBufferUsage_CopyDst | + WGPUBufferUsage_CopySrc, + .size = numbers_size, + .mappedAtCreation = false, + }); + assert(storage_buffer); + + WGPUBuffer staging_buffer = wgpuDeviceCreateBuffer( + device, &(const WGPUBufferDescriptor){ + .label = "staging_buffer", + .usage = WGPUBufferUsage_MapRead | WGPUBufferUsage_CopyDst, + .size = numbers_size, + .mappedAtCreation = false, + }); + assert(staging_buffer); + + WGPUPushConstantRange push_constant_range = { + .stages = WGPUShaderStage_Compute, + .start = 0, + .end = sizeof(uint32_t), + }; + + WGPUPipelineLayoutExtras pipeline_layout_extras = { + .chain = + { + .sType = WGPUSType_PipelineLayoutExtras, + }, + .pushConstantRangeCount = 1, + .pushConstantRanges = &push_constant_range, + }; + + WGPUBindGroupLayoutEntry bind_group_layout_entries[] = { + { + .binding = 0, + .visibility = WGPUShaderStage_Compute, + .buffer = + { + .type = WGPUBufferBindingType_Storage, + }, + }, + }; + WGPUBindGroupLayoutDescriptor bind_group_layout_desc = { + .label = "bind_group_layout", + .nextInChain = NULL, + .entryCount = 1, + .entries = bind_group_layout_entries, + }; + WGPUBindGroupLayout bind_group_layout = + wgpuDeviceCreateBindGroupLayout(device, &bind_group_layout_desc); + assert(bind_group_layout); + + WGPUPipelineLayoutDescriptor pipeline_layout_desc = { + .label = "pipeline_layout", + .nextInChain = &pipeline_layout_extras.chain, + .bindGroupLayouts = &bind_group_layout, + .bindGroupLayoutCount = 1, + }; + WGPUPipelineLayout pipeline_layout = + wgpuDeviceCreatePipelineLayout(device, &pipeline_layout_desc); + assert(pipeline_layout); + + WGPUComputePipeline compute_pipeline = wgpuDeviceCreateComputePipeline( + device, &(const WGPUComputePipelineDescriptor){ + .label = "compute_pipeline", + .compute = + (const WGPUProgrammableStageDescriptor){ + .module = shader_module, + .entryPoint = "main", + }, + .layout = pipeline_layout, + }); + assert(compute_pipeline); + + WGPUBindGroup bind_group = wgpuDeviceCreateBindGroup( + device, &(const WGPUBindGroupDescriptor){ + .label = "bind_group", + .layout = bind_group_layout, + .entryCount = 1, + .entries = + (const WGPUBindGroupEntry[]){ + (const WGPUBindGroupEntry){ + .binding = 0, + .buffer = storage_buffer, + .offset = 0, + .size = numbers_size, + }, + }, + }); + assert(bind_group); + + WGPUCommandEncoder command_encoder = wgpuDeviceCreateCommandEncoder( + device, &(const WGPUCommandEncoderDescriptor){ + .label = "command_encoder", + }); + assert(command_encoder); + + WGPUComputePassEncoder compute_pass_encoder = + wgpuCommandEncoderBeginComputePass(command_encoder, + &(const WGPUComputePassDescriptor){ + .label = "compute_pass", + }); + assert(compute_pass_encoder); + + wgpuComputePassEncoderSetPipeline(compute_pass_encoder, compute_pipeline); + wgpuComputePassEncoderSetBindGroup(compute_pass_encoder, 0, bind_group, 0, + NULL); + + for (uint32_t i = 0; i < numbers_length; i++) { + uint32_t pushConst = i; + wgpuComputePassEncoderSetPushConstants(compute_pass_encoder, 0, + sizeof(uint32_t), &pushConst); + + wgpuComputePassEncoderDispatchWorkgroups(compute_pass_encoder, + numbers_length, 1, 1); + } + + wgpuComputePassEncoderEnd(compute_pass_encoder); + wgpuComputePassEncoderRelease(compute_pass_encoder); + + wgpuCommandEncoderCopyBufferToBuffer(command_encoder, storage_buffer, 0, + staging_buffer, 0, numbers_size); + + WGPUCommandBuffer command_buffer = wgpuCommandEncoderFinish( + command_encoder, &(const WGPUCommandBufferDescriptor){ + .label = "command_buffer", + }); + assert(command_buffer); + + wgpuQueueWriteBuffer(queue, storage_buffer, 0, &numbers, numbers_size); + wgpuQueueSubmit(queue, 1, &command_buffer); + + wgpuBufferMapAsync(staging_buffer, WGPUMapMode_Read, 0, numbers_size, + handle_buffer_map, NULL); + wgpuDevicePoll(device, true, NULL); + + uint32_t *buf = + (uint32_t *)wgpuBufferGetMappedRange(staging_buffer, 0, numbers_size); + assert(buf); + + printf("times: [%d, %d, %d, %d]\n", buf[0], buf[1], buf[2], buf[3]); + + wgpuBufferUnmap(staging_buffer); + wgpuCommandBufferRelease(command_buffer); + wgpuCommandEncoderRelease(command_encoder); + wgpuBindGroupRelease(bind_group); + wgpuBindGroupLayoutRelease(bind_group_layout); + wgpuComputePipelineRelease(compute_pipeline); + wgpuBufferRelease(storage_buffer); + wgpuBufferRelease(staging_buffer); + wgpuShaderModuleRelease(shader_module); + wgpuQueueRelease(queue); + wgpuDeviceRelease(device); + wgpuAdapterRelease(adapter); + wgpuInstanceRelease(instance); + return EXIT_SUCCESS; +} diff --git a/examples/push_constants/shader.wgsl b/examples/push_constants/shader.wgsl new file mode 100644 index 00000000..e0b86e42 --- /dev/null +++ b/examples/push_constants/shader.wgsl @@ -0,0 +1,15 @@ +@group(0) +@binding(0) +var buffer: array; + +struct PushConstants { + i: u32, +} +var push_constants: PushConstants; + +@compute +@workgroup_size(1) +fn main(@builtin(global_invocation_id) global_id: vec3) { + let i = push_constants.i; + buffer[i] = i * 2; +}