enum Index {
Expression(Handle<crate::Expression>), Static(u32),
}
pub(super) struct EpStructMember { pub(super) name: String, pub(super) ty: Handle<crate::Type>, // technically, this should always be `Some` // (we `debug_assert!` this in `write_interface_struct`) pub(super) binding: Option<crate::Binding>, pub(super) index: u32,
}
/// Structure contains information required for generating /// wrapped structure of all entry points arguments pub(super) struct EntryPointBinding { /// Name of the fake EP argument that contains the struct /// with all the flattened input data. pub(super) arg_name: String, /// Generated structure name pub(super) ty_name: String, /// Members of generated structure pub(super) members: Vec<EpStructMember>, pub(super) local_invocation_index_name: Option<String>,
}
pub(super) struct EntryPointInterface { /// If `Some`, the input of an entry point is gathered in a special /// struct with members sorted by binding. /// The `EntryPointBinding::members` array is sorted by index, /// so that we can walk it in `write_ep_arguments_initialization`. pub(crate) input: Option<EntryPointBinding>, /// If `Some`, the output of an entry point is flattened. /// The `EntryPointBinding::members` array is sorted by binding, /// So that we can walk it in `Statement::Return` handler. pub(crate) output: Option<EntryPointBinding>, pub(crate) mesh_vertices: Option<EntryPointBinding>, pub(crate) mesh_primitives: Option<EntryPointBinding>, pub(crate) mesh_indices: Option<EntryPointBinding>,
}
/// Information for how to generate a `binding_array<sampler>` access. struct BindingArraySamplerInfo { /// Variable name of the sampler heap
sampler_heap_name: &'static str, /// Variable name of the sampler index buffer
sampler_index_buffer_name: String, /// Variable name of the base index _into_ the sampler index buffer
binding_array_base_index_name: String,
}
/// Generates statements to be inserted immediately before and at the very /// start of the body of each loop, to defeat infinite loop reasoning. /// The 0th item of the returned tuple should be inserted immediately prior /// to the loop and the 1st item should be inserted at the very start of /// the loop body. /// /// See [`back::msl::Writer::gen_force_bounded_loop_statements`] for details. fn gen_force_bounded_loop_statements(
&mutself,
level: back::Level,
) -> Option<(String, String)> { if !self.options.force_loop_bounding { return None;
}
let loop_bound_name = self.namer.call("loop_bound"); let max = u32::MAX; // Count down from u32::MAX rather than up from 0 to avoid hang on // certain Intel drivers. See <https://github.com/gfx-rs/wgpu/issues/7319>. let decl = format!("{level}uint2 {loop_bound_name} = uint2({max}u, {max}u);"); let level = level.next(); let break_and_inc = format!( "{level}if (all({loop_bound_name} == uint2(0u, 0u))) {{ break; }}
{level}{loop_bound_name} -= uint2({loop_bound_name}.y == 0u, 1u);"
);
Some((decl, break_and_inc))
}
/// Helper method used to find which expressions of a given function require baking /// /// # Notes /// Clears `need_bake_expressions` set before adding to it fn update_expressions_to_bake(
&mutself,
module: &Module,
func: &crate::Function,
info: &valid::FunctionInfo,
) { usecrate::Expression; self.need_bake_expressions.clear(); for (exp_handle, expr) in func.expressions.iter() { let expr_info = &info[exp_handle]; let min_ref_count = func.expressions[exp_handle].bake_ref_count(); if min_ref_count <= expr_info.ref_count { self.need_bake_expressions.insert(exp_handle);
} iflet Expression::Load { pointer } = *expr { if info[pointer]
.ty
.inner_with(&module.types)
.is_atomic_pointer(&module.types)
{ self.need_bake_expressions.insert(exp_handle);
}
}
// Extra newline for readability
writeln!(self.out)?;
}
// Save all entry point output types let ep_results = module
.entry_points
.iter()
.map(|ep| (ep.stage, ep.function.result.clone()))
.collect::<Vec<(ShaderStage, Option<crate::FunctionResult>)>>();
// Write all structs for (handle, ty) in module.types.iter() { iflet TypeInner::Struct { ref members, span } = ty.inner { if module.types[members.last().unwrap().ty]
.inner
.is_dynamically_sized(&module.types)
{ // unsized arrays can only be in storage buffers, // for which we use `ByteAddressBuffer` anyway. continue;
}
// Write all named constants letmut constants = module
.constants
.iter()
.filter(|&(_, c)| c.name.is_some())
.peekable(); whilelet Some((handle, _)) = constants.next() { self.write_global_constant(module, handle)?; // Add extra newline for readability on last iteration if constants.peek().is_none() {
writeln!(self.out)?;
}
}
// Write all globals for (global, _) in module.global_variables.iter() { self.write_global(module, global)?;
}
if !module.global_variables.is_empty() { // Add extra newline for readability
writeln!(self.out)?;
}
let ep_range = get_entry_points(module, self.pipeline_options.entry_point.as_ref())
.map_err(|(stage, name)| Error::EntryPointNotFound(stage, name))?;
// Write all entry points wrapped structs for index in ep_range.clone() { let ep = &module.entry_points[index]; let ep_name = self.names[&NameKey::EntryPoint(index as u16)].clone(); let ep_io = self.write_ep_interface(module, ep, &ep_name, fragment_entry_point)?; self.entry_point_io.insert(index, ep_io);
}
// Write all regular functions for (handle, function) in module.functions.iter() { let info = &module_info[handle];
// Check if all of the globals are accessible if !self.options.fake_missing_bindings { iflet Some((var_handle, _)) =
module
.global_variables
.iter()
.find(|&(var_handle, var)| match var.binding {
Some(ref binding) if !info[var_handle].is_empty() => { self.options.resolve_resource_binding(binding).is_err()
&& self
.options
.resolve_external_texture_resource_binding(binding)
.is_err()
}
_ => false,
})
{
log::debug!( "Skipping function {:?} (name {:?}) because global {:?} is inaccessible",
handle,
function.name,
var_handle
); continue;
}
}
let ctx = back::FunctionCtx {
ty: back::FunctionType::Function(handle),
info,
expressions: &function.expressions,
named_expressions: &function.named_expressions,
}; let name = self.names[&NameKey::Function(handle)].clone();
let ctx = back::FunctionCtx {
ty: back::FunctionType::EntryPoint(index as u16),
info,
expressions: &ep.function.expressions,
named_expressions: &ep.function.named_expressions,
};
self.write_wrapped_functions(module, &ctx)?;
// Mesh/task shaders have a wrapper entry point which is declared after the "main" // user-written function. We therefore cannot always just document the next function. letmut attribute_string = String::new(); if ep.stage.compute_like() { // HLSL is calling workgroup size "num threads" let num_threads = ep.workgroup_size;
writeln!(
attribute_string, "[numthreads({}, {}, {})]",
num_threads[0], num_threads[1], num_threads[2]
)?;
} iflet Some(ref info) = ep.mesh_info { let topology_str = match info.topology { crate::MeshOutputTopology::Points => unreachable!(), crate::MeshOutputTopology::Lines => "line", crate::MeshOutputTopology::Triangles => "triangle",
};
writeln!(attribute_string, "[outputtopology(\"{topology_str}\")]")?;
}
let name = self.names[&NameKey::EntryPoint(index as u16)].clone(); self.write_function(module, &name, &ep.function, &ctx, info, attribute_string)?;
if index < module.entry_points.len() - 1 {
writeln!(self.out)?;
}
//TODO: we could force fragment outputs to always go through `entry_point_io.output` path // if they are struct, so that the `stage` argument here could be omitted. pub(super) fn write_semantic(
&mutself,
binding: &Option<crate::Binding>,
stage: Option<(ShaderStage, Io)>,
) -> BackendResult { let is_per_primitive = match *binding {
Some(crate::Binding::BuiltIn(builtin)) if !is_subgroup_builtin_binding(binding) => { if builtin == crate::BuiltIn::ViewIndex
&& self.options.shader_model < ShaderModel::V6_1
{ return Err(Error::ShaderModelTooLow( "used @builtin(view_index) or SV_ViewID".to_string(),
ShaderModel::V6_1,
));
} iflet Some(builtin_str) = builtin.to_hlsl_str()? {
write!(self.out, " : {builtin_str}")?;
} false
}
Some(crate::Binding::Location {
blend_src: Some(1),
per_primitive,
..
}) => {
write!(self.out, " : SV_Target1")?;
per_primitive
}
Some(crate::Binding::Location {
location,
per_primitive,
..
}) => { if stage == Some((ShaderStage::Fragment, Io::Output)) {
write!(self.out, " : SV_Target{location}")?;
} else {
write!(self.out, " : {LOCATION_SEMANTIC}{location}")?;
}
per_primitive
}
_ => false,
}; if is_per_primitive {
write!(self.out, " : primitive")?;
}
Ok(())
}
pub(super) fn write_interface_struct(
&mutself,
module: &Module,
shader_stage: (ShaderStage, Io),
struct_name: String,
var_name: Option<&str>, mut members: Vec<EpStructMember>,
) -> Result<EntryPointBinding, Error> { let struct_name = self.namer.call(&struct_name); // Sort the members so that first come the user-defined varyings // in ascending locations, and then built-ins. This allows VS and FS // interfaces to match with regards to order.
members.sort_by_key(|m| InterfaceKey::new(m.binding.as_ref()));
write!(self.out, "struct {struct_name}")?;
writeln!(self.out, " {{")?; letmut local_invocation_index_name = None; letmut subgroup_id_used = false; for m in members.iter() { // Sanity check that each IO member is a built-in or is assigned a // location. Also see note about nesting in `write_ep_input_struct`.
debug_assert!(m.binding.is_some());
// See ordering notes on EntryPointInterface fields match shader_stage.1 {
Io::Input => { // bring back the original order
members.sort_by_key(|m| m.index);
}
Io::Output | Io::MeshVertices | Io::MeshPrimitives => { // keep it sorted by binding
}
}
/// Flatten all entry point arguments into a single struct. /// This is needed since we need to re-order them: first placing user locations, /// then built-ins. fn write_ep_input_struct(
&mutself,
module: &Module,
func: &crate::Function,
stage: ShaderStage,
entry_point_name: &str,
) -> Result<EntryPointBinding, Error> { let struct_name = format!("{stage:?}Input_{entry_point_name}");
letmut fake_members = Vec::new(); for arg in func.arguments.iter() { // NOTE: We don't need to handle nesting structs. All members must // be either built-ins or assigned a location. I.E. `binding` is // `Some`. This is checked in `VaryingContext::validate`. See: // https://gpuweb.github.io/gpuweb/wgsl/#input-output-locations match module.types[arg.ty].inner {
TypeInner::Struct { ref members, .. } => { for member in members.iter() { let name = self.namer.call_or(&member.name, "member"); let index = fake_members.len() as u32;
fake_members.push(EpStructMember {
name,
ty: member.ty,
binding: member.binding.clone(),
index,
});
}
}
_ => { let member_name = self.namer.call_or(&arg.name, "member"); let index = fake_members.len() as u32;
fake_members.push(EpStructMember {
name: member_name,
ty: arg.ty,
binding: arg.binding.clone(),
index,
});
}
}
}
/// Flatten all entry point results into a single struct. /// This is needed since we need to re-order them: first placing user locations, /// then built-ins. fn write_ep_output_struct(
&mutself,
module: &Module,
result: &crate::FunctionResult,
stage: ShaderStage,
entry_point_name: &str,
frag_ep: Option<&FragmentEntryPoint<'_>>,
) -> Result<EntryPointBinding, Error> { let struct_name = format!("{stage:?}Output_{entry_point_name}");
let empty = []; let members = match module.types[result.ty].inner {
TypeInner::Struct { ref members, .. } => members, ref other => {
log::error!("Unexpected {other:?} output type without a binding");
&empty[..]
}
};
// Gather list of fragment input locations. We use this below to remove user-defined // varyings from VS outputs that aren't in the FS inputs. This makes the VS interface match // as long as the FS inputs are a subset of the VS outputs. This is only applied if the // writer is supplied with information about the fragment entry point. let fs_input_locs = iflet (Some(frag_ep), ShaderStage::Vertex) = (frag_ep, stage) { letmut fs_input_locs = Vec::new(); for arg in frag_ep.func.arguments.iter() { letmut push_if_location = |binding: &Option<crate::Binding>| match *binding {
Some(crate::Binding::Location { location, .. }) => fs_input_locs.push(location),
Some(crate::Binding::BuiltIn(_)) | None => {}
};
// NOTE: We don't need to handle struct nesting. See note in // `write_ep_input_struct`. match frag_ep.module.types[arg.ty].inner {
TypeInner::Struct { ref members, .. } => { for member in members.iter() {
push_if_location(&member.binding);
}
}
_ => push_if_location(&arg.binding),
}
}
fs_input_locs.sort();
Some(fs_input_locs)
} else {
None
};
letmut fake_members = Vec::new(); for (index, member) in members.iter().enumerate() { iflet Some(ref fs_input_locs) = fs_input_locs { match member.binding {
Some(crate::Binding::Location { location, .. }) => { if fs_input_locs.binary_search(&location).is_err() { continue;
}
}
Some(crate::Binding::BuiltIn(_)) | None => {}
}
}
let member_name = self.namer.call_or(&member.name, "member");
fake_members.push(EpStructMember {
name: member_name,
ty: member.ty,
binding: member.binding.clone(),
index: index as u32,
});
}
/// Write an entry point preface that initializes the arguments as specified in IR. fn write_ep_arguments_initialization(
&mutself,
module: &Module,
func: &crate::Function,
ep_index: u16,
) -> BackendResult { let ep = &module.entry_points[ep_index as usize]; let ep_input = matchself
.entry_point_io
.get_mut(&(ep_index as usize))
.unwrap()
.input
.take()
{
Some(ep_input) => ep_input,
None => return Ok(()),
}; letmut fake_iter = ep_input.members.iter(); for (arg_index, arg) in func.arguments.iter().enumerate() {
write!(self.out, "{}", back::INDENT)?; self.write_type(module, arg.ty)?; let arg_name = &self.names[&NameKey::EntryPointArgument(ep_index, arg_index as u32)];
write!(self.out, " {arg_name}")?; match module.types[arg.ty].inner {
TypeInner::Array { base, size, .. } => { self.write_array_size(module, base, size)?;
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, " = {{ ")?; for index in0..members.len() { if index != 0 {
write!(self.out, ", ")?;
} self.write_ep_argument_initialization(
ep,
&ep_input,
fake_iter.next().unwrap(),
)?;
}
writeln!(self.out, " }};")?;
}
_ => {
write!(self.out, " = ")?; self.write_ep_argument_initialization(
ep,
&ep_input,
fake_iter.next().unwrap(),
)?;
writeln!(self.out, ";")?;
}
}
}
assert!(fake_iter.next().is_none());
Ok(())
}
/// Helper method used to write global variables /// # Notes /// Always adds a newline fn write_global(
&mutself,
module: &Module,
handle: Handle<crate::GlobalVariable>,
) -> BackendResult { let global = &module.global_variables[handle]; let inner = &module.types[global.ty].inner;
let handle_ty = match *inner {
TypeInner::BindingArray { ref base, .. } => &module.types[*base].inner,
_ => inner,
};
// External textures are handled entirely differently, so defer entirely to that method. // We do so prior to calling resolve_resource_binding() below, as we even need to resolve // their bindings separately. let is_external_texture = matches!(
*handle_ty,
TypeInner::Image {
class: crate::ImageClass::External,
..
}
); if is_external_texture { returnself.write_global_external_texture(module, handle, global);
}
iflet Some(ref binding) = global.binding { iflet Err(err) = self.options.resolve_resource_binding(binding) {
log::debug!( "Skipping global {:?} (name {:?}) for being inaccessible: {}",
handle,
global.name,
err,
); return Ok(());
}
}
// Samplers are handled entirely differently, so defer entirely to that method. let is_sampler = matches!(*handle_ty, TypeInner::Sampler { .. });
if is_sampler { returnself.write_global_sampler(module, handle, global);
}
// https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-variable-register let register_ty = match global.space { crate::AddressSpace::Function => unreachable!("Function address space"), crate::AddressSpace::Private => {
write!(self.out, "static ")?; self.write_type(module, global.ty)?; ""
} crate::AddressSpace::WorkGroup | crate::AddressSpace::TaskPayload => {
write!(self.out, "groupshared ")?; self.write_type(module, global.ty)?; ""
} crate::AddressSpace::Uniform => { // constant buffer declarations are expected to be inlined, e.g. // `cbuffer foo: register(b0) { field1: type1; }`
write!(self.out, "cbuffer")?; "b"
} crate::AddressSpace::Storage { access } => { if global
.memory_decorations
.contains(crate::MemoryDecorations::COHERENT)
{
write!(self.out, "globallycoherent ")?;
} let (prefix, register) = if access.contains(crate::StorageAccess::STORE) {
("RW", "u")
} else {
("", "t")
};
write!(self.out, "{prefix}ByteAddressBuffer")?;
register
} crate::AddressSpace::Handle => { let register = match *handle_ty { // all storage textures are UAV, unconditionally
TypeInner::Image {
class: crate::ImageClass::Storage { .. },
..
} => "u",
_ => "t",
}; self.write_type(module, global.ty)?;
register
} crate::AddressSpace::Immediate => { // The type of the immediates will be wrapped in `ConstantBuffer`
write!(self.out, "ConstantBuffer<")?; "b"
} crate::AddressSpace::RayPayload | crate::AddressSpace::IncomingRayPayload => {
unimplemented!()
}
};
// If the global is a immediate data write the type now because it will be a // generic argument to `ConstantBuffer` if global.space == crate::AddressSpace::Immediate { self.write_global_type(module, global.ty)?;
// need to write the array size if the type was emitted with `write_type` iflet TypeInner::Array { base, size, .. } = module.types[global.ty].inner { self.write_array_size(module, base, size)?;
}
// Close the angled brackets for the generic argument
write!(self.out, ">")?;
}
let name = &self.names[&NameKey::GlobalVariable(handle)];
write!(self.out, " {name}")?;
// Immediates need to be assigned a binding explicitly by the consumer // since naga has no way to know the binding from the shader alone if global.space == crate::AddressSpace::Immediate { match module.types[global.ty].inner {
TypeInner::Struct { .. } => {}
_ => { return Err(Error::Unimplemented(format!( "push-constant '{name}' has non-struct type; tracked by: https://github.com/gfx-rs/wgpu/issues/5683"
)));
}
}
let target = self
.options
.immediates_target
.as_ref()
.expect("No bind target was defined for the immediates block");
write!(self.out, ": register(b{}", target.register)?; if target.space != 0 {
write!(self.out, ", space{}", target.space)?;
}
write!(self.out, ")")?;
}
iflet Some(ref binding) = global.binding { // this was already resolved earlier when we started evaluating an entry point. let bt = self.options.resolve_resource_binding(binding).unwrap();
// need to write the binding array size if the type was emitted with `write_type` iflet TypeInner::BindingArray { base, size, .. } = module.types[global.ty].inner { iflet Some(overridden_size) = bt.binding_array_size {
write!(self.out, "[{overridden_size}]")?;
} else { self.write_array_size(module, base, size)?;
}
}
write!(self.out, " : register({}{}", register_ty, bt.register)?; if bt.space != 0 {
write!(self.out, ", space{}", bt.space)?;
}
write!(self.out, ")")?;
} else { // need to write the array size if the type was emitted with `write_type` iflet TypeInner::Array { base, size, .. } = module.types[global.ty].inner { self.write_array_size(module, base, size)?;
} if global.space == crate::AddressSpace::Private {
write!(self.out, " = ")?; iflet Some(init) = global.init { self.write_const_expression(module, init, &module.global_expressions)?;
} else { self.write_default_init(module, global.ty)?;
}
}
}
if global.space == crate::AddressSpace::Uniform {
write!(self.out, " {{ ")?;
// need to write the array size if the type was emitted with `write_type` iflet TypeInner::Array { base, size, .. } = module.types[global.ty].inner { self.write_array_size(module, base, size)?;
}
let key = super::SamplerIndexBufferKey {
group: binding.group,
}; self.write_wrapped_sampler_buffer(key)?;
// This was already validated, so we can confidently unwrap it. let bt = self.options.resolve_resource_binding(&binding).unwrap();
match module.types[global.ty].inner {
TypeInner::Sampler { comparison } => { // If we are generating a static access, we create a variable for the sampler. // // This prevents the DXIL from containing multiple lookups for the sampler, which // the backend compiler will then have to eliminate. AMD does seem to be able to // eliminate these, but better safe than sorry.
let heap_var = if comparison {
COMPARISON_SAMPLER_HEAP_VAR
} else {
SAMPLER_HEAP_VAR
};
let index_buffer_name = &self.wrapped.sampler_index_buffers[&key]; let name = &self.names[&NameKey::GlobalVariable(handle)];
writeln!( self.out, " {name} = {heap_var}[{index_buffer_name}[{register}]];",
register = bt.register
)?;
}
TypeInner::BindingArray { .. } => { // If we are generating a binding array, we cannot directly access the sampler as the index // into the sampler index buffer is unknown at compile time. Instead we generate a constant // that represents the "base" index into the sampler index buffer. This constant is added // to the user provided index to get the final index into the sampler index buffer.
/// Write the declarations for an external texture global variable. /// These are emitted as multiple global variables: Three `Texture2D`s /// (one for each plane) and a parameters cbuffer. fn write_global_external_texture(
&mutself,
module: &Module,
handle: Handle<crate::GlobalVariable>,
global: &crate::GlobalVariable,
) -> BackendResult { let res_binding = global
.binding
.as_ref()
.expect("External texture global variables must have a resource binding"); let ext_tex_bindings = matchself
.options
.resolve_external_texture_resource_binding(res_binding)
{
Ok(bindings) => bindings,
Err(err) => {
log::debug!( "Skipping global {:?} (name {:?}) for being inaccessible: {}",
handle,
global.name,
err,
); return Ok(());
}
};
/// Helper method used to write structs /// /// # Notes /// Ends in a newline fn write_struct(
&mutself,
module: &Module,
handle: Handle<crate::Type>,
members: &[crate::StructMember],
span: u32,
shader_stage: Option<(ShaderStage, Io)>,
) -> BackendResult { // Write struct name let struct_name = &self.names[&NameKey::Type(handle)];
writeln!(self.out, "struct {struct_name} {{")?;
letmut last_offset = 0; for (index, member) in members.iter().enumerate() { if member.binding.is_none() && member.offset > last_offset { // using int as padding should work as long as the backend // doesn't support a type that's less than 4 bytes in size // (Error::UnsupportedScalar catches this) let padding = (member.offset - last_offset) / 4; for i in0..padding {
writeln!(self.out, "{}int _pad{}_{};", back::INDENT, index, i)?;
}
} let ty_inner = &module.types[member.ty].inner;
last_offset = member.offset + ty_inner.size_hlsl(module.to_ctx())?;
// The indentation is only for readability
write!(self.out, "{}", back::INDENT)?;
match module.types[member.ty].inner {
TypeInner::Array { base, size, .. } => { // HLSL arrays are written as `type name[size]`
self.write_global_type(module, member.ty)?;
// Write `name`
write!( self.out, " {}",
&self.names[&NameKey::StructMember(handle, index as u32)]
)?; // Write [size] self.write_array_size(module, base, size)?;
} // We treat matrices of the form `matCx2` as a sequence of C `vec2`s. // See the module-level block comment in mod.rs for details.
TypeInner::Matrix {
rows,
columns,
scalar,
} if member.binding.is_none() && rows == crate::VectorSize::Bi => { let vec_ty = TypeInner::Vector { size: rows, scalar }; let field_name_key = NameKey::StructMember(handle, index as u32);
for i in0..columns as u8 { if i != 0 {
write!(self.out, "; ")?;
} self.write_value_type(module, &vec_ty)?;
write!(self.out, " {}_{}", &self.names[&field_name_key], i)?;
}
}
_ => { // Write modifier before type iflet Some(ref binding) = member.binding { self.write_modifier(binding)?;
}
// Even though Naga IR matrices are column-major, we must describe // matrices passed from the CPU as being in row-major order. // See the module-level block comment in mod.rs for details. iflet TypeInner::Matrix { .. } = module.types[member.ty].inner {
write!(self.out, "row_major ")?;
}
// Write the member type and name self.write_type(module, member.ty)?;
write!( self.out, " {}",
&self.names[&NameKey::StructMember(handle, index as u32)]
)?;
}
}
// add padding at the end since sizes of types don't get rounded up to their alignment in HLSL if members.last().unwrap().binding.is_none() && span > last_offset { let padding = (span - last_offset) / 4; for i in0..padding {
writeln!(self.out, "{}int _end_pad_{};", back::INDENT, i)?;
}
}
writeln!(self.out, "}};")?;
Ok(())
}
/// Helper method used to write global/structs non image/sampler types /// /// # Notes /// Adds no trailing or leading whitespace pub(super) fn write_global_type(
&mutself,
module: &Module,
ty: Handle<crate::Type>,
) -> BackendResult { let matrix_data = get_inner_matrix_data(module, ty);
// We treat matrices of the form `matCx2` as a sequence of C `vec2`s. // See the module-level block comment in mod.rs for details. iflet Some(MatrixType {
columns,
rows: crate::VectorSize::Bi,
width,
}) = matrix_data
{
write!(self.out, "__mat{}x2_f{}", columns as u8, width * 8)?;
} else { // Even though Naga IR matrices are column-major, we must describe // matrices passed from the CPU as being in row-major order. // See the module-level block comment in mod.rs for details. if matrix_data.is_some() {
write!(self.out, "row_major ")?;
}
self.write_type(module, ty)?;
}
Ok(())
}
/// Helper method used to write non image/sampler types /// /// # Notes /// Adds no trailing or leading whitespace pub(super) fn write_type(&mutself, module: &Module, ty: Handle<crate::Type>) -> BackendResult { let inner = &module.types[ty].inner; match *inner {
TypeInner::Struct { .. } => write!(self.out, "{}", self.names[&NameKey::Type(ty)])?, // hlsl array has the size separated from the base type
TypeInner::Array { base, .. } | TypeInner::BindingArray { base, .. } => { self.write_type(module, base)?
} ref other => self.write_value_type(module, other)?,
}
Ok(())
}
/// Helper method used to write value types /// /// # Notes /// Adds no trailing or leading whitespace pub(super) fn write_value_type(&mutself, module: &Module, inner: &TypeInner) -> BackendResult { match *inner {
TypeInner::Scalar(scalar) | TypeInner::Atomic(scalar) => {
write!(self.out, "{}", scalar.to_hlsl_str()?)?;
}
TypeInner::Vector { size, scalar } => {
write!( self.out, "{}{}",
scalar.to_hlsl_str()?,
common::vector_size_str(size)
)?;
}
TypeInner::Matrix {
columns,
rows,
scalar,
} => { // The IR supports only float matrix // https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-matrix
// Because of the implicit transpose all matrices have in HLSL, we need to transpose the size as well.
write!( self.out, "{}{}x{}",
scalar.to_hlsl_str()?,
common::vector_size_str(columns),
common::vector_size_str(rows),
)?;
}
TypeInner::Image {
dim,
arrayed,
class,
} => { self.write_image_type(dim, arrayed, class)?;
}
TypeInner::Sampler { comparison } => { let sampler = if comparison { "SamplerComparisonState"
} else { "SamplerState"
};
write!(self.out, "{sampler}")?;
} // HLSL arrays are written as `type name[size]` // Current code is written arrays only as `[size]` // Base `type` and `name` should be written outside
TypeInner::Array { base, size, .. } | TypeInner::BindingArray { base, size } => { self.write_array_size(module, base, size)?;
}
TypeInner::AccelerationStructure { .. } => {
write!(self.out, "RaytracingAccelerationStructure")?;
}
TypeInner::RayQuery { .. } => { // these are constant flags, there are dynamic flags also but constant flags are not supported by naga
write!(self.out, "RayQuery<RAY_FLAG_NONE>")?;
}
_ => return Err(Error::Unimplemented(format!("write_value_type {inner:?}"))),
}
Ok(())
}
/// Helper method used to write functions /// # Notes /// Ends in a newline fn write_function(
&mutself,
module: &Module,
name: &str,
func: &crate::Function,
func_ctx: &back::FunctionCtx<'_>,
info: &valid::FunctionInfo,
header: String,
) -> BackendResult { // Function Declaration Syntax - https://docs.microsoft.com/en-us/windows/win32/direct3dhlsl/dx-graphics-hlsl-function-syntax
self.update_expressions_to_bake(module, func, info); let ep = match func_ctx.ty {
back::FunctionType::EntryPoint(idx) => Some(&module.entry_points[idx as usize]),
back::FunctionType::Function(_) => None,
};
let nested = matches!(
ep,
Some(crate::EntryPoint {
stage: ShaderStage::Task | ShaderStage::Mesh,
..
})
); if !nested {
write!(self.out, "{header}")?;
}
iflet Some(ref result) = func.result { // Write typedef if return type is an array let array_return_type = match module.types[result.ty].inner {
TypeInner::Array { base, size, .. } => { let array_return_type = self.namer.call(&format!("ret_{name}"));
write!(self.out, "typedef ")?; self.write_type(module, result.ty)?;
write!(self.out, " {array_return_type}")?; self.write_array_size(module, base, size)?;
writeln!(self.out, ";")?;
Some(array_return_type)
}
_ => None,
};
let needs_local_invocation_index_name = need_workgroup_variables_initialization || nested; letmut local_invocation_index_name = None; // For nested entry points, collect arg names as we write them so that // write_nested_function_outer can pass the exact same names to the call site. letmut nested_wgsl_args: Vec<String> = Vec::new(); letmut nested_task_payload_name: Option<String> = None; // Write function arguments for non entry point functions match func_ctx.ty {
back::FunctionType::Function(handle) => { for (index, arg) in func.arguments.iter().enumerate() {
write!(self.out, "{}", separator())?; self.write_function_argument(module, handle, arg, index)?;
} // If this reads a task payload variable the variable needs to be passed as an `in` argument for (var_handle, var) in module.global_variables.iter() { let uses = info[var_handle]; if uses.contains(valid::GlobalUse::READ)
&& !uses.contains(valid::GlobalUse::WRITE)
&& var.space == crate::AddressSpace::TaskPayload
{ self.function_task_payload_var.insert(handle, var_handle);
write!(self.out, "{}in ", separator())?;
self.write_type(module, var.ty)?; let name = &self.names[&NameKey::GlobalVariable(var_handle)];
write!(self.out, " {name}")?; break;
}
}
}
back::FunctionType::EntryPoint(ep_index) => { let ep = &module.entry_points[ep_index as usize]; iflet Some(ref ep_input) = self.entry_point_io.get(&(ep_index as usize)).unwrap().input
{
write!(self.out, "{} {}", ep_input.ty_name, ep_input.arg_name)?;
separator();
nested_wgsl_args.push(ep_input.arg_name.clone());
} else { let stage = ep.stage; for (index, arg) in func.arguments.iter().enumerate() {
write!(self.out, "{}", separator())?; self.write_type(module, arg.ty)?;
let argument_name =
&self.names[&NameKey::EntryPointArgument(ep_index, index as u32)];
self.write_semantic(&arg.binding, Some((stage, Io::Input)))?;
}
} if ep.stage == ShaderStage::Mesh { iflet Some(var_handle) = ep.task_payload { let var = &module.global_variables[var_handle];
write!(self.out, "{}in ", separator())?; self.write_type(module, var.ty)?; let arg_name = &self.names[&NameKey::GlobalVariable(var_handle)];
write!(self.out, " {arg_name}")?;
nested_task_payload_name = Some(arg_name.clone()); iflet TypeInner::Array { base, size, .. } = module.types[var.ty].inner { self.write_array_size(module, base, size)?;
}
}
} if needs_local_invocation_index_name && local_invocation_index_name.is_none() { let name = self.namer.call("local_invocation_index");
write!(self.out, "{}uint {name}", separator())?;
write!(self.out, " : SV_GroupIndex")?;
local_invocation_index_name = Some(name);
}
}
} // Ends of arguments
write!(self.out, ")")?;
// Write semantic if it present iflet back::FunctionType::EntryPoint(index) = func_ctx.ty { let stage = module.entry_points[index as usize].stage; iflet Some(crate::FunctionResult { ref binding, .. }) = func.result { self.write_semantic(binding, Some((stage, Io::Output)))?;
}
}
// Function body start
writeln!(self.out)?;
writeln!(self.out, "{{")?;
if need_workgroup_variables_initialization && !nested { let back::FunctionType::EntryPoint(index) = func_ctx.ty else {
unreachable!();
};
writeln!( self.out, "{}if ({} == 0) {{",
back::INDENT, // need_workgroup_variables_initialization forces this to be written // if the user doesn't specify it (so this must be Some())
local_invocation_index_name.as_ref().unwrap(),
)?; self.write_workgroup_variables_initialization(
func_ctx,
module,
module.entry_points[index as usize].stage,
)?;
// Write function local variables for (handle, local) in func.local_variables.iter() { // Write indentation (only for readability)
write!(self.out, "{}", back::INDENT)?;
// Write the local name // The leading space is important self.write_type(module, local.ty)?;
write!(self.out, " {}", self.names[&func_ctx.name_key(handle)])?; // Write size for array type iflet TypeInner::Array { base, size, .. } = module.types[local.ty].inner { self.write_array_size(module, base, size)?;
}
let is_ray_query = match module.types[local.ty].inner { // from https://microsoft.github.io/DirectX-Specs/d3d/Raytracing.html#tracerayinline-example-1 it seems that ray queries shouldn't be zeroed
TypeInner::RayQuery { .. } => true,
_ => {
write!(self.out, " = ")?; // Write the local initializer if needed iflet Some(init) = local.init { self.write_expr(module, init, func_ctx)?;
} else { // Zero initialize local variables self.write_default_init(module, local.ty)?;
} false
}
}; // Finish the local with `;` and add a newline (only for readability)
writeln!(self.out, ";")?; // If it's a ray query, we also want a tracker variable if is_ray_query {
write!(self.out, "{}", back::INDENT)?; self.write_value_type(module, &TypeInner::Scalar(Scalar::U32))?;
writeln!( self.out, " {RAY_QUERY_TRACKER_VARIABLE_PREFIX}{} = 0;", self.names[&func_ctx.name_key(handle)]
)?;
}
}
if !func.local_variables.is_empty() {
writeln!(self.out)?;
}
// Write the function body (statement list) for sta in func.body.iter() { // The indentation should always be 1 when writing the function body self.write_stmt(module, sta, func_ctx, back::Level(1))?;
}
writeln!(self.out, "}}")?;
if nested { self.write_nested_function_outer(
module,
func_ctx,
&header,
name,
need_workgroup_variables_initialization,
&nested_name,
ep.unwrap(),
NestedEntryPointArgs {
user_args: nested_wgsl_args,
task_payload: nested_task_payload_name, // guaranteed to be set for nested functions (task/mesh shaders)
local_invocation_index: local_invocation_index_name.unwrap(),
},
)?;
}
self.named_expressions.clear();
Ok(())
}
fn write_function_argument(
&mutself,
module: &Module,
handle: Handle<crate::Function>,
arg: &crate::FunctionArgument,
index: usize,
) -> BackendResult { // External texture arguments must be expanded into separate // arguments for each plane and the params buffer. iflet TypeInner::Image {
class: crate::ImageClass::External,
..
} = module.types[arg.ty].inner
{ returnself.write_function_external_texture_argument(module, handle, index);
}
// Write argument type let arg_ty = match module.types[arg.ty].inner { // pointers in function arguments are expected and resolve to `inout`
TypeInner::Pointer { base, .. } => { //TODO: can we narrow this down to just `in` when possible?
write!(self.out, "inout ")?;
base
}
_ => arg.ty,
}; self.write_type(module, arg_ty)?;
let argument_name = &self.names[&NameKey::FunctionArgument(handle, index as u32)];
for (handle, var) in vars { let name = &self.names[&NameKey::GlobalVariable(handle)];
write!(self.out, "{}{} = ", back::Level(2), name)?; self.write_default_init(module, var.ty)?;
writeln!(self.out, ";")?;
}
Ok(())
}
/// Helper method used to write switches fn write_switch(
&mutself,
module: &Module,
func_ctx: &back::FunctionCtx<'_>,
level: back::Level,
selector: Handle<crate::Expression>,
cases: &[crate::SwitchCase],
) -> BackendResult { // Write all cases let indent_level_1 = level.next(); let indent_level_2 = indent_level_1.next();
// See docs of `back::continue_forward` module. iflet Some(variable) = self.continue_ctx.enter_switch(&mutself.namer) {
writeln!(self.out, "{level}bool {variable} = false;",)?;
};
// Check if there is only one body, by seeing if all except the last case are fall through // with empty bodies. FXC doesn't handle these switches correctly, so // we generate a `do {} while(false);` loop instead. There must be a default case, so there // is no need to check if one of the cases would have matched. let one_body = cases
.iter()
.rev()
.skip(1)
.all(|case| case.fall_through && case.body.is_empty()); if one_body { // Start the do-while
writeln!(self.out, "{level}do {{")?; // Note: Expressions have no side-effects so we don't need to emit selector expression.
// Body iflet Some(case) = cases.last() { for sta in case.body.iter() { self.write_stmt(module, sta, func_ctx, indent_level_1)?;
}
} // End do-while
writeln!(self.out, "{level}}} while(false);")?;
} else { // Start the switch
write!(self.out, "{level}")?;
write!(self.out, "switch(")?; self.write_expr(module, selector, func_ctx)?;
writeln!(self.out, ") {{")?;
for (i, case) in cases.iter().enumerate() { match case.value { crate::SwitchValue::I32(value) => {
write!(self.out, "{indent_level_1}case {value}:")?
} crate::SwitchValue::U32(value) => {
write!(self.out, "{indent_level_1}case {value}u:")?
} crate::SwitchValue::Default => write!(self.out, "{indent_level_1}default:")?,
}
// The new block is not only stylistic, it plays a role here: // We might end up having to write the same case body // multiple times due to FXC not supporting fallthrough. // Therefore, some `Expression`s written by `Statement::Emit` // will end up having the same name (`_expr<handle_index>`). // So we need to put each case in its own scope. let write_block_braces = !(case.fall_through && case.body.is_empty()); if write_block_braces {
writeln!(self.out, " {{")?;
} else {
writeln!(self.out)?;
}
// Although FXC does support a series of case clauses before // a block[^yes], it does not support fallthrough from a // non-empty case block to the next[^no]. If this case has a // non-empty body with a fallthrough, emulate that by // duplicating the bodies of all the cases it would fall // into as extensions of this case's own body. This makes // the HLSL output potentially quadratic in the size of the // Naga IR. // // [^yes]: ```hlsl // case 1: // case 2: do_stuff() // ``` // [^no]: ```hlsl // case 1: do_this(); // case 2: do_that(); // ``` if case.fall_through && !case.body.is_empty() { let curr_len = i + 1; let end_case_idx = curr_len
+ cases
.iter()
.skip(curr_len)
.position(|case| !case.fall_through)
.unwrap(); let indent_level_3 = indent_level_2.next(); for case in &cases[i..=end_case_idx] {
writeln!(self.out, "{indent_level_2}{{")?; let prev_len = self.named_expressions.len(); for sta in case.body.iter() { self.write_stmt(module, sta, func_ctx, indent_level_3)?;
} // Clear all named expressions that were previously inserted by the statements in the block self.named_expressions.truncate(prev_len);
writeln!(self.out, "{indent_level_2}}}")?;
}
let last_case = &cases[end_case_idx]; if last_case.body.last().is_none_or(|s| !s.is_terminator()) {
writeln!(self.out, "{indent_level_2}break;")?;
}
} else { for sta in case.body.iter() { self.write_stmt(module, sta, func_ctx, indent_level_2)?;
} if !case.fall_through && case.body.last().is_none_or(|s| !s.is_terminator()) {
writeln!(self.out, "{indent_level_2}break;")?;
}
}
if write_block_braces {
writeln!(self.out, "{indent_level_1}}}")?;
}
}
writeln!(self.out, "{level}}}")?;
}
// Handle any forwarded continue statements. use back::continue_forward::ExitControlFlow; let op = matchself.continue_ctx.exit_switch() {
ExitControlFlow::None => None,
ExitControlFlow::Continue { variable } => Some(("continue", variable)),
ExitControlFlow::Break { variable } => Some(("break", variable)),
}; iflet Some((control_flow, variable)) = op {
writeln!(self.out, "{level}if ({variable}) {{")?;
writeln!(self.out, "{indent_level_1}{control_flow};")?;
writeln!(self.out, "{level}}}")?;
}
match *stmt {
Statement::Emit(ref range) => { for handle in range.clone() { let ptr_class = func_ctx.resolve_type(handle, &module.types).pointer_space(); let expr_name = if ptr_class.is_some() { // HLSL can't save a pointer-valued expression in a variable, // but we shouldn't ever need to: they should never be named expressions, // and none of the expression types flagged by bake_ref_count can be pointer-valued.
None
} elseiflet Some(name) = func_ctx.named_expressions.get(&handle) { // Front end provides names for all variables at the start of writing. // But we write them to step by step. We need to recache them // Otherwise, we could accidentally write variable name instead of full expression. // Also, we use sanitized names! It defense backend from generating variable with name from reserved keywords.
Some(self.namer.call(name))
} elseifself.need_bake_expressions.contains(&handle) {
Some(Baked(handle).to_string())
} else {
None
};
iflet Some(name) = expr_name {
write!(self.out, "{level}")?; self.write_named_expr(module, handle, name, handle, func_ctx)?;
}
}
} // TODO: copy-paste from glsl-out
Statement::Block(ref block) => {
write!(self.out, "{level}")?;
writeln!(self.out, "{{")?; for sta in block.iter() { // Increase the indentation to help with readability self.write_stmt(module, sta, func_ctx, level.next())?
}
writeln!(self.out, "{level}}}")?
} // TODO: copy-paste from glsl-out
Statement::If {
condition, ref accept, ref reject,
} => {
write!(self.out, "{level}")?;
write!(self.out, "if (")?; self.write_expr(module, condition, func_ctx)?;
writeln!(self.out, ") {{")?;
let l2 = level.next(); for sta in accept { // Increase indentation to help with readability self.write_stmt(module, sta, func_ctx, l2)?;
}
// If there are no statements in the reject block we skip writing it // This is only for readability if !reject.is_empty() {
writeln!(self.out, "{level}}} else {{")?;
for sta in reject { // Increase indentation to help with readability self.write_stmt(module, sta, func_ctx, l2)?;
}
}
iflet TypeInner::Struct { .. } = *resolved { // We can safely unwrap here, since we now we working with struct let ty = base_ty_res.handle().unwrap(); let struct_name = &self.names[&NameKey::Type(ty)]; let variable_name = self.namer.call(&struct_name.to_lowercase());
write!(self.out, "{level}const {struct_name} {variable_name} = ",)?; self.write_expr(module, expr, func_ctx)?;
writeln!(self.out, ";")?;
// for entry point returns, we may need to reshuffle the outputs into a different struct let ep_output = match func_ctx.ty {
back::FunctionType::Function(_) => None,
back::FunctionType::EntryPoint(index) => self
.entry_point_io
.get(&(index as usize))
.unwrap()
.output
.as_ref(),
}; let final_name = match ep_output {
Some(ep_output) => { let final_name = self.namer.call(&variable_name);
write!( self.out, "{}const {} {} = {{ ",
level, ep_output.ty_name, final_name,
)?; for (index, m) in ep_output.members.iter().enumerate() { if index != 0 {
write!(self.out, ", ")?;
} let member_name = &self.names[&NameKey::StructMember(ty, m.index)];
write!(self.out, "{variable_name}.{member_name}")?;
}
writeln!(self.out, " }};")?;
final_name
}
None => variable_name,
};
writeln!(self.out, "{level}return {final_name};")?;
} else {
write!(self.out, "{level}return ")?; self.write_expr(module, expr, func_ctx)?;
writeln!(self.out, ";")?
}
}
Statement::Store { pointer, value } => { let ty_inner = func_ctx.resolve_type(pointer, &module.types); if ty_inner.is_atomic_pointer(&module.types) { let pointer_space = ty_inner.pointer_space().unwrap(); let dummy = self.namer.call("dummy");
write!(self.out, "{level}{{ ")?; iflet TypeInner::Pointer { base, .. } = *ty_inner { self.write_value_type(module, &module.types[base].inner)?;
}
write!(self.out, " {dummy} = 0; ")?; match pointer_space { crate::AddressSpace::WorkGroup => {
write!(self.out, "InterlockedExchange(")?; self.write_expr(module, pointer, func_ctx)?;
} crate::AddressSpace::Storage { .. } => { let var_handle = self.fill_access_chain(module, pointer, func_ctx)?; let var_name = &self.names[&NameKey::GlobalVariable(var_handle)];
write!(self.out, "{var_name}.InterlockedExchange(")?; let chain = mem::take(&mutself.temp_access_chain); self.write_storage_address(module, &chain, func_ctx)?; self.temp_access_chain = chain;
}
_ => unreachable!(),
}
write!(self.out, ", ")?; self.write_expr(module, value, func_ctx)?;
writeln!(self.out, ", {dummy}); }}")?;
} elseiflet Some(crate::AddressSpace::Storage { .. }) = ty_inner.pointer_space() { let var_handle = self.fill_access_chain(module, pointer, func_ctx)?; self.write_storage_store(
module,
var_handle,
StoreValue::Expression(value),
func_ctx,
level,
None,
)?;
} else { // We treat matrices of the form `matCx2` as a sequence of C `vec2`s. // See the module-level block comment in mod.rs for details. // // We handle matrix Stores here directly (including sub accesses for Vectors and Scalars). // Loads are handled by `Expression::AccessIndex` (since sub accesses work fine for Loads). enum MatrixAccess {
Direct {
base: Handle<crate::Expression>,
index: u32,
}, Struct {
columns: crate::VectorSize,
width: u8,
base: Handle<crate::Expression>,
},
}
let get_members = |expr: Handle<crate::Expression>| { let resolved = func_ctx.resolve_type(expr, &module.types); match *resolved {
TypeInner::Pointer { base, .. } => match module.types[base].inner {
TypeInner::Struct { ref members, .. } => Some(members),
_ => None,
},
_ => None,
}
};
// We cast the RHS of this store in cases where the LHS // is a struct member with type: // - matCx2 or // - a (possibly nested) array of matCx2's iflet Some(MatrixType {
columns,
rows: crate::VectorSize::Bi,
width,
}) = get_inner_matrix_of_struct_array_member(
module, pointer, func_ctx, false,
) { letmut resolved = func_ctx.resolve_type(pointer, &module.types); iflet TypeInner::Pointer { base, .. } = *resolved {
resolved = &module.types[base].inner;
}
self.write_expr(module, value, func_ctx)?;
writeln!(self.out, ");")?;
}
Statement::WorkGroupUniformLoad { pointer, result } => { self.write_control_barrier(crate::Barrier::WORK_GROUP, level)?;
write!(self.out, "{level}")?; let name = Baked(result).to_string(); self.write_named_expr(module, pointer, name, result, func_ctx)?;
self.write_control_barrier(crate::Barrier::WORK_GROUP, level)?;
}
Statement::Switch {
selector, ref cases,
} => { self.write_switch(module, func_ctx, level, selector, cases)?;
}
Statement::RayQuery { query, ref fun } => { // There are three possibilities for a ptr to be: // 1. A variable // 2. A function argument // 3. part of a struct // // 2 and 3 are not possible, a ray query (in naga IR) // is not allowed to be passed into a function, and // all languages disallow it in a struct (you get fun results if // you try it :) ). // // Therefore, the ray query expression must be a variable. letcrate::Expression::LocalVariable(query_var) = func_ctx.expressions[query] else {
unreachable!()
};
let tracker_expr_name = format!( "{RAY_QUERY_TRACKER_VARIABLE_PREFIX}{}", self.names[&func_ctx.name_key(query_var)]
);
pub(super) fn write_literal(&mutself, literal: crate::Literal) -> BackendResult { match literal { crate::Literal::F64(value) => write!(self.out, "{value:?}L")?, crate::Literal::F32(value) => write!(self.out, "{value:?}")?, crate::Literal::F16(value) => write!(self.out, "{value:?}h")?, crate::Literal::U16(value) => write!(self.out, "uint16_t({value})")?, crate::Literal::I16(value) => write!(self.out, "int16_t({value})")?, crate::Literal::U32(value) => write!(self.out, "{value}u")?, // `-2147483648` is parsed by some compilers as unary negation of // positive 2147483648, which is too large for an int, causing // issues for some compilers. Neither DXC nor FXC appear to have // this problem, but this is not specified and could change. We // therefore use `-2147483647 - 1` as a precaution. crate::Literal::I32(value) if value == i32::MIN => {
write!(self.out, "int({} - 1)", value + 1)?
} // HLSL has no suffix for explicit i32 literals, but not using any suffix // makes the type ambiguous which prevents overload resolution from // working. So we explicitly use the int() constructor syntax. crate::Literal::I32(value) => write!(self.out, "int({value})")?, crate::Literal::U64(value) => write!(self.out, "{value}uL")?, // I64 version of the minimum I32 value issue described above. crate::Literal::I64(value) if value == i64::MIN => {
write!(self.out, "({}L - 1L)", value + 1)?;
} crate::Literal::I64(value) => write!(self.out, "{value}L")?, crate::Literal::Bool(value) => write!(self.out, "{value}")?, crate::Literal::AbstractInt(_) | crate::Literal::AbstractFloat(_) => { return Err(Error::Custom( "Abstract types should not appear in IR presented to backends".into(),
));
}
}
Ok(())
}
// Handle the special semantics of vertex_index/instance_index let ff_input = ifself.options.special_constants_binding.is_some() {
func_ctx.is_fixed_function_input(expr, module)
} else {
None
}; let closing_bracket = match ff_input {
Some(crate::BuiltIn::VertexIndex) => {
write!(self.out, "({SPECIAL_CBUF_VAR}.{SPECIAL_FIRST_VERTEX} + ")?; ")"
}
Some(crate::BuiltIn::InstanceIndex) => {
write!(self.out, "({SPECIAL_CBUF_VAR}.{SPECIAL_FIRST_INSTANCE} + ",)?; ")"
}
Some(crate::BuiltIn::NumWorkGroups) => { // Note: despite their names (`FIRST_VERTEX` and `FIRST_INSTANCE`), // in compute shaders the special constants contain the number // of workgroups, which we are using here.
write!( self.out, "uint3({SPECIAL_CBUF_VAR}.{SPECIAL_FIRST_VERTEX}, {SPECIAL_CBUF_VAR}.{SPECIAL_FIRST_INSTANCE}, {SPECIAL_CBUF_VAR}.{SPECIAL_OTHER})",
)?; return Ok(());
}
_ => "",
};
match *expression {
Expression::Literal(_)
| Expression::Constant(_)
| Expression::ZeroValue(_)
| Expression::Compose { .. }
| Expression::Splat { .. } => { self.write_possibly_const_expression(
module,
expr,
func_ctx.expressions,
|writer, expr| writer.write_expr(module, expr, func_ctx),
)?;
}
Expression::Override(_) => return Err(Error::Override), // Avoid undefined behaviour for addition, subtraction, and // multiplication of signed integers by casting operands to // unsigned, performing the operation, then casting the result back // to signed. // TODO(#7109): This relies on the asint()/asuint() functions which only work // for 32-bit types, so we must find another solution for different bit widths.
Expression::Binary {
op:
op @ crate::BinaryOperator::Add
| op @ crate::BinaryOperator::Subtract
| op @ crate::BinaryOperator::Multiply,
left,
right,
} if matches!(
func_ctx.resolve_type(expr, &module.types).scalar(),
Some(Scalar::I32)
) =>
{
write!(self.out, "asint(asuint(",)?; self.write_expr(module, left, func_ctx)?;
write!(self.out, ") {} asuint(", back::binary_operation_str(op))?; self.write_expr(module, right, func_ctx)?;
write!(self.out, "))")?;
} // All of the multiplication can be expressed as `mul`, // except vector * vector, which needs to use the "*" operator.
Expression::Binary {
op: crate::BinaryOperator::Multiply,
left,
right,
} if func_ctx.resolve_type(left, &module.types).is_matrix()
|| func_ctx.resolve_type(right, &module.types).is_matrix() =>
{ // We intentionally flip the order of multiplication as our matrices are implicitly transposed.
write!(self.out, "mul(")?; self.write_expr(module, right, func_ctx)?;
write!(self.out, ", ")?; self.write_expr(module, left, func_ctx)?;
write!(self.out, ")")?;
}
// WGSL says that floating-point division by zero should return // infinity. Microsoft's Direct3D 11 functional specification // (https://microsoft.github.io/DirectX-Specs/d3d/archive/D3D11_3_FunctionalSpec.htm) // says: // // Divide by 0 produces +/- INF, except 0/0 which results in NaN. // // which is what we want. The DXIL specification for the FDiv // instruction corroborates this: // // https://github.com/microsoft/DirectXShaderCompiler/blob/main/docs/DXIL.rst#fdiv
Expression::Binary {
op: crate::BinaryOperator::Divide,
left,
right,
} if matches!(
func_ctx.resolve_type(expr, &module.types).scalar_kind(),
Some(ScalarKind::Sint | ScalarKind::Uint)
) =>
{
write!(self.out, "{DIV_FUNCTION}(")?; self.write_expr(module, left, func_ctx)?;
write!(self.out, ", ")?; self.write_expr(module, right, func_ctx)?;
write!(self.out, ")")?;
}
Expression::Binary { op, left, right } => {
write!(self.out, "(")?; self.write_expr(module, left, func_ctx)?;
write!(self.out, " {} ", back::binary_operation_str(op))?; self.write_expr(module, right, func_ctx)?;
write!(self.out, ")")?;
}
Expression::Access { base, index } => { iflet Some(crate::AddressSpace::Storage { .. }) =
func_ctx.resolve_type(expr, &module.types).pointer_space()
{ // do nothing, the chain is written on `Load`/`Store`
} else { // We use the function __get_col_of_matCx2 here in cases // where `base`s type resolves to a matCx2 and is part of a // struct member with type of (possibly nested) array of matCx2's. // // Note that this only works for `Load`s and we handle // `Store`s differently in `Statement::Store`. iflet Some(MatrixType {
columns,
rows: crate::VectorSize::Bi,
width,
}) = get_inner_matrix_of_struct_array_member(module, base, func_ctx, true)
.or_else(|| {
get_inner_matrix_of_global_uniform(module, base, func_ctx, true)
})
{
write!( self.out, "__get_col_of_mat{}x2_f{}(",
columns as u8,
width * 8
)?; self.write_expr(module, base, func_ctx)?;
write!(self.out, ", ")?; self.write_expr(module, index, func_ctx)?;
write!(self.out, ")")?; return Ok(());
}
let resolved = func_ctx.resolve_type(base, &module.types);
let (indexing_binding_array, non_uniform_qualifier) = match *resolved {
TypeInner::BindingArray { .. } => { let uniformity = &func_ctx.info[index].uniformity;
let needs_bound_check = self.options.restrict_indexing
&& !indexing_binding_array
&& match resolved.pointer_space() {
Some( crate::AddressSpace::Function
| crate::AddressSpace::Private
| crate::AddressSpace::WorkGroup
| crate::AddressSpace::Immediate
| crate::AddressSpace::TaskPayload
| crate::AddressSpace::RayPayload
| crate::AddressSpace::IncomingRayPayload,
)
| None => true,
Some(crate::AddressSpace::Uniform) => { // check if BindTarget.restrict_indexing is set, this is used for dynamic buffers let var_handle = self.fill_access_chain(module, base, func_ctx)?; let bind_target = self
.options
.resolve_resource_binding(
module.global_variables[var_handle]
.binding
.as_ref()
.unwrap(),
)
.unwrap();
bind_target.restrict_indexing
}
Some( crate::AddressSpace::Handle | crate::AddressSpace::Storage { .. },
) => unreachable!(),
}; // Decide whether this index needs to be clamped to fall within range. let restriction_needed = if needs_bound_check {
index::access_needs_check(
base,
index::GuardedIndex::Expression(index),
module,
func_ctx.expressions,
func_ctx.info,
)
} else {
None
}; iflet Some(limit) = restriction_needed {
write!(self.out, "min(uint(")?; self.write_expr(module, index, func_ctx)?;
write!(self.out, "), ")?; match limit {
index::IndexableLength::Known(limit) => {
write!(self.out, "{}u", limit - 1)?;
}
index::IndexableLength::Dynamic => unreachable!(),
}
write!(self.out, ")")?;
} else { if non_uniform_qualifier {
write!(self.out, "NonUniformResourceIndex(")?;
} iflet Some(ref info) = array_sampler_info {
write!( self.out, "{}[{} + ",
info.sampler_index_buffer_name, info.binding_array_base_index_name,
)?;
} self.write_expr(module, index, func_ctx)?; if array_sampler_info.is_some() {
write!(self.out, "]")?;
} if non_uniform_qualifier {
write!(self.out, ")")?;
}
}
write!(self.out, "]")?;
}
}
Expression::AccessIndex { base, index } => { iflet Some(crate::AddressSpace::Storage { .. }) =
func_ctx.resolve_type(expr, &module.types).pointer_space()
{ // do nothing, the chain is written on `Load`/`Store`
} else { // See if we need to write the matrix column access in a // special way since the type of `base` is our special // __matCx2 struct. iflet Some(MatrixType {
rows: crate::VectorSize::Bi,
..
}) = get_inner_matrix_of_struct_array_member(module, base, func_ctx, true)
.or_else(|| {
get_inner_matrix_of_global_uniform(module, base, func_ctx, true)
})
{ self.write_expr(module, base, func_ctx)?;
write!(self.out, "._{index}")?; return Ok(());
}
let base_ty_res = &func_ctx.info[base].ty; letmut resolved = base_ty_res.inner_with(&module.types); let base_ty_handle = match *resolved {
TypeInner::Pointer { base, .. } => {
resolved = &module.types[base].inner;
Some(base)
}
_ => base_ty_res.handle(),
};
// We treat matrices of the form `matCx2` as a sequence of C `vec2`s. // See the module-level block comment in mod.rs for details. // // We handle matrix reconstruction here for Loads. // Stores are handled directly by `Statement::Store`. iflet TypeInner::Struct { ref members, .. } = *resolved { let member = &members[index as usize];
match module.types[member.ty].inner {
TypeInner::Matrix {
rows: crate::VectorSize::Bi,
..
} if member.binding.is_none() => { let ty = base_ty_handle.unwrap(); self.write_wrapped_struct_matrix_get_function_name(
WrappedStructMatrixAccess { ty, index },
)?;
write!(self.out, "(")?; self.write_expr(module, base, func_ctx)?;
write!(self.out, ")")?; return Ok(());
}
_ => {}
}
}
let array_sampler_info = self.sampler_binding_array_info_from_expression(
module, func_ctx, base, resolved,
);
match *resolved { // We specifically lift the ValuePointer to this case. While `[0]` is valid // HLSL for any vector behind a value pointer, FXC completely miscompiles // it and generates completely nonsensical DXBC. // // See https://github.com/gfx-rs/naga/issues/2095 for more details.
TypeInner::Vector { .. } | TypeInner::ValuePointer { .. } => { // Write vector access as a swizzle
write!(self.out, ".{}", back::COMPONENTS[index as usize])?
}
TypeInner::Matrix { .. }
| TypeInner::Array { .. }
| TypeInner::BindingArray { .. } => { iflet Some(ref info) = array_sampler_info {
write!( self.out, "[{} + {index}]",
info.binding_array_base_index_name
)?;
} else {
write!(self.out, "[{index}]")?;
}
}
TypeInner::Struct { .. } => { // This will never panic in case the type is a `Struct`, this is not true // for other types so we can only check while inside this match arm let ty = base_ty_handle.unwrap();
write!( self.out, ".{}",
&self.names[&NameKey::StructMember(ty, index)]
)?
} ref other => return Err(Error::Custom(format!("Cannot index {other:?}"))),
}
if array_sampler_info.is_some() {
write!(self.out, "]")?;
}
}
}
Expression::FunctionArgument(pos) => { let ty = func_ctx.resolve_type(expr, &module.types);
// We know that any external texture function argument has been expanded into // separate consecutive arguments for each plane and the parameters buffer. And we // also know that external textures can only ever be used as an argument to another // function. Therefore we can simply emit each of the expanded arguments in a // consecutive comma-separated list. iflet TypeInner::Image {
class: crate::ImageClass::External,
..
} = *ty
{ let plane_names = [0, 1, 2].map(|i| {
&self.names[&func_ctx
.external_texture_argument_key(pos, ExternalTextureNameKey::Plane(i))]
}); let params_name = &self.names[&func_ctx
.external_texture_argument_key(pos, ExternalTextureNameKey::Params)];
write!( self.out, "{}, {}, {}, {}",
plane_names[0], plane_names[1], plane_names[2], params_name
)?;
} else { let key = func_ctx.argument_key(pos); let name = &self.names[&key];
write!(self.out, "{name}")?;
}
}
Expression::ImageSample {
coordinate,
image,
sampler,
clamp_to_edge: true,
gather: None,
array_index: None,
offset: None,
level: crate::SampleLevel::Zero,
depth_ref: None,
} => {
write!(self.out, "{IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION}(")?; self.write_expr(module, image, func_ctx)?;
write!(self.out, ", ")?; self.write_expr(module, sampler, func_ctx)?;
write!(self.out, ", ")?; self.write_expr(module, coordinate, func_ctx)?;
write!(self.out, ")")?;
}
Expression::ImageSample {
image,
sampler,
gather,
coordinate,
array_index,
offset,
level,
depth_ref,
clamp_to_edge,
} => { if clamp_to_edge { return Err(Error::Custom( "ImageSample::clamp_to_edge should have been validated out".to_string(),
));
}
// In the case of binding arrays of samplers, we need to not write anything // as the we are in the wrong position to fully write the expression. // // The entire writing is done by AccessIndex. let is_binding_array_of_samplers = match *ty {
TypeInner::BindingArray { base, .. } => { let base_ty = &module.types[base].inner;
matches!(*base_ty, TypeInner::Sampler { .. })
}
_ => false,
};
let is_storage_space =
matches!(global_variable.space, crate::AddressSpace::Storage { .. });
// Our external texture global variable has been expanded into multiple // global variables, one for each plane and the parameters buffer. // External textures can only ever be used as arguments to a function // call, and we know that an external texture argument to any function // will have been expanded to separate consecutive arguments for each // plane and the parameters buffer. Therefore we can simply emit each of // the expanded global variables in a consecutive comma-separated list. iflet TypeInner::Image {
class: crate::ImageClass::External,
..
} = *ty
{ let plane_names = [0, 1, 2].map(|i| {
&self.names[&NameKey::ExternalTextureGlobalVariable(
handle,
ExternalTextureNameKey::Plane(i),
)]
}); let params_name = &self.names[&NameKey::ExternalTextureGlobalVariable(
handle,
ExternalTextureNameKey::Params,
)];
write!( self.out, "{}, {}, {}, {}",
plane_names[0], plane_names[1], plane_names[2], params_name
)?;
} elseif !is_binding_array_of_samplers && !is_storage_space { let name = &self.names[&NameKey::GlobalVariable(handle)];
write!(self.out, "{name}")?;
}
}
Expression::LocalVariable(handle) => {
write!(self.out, "{}", self.names[&func_ctx.name_key(handle)])?
}
Expression::Load { pointer } => { match func_ctx
.resolve_type(pointer, &module.types)
.pointer_space()
{
Some(crate::AddressSpace::Storage { .. }) => { let var_handle = self.fill_access_chain(module, pointer, func_ctx)?; let result_ty = func_ctx.info[expr].ty.clone(); self.write_storage_load(module, var_handle, result_ty, func_ctx)?;
}
_ => { letmut close_paren = false;
// We cast the value loaded to a native HLSL floatCx2 // in cases where it is of type: // - __matCx2 or // - a (possibly nested) array of __matCx2's iflet Some(MatrixType {
rows: crate::VectorSize::Bi,
..
}) = get_inner_matrix_of_struct_array_member(
module, pointer, func_ctx, false,
)
.or_else(|| {
get_inner_matrix_of_global_uniform(module, pointer, func_ctx, false)
}) { letmut resolved = func_ctx.resolve_type(pointer, &module.types); let ptr_tr = resolved.pointer_base_type(); iflet Some(ptr_ty) =
ptr_tr.as_ref().map(|tr| tr.inner_with(&module.types))
{
resolved = ptr_ty;
}
ifself.options.shader_model >= ShaderModel::V6_4 { // Intrinsics `dot4add_{i, u}8packed` are available in SM 6.4 and later. let function_name = match fun {
Function::Dot4I8Packed => "dot4add_i8packed",
Function::Dot4U8Packed => "dot4add_u8packed",
_ => unreachable!(),
};
write!(self.out, "{function_name}(")?; self.write_expr(module, arg, func_ctx)?;
write!(self.out, ", ")?; self.write_expr(module, arg1, func_ctx)?;
write!(self.out, ", 0)")?;
} else { // Fall back to a polyfill as `dot4add_u8packed` is not available.
write!(self.out, "dot(")?;
// close bracket for Load function
write!(self.out, ")")?;
if wrapping_type.is_some() {
write!(self.out, ")")?;
}
// return x component if return type is scalar iflet TypeInner::Scalar(_) = *func_ctx.resolve_type(expr, &module.types) {
write!(self.out, ".x")?;
}
Ok(())
}
/// Find the [`BindingArraySamplerInfo`] from an expression so that such an access /// can be generated later. fn sampler_binding_array_info_from_expression(
&mutself,
module: &Module,
func_ctx: &back::FunctionCtx<'_>,
base: Handle<crate::Expression>,
resolved: &TypeInner,
) -> Option<BindingArraySamplerInfo> { iflet TypeInner::BindingArray {
base: base_ty_handle,
..
} = *resolved
{ let base_ty = &module.types[base_ty_handle].inner; iflet TypeInner::Sampler { comparison, .. } = *base_ty { let base = &func_ctx.expressions[base];
ifletcrate::Expression::GlobalVariable(handle) = *base { let variable = &module.global_variables[handle];
let sampler_heap_name = match comparison { true => COMPARISON_SAMPLER_HEAP_VAR, false => SAMPLER_HEAP_VAR,
};
/// If `base` is an access chain of the form `mat`, `mat[col]`, or `mat[col][row]`, /// returns a tuple of the matrix, the column (vector) index (if present), and /// the row (scalar) index (if present). fn find_matrix_in_access_chain(
module: &Module,
base: Handle<crate::Expression>,
func_ctx: &back::FunctionCtx<'_>,
) -> Option<(Handle<crate::Expression>, Option<Index>, Option<Index>)> { letmut current_base = base; letmut vector = None; letmut scalar = None; loop { let resolved_tr = func_ctx
.resolve_type(current_base, &module.types)
.pointer_base_type(); let resolved = resolved_tr.as_ref()?.inner_with(&module.types);
/// Returns the matrix data if the access chain starting at `base`: /// - starts with an expression with resolved type of [`TypeInner::Matrix`] if `direct = true` /// - contains one or more expressions with resolved type of [`TypeInner::Array`] of [`TypeInner::Matrix`] /// - ends at an expression with resolved type of [`TypeInner::Struct`] pub(super) fn get_inner_matrix_of_struct_array_member(
module: &Module,
base: Handle<crate::Expression>,
func_ctx: &back::FunctionCtx<'_>,
direct: bool,
) -> Option<MatrixType> { letmut mat_data = None; letmut array_base = None;
/// Returns the matrix data if the access chain starting at `base`: /// - starts with an expression with resolved type of [`TypeInner::Matrix`], or /// - contains zero or more expressions with resolved type of [`TypeInner::Array`] of [`TypeInner::Matrix`] if `direct = false` /// - and ends with an [`Expression::GlobalVariable`](crate::Expression::GlobalVariable) in [`AddressSpace::Uniform`](crate::AddressSpace::Uniform) fn get_inner_matrix_of_global_uniform(
module: &Module,
base: Handle<crate::Expression>,
func_ctx: &back::FunctionCtx<'_>,
direct: bool,
) -> Option<MatrixType> { letmut mat_data = None; letmut array_base = None;
¤ Diese beiden folgenden Angebotsgruppen bietet das Unternehmen0.187Angebot
(Wie Sie bei der Firma Beratungs- und Dienstleistungen beauftragen können 2026-10-09)
¤
Die Informationen auf dieser Webseite wurden
nach bestem Wissen sorgfältig zusammengestellt. Es wird jedoch weder Vollständigkeit, noch Richtigkeit,
noch Qualität der bereit gestellten Informationen zugesichert.
Bemerkung:
Die farbliche Syntaxdarstellung und die Messung sind noch experimentell.
