Add a vertex_pulling_transform to compilation_options, such that

pipelines can opt-in to using the technique. This also updates the
vertex_indices tests to run all of them both with and without this
option.

deno_webgpu/pipeline.rs

@@ -114,6 +114,7 @@ pub fn op_webgpu_create_compute_pipeline(
             entry_point: compute.entry_point.map(Cow::from),
             constants: Cow::Owned(compute.constants.unwrap_or_default()),
             zero_initialize_workgroup_memory: true,
+            vertex_pulling_transform: false,
         },
         cache: None,
     };
@@ -363,6 +364,7 @@ pub fn op_webgpu_create_render_pipeline(
                 constants: Cow::Owned(fragment.constants.unwrap_or_default()),
                 // Required to be true for WebGPU
                 zero_initialize_workgroup_memory: true,
+                vertex_pulling_transform: false,
             },
             targets: Cow::Owned(fragment.targets),
         })
@@ -388,6 +390,7 @@ pub fn op_webgpu_create_render_pipeline(
                 constants: Cow::Owned(args.vertex.constants.unwrap_or_default()),
                 // Required to be true for WebGPU
                 zero_initialize_workgroup_memory: true,
+                vertex_pulling_transform: false,
             },
             buffers: Cow::Owned(vertex_buffers),
         },

player/tests/data/bind-group.ron

@@ -59,6 +59,7 @@
                     entry_point: None,
                     constants: {},
                     zero_initialize_workgroup_memory: true,
+                    vertex_pulling_transform: false,
                 ),
             ),
         ),

player/tests/data/pipeline-statistics-query.ron

@@ -32,6 +32,7 @@
                     entry_point: None,
                     constants: {},
                     zero_initialize_workgroup_memory: true,
+                    vertex_pulling_transform: false,
                 ),
             ),
         ),

player/tests/data/quad.ron

@@ -60,6 +60,7 @@
                         entry_point: None,
                         constants: {},
                         zero_initialize_workgroup_memory: true,
+                        vertex_pulling_transform: false,
                     ),
                     buffers: [],
                 ),
@@ -69,6 +70,7 @@
                         entry_point: None,
                         constants: {},
                         zero_initialize_workgroup_memory: true,
+                        vertex_pulling_transform: false,
                     ),
                     targets: [
                         Some((

player/tests/data/zero-init-buffer.ron

@@ -136,6 +136,7 @@
                     entry_point: None,
                     constants: {},
                     zero_initialize_workgroup_memory: true,
+                    vertex_pulling_transform: false,
                 ),
             ),
         ),

player/tests/data/zero-init-texture-binding.ron

@@ -137,6 +137,7 @@
                     entry_point: None,
                     constants: {},
                     zero_initialize_workgroup_memory: true,
+                    vertex_pulling_transform: false,
                 ),
             ),
         ),

tests/tests/vertex_indices/mod.rs

@@ -185,6 +185,7 @@ struct Test {
     id_source: IdSource,
     draw_call_kind: DrawCallKind,
     encoder_kind: EncoderKind,
+    vertex_pulling_transform: bool,
 }
 
 impl Test {
@@ -298,6 +299,16 @@ async fn vertex_index_common(ctx: TestingContext) {
         cache: None,
     };
     let builtin_pipeline = ctx.device.create_render_pipeline(&pipeline_desc);
+    pipeline_desc
+        .vertex
+        .compilation_options
+        .vertex_pulling_transform = true;
+    let builtin_pipeline_vpt = ctx.device.create_render_pipeline(&pipeline_desc);
+    pipeline_desc
+        .vertex
+        .compilation_options
+        .vertex_pulling_transform = false;
+
     pipeline_desc.vertex.entry_point = "vs_main_buffers";
     pipeline_desc.vertex.buffers = &[
         wgpu::VertexBufferLayout {
@@ -312,6 +323,15 @@ async fn vertex_index_common(ctx: TestingContext) {
         },
     ];
     let buffer_pipeline = ctx.device.create_render_pipeline(&pipeline_desc);
+    pipeline_desc
+        .vertex
+        .compilation_options
+        .vertex_pulling_transform = true;
+    let buffer_pipeline_vpt = ctx.device.create_render_pipeline(&pipeline_desc);
+    pipeline_desc
+        .vertex
+        .compilation_options
+        .vertex_pulling_transform = false;
 
     let dummy = ctx
         .device
@@ -336,17 +356,20 @@ async fn vertex_index_common(ctx: TestingContext) {
         )
         .create_view(&wgpu::TextureViewDescriptor::default());
 
-    let mut tests = Vec::with_capacity(5 * 2 * 2);
+    let mut tests = Vec::with_capacity(5 * 2 * 2 * 2);
     for case in TestCase::ARRAY {
         for id_source in IdSource::ARRAY {
             for draw_call_kind in DrawCallKind::ARRAY {
                 for encoder_kind in EncoderKind::ARRAY {
-                    tests.push(Test {
-                        case,
-                        id_source,
-                        draw_call_kind,
-                        encoder_kind,
-                    })
+                    for vertex_pulling_transform in [false, true] {
+                        tests.push(Test {
+                            case,
+                            id_source,
+                            draw_call_kind,
+                            encoder_kind,
+                            vertex_pulling_transform,
+                        })
+                    }
                 }
             }
         }
@@ -357,8 +380,20 @@ async fn vertex_index_common(ctx: TestingContext) {
     let mut failed = false;
     for test in tests {
         let pipeline = match test.id_source {
-            IdSource::Buffers => &buffer_pipeline,
-            IdSource::Builtins => &builtin_pipeline,
+            IdSource::Buffers => {
+                if test.vertex_pulling_transform {
+                    &buffer_pipeline_vpt
+                } else {
+                    &buffer_pipeline
+                }
+            }
+            IdSource::Builtins => {
+                if test.vertex_pulling_transform {
+                    &builtin_pipeline_vpt
+                } else {
+                    &builtin_pipeline
+                }
+            }
         };
 
         let expected = test.expectation(&ctx);

wgpu-core/src/device/resource.rs

@@ -2737,6 +2737,7 @@ impl<A: HalApi> Device<A> {
                 entry_point: final_entry_point_name.as_ref(),
                 constants: desc.stage.constants.as_ref(),
                 zero_initialize_workgroup_memory: desc.stage.zero_initialize_workgroup_memory,
+                vertex_pulling_transform: false,
             },
             cache: cache.as_ref().and_then(|it| it.raw.as_ref()),
         };
@@ -3165,6 +3166,7 @@ impl<A: HalApi> Device<A> {
                 entry_point: &vertex_entry_point_name,
                 constants: stage_desc.constants.as_ref(),
                 zero_initialize_workgroup_memory: stage_desc.zero_initialize_workgroup_memory,
+                vertex_pulling_transform: stage_desc.vertex_pulling_transform,
             }
         };
 
@@ -3228,6 +3230,7 @@ impl<A: HalApi> Device<A> {
                     zero_initialize_workgroup_memory: fragment_state
                         .stage
                         .zero_initialize_workgroup_memory,
+                    vertex_pulling_transform: false,
                 })
             }
             None => None,

wgpu-core/src/pipeline.rs

@@ -166,6 +166,8 @@ pub struct ProgrammableStageDescriptor<'a> {
     /// This is required by the WebGPU spec, but may have overhead which can be avoided
     /// for cross-platform applications
     pub zero_initialize_workgroup_memory: bool,
+    /// Should the pipeline attempt to transform vertex shaders to use vertex pulling.
+    pub vertex_pulling_transform: bool,
 }
 
 /// Number of implicit bind groups derived at pipeline creation.

wgpu-hal/examples/halmark/main.rs

