Verify debugPrintf specifiers

CHANGELOG.md

@@ -123,6 +123,7 @@ By @beholdnec in [#8505](https://github.com/gfx-rs/wgpu/pull/8505).
 - Make `wgpu_types::texture::format::TextureChannel` accessible as `wgpu::TextureChannel`. By @TornaxO7 in [#9394](https://github.com/gfx-rs/wgpu/pull/9349).
 - Add support for `per_vertex` in Metal and DX12, as well as some validation for `per_vertex`, and a new enable extension, `wgpu_per_vertex`. By @inner-daemons in [#9219](https://github.com/gfx-rs/wgpu/pull/9219).
 - Add `ComputePass` version of `CommandEncoder::transition_resources` that allows intra-pass transitions. By @wingertge in [#9371](https://github.com/gfx-rs/wgpu/pull/9371).
+- Add support for shader `debugPrintf` in Metal and Vulkan, behind a wgsl extension. By @39ali [#9389](https://github.com/gfx-rs/wgpu/pull/9389).
 - `Device::create_texture_from_hal` now takes an explicit `initial_state: wgt::TextureUses` parameter declaring the state the wrapped foreign resource is already in. Previously the tracker hard-coded `TextureUses::UNINITIALIZED` for the wrapped texture, which is a content-discarding transition under the Vulkan spec. This affected zero-copy hardware-decoded video imports on the platforms where compressed modifiers are used. To migrate, pass `wgpu::TextureUses::UNINITIALIZED` to preserve the previous behaviour:
   ```diff
     let texture = unsafe {
@@ -161,6 +162,7 @@ By @beholdnec in [#8505](https://github.com/gfx-rs/wgpu/pull/8505).
 - spirv-out ray tracing pipelines. By @Vecvec in [#9085](https://github.com/gfx-rs/wgpu/pull/9085).
 - Add `spirv-out` ray tracing pipelines. By @Vecvec in [#9085](https://github.com/gfx-rs/wgpu/pull/9085).
 - Add `naga::front::wgsl::ParseError::notes()`. By @kwillemsen in [#9572](https://github.com/gfx-rs/wgpu/pull/9572).
+- Implement `debugPrintf()` in Vulkan backend. By @39ali [#9389](https://github.com/gfx-rs/wgpu/pull/9389).
 
 ### Changes
 

Cargo.toml

@@ -254,6 +254,7 @@ objc2-metal = { version = "0.3.2", default-features = false, features = [
     "MTLAccelerationStructureTypes",
     "MTLAccelerationStructureCommandEncoder",
     "MTLResidencySet",
+    "MTLLogState",
 ] }
 objc2-quartz-core = { version = "0.3.2", default-features = false, features = [
     "std",

README.md

@@ -91,6 +91,12 @@ See the [documentation](https://docs.rs/wgpu/latest/wgpu/index.html?search=env)
 
 When running the CTS, use the variables `DENO_WEBGPU_ADAPTER_NAME`, `DENO_WEBGPU_BACKEND`, `DENO_WEBGPU_POWER_PREFERENCE`, and `DENO_WEBGPU_DX12_COMPILER`.
 
+## Extension Specifications
+
+Some native-only functionality extends beyond core WebGPU. See the API specs for details:
+
+- [Shader `debugPrintf`](./docs/api-specs/debug_printf.md)
+
 ## Repo Overview
 
 For an overview of all the components in the gfx-rs ecosystem, see [the big picture](./docs/big-picture.png).

docs/api-specs/debug_printf.md

@@ -0,0 +1,43 @@
+# Shader `debugPrintf`
+
+`wgpu` supports shader debug printing on native backends when `Features::DEBUG_PRINTF` is enabled.
+This is a debugging extension and is not part of core WebGPU.
+
+## Requirements
+
+- Request `Features::DEBUG_PRINTF` when creating the device.
+- Add `enable wgpu_debug_printf;` to each WGSL module that calls `debugPrintf`.
+- Use a backend that advertises the feature:
+  - Metal with shader logging support, available in Metal 3.2 and later.
+  - Vulkan with `VK_KHR_shader_non_semantic_info` support.
+
+On Vulkan, `debugPrintf` output is produced through the validation layer debug-printf path. It is not enabled when GPU-assisted validation is enabled, because the two validation features are mutually exclusive.
+
+## WGSL Syntax
+
+`debugPrintf` is a statement-like built-in:
+
+```wgsl
+enable wgpu_debug_printf;
+
+@compute @workgroup_size(1)
+fn main(@builtin(global_invocation_id) id: vec3<u32>) {
+    debugPrintf("invocation: %u %u %u", id.x, id.y, id.z);
+}
+```
+
+The first argument must be a string literal. String literals are currently only accepted as the format argument to `debugPrintf`.
+
+Remaining arguments must currently be scalar values. Vector and matrix arguments may be supported in the future, but for now vector components should be passed individually.
+
+Format string interpretation follows the active backend's shader logging implementation. The supported format syntax is therefore intentionally limited to the common C-style debug printf forms accepted by Metal shader logging and Vulkan shader debug printf.
+
+## Backend Notes
+
+- Metal lowers `debugPrintf` to `metal::os_log_default.log_info`.
+- Vulkan lowers `debugPrintf` through SPIR-V `NonSemantic.DebugPrintf`.
+
+## References
+
+- Apple Metal shader logging: https://developer.apple.com/documentation/metal/logging-shader-debug-messages
+- Vulkan shader debug printf sample: https://docs.vulkan.org/samples/latest/samples/extensions/shader_debugprintf/README.html

naga/src/back/dot/mod.rs

@@ -429,6 +429,15 @@ impl StatementGraph {
                         "TraceRay"
                     }
                 },
+                S::DebugPrintf {
+                    format: _,
+                    ref arguments,
+                } => {
+                    for &expr in arguments {
+                        self.dependencies.push((id, expr, "arg"));
+                    }
+                    "DebugPrintf"
+                }
             };
             // Set the last node to the merge node
             last_node = merge_id;

naga/src/back/glsl/writer.rs

@@ -2264,6 +2264,7 @@ impl<'a, W: Write> Writer<'a, W> {
             }
             Statement::CooperativeStore { .. } => unimplemented!(),
             Statement::RayPipelineFunction(_) => unimplemented!(),
+            Statement::DebugPrintf { .. } => unimplemented!(),
         }
 
         Ok(())

naga/src/back/hlsl/writer.rs

@@ -3063,6 +3063,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
             }
             Statement::CooperativeStore { .. } => unimplemented!(),
             Statement::RayPipelineFunction(_) => unreachable!(),
+            Statement::DebugPrintf { .. } => unimplemented!(),
         }
 
         Ok(())

naga/src/back/msl/mod.rs

@@ -931,6 +931,7 @@ pub fn supported_capabilities() -> crate::valid::Capabilities {
         // No DRAW_INDEX
         // No MEMORY_DECORATION_VOLATILE
         | Caps::MEMORY_DECORATION_COHERENT
+        | Caps::DEBUG_PRINTF
 }
 
 #[test]

naga/src/back/msl/writer.rs

@@ -4288,6 +4288,21 @@ impl<W: Write> Writer<W> {
                     writeln!(self.out, ");")?;
                 }
                 crate::Statement::RayPipelineFunction(_) => unreachable!(),
+                crate::Statement::DebugPrintf {
+                    ref format,
+                    ref arguments,
+                } => {
+                    write!(
+                        self.out,
+                        "{level}metal::os_log_default.log_info(\"{}\"",
+                        format
+                    )?;
+                    for &arg in arguments {
+                        write!(self.out, ", ")?;
+                        self.put_expression(arg, &context.expression, true)?;
+                    }
+                    writeln!(self.out, ");")?;
+                }
             }
         }
 

naga/src/back/pipeline_constants.rs

@@ -921,6 +921,14 @@ fn adjust_stmt(new_pos: &HandleVec<Expression, Handle<Expression>>, stmt: &mut S
                 adjust(payload);
             }
         },
+        Statement::DebugPrintf {
+            format: _,
+            ref mut arguments,
+        } => {
+            for argument in arguments.iter_mut() {
+                adjust(argument);
+            }
+        }
         Statement::Break
         | Statement::Continue
         | Statement::Kill

naga/src/back/spv/block.rs

@@ -4180,6 +4180,19 @@ impl BlockContext<'_> {
                 Statement::RayPipelineFunction(ref fun) => {
                     self.write_ray_tracing_pipeline_function(fun, &mut block);
                 }
+                Statement::DebugPrintf {
+                    ref format,
+                    ref arguments,
+                } => {
+                    let mut format_params = Vec::with_capacity(arguments.len());
+                    for &arg in arguments {
+                        let word = self.cached[arg];
+                        format_params.push(word);
+                    }
+
+                    self.writer
+                        .write_debug_printf(&mut block, format, &format_params);
+                }
             }
         }
 

naga/src/back/spv/mod.rs

@@ -1222,4 +1222,5 @@ pub fn supported_capabilities() -> crate::valid::Capabilities {
         | Caps::DRAW_INDEX
         | Caps::MEMORY_DECORATION_COHERENT
         | Caps::MEMORY_DECORATION_VOLATILE
+        | Caps::DEBUG_PRINTF
 }

naga/src/back/wgsl/writer.rs

@@ -283,6 +283,38 @@ impl<W: Write> Writer<W> {
         Ok(())
     }
 
+    /// Returns `true` if a `debugPrintf` is found anywhere within this statement, `false` otherwise.
+    ///
+    /// Does not traverse into function calls.
+    fn find_debug_printf(stmt: &crate::Statement) -> bool {
+        match *stmt {
+            crate::Statement::DebugPrintf { .. } => true,
+            crate::Statement::Block(ref b) => b.iter().any(Self::find_debug_printf),
+            crate::Statement::If {
+                ref accept,
+                ref reject,
+                ..
+            } => {
+                accept.iter().any(Self::find_debug_printf)
+                    || reject.iter().any(Self::find_debug_printf)
+            }
+            crate::Statement::Loop {
+                ref body,
+                ref continuing,
+                ..
+            } => {
+                body.iter().any(Self::find_debug_printf)
+                    || continuing.iter().any(Self::find_debug_printf)
+            }
+            crate::Statement::Switch { ref cases, .. } => cases
+                .iter()
+                .any(|c| c.body.iter().any(Self::find_debug_printf)),
+            // Note: does not match on function `Call`s to look inside function bodies recursively.
+            // This is fine because this is called on all functions and entrypoints in a module to check for `debugPrintf` usage.
+            _ => false,
+        }
+    }
+
     /// Helper method which writes all the `enable` declarations
     /// needed for a module.
     fn write_enable_declarations(&mut self, module: &Module) -> BackendResult {
@@ -299,6 +331,7 @@ impl<W: Write> Writer<W> {
             ray_tracing_pipeline: bool,
             per_vertex: bool,
             binding_array: bool,
+            debug_printf: bool,
         }
         let mut needed = RequiredEnabled {
             mesh_shaders: module.uses_mesh_shaders(),
@@ -429,6 +462,19 @@ impl<W: Write> Writer<W> {
             needed.ray_tracing_pipeline = true;
         }
 
+        // search for debugPrintf statements in all functions and entry points
+        let mut functions_to_check = module.functions.iter().map(|(_, f)| f).collect::<Vec<_>>();
+        functions_to_check.extend(module.entry_points.iter().map(|ep| &ep.function));
+
+        'outer: for func in functions_to_check {
+            for stmt in func.body.iter() {
+                if Self::find_debug_printf(stmt) {
+                    needed.debug_printf = true;
+                    break 'outer;
+                }
+            }
+        }
+
         // Write required declarations
         let mut any_written = false;
         if needed.f16 {
@@ -475,6 +521,10 @@ impl<W: Write> Writer<W> {
             writeln!(self.out, "enable wgpu_per_vertex;")?;
             any_written = true;
         }
+        if needed.debug_printf {
+            writeln!(self.out, "enable wgpu_debug_printf;")?;
+            any_written = true;
+        }
         if any_written {
             // Empty line for readability
             writeln!(self.out)?;
@@ -1216,6 +1266,17 @@ impl<W: Write> Writer<W> {
                     writeln!(self.out, ");")?
                 }
             },
+            Statement::DebugPrintf {
+                ref format,
+                ref arguments,
+            } => {
+                write!(self.out, "{level}debugPrintf(\"{}\"", format)?;
+                for &arg in arguments {
+                    write!(self.out, ", ")?;
+                    self.write_expr(module, arg, func_ctx)?;
+                }
+                writeln!(self.out, ");")?;
+            }
         }
 
         Ok(())

naga/src/compact/statements.rs

@@ -168,6 +168,14 @@ impl FunctionTracer<'_> {
                             self.expressions_used.insert(payload);
                         }
                     },
+                    St::DebugPrintf {
+                        format: _,
+                        ref arguments,
+                    } => {
+                        for &expr in arguments {
+                            self.expressions_used.insert(expr);
+                        }
+                    }
 
                     // Trivial statements.
                     St::Break
@@ -406,6 +414,14 @@ impl FunctionMap {
                             adjust(payload);
                         }
                     },
+                    St::DebugPrintf {
+                        format: _,
+                        ref mut arguments,
+                    } => {
+                        for expr in arguments {
+                            adjust(expr);
+                        }
+                    }
 
                     // Trivial statements.
                     St::Break

naga/src/front/spv/mod.rs

@@ -1664,6 +1664,7 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                 }
                 S::WorkGroupUniformLoad { .. } => unreachable!(),
                 S::CooperativeStore { .. } => unreachable!(),
+                S::DebugPrintf { .. } => unreachable!(),
             }
             i += 1;
         }

naga/src/front/wgsl/error.rs

@@ -398,6 +398,10 @@ pub(crate) enum Error<'a> {
     FunctionMustUseReturnsVoid(Span, Span),
     FunctionMustUseOnNonFunction(Span),
     InvalidWorkGroupUniformLoad(Span),
+    InvalidStringLiteral {
+        span: Span,
+        description: &'static str,
+    },
     Internal(&'static str),
     ExpectedConstExprConcreteIntegerScalar(Span),
     ExpectedNonNegative(Span),
@@ -535,6 +539,7 @@ impl<'a> Error<'a> {
                         Token::Attribute => "@".to_string(),
                         Token::Number(_) => "number".to_string(),
                         Token::Word(s) => s.to_string(),
+                        Token::String(_) => "string literal".to_string(),
                         Token::Operation(c) => format!("operation (`{c}`)"),
                         Token::LogicalOperation(c) => format!("logical operation (`{c}`)"),
                         Token::ShiftOperation(c) => format!("bitshift (`{c}{c}`)"),
@@ -1124,6 +1129,11 @@ impl<'a> Error<'a> {
                 labels: vec![(span, "".into())],
                 notes: vec!["passed type must be a workgroup pointer".into()],
             },
+            Error::InvalidStringLiteral { span, description } => ParseError {
+                message: description.to_string(),
+                labels: vec![(span, "invalid string literal".into())],
+                notes: vec![],
+            },
             Error::Internal(message) => ParseError {
                 message: "internal WGSL front end error".to_string(),
                 labels: vec![],