From 16828d8058c6e818c33c5f5c01d9df6154e706b7 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Sun, 8 Oct 2023 14:44:14 -0400 Subject: [PATCH 01/17] spv: prepare for multiple ext inst sets --- naga/src/back/spv/block.rs | 15 +-- naga/src/back/spv/image.rs | 4 +- naga/src/back/spv/index.rs | 2 +- naga/src/back/spv/mod.rs | 2 +- naga/src/back/spv/writer.rs | 25 +++-- naga/src/front/spv/ext_inst.rs | 182 +++++++++++++++++++++++++++++++++ naga/src/front/spv/mod.rs | 152 +++------------------------ 7 files changed, 228 insertions(+), 154 deletions(-) create mode 100644 naga/src/front/spv/ext_inst.rs diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index 6c96fa09e3..46e96cbde8 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -133,8 +133,9 @@ impl Writer { )); let clamp_id = self.id_gen.next(); + body.push(Instruction::ext_inst( - self.gl450_ext_inst_id, + self.extension_inst_import("GLSL.std.450"), spirv::GLOp::FClamp, float_type_id, clamp_id, @@ -766,7 +767,7 @@ impl<'w> BlockContext<'w> { } MathOp::Custom(Instruction::ext_inst( - self.writer.gl450_ext_inst_id, + self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::FClamp, result_type_id, id, @@ -885,7 +886,7 @@ impl<'w> BlockContext<'w> { )); MathOp::Custom(Instruction::ext_inst( - self.writer.gl450_ext_inst_id, + self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::FMix, result_type_id, id, @@ -941,7 +942,7 @@ impl<'w> BlockContext<'w> { let lsb_id = self.gen_id(); block.body.push(Instruction::ext_inst( - self.writer.gl450_ext_inst_id, + self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::FindILsb, result_type_id, lsb_id, @@ -949,7 +950,7 @@ impl<'w> BlockContext<'w> { )); MathOp::Custom(Instruction::ext_inst( - self.writer.gl450_ext_inst_id, + self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::UMin, result_type_id, id, @@ -994,7 +995,7 @@ impl<'w> BlockContext<'w> { let msb_id = self.gen_id(); block.body.push(Instruction::ext_inst( - self.writer.gl450_ext_inst_id, + self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::FindUMsb, int_type_id, msb_id, @@ -1059,7 +1060,7 @@ impl<'w> BlockContext<'w> { block.body.push(match math_op { MathOp::Ext(op) => Instruction::ext_inst( - self.writer.gl450_ext_inst_id, + self.writer.extension_inst_import("GLSL.std.450"), op, result_type_id, id, diff --git a/naga/src/back/spv/image.rs b/naga/src/back/spv/image.rs index 460c906d47..052929280b 100644 --- a/naga/src/back/spv/image.rs +++ b/naga/src/back/spv/image.rs @@ -453,7 +453,7 @@ impl<'w> BlockContext<'w> { // `input_id` get treated as very large positive values. let restricted_id = self.gen_id(); block.body.push(Instruction::ext_inst( - self.writer.gl450_ext_inst_id, + self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::UMin, type_id, restricted_id, @@ -591,7 +591,7 @@ impl<'w> BlockContext<'w> { // `coordinates` get treated as very large positive values. let restricted_coordinates_id = self.gen_id(); block.body.push(Instruction::ext_inst( - self.writer.gl450_ext_inst_id, + self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::UMin, coordinates.type_id, restricted_coordinates_id, diff --git a/naga/src/back/spv/index.rs b/naga/src/back/spv/index.rs index 92e0f88d9a..a315bc356b 100644 --- a/naga/src/back/spv/index.rs +++ b/naga/src/back/spv/index.rs @@ -200,7 +200,7 @@ impl<'w> BlockContext<'w> { // BoundsCheckPolicy::Restrict. let restricted_index_id = self.gen_id(); block.body.push(Instruction::ext_inst( - self.writer.gl450_ext_inst_id, + self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::UMin, self.writer.get_uint_type_id(), restricted_index_id, diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index b7d57be0d4..c43691e307 100644 --- a/naga/src/back/spv/mod.rs +++ b/naga/src/back/spv/mod.rs @@ -623,7 +623,7 @@ pub struct Writer { // retain the table here between functions to save heap allocations. saved_cached: CachedExpressions, - gl450_ext_inst_id: Word, + ext_inst_ids: crate::FastHashMap<&'static str, Word>, // Just a temporary list of SPIR-V ids temp_list: Vec, diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index 4db86c93a7..ea2cd57abd 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -52,7 +52,8 @@ impl Writer { capabilities_used.insert(spirv::Capability::Shader); let mut id_gen = IdGenerator::default(); - let gl450_ext_inst_id = id_gen.next(); + let mut ext_inst_ids = crate::FastHashMap::default(); + ext_inst_ids.insert("GLSL.std.450", id_gen.next()); let void_type = id_gen.next(); Ok(Writer { @@ -76,7 +77,7 @@ impl Writer { global_variables: Vec::new(), binding_map: options.binding_map.clone(), saved_cached: CachedExpressions::default(), - gl450_ext_inst_id, + ext_inst_ids, temp_list: Vec::new(), }) } @@ -95,7 +96,8 @@ impl Writer { use std::mem::take; let mut id_gen = IdGenerator::default(); - let gl450_ext_inst_id = id_gen.next(); + let mut ext_inst_ids = take(&mut self.ext_inst_ids).recycle(); + ext_inst_ids.insert("GLSL.std.450", id_gen.next()); let void_type = id_gen.next(); // Every field of the old writer that is not determined by the `Options` @@ -111,7 +113,6 @@ impl Writer { // Initialized afresh: id_gen, void_type, - gl450_ext_inst_id, // Recycled: capabilities_used: take(&mut self.capabilities_used).recycle(), @@ -127,6 +128,7 @@ impl Writer { cached_constants: take(&mut self.cached_constants).recycle(), global_variables: take(&mut self.global_variables).recycle(), saved_cached: take(&mut self.saved_cached).recycle(), + ext_inst_ids, temp_list: take(&mut self.temp_list).recycle(), }; @@ -181,6 +183,13 @@ impl Writer { self.extensions_used.insert(extension); } + pub(super) fn extension_inst_import(&mut self, extension: &'static str) -> Word { + *self + .ext_inst_ids + .entry(extension) + .or_insert_with(|| self.id_gen.next()) + } + pub(super) fn get_type_id(&mut self, lookup_ty: LookupType) -> Word { match self.lookup_type.entry(lookup_ty) { Entry::Occupied(e) => *e.get(), @@ -1872,8 +1881,6 @@ impl Writer { .to_words(&mut self.logical_layout.extensions) } Instruction::type_void(self.void_type).to_words(&mut self.logical_layout.declarations); - Instruction::ext_inst_import(self.gl450_ext_inst_id, "GLSL.std.450") - .to_words(&mut self.logical_layout.ext_inst_imports); let mut debug_info_inner = None; if self.flags.contains(WriterFlags::DEBUG) { @@ -1983,6 +1990,12 @@ impl Writer { for extension in self.extensions_used.iter() { Instruction::extension(extension).to_words(&mut self.logical_layout.extensions); } + + for (ext, id) in self.ext_inst_ids.iter() { + Instruction::ext_inst_import(*id, ext) + .to_words(&mut self.logical_layout.ext_inst_imports); + } + if ir_module.entry_points.is_empty() { // SPIR-V doesn't like modules without entry points Instruction::capability(spirv::Capability::Linkage) diff --git a/naga/src/front/spv/ext_inst.rs b/naga/src/front/spv/ext_inst.rs new file mode 100644 index 0000000000..c1373e0bdc --- /dev/null +++ b/naga/src/front/spv/ext_inst.rs @@ -0,0 +1,182 @@ +use super::{Error, LookupExpression, LookupHelper as _}; + +struct ExtInst { + result_type_id: spirv::Word, + result_id: spirv::Word, + set_id: spirv::Word, + inst_id: spirv::Word, +} + +impl> super::Frontend { + pub(super) fn parse_ext_inst( + &mut self, + inst: super::Instruction, + span: crate::Span, + ctx: &mut super::BlockContext, + emitter: &mut crate::proc::Emitter, + block: &mut crate::Block, + block_id: spirv::Word, + body_idx: usize, + ) -> Result<(), Error> { + let base_wc = 5; + inst.expect_at_least(base_wc)?; + + let ext_inst = ExtInst { + result_type_id: self.next()?, + result_id: self.next()?, + set_id: self.next()?, + inst_id: self.next()?, + }; + let ext_name = if let Some(name) = self.ext_inst_imports.get(&ext_inst.set_id) { + name + } else { + return Err(Error::UnsupportedExtInstSet(ext_inst.set_id)); + }; + + match *ext_name { + "GLSL.std.450" => self.parse_ext_inst_glsl_std( + inst, ext_inst, span, ctx, emitter, block, block_id, body_idx, + ), + _ => { + return Err(Error::UnsupportedExtInstSet(ext_inst.set_id)); + } + } + } + fn parse_ext_inst_glsl_std( + &mut self, + inst: super::Instruction, + ext_inst: ExtInst, + span: crate::Span, + ctx: &mut super::BlockContext, + emitter: &mut crate::proc::Emitter, + block: &mut crate::Block, + block_id: spirv::Word, + body_idx: usize, + ) -> Result<(), Error> { + use crate::MathFunction as Mf; + use spirv::GLOp as Glo; + + let base_wc = 5; + + let gl_op = + Glo::from_u32(ext_inst.inst_id).ok_or(Error::UnsupportedExtInst(ext_inst.inst_id))?; + + let fun = match gl_op { + Glo::Round => Mf::Round, + Glo::RoundEven => Mf::Round, + Glo::Trunc => Mf::Trunc, + Glo::FAbs | Glo::SAbs => Mf::Abs, + Glo::FSign | Glo::SSign => Mf::Sign, + Glo::Floor => Mf::Floor, + Glo::Ceil => Mf::Ceil, + Glo::Fract => Mf::Fract, + Glo::Sin => Mf::Sin, + Glo::Cos => Mf::Cos, + Glo::Tan => Mf::Tan, + Glo::Asin => Mf::Asin, + Glo::Acos => Mf::Acos, + Glo::Atan => Mf::Atan, + Glo::Sinh => Mf::Sinh, + Glo::Cosh => Mf::Cosh, + Glo::Tanh => Mf::Tanh, + Glo::Atan2 => Mf::Atan2, + Glo::Asinh => Mf::Asinh, + Glo::Acosh => Mf::Acosh, + Glo::Atanh => Mf::Atanh, + Glo::Radians => Mf::Radians, + Glo::Degrees => Mf::Degrees, + Glo::Pow => Mf::Pow, + Glo::Exp => Mf::Exp, + Glo::Log => Mf::Log, + Glo::Exp2 => Mf::Exp2, + Glo::Log2 => Mf::Log2, + Glo::Sqrt => Mf::Sqrt, + Glo::InverseSqrt => Mf::InverseSqrt, + Glo::MatrixInverse => Mf::Inverse, + Glo::Determinant => Mf::Determinant, + Glo::ModfStruct => Mf::Modf, + Glo::FMin | Glo::UMin | Glo::SMin | Glo::NMin => Mf::Min, + Glo::FMax | Glo::UMax | Glo::SMax | Glo::NMax => Mf::Max, + Glo::FClamp | Glo::UClamp | Glo::SClamp | Glo::NClamp => Mf::Clamp, + Glo::FMix => Mf::Mix, + Glo::Step => Mf::Step, + Glo::SmoothStep => Mf::SmoothStep, + Glo::Fma => Mf::Fma, + Glo::FrexpStruct => Mf::Frexp, + Glo::Ldexp => Mf::Ldexp, + Glo::Length => Mf::Length, + Glo::Distance => Mf::Distance, + Glo::Cross => Mf::Cross, + Glo::Normalize => Mf::Normalize, + Glo::FaceForward => Mf::FaceForward, + Glo::Reflect => Mf::Reflect, + Glo::Refract => Mf::Refract, + Glo::PackUnorm4x8 => Mf::Pack4x8unorm, + Glo::PackSnorm4x8 => Mf::Pack4x8snorm, + Glo::PackHalf2x16 => Mf::Pack2x16float, + Glo::PackUnorm2x16 => Mf::Pack2x16unorm, + Glo::PackSnorm2x16 => Mf::Pack2x16snorm, + Glo::UnpackUnorm4x8 => Mf::Unpack4x8unorm, + Glo::UnpackSnorm4x8 => Mf::Unpack4x8snorm, + Glo::UnpackHalf2x16 => Mf::Unpack2x16float, + Glo::UnpackUnorm2x16 => Mf::Unpack2x16unorm, + Glo::UnpackSnorm2x16 => Mf::Unpack2x16snorm, + Glo::FindILsb => Mf::FindLsb, + Glo::FindUMsb | Glo::FindSMsb => Mf::FindMsb, + // TODO: https://github.com/gfx-rs/naga/issues/2526 + Glo::Modf | Glo::Frexp => return Err(Error::UnsupportedExtInst(ext_inst.inst_id)), + Glo::IMix + | Glo::PackDouble2x32 + | Glo::UnpackDouble2x32 + | Glo::InterpolateAtCentroid + | Glo::InterpolateAtSample + | Glo::InterpolateAtOffset => return Err(Error::UnsupportedExtInst(ext_inst.inst_id)), + }; + + let arg_count = fun.argument_count(); + inst.expect(base_wc + arg_count as u16)?; + let arg = { + let arg_id = self.next()?; + let lexp = self.lookup_expression.lookup(arg_id)?; + self.get_expr_handle(arg_id, lexp, ctx, emitter, block, body_idx) + }; + let arg1 = if arg_count > 1 { + let arg_id = self.next()?; + let lexp = self.lookup_expression.lookup(arg_id)?; + Some(self.get_expr_handle(arg_id, lexp, ctx, emitter, block, body_idx)) + } else { + None + }; + let arg2 = if arg_count > 2 { + let arg_id = self.next()?; + let lexp = self.lookup_expression.lookup(arg_id)?; + Some(self.get_expr_handle(arg_id, lexp, ctx, emitter, block, body_idx)) + } else { + None + }; + let arg3 = if arg_count > 3 { + let arg_id = self.next()?; + let lexp = self.lookup_expression.lookup(arg_id)?; + Some(self.get_expr_handle(arg_id, lexp, ctx, emitter, block, body_idx)) + } else { + None + }; + + let expr = crate::Expression::Math { + fun, + arg, + arg1, + arg2, + arg3, + }; + self.lookup_expression.insert( + ext_inst.result_id, + LookupExpression { + handle: ctx.expressions.append(expr, span), + type_id: ext_inst.result_type_id, + block_id, + }, + ); + Ok(()) + } +} diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs index d347302825..c21b8ed7fa 100644 --- a/naga/src/front/spv/mod.rs +++ b/naga/src/front/spv/mod.rs @@ -29,6 +29,7 @@ This value then gets used instead of `OpLoad` result later on. mod convert; mod error; +mod ext_inst; mod function; mod image; mod null; @@ -559,7 +560,7 @@ pub struct Frontend { state: ModuleState, layouter: Layouter, temp_bytes: Vec, - ext_glsl_id: Option, + ext_inst_imports: FastHashMap, future_decor: FastHashMap, future_member_decor: FastHashMap<(spirv::Word, MemberIndex), Decoration>, lookup_member: FastHashMap<(Handle, MemberIndex), LookupMember>, @@ -612,7 +613,7 @@ impl> Frontend { state: ModuleState::Empty, layouter: Layouter::default(), temp_bytes: Vec::new(), - ext_glsl_id: None, + ext_inst_imports: FastHashMap::default(), future_decor: FastHashMap::default(), future_member_decor: FastHashMap::default(), handle_sampling: FastHashMap::default(), @@ -2900,139 +2901,15 @@ impl> Frontend { emitter.start(ctx.expressions); } Op::ExtInst => { - use crate::MathFunction as Mf; - use spirv::GLOp as Glo; - - let base_wc = 5; - inst.expect_at_least(base_wc)?; - - let result_type_id = self.next()?; - let result_id = self.next()?; - let set_id = self.next()?; - if Some(set_id) != self.ext_glsl_id { - return Err(Error::UnsupportedExtInstSet(set_id)); - } - let inst_id = self.next()?; - let gl_op = Glo::from_u32(inst_id).ok_or(Error::UnsupportedExtInst(inst_id))?; - - let fun = match gl_op { - Glo::Round => Mf::Round, - Glo::RoundEven => Mf::Round, - Glo::Trunc => Mf::Trunc, - Glo::FAbs | Glo::SAbs => Mf::Abs, - Glo::FSign | Glo::SSign => Mf::Sign, - Glo::Floor => Mf::Floor, - Glo::Ceil => Mf::Ceil, - Glo::Fract => Mf::Fract, - Glo::Sin => Mf::Sin, - Glo::Cos => Mf::Cos, - Glo::Tan => Mf::Tan, - Glo::Asin => Mf::Asin, - Glo::Acos => Mf::Acos, - Glo::Atan => Mf::Atan, - Glo::Sinh => Mf::Sinh, - Glo::Cosh => Mf::Cosh, - Glo::Tanh => Mf::Tanh, - Glo::Atan2 => Mf::Atan2, - Glo::Asinh => Mf::Asinh, - Glo::Acosh => Mf::Acosh, - Glo::Atanh => Mf::Atanh, - Glo::Radians => Mf::Radians, - Glo::Degrees => Mf::Degrees, - Glo::Pow => Mf::Pow, - Glo::Exp => Mf::Exp, - Glo::Log => Mf::Log, - Glo::Exp2 => Mf::Exp2, - Glo::Log2 => Mf::Log2, - Glo::Sqrt => Mf::Sqrt, - Glo::InverseSqrt => Mf::InverseSqrt, - Glo::MatrixInverse => Mf::Inverse, - Glo::Determinant => Mf::Determinant, - Glo::ModfStruct => Mf::Modf, - Glo::FMin | Glo::UMin | Glo::SMin | Glo::NMin => Mf::Min, - Glo::FMax | Glo::UMax | Glo::SMax | Glo::NMax => Mf::Max, - Glo::FClamp | Glo::UClamp | Glo::SClamp | Glo::NClamp => Mf::Clamp, - Glo::FMix => Mf::Mix, - Glo::Step => Mf::Step, - Glo::SmoothStep => Mf::SmoothStep, - Glo::Fma => Mf::Fma, - Glo::FrexpStruct => Mf::Frexp, - Glo::Ldexp => Mf::Ldexp, - Glo::Length => Mf::Length, - Glo::Distance => Mf::Distance, - Glo::Cross => Mf::Cross, - Glo::Normalize => Mf::Normalize, - Glo::FaceForward => Mf::FaceForward, - Glo::Reflect => Mf::Reflect, - Glo::Refract => Mf::Refract, - Glo::PackUnorm4x8 => Mf::Pack4x8unorm, - Glo::PackSnorm4x8 => Mf::Pack4x8snorm, - Glo::PackHalf2x16 => Mf::Pack2x16float, - Glo::PackUnorm2x16 => Mf::Pack2x16unorm, - Glo::PackSnorm2x16 => Mf::Pack2x16snorm, - Glo::UnpackUnorm4x8 => Mf::Unpack4x8unorm, - Glo::UnpackSnorm4x8 => Mf::Unpack4x8snorm, - Glo::UnpackHalf2x16 => Mf::Unpack2x16float, - Glo::UnpackUnorm2x16 => Mf::Unpack2x16unorm, - Glo::UnpackSnorm2x16 => Mf::Unpack2x16snorm, - Glo::FindILsb => Mf::FindLsb, - Glo::FindUMsb | Glo::FindSMsb => Mf::FindMsb, - // TODO: https://github.com/gfx-rs/naga/issues/2526 - Glo::Modf | Glo::Frexp => return Err(Error::UnsupportedExtInst(inst_id)), - Glo::IMix - | Glo::PackDouble2x32 - | Glo::UnpackDouble2x32 - | Glo::InterpolateAtCentroid - | Glo::InterpolateAtSample - | Glo::InterpolateAtOffset => { - return Err(Error::UnsupportedExtInst(inst_id)) - } - }; - - let arg_count = fun.argument_count(); - inst.expect(base_wc + arg_count as u16)?; - let arg = { - let arg_id = self.next()?; - let lexp = self.lookup_expression.lookup(arg_id)?; - get_expr_handle!(arg_id, lexp) - }; - let arg1 = if arg_count > 1 { - let arg_id = self.next()?; - let lexp = self.lookup_expression.lookup(arg_id)?; - Some(get_expr_handle!(arg_id, lexp)) - } else { - None - }; - let arg2 = if arg_count > 2 { - let arg_id = self.next()?; - let lexp = self.lookup_expression.lookup(arg_id)?; - Some(get_expr_handle!(arg_id, lexp)) - } else { - None - }; - let arg3 = if arg_count > 3 { - let arg_id = self.next()?; - let lexp = self.lookup_expression.lookup(arg_id)?; - Some(get_expr_handle!(arg_id, lexp)) - } else { - None - }; - - let expr = crate::Expression::Math { - fun, - arg, - arg1, - arg2, - arg3, - }; - self.lookup_expression.insert( - result_id, - LookupExpression { - handle: ctx.expressions.append(expr, span), - type_id: result_type_id, - block_id, - }, - ); + self.parse_ext_inst( + inst, + span, + ctx, + &mut emitter, + &mut block, + block_id, + body_idx, + )?; } // Relational and Logical Instructions Op::LogicalNot => { @@ -4044,10 +3921,11 @@ impl> Frontend { if left != 0 { return Err(Error::InvalidOperand); } - if !SUPPORTED_EXT_SETS.contains(&name.as_str()) { + if let Some(ext) = SUPPORTED_EXT_SETS.iter().find(|ext| **ext == name.as_str()) { + self.ext_inst_imports.insert(result_id, ext); + } else { return Err(Error::UnsupportedExtSet(name)); } - self.ext_glsl_id = Some(result_id); Ok(()) } From fa88dcf3a2955f4bf951a07c6e7fb0e8c7a67028 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Mon, 16 Oct 2023 09:50:23 -0400 Subject: [PATCH 02/17] debug_printf: wgsl-in, wgsl-out, spv-out --- naga/src/back/dot/mod.rs | 6 +++++ naga/src/back/glsl/mod.rs | 3 +++ naga/src/back/hlsl/writer.rs | 3 +++ naga/src/back/msl/writer.rs | 3 +++ naga/src/back/spv/block.rs | 40 ++++++++++++++++++++++++------ naga/src/back/spv/image.rs | 4 +-- naga/src/back/spv/index.rs | 2 +- naga/src/back/spv/instructions.rs | 14 +++++++++-- naga/src/back/spv/mod.rs | 1 + naga/src/back/spv/writer.rs | 7 ++++++ naga/src/back/wgsl/writer.rs | 14 +++++++++++ naga/src/compact/statements.rs | 16 ++++++++++++ naga/src/front/spv/mod.rs | 3 ++- naga/src/front/wgsl/error.rs | 11 ++++++++ naga/src/front/wgsl/lower/mod.rs | 33 ++++++++++++++++++++++++ naga/src/front/wgsl/parse/ast.rs | 5 ++-- naga/src/front/wgsl/parse/lexer.rs | 8 ++++++ naga/src/front/wgsl/parse/mod.rs | 4 +++ naga/src/lib.rs | 4 +++ naga/src/proc/terminator.rs | 3 ++- naga/src/valid/analyzer.rs | 9 +++++++ naga/src/valid/function.rs | 3 +++ naga/src/valid/handles.rs | 6 +++++ 23 files changed, 186 insertions(+), 16 deletions(-) diff --git a/naga/src/back/dot/mod.rs b/naga/src/back/dot/mod.rs index 1556371df1..dc43af7dfb 100644 --- a/naga/src/back/dot/mod.rs +++ b/naga/src/back/dot/mod.rs @@ -279,6 +279,12 @@ impl StatementGraph { crate::RayQueryFunction::Terminate => "RayQueryTerminate", } } + S::DebugPrintf { ref arguments, .. } => { + for &arg in arguments { + self.dependencies.push((id, arg, "arg")); + } + "DebugPrintf" + } }; // Set the last node to the merge node last_node = merge_id; diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index e1dc906630..66f553b1a2 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -2379,6 +2379,9 @@ impl<'a, W: Write> Writer<'a, W> { writeln!(self.out, ");")?; } Statement::RayQuery { .. } => unreachable!(), + Statement::DebugPrintf { .. } => { + return Err(Error::Custom("debugPrintf is not implemented".to_string())); + } } Ok(()) diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index 0dd60c6ad7..b308e6f973 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/naga/src/back/hlsl/writer.rs @@ -2000,6 +2000,9 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { writeln!(self.out, "{level}}}")? } Statement::RayQuery { .. } => unreachable!(), + Statement::DebugPrintf { .. } => { + return Err(Error::Unimplemented("debug printf".to_string())); + } } Ok(()) diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index f900add71e..7a6e0e2960 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -3040,6 +3040,9 @@ impl Writer { } } } + crate::Statement::DebugPrintf { .. } => { + return Err(Error::FeatureNotImplemented("debug printf".to_string())); + } } } diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index 46e96cbde8..e60ecc1890 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -134,7 +134,7 @@ impl Writer { let clamp_id = self.id_gen.next(); - body.push(Instruction::ext_inst( + body.push(Instruction::ext_inst_glsl_std( self.extension_inst_import("GLSL.std.450"), spirv::GLOp::FClamp, float_type_id, @@ -766,7 +766,7 @@ impl<'w> BlockContext<'w> { arg2_id = self.writer.get_constant_composite(ty, &self.temp_list); } - MathOp::Custom(Instruction::ext_inst( + MathOp::Custom(Instruction::ext_inst_glsl_std( self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::FClamp, result_type_id, @@ -885,7 +885,7 @@ impl<'w> BlockContext<'w> { &self.temp_list, )); - MathOp::Custom(Instruction::ext_inst( + MathOp::Custom(Instruction::ext_inst_glsl_std( self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::FMix, result_type_id, @@ -941,7 +941,7 @@ impl<'w> BlockContext<'w> { }; let lsb_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_glsl_std( self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::FindILsb, result_type_id, @@ -949,7 +949,7 @@ impl<'w> BlockContext<'w> { &[arg0_id], )); - MathOp::Custom(Instruction::ext_inst( + MathOp::Custom(Instruction::ext_inst_glsl_std( self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::UMin, result_type_id, @@ -994,7 +994,7 @@ impl<'w> BlockContext<'w> { }; let msb_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_glsl_std( self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::FindUMsb, int_type_id, @@ -1059,7 +1059,7 @@ impl<'w> BlockContext<'w> { }; block.body.push(match math_op { - MathOp::Ext(op) => Instruction::ext_inst( + MathOp::Ext(op) => Instruction::ext_inst_glsl_std( self.writer.extension_inst_import("GLSL.std.450"), op, result_type_id, @@ -2334,6 +2334,32 @@ impl<'w> BlockContext<'w> { crate::Statement::RayQuery { query, ref fun } => { self.write_ray_query_function(query, fun, &mut block); } + crate::Statement::DebugPrintf { + ref format, + ref arguments, + } => { + self.writer.use_extension("SPV_KHR_non_semantic_info"); + let format_id = self.gen_id(); + self.writer + .strings + .push(Instruction::string(format, format_id)); + let id = self.gen_id(); + + self.temp_list.clear(); + self.temp_list.push(format_id); + for &argument in arguments { + self.temp_list.push(self.cached[argument]); + } + + let set_id = self.writer.extension_inst_import("NonSemantic.DebugPrintf"); + block.body.push(Instruction::ext_inst( + set_id, + 1, + self.writer.void_type, + id, + &self.temp_list, + )); + } } } diff --git a/naga/src/back/spv/image.rs b/naga/src/back/spv/image.rs index 052929280b..c9abdaad4d 100644 --- a/naga/src/back/spv/image.rs +++ b/naga/src/back/spv/image.rs @@ -452,7 +452,7 @@ impl<'w> BlockContext<'w> { // and negative values in a single instruction: negative values of // `input_id` get treated as very large positive values. let restricted_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_glsl_std( self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::UMin, type_id, @@ -590,7 +590,7 @@ impl<'w> BlockContext<'w> { // and negative values in a single instruction: negative values of // `coordinates` get treated as very large positive values. let restricted_coordinates_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_glsl_std( self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::UMin, coordinates.type_id, diff --git a/naga/src/back/spv/index.rs b/naga/src/back/spv/index.rs index a315bc356b..18249da94b 100644 --- a/naga/src/back/spv/index.rs +++ b/naga/src/back/spv/index.rs @@ -199,7 +199,7 @@ impl<'w> BlockContext<'w> { // One or the other of the index or length is dynamic, so emit code for // BoundsCheckPolicy::Restrict. let restricted_index_id = self.gen_id(); - block.body.push(Instruction::ext_inst( + block.body.push(Instruction::ext_inst_glsl_std( self.writer.extension_inst_import("GLSL.std.450"), spirv::GLOp::UMin, self.writer.get_uint_type_id(), diff --git a/naga/src/back/spv/instructions.rs b/naga/src/back/spv/instructions.rs index b963793ad3..b8e92b7624 100644 --- a/naga/src/back/spv/instructions.rs +++ b/naga/src/back/spv/instructions.rs @@ -121,18 +121,28 @@ impl super::Instruction { instruction } - pub(super) fn ext_inst( + pub(super) fn ext_inst_glsl_std( set_id: Word, op: spirv::GLOp, result_type_id: Word, id: Word, operands: &[Word], + ) -> Self { + Self::ext_inst(set_id, op as Word, result_type_id, id, operands) + } + + pub(super) fn ext_inst( + set_id: Word, + op: Word, + result_type_id: Word, + id: Word, + operands: &[Word], ) -> Self { let mut instruction = Self::new(Op::ExtInst); instruction.set_type(result_type_id); instruction.set_result(id); instruction.add_operand(set_id); - instruction.add_operand(op as u32); + instruction.add_operand(op); for operand in operands { instruction.add_operand(*operand) } diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index c43691e307..1c2446d716 100644 --- a/naga/src/back/spv/mod.rs +++ b/naga/src/back/spv/mod.rs @@ -603,6 +603,7 @@ pub struct Writer { /// The set of spirv extensions used. extensions_used: crate::FastIndexSet<&'static str>, + strings: Vec, debugs: Vec, annotations: Vec, flags: WriterFlags, diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index ea2cd57abd..812281b36e 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -63,6 +63,7 @@ impl Writer { capabilities_available: options.capabilities.clone(), capabilities_used, extensions_used: crate::FastIndexSet::default(), + strings: vec![], debugs: vec![], annotations: vec![], flags: options.flags, @@ -119,6 +120,7 @@ impl Writer { extensions_used: take(&mut self.extensions_used).recycle(), physical_layout: self.physical_layout.clone().recycle(), logical_layout: take(&mut self.logical_layout).recycle(), + strings: take(&mut self.strings).recycle(), debugs: take(&mut self.debugs).recycle(), annotations: take(&mut self.annotations).recycle(), lookup_type: take(&mut self.lookup_type).recycle(), @@ -2010,6 +2012,11 @@ impl Writer { Instruction::memory_model(addressing_model, memory_model) .to_words(&mut self.logical_layout.memory_model); + // Strings come before other debug instructions + for string in self.strings.iter() { + string.to_words(&mut self.logical_layout.debugs); + } + if self.flags.contains(WriterFlags::DEBUG) { for debug in self.debugs.iter() { debug.to_words(&mut self.logical_layout.debugs); diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs index c737934f5e..cbd24ed7f7 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -919,6 +919,20 @@ impl Writer { } } Statement::RayQuery { .. } => unreachable!(), + Statement::DebugPrintf { + ref format, + ref arguments, + } => { + write!(self.out, "{level}")?; + write!(self.out, "debugPrintf(\"{format}\",")?; + for (index, &argument) in arguments.iter().enumerate() { + if index != 0 { + write!(self.out, ", ")?; + } + self.write_expr(module, argument, func_ctx)?; + } + writeln!(self.out, ");")? + } } Ok(()) diff --git a/naga/src/compact/statements.rs b/naga/src/compact/statements.rs index 0698b57258..e05fdc13e9 100644 --- a/naga/src/compact/statements.rs +++ b/naga/src/compact/statements.rs @@ -97,6 +97,14 @@ impl FunctionTracer<'_> { self.expressions_used.insert(query); self.trace_ray_query_function(fun); } + St::DebugPrintf { + format: _, + ref arguments, + } => { + for expr in arguments { + self.expressions_used.insert(*expr); + } + } // Trivial statements. St::Break @@ -250,6 +258,14 @@ impl FunctionMap { adjust(query); self.adjust_ray_query_function(fun); } + St::DebugPrintf { + format: _, + ref mut arguments, + } => { + for expr in arguments { + adjust(expr); + } + } // Trivial statements. St::Break diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs index c21b8ed7fa..91297b4cf7 100644 --- a/naga/src/front/spv/mod.rs +++ b/naga/src/front/spv/mod.rs @@ -3687,7 +3687,8 @@ impl> Frontend { | S::Store { .. } | S::ImageStore { .. } | S::Atomic { .. } - | S::RayQuery { .. } => {} + | S::RayQuery { .. } + | S::DebugPrintf { .. } => {} S::Call { function: ref mut callee, ref arguments, diff --git a/naga/src/front/wgsl/error.rs b/naga/src/front/wgsl/error.rs index 5b3657f1f1..aca32c5d21 100644 --- a/naga/src/front/wgsl/error.rs +++ b/naga/src/front/wgsl/error.rs @@ -269,6 +269,8 @@ pub enum Error<'a> { scalar: String, inner: ConstantEvaluatorError, }, + /// String literals are only used with debugPrintf + UnexpectedStringLiteral(Span), } impl<'a> Error<'a> { @@ -283,6 +285,7 @@ impl<'a> Error<'a> { Token::Attribute => "@".to_string(), Token::Number(_) => "number".to_string(), Token::Word(s) => s.to_string(), + Token::String(_) => "string".to_string(), Token::Operation(c) => format!("operation ('{c}')"), Token::LogicalOperation(c) => format!("logical operation ('{c}')"), Token::ShiftOperation(c) => format!("bitshift ('{c}{c}')"), @@ -770,6 +773,14 @@ impl<'a> Error<'a> { format!("the expression should have been converted to have {} scalar type", scalar), ] }, + Error::UnexpectedStringLiteral(span) => ParseError { + message: "unexpected string literal".to_string(), + labels: vec![( + span, + "string literals can only be used as the first argument to debugPrintf".into(), + )], + notes: vec![], + }, } } } diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index ba9b49e135..30246dd621 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -1536,6 +1536,9 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { crate::Literal::AbstractFloat(f) } ast::Literal::Bool(b) => crate::Literal::Bool(b), + ast::Literal::String(_) => { + return Err(Error::UnexpectedStringLiteral(span)); + } }; let handle = ctx.interrupt_emitter(crate::Expression::Literal(literal), span)?; return Ok(Typed::Plain(handle)); @@ -2309,6 +2312,36 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { )?; return Ok(Some(handle)); } + "debugPrintf" => { + let format = arguments.first().ok_or(Error::WrongArgumentCount { + span, + expected: 1..16, + found: 0, + })?; + + let format = match ctx.ast_expressions[*format] { + ast::Expression::Literal(ast::Literal::String(format)) => { + format.to_string() + } + _ => return Err(Error::Internal("Expected format string")), + }; + + let arguments = arguments + .iter() + .skip(1) + .map(|&arg| self.expression(arg, ctx)) + .collect::, _>>()?; + let rctx = ctx.runtime_expression_ctx(span)?; + + rctx.block + .extend(rctx.emitter.finish(&rctx.function.expressions)); + + rctx.emitter.start(&rctx.function.expressions); + rctx.block + .push(crate::Statement::DebugPrintf { format, arguments }, span); + + return Ok(None); + } _ => return Err(Error::UnknownIdent(function.span, function.name)), } }; diff --git a/naga/src/front/wgsl/parse/ast.rs b/naga/src/front/wgsl/parse/ast.rs index dbaac523cb..72771c44ed 100644 --- a/naga/src/front/wgsl/parse/ast.rs +++ b/naga/src/front/wgsl/parse/ast.rs @@ -385,9 +385,10 @@ pub enum ConstructorType<'a> { } #[derive(Debug, Copy, Clone)] -pub enum Literal { +pub enum Literal<'a> { Bool(bool), Number(Number), + String(&'a str), } #[cfg(doc)] @@ -395,7 +396,7 @@ use crate::front::wgsl::lower::Lowerer; #[derive(Debug)] pub enum Expression<'a> { - Literal(Literal), + Literal(Literal<'a>), Ident(IdentExpr<'a>), /// A type constructor expression. diff --git a/naga/src/front/wgsl/parse/lexer.rs b/naga/src/front/wgsl/parse/lexer.rs index d03a448561..e30df05403 100644 --- a/naga/src/front/wgsl/parse/lexer.rs +++ b/naga/src/front/wgsl/parse/lexer.rs @@ -13,6 +13,7 @@ pub enum Token<'a> { Attribute, Number(Result), Word(&'a str), + String(&'a str), Operation(char), LogicalOperation(char), ShiftOperation(char), @@ -153,6 +154,13 @@ fn consume_token(input: &str, generic: bool) -> (Token<'_>, &str) { _ => (Token::Operation(cur), og_chars), } } + '"' => { + let (string, mut rest) = consume_any(&input[1..], |c| c != '"'); + if !rest.is_empty() { + rest = &rest[1..]; + } + (Token::String(string), rest) + } _ if is_blankspace(cur) => { let (_, rest) = consume_any(input, is_blankspace); (Token::Trivia, rest) diff --git a/naga/src/front/wgsl/parse/mod.rs b/naga/src/front/wgsl/parse/mod.rs index 51fc2f013b..312ee34e69 100644 --- a/naga/src/front/wgsl/parse/mod.rs +++ b/naga/src/front/wgsl/parse/mod.rs @@ -633,6 +633,10 @@ impl Parser { let num = res.map_err(|err| Error::BadNumber(span, err))?; ast::Expression::Literal(ast::Literal::Number(num)) } + (Token::String(string), _) => { + let _ = lexer.next(); + ast::Expression::Literal(ast::Literal::String(string)) + } (Token::Word("RAY_FLAG_NONE"), _) => { let _ = lexer.next(); ast::Expression::Literal(ast::Literal::Number(Number::U32(0))) diff --git a/naga/src/lib.rs b/naga/src/lib.rs index 5c63e7db4a..1aa42744a0 100644 --- a/naga/src/lib.rs +++ b/naga/src/lib.rs @@ -1865,6 +1865,10 @@ pub enum Statement { /// The specific operation we're performing on `query`. fun: RayQueryFunction, }, + DebugPrintf { + format: String, + arguments: Vec>, + }, } /// A function argument. diff --git a/naga/src/proc/terminator.rs b/naga/src/proc/terminator.rs index a5239d4eca..e82b513fd7 100644 --- a/naga/src/proc/terminator.rs +++ b/naga/src/proc/terminator.rs @@ -37,7 +37,8 @@ pub fn ensure_block_returns(block: &mut crate::Block) { | S::RayQuery { .. } | S::Atomic { .. } | S::WorkGroupUniformLoad { .. } - | S::Barrier(_)), + | S::Barrier(_) + | S::DebugPrintf { .. }), ) | None => block.push(S::Return { value: None }, Default::default()), } diff --git a/naga/src/valid/analyzer.rs b/naga/src/valid/analyzer.rs index df6fc5e9b0..c24cfc20fa 100644 --- a/naga/src/valid/analyzer.rs +++ b/naga/src/valid/analyzer.rs @@ -982,6 +982,15 @@ impl FunctionInfo { } FunctionUniformity::new() } + S::DebugPrintf { + format: _, + ref arguments, + } => { + for &argument in arguments { + let _ = self.add_ref(argument); + } + FunctionUniformity::new() + } }; disruptor = disruptor.or(uniformity.exit_disruptor()); diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index 3b12e59067..72e5e5db6b 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -904,6 +904,9 @@ impl super::Validator { crate::RayQueryFunction::Terminate => {} } } + S::DebugPrintf { .. } => { + // FIXME + } } } Ok(BlockInfo { stages, finished }) diff --git a/naga/src/valid/handles.rs b/naga/src/valid/handles.rs index e482f293bb..a59763f9bb 100644 --- a/naga/src/valid/handles.rs +++ b/naga/src/valid/handles.rs @@ -535,6 +535,12 @@ impl super::Validator { } Ok(()) } + crate::Statement::DebugPrintf { ref arguments, .. } => { + for arg in arguments.iter().copied() { + validate_expr(arg)?; + } + Ok(()) + } crate::Statement::Break | crate::Statement::Continue | crate::Statement::Kill From dfa93e7e609f15407dfb515b9a7f6ff908f8f3e4 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Mon, 16 Oct 2023 11:37:39 -0400 Subject: [PATCH 03/17] debug_printf: spv-in --- naga/src/front/spv/error.rs | 10 ++++--- naga/src/front/spv/ext_inst.rs | 54 +++++++++++++++++++++++++++++----- naga/src/front/spv/mod.rs | 10 +++++-- 3 files changed, 59 insertions(+), 15 deletions(-) diff --git a/naga/src/front/spv/error.rs b/naga/src/front/spv/error.rs index af025636c0..f4d51e414d 100644 --- a/naga/src/front/spv/error.rs +++ b/naga/src/front/spv/error.rs @@ -23,10 +23,12 @@ pub enum Error { UnsupportedExtension(String), #[error("unsupported extension set {0}")] UnsupportedExtSet(String), - #[error("unsupported extension instantiation set %{0}")] - UnsupportedExtInstSet(spirv::Word), - #[error("unsupported extension instantiation %{0}")] - UnsupportedExtInst(spirv::Word), + #[error("unsupported extension instantiation instruction id %{0} from set %{1}")] + UnsupportedExtInst(spirv::Word, &'static str), + #[error( + "extension instantiation references id %{0} which is not an imported extension instantation set" + )] + InvalidExtInst(spirv::Word), #[error("unsupported type {0:?}")] UnsupportedType(Handle), #[error("unsupported execution model %{0}")] diff --git a/naga/src/front/spv/ext_inst.rs b/naga/src/front/spv/ext_inst.rs index c1373e0bdc..388702d4bc 100644 --- a/naga/src/front/spv/ext_inst.rs +++ b/naga/src/front/spv/ext_inst.rs @@ -8,6 +8,7 @@ struct ExtInst { } impl> super::Frontend { + #[allow(clippy::too_many_arguments)] pub(super) fn parse_ext_inst( &mut self, inst: super::Instruction, @@ -30,20 +31,53 @@ impl> super::Frontend { let ext_name = if let Some(name) = self.ext_inst_imports.get(&ext_inst.set_id) { name } else { - return Err(Error::UnsupportedExtInstSet(ext_inst.set_id)); + // We get here only if the set_id doesn't point to an earlier OpExtInstImport. + // If the earlier ExtInstSet was unsupported we would have emitted an error then. + return Err(Error::InvalidExtInst(ext_inst.set_id)); }; match *ext_name { "GLSL.std.450" => self.parse_ext_inst_glsl_std( - inst, ext_inst, span, ctx, emitter, block, block_id, body_idx, + ext_name, inst, ext_inst, span, ctx, emitter, block, block_id, body_idx, ), - _ => { - return Err(Error::UnsupportedExtInstSet(ext_inst.set_id)); + "NonSemantic.DebugPrintf" if ext_inst.inst_id == 1 => { + self.parse_ext_inst_debug_printf(inst, span, ctx, emitter, block, body_idx) } + _ => Err(Error::UnsupportedExtInst(ext_inst.inst_id, ext_name)), } } + fn parse_ext_inst_debug_printf( + &mut self, + inst: super::Instruction, + span: crate::Span, + ctx: &mut super::BlockContext, + emitter: &mut crate::proc::Emitter, + block: &mut crate::Block, + body_idx: usize, + ) -> Result<(), Error> { + let base_wc = 5; + inst.expect_at_least(base_wc + 1)?; + let format_id = self.next()?; + let format = self.strings.lookup(format_id)?.clone(); + + block.extend(emitter.finish(ctx.expressions)); + + let mut arguments = Vec::with_capacity(inst.wc as usize - (base_wc as usize + 1)); + for _ in 0..arguments.capacity() { + let arg_id = self.next()?; + let lexp = self.lookup_expression.lookup(arg_id)?; + arguments.push(self.get_expr_handle(arg_id, lexp, ctx, emitter, block, body_idx)); + } + + block.push(crate::Statement::DebugPrintf { format, arguments }, span); + emitter.start(ctx.expressions); + + Ok(()) + } + #[allow(clippy::too_many_arguments)] fn parse_ext_inst_glsl_std( &mut self, + set_name: &'static str, inst: super::Instruction, ext_inst: ExtInst, span: crate::Span, @@ -58,8 +92,8 @@ impl> super::Frontend { let base_wc = 5; - let gl_op = - Glo::from_u32(ext_inst.inst_id).ok_or(Error::UnsupportedExtInst(ext_inst.inst_id))?; + let gl_op = Glo::from_u32(ext_inst.inst_id) + .ok_or(Error::UnsupportedExtInst(ext_inst.inst_id, set_name))?; let fun = match gl_op { Glo::Round => Mf::Round, @@ -124,13 +158,17 @@ impl> super::Frontend { Glo::FindILsb => Mf::FindLsb, Glo::FindUMsb | Glo::FindSMsb => Mf::FindMsb, // TODO: https://github.com/gfx-rs/naga/issues/2526 - Glo::Modf | Glo::Frexp => return Err(Error::UnsupportedExtInst(ext_inst.inst_id)), + Glo::Modf | Glo::Frexp => { + return Err(Error::UnsupportedExtInst(ext_inst.inst_id, set_name)) + } Glo::IMix | Glo::PackDouble2x32 | Glo::UnpackDouble2x32 | Glo::InterpolateAtCentroid | Glo::InterpolateAtSample - | Glo::InterpolateAtOffset => return Err(Error::UnsupportedExtInst(ext_inst.inst_id)), + | Glo::InterpolateAtOffset => { + return Err(Error::UnsupportedExtInst(ext_inst.inst_id, set_name)) + } }; let arg_count = fun.argument_count(); diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs index 91297b4cf7..81a2d1298d 100644 --- a/naga/src/front/spv/mod.rs +++ b/naga/src/front/spv/mod.rs @@ -76,8 +76,9 @@ pub const SUPPORTED_EXTENSIONS: &[&str] = &[ "SPV_KHR_storage_buffer_storage_class", "SPV_KHR_vulkan_memory_model", "SPV_KHR_multiview", + "SPV_KHR_non_semantic_info", ]; -pub const SUPPORTED_EXT_SETS: &[&str] = &["GLSL.std.450"]; +pub const SUPPORTED_EXT_SETS: &[&str] = &["GLSL.std.450", "NonSemantic.DebugPrintf"]; #[derive(Copy, Clone)] pub struct Instruction { @@ -561,6 +562,7 @@ pub struct Frontend { layouter: Layouter, temp_bytes: Vec, ext_inst_imports: FastHashMap, + strings: FastHashMap, future_decor: FastHashMap, future_member_decor: FastHashMap<(spirv::Word, MemberIndex), Decoration>, lookup_member: FastHashMap<(Handle, MemberIndex), LookupMember>, @@ -614,6 +616,7 @@ impl> Frontend { layouter: Layouter::default(), temp_bytes: Vec::new(), ext_inst_imports: FastHashMap::default(), + strings: FastHashMap::default(), future_decor: FastHashMap::default(), future_member_decor: FastHashMap::default(), handle_sampling: FastHashMap::default(), @@ -4020,8 +4023,9 @@ impl> Frontend { fn parse_string(&mut self, inst: Instruction) -> Result<(), Error> { self.switch(ModuleState::Source, inst.op)?; inst.expect_at_least(3)?; - let _id = self.next()?; - let (_name, _) = self.next_string(inst.wc - 2)?; + let id = self.next()?; + let (name, _) = self.next_string(inst.wc - 2)?; + self.strings.entry(id).or_insert(name); Ok(()) } From 67f59297012b680af47071e33632f1ce27da20da Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Mon, 16 Oct 2023 09:51:19 -0400 Subject: [PATCH 04/17] debug_printf: glsl-out --- naga/src/back/glsl/features.rs | 22 ++++++++++++++++++++++ naga/src/back/glsl/mod.rs | 10 ++++++++-- 2 files changed, 30 insertions(+), 2 deletions(-) diff --git a/naga/src/back/glsl/features.rs b/naga/src/back/glsl/features.rs index aaebfde9cb..1e610aff60 100644 --- a/naga/src/back/glsl/features.rs +++ b/naga/src/back/glsl/features.rs @@ -48,6 +48,8 @@ bitflags::bitflags! { /// /// We can always support this, either through the language or a polyfill const INSTANCE_INDEX = 1 << 22; + /// Debug Printf + const DEBUG_PRINTF = 1 << 23; } } @@ -243,6 +245,10 @@ impl FeaturesManager { // https://registry.khronos.org/OpenGL/extensions/EXT/EXT_blend_func_extended.txt writeln!(out, "#extension GL_EXT_blend_func_extended : require")?; } + if self.0.contains(Features::DEBUG_PRINTF) { + // https://github.com/KhronosGroup/GLSL/blob/master/extensions/ext/GLSL_EXT_debug_printf.txt + writeln!(out, "#extension GL_EXT_debug_printf : enable")?; + } if self.0.contains(Features::INSTANCE_INDEX) { if options.writer_flags.contains(WriterFlags::DRAW_PARAMETERS) { @@ -420,6 +426,22 @@ impl<'a, W> Writer<'a, W> { .. } = self; + for block in module + .functions + .iter() + .map(|(_, f)| &f.body) + .chain(std::iter::once(&entry_point.function.body)) + { + for statement in block.iter() { + match *statement { + crate::Statement::DebugPrintf { .. } => { + features.request(Features::DEBUG_PRINTF) + } + _ => {} + } + } + } + // Loop trough all expressions in both functions and the entry point // to check for needed features for (expressions, info) in module diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index 66f553b1a2..923fec5451 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -2379,8 +2379,14 @@ impl<'a, W: Write> Writer<'a, W> { writeln!(self.out, ");")?; } Statement::RayQuery { .. } => unreachable!(), - Statement::DebugPrintf { .. } => { - return Err(Error::Custom("debugPrintf is not implemented".to_string())); + Statement::DebugPrintf { + ref format, + ref arguments, + } => { + write!(self.out, "{level}")?; + write!(self.out, "debugPrintfEXT(\"{format}\",")?; + self.write_slice(arguments, |this, _, arg| this.write_expr(*arg, ctx))?; + writeln!(self.out, ");")? } } From 396067d673f6f29032386f7b1c9d1a1d33a90e32 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Mon, 16 Oct 2023 11:38:42 -0400 Subject: [PATCH 05/17] debug_printf: hlsl-out --- naga/src/back/hlsl/writer.rs | 15 +++++++++++++-- 1 file changed, 13 insertions(+), 2 deletions(-) diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index b308e6f973..da02ce0069 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/naga/src/back/hlsl/writer.rs @@ -2000,8 +2000,19 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { writeln!(self.out, "{level}}}")? } Statement::RayQuery { .. } => unreachable!(), - Statement::DebugPrintf { .. } => { - return Err(Error::Unimplemented("debug printf".to_string())); + Statement::DebugPrintf { + ref format, + ref arguments, + } => { + write!(self.out, "{level}")?; + write!(self.out, "printf(\"{format}\",")?; + for (index, argument) in arguments.iter().enumerate() { + if index != 0 { + write!(self.out, ", ")?; + } + self.write_expr(module, *argument, func_ctx)?; + } + writeln!(self.out, ");")? } } From 9737edc3cfcacdceb00da455925f280ba00579cb Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Mon, 16 Oct 2023 15:59:07 -0400 Subject: [PATCH 06/17] debug_printf: simple validation of debugPrintf arguments --- naga/src/valid/function.rs | 16 ++++++++++++++-- 1 file changed, 14 insertions(+), 2 deletions(-) diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index 72e5e5db6b..85de8f31e5 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -135,6 +135,11 @@ pub enum FunctionError { InvalidRayDescriptor(Handle), #[error("Ray Query {0:?} does not have a matching type")] InvalidRayQueryType(Handle), + #[error("Printf value argument {index} expression is invalid")] + InvalidPrintfArgument { + index: usize, + source: ExpressionError, + }, #[error( "Required uniformity of control flow for {0:?} in {1:?} is not fulfilled because of {2:?}" )] @@ -904,8 +909,15 @@ impl super::Validator { crate::RayQueryFunction::Terminate => {} } } - S::DebugPrintf { .. } => { - // FIXME + S::DebugPrintf { ref arguments, .. } => { + for (index, &expr) in arguments.iter().enumerate() { + context + .resolve_type_impl(expr, &self.valid_expression_set) + .map_err_inner(|source| { + FunctionError::InvalidPrintfArgument { index, source } + .with_span_handle(expr, context.expressions) + })?; + } } } } From d33bbdb4715529f2079682ab9bfc549e0e34fe20 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Mon, 16 Oct 2023 13:39:21 -0400 Subject: [PATCH 07/17] debug_printf: add tests --- naga/tests/in/debug-printf.wgsl | 4 +++ naga/tests/in/spv/debug-printf-s.spv | Bin 0 -> 348 bytes naga/tests/in/spv/debug-printf-s.spvasm | 28 ++++++++++++++++++ .../out/glsl/debug-printf-s.main.Compute.glsl | 18 +++++++++++ .../out/glsl/debug-printf.main.Compute.glsl | 14 +++++++++ naga/tests/out/hlsl/debug-printf-s.hlsl | 11 +++++++ naga/tests/out/hlsl/debug-printf-s.ron | 12 ++++++++ naga/tests/out/spv/debug-printf.spvasm | 23 ++++++++++++++ naga/tests/out/wgsl/debug-printf-s.wgsl | 9 ++++++ naga/tests/out/wgsl/debug-printf.wgsl | 5 ++++ naga/tests/snapshots.rs | 9 ++++++ 11 files changed, 133 insertions(+) create mode 100644 naga/tests/in/debug-printf.wgsl create mode 100644 naga/tests/in/spv/debug-printf-s.spv create mode 100644 naga/tests/in/spv/debug-printf-s.spvasm create mode 100644 naga/tests/out/glsl/debug-printf-s.main.Compute.glsl create mode 100644 naga/tests/out/glsl/debug-printf.main.Compute.glsl create mode 100644 naga/tests/out/hlsl/debug-printf-s.hlsl create mode 100644 naga/tests/out/hlsl/debug-printf-s.ron create mode 100644 naga/tests/out/spv/debug-printf.spvasm create mode 100644 naga/tests/out/wgsl/debug-printf-s.wgsl create mode 100644 naga/tests/out/wgsl/debug-printf.wgsl diff --git a/naga/tests/in/debug-printf.wgsl b/naga/tests/in/debug-printf.wgsl new file mode 100644 index 0000000000..1adcf9aa64 --- /dev/null +++ b/naga/tests/in/debug-printf.wgsl @@ -0,0 +1,4 @@ +@compute @workgroup_size(1) +fn main() { + debugPrintf("%d", 42); +} diff --git a/naga/tests/in/spv/debug-printf-s.spv b/naga/tests/in/spv/debug-printf-s.spv new file mode 100644 index 0000000000000000000000000000000000000000..96ca19fdcd161124227d7b84ad9a9b5a48308585 GIT binary patch literal 348 zcmY*Uy9&ZU5FF#%_(be16tPhV7Ak&#prD8tunoRR;RQ7QyB2~oF&5r2?9J}Z>|WC9 zWPt>l07`^T4Hh#%0Xewqt+$wOJQEo2e!A{V=pQ`aoCDTvWTL*=(p@_D@aPOj1L!7h^u=ibYn{Ye(;|w|0}) Io0BHQ2h~|1>Hq)$ literal 0 HcmV?d00001 diff --git a/naga/tests/in/spv/debug-printf-s.spvasm b/naga/tests/in/spv/debug-printf-s.spvasm new file mode 100644 index 0000000000..992f133ff8 --- /dev/null +++ b/naga/tests/in/spv/debug-printf-s.spvasm @@ -0,0 +1,28 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 18 +OpCapability Shader +OpCapability Linkage +OpExtension "SPV_KHR_non_semantic_info" +%1 = OpExtInstImport "GLSL.std.450" +%15 = OpExtInstImport "NonSemantic.DebugPrintf" +OpMemoryModel Logical GLSL450 +%13 = OpString "%d" +%16 = OpString "%v4f" +%2 = OpTypeVoid +%5 = OpTypeFunction %2 +%6 = OpTypeInt 32 1 +%7 = OpConstant %6 42 +%8 = OpTypeFloat 32 +%9 = OpConstant %8 3.3 +%10 = OpTypeVector %8 4 +%11 = OpConstantComposite %10 %9 %9 %9 %9 +%4 = OpFunction %2 None %5 +%3 = OpLabel +OpBranch %12 +%12 = OpLabel +%14 = OpExtInst %2 %15 1 %13 %7 +%17 = OpExtInst %2 %15 1 %16 %11 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/glsl/debug-printf-s.main.Compute.glsl b/naga/tests/out/glsl/debug-printf-s.main.Compute.glsl new file mode 100644 index 0000000000..2cd87e2b25 --- /dev/null +++ b/naga/tests/out/glsl/debug-printf-s.main.Compute.glsl @@ -0,0 +1,18 @@ +#version 310 es +#extension GL_EXT_debug_printf : enable + +precision highp float; +precision highp int; + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + + +void main_1() { + debugPrintfEXT("%d",42); + return; +} + +void main() { + main_1(); +} + diff --git a/naga/tests/out/glsl/debug-printf.main.Compute.glsl b/naga/tests/out/glsl/debug-printf.main.Compute.glsl new file mode 100644 index 0000000000..0923ee6cdc --- /dev/null +++ b/naga/tests/out/glsl/debug-printf.main.Compute.glsl @@ -0,0 +1,14 @@ +#version 310 es +#extension GL_EXT_debug_printf : enable + +precision highp float; +precision highp int; + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + + +void main() { + debugPrintfEXT("%d",42); + return; +} + diff --git a/naga/tests/out/hlsl/debug-printf-s.hlsl b/naga/tests/out/hlsl/debug-printf-s.hlsl new file mode 100644 index 0000000000..638bbf0d9a --- /dev/null +++ b/naga/tests/out/hlsl/debug-printf-s.hlsl @@ -0,0 +1,11 @@ +void main_1() +{ + printf("%d",42); + return; +} + +[numthreads(1, 1, 1)] +void main() +{ + main_1(); +} diff --git a/naga/tests/out/hlsl/debug-printf-s.ron b/naga/tests/out/hlsl/debug-printf-s.ron new file mode 100644 index 0000000000..a07b03300b --- /dev/null +++ b/naga/tests/out/hlsl/debug-printf-s.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_5_1", + ), + ], +) diff --git a/naga/tests/out/spv/debug-printf.spvasm b/naga/tests/out/spv/debug-printf.spvasm new file mode 100644 index 0000000000..d0d9b89abb --- /dev/null +++ b/naga/tests/out/spv/debug-printf.spvasm @@ -0,0 +1,23 @@ +; SPIR-V +; Version: 1.1 +; Generator: rspirv +; Bound: 12 +OpCapability Shader +OpExtension "SPV_KHR_non_semantic_info" +%1 = OpExtInstImport "GLSL.std.450" +%11 = OpExtInstImport "NonSemantic.DebugPrintf" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %4 "main" +OpExecutionMode %4 LocalSize 1 1 1 +%9 = OpString "%d" +%2 = OpTypeVoid +%5 = OpTypeFunction %2 +%6 = OpTypeInt 32 1 +%7 = OpConstant %6 42 +%4 = OpFunction %2 None %5 +%3 = OpLabel +OpBranch %8 +%8 = OpLabel +%10 = OpExtInst %2 %11 1 %9 %7 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/debug-printf-s.wgsl b/naga/tests/out/wgsl/debug-printf-s.wgsl new file mode 100644 index 0000000000..3e5f740ce5 --- /dev/null +++ b/naga/tests/out/wgsl/debug-printf-s.wgsl @@ -0,0 +1,9 @@ +fn main_1() { + debugPrintf("%d",42i); + return; +} + +@compute @workgroup_size(1, 1, 1) +fn main() { + main_1(); +} diff --git a/naga/tests/out/wgsl/debug-printf.wgsl b/naga/tests/out/wgsl/debug-printf.wgsl new file mode 100644 index 0000000000..d146290537 --- /dev/null +++ b/naga/tests/out/wgsl/debug-printf.wgsl @@ -0,0 +1,5 @@ +@compute @workgroup_size(1, 1, 1) +fn main() { + debugPrintf("%d",42i); + return; +} diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index 6b934de55b..1e4fe5fd33 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -802,6 +802,10 @@ fn convert_wgsl() { "abstract-types-operators", Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL, ), + ( + "debug-printf", + Targets::WGSL | Targets::GLSL | Targets::SPIRV | Targets::HLSL, + ), ]; for &(name, targets) in inputs.iter() { @@ -881,6 +885,11 @@ fn convert_spv_all() { true, Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ); + convert_spv( + "debug-printf-s", + false, + Targets::GLSL | Targets::HLSL | Targets::WGSL, + ); } #[cfg(feature = "glsl-in")] From a841cd494caee44426557e9a3164f238d613f690 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Mon, 16 Oct 2023 21:10:12 -0400 Subject: [PATCH 08/17] debug_printf: Disable hlsl-out test DXC doesn't support printf when targeting dxil, disable the test to allow CI to pass. --- naga/tests/snapshots.rs | 8 ++------ 1 file changed, 2 insertions(+), 6 deletions(-) diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index 1e4fe5fd33..f183adbf2e 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -804,7 +804,7 @@ fn convert_wgsl() { ), ( "debug-printf", - Targets::WGSL | Targets::GLSL | Targets::SPIRV | Targets::HLSL, + Targets::WGSL | Targets::GLSL | Targets::SPIRV, ), ]; @@ -885,11 +885,7 @@ fn convert_spv_all() { true, Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ); - convert_spv( - "debug-printf-s", - false, - Targets::GLSL | Targets::HLSL | Targets::WGSL, - ); + convert_spv("debug-printf-s", false, Targets::GLSL | Targets::WGSL); } #[cfg(feature = "glsl-in")] From ba63f1f9fd8fcdf17b736d59631e24dc0806c367 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Fri, 20 Oct 2023 21:04:24 -0400 Subject: [PATCH 09/17] debug_printf: Add DEBUG_PRINTF capability --- naga/src/valid/function.rs | 11 +++++++++++ naga/src/valid/mod.rs | 2 ++ naga/tests/in/debug-printf.param.ron | 3 +++ naga/tests/in/spv/debug-printf-s.param.ron | 3 +++ 4 files changed, 19 insertions(+) create mode 100644 naga/tests/in/debug-printf.param.ron create mode 100644 naga/tests/in/spv/debug-printf-s.param.ron diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index 85de8f31e5..ca7fb2592e 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -140,6 +140,8 @@ pub enum FunctionError { index: usize, source: ExpressionError, }, + #[error("Capability {0:?} is required")] + MissingCapability(super::Capabilities), #[error( "Required uniformity of control flow for {0:?} in {1:?} is not fulfilled because of {2:?}" )] @@ -910,6 +912,15 @@ impl super::Validator { } } S::DebugPrintf { ref arguments, .. } => { + if !self + .capabilities + .contains(super::Capabilities::DEBUG_PRINTF) + { + return Err(FunctionError::MissingCapability( + super::Capabilities::DEBUG_PRINTF, + ) + .with_span_static(span, "debugPrintf")); + } for (index, &expr) in arguments.iter().enumerate() { context .resolve_type_impl(expr, &self.valid_expression_set) diff --git a/naga/src/valid/mod.rs b/naga/src/valid/mod.rs index 70a4d39d2a..b5a3fca5bd 100644 --- a/naga/src/valid/mod.rs +++ b/naga/src/valid/mod.rs @@ -108,6 +108,8 @@ bitflags::bitflags! { const DUAL_SOURCE_BLENDING = 0x2000; /// Support for arrayed cube textures. const CUBE_ARRAY_TEXTURES = 0x4000; + /// Support for `debugPrintf` + const DEBUG_PRINTF = 0x8000; } } diff --git a/naga/tests/in/debug-printf.param.ron b/naga/tests/in/debug-printf.param.ron new file mode 100644 index 0000000000..a0f37d5325 --- /dev/null +++ b/naga/tests/in/debug-printf.param.ron @@ -0,0 +1,3 @@ +( + god_mode: true, +) diff --git a/naga/tests/in/spv/debug-printf-s.param.ron b/naga/tests/in/spv/debug-printf-s.param.ron new file mode 100644 index 0000000000..a0f37d5325 --- /dev/null +++ b/naga/tests/in/spv/debug-printf-s.param.ron @@ -0,0 +1,3 @@ +( + god_mode: true, +) From 39fd57725b43c80c75966c4ccb969616cd8c96c2 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Wed, 25 Oct 2023 11:42:46 -0700 Subject: [PATCH 10/17] debug_printf: Add DEBUG_PRINTF wgpu feature --- wgpu-core/src/device/resource.rs | 4 ++++ wgpu-hal/src/vulkan/adapter.rs | 10 ++++++++++ wgpu-types/src/lib.rs | 12 ++++++++++-- 3 files changed, 24 insertions(+), 2 deletions(-) diff --git a/wgpu-core/src/device/resource.rs b/wgpu-core/src/device/resource.rs index 06aedc00f6..ab2cfa9a5c 100644 --- a/wgpu-core/src/device/resource.rs +++ b/wgpu-core/src/device/resource.rs @@ -1431,6 +1431,10 @@ impl Device { .flags .contains(wgt::DownlevelFlags::CUBE_ARRAY_TEXTURES), ); + caps.set( + Caps::DEBUG_PRINTF, + self.features.contains(wgt::Features::DEBUG_PRINTF), + ); let debug_source = if self.instance_flags.contains(wgt::InstanceFlags::DEBUG) && !source.is_empty() { diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs index 85e620d23c..2fe9ff2e44 100644 --- a/wgpu-hal/src/vulkan/adapter.rs +++ b/wgpu-hal/src/vulkan/adapter.rs @@ -482,6 +482,11 @@ impl PhysicalDeviceFeatures { caps.supports_extension(vk::ExtConservativeRasterizationFn::name()), ); + features.set( + F::DEBUG_PRINTF, + caps.supports_extension(vk::KhrShaderNonSemanticInfoFn::name()), + ); + let intel_windows = caps.properties.vendor_id == db::intel::VENDOR && cfg!(windows); if let Some(ref descriptor_indexing) = self.descriptor_indexing { @@ -798,6 +803,11 @@ impl PhysicalDeviceCapabilities { extensions.push(vk::ExtConservativeRasterizationFn::name()); } + // Require `VK_KHR_shader_non_semantic_info` if the associated feature was requested + if requested_features.contains(wgt::Features::DEBUG_PRINTF) { + extensions.push(vk::KhrShaderNonSemanticInfoFn::name()); + } + // Require `VK_KHR_portability_subset` on macOS/iOS #[cfg(any(target_os = "macos", target_os = "ios"))] extensions.push(vk::KhrPortabilitySubsetFn::name()); diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs index fa2a8df5f8..6b4b9d1aff 100644 --- a/wgpu-types/src/lib.rs +++ b/wgpu-types/src/lib.rs @@ -793,10 +793,18 @@ bitflags::bitflags! { /// This is a native-only feature. const RAY_TRACING_ACCELERATION_STRUCTURE = 1 << 56; - // 57 available - // Shader: + /// Enables support for debugPrintf in WGSL shaders. + /// + /// Supported Platforms: + /// - DX11 (fxc only) + /// - Dx12 (fxc only) + /// - Vulkan + /// - OpenGL + /// + /// This is a native only feature + const DEBUG_PRINTF = 1 << 57; /// Allows for the creation of ray-tracing queries within shaders. /// /// Supported platforms: From f1ee02de3cacce6eb0feab9a29ca198bc6169599 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Wed, 25 Oct 2023 18:22:53 -0700 Subject: [PATCH 11/17] debug_printf: Add gpu test --- tests/tests/shader/debug_printf.rs | 51 ++++++++++++++++++++++++++++ tests/tests/shader/debug_printf.wgsl | 4 +++ tests/tests/shader/mod.rs | 1 + 3 files changed, 56 insertions(+) create mode 100644 tests/tests/shader/debug_printf.rs create mode 100644 tests/tests/shader/debug_printf.wgsl diff --git a/tests/tests/shader/debug_printf.rs b/tests/tests/shader/debug_printf.rs new file mode 100644 index 0000000000..4388f1551d --- /dev/null +++ b/tests/tests/shader/debug_printf.rs @@ -0,0 +1,51 @@ +use wgpu::{ + include_wgsl, CommandEncoderDescriptor, ComputePassDescriptor, ComputePipelineDescriptor, + Features, Limits, Maintain, PipelineLayoutDescriptor, +}; + +use wgpu_test::{gpu_test, GpuTestConfiguration, TestParameters}; + +#[gpu_test] +static DEBUG_PRINTF: GpuTestConfiguration = GpuTestConfiguration::new() + .parameters( + TestParameters::default() + .features(Features::DEBUG_PRINTF) + .limits(Limits::default()), + ) + .run_sync(|ctx| { + let pll = ctx + .device + .create_pipeline_layout(&PipelineLayoutDescriptor { + label: None, + bind_group_layouts: &[], + push_constant_ranges: &[], + }); + + let sm = ctx + .device + .create_shader_module(include_wgsl!("debug_printf.wgsl")); + + let pipeline = ctx + .device + .create_compute_pipeline(&ComputePipelineDescriptor { + label: Some("debugprintf"), + layout: Some(&pll), + module: &sm, + entry_point: "main", + }); + + // -- Run test -- + + let mut encoder = ctx + .device + .create_command_encoder(&CommandEncoderDescriptor::default()); + + { + let mut cpass = encoder.begin_compute_pass(&ComputePassDescriptor::default()); + cpass.set_pipeline(&pipeline); + } + + ctx.queue.submit(Some(encoder.finish())); + + ctx.device.poll(Maintain::Wait); + }); diff --git a/tests/tests/shader/debug_printf.wgsl b/tests/tests/shader/debug_printf.wgsl new file mode 100644 index 0000000000..d560bd7c08 --- /dev/null +++ b/tests/tests/shader/debug_printf.wgsl @@ -0,0 +1,4 @@ +@compute @workgroup_size(8,8,1) +fn main() { + debugPrintf("Hello world %d", 1); +} diff --git a/tests/tests/shader/mod.rs b/tests/tests/shader/mod.rs index 1a981971f7..a41684c9b2 100644 --- a/tests/tests/shader/mod.rs +++ b/tests/tests/shader/mod.rs @@ -15,6 +15,7 @@ use wgpu::{ use wgpu_test::TestingContext; +pub mod debug_printf; pub mod numeric_builtins; pub mod struct_layout; pub mod zero_init_workgroup_mem; From a05b30a1868ce1294424d6e6b23cd859fe6db612 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Wed, 1 Nov 2023 05:18:39 -0700 Subject: [PATCH 12/17] debug_printf: Add writer flags to conditionally emit/ignore debugPrintf statements --- naga/src/back/glsl/mod.rs | 16 ++++++-- naga/src/back/hlsl/mod.rs | 13 ++++++ naga/src/back/hlsl/writer.rs | 18 +++++---- naga/src/back/msl/writer.rs | 2 +- naga/src/back/spv/block.rs | 40 ++++++++++--------- naga/src/back/spv/mod.rs | 6 ++- naga/src/back/wgsl/writer.rs | 20 ++++++---- naga/tests/in/binding-arrays.param.ron | 1 + naga/tests/in/debug-printf.param.ron | 13 ++++++ naga/tests/in/interface.param.ron | 1 + naga/tests/in/push-constants.param.ron | 1 + naga/tests/in/skybox.param.ron | 1 + naga/tests/in/spv/debug-printf-s.param.ron | 13 ++++++ .../out/glsl/debug-printf-s.main.Compute.glsl | 7 +--- .../out/glsl/debug-printf.main.Compute.glsl | 7 +--- naga/tests/snapshots.rs | 12 ++++++ wgpu-hal/src/dx12/device.rs | 1 + wgpu-hal/src/vulkan/adapter.rs | 4 ++ 18 files changed, 125 insertions(+), 51 deletions(-) diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index 923fec5451..1a9a610363 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -252,6 +252,8 @@ bitflags::bitflags! { /// The variable gl_PointSize is intended for a shader to write the size of the point to be rasterized. It is measured in pixels. /// If gl_PointSize is not written to, its value is undefined in subsequent pipe stages. const FORCE_POINT_SIZE = 0x20; + /// Emit debug printf statements + const EMIT_DEBUG_PRINTF = 0x40; } } @@ -2383,10 +2385,16 @@ impl<'a, W: Write> Writer<'a, W> { ref format, ref arguments, } => { - write!(self.out, "{level}")?; - write!(self.out, "debugPrintfEXT(\"{format}\",")?; - self.write_slice(arguments, |this, _, arg| this.write_expr(*arg, ctx))?; - writeln!(self.out, ");")? + if self + .options + .writer_flags + .contains(WriterFlags::EMIT_DEBUG_PRINTF) + { + write!(self.out, "{level}")?; + write!(self.out, "debugPrintfEXT(\"{format}\",")?; + self.write_slice(arguments, |this, _, arg| this.write_expr(*arg, ctx))?; + writeln!(self.out, ");")? + } } } diff --git a/naga/src/back/hlsl/mod.rs b/naga/src/back/hlsl/mod.rs index 37ddbd3d67..6e9883e4e4 100644 --- a/naga/src/back/hlsl/mod.rs +++ b/naga/src/back/hlsl/mod.rs @@ -175,11 +175,23 @@ pub enum EntryPointError { MissingBinding(crate::ResourceBinding), } +bitflags::bitflags! { + #[cfg_attr(feature = "serialize", derive(serde::Serialize))] + #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))] + #[derive(Clone, Copy, Debug, Hash, PartialEq, Eq, Default)] + pub struct WriterFlags: u32 { + /// Emit debug printf statements + const EMIT_DEBUG_PRINTF = 0x1; + } +} + /// Configuration used in the [`Writer`]. #[derive(Clone, Debug, Hash, PartialEq, Eq)] #[cfg_attr(feature = "serialize", derive(serde::Serialize))] #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))] pub struct Options { + /// Configuration flags for the writer. + pub flags: WriterFlags, /// The hlsl shader model to be used pub shader_model: ShaderModel, /// Map of resources association to binding locations. @@ -198,6 +210,7 @@ pub struct Options { impl Default for Options { fn default() -> Self { Options { + flags: WriterFlags::empty(), shader_model: ShaderModel::V5_1, binding_map: BindingMap::default(), fake_missing_bindings: true, diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index da02ce0069..3711f4d9de 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/naga/src/back/hlsl/writer.rs @@ -4,7 +4,7 @@ use super::{ BackendResult, Error, Options, }; use crate::{ - back, + back::{self, hlsl::WriterFlags}, proc::{self, NameKey}, valid, Handle, Module, ScalarKind, ShaderStage, TypeInner, }; @@ -2004,15 +2004,17 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { ref format, ref arguments, } => { - write!(self.out, "{level}")?; - write!(self.out, "printf(\"{format}\",")?; - for (index, argument) in arguments.iter().enumerate() { - if index != 0 { - write!(self.out, ", ")?; + if self.options.flags.contains(WriterFlags::EMIT_DEBUG_PRINTF) { + write!(self.out, "{level}")?; + write!(self.out, "printf(\"{format}\",")?; + for (index, argument) in arguments.iter().enumerate() { + if index != 0 { + write!(self.out, ", ")?; + } + self.write_expr(module, *argument, func_ctx)?; } - self.write_expr(module, *argument, func_ctx)?; + writeln!(self.out, ");")? } - writeln!(self.out, ");")? } } diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index 7a6e0e2960..a40cbf0540 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -3041,7 +3041,7 @@ impl Writer { } } crate::Statement::DebugPrintf { .. } => { - return Err(Error::FeatureNotImplemented("debug printf".to_string())); + // metal doesn't provide a debug printf implementation } } } diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index e60ecc1890..bfc433a7be 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -2338,27 +2338,29 @@ impl<'w> BlockContext<'w> { ref format, ref arguments, } => { - self.writer.use_extension("SPV_KHR_non_semantic_info"); - let format_id = self.gen_id(); - self.writer - .strings - .push(Instruction::string(format, format_id)); - let id = self.gen_id(); + if self.writer.flags.contains(WriterFlags::EMIT_DEBUG_PRINTF) { + self.writer.use_extension("SPV_KHR_non_semantic_info"); + let format_id = self.gen_id(); + self.writer + .strings + .push(Instruction::string(format, format_id)); + let id = self.gen_id(); - self.temp_list.clear(); - self.temp_list.push(format_id); - for &argument in arguments { - self.temp_list.push(self.cached[argument]); - } + self.temp_list.clear(); + self.temp_list.push(format_id); + for &argument in arguments { + self.temp_list.push(self.cached[argument]); + } - let set_id = self.writer.extension_inst_import("NonSemantic.DebugPrintf"); - block.body.push(Instruction::ext_inst( - set_id, - 1, - self.writer.void_type, - id, - &self.temp_list, - )); + let set_id = self.writer.extension_inst_import("NonSemantic.DebugPrintf"); + block.body.push(Instruction::ext_inst( + set_id, + 1, + self.writer.void_type, + id, + &self.temp_list, + )); + } } } } diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index 1c2446d716..c177760fdf 100644 --- a/naga/src/back/spv/mod.rs +++ b/naga/src/back/spv/mod.rs @@ -631,7 +631,9 @@ pub struct Writer { } bitflags::bitflags! { - #[derive(Clone, Copy, Debug, Eq, PartialEq)] + #[cfg_attr(feature = "serialize", derive(serde::Serialize))] + #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))] + #[derive(Clone, Copy, Debug, Eq, PartialEq, Default)] pub struct WriterFlags: u32 { /// Include debug labels for everything. const DEBUG = 0x1; @@ -646,6 +648,8 @@ bitflags::bitflags! { const FORCE_POINT_SIZE = 0x8; /// Clamp `BuiltIn::FragDepth` output between 0 and 1. const CLAMP_FRAG_DEPTH = 0x10; + /// Emit debug printf statements + const EMIT_DEBUG_PRINTF = 0x20; } } diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs index cbd24ed7f7..01b7aa8170 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -54,10 +54,12 @@ enum Indirection { bitflags::bitflags! { #[cfg_attr(feature = "serialize", derive(serde::Serialize))] #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))] - #[derive(Clone, Copy, Debug, Eq, PartialEq)] + #[derive(Clone, Copy, Debug, Eq, PartialEq, Default)] pub struct WriterFlags: u32 { /// Always annotate the type information instead of inferring. const EXPLICIT_TYPES = 0x1; + /// Emit debug printf statements + const EMIT_DEBUG_PRINTF = 0x2; } } @@ -923,15 +925,17 @@ impl Writer { ref format, ref arguments, } => { - write!(self.out, "{level}")?; - write!(self.out, "debugPrintf(\"{format}\",")?; - for (index, &argument) in arguments.iter().enumerate() { - if index != 0 { - write!(self.out, ", ")?; + if self.flags.contains(WriterFlags::EMIT_DEBUG_PRINTF) { + write!(self.out, "{level}")?; + write!(self.out, "debugPrintf(\"{format}\",")?; + for (index, &argument) in arguments.iter().enumerate() { + if index != 0 { + write!(self.out, ", ")?; + } + self.write_expr(module, argument, func_ctx)?; } - self.write_expr(module, argument, func_ctx)?; + writeln!(self.out, ");")? } - writeln!(self.out, ");")? } } diff --git a/naga/tests/in/binding-arrays.param.ron b/naga/tests/in/binding-arrays.param.ron index 39d6c03664..652bebdf24 100644 --- a/naga/tests/in/binding-arrays.param.ron +++ b/naga/tests/in/binding-arrays.param.ron @@ -2,6 +2,7 @@ god_mode: true, hlsl: ( shader_model: V5_1, + flags: (""), binding_map: { (group: 0, binding: 0): (space: 0, register: 0, binding_array_size: Some(10)), (group: 0, binding: 1): (space: 1, register: 0), diff --git a/naga/tests/in/debug-printf.param.ron b/naga/tests/in/debug-printf.param.ron index a0f37d5325..6167f322a4 100644 --- a/naga/tests/in/debug-printf.param.ron +++ b/naga/tests/in/debug-printf.param.ron @@ -1,3 +1,16 @@ ( god_mode: true, + spv: ( + version: (1, 1), + emit_debug_printf: true, + ), + wgsl: ( + emit_debug_printf: true, + ), + glsl: ( + version: Desktop(450), + writer_flags: ("EMIT_DEBUG_PRINTF"), + binding_map: {}, + zero_initialize_workgroup_memory: true, + ), ) diff --git a/naga/tests/in/interface.param.ron b/naga/tests/in/interface.param.ron index 4d85661767..75adcc8e2e 100644 --- a/naga/tests/in/interface.param.ron +++ b/naga/tests/in/interface.param.ron @@ -9,6 +9,7 @@ ), hlsl: ( shader_model: V5_1, + flags: (""), binding_map: {}, fake_missing_bindings: false, special_constants_binding: Some((space: 1, register: 0)), diff --git a/naga/tests/in/push-constants.param.ron b/naga/tests/in/push-constants.param.ron index 083d028bbf..8f89361012 100644 --- a/naga/tests/in/push-constants.param.ron +++ b/naga/tests/in/push-constants.param.ron @@ -11,6 +11,7 @@ ), hlsl: ( shader_model: V5_1, + flags: (""), binding_map: {}, fake_missing_bindings: true, special_constants_binding: Some((space: 1, register: 0)), diff --git a/naga/tests/in/skybox.param.ron b/naga/tests/in/skybox.param.ron index 4d7fdf7347..a9e8a2ccee 100644 --- a/naga/tests/in/skybox.param.ron +++ b/naga/tests/in/skybox.param.ron @@ -51,6 +51,7 @@ ), hlsl: ( shader_model: V5_1, + flags: (""), binding_map: { (group: 0, binding: 0): (space: 0, register: 0), (group: 0, binding: 1): (space: 0, register: 0), diff --git a/naga/tests/in/spv/debug-printf-s.param.ron b/naga/tests/in/spv/debug-printf-s.param.ron index a0f37d5325..6167f322a4 100644 --- a/naga/tests/in/spv/debug-printf-s.param.ron +++ b/naga/tests/in/spv/debug-printf-s.param.ron @@ -1,3 +1,16 @@ ( god_mode: true, + spv: ( + version: (1, 1), + emit_debug_printf: true, + ), + wgsl: ( + emit_debug_printf: true, + ), + glsl: ( + version: Desktop(450), + writer_flags: ("EMIT_DEBUG_PRINTF"), + binding_map: {}, + zero_initialize_workgroup_memory: true, + ), ) diff --git a/naga/tests/out/glsl/debug-printf-s.main.Compute.glsl b/naga/tests/out/glsl/debug-printf-s.main.Compute.glsl index 2cd87e2b25..7a21763d6f 100644 --- a/naga/tests/out/glsl/debug-printf-s.main.Compute.glsl +++ b/naga/tests/out/glsl/debug-printf-s.main.Compute.glsl @@ -1,9 +1,6 @@ -#version 310 es +#version 450 core +#extension GL_ARB_compute_shader : require #extension GL_EXT_debug_printf : enable - -precision highp float; -precision highp int; - layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/naga/tests/out/glsl/debug-printf.main.Compute.glsl b/naga/tests/out/glsl/debug-printf.main.Compute.glsl index 0923ee6cdc..58abf6dd16 100644 --- a/naga/tests/out/glsl/debug-printf.main.Compute.glsl +++ b/naga/tests/out/glsl/debug-printf.main.Compute.glsl @@ -1,9 +1,6 @@ -#version 310 es +#version 450 core +#extension GL_ARB_compute_shader : require #extension GL_EXT_debug_printf : enable - -precision highp float; -precision highp int; - layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index f183adbf2e..ac122a3a6c 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -49,6 +49,8 @@ struct SpirvOutParameters { #[serde(default)] separate_entry_points: bool, #[serde(default)] + emit_debug_printf: bool, + #[serde(default)] #[cfg(all(feature = "deserialize", feature = "spv-out"))] binding_map: naga::back::spv::BindingMap, } @@ -57,6 +59,8 @@ struct SpirvOutParameters { struct WgslOutParameters { #[serde(default)] explicit_types: bool, + #[serde(default)] + emit_debug_printf: bool, } #[derive(Default, serde::Deserialize)] @@ -408,6 +412,10 @@ fn write_output_spv( ); flags.set(spv::WriterFlags::FORCE_POINT_SIZE, params.force_point_size); flags.set(spv::WriterFlags::CLAMP_FRAG_DEPTH, params.clamp_frag_depth); + flags.set( + spv::WriterFlags::EMIT_DEBUG_PRINTF, + params.emit_debug_printf, + ); let options = spv::Options { lang_version: (params.version.0, params.version.1), @@ -592,6 +600,10 @@ fn write_output_wgsl( let mut flags = wgsl::WriterFlags::empty(); flags.set(wgsl::WriterFlags::EXPLICIT_TYPES, params.explicit_types); + flags.set( + wgsl::WriterFlags::EMIT_DEBUG_PRINTF, + params.emit_debug_printf, + ); let string = wgsl::write_string(module, info, flags).expect("WGSL write failed"); diff --git a/wgpu-hal/src/dx12/device.rs b/wgpu-hal/src/dx12/device.rs index bb128b2a6d..cb55657aac 100644 --- a/wgpu-hal/src/dx12/device.rs +++ b/wgpu-hal/src/dx12/device.rs @@ -1070,6 +1070,7 @@ impl crate::Device for super::Device { // FXC doesn't support SM 6.0 None => hlsl::ShaderModel::V5_1, }, + flags: hlsl::WriterFlags::default(), binding_map, fake_missing_bindings: false, special_constants_binding, diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs index 2fe9ff2e44..27049e03bb 100644 --- a/wgpu-hal/src/vulkan/adapter.rs +++ b/wgpu-hal/src/vulkan/adapter.rs @@ -1470,6 +1470,10 @@ impl super::Adapter { // But this requires cloning the `spv::Options` struct, which has heap allocations. true, // could check `super::Workarounds::SEPARATE_ENTRY_POINTS` ); + flags.set( + spv::WriterFlags::EMIT_DEBUG_PRINTF, + features.contains(wgt::Features::DEBUG_PRINTF), + ); spv::Options { lang_version: (1, 0), flags, From 65d1b500ca58c3d4589ce244cd6b5b302344c8ef Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Wed, 1 Nov 2023 05:20:20 -0700 Subject: [PATCH 13/17] debug_printf: Enable debug_printf validation feature --- wgpu-hal/src/vulkan/instance.rs | 34 ++++++++++++++++++++++++++++++++- 1 file changed, 33 insertions(+), 1 deletion(-) diff --git a/wgpu-hal/src/vulkan/instance.rs b/wgpu-hal/src/vulkan/instance.rs index 1f0159413f..3d81a8e917 100644 --- a/wgpu-hal/src/vulkan/instance.rs +++ b/wgpu-hal/src/vulkan/instance.rs @@ -620,7 +620,7 @@ impl crate::Instance for super::Instance { }, ); - let extensions = Self::desired_extensions(&entry, instance_api_version, desc.flags)?; + let mut extensions = Self::desired_extensions(&entry, instance_api_version, desc.flags)?; let instance_layers = { profiling::scope!("vkEnumerateInstanceLayerProperties"); @@ -653,6 +653,7 @@ impl crate::Instance for super::Instance { // Request validation layer if asked. let mut debug_utils = None; + let mut validation_features = None; if desc.flags.intersects(wgt::InstanceFlags::VALIDATION) { let validation_layer_name = CStr::from_bytes_with_nul(b"VK_LAYER_KHRONOS_validation\0").unwrap(); @@ -698,6 +699,33 @@ impl crate::Instance for super::Instance { debug_utils = Some((create_info, vk_create_info)); } + + let validation_features_name = vk::ExtValidationFeaturesFn::name(); + match entry.enumerate_instance_extension_properties(Some(validation_layer_name)) { + Ok(validation_extensions) => { + if validation_extensions.iter().any(|inst_ext| { + cstr_from_bytes_until_nul(&inst_ext.extension_name) + == Some(validation_features_name) + }) { + extensions.push(validation_features_name); + validation_features = Some( + vk::ValidationFeaturesEXT::builder().enabled_validation_features( + &[vk::ValidationFeatureEnableEXT::DEBUG_PRINTF], + ), + ); + } else { + log::info!( + "Unable to find validation layer extension {}, not enabling DEBUG_PRINTF", + validation_features_name.to_string_lossy() + ) + } + } + Err(e) => { + log::warn!( + "enumerate_instance_extension_properties() failed for validation layer: {:?}", e + ) + } + } } else { log::warn!( "InstanceFlags::VALIDATION requested, but unable to find layer: {}", @@ -752,6 +780,10 @@ impl crate::Instance for super::Instance { .enabled_layer_names(&str_pointers[..layers.len()]) .enabled_extension_names(&str_pointers[layers.len()..]); + if let Some(validation_features) = validation_features.as_mut() { + create_info = create_info.push_next(validation_features); + } + if let Some(&mut (_, ref mut vk_create_info)) = debug_utils.as_mut() { create_info = create_info.push_next(vk_create_info); } From 11df7f3d21ef88f8c7b2cf0d46df5228bfa673f1 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Wed, 15 Nov 2023 18:48:40 -0800 Subject: [PATCH 14/17] debug_printf: fix feature logic for DX12 --- wgpu-hal/src/dx12/adapter.rs | 4 ++++ wgpu-hal/src/dx12/device.rs | 6 +++++- wgpu-types/src/lib.rs | 2 +- 3 files changed, 10 insertions(+), 2 deletions(-) diff --git a/wgpu-hal/src/dx12/adapter.rs b/wgpu-hal/src/dx12/adapter.rs index f6027014d2..9a7598b685 100644 --- a/wgpu-hal/src/dx12/adapter.rs +++ b/wgpu-hal/src/dx12/adapter.rs @@ -297,6 +297,10 @@ impl super::Adapter { // float32-filterable should always be available on d3d12 features.set(wgt::Features::FLOAT32_FILTERABLE, true); + // DXC lacks support for printf, so only enable it on FXC + // see https://github.com/microsoft/DirectXShaderCompiler/issues/357 + features.set(wgt::Features::DEBUG_PRINTF, dxc_container.is_none()); + // TODO: Determine if IPresentationManager is supported let presentation_timer = auxil::dxgi::time::PresentationTimer::new_dxgi(); diff --git a/wgpu-hal/src/dx12/device.rs b/wgpu-hal/src/dx12/device.rs index cb55657aac..4763cd830e 100644 --- a/wgpu-hal/src/dx12/device.rs +++ b/wgpu-hal/src/dx12/device.rs @@ -1070,7 +1070,11 @@ impl crate::Device for super::Device { // FXC doesn't support SM 6.0 None => hlsl::ShaderModel::V5_1, }, - flags: hlsl::WriterFlags::default(), + flags: match self.dxc_container { + // DXC doesn't support printf: https://github.com/microsoft/DirectXShaderCompiler/issues/357 + Some(_) => hlsl::WriterFlags::empty(), + None => hlsl::WriterFlags::EMIT_DEBUG_PRINTF, + }, binding_map, fake_missing_bindings: false, special_constants_binding, diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs index 6b4b9d1aff..db8c238012 100644 --- a/wgpu-types/src/lib.rs +++ b/wgpu-types/src/lib.rs @@ -799,7 +799,7 @@ bitflags::bitflags! { /// /// Supported Platforms: /// - DX11 (fxc only) - /// - Dx12 (fxc only) + /// - DX12 (fxc only) /// - Vulkan /// - OpenGL /// From 9ee15532abd785d2a9a0ace217169c73d40edf29 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Wed, 15 Nov 2023 19:11:46 -0800 Subject: [PATCH 15/17] debug_printf: Flatten validation layer extension logic --- wgpu-hal/src/vulkan/instance.rs | 46 +++++++++++++++++---------------- 1 file changed, 24 insertions(+), 22 deletions(-) diff --git a/wgpu-hal/src/vulkan/instance.rs b/wgpu-hal/src/vulkan/instance.rs index 3d81a8e917..104e3900ff 100644 --- a/wgpu-hal/src/vulkan/instance.rs +++ b/wgpu-hal/src/vulkan/instance.rs @@ -701,30 +701,32 @@ impl crate::Instance for super::Instance { } let validation_features_name = vk::ExtValidationFeaturesFn::name(); - match entry.enumerate_instance_extension_properties(Some(validation_layer_name)) { - Ok(validation_extensions) => { - if validation_extensions.iter().any(|inst_ext| { - cstr_from_bytes_until_nul(&inst_ext.extension_name) - == Some(validation_features_name) - }) { - extensions.push(validation_features_name); - validation_features = Some( - vk::ValidationFeaturesEXT::builder().enabled_validation_features( - &[vk::ValidationFeatureEnableEXT::DEBUG_PRINTF], - ), - ); - } else { - log::info!( - "Unable to find validation layer extension {}, not enabling DEBUG_PRINTF", - validation_features_name.to_string_lossy() - ) + 'validation_exts: { + let validation_extensions = match entry + .enumerate_instance_extension_properties(Some(validation_layer_name)) + { + Ok(e) => e, + Err(e) => { + log::warn!( "enumerate_instance_extension_properties() failed for validation layer: {:?}", e ); + break 'validation_exts; } + }; + + let extension_found = validation_extensions.iter().any(|inst_ext| { + cstr_from_bytes_until_nul(&inst_ext.extension_name) + == Some(validation_features_name) + }); + if !extension_found { + log::info!( "Unable to find validation layer extension {}, not enabling DEBUG_PRINTF", validation_features_name.to_string_lossy() ); + break 'validation_exts; } - Err(e) => { - log::warn!( - "enumerate_instance_extension_properties() failed for validation layer: {:?}", e - ) - } + + extensions.push(validation_features_name); + validation_features = Some( + vk::ValidationFeaturesEXT::builder().enabled_validation_features(&[ + vk::ValidationFeatureEnableEXT::DEBUG_PRINTF, + ]), + ); } } else { log::warn!( From b01083453f32b952fa51d98ab14d0d76a2882d56 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Thu, 28 Dec 2023 07:03:22 -0800 Subject: [PATCH 16/17] debug_printf: refactor debug_printf ext_inst handling --- naga/src/front/spv/ext_inst.rs | 23 ++++++++++++++++------- 1 file changed, 16 insertions(+), 7 deletions(-) diff --git a/naga/src/front/spv/ext_inst.rs b/naga/src/front/spv/ext_inst.rs index 388702d4bc..13f99ecd42 100644 --- a/naga/src/front/spv/ext_inst.rs +++ b/naga/src/front/spv/ext_inst.rs @@ -40,15 +40,19 @@ impl> super::Frontend { "GLSL.std.450" => self.parse_ext_inst_glsl_std( ext_name, inst, ext_inst, span, ctx, emitter, block, block_id, body_idx, ), - "NonSemantic.DebugPrintf" if ext_inst.inst_id == 1 => { - self.parse_ext_inst_debug_printf(inst, span, ctx, emitter, block, body_idx) - } + "NonSemantic.DebugPrintf" => self.parse_ext_inst_debug_printf( + ext_name, inst, ext_inst, span, ctx, emitter, block, body_idx, + ), + _ => Err(Error::UnsupportedExtInst(ext_inst.inst_id, ext_name)), } } + #[allow(clippy::too_many_arguments)] fn parse_ext_inst_debug_printf( &mut self, + ext_name: &'static str, inst: super::Instruction, + ext_inst: ExtInst, span: crate::Span, ctx: &mut super::BlockContext, emitter: &mut crate::proc::Emitter, @@ -56,6 +60,11 @@ impl> super::Frontend { body_idx: usize, ) -> Result<(), Error> { let base_wc = 5; + + if ext_inst.inst_id != 1 { + return Err(Error::UnsupportedExtInst(ext_inst.inst_id, ext_name)); + } + inst.expect_at_least(base_wc + 1)?; let format_id = self.next()?; let format = self.strings.lookup(format_id)?.clone(); @@ -77,7 +86,7 @@ impl> super::Frontend { #[allow(clippy::too_many_arguments)] fn parse_ext_inst_glsl_std( &mut self, - set_name: &'static str, + ext_name: &'static str, inst: super::Instruction, ext_inst: ExtInst, span: crate::Span, @@ -93,7 +102,7 @@ impl> super::Frontend { let base_wc = 5; let gl_op = Glo::from_u32(ext_inst.inst_id) - .ok_or(Error::UnsupportedExtInst(ext_inst.inst_id, set_name))?; + .ok_or(Error::UnsupportedExtInst(ext_inst.inst_id, ext_name))?; let fun = match gl_op { Glo::Round => Mf::Round, @@ -159,7 +168,7 @@ impl> super::Frontend { Glo::FindUMsb | Glo::FindSMsb => Mf::FindMsb, // TODO: https://github.com/gfx-rs/naga/issues/2526 Glo::Modf | Glo::Frexp => { - return Err(Error::UnsupportedExtInst(ext_inst.inst_id, set_name)) + return Err(Error::UnsupportedExtInst(ext_inst.inst_id, ext_name)) } Glo::IMix | Glo::PackDouble2x32 @@ -167,7 +176,7 @@ impl> super::Frontend { | Glo::InterpolateAtCentroid | Glo::InterpolateAtSample | Glo::InterpolateAtOffset => { - return Err(Error::UnsupportedExtInst(ext_inst.inst_id, set_name)) + return Err(Error::UnsupportedExtInst(ext_inst.inst_id, ext_name)) } }; From 3a78dcaf2bd0f51e95f63fdf4995b40804c93e28 Mon Sep 17 00:00:00 2001 From: Jacob Hughes Date: Wed, 24 Jan 2024 22:24:06 -0500 Subject: [PATCH 17/17] debug_printf: Skip naga tests on DXC --- naga/hlsl-snapshots/src/lib.rs | 1 + naga/tests/out/hlsl/access.ron | 3 ++ naga/tests/out/hlsl/array-in-ctor.ron | 1 + naga/tests/out/hlsl/atomicOps.ron | 1 + naga/tests/out/hlsl/binding-arrays.ron | 1 + naga/tests/out/hlsl/bitcast.ron | 1 + naga/tests/out/hlsl/bits.ron | 1 + naga/tests/out/hlsl/boids.ron | 1 + naga/tests/out/hlsl/break-if.ron | 1 + naga/tests/out/hlsl/collatz.ron | 1 + naga/tests/out/hlsl/const-exprs.ron | 1 + naga/tests/out/hlsl/constructors.ron | 1 + naga/tests/out/hlsl/control-flow.ron | 1 + naga/tests/out/hlsl/debug-printf-s.hlsl | 1 - naga/tests/out/hlsl/debug-printf-s.ron | 1 + naga/tests/out/hlsl/debug-printf.hlsl | 5 +++ naga/tests/out/hlsl/debug-printf.ron | 13 +++++++ naga/tests/out/hlsl/do-while.ron | 1 + naga/tests/out/hlsl/dualsource.ron | 1 + naga/tests/out/hlsl/empty-global-name.ron | 1 + naga/tests/out/hlsl/empty.ron | 1 + naga/tests/out/hlsl/f64.ron | 1 + naga/tests/out/hlsl/fragment-output.ron | 2 + naga/tests/out/hlsl/functions.ron | 1 + naga/tests/out/hlsl/globals.ron | 1 + naga/tests/out/hlsl/hlsl-keyword.ron | 1 + naga/tests/out/hlsl/image.ron | 8 ++++ naga/tests/out/hlsl/interface.ron | 4 ++ naga/tests/out/hlsl/interpolate.ron | 2 + .../hlsl/inv-hyperbolic-trig-functions.ron | 1 + naga/tests/out/hlsl/math-functions.ron | 1 + naga/tests/out/hlsl/operators.ron | 1 + naga/tests/out/hlsl/padding.ron | 1 + naga/tests/out/hlsl/push-constants.ron | 2 + naga/tests/out/hlsl/quad-vert.ron | 1 + naga/tests/out/hlsl/quad.ron | 3 ++ naga/tests/out/hlsl/shadow.ron | 3 ++ naga/tests/out/hlsl/skybox.ron | 2 + naga/tests/out/hlsl/standard.ron | 1 + naga/tests/out/hlsl/struct-layout.ron | 6 +++ naga/tests/out/hlsl/texture-arg.ron | 1 + .../tests/out/hlsl/workgroup-uniform-load.ron | 1 + naga/tests/out/hlsl/workgroup-var-init.ron | 1 + naga/tests/snapshots.rs | 12 ++++-- naga/xtask/src/validate.rs | 39 ++++++++++++------- 45 files changed, 116 insertions(+), 19 deletions(-) create mode 100644 naga/tests/out/hlsl/debug-printf.hlsl create mode 100644 naga/tests/out/hlsl/debug-printf.ron diff --git a/naga/hlsl-snapshots/src/lib.rs b/naga/hlsl-snapshots/src/lib.rs index 616aa73f01..b67ee1c1f6 100644 --- a/naga/hlsl-snapshots/src/lib.rs +++ b/naga/hlsl-snapshots/src/lib.rs @@ -94,4 +94,5 @@ pub struct ConfigItem { /// See also /// . pub target_profile: String, + pub debug_printf: bool, } diff --git a/naga/tests/out/hlsl/access.ron b/naga/tests/out/hlsl/access.ron index 73c9e44448..192821730c 100644 --- a/naga/tests/out/hlsl/access.ron +++ b/naga/tests/out/hlsl/access.ron @@ -3,18 +3,21 @@ ( entry_point:"foo_vert", target_profile:"vs_5_1", + debug_printf:false, ), ], fragment:[ ( entry_point:"foo_frag", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ ( entry_point:"assign_through_ptr", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/array-in-ctor.ron b/naga/tests/out/hlsl/array-in-ctor.ron index 5c261e59b2..5d12c87175 100644 --- a/naga/tests/out/hlsl/array-in-ctor.ron +++ b/naga/tests/out/hlsl/array-in-ctor.ron @@ -7,6 +7,7 @@ ( entry_point:"cs_main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/atomicOps.ron b/naga/tests/out/hlsl/atomicOps.ron index 5c261e59b2..5d12c87175 100644 --- a/naga/tests/out/hlsl/atomicOps.ron +++ b/naga/tests/out/hlsl/atomicOps.ron @@ -7,6 +7,7 @@ ( entry_point:"cs_main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/binding-arrays.ron b/naga/tests/out/hlsl/binding-arrays.ron index 341a4c528e..af020fcb28 100644 --- a/naga/tests/out/hlsl/binding-arrays.ron +++ b/naga/tests/out/hlsl/binding-arrays.ron @@ -5,6 +5,7 @@ ( entry_point:"main", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/bitcast.ron b/naga/tests/out/hlsl/bitcast.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/bitcast.ron +++ b/naga/tests/out/hlsl/bitcast.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/bits.ron b/naga/tests/out/hlsl/bits.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/bits.ron +++ b/naga/tests/out/hlsl/bits.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/boids.ron b/naga/tests/out/hlsl/boids.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/boids.ron +++ b/naga/tests/out/hlsl/boids.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/break-if.ron b/naga/tests/out/hlsl/break-if.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/break-if.ron +++ b/naga/tests/out/hlsl/break-if.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/collatz.ron b/naga/tests/out/hlsl/collatz.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/collatz.ron +++ b/naga/tests/out/hlsl/collatz.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/const-exprs.ron b/naga/tests/out/hlsl/const-exprs.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/const-exprs.ron +++ b/naga/tests/out/hlsl/const-exprs.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/constructors.ron b/naga/tests/out/hlsl/constructors.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/constructors.ron +++ b/naga/tests/out/hlsl/constructors.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/control-flow.ron b/naga/tests/out/hlsl/control-flow.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/control-flow.ron +++ b/naga/tests/out/hlsl/control-flow.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/debug-printf-s.hlsl b/naga/tests/out/hlsl/debug-printf-s.hlsl index 638bbf0d9a..a7ee013445 100644 --- a/naga/tests/out/hlsl/debug-printf-s.hlsl +++ b/naga/tests/out/hlsl/debug-printf-s.hlsl @@ -1,6 +1,5 @@ void main_1() { - printf("%d",42); return; } diff --git a/naga/tests/out/hlsl/debug-printf-s.ron b/naga/tests/out/hlsl/debug-printf-s.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/debug-printf-s.ron +++ b/naga/tests/out/hlsl/debug-printf-s.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/debug-printf.hlsl b/naga/tests/out/hlsl/debug-printf.hlsl new file mode 100644 index 0000000000..e79d36b9c8 --- /dev/null +++ b/naga/tests/out/hlsl/debug-printf.hlsl @@ -0,0 +1,5 @@ +[numthreads(1, 1, 1)] +void main() +{ + return; +} diff --git a/naga/tests/out/hlsl/debug-printf.ron b/naga/tests/out/hlsl/debug-printf.ron new file mode 100644 index 0000000000..a3f6571767 --- /dev/null +++ b/naga/tests/out/hlsl/debug-printf.ron @@ -0,0 +1,13 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"main", + target_profile:"cs_5_1", + debug_printf:false, + ), + ], +) diff --git a/naga/tests/out/hlsl/do-while.ron b/naga/tests/out/hlsl/do-while.ron index 341a4c528e..af020fcb28 100644 --- a/naga/tests/out/hlsl/do-while.ron +++ b/naga/tests/out/hlsl/do-while.ron @@ -5,6 +5,7 @@ ( entry_point:"main", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/dualsource.ron b/naga/tests/out/hlsl/dualsource.ron index 341a4c528e..af020fcb28 100644 --- a/naga/tests/out/hlsl/dualsource.ron +++ b/naga/tests/out/hlsl/dualsource.ron @@ -5,6 +5,7 @@ ( entry_point:"main", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/empty-global-name.ron b/naga/tests/out/hlsl/empty-global-name.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/empty-global-name.ron +++ b/naga/tests/out/hlsl/empty-global-name.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/empty.ron b/naga/tests/out/hlsl/empty.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/empty.ron +++ b/naga/tests/out/hlsl/empty.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/f64.ron b/naga/tests/out/hlsl/f64.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/f64.ron +++ b/naga/tests/out/hlsl/f64.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/fragment-output.ron b/naga/tests/out/hlsl/fragment-output.ron index 9dfaf7393b..e894433476 100644 --- a/naga/tests/out/hlsl/fragment-output.ron +++ b/naga/tests/out/hlsl/fragment-output.ron @@ -5,10 +5,12 @@ ( entry_point:"main_vec4vec3_", target_profile:"ps_5_1", + debug_printf:false, ), ( entry_point:"main_vec2scalar", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/functions.ron b/naga/tests/out/hlsl/functions.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/functions.ron +++ b/naga/tests/out/hlsl/functions.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/globals.ron b/naga/tests/out/hlsl/globals.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/globals.ron +++ b/naga/tests/out/hlsl/globals.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/hlsl-keyword.ron b/naga/tests/out/hlsl/hlsl-keyword.ron index eac1b945d2..a21e73bcf7 100644 --- a/naga/tests/out/hlsl/hlsl-keyword.ron +++ b/naga/tests/out/hlsl/hlsl-keyword.ron @@ -5,6 +5,7 @@ ( entry_point:"fs_main", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/image.ron b/naga/tests/out/hlsl/image.ron index f5ca4931d4..5eedb86818 100644 --- a/naga/tests/out/hlsl/image.ron +++ b/naga/tests/out/hlsl/image.ron @@ -3,38 +3,46 @@ ( entry_point:"queries", target_profile:"vs_5_1", + debug_printf:false, ), ( entry_point:"levels_queries", target_profile:"vs_5_1", + debug_printf:false, ), ], fragment:[ ( entry_point:"texture_sample", target_profile:"ps_5_1", + debug_printf:false, ), ( entry_point:"texture_sample_comparison", target_profile:"ps_5_1", + debug_printf:false, ), ( entry_point:"gather", target_profile:"ps_5_1", + debug_printf:false, ), ( entry_point:"depth_no_comparison", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ( entry_point:"depth_load", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/interface.ron b/naga/tests/out/hlsl/interface.ron index 948962b991..3fbe892516 100644 --- a/naga/tests/out/hlsl/interface.ron +++ b/naga/tests/out/hlsl/interface.ron @@ -3,22 +3,26 @@ ( entry_point:"vertex", target_profile:"vs_5_1", + debug_printf:false, ), ( entry_point:"vertex_two_structs", target_profile:"vs_5_1", + debug_printf:false, ), ], fragment:[ ( entry_point:"fragment", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ ( entry_point:"compute", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/interpolate.ron b/naga/tests/out/hlsl/interpolate.ron index d0046b04dd..88bc23f276 100644 --- a/naga/tests/out/hlsl/interpolate.ron +++ b/naga/tests/out/hlsl/interpolate.ron @@ -3,12 +3,14 @@ ( entry_point:"vert_main", target_profile:"vs_5_1", + debug_printf:false, ), ], fragment:[ ( entry_point:"frag_main", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/inv-hyperbolic-trig-functions.ron b/naga/tests/out/hlsl/inv-hyperbolic-trig-functions.ron index 341a4c528e..af020fcb28 100644 --- a/naga/tests/out/hlsl/inv-hyperbolic-trig-functions.ron +++ b/naga/tests/out/hlsl/inv-hyperbolic-trig-functions.ron @@ -5,6 +5,7 @@ ( entry_point:"main", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/math-functions.ron b/naga/tests/out/hlsl/math-functions.ron index 341a4c528e..af020fcb28 100644 --- a/naga/tests/out/hlsl/math-functions.ron +++ b/naga/tests/out/hlsl/math-functions.ron @@ -5,6 +5,7 @@ ( entry_point:"main", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/operators.ron b/naga/tests/out/hlsl/operators.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/operators.ron +++ b/naga/tests/out/hlsl/operators.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/padding.ron b/naga/tests/out/hlsl/padding.ron index 46dfdd83e3..0c164d4243 100644 --- a/naga/tests/out/hlsl/padding.ron +++ b/naga/tests/out/hlsl/padding.ron @@ -3,6 +3,7 @@ ( entry_point:"vertex", target_profile:"vs_5_1", + debug_printf:false, ), ], fragment:[ diff --git a/naga/tests/out/hlsl/push-constants.ron b/naga/tests/out/hlsl/push-constants.ron index e444486559..90d92ee5b2 100644 --- a/naga/tests/out/hlsl/push-constants.ron +++ b/naga/tests/out/hlsl/push-constants.ron @@ -3,12 +3,14 @@ ( entry_point:"vert_main", target_profile:"vs_5_1", + debug_printf:false, ), ], fragment:[ ( entry_point:"main", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/quad-vert.ron b/naga/tests/out/hlsl/quad-vert.ron index 8240856a5c..3142b9b2f7 100644 --- a/naga/tests/out/hlsl/quad-vert.ron +++ b/naga/tests/out/hlsl/quad-vert.ron @@ -3,6 +3,7 @@ ( entry_point:"main", target_profile:"vs_5_1", + debug_printf:false, ), ], fragment:[ diff --git a/naga/tests/out/hlsl/quad.ron b/naga/tests/out/hlsl/quad.ron index de90552356..3c7e251a95 100644 --- a/naga/tests/out/hlsl/quad.ron +++ b/naga/tests/out/hlsl/quad.ron @@ -3,16 +3,19 @@ ( entry_point:"vert_main", target_profile:"vs_5_1", + debug_printf:false, ), ], fragment:[ ( entry_point:"frag_main", target_profile:"ps_5_1", + debug_printf:false, ), ( entry_point:"fs_extra", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/shadow.ron b/naga/tests/out/hlsl/shadow.ron index 69be5b25e0..3b25fe1aaa 100644 --- a/naga/tests/out/hlsl/shadow.ron +++ b/naga/tests/out/hlsl/shadow.ron @@ -3,16 +3,19 @@ ( entry_point:"vs_main", target_profile:"vs_5_1", + debug_printf:false, ), ], fragment:[ ( entry_point:"fs_main", target_profile:"ps_5_1", + debug_printf:false, ), ( entry_point:"fs_main_without_storage", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/skybox.ron b/naga/tests/out/hlsl/skybox.ron index 27b0c4af4d..2f126ce7c1 100644 --- a/naga/tests/out/hlsl/skybox.ron +++ b/naga/tests/out/hlsl/skybox.ron @@ -3,12 +3,14 @@ ( entry_point:"vs_main", target_profile:"vs_5_1", + debug_printf:false, ), ], fragment:[ ( entry_point:"fs_main", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/standard.ron b/naga/tests/out/hlsl/standard.ron index 82373299d8..b2a703049b 100644 --- a/naga/tests/out/hlsl/standard.ron +++ b/naga/tests/out/hlsl/standard.ron @@ -5,6 +5,7 @@ ( entry_point:"derivatives", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/struct-layout.ron b/naga/tests/out/hlsl/struct-layout.ron index 04fe25e38a..16469b2de6 100644 --- a/naga/tests/out/hlsl/struct-layout.ron +++ b/naga/tests/out/hlsl/struct-layout.ron @@ -3,30 +3,36 @@ ( entry_point:"no_padding_vert", target_profile:"vs_5_1", + debug_printf:false, ), ( entry_point:"needs_padding_vert", target_profile:"vs_5_1", + debug_printf:false, ), ], fragment:[ ( entry_point:"no_padding_frag", target_profile:"ps_5_1", + debug_printf:false, ), ( entry_point:"needs_padding_frag", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ ( entry_point:"no_padding_comp", target_profile:"cs_5_1", + debug_printf:false, ), ( entry_point:"needs_padding_comp", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/texture-arg.ron b/naga/tests/out/hlsl/texture-arg.ron index 341a4c528e..af020fcb28 100644 --- a/naga/tests/out/hlsl/texture-arg.ron +++ b/naga/tests/out/hlsl/texture-arg.ron @@ -5,6 +5,7 @@ ( entry_point:"main", target_profile:"ps_5_1", + debug_printf:false, ), ], compute:[ diff --git a/naga/tests/out/hlsl/workgroup-uniform-load.ron b/naga/tests/out/hlsl/workgroup-uniform-load.ron index 17e926cdeb..4956d5bda8 100644 --- a/naga/tests/out/hlsl/workgroup-uniform-load.ron +++ b/naga/tests/out/hlsl/workgroup-uniform-load.ron @@ -7,6 +7,7 @@ ( entry_point:"test_workgroupUniformLoad", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/out/hlsl/workgroup-var-init.ron b/naga/tests/out/hlsl/workgroup-var-init.ron index a07b03300b..a3f6571767 100644 --- a/naga/tests/out/hlsl/workgroup-var-init.ron +++ b/naga/tests/out/hlsl/workgroup-var-init.ron @@ -7,6 +7,7 @@ ( entry_point:"main", target_profile:"cs_5_1", + debug_printf:false, ), ], ) diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index ac122a3a6c..311c204e43 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -549,7 +549,7 @@ fn write_output_hlsl( info: &naga::valid::ModuleInfo, options: &naga::back::hlsl::Options, ) { - use naga::back::hlsl; + use naga::back::hlsl::{self, WriterFlags}; use std::fmt::Write as _; println!("generating HLSL"); @@ -576,6 +576,8 @@ fn write_output_hlsl( } .push(hlsl_snapshots::ConfigItem { entry_point: name.clone(), + // Skip DXC until it supports debug printf + debug_printf: options.flags.contains(WriterFlags::EMIT_DEBUG_PRINTF), target_profile: format!( "{}_{}", ep.stage.to_hlsl_str(), @@ -816,7 +818,7 @@ fn convert_wgsl() { ), ( "debug-printf", - Targets::WGSL | Targets::GLSL | Targets::SPIRV, + Targets::WGSL | Targets::GLSL | Targets::SPIRV | Targets::HLSL, ), ]; @@ -897,7 +899,11 @@ fn convert_spv_all() { true, Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL, ); - convert_spv("debug-printf-s", false, Targets::GLSL | Targets::WGSL); + convert_spv( + "debug-printf-s", + false, + Targets::GLSL | Targets::WGSL | Targets::HLSL, + ); } #[cfg(feature = "glsl-in")] diff --git a/naga/xtask/src/validate.rs b/naga/xtask/src/validate.rs index 394b7b00d4..2a379e070c 100644 --- a/naga/xtask/src/validate.rs +++ b/naga/xtask/src/validate.rs @@ -311,21 +311,29 @@ fn validate_hlsl_with_dxc( config_item: hlsl_snapshots::ConfigItem, dxc: &str, ) -> anyhow::Result<()> { - // Reference: - // . - validate_hlsl( - file, - dxc, - config_item, - &[ - "-Wno-parentheses-equality", - "-Zi", - "-Qembed_debug", - "-Od", - "-HV", - "2018", - ], - ) + if config_item.debug_printf { + log::debug!( + "skipping config. item {config_item:?} because it \ + uses debug printf which is not supported on DXC" + ); + Ok(()) + } else { + // Reference: + // . + validate_hlsl( + file, + dxc, + config_item, + &[ + "-Wno-parentheses-equality", + "-Zi", + "-Qembed_debug", + "-Od", + "-HV", + "2018", + ], + ) + } } fn validate_hlsl_with_fxc( @@ -370,6 +378,7 @@ fn validate_hlsl( let hlsl_snapshots::ConfigItem { entry_point, target_profile, + .. } = config_item; EasyCommand::new(bin, |cmd| { cmd.arg(file)