compute example uses push constants for testing

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)

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",

examples/compute/shader.wgsl

@@ -1,38 +1,15 @@
 @group(0)
 @binding(0)
-var<storage, read_write> v_indices: array<u32>; // this is used as both input and output for convenience
+var<storage, read_write> buffer: array<u32>;
 
-// 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_constant> push_constants: PushConstants;
 
 @compute
 @workgroup_size(1)
 fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
-    v_indices[global_id.x] = collatz_iterations(v_indices[global_id.x]);
+    let i = push_constants.i;
+    buffer[i] = i * 2;
 }