From 8dd16259287f58f9273002717ec4d27e97127719 Mon Sep 17 00:00:00 2001 From: Daniel Baumann Date: Wed, 12 Jun 2024 07:43:14 +0200 Subject: Merging upstream version 127.0. Signed-off-by: Daniel Baumann --- third_party/rust/naga/src/back/hlsl/conv.rs | 5 + third_party/rust/naga/src/back/hlsl/help.rs | 94 +++++++- third_party/rust/naga/src/back/hlsl/mod.rs | 17 ++ third_party/rust/naga/src/back/hlsl/writer.rs | 315 ++++++++++++++++++++++---- 4 files changed, 387 insertions(+), 44 deletions(-) (limited to 'third_party/rust/naga/src/back/hlsl') diff --git a/third_party/rust/naga/src/back/hlsl/conv.rs b/third_party/rust/naga/src/back/hlsl/conv.rs index 2a6db35db8..7d15f43f6c 100644 --- a/third_party/rust/naga/src/back/hlsl/conv.rs +++ b/third_party/rust/naga/src/back/hlsl/conv.rs @@ -179,6 +179,11 @@ impl crate::BuiltIn { // to this field will get replaced with references to `SPECIAL_CBUF_VAR` // in `Writer::write_expr`. Self::NumWorkGroups => "SV_GroupID", + // These builtins map to functions + Self::SubgroupSize + | Self::SubgroupInvocationId + | Self::NumSubgroups + | Self::SubgroupId => unreachable!(), Self::BaseInstance | Self::BaseVertex | Self::WorkGroupSize => { return Err(Error::Unimplemented(format!("builtin {self:?}"))) } diff --git a/third_party/rust/naga/src/back/hlsl/help.rs b/third_party/rust/naga/src/back/hlsl/help.rs index 4dd9ea5987..d3bb1ce7f5 100644 --- a/third_party/rust/naga/src/back/hlsl/help.rs +++ b/third_party/rust/naga/src/back/hlsl/help.rs @@ -70,6 +70,11 @@ pub(super) struct WrappedMath { pub(super) components: Option, } +#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)] +pub(super) struct WrappedZeroValue { + pub(super) ty: Handle, +} + /// HLSL backend requires its own `ImageQuery` enum. /// /// It is used inside `WrappedImageQuery` and should be unique per ImageQuery function. @@ -359,7 +364,7 @@ impl<'a, W: Write> super::Writer<'a, W> { } /// Helper function that write wrapped function for `Expression::Compose` for structures. - pub(super) fn write_wrapped_constructor_function( + fn write_wrapped_constructor_function( &mut self, module: &crate::Module, constructor: WrappedConstructor, @@ -862,6 +867,25 @@ impl<'a, W: Write> super::Writer<'a, W> { Ok(()) } + // TODO: we could merge this with iteration in write_wrapped_compose_functions... + // + /// Helper function that writes zero value wrapped functions + pub(super) fn write_wrapped_zero_value_functions( + &mut self, + module: &crate::Module, + expressions: &crate::Arena, + ) -> BackendResult { + for (handle, _) in expressions.iter() { + if let crate::Expression::ZeroValue(ty) = expressions[handle] { + let zero_value = WrappedZeroValue { ty }; + if self.wrapped.zero_values.insert(zero_value) { + self.write_wrapped_zero_value_function(module, zero_value)?; + } + } + } + Ok(()) + } + pub(super) fn write_wrapped_math_functions( &mut self, module: &crate::Module, @@ -1006,6 +1030,7 @@ impl<'a, W: Write> super::Writer<'a, W> { ) -> BackendResult { self.write_wrapped_math_functions(module, func_ctx)?; self.write_wrapped_compose_functions(module, func_ctx.expressions)?; + self.write_wrapped_zero_value_functions(module, func_ctx.expressions)?; for (handle, _) in func_ctx.expressions.iter() { match func_ctx.expressions[handle] { @@ -1283,4 +1308,71 @@ impl<'a, W: Write> super::Writer<'a, W> { Ok(()) } + + pub(super) fn write_wrapped_zero_value_function_name( + &mut self, + module: &crate::Module, + zero_value: WrappedZeroValue, + ) -> BackendResult { + let name = crate::TypeInner::hlsl_type_id(zero_value.ty, module.to_ctx(), &self.names)?; + write!(self.out, "ZeroValue{name}")?; + Ok(()) + } + + /// Helper function that write wrapped function for `Expression::ZeroValue` + /// + /// This is necessary since we might have a member access after the zero value expression, e.g. + /// `.y` (in practice this can come up when consuming SPIRV that's been produced by glslc). + /// + /// So we can't just write `(float4)0` since `(float4)0.y` won't parse correctly. + /// + /// Parenthesizing the expression like `((float4)0).y` would work... except DXC can't handle + /// cases like: + /// + /// ```ignore + /// tests\out\hlsl\access.hlsl:183:41: error: cannot compile this l-value expression yet + /// t_1.am = (__mat4x2[2])((float4x2[2])0); + /// ^ + /// ``` + fn write_wrapped_zero_value_function( + &mut self, + module: &crate::Module, + zero_value: WrappedZeroValue, + ) -> BackendResult { + use crate::back::INDENT; + + const RETURN_VARIABLE_NAME: &str = "ret"; + + // Write function return type and name + if let crate::TypeInner::Array { base, size, .. } = module.types[zero_value.ty].inner { + write!(self.out, "typedef ")?; + self.write_type(module, zero_value.ty)?; + write!(self.out, " ret_")?; + self.write_wrapped_zero_value_function_name(module, zero_value)?; + self.write_array_size(module, base, size)?; + writeln!(self.out, ";")?; + + write!(self.out, "ret_")?; + self.write_wrapped_zero_value_function_name(module, zero_value)?; + } else { + self.write_type(module, zero_value.ty)?; + } + write!(self.out, " ")?; + self.write_wrapped_zero_value_function_name(module, zero_value)?; + + // Write function parameters (none) and start function body + writeln!(self.out, "() {{")?; + + // Write `ZeroValue` function. + write!(self.out, "{INDENT}return ")?; + self.write_default_init(module, zero_value.ty)?; + writeln!(self.out, ";")?; + + // End of function body + writeln!(self.out, "}}")?; + // Write extra new line + writeln!(self.out)?; + + Ok(()) + } } diff --git a/third_party/rust/naga/src/back/hlsl/mod.rs b/third_party/rust/naga/src/back/hlsl/mod.rs index f37a223f47..28edbf70e1 100644 --- a/third_party/rust/naga/src/back/hlsl/mod.rs +++ b/third_party/rust/naga/src/back/hlsl/mod.rs @@ -131,6 +131,13 @@ pub enum ShaderModel { V5_0, V5_1, V6_0, + V6_1, + V6_2, + V6_3, + V6_4, + V6_5, + V6_6, + V6_7, } impl ShaderModel { @@ -139,6 +146,13 @@ impl ShaderModel { Self::V5_0 => "5_0", Self::V5_1 => "5_1", Self::V6_0 => "6_0", + Self::V6_1 => "6_1", + Self::V6_2 => "6_2", + Self::V6_3 => "6_3", + Self::V6_4 => "6_4", + Self::V6_5 => "6_5", + Self::V6_6 => "6_6", + Self::V6_7 => "6_7", } } } @@ -247,10 +261,13 @@ pub enum Error { Unimplemented(String), // TODO: Error used only during development #[error("{0}")] Custom(String), + #[error("overrides should not be present at this stage")] + Override, } #[derive(Default)] struct Wrapped { + zero_values: crate::FastHashSet, array_lengths: crate::FastHashSet, image_queries: crate::FastHashSet, constructors: crate::FastHashSet, diff --git a/third_party/rust/naga/src/back/hlsl/writer.rs b/third_party/rust/naga/src/back/hlsl/writer.rs index 4ba856946b..86d8f89035 100644 --- a/third_party/rust/naga/src/back/hlsl/writer.rs +++ b/third_party/rust/naga/src/back/hlsl/writer.rs @@ -1,5 +1,8 @@ use super::{ - help::{WrappedArrayLength, WrappedConstructor, WrappedImageQuery, WrappedStructMatrixAccess}, + help::{ + WrappedArrayLength, WrappedConstructor, WrappedImageQuery, WrappedStructMatrixAccess, + WrappedZeroValue, + }, storage::StoreValue, BackendResult, Error, Options, }; @@ -77,6 +80,19 @@ enum Io { Output, } +const fn is_subgroup_builtin_binding(binding: &Option) -> bool { + let &Some(crate::Binding::BuiltIn(builtin)) = binding else { + return false; + }; + matches!( + builtin, + crate::BuiltIn::SubgroupSize + | crate::BuiltIn::SubgroupInvocationId + | crate::BuiltIn::NumSubgroups + | crate::BuiltIn::SubgroupId + ) +} + impl<'a, W: fmt::Write> super::Writer<'a, W> { pub fn new(out: W, options: &'a Options) -> Self { Self { @@ -161,6 +177,19 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } } } + for statement in func.body.iter() { + match *statement { + crate::Statement::SubgroupCollectiveOperation { + op: _, + collective_op: crate::CollectiveOperation::InclusiveScan, + argument, + result: _, + } => { + self.need_bake_expressions.insert(argument); + } + _ => {} + } + } } pub fn write( @@ -168,6 +197,10 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { module: &Module, module_info: &valid::ModuleInfo, ) -> Result { + if !module.overrides.is_empty() { + return Err(Error::Override); + } + self.reset(module); // Write special constants, if needed @@ -233,7 +266,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { self.write_special_functions(module)?; - self.write_wrapped_compose_functions(module, &module.const_expressions)?; + self.write_wrapped_compose_functions(module, &module.global_expressions)?; + self.write_wrapped_zero_value_functions(module, &module.global_expressions)?; // Write all named constants let mut constants = module @@ -397,31 +431,32 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // if they are struct, so that the `stage` argument here could be omitted. fn write_semantic( &mut self, - binding: &crate::Binding, + binding: &Option, stage: Option<(ShaderStage, Io)>, ) -> BackendResult { match *binding { - crate::Binding::BuiltIn(builtin) => { + Some(crate::Binding::BuiltIn(builtin)) if !is_subgroup_builtin_binding(binding) => { let builtin_str = builtin.to_hlsl_str()?; write!(self.out, " : {builtin_str}")?; } - crate::Binding::Location { + Some(crate::Binding::Location { second_blend_source: true, .. - } => { + }) => { write!(self.out, " : SV_Target1")?; } - crate::Binding::Location { + Some(crate::Binding::Location { location, second_blend_source: false, .. - } => { + }) => { if stage == Some((crate::ShaderStage::Fragment, Io::Output)) { write!(self.out, " : SV_Target{location}")?; } else { write!(self.out, " : {LOCATION_SEMANTIC}{location}")?; } } + _ => {} } Ok(()) @@ -442,17 +477,30 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { write!(self.out, "struct {struct_name}")?; writeln!(self.out, " {{")?; for m in members.iter() { + if is_subgroup_builtin_binding(&m.binding) { + continue; + } write!(self.out, "{}", back::INDENT)?; if let Some(ref binding) = m.binding { self.write_modifier(binding)?; } self.write_type(module, m.ty)?; write!(self.out, " {}", &m.name)?; - if let Some(ref binding) = m.binding { - self.write_semantic(binding, Some(shader_stage))?; - } + self.write_semantic(&m.binding, Some(shader_stage))?; writeln!(self.out, ";")?; } + if members.iter().any(|arg| { + matches!( + arg.binding, + Some(crate::Binding::BuiltIn(crate::BuiltIn::SubgroupId)) + ) + }) { + writeln!( + self.out, + "{}uint __local_invocation_index : SV_GroupIndex;", + back::INDENT + )?; + } writeln!(self.out, "}};")?; writeln!(self.out)?; @@ -553,8 +601,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } /// Writes special interface structures for an entry point. The special structures have - /// all the fields flattened into them and sorted by binding. They are only needed for - /// VS outputs and FS inputs, so that these interfaces match. + /// all the fields flattened into them and sorted by binding. They are needed to emulate + /// subgroup built-ins and to make the interfaces between VS outputs and FS inputs match. fn write_ep_interface( &mut self, module: &Module, @@ -563,7 +611,13 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { ep_name: &str, ) -> Result { Ok(EntryPointInterface { - input: if !func.arguments.is_empty() && stage == ShaderStage::Fragment { + input: if !func.arguments.is_empty() + && (stage == ShaderStage::Fragment + || func + .arguments + .iter() + .any(|arg| is_subgroup_builtin_binding(&arg.binding))) + { Some(self.write_ep_input_struct(module, func, stage, ep_name)?) } else { None @@ -577,6 +631,38 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { }) } + fn write_ep_argument_initialization( + &mut self, + ep: &crate::EntryPoint, + ep_input: &EntryPointBinding, + fake_member: &EpStructMember, + ) -> BackendResult { + match fake_member.binding { + Some(crate::Binding::BuiltIn(crate::BuiltIn::SubgroupSize)) => { + write!(self.out, "WaveGetLaneCount()")? + } + Some(crate::Binding::BuiltIn(crate::BuiltIn::SubgroupInvocationId)) => { + write!(self.out, "WaveGetLaneIndex()")? + } + Some(crate::Binding::BuiltIn(crate::BuiltIn::NumSubgroups)) => write!( + self.out, + "({}u + WaveGetLaneCount() - 1u) / WaveGetLaneCount()", + ep.workgroup_size[0] * ep.workgroup_size[1] * ep.workgroup_size[2] + )?, + Some(crate::Binding::BuiltIn(crate::BuiltIn::SubgroupId)) => { + write!( + self.out, + "{}.__local_invocation_index / WaveGetLaneCount()", + ep_input.arg_name + )?; + } + _ => { + write!(self.out, "{}.{}", ep_input.arg_name, fake_member.name)?; + } + } + Ok(()) + } + /// Write an entry point preface that initializes the arguments as specified in IR. fn write_ep_arguments_initialization( &mut self, @@ -584,6 +670,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { func: &crate::Function, ep_index: u16, ) -> BackendResult { + let ep = &module.entry_points[ep_index as usize]; let ep_input = match self.entry_point_io[ep_index as usize].input.take() { Some(ep_input) => ep_input, None => return Ok(()), @@ -597,8 +684,13 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { match module.types[arg.ty].inner { TypeInner::Array { base, size, .. } => { self.write_array_size(module, base, size)?; - let fake_member = fake_iter.next().unwrap(); - writeln!(self.out, " = {}.{};", ep_input.arg_name, fake_member.name)?; + write!(self.out, " = ")?; + self.write_ep_argument_initialization( + ep, + &ep_input, + fake_iter.next().unwrap(), + )?; + writeln!(self.out, ";")?; } TypeInner::Struct { ref members, .. } => { write!(self.out, " = {{ ")?; @@ -606,14 +698,22 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { if index != 0 { write!(self.out, ", ")?; } - let fake_member = fake_iter.next().unwrap(); - write!(self.out, "{}.{}", ep_input.arg_name, fake_member.name)?; + self.write_ep_argument_initialization( + ep, + &ep_input, + fake_iter.next().unwrap(), + )?; } writeln!(self.out, " }};")?; } _ => { - let fake_member = fake_iter.next().unwrap(); - writeln!(self.out, " = {}.{};", ep_input.arg_name, fake_member.name)?; + write!(self.out, " = ")?; + self.write_ep_argument_initialization( + ep, + &ep_input, + fake_iter.next().unwrap(), + )?; + writeln!(self.out, ";")?; } } } @@ -928,9 +1028,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } } - if let Some(ref binding) = member.binding { - self.write_semantic(binding, shader_stage)?; - }; + self.write_semantic(&member.binding, shader_stage)?; writeln!(self.out, ";")?; } @@ -1143,7 +1241,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } back::FunctionType::EntryPoint(ep_index) => { if let Some(ref ep_input) = self.entry_point_io[ep_index as usize].input { - write!(self.out, "{} {}", ep_input.ty_name, ep_input.arg_name,)?; + write!(self.out, "{} {}", ep_input.ty_name, ep_input.arg_name)?; } else { let stage = module.entry_points[ep_index as usize].stage; for (index, arg) in func.arguments.iter().enumerate() { @@ -1160,17 +1258,16 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { self.write_array_size(module, base, size)?; } - if let Some(ref binding) = arg.binding { - self.write_semantic(binding, Some((stage, Io::Input)))?; - } + self.write_semantic(&arg.binding, Some((stage, Io::Input)))?; } - - if need_workgroup_variables_initialization { - if !func.arguments.is_empty() { - write!(self.out, ", ")?; - } - write!(self.out, "uint3 __local_invocation_id : SV_GroupThreadID")?; + } + if need_workgroup_variables_initialization { + if self.entry_point_io[ep_index as usize].input.is_some() + || !func.arguments.is_empty() + { + write!(self.out, ", ")?; } + write!(self.out, "uint3 __local_invocation_id : SV_GroupThreadID")?; } } } @@ -1180,11 +1277,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // Write semantic if it present if let back::FunctionType::EntryPoint(index) = func_ctx.ty { let stage = module.entry_points[index as usize].stage; - if let Some(crate::FunctionResult { - binding: Some(ref binding), - .. - }) = func.result - { + if let Some(crate::FunctionResult { ref binding, .. }) = func.result { self.write_semantic(binding, Some((stage, Io::Output)))?; } } @@ -1984,6 +2077,129 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { writeln!(self.out, "{level}}}")? } Statement::RayQuery { .. } => unreachable!(), + Statement::SubgroupBallot { result, predicate } => { + write!(self.out, "{level}")?; + let name = format!("{}{}", back::BAKE_PREFIX, result.index()); + write!(self.out, "const uint4 {name} = ")?; + self.named_expressions.insert(result, name); + + write!(self.out, "WaveActiveBallot(")?; + match predicate { + Some(predicate) => self.write_expr(module, predicate, func_ctx)?, + None => write!(self.out, "true")?, + } + writeln!(self.out, ");")?; + } + Statement::SubgroupCollectiveOperation { + op, + collective_op, + argument, + result, + } => { + write!(self.out, "{level}")?; + write!(self.out, "const ")?; + let name = format!("{}{}", back::BAKE_PREFIX, result.index()); + match func_ctx.info[result].ty { + proc::TypeResolution::Handle(handle) => self.write_type(module, handle)?, + proc::TypeResolution::Value(ref value) => { + self.write_value_type(module, value)? + } + }; + write!(self.out, " {name} = ")?; + self.named_expressions.insert(result, name); + + match (collective_op, op) { + (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::All) => { + write!(self.out, "WaveActiveAllTrue(")? + } + (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Any) => { + write!(self.out, "WaveActiveAnyTrue(")? + } + (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Add) => { + write!(self.out, "WaveActiveSum(")? + } + (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Mul) => { + write!(self.out, "WaveActiveProduct(")? + } + (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Max) => { + write!(self.out, "WaveActiveMax(")? + } + (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Min) => { + write!(self.out, "WaveActiveMin(")? + } + (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::And) => { + write!(self.out, "WaveActiveBitAnd(")? + } + (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Or) => { + write!(self.out, "WaveActiveBitOr(")? + } + (crate::CollectiveOperation::Reduce, crate::SubgroupOperation::Xor) => { + write!(self.out, "WaveActiveBitXor(")? + } + (crate::CollectiveOperation::ExclusiveScan, crate::SubgroupOperation::Add) => { + write!(self.out, "WavePrefixSum(")? + } + (crate::CollectiveOperation::ExclusiveScan, crate::SubgroupOperation::Mul) => { + write!(self.out, "WavePrefixProduct(")? + } + (crate::CollectiveOperation::InclusiveScan, crate::SubgroupOperation::Add) => { + self.write_expr(module, argument, func_ctx)?; + write!(self.out, " + WavePrefixSum(")?; + } + (crate::CollectiveOperation::InclusiveScan, crate::SubgroupOperation::Mul) => { + self.write_expr(module, argument, func_ctx)?; + write!(self.out, " * WavePrefixProduct(")?; + } + _ => unimplemented!(), + } + self.write_expr(module, argument, func_ctx)?; + writeln!(self.out, ");")?; + } + Statement::SubgroupGather { + mode, + argument, + result, + } => { + write!(self.out, "{level}")?; + write!(self.out, "const ")?; + let name = format!("{}{}", back::BAKE_PREFIX, result.index()); + match func_ctx.info[result].ty { + proc::TypeResolution::Handle(handle) => self.write_type(module, handle)?, + proc::TypeResolution::Value(ref value) => { + self.write_value_type(module, value)? + } + }; + write!(self.out, " {name} = ")?; + self.named_expressions.insert(result, name); + + if matches!(mode, crate::GatherMode::BroadcastFirst) { + write!(self.out, "WaveReadLaneFirst(")?; + self.write_expr(module, argument, func_ctx)?; + } else { + write!(self.out, "WaveReadLaneAt(")?; + self.write_expr(module, argument, func_ctx)?; + write!(self.out, ", ")?; + match mode { + crate::GatherMode::BroadcastFirst => unreachable!(), + crate::GatherMode::Broadcast(index) | crate::GatherMode::Shuffle(index) => { + self.write_expr(module, index, func_ctx)?; + } + crate::GatherMode::ShuffleDown(index) => { + write!(self.out, "WaveGetLaneIndex() + ")?; + self.write_expr(module, index, func_ctx)?; + } + crate::GatherMode::ShuffleUp(index) => { + write!(self.out, "WaveGetLaneIndex() - ")?; + self.write_expr(module, index, func_ctx)?; + } + crate::GatherMode::ShuffleXor(index) => { + write!(self.out, "WaveGetLaneIndex() ^ ")?; + self.write_expr(module, index, func_ctx)?; + } + } + } + writeln!(self.out, ");")?; + } } Ok(()) @@ -1997,7 +2213,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { self.write_possibly_const_expression( module, expr, - &module.const_expressions, + &module.global_expressions, |writer, expr| writer.write_const_expression(module, expr), ) } @@ -2039,7 +2255,10 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { self.write_const_expression(module, constant.init)?; } } - Expression::ZeroValue(ty) => self.write_default_init(module, ty)?, + Expression::ZeroValue(ty) => { + self.write_wrapped_zero_value_function_name(module, WrappedZeroValue { ty })?; + write!(self.out, "()")?; + } Expression::Compose { ty, ref components } => { match module.types[ty].inner { TypeInner::Struct { .. } | TypeInner::Array { .. } => { @@ -2140,6 +2359,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { |writer, expr| writer.write_expr(module, expr, func_ctx), )?; } + Expression::Override(_) => return Err(Error::Override), // All of the multiplication can be expressed as `mul`, // except vector * vector, which needs to use the "*" operator. Expression::Binary { @@ -2588,7 +2808,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { true } None => { - if inner.scalar_width() == Some(64) { + if inner.scalar_width() == Some(8) { false } else { write!(self.out, "{}(", kind.to_hlsl_cast(),)?; @@ -3129,7 +3349,9 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { Expression::CallResult(_) | Expression::AtomicResult { .. } | Expression::WorkGroupUniformLoadResult { .. } - | Expression::RayQueryProceedResult => {} + | Expression::RayQueryProceedResult + | Expression::SubgroupBallotResult + | Expression::SubgroupOperationResult { .. } => {} } if !closing_bracket.is_empty() { @@ -3179,7 +3401,11 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } /// Helper function that write default zero initialization - fn write_default_init(&mut self, module: &Module, ty: Handle) -> BackendResult { + pub(super) fn write_default_init( + &mut self, + module: &Module, + ty: Handle, + ) -> BackendResult { write!(self.out, "(")?; self.write_type(module, ty)?; if let TypeInner::Array { base, size, .. } = module.types[ty].inner { @@ -3196,6 +3422,9 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { if barrier.contains(crate::Barrier::WORK_GROUP) { writeln!(self.out, "{level}GroupMemoryBarrierWithGroupSync();")?; } + if barrier.contains(crate::Barrier::SUB_GROUP) { + // Does not exist in DirectX + } Ok(()) } } -- cgit v1.2.3