From 2befee3c663225b354d82a8e0d1e017010339e3b Mon Sep 17 00:00:00 2001 From: Zack Gomez Date: Thu, 10 Oct 2024 03:05:10 -0400 Subject: [PATCH] compute example uses push constants for testing --- examples/CMakeLists.txt | 2 + examples/compute/main.c | 116 +++++++++++++++++++++++++++++------ examples/compute/shader.wgsl | 35 ++--------- 3 files changed, 106 insertions(+), 47 deletions(-) diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 23b87380..9793f9ef 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -36,4 +36,6 @@ add_subdirectory(texture_arrays) add_subdirectory(triangle) set(GLFW_BUILD_DOCS OFF) +set(GLFW_BUILD_X11 OFF) +set(GLFW_BUILD_WAYLAND OFF) add_subdirectory(vendor/glfw) diff --git a/examples/compute/main.c b/examples/compute/main.c index 03dcaeba..2093264f 100644 --- a/examples/compute/main.c +++ b/examples/compute/main.c @@ -30,7 +30,7 @@ int main(int argc, char *argv[]) { UNUSED(argv) frmwrk_setup_logging(WGPULogLevel_Warn); - uint32_t numbers[] = {1, 2, 3, 4}; + uint32_t numbers[] = {0, 0, 0, 0}; uint32_t numbers_size = sizeof(numbers); uint32_t numbers_length = numbers_size / sizeof(uint32_t); @@ -41,9 +41,43 @@ int main(int argc, char *argv[]) { 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, NULL, handle_request_device, + wgpuAdapterRequestDevice(adapter, &device_desc, handle_request_device, (void *)&device); assert(device); @@ -54,15 +88,6 @@ int main(int argc, char *argv[]) { frmwrk_load_shader_module(device, "shader.wgsl"); assert(shader_module); - WGPUBuffer staging_buffer = wgpuDeviceCreateBuffer( - device, &(const WGPUBufferDescriptor){ - .label = "staging_buffer", - .usage = WGPUBufferUsage_MapRead | WGPUBufferUsage_CopyDst, - .size = numbers_size, - .mappedAtCreation = false, - }); - assert(staging_buffer); - WGPUBuffer storage_buffer = wgpuDeviceCreateBuffer( device, &(const WGPUBufferDescriptor){ .label = "storage_buffer", @@ -73,6 +98,56 @@ int main(int argc, char *argv[]) { }); 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", @@ -81,13 +156,10 @@ int main(int argc, char *argv[]) { .module = shader_module, .entryPoint = "main", }, + .layout = pipeline_layout, }); assert(compute_pipeline); - WGPUBindGroupLayout bind_group_layout = - wgpuComputePipelineGetBindGroupLayout(compute_pipeline, 0); - assert(bind_group_layout); - WGPUBindGroup bind_group = wgpuDeviceCreateBindGroup( device, &(const WGPUBindGroupDescriptor){ .label = "bind_group", @@ -121,14 +193,22 @@ int main(int argc, char *argv[]) { wgpuComputePassEncoderSetPipeline(compute_pass_encoder, compute_pipeline); wgpuComputePassEncoderSetBindGroup(compute_pass_encoder, 0, bind_group, 0, NULL); - wgpuComputePassEncoderDispatchWorkgroups(compute_pass_encoder, numbers_length, - 1, 1); + + 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", diff --git a/examples/compute/shader.wgsl b/examples/compute/shader.wgsl index 41af4363..e0b86e42 100644 --- a/examples/compute/shader.wgsl +++ b/examples/compute/shader.wgsl @@ -1,38 +1,15 @@ @group(0) @binding(0) -var v_indices: array; // this is used as both input and output for convenience +var buffer: array; -// The Collatz Conjecture states that for any integer n: -// If n is even, n = n/2 -// If n is odd, n = 3n+1 -// And repeat this process for each new n, you will always eventually reach 1. -// Though the conjecture has not been proven, no counterexample has ever been found. -// This function returns how many times this recurrence needs to be applied to reach 1. -fn collatz_iterations(n_base: u32) -> u32{ - var n: u32 = n_base; - var i: u32 = 0u; - loop { - if (n <= 1u) { - break; - } - if (n % 2u == 0u) { - n = n / 2u; - } - else { - // Overflow? (i.e. 3*n + 1 > 0xffffffffu?) - if (n >= 1431655765u) { // 0x55555555u - return 4294967295u; // 0xffffffffu - } - - n = 3u * n + 1u; - } - i = i + 1u; - } - return i; +struct PushConstants { + i: u32, } +var push_constants: PushConstants; @compute @workgroup_size(1) fn main(@builtin(global_invocation_id) global_id: vec3) { - v_indices[global_id.x] = collatz_iterations(v_indices[global_id.x]); + let i = push_constants.i; + buffer[i] = i * 2; }