removed clear heads shader (not needed)

Author: Geolm

lib/binning.h

@@ -3,7 +3,7 @@
 
 #include <stddef.h>
 
-static const size_t binning_shader_size = 26020;
+static const size_t binning_shader_size = 25989;
 static const char binning_shader[] =
     "// ---------------------------------------------------------------------------------------------------------------------------\n"
     "// Structures\n"
@@ -814,7 +814,7 @@ static const char binning_shader[] =
     "@compute @workgroup_size(1, 1, 1)\n"
     "fn write_indirect_args()\n"
     "{\n"
-    "    g_indirect_draw.vertex_count   = 4u; // 4 vertices for a quad/strip\n"
+    "    g_indirect_draw.vertex_count   = 4u;\n"
     "    g_indirect_draw.instance_count = atomicLoad(&g_counters.num_tiles);\n"
     "    g_indirect_draw.first_vertex   = 0u;\n"
     "    g_indirect_draw.first_instance = 0u;\n"

lib/onedraw.c

@@ -144,7 +144,6 @@ struct onedraw
     struct 
     {
         WGPUBuffer heads;
-        WGPUComputePipeline clear_pso;
         WGPUComputePipeline binning_pso;
         WGPUComputePipeline write_indirect_buffer_pso;
         WGPUBuffer counters;
@@ -933,21 +932,7 @@ void build_pso(struct onedraw* r, WGPUTextureFormat surface_format)
         .label = WGPU_STRING_VIEW("binning")
     });
 
-    WGPUComputePipelineDescriptor compute_pipeline_desc = 
-    {
-        .layout = layout,
-        .compute = 
-        {
-            .module = shader,
-            .entryPoint = WGPU_STRING_VIEW("clear_heads"),
-            .constants = NULL
-        }
-    };
-
-    r->tiles.clear_pso = wgpuDeviceCreateComputePipeline(r->device, &compute_pipeline_desc);
-    assert_msg(r->tiles.clear_pso != NULL, "can't create binning pso");
-
-    compute_pipeline_desc = (WGPUComputePipelineDescriptor)
+    WGPUComputePipelineDescriptor compute_pipeline_desc = (WGPUComputePipelineDescriptor)
     {
         .layout = layout,
         .compute = 
@@ -1310,18 +1295,6 @@ void od_end_frame(struct onedraw* r, WGPUTextureView target_view)
     // clear counter
     wgpuCommandEncoderClearBuffer(encoder, r->tiles.counters, 0, sizeof(uint32_t) * 2);
 
-    // clear heads pass
-    {
-        WGPUComputePassDescriptor desc = {.label = WGPU_STRING_VIEW("clear_heads pass")};
-        WGPUComputePassEncoder pass = wgpuCommandEncoderBeginComputePass(encoder, &desc);
-        wgpuComputePassEncoderSetPipeline(pass, r->tiles.clear_pso);
-        wgpuComputePassEncoderSetBindGroup(pass, 0, r->binding.binning_bindgroup, 0, NULL);
-        wgpuComputePassEncoderSetBindGroup(pass, 1, r->binding.frame_bindgroup[buffer_index], 0, NULL);
-        wgpuComputePassEncoderDispatchWorkgroups(pass, r->tiles.num_width * r->tiles.num_height, 1, 1);
-        wgpuComputePassEncoderEnd(pass);
-        wgpuComputePassEncoderRelease(pass);
-    }
-
     // binning pass
     {
         WGPUComputePassDescriptor desc = {.label = WGPU_STRING_VIEW("binning pass")};
@@ -1418,6 +1391,8 @@ void od_terminate(struct onedraw* r)
     SAFE_RELEASE(r->tiles.nodes, wgpuBufferRelease);
     SAFE_RELEASE(r->tiles.heads, wgpuBufferRelease);
     SAFE_RELEASE(r->rasterizer.pso, wgpuRenderPipelineRelease);
+    SAFE_RELEASE(r->tiles.binning_pso, wgpuComputePipelineRelease);
+    SAFE_RELEASE(r->tiles.write_indirect_buffer_pso, wgpuComputePipelineRelease);
     r->mem_interface.free_fn(r, r->mem_interface.user);
 }
 

src/shaders/binning.wgsl

@@ -613,7 +613,7 @@ fn tile_bin(@builtin(global_invocation_id) global_id: vec3<u32>)
 @compute @workgroup_size(1, 1, 1)
 fn write_indirect_args()
 {
-    g_indirect_draw.vertex_count   = 4u; // 4 vertices for a quad/strip
+    g_indirect_draw.vertex_count   = 4u;
     g_indirect_draw.instance_count = atomicLoad(&g_counters.num_tiles);
     g_indirect_draw.first_vertex   = 0u;
     g_indirect_draw.first_instance = 0u;