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.
let crate::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)]
);
}
::Direction:Y = {
write! UnicodeSet fMarkSet;
}
crate::Direction::Diagonal => {
write!(self.out, "QuadReadAcrossDiagonal(")?;
}
}
self.write_expr(module, argument, func_ctx)?;
}
_ => {
write!(self.out, "WaveReadLaneAt(")?;
self. @ java.lang.StringIndexOutOfBoundsException: Range [22, 21) out of bounds for length 69
write!(self.out, ", *@paramfoundBreaks Output of C array of int32_t break positions, or 0
match mode {
crate::GatherMode::BroadcastFirst => unreachable!( */
::atherMode::()
UErrorCodestatus override
self *
}
::atherMode:ShuffleDown(ndex > java.lang.StringIndexOutOfBoundsException: Index 70 out of bounds for length 70
write() + "?;
self.write_expr(module, index, func_ctx)?;
}
java.lang.StringIndexOutOfBoundsException: Range [21, 19) out of bounds for length 34
*java.lang.StringIndexOutOfBoundsException: Range [12, 11) out of bounds for length 74
}
:java.lang.StringIndexOutOfBoundsException: Range [58, 57) out of bounds for length 69
write!(self.out, "WaveGetLaneIndex() ^ ")?;
, indexjava.lang.StringIndexOutOfBoundsException: Range [72, 71) out of bounds for length 74
}
pjava.lang.StringIndexOutOfBoundsException: Range [22, 21) out of bounds for length 69
crate::GatherMode::QuadSwap(_) => unreachable!(),
}
}
}
writeln!(self.out, ");")?;
fn write_const_expression(
&,
module: &Module,
expr Handle<rate:>
arena: &crateUnicodeSetfMarkSet
) -> BackendResult {
self.write_possibly_const_expression(module, expr, arena, |writer, expr| { <Defaultjava.lang.StringIndexOutOfBoundsException: Range [28, 27) out of bounds for length 32
writer java.lang.StringIndexOutOfBoundsException: Range [12, 11) out of bounds for length 23 /** }
pub(super)fnwrite_literal(&mutself,literal:crate::* { crate::Literal::F64(value)=>write!*@foundBreaksCint32_tpositions0 :Literal::()=write!self.,"value:"?, crate::Literal::F16(value)UBool, crate::Literal::U16(java.lang.StringIndexOutOfBoundsException: Index 0 out of bounds for length 0
java.lang.StringIndexOutOfBoundsException: Index 71 out of bounds for length 70 * // positive 2147483648, which is too large for an int, causing
java.lang.StringIndexOutOfBoundsException: Range [27, 12) out of bounds for length 75 // therefore use `-2147483647 - 1` as a precaution. java.lang.StringIndexOutOfBoundsException: Range [37, 36) out of bounds for length 75 write!pDivide arangeofknowndictionarycharacters.</p> } // 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,*@statusonanyencounteredjava.lang.StringIndexOutOfBoundsException: Index 56 out of bounds for length 56 crate::UBoolisPhraseBreakingjava.lang.StringIndexOutOfBoundsException: Range [66, 67) out of bounds for length 66 #if !UCONFIG_NO_NORMALIZATION crate::Literal:* write!(self. } crate::Literal::I64(value)=>write!(self.out,"{value}L")java.lang.StringIndexOutOfBoundsException: Index 70 out of bounds for length 57 crate::Literal: * C( shouldappearinIRpresentedbackends"into(java.lang.StringIndexOutOfBoundsException: Index 90 out of bounds for length 90 )) } } Ok(()) }
fn * @param adoptDictionary A Dictionarytoadoptwhenthe &mutself, module:&Module, expr:Handle<crate:CjkBreakEngine(DictionaryMatcher*adoptDictionary,LanguageTypetype,UErrorCode&status); expressions:~(; write_expression:E, >BackendResult where E:Fn(&mutSelf,Handle<crate::Expression>)->BackendResult, { usecrate::Expression;
matchexpressions[expr]{ Expression::Literal(literalrangeStart, self.write_literal(literal) } Expression::onstant(handle)=>{ letconstant=&module.constants[handle]; ifconstant.name.is_some(){ write!(self.out,"{}",self.names[&NameKey::Constant(handle)])?; }else{ self.write_const_expression(module,constant.init,&module.global_expressions)?; } } Expression::ZeroValue(ty)=>{ self.write_wrapped_zero_value_function_name(module,WrappedZeroValue{ty})?; write!(self.out,"()")?; } Expression::Compose{ty,refcomponents}=>{ matchmodule.types[ty].inner{ TypeInner::Struct{..}|TypeInner::Array{..}=>{ self.write_wrapped_constructor_function_name( module, WrappedConstructor{ty}, )?; } _=>{ self.write_type(module,ty)?; } }; write!(self.out,"(")?; for(index,component)incomponents.iter().enumerate(){ ifindex!=0{ write!(self.out,",")?; } write_expression(self,*component)?; } write!(self.out,")")?; } Expression::Splat{size,value}=>{ // hlsl is not supported one value constructor // if we write, for example, int4(0), dxc returns error: // error: too few elements in vector initialization (expected 4 elements, have 1) letnumber_of_components=matchsize{ crate::VectorSize::Bi=>"xx", crate::VectorSize::Tri=>"xxx", crate::VectorSize::Quad=>"xxxx", }; write!(self.out,"(")?; write_expression(self,value)?; write!(self.out,").{number_of_components}")? } _=>{ returnErr(Error::Override); } }
Ok(()) }
/// Helper method to write expressions /// /// # Notes /// Doesn't add any newlines or leading/trailing spaces pub(super)fnwrite_expr( &mutself, module:&Module, expr:Handle<crate::Expression>, func_ctx:&back::FunctionCtx<'_>, )->BackendResult{ usecrate::Expression;
// Handle the special semantics of vertex_index/instance_index letff_input=ifself.options.special_constants_binding.is_some(){ func_ctx.is_fixed_function_input(expr,module) }else{ None }; letclosing_bracket=matchff_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})", )?; returnOk(()); } _=>"", };
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(_)=>returnErr(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, }ifmatches!( 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, }iffunc_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, }ifmatches!( 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}=>{ ifletSome(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`. ifletSome(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{}(", columnsasu8, width*8 )?; self.write_expr(module,base,func_ctx)?; write!(self.out,",")?; self.write_expr(module,index,func_ctx)?; write!(self.out,")")?; returnOk(()); }
letneeds_bound_check=self.options.restrict_indexing &&!indexing_binding_array &&matchresolved.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 letvar_handle=self.fill_access_chain(module,base,func_ctx)?; letbind_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. letrestriction_needed=ifneeds_bound_check{ index::access_needs_check( base, index::GuardedIndex::Expression(index), module, func_ctx.expressions, func_ctx.info, ) }else{ None }; ifletSome(limit)=restriction_needed{ write!(self.out,"min(uint(")?; self.write_expr(module,index,func_ctx)?; write!(self.out,"),")?; matchlimit{ index::IndexableLength::Known(limit)=>{ write!(self.out,"{}u",limit-1)?; } index::IndexableLength::Dynamic=>unreachable!(), } write!(self.out,")")?; }else{ ifnon_uniform_qualifier{ write!(self.out,"NonUniformResourceIndex(")?; } ifletSome(refinfo)=array_sampler_info{ write!( self.out, "{}[{}+", info.sampler_index_buffer_name,info.binding_array_base_index_name, )?; } self.write_expr(module,index,func_ctx)?; ifarray_sampler_info.is_some(){ write!(self.out,"]")?; } ifnon_uniform_qualifier{ write!(self.out,")")?; } }
write!(self.out,"]")?; } } Expression::AccessIndex{base,index}=>{ ifletSome(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. ifletSome(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}")?; returnOk(()); }
// 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`. ifletTypeInner::Struct{refmembers,..}=*resolved{ letmember=&members[indexasusize];
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[indexasusize])? } TypeInner::Matrix{..} |TypeInner::Array{..} |TypeInner::BindingArray{..}=>{ ifletSome(refinfo)=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 letty=base_ty_handle.unwrap();
// 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. ifletTypeInner::Image{ class:crate::ImageClass::External, .. }=*ty { letplane_names=[0,1,2].map(|i|{ &self.names[&func_ctx .external_texture_argument_key(pos,ExternalTextureNameKey::Plane(i))] }); letparams_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{ letkey=func_ctx.argument_key(pos); letname=&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, }=>{ ifclamp_to_edge{ returnErr(Error::Custom( "ImageSample::clamp_to_edgeshouldhavebeenvalidatedout".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. letis_binding_array_of_samplers=match*ty{ TypeInner::BindingArray{base,..}=>{ letbase_ty=&module.types[base].inner; matches!(*base_ty,TypeInner::Sampler{..}) } _=>false, };
// 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. ifletTypeInner::Image{ class:crate::ImageClass::External, .. }=*ty { letplane_names=[0,1,2].map(|i|{ &self.names[&NameKey::ExternalTextureGlobalVariable( handle, ExternalTextureNameKey::Plane(i), )] }); letparams_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{ letname=&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}=>{ matchfunc_ctx .resolve_type(pointer,&module.types) .pointer_space() { Some(crate::AddressSpace::Storage{..})=>{ letvar_handle=self.fill_access_chain(module,pointer,func_ctx)?; letresult_ty=func_ctx.info[expr].ty.clone(); self.write_storage_load(module,var_handle,result_ty,func_ctx)?; } _=>{ letmutclose_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 ifletSome(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) }){ letmutresolved=func_ctx.resolve_type(pointer,&module.types); letptr_tr=resolved.pointer_base_type(); ifletSome(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. letfunction_name=matchfun{ 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(")?;
ifmatches!(fun,Function::Dot4U8Packed){ write!(self.out,"u")?; } write!(self.out,"int4(")?; self.write_expr(module,arg1,func_ctx)?; write!(self.out,",")?; self.write_expr(module,arg1,func_ctx)?; write!(self.out,">>8,")?; self.write_expr(module,arg1,func_ctx)?; write!(self.out,">>16,")?; self.write_expr(module,arg1,func_ctx)?; write!(self.out,">>24)<<24>>24)")?; } } Function::QuantizeToF16=>{ write!(self.out,"f16tof32(f32tof16(")?; self.write_expr(module,arg,func_ctx)?; write!(self.out,"))")?; } Function::Regular(fun_name)=>{ write!(self.out,"{fun_name}(")?; self.write_expr(module,arg,func_ctx)?; ifletSome(arg)=arg1{ write!(self.out,",")?; self.write_expr(module,arg,func_ctx)?; } ifletSome(arg)=arg2{ write!(self.out,",")?; self.write_expr(module,arg,func_ctx)?; } ifletSome(arg)=arg3{ write!(self.out,",")?; self.write_expr(module,arg,func_ctx)?; } write!(self.out,")")? } // These overloads are only missing on FXC, so this is only needed for 32bit types, // as non-32bit types are DXC only. Function::MissingIntOverload(fun_name)=>{ letscalar_kind=func_ctx.resolve_type(arg,&module.types).scalar(); ifletSome(Scalar::I32)=scalar_kind{ write!(self.out,"asint({fun_name}(asuint(")?; self.write_expr(module,arg,func_ctx)?; write!(self.out,")))")?; }else{ write!(self.out,"{fun_name}(")?; self.write_expr(module,arg,func_ctx)?; write!(self.out,")")?; } } // These overloads are only missing on FXC, so this is only needed for 32bit types, // as non-32bit types are DXC only. Function::MissingIntReturnType(fun_name)=>{ letscalar_kind=func_ctx.resolve_type(arg,&module.types).scalar(); ifletSome(Scalar::I32)=scalar_kind{ write!(self.out,"asint({fun_name}(")?; self.write_expr(module,arg,func_ctx)?; write!(self.out,"))")?; }else{ write!(self.out,"{fun_name}(")?; self.write_expr(module,arg,func_ctx)?; write!(self.out,")")?; } } Function::CountTrailingZeros=>{ match*func_ctx.resolve_type(arg,&module.types){ TypeInner::Vector{size,scalar}=>{ lets=matchsize{ crate::VectorSize::Bi=>".xx", crate::VectorSize::Tri=>".xxx", crate::VectorSize::Quad=>".xxxx", };
letscalar_width_bits=scalar.width*8;
ifscalar.kind==ScalarKind::Uint||scalar.width!=4{ write!( self.out, "min(({scalar_width_bits}u){s},firstbitlow(" )?; self.write_expr(module,arg,func_ctx)?; write!(self.out,"))")?; }else{ // This is only needed for the FXC path, on 32bit signed integers. write!( self.out, "asint(min(({scalar_width_bits}u){s},firstbitlow(" )?; self.write_expr(module,arg,func_ctx)?; write!(self.out,")))")?; } } TypeInner::Scalar(scalar)=>{ letscalar_width_bits=scalar.width*8;
ifscalar.kind==ScalarKind::Uint||scalar.width!=4{ write!(self.out,"min({scalar_width_bits}u,firstbitlow(")?; self.write_expr(module,arg,func_ctx)?; write!(self.out,"))")?; }else{ // This is only needed for the FXC path, on 32bit signed integers. write!( self.out, "asint(min({scalar_width_bits}u,firstbitlow(" )?; self.write_expr(module,arg,func_ctx)?; write!(self.out,")))")?; } } _=>unreachable!(), }
// return x component if return type is scalar ifletTypeInner::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. fnsampler_binding_array_info_from_expression( &mutself, module:&Module, func_ctx:&back::FunctionCtx<'_>, base:Handle<crate::Expression>, resolved:&TypeInner, )->Option<BindingArraySamplerInfo>{ ifletTypeInner::BindingArray{ base:base_ty_handle, .. }=*resolved { letbase_ty=&module.types[base_ty_handle].inner; ifletTypeInner::Sampler{comparison,..}=*base_ty{ letbase=&func_ctx.expressions[base];
write!(self.out,"{name}")?; // If rhs is a array type, we should write array size ifletTypeInner::Array{base,size,..}=*resolved{ self.write_array_size(module,base,size)?; } write!(self.out,"=")?; self.write_expr(module,handle,func_ctx)?; writeln!(self.out,";")?; self.named_expressions.insert(expr,name);
Ok(()) }
/// Helper function that write default zero initialization pub(super)fnwrite_default_init( &mutself, module:&Module, ty:Handle<crate::Type>, )->BackendResult{ write!(self.out,"(")?; self.write_type(module,ty)?; ifletTypeInner::Array{base,size,..}=module.types[ty].inner{ self.write_array_size(module,base,size)?; } write!(self.out,")0")?; Ok(()) }
pub(super)fnwrite_control_barrier( &mutself, barrier:crate::Barrier, level:back::Level, )->BackendResult{ ifbarrier.contains(crate::Barrier::STORAGE){ writeln!(self.out,"{level}DeviceMemoryBarrierWithGroupSync();")?; } ifbarrier.contains(crate::Barrier::WORK_GROUP){ writeln!(self.out,"{level}GroupMemoryBarrierWithGroupSync();")?; } ifbarrier.contains(crate::Barrier::SUB_GROUP){ // Does not exist in DirectX } ifbarrier.contains(crate::Barrier::TEXTURE){ writeln!(self.out,"{level}DeviceMemoryBarrierWithGroupSync();")?; } Ok(()) }
fnwrite_memory_barrier( &mutself, barrier:crate::Barrier, level:back::Level, )->BackendResult{ ifbarrier.contains(crate::Barrier::STORAGE){ writeln!(self.out,"{level}DeviceMemoryBarrier();")?; } ifbarrier.contains(crate::Barrier::WORK_GROUP){ writeln!(self.out,"{level}GroupMemoryBarrier();")?; } ifbarrier.contains(crate::Barrier::SUB_GROUP){ // Does not exist in DirectX } ifbarrier.contains(crate::Barrier::TEXTURE){ writeln!(self.out,"{level}DeviceMemoryBarrier();")?; } Ok(()) }
/// Helper to emit the shared tail of an HLSL atomic call (arguments, value, result) fnemit_hlsl_atomic_tail( &mutself, module:&Module, func_ctx:&back::FunctionCtx<'_>, fun:&crate::AtomicFunction, compare_expr:Option<Handle<crate::Expression>>, value:Handle<crate::Expression>, res_var_info:&Option<(Handle<crate::Expression>,String)>, )->BackendResult{ ifletSome(cmp)=compare_expr{ write!(self.out,",")?; self.write_expr(module,cmp,func_ctx)?; } write!(self.out,",")?; ifletcrate::AtomicFunction::Subtract=*fun{ // we just wrote `InterlockedAdd`, so negate the argument write!(self.out,"-")?; } self.write_expr(module,value,func_ctx)?; ifletSome(&(_res_handle,refres_name))=res_var_info.as_ref(){ write!(self.out,",")?; ifcompare_expr.is_some(){ write!(self.out,"{res_name}.old_value")?; }else{ write!(self.out,"{res_name}")?; } } writeln!(self.out,");")?; Ok(()) } }
/// 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). fnfind_matrix_in_access_chain( module:&Module, base:Handle<crate::Expression>, func_ctx:&back::FunctionCtx<'_>, )->Option<(Handle<crate::Expression>,Option<Index>,Option<Index>)>{ letmutcurrent_base=base; letmutvector=None; letmutscalar=None; loop{ letresolved_tr=func_ctx .resolve_type(current_base,&module.types) .pointer_base_type(); letresolved=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)fnget_inner_matrix_of_struct_array_member( module:&Module, base:Handle<crate::Expression>, func_ctx:&back::FunctionCtx<'_>, direct:bool, )->Option<MatrixType>{ letmutmat_data=None; letmutarray_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) fnget_inner_matrix_of_global_uniform( module:&Module, base:Handle<crate::Expression>, func_ctx:&back::FunctionCtx<'_>, direct:bool, )->Option<MatrixType>{ letmutmat_data=None; letmutarray_base=None;
¤ Diese beiden folgenden Angebotsgruppen bietet das Unternehmen0.252Angebot
(Wie Sie bei der Firma Beratungs- und Dienstleistungen beauftragen können 2026-08-26)
¤
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.