@@ -254,13 +254,15 @@ impl<A: hal::Api> Example<A> {
                 entry_point: "vs_main",
                 constants: &constants,
                 zero_initialize_workgroup_memory: true,
+                vertex_pulling_transform: false,
             },
             vertex_buffers: &[],
             fragment_stage: Some(hal::ProgrammableStage {
                 module: &shader,
                 entry_point: "fs_main",
                 constants: &constants,
                 zero_initialize_workgroup_memory: true,
+                vertex_pulling_transform: false,
             }),
             primitive: wgt::PrimitiveState {
                 topology: wgt::PrimitiveTopology::TriangleStrip,

wgpu-hal/examples/ray-traced-triangle/main.rs

@@ -373,6 +373,7 @@ impl<A: hal::Api> Example<A> {
                     entry_point: "main",
                     constants: &Default::default(),
                     zero_initialize_workgroup_memory: true,
+                    vertex_pulling_transform: false,
                 },
                 cache: None,
             })

wgpu-hal/src/lib.rs

@@ -1637,6 +1637,8 @@ pub struct ProgrammableStage<'a, A: Api> {
     /// This is required by the WebGPU spec, but may have overhead which can be avoided
     /// for cross-platform applications
     pub zero_initialize_workgroup_memory: bool,
+    /// Should the pipeline attempt to transform vertex shaders to use vertex pulling.
+    pub vertex_pulling_transform: bool,
 }
 
 // Rust gets confused about the impl requirements for `A`
@@ -1647,6 +1649,7 @@ impl<A: Api> Clone for ProgrammableStage<'_, A> {
             entry_point: self.entry_point,
             constants: self.constants,
             zero_initialize_workgroup_memory: self.zero_initialize_workgroup_memory,
+            vertex_pulling_transform: self.vertex_pulling_transform,
         }
     }
 }

wgpu-hal/src/metal/device.rs

@@ -158,7 +158,7 @@ impl super::Device {
                 metal::MTLPrimitiveTopologyClass::Point => true,
                 _ => false,
             },
-            vertex_pulling_transform: false,
+            vertex_pulling_transform: stage.vertex_pulling_transform,
             vertex_buffer_mappings: vertex_buffer_mappings.to_vec(),
         };
 

wgpu/src/backend/wgpu_core.rs

@@ -1175,6 +1175,10 @@ impl crate::Context for ContextWgpuCore {
                         .vertex
                         .compilation_options
                         .zero_initialize_workgroup_memory,
+                    vertex_pulling_transform: desc
+                        .vertex
+                        .compilation_options
+                        .vertex_pulling_transform,
                 },
                 buffers: Borrowed(&vertex_buffers),
             },
@@ -1189,6 +1193,7 @@ impl crate::Context for ContextWgpuCore {
                     zero_initialize_workgroup_memory: frag
                         .compilation_options
                         .zero_initialize_workgroup_memory,
+                    vertex_pulling_transform: false,
                 },
                 targets: Borrowed(frag.targets),
             }),
@@ -1242,6 +1247,7 @@ impl crate::Context for ContextWgpuCore {
                 zero_initialize_workgroup_memory: desc
                     .compilation_options
                     .zero_initialize_workgroup_memory,
+                vertex_pulling_transform: false,
             },
             cache: desc.cache.map(|c| c.id.into()),
         };

wgpu/src/lib.rs

@@ -1987,6 +1987,8 @@ pub struct PipelineCompilationOptions<'a> {
     /// This is required by the WebGPU spec, but may have overhead which can be avoided
     /// for cross-platform applications
     pub zero_initialize_workgroup_memory: bool,
+    /// Should the pipeline attempt to transform vertex shaders to use vertex pulling.
+    pub vertex_pulling_transform: bool,
 }
 
 impl<'a> Default for PipelineCompilationOptions<'a> {
@@ -2000,6 +2002,7 @@ impl<'a> Default for PipelineCompilationOptions<'a> {
         Self {
             constants,
             zero_initialize_workgroup_memory: true,
+            vertex_pulling_transform: false,
         }
     }
 }