From f038e59f452ef4f2fcaaa6e092889270ce4b197e Mon Sep 17 00:00:00 2001 From: Vecvec Date: Sun, 16 Nov 2025 19:39:24 +1300 Subject: [PATCH 01/15] rt pipelines in naga IR (w/ partial validation). --- naga/src/back/dot/mod.rs | 16 ++ naga/src/back/glsl/mod.rs | 51 +++++- naga/src/back/hlsl/conv.rs | 13 ++ naga/src/back/hlsl/mod.rs | 1 + naga/src/back/hlsl/writer.rs | 8 +- naga/src/back/msl/mod.rs | 15 +- naga/src/back/msl/writer.rs | 19 ++- naga/src/back/pipeline_constants.rs | 11 ++ naga/src/back/spv/block.rs | 1 + naga/src/back/spv/helpers.rs | 1 + naga/src/back/spv/writer.rs | 22 ++- naga/src/back/wgsl/writer.rs | 15 +- naga/src/common/wgsl/to_wgsl.rs | 16 +- naga/src/compact/statements.rs | 22 +++ naga/src/front/glsl/functions.rs | 1 + naga/src/front/spv/function.rs | 1 + naga/src/front/spv/mod.rs | 3 +- naga/src/front/wgsl/lower/mod.rs | 1 + naga/src/ir/mod.rs | 93 +++++++++++ naga/src/proc/mod.rs | 4 + naga/src/proc/namer.rs | 7 +- naga/src/proc/terminator.rs | 3 +- naga/src/valid/analyzer.rs | 16 +- naga/src/valid/function.rs | 77 +++++++++ naga/src/valid/handles.rs | 12 ++ naga/src/valid/interface.rs | 240 +++++++++++++++++++++++++++- naga/src/valid/mod.rs | 23 ++- naga/src/valid/type.rs | 2 +- naga/tests/naga/snapshots.rs | 2 +- wgpu-core/src/validation.rs | 2 +- wgpu-hal/src/auxil/mod.rs | 4 + wgpu-hal/src/gles/device.rs | 4 +- wgpu-types/src/lib.rs | 8 + 33 files changed, 685 insertions(+), 29 deletions(-) diff --git a/naga/src/back/dot/mod.rs b/naga/src/back/dot/mod.rs index 826dad1c219..6feb1ca73af 100644 --- a/naga/src/back/dot/mod.rs +++ b/naga/src/back/dot/mod.rs @@ -403,6 +403,22 @@ impl StatementGraph { }, } } + S::RayPipelineFunction(func) => match func { + crate::RayPipelineFunction::TraceRay { + acceleration_structure, + descriptor, + payload, + } => { + self.dependencies.push(( + id, + acceleration_structure, + "acceleration_structure", + )); + self.dependencies.push((id, descriptor, "descriptor")); + self.dependencies.push((id, payload, "payload")); + "TraceRay" + } + }, }; // Set the last node to the merge node last_node = merge_id; diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index 062734b049e..b1e95d15bb1 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -140,7 +140,10 @@ impl crate::AddressSpace { | crate::AddressSpace::Storage { .. } | crate::AddressSpace::Handle | crate::AddressSpace::PushConstant - | crate::AddressSpace::TaskPayload => false, + | crate::AddressSpace::TaskPayload + // just a default impl, not really supported + | crate::AddressSpace::RayPayload + | crate::AddressSpace::IncomingRayPayload => false, } } } @@ -504,7 +507,15 @@ impl fmt::Display for VaryingName<'_> { (ShaderStage::Vertex, true) | (ShaderStage::Fragment, false) => "vs2fs", // fragment to pipeline (ShaderStage::Fragment, true) => "fs2p", - (ShaderStage::Task | ShaderStage::Mesh, _) => unreachable!(), + ( + ShaderStage::Task + | ShaderStage::Mesh + | ShaderStage::RayGeneration + | ShaderStage::AnyHit + | ShaderStage::ClosestHit + | ShaderStage::Miss, + _, + ) => unreachable!(), }; write!(f, "_{prefix}_location{location}",) } @@ -521,7 +532,12 @@ impl ShaderStage { ShaderStage::Compute => "cs", ShaderStage::Fragment => "fs", ShaderStage::Vertex => "vs", - ShaderStage::Task | ShaderStage::Mesh => unreachable!(), + ShaderStage::Task + | ShaderStage::Mesh + | ShaderStage::RayGeneration + | ShaderStage::AnyHit + | ShaderStage::ClosestHit + | ShaderStage::Miss => unreachable!(), } } } @@ -1309,6 +1325,10 @@ impl<'a, W: Write> Writer<'a, W> { crate::AddressSpace::Function => unreachable!(), // Textures and samplers are handled directly in `Writer::write`. crate::AddressSpace::Handle => unreachable!(), + // ray tracing pipelines unsupported + crate::AddressSpace::RayPayload | crate::AddressSpace::IncomingRayPayload => { + unreachable!() + } } Ok(()) @@ -1675,7 +1695,12 @@ impl<'a, W: Write> Writer<'a, W> { ShaderStage::Vertex => output, ShaderStage::Fragment => !output, ShaderStage::Compute => false, - ShaderStage::Task | ShaderStage::Mesh => unreachable!(), + ShaderStage::Task + | ShaderStage::Mesh + | ShaderStage::RayGeneration + | ShaderStage::AnyHit + | ShaderStage::ClosestHit + | ShaderStage::Miss => unreachable!(), }; // Write the I/O locations, if allowed @@ -2810,6 +2835,7 @@ impl<'a, W: Write> Writer<'a, W> { } writeln!(self.out, ");")?; } + Statement::RayPipelineFunction(_) => unimplemented!(), } Ok(()) @@ -5269,7 +5295,20 @@ const fn glsl_built_in(built_in: crate::BuiltIn, options: VaryingOptions) -> &'s | Bi::VertexCount | Bi::PrimitiveCount | Bi::Vertices - | Bi::Primitives => { + | Bi::Primitives + | Bi::RayInvocationId + | Bi::NumRayInvocations + | Bi::InstanceCustomData + | Bi::GeometryIndex + | Bi::WorldRayOrigin + | Bi::WorldRayDirection + | Bi::ObjectRayOrigin + | Bi::ObjectRayDirection + | Bi::RayTmin + | Bi::RayTCurrentMax + | Bi::ObjectToWorld + | Bi::WorldToObject + | Bi::HitKind => { unimplemented!() } } @@ -5287,7 +5326,7 @@ const fn glsl_storage_qualifier(space: crate::AddressSpace) -> Option<&'static s As::Handle => Some("uniform"), As::WorkGroup => Some("shared"), As::PushConstant => Some("uniform"), - As::TaskPayload => unreachable!(), + As::TaskPayload | As::RayPayload | As::IncomingRayPayload => unreachable!(), } } diff --git a/naga/src/back/hlsl/conv.rs b/naga/src/back/hlsl/conv.rs index 6cd3679e817..ab534019dc0 100644 --- a/naga/src/back/hlsl/conv.rs +++ b/naga/src/back/hlsl/conv.rs @@ -192,6 +192,19 @@ impl crate::BuiltIn { | Self::PrimitiveCount | Self::Vertices | Self::Primitives => unreachable!(), + Self::RayInvocationId + | Self::NumRayInvocations + | Self::InstanceCustomData + | Self::GeometryIndex + | Self::WorldRayOrigin + | Self::WorldRayDirection + | Self::ObjectRayOrigin + | Self::ObjectRayDirection + | Self::RayTmin + | Self::RayTCurrentMax + | Self::ObjectToWorld + | Self::WorldToObject + | Self::HitKind => unreachable!(), }) } } diff --git a/naga/src/back/hlsl/mod.rs b/naga/src/back/hlsl/mod.rs index 7fca9670921..a745d153f89 100644 --- a/naga/src/back/hlsl/mod.rs +++ b/naga/src/back/hlsl/mod.rs @@ -285,6 +285,7 @@ impl crate::ShaderStage { Self::Compute => "cs", Self::Task => "as", Self::Mesh => "ms", + Self::RayGeneration | Self::AnyHit | Self::ClosestHit | Self::Miss => "lib", } } } diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index e55472460e4..c8bf2d8c678 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/naga/src/back/hlsl/writer.rs @@ -1008,6 +1008,9 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { write!(self.out, "ConstantBuffer<")?; "b" } + crate::AddressSpace::RayPayload | crate::AddressSpace::IncomingRayPayload => { + unimplemented!() + } }; // If the global is a push constant write the type now because it will be a @@ -2756,6 +2759,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } writeln!(self.out, ");")?; } + Statement::RayPipelineFunction(_) => unreachable!(), } Ok(()) @@ -3086,7 +3090,9 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { | crate::AddressSpace::Private | crate::AddressSpace::WorkGroup | crate::AddressSpace::PushConstant - | crate::AddressSpace::TaskPayload, + | crate::AddressSpace::TaskPayload + | crate::AddressSpace::RayPayload + | crate::AddressSpace::IncomingRayPayload, ) | None => true, Some(crate::AddressSpace::Uniform) => { diff --git a/naga/src/back/msl/mod.rs b/naga/src/back/msl/mod.rs index 2456cdbae8b..a160f37d584 100644 --- a/naga/src/back/msl/mod.rs +++ b/naga/src/back/msl/mod.rs @@ -718,7 +718,20 @@ impl ResolvedBinding { | Bi::VertexCount | Bi::PrimitiveCount | Bi::Vertices - | Bi::Primitives => unreachable!(), + | Bi::Primitives + | Bi::RayInvocationId + | Bi::NumRayInvocations + | Bi::InstanceCustomData + | Bi::GeometryIndex + | Bi::WorldRayOrigin + | Bi::WorldRayDirection + | Bi::ObjectRayOrigin + | Bi::ObjectRayDirection + | Bi::RayTmin + | Bi::RayTCurrentMax + | Bi::ObjectToWorld + | Bi::WorldToObject + | Bi::HitKind => unreachable!(), }; write!(out, "{name}")?; } diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index 8c21c944718..fd6e27c7ace 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -596,7 +596,9 @@ impl crate::AddressSpace { | Self::WorkGroup | Self::PushConstant | Self::Handle - | Self::TaskPayload => true, + | Self::TaskPayload + | Self::RayPayload + | Self::IncomingRayPayload => true, Self::Function => false, } } @@ -609,7 +611,7 @@ impl crate::AddressSpace { // may end up with "const" even if the binding is read-write, // and that should be OK. Self::Storage { .. } => true, - Self::TaskPayload => unimplemented!(), + Self::TaskPayload | Self::RayPayload | Self::IncomingRayPayload => unimplemented!(), // These should always be read-write. Self::Private | Self::WorkGroup => false, // These translate to `constant` address space, no need for qualifiers. @@ -624,9 +626,13 @@ impl crate::AddressSpace { Self::Handle => None, Self::Uniform | Self::PushConstant => Some("constant"), Self::Storage { .. } => Some("device"), - Self::Private | Self::Function => Some("thread"), + // note for `RayPayload`, this probably needs to be emulated as a + // private variable, as metal has essentially an inout input + // for where it is passed. + Self::Private | Self::Function | Self::RayPayload => Some("thread"), Self::WorkGroup => Some("threadgroup"), Self::TaskPayload => Some("object_data"), + Self::IncomingRayPayload => Some("ray_data"), } } } @@ -4194,6 +4200,7 @@ impl Writer { } writeln!(self.out, ");")?; } + crate::Statement::RayPipelineFunction(_) => unreachable!(), } } @@ -6672,6 +6679,10 @@ template false, ), crate::ShaderStage::Task | crate::ShaderStage::Mesh => unimplemented!(), + crate::ShaderStage::RayGeneration + | crate::ShaderStage::AnyHit + | crate::ShaderStage::ClosestHit + | crate::ShaderStage::Miss => unimplemented!(), }; // Should this entry point be modified to do vertex pulling? @@ -6744,6 +6755,8 @@ template crate::AddressSpace::Function | crate::AddressSpace::Private | crate::AddressSpace::WorkGroup => {} + crate::AddressSpace::RayPayload + | crate::AddressSpace::IncomingRayPayload => unimplemented!(), } } if needs_buffer_sizes { diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs index de643b82fab..8bdc1349d56 100644 --- a/naga/src/back/pipeline_constants.rs +++ b/naga/src/back/pipeline_constants.rs @@ -860,6 +860,17 @@ fn adjust_stmt(new_pos: &HandleVec>, stmt: &mut S crate::RayQueryFunction::Terminate => {} } } + Statement::RayPipelineFunction(ref mut func) => match func { + crate::RayPipelineFunction::TraceRay { + ref mut acceleration_structure, + ref mut descriptor, + ref mut payload, + } => { + adjust(acceleration_structure); + adjust(descriptor); + adjust(payload); + } + }, Statement::Break | Statement::Continue | Statement::Kill diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index ace9565a1cd..1a95b3eb03e 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -3679,6 +3679,7 @@ impl BlockContext<'_> { } => { self.write_subgroup_gather(mode, argument, result, &mut block)?; } + Statement::RayPipelineFunction(_) => unreachable!(), } } diff --git a/naga/src/back/spv/helpers.rs b/naga/src/back/spv/helpers.rs index 6522a0970d7..366aacd0f3a 100644 --- a/naga/src/back/spv/helpers.rs +++ b/naga/src/back/spv/helpers.rs @@ -55,6 +55,7 @@ pub(super) const fn map_storage_class(space: crate::AddressSpace) -> spirv::Stor crate::AddressSpace::WorkGroup => spirv::StorageClass::Workgroup, crate::AddressSpace::PushConstant => spirv::StorageClass::PushConstant, crate::AddressSpace::TaskPayload => unreachable!(), + crate::AddressSpace::IncomingRayPayload | crate::AddressSpace::RayPayload => unreachable!(), } } diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index 6180f8a599f..bd66e7eb764 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -1286,7 +1286,12 @@ impl Writer { .to_words(&mut self.logical_layout.execution_modes); spirv::ExecutionModel::GLCompute } - crate::ShaderStage::Task | crate::ShaderStage::Mesh => unreachable!(), + crate::ShaderStage::Task + | crate::ShaderStage::Mesh + | crate::ShaderStage::RayGeneration + | crate::ShaderStage::AnyHit + | crate::ShaderStage::ClosestHit + | crate::ShaderStage::Miss => unreachable!(), }; //self.check(exec_model.required_capabilities())?; @@ -2223,7 +2228,20 @@ impl Writer { | Bi::VertexCount | Bi::PrimitiveCount | Bi::Vertices - | Bi::Primitives => unreachable!(), + | Bi::Primitives + | Bi::RayInvocationId + | Bi::NumRayInvocations + | Bi::InstanceCustomData + | Bi::GeometryIndex + | Bi::WorldRayOrigin + | Bi::WorldRayDirection + | Bi::ObjectRayOrigin + | Bi::ObjectRayDirection + | Bi::RayTmin + | Bi::RayTCurrentMax + | Bi::ObjectToWorld + | Bi::WorldToObject + | Bi::HitKind => unreachable!(), }; self.decorate(id, Decoration::BuiltIn, &[built_in as u32]); diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs index daf32a7116f..253e59f6d1b 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -207,7 +207,12 @@ impl Writer { Attribute::Stage(ShaderStage::Compute), Attribute::WorkGroupSize(ep.workgroup_size), ], - ShaderStage::Mesh | ShaderStage::Task => unreachable!(), + ShaderStage::Mesh + | ShaderStage::Task + | ShaderStage::RayGeneration + | ShaderStage::AnyHit + | ShaderStage::ClosestHit + | ShaderStage::Miss => unreachable!(), }; self.write_attributes(&attributes)?; @@ -403,7 +408,12 @@ impl Writer { ShaderStage::Vertex => "vertex", ShaderStage::Fragment => "fragment", ShaderStage::Compute => "compute", - ShaderStage::Task | ShaderStage::Mesh => unreachable!(), + ShaderStage::Task + | ShaderStage::Mesh + | ShaderStage::RayGeneration + | ShaderStage::AnyHit + | ShaderStage::ClosestHit + | ShaderStage::Miss => unreachable!(), }; write!(self.out, "@{stage_str} ")?; } @@ -984,6 +994,7 @@ impl Writer { } writeln!(self.out, ");")?; } + Statement::RayPipelineFunction(_) => unreachable!(), } Ok(()) diff --git a/naga/src/common/wgsl/to_wgsl.rs b/naga/src/common/wgsl/to_wgsl.rs index 5e6178c049c..4be75c0551d 100644 --- a/naga/src/common/wgsl/to_wgsl.rs +++ b/naga/src/common/wgsl/to_wgsl.rs @@ -198,7 +198,20 @@ impl TryToWgsl for crate::BuiltIn { | Bi::VertexCount | Bi::PrimitiveCount | Bi::Vertices - | Bi::Primitives => return None, + | Bi::Primitives + | Bi::RayInvocationId + | Bi::NumRayInvocations + | Bi::InstanceCustomData + | Bi::GeometryIndex + | Bi::WorldRayOrigin + | Bi::WorldRayDirection + | Bi::ObjectRayOrigin + | Bi::ObjectRayDirection + | Bi::RayTmin + | Bi::RayTCurrentMax + | Bi::ObjectToWorld + | Bi::WorldToObject + | Bi::HitKind => return None, }) } } @@ -363,6 +376,7 @@ pub const fn address_space_str( As::Handle => return (None, None), As::Function => "function", As::TaskPayload => return (None, None), + As::IncomingRayPayload | As::RayPayload => return (None, None), }), None, ) diff --git a/naga/src/compact/statements.rs b/naga/src/compact/statements.rs index 39d6065f5f0..18e477c97ab 100644 --- a/naga/src/compact/statements.rs +++ b/naga/src/compact/statements.rs @@ -152,6 +152,17 @@ impl FunctionTracer<'_> { self.expressions_used.insert(argument); self.expressions_used.insert(result); } + St::RayPipelineFunction(func) => match func { + crate::RayPipelineFunction::TraceRay { + acceleration_structure, + descriptor, + payload, + } => { + self.expressions_used.insert(acceleration_structure); + self.expressions_used.insert(descriptor); + self.expressions_used.insert(payload); + } + }, // Trivial statements. St::Break @@ -371,6 +382,17 @@ impl FunctionMap { adjust(argument); adjust(result); } + St::RayPipelineFunction(ref mut func) => match func { + crate::RayPipelineFunction::TraceRay { + ref mut acceleration_structure, + ref mut descriptor, + ref mut payload, + } => { + adjust(acceleration_structure); + adjust(descriptor); + adjust(payload); + } + }, // Trivial statements. St::Break diff --git a/naga/src/front/glsl/functions.rs b/naga/src/front/glsl/functions.rs index ba096a82b3b..5d1f930452b 100644 --- a/naga/src/front/glsl/functions.rs +++ b/naga/src/front/glsl/functions.rs @@ -1379,6 +1379,7 @@ impl Frontend { }, mesh_info: None, task_payload: None, + ray_incoming_payload: None, }); Ok(()) diff --git a/naga/src/front/spv/function.rs b/naga/src/front/spv/function.rs index 48b23e7c4c4..a88bcf306bb 100644 --- a/naga/src/front/spv/function.rs +++ b/naga/src/front/spv/function.rs @@ -598,6 +598,7 @@ impl> super::Frontend { function, mesh_info: None, task_payload: None, + ray_incoming_payload: None, }); Ok(()) diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs index ac9eaf8306f..48fb60d7640 100644 --- a/naga/src/front/spv/mod.rs +++ b/naga/src/front/spv/mod.rs @@ -4663,7 +4663,8 @@ impl> Frontend { | S::RayQuery { .. } | S::SubgroupBallot { .. } | S::SubgroupCollectiveOperation { .. } - | S::SubgroupGather { .. } => {} + | S::SubgroupGather { .. } + | S::RayPipelineFunction(..) => {} S::Call { function: ref mut callee, ref arguments, diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index c47941f0f86..a9d34646b9a 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -1638,6 +1638,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { function, mesh_info, task_payload, + ray_incoming_payload: None, }); Ok(LoweredGlobalDecl::EntryPoint( ctx.module.entry_points.len() - 1, diff --git a/naga/src/ir/mod.rs b/naga/src/ir/mod.rs index 5aa411b72e5..01c6057fb89 100644 --- a/naga/src/ir/mod.rs +++ b/naga/src/ir/mod.rs @@ -335,6 +335,17 @@ pub enum ShaderStage { /// Compute pipeline shader. Compute, + + /// A ray generation shader, in a ray tracing pipeline. + RayGeneration, + + /// A miss shader, in a ray tracing pipeline. + Miss, + + /// A any hit shader, in a ray tracing pipeline. + AnyHit, + /// A closest hit shader, in a ray tracing pipeline. + ClosestHit, } /// Addressing space of variables. @@ -373,6 +384,14 @@ pub enum AddressSpace { PushConstant, /// Task shader to mesh shader payload TaskPayload, + + /// Ray tracing payload, for inputting in TraceRays + RayPayload, + /// Ray tracing payload, for entrypoints invoked by a TraceRays call + /// + /// Each entrypoint may reference only one variable in this scope, as + /// only one may be passed as a payload. + IncomingRayPayload, } /// Built-in inputs and outputs. @@ -459,6 +478,46 @@ pub enum BuiltIn { PrimitiveCount, /// Written to a workgroup variable in mesh shaders Primitives, + + /// Read in all ray tracing pipeline shaders, the id within the number of + /// rays that this current ray is. + RayInvocationId, + /// Read in all ray tracing pipeline shaders, the number of rays created. + NumRayInvocations, + /// Read in closest hit and any hit shaders, the custom data in the tlas + /// instance + InstanceCustomData, + /// Read in closest hit and any hit shaders, the index of the geometry in + /// the blas. + GeometryIndex, + /// Read in closest hit, any hit, and miss shaders, the origin of the ray. + WorldRayOrigin, + /// Read in closest hit, any hit, and miss shaders, the direction of the + /// ray. + WorldRayDirection, + /// Read in closest hit and any hit shaders, the direction of the ray in + /// object space. + ObjectRayOrigin, + /// Read in closest hit and any hit shaders, the direction of the ray in + /// object space. + ObjectRayDirection, + /// Read in closest hit, any hit, and miss shaders, the t min provided by + /// in the ray desc. + RayTmin, + /// Read in closest hit, any hit, and miss shaders, the final bounds at which + /// a hit is accepted (the closest committed hit if there is one otherwise, t + /// max provided in the ray desc). + RayTCurrentMax, + /// Read in closest hit and any hit shaders, the matrix for converting from + /// object space to world space + ObjectToWorld, + /// Read in closest hit and any hit shaders, the matrix for converting from + /// world space to object space + WorldToObject, + /// Read in closest hit and any hit shaders, the type of hit as provided by + /// the intersection function if any, otherwise this is 254 (0xFE) for a + /// front facing triangle and 255 (0xFF) for a back facing triangle + HitKind, } /// Number of bytes per scalar. @@ -2220,6 +2279,8 @@ pub enum Statement { /// The specific operation we're performing on `query`. fun: RayQueryFunction, }, + /// A ray tracing pipeline shader intrinsic. + RayPipelineFunction(RayPipelineFunction), /// Calculate a bitmask using a boolean from each active thread in the subgroup SubgroupBallot { /// The [`SubgroupBallotResult`] expression representing this load's result. @@ -2399,6 +2460,9 @@ pub struct EntryPoint { pub mesh_info: Option, /// The unique global variable used as a task payload from task shader to mesh shader pub task_payload: Option>, + /// The unique global variable used as an incoming ray payload going into any hit, closest hit and miss shaders. + /// Unlike the outgoing ray payload, an incoming ray payload must be unique + pub ray_incoming_payload: Option>, } /// Return types predeclared for the frexp, modf, and atomicCompareExchangeWeak built-in functions. @@ -2614,6 +2678,35 @@ pub struct MeshStageInfo { pub output_variable: Handle, } +/// Ray tracing pipeine intrinsics +#[derive(Debug, Clone, Copy)] +#[cfg_attr(feature = "serialize", derive(Serialize))] +#[cfg_attr(feature = "deserialize", derive(Deserialize))] +#[cfg_attr(feature = "arbitrary", derive(Arbitrary))] +pub enum RayPipelineFunction { + /// Traces a ray through the given acceleration structure + TraceRay { + /// The acceleration structure within which this ray should search for hits. + /// + /// The expression must be an [`AccelerationStructure`]. + /// + /// [`AccelerationStructure`]: TypeInner::AccelerationStructure + acceleration_structure: Handle, + + #[allow(rustdoc::private_intra_doc_links)] + /// A struct of detailed parameters for the ray query. + /// + /// This expression should have the struct type given in + /// [`SpecialTypes::ray_desc`]. This is available in the WGSL + /// front end as the `RayDesc` type. + descriptor: Handle, + + /// A pointer in the ray_payload or incoming_ray_payload address spaces + payload: Handle, + // Do we want miss index? What about sbt offset and sbt stride (could be hard to validate)? + }, +} + /// Shader module. /// /// A module is a set of constants, global variables and functions, as well as diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs index cebd98f2e47..6c929d7d4a3 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -186,6 +186,9 @@ impl super::AddressSpace { // TaskPayload isn't always writable, but this is checked for elsewhere, // when not using multiple payloads and matching the entry payload is checked. crate::AddressSpace::TaskPayload => Sa::LOAD | Sa::STORE, + crate::AddressSpace::RayPayload | crate::AddressSpace::IncomingRayPayload => { + Sa::LOAD | Sa::STORE + } } } } @@ -640,6 +643,7 @@ impl super::ShaderStage { match self { Self::Vertex | Self::Fragment => false, Self::Compute | Self::Task | Self::Mesh => true, + Self::RayGeneration | Self::AnyHit | Self::ClosestHit | Self::Miss => false, } } } diff --git a/naga/src/proc/namer.rs b/naga/src/proc/namer.rs index 127e346f3a1..108b377c3d4 100644 --- a/naga/src/proc/namer.rs +++ b/naga/src/proc/namer.rs @@ -255,7 +255,12 @@ impl Namer { crate::ShaderStage::Vertex => "VertexOutput", crate::ShaderStage::Fragment => "FragmentOutput", crate::ShaderStage::Compute => "ComputeOutput", - crate::ShaderStage::Task | crate::ShaderStage::Mesh => unreachable!(), + crate::ShaderStage::Task + | crate::ShaderStage::Mesh + | crate::ShaderStage::RayGeneration + | crate::ShaderStage::ClosestHit + | crate::ShaderStage::AnyHit + | crate::ShaderStage::Miss => unreachable!(), }; entrypoint_type_fallbacks.insert(result.ty, label); } diff --git a/naga/src/proc/terminator.rs b/naga/src/proc/terminator.rs index b29ccb054a3..f15e5ecc837 100644 --- a/naga/src/proc/terminator.rs +++ b/naga/src/proc/terminator.rs @@ -43,7 +43,8 @@ pub fn ensure_block_returns(block: &mut crate::Block) { | S::SubgroupCollectiveOperation { .. } | S::SubgroupGather { .. } | S::ControlBarrier(_) - | S::MemoryBarrier(_)), + | S::MemoryBarrier(_) + | S::RayPipelineFunction(_)), ) | None => block.push(S::Return { value: None }, Default::default()), } diff --git a/naga/src/valid/analyzer.rs b/naga/src/valid/analyzer.rs index e01a7b0b735..d5f416092f5 100644 --- a/naga/src/valid/analyzer.rs +++ b/naga/src/valid/analyzer.rs @@ -649,7 +649,7 @@ impl FunctionInfo { let var = &resolve_context.global_vars[gh]; let uniform = match var.space { // local data is non-uniform - As::Function | As::Private => false, + As::Function | As::Private | As::RayPayload | As::IncomingRayPayload => false, // workgroup memory is exclusively accessed by the group // task payload memory is very similar to workgroup memory As::WorkGroup | As::TaskPayload => true, @@ -1168,6 +1168,20 @@ impl FunctionInfo { } FunctionUniformity::new() } + S::RayPipelineFunction(ref fun) => { + match fun { + crate::RayPipelineFunction::TraceRay { + acceleration_structure, + descriptor, + payload, + } => { + let _ = self.add_ref(*acceleration_structure); + let _ = self.add_ref(*descriptor); + let _ = self.add_ref(*payload); + } + } + FunctionUniformity::new() + } }; disruptor = disruptor.or(uniformity.exit_disruptor()); diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index abf6bc430a6..6f00849291b 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -225,6 +225,12 @@ pub enum FunctionError { ConflictingTaskPayloadVariables(Handle, Handle), #[error("Mesh shader output at {0:?} is not a user-defined struct")] InvalidMeshShaderOutputType(Handle), + #[error("The payload type passed to `traceRay` must be a pointer")] + InvalidPayloadType, + #[error("The payload type passed to `traceRay` must be a pointer with an adress space of `ray_payload` or `incoming_ray_payload`, instead got {0:?}")] + InvalidPayloadAdressSpace(crate::AddressSpace), + #[error("The payload type ({0:?}) passed to `traceRay` does not match the previous one {1:?}")] + MismatchedPayloadType(Handle, Handle), } bitflags::bitflags! { @@ -1626,6 +1632,77 @@ impl super::Validator { } self.validate_subgroup_gather(mode, argument, result, context)?; } + S::RayPipelineFunction(ref fun) => match *fun { + crate::RayPipelineFunction::TraceRay { + acceleration_structure, + descriptor, + payload, + } => { + match *context.resolve_type_inner( + acceleration_structure, + &self.valid_expression_set, + )? { + crate::TypeInner::AccelerationStructure { vertex_return } => { + if !vertex_return { + self.trace_rays_no_vertex_return = Some(Some(span)); + } else { + if let None = self.trace_rays_no_vertex_return { + self.trace_rays_no_vertex_return = Some(None); + } + } + } + _ => { + return Err(FunctionError::InvalidAccelerationStructure( + acceleration_structure, + ) + .with_span_handle(acceleration_structure, context.expressions)) + } + } + + let current_payload_ty = match *context + .resolve_type_inner(payload, &self.valid_expression_set)? + { + crate::TypeInner::Pointer { base, space } => { + match space { + AddressSpace::RayPayload | AddressSpace::IncomingRayPayload => { + } + space => { + return Err(FunctionError::InvalidPayloadAdressSpace(space) + .with_span_handle(payload, context.expressions)) + } + } + base + } + _ => { + return Err(FunctionError::InvalidPayloadType + .with_span_handle(payload, context.expressions)) + } + }; + + let ty = *self + .trace_rays_payload_type + .get_or_insert(current_payload_ty); + + if ty != current_payload_ty { + return Err(FunctionError::MismatchedPayloadType( + current_payload_ty, + ty, + ) + .with_span_handle(ty, context.types)); + } + + let desc_ty_given = + context.resolve_type_inner(descriptor, &self.valid_expression_set)?; + let desc_ty_expected = context + .special_types + .ray_desc + .map(|handle| &context.types[handle].inner); + if Some(desc_ty_given) != desc_ty_expected { + return Err(FunctionError::InvalidRayDescriptor(descriptor) + .with_span_static(span, "invalid ray descriptor")); + } + } + }, } } Ok(BlockInfo { stages }) diff --git a/naga/src/valid/handles.rs b/naga/src/valid/handles.rs index 5b7fb3fab75..2e61a3b9ad5 100644 --- a/naga/src/valid/handles.rs +++ b/naga/src/valid/handles.rs @@ -850,6 +850,18 @@ impl super::Validator { validate_expr(result)?; Ok(()) } + crate::Statement::RayPipelineFunction(fun) => match fun { + crate::RayPipelineFunction::TraceRay { + acceleration_structure, + descriptor, + payload, + } => { + validate_expr(acceleration_structure)?; + validate_expr(descriptor)?; + validate_expr(payload)?; + Ok(()) + } + }, crate::Statement::Break | crate::Statement::Continue | crate::Statement::Kill diff --git a/naga/src/valid/interface.rs b/naga/src/valid/interface.rs index a18b3738e55..4ca46be6c05 100644 --- a/naga/src/valid/interface.rs +++ b/naga/src/valid/interface.rs @@ -164,6 +164,10 @@ pub enum EntryPointError { WrongMeshOutputAddressSpace, #[error("Task payload must be at least 4 bytes, but is {0} bytes")] TaskPayloadTooSmall(u32), + #[error("Only the `ray_generation`, `closest_hit`, and `any_hit` shader stages can access a global variable in the `ray_payload` address space")] + RayPayloadInInvalidStage, + #[error("Only the `closest_hit`, `any_hit`, and `miss` shader stages can access a global variable in the `incoming_ray_payload` address space")] + IncomingRayPayloadInInvalidStage, } fn storage_usage(access: crate::StorageAccess) -> GlobalUse { @@ -294,6 +298,7 @@ impl VaryingContext<'_> { St::Vertex | St::Mesh => self.output, St::Fragment => !self.output, St::Compute | St::Task => false, + St::RayGeneration | St::AnyHit | St::ClosestHit | St::Miss => false, }, *ty_inner == Ti::Vector { @@ -304,7 +309,11 @@ impl VaryingContext<'_> { Bi::ViewIndex => ( match self.stage { St::Vertex | St::Fragment | St::Task | St::Mesh => !self.output, - St::Compute => false, + St::Compute + | St::RayGeneration + | St::AnyHit + | St::ClosestHit + | St::Miss => false, }, *ty_inner == Ti::Scalar(crate::Scalar::U32), ), @@ -361,7 +370,14 @@ impl VaryingContext<'_> { ), Bi::SubgroupSize | Bi::SubgroupInvocationId => ( match self.stage { - St::Compute | St::Fragment | St::Task | St::Mesh => !self.output, + St::Compute + | St::Fragment + | St::Task + | St::Mesh + | St::RayGeneration + | St::AnyHit + | St::ClosestHit + | St::Miss => !self.output, St::Vertex => false, }, *ty_inner == Ti::Scalar(crate::Scalar::U32), @@ -398,6 +414,193 @@ impl VaryingContext<'_> { scalar: crate::Scalar::U32, }, ), + Bi::RayInvocationId => ( + match self.stage { + St::Vertex | St::Fragment | St::Compute | St::Mesh | St::Task => false, + St::RayGeneration | St::AnyHit | St::ClosestHit | St::Miss => true, + }, + *ty_inner + == Ti::Vector { + size: Vs::Tri, + scalar: crate::Scalar::U32, + }, + ), + Bi::NumRayInvocations => ( + match self.stage { + St::Vertex | St::Fragment | St::Compute | St::Mesh | St::Task => false, + St::RayGeneration | St::AnyHit | St::ClosestHit | St::Miss => true, + }, + *ty_inner + == Ti::Vector { + size: Vs::Tri, + scalar: crate::Scalar::U32, + }, + ), + Bi::InstanceCustomData => ( + match self.stage { + St::RayGeneration + | St::Miss + | St::Vertex + | St::Fragment + | St::Compute + | St::Mesh + | St::Task => false, + St::AnyHit | St::ClosestHit => true, + }, + *ty_inner == Ti::Scalar(crate::Scalar::U32), + ), + Bi::GeometryIndex => ( + match self.stage { + St::RayGeneration + | St::Miss + | St::Vertex + | St::Fragment + | St::Compute + | St::Mesh + | St::Task => false, + St::AnyHit | St::ClosestHit => true, + }, + *ty_inner == Ti::Scalar(crate::Scalar::U32), + ), + Bi::WorldRayOrigin => ( + match self.stage { + St::RayGeneration + | St::Vertex + | St::Fragment + | St::Compute + | St::Mesh + | St::Task => false, + St::AnyHit | St::ClosestHit | St::Miss => true, + }, + *ty_inner + == Ti::Vector { + size: Vs::Tri, + scalar: crate::Scalar::F32, + }, + ), + Bi::WorldRayDirection => ( + match self.stage { + St::RayGeneration + | St::Vertex + | St::Fragment + | St::Compute + | St::Mesh + | St::Task => false, + St::AnyHit | St::ClosestHit | St::Miss => true, + }, + *ty_inner + == Ti::Vector { + size: Vs::Tri, + scalar: crate::Scalar::F32, + }, + ), + Bi::ObjectRayOrigin => ( + match self.stage { + St::RayGeneration + | St::Miss + | St::Vertex + | St::Fragment + | St::Compute + | St::Mesh + | St::Task => false, + St::AnyHit | St::ClosestHit => true, + }, + *ty_inner + == Ti::Vector { + size: Vs::Tri, + scalar: crate::Scalar::F32, + }, + ), + Bi::ObjectRayDirection => ( + match self.stage { + St::RayGeneration + | St::Miss + | St::Vertex + | St::Fragment + | St::Compute + | St::Mesh + | St::Task => false, + St::AnyHit | St::ClosestHit => true, + }, + *ty_inner + == Ti::Vector { + size: Vs::Tri, + scalar: crate::Scalar::F32, + }, + ), + Bi::RayTmin => ( + match self.stage { + St::RayGeneration + | St::Vertex + | St::Fragment + | St::Compute + | St::Mesh + | St::Task => false, + St::AnyHit | St::ClosestHit | St::Miss => true, + }, + *ty_inner == Ti::Scalar(crate::Scalar::F32), + ), + Bi::RayTCurrentMax => ( + match self.stage { + St::RayGeneration + | St::Vertex + | St::Fragment + | St::Compute + | St::Mesh + | St::Task => false, + St::AnyHit | St::ClosestHit | St::Miss => true, + }, + *ty_inner == Ti::Scalar(crate::Scalar::F32), + ), + Bi::ObjectToWorld => ( + match self.stage { + St::RayGeneration + | St::Miss + | St::Vertex + | St::Fragment + | St::Compute + | St::Mesh + | St::Task => false, + St::AnyHit | St::ClosestHit => true, + }, + *ty_inner + == Ti::Matrix { + columns: crate::VectorSize::Quad, + rows: crate::VectorSize::Tri, + scalar: crate::Scalar::F32, + }, + ), + Bi::WorldToObject => ( + match self.stage { + St::RayGeneration + | St::Miss + | St::Vertex + | St::Fragment + | St::Compute + | St::Mesh + | St::Task => false, + St::AnyHit | St::ClosestHit => true, + }, + *ty_inner + == Ti::Matrix { + columns: crate::VectorSize::Quad, + rows: crate::VectorSize::Tri, + scalar: crate::Scalar::F32, + }, + ), + Bi::HitKind => ( + match self.stage { + St::RayGeneration + | St::Miss + | St::Vertex + | St::Fragment + | St::Compute + | St::Mesh + | St::Task => false, + St::AnyHit | St::ClosestHit => true, + }, + *ty_inner == Ti::Scalar(crate::Scalar::U32), + ), // Validated elsewhere, shouldn't be here Bi::VertexCount | Bi::PrimitiveCount | Bi::Vertices | Bi::Primitives => { (false, true) @@ -532,7 +735,7 @@ impl VaryingContext<'_> { let needs_interpolation = match self.stage { crate::ShaderStage::Vertex => self.output, crate::ShaderStage::Fragment => !self.output && !per_primitive, - crate::ShaderStage::Compute | crate::ShaderStage::Task => false, + crate::ShaderStage::Compute | crate::ShaderStage::Task | crate::ShaderStage::RayGeneration | crate::ShaderStage::AnyHit | crate::ShaderStage::ClosestHit | crate::ShaderStage::Miss => false, crate::ShaderStage::Mesh => self.output, }; @@ -758,6 +961,17 @@ impl super::Validator { false, ) } + crate::AddressSpace::RayPayload | crate::AddressSpace::IncomingRayPayload => { + if !self.capabilities.contains(Capabilities::RAY_TRACING_PIPELINE) { + return Err(GlobalVariableError::UnsupportedCapability( + Capabilities::PUSH_CONSTANT, + )); + } + ( + TypeFlags::DATA | TypeFlags::SIZED, + false, + ) + } }; if !type_info.flags.contains(required_type_flags) { @@ -959,6 +1173,10 @@ impl super::Validator { crate::ShaderStage::Compute => ShaderStages::COMPUTE, crate::ShaderStage::Mesh => ShaderStages::MESH, crate::ShaderStage::Task => ShaderStages::TASK, + crate::ShaderStage::RayGeneration => ShaderStages::RAY_GENERATION, + crate::ShaderStage::AnyHit => ShaderStages::ANY_HIT, + crate::ShaderStage::ClosestHit => ShaderStages::CLOSEST_HIT, + crate::ShaderStage::Miss => ShaderStages::MISS, }; if !info.available_stages.contains(stage_bit) { @@ -1096,6 +1314,22 @@ impl super::Validator { } } crate::AddressSpace::PushConstant => GlobalUse::READ, + crate::AddressSpace::RayPayload => { + if matches!(ep.stage, crate::ShaderStage::RayGeneration | crate::ShaderStage::ClosestHit | crate::ShaderStage::Miss) { + return Err(EntryPointError::RayPayloadInInvalidStage.with_span_handle(var_handle, &module.global_variables)); + } + GlobalUse::READ + | GlobalUse::QUERY + | GlobalUse::WRITE + } + crate::AddressSpace::IncomingRayPayload => { + if !matches!(ep.stage, crate::ShaderStage::AnyHit | crate::ShaderStage::ClosestHit | crate::ShaderStage::Miss) { + return Err(EntryPointError::IncomingRayPayloadInInvalidStage.with_span_handle(var_handle, &module.global_variables)); + } + GlobalUse::READ + | GlobalUse::QUERY + | GlobalUse::WRITE + } }; if !allowed_usage.contains(usage) { log::warn!("\tUsage error for: {var:?}"); diff --git a/naga/src/valid/mod.rs b/naga/src/valid/mod.rs index 8dcea0acffc..49fe6617871 100644 --- a/naga/src/valid/mod.rs +++ b/naga/src/valid/mod.rs @@ -192,6 +192,8 @@ bitflags::bitflags! { const MESH_SHADER = 1 << 30; /// Support for mesh shaders which output points. const MESH_SHADER_POINT_TOPOLOGY = 1 << 30; + /// Support for ray generation, any hit, closest hit, and miss shaders. + const RAY_TRACING_PIPELINE = 1 << 31; } } @@ -282,12 +284,16 @@ bitflags::bitflags! { #[cfg_attr(feature = "serialize", derive(serde::Serialize))] #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))] #[derive(Clone, Copy, Debug, Eq, PartialEq)] - pub struct ShaderStages: u8 { + pub struct ShaderStages: u16 { const VERTEX = 0x1; const FRAGMENT = 0x2; const COMPUTE = 0x4; const MESH = 0x8; const TASK = 0x10; + const RAY_GENERATION = 0x20; + const ANY_HIT = 0x40; + const CLOSEST_HIT = 0x80; + const MISS = 0x100; } } @@ -362,6 +368,19 @@ pub struct Validator { /// [`Expression`]: crate::Expression /// [`Statement`]: crate::Statement needs_visit: HandleSet, + + /// Whether any trace rays call doesn't get called with an acceleration structure + /// with vertex return. This is in case another shader uses a vertex return only + /// builtin. If inner option is `Some`, then the span is of one of the `traceRay` + /// calls with a acceleration structure without vertex return. If the inner option + /// is `None` then the shader only uses acceleration structures with vertex return + /// in its trace ray calls. If the outer option is `None`, there are no `traceRay` + /// calls. + trace_rays_no_vertex_return: Option>, + + /// The type of the ray payload, this must always be the same type in a particular + /// entrypoint + trace_rays_payload_type: Option>, } #[derive(Clone, Debug, thiserror::Error)] @@ -560,6 +579,8 @@ impl Validator { override_ids: FastHashSet::default(), overrides_resolved: false, needs_visit: HandleSet::new(), + trace_rays_no_vertex_return: None, + trace_rays_payload_type: None, } } diff --git a/naga/src/valid/type.rs b/naga/src/valid/type.rs index aa0633e1852..55885eff1eb 100644 --- a/naga/src/valid/type.rs +++ b/naga/src/valid/type.rs @@ -219,7 +219,7 @@ fn check_member_layout( const fn ptr_space_argument_flag(space: crate::AddressSpace) -> TypeFlags { use crate::AddressSpace as As; match space { - As::Function | As::Private => TypeFlags::ARGUMENT, + As::Function | As::Private | As::RayPayload | As::IncomingRayPayload => TypeFlags::ARGUMENT, As::Uniform | As::Storage { .. } | As::Handle diff --git a/naga/tests/naga/snapshots.rs b/naga/tests/naga/snapshots.rs index e01610f29a8..f324a5029ed 100644 --- a/naga/tests/naga/snapshots.rs +++ b/naga/tests/naga/snapshots.rs @@ -351,7 +351,7 @@ fn write_output_hlsl( naga::ShaderStage::Vertex => &mut config.vertex, naga::ShaderStage::Fragment => &mut config.fragment, naga::ShaderStage::Compute => &mut config.compute, - naga::ShaderStage::Task | naga::ShaderStage::Mesh => unreachable!(), + naga::ShaderStage::Task | naga::ShaderStage::Mesh | naga::ShaderStage::RayGeneration | naga::ShaderStage::AnyHit | naga::ShaderStage::ClosestHit | naga::ShaderStage::Miss => unreachable!(), } .push(hlsl_snapshots::ConfigItem { entry_point: name.clone(), diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs index ffe2c7e7572..5c499af749d 100644 --- a/wgpu-core/src/validation.rs +++ b/wgpu-core/src/validation.rs @@ -1301,7 +1301,7 @@ impl Interface { } naga::ShaderStage::Compute => (false, 0), // TODO: add validation for these, see https://github.com/gfx-rs/wgpu/issues/8003 - naga::ShaderStage::Task | naga::ShaderStage::Mesh => { + naga::ShaderStage::Task | naga::ShaderStage::Mesh | naga::ShaderStage::RayGeneration | naga::ShaderStage::AnyHit | naga::ShaderStage::ClosestHit | naga::ShaderStage::Miss => { unreachable!() } }; diff --git a/wgpu-hal/src/auxil/mod.rs b/wgpu-hal/src/auxil/mod.rs index 3fef171d9bd..c4d3c3fc86f 100644 --- a/wgpu-hal/src/auxil/mod.rs +++ b/wgpu-hal/src/auxil/mod.rs @@ -62,6 +62,10 @@ pub fn map_naga_stage(stage: naga::ShaderStage) -> wgt::ShaderStages { naga::ShaderStage::Compute => wgt::ShaderStages::COMPUTE, naga::ShaderStage::Task => wgt::ShaderStages::TASK, naga::ShaderStage::Mesh => wgt::ShaderStages::MESH, + naga::ShaderStage::RayGeneration => wgt::ShaderStages::RAY_GENERATION, + naga::ShaderStage::AnyHit => wgt::ShaderStages::ANY_HIT, + naga::ShaderStage::ClosestHit => wgt::ShaderStages::CLOSEST_HIT, + naga::ShaderStage::Miss => wgt::ShaderStages::MISS, } } diff --git a/wgpu-hal/src/gles/device.rs b/wgpu-hal/src/gles/device.rs index 05741c18f74..dbdd8d3f0bc 100644 --- a/wgpu-hal/src/gles/device.rs +++ b/wgpu-hal/src/gles/device.rs @@ -99,7 +99,7 @@ impl CompilationContext<'_> { unsafe { gl.bind_frag_data_location(program, location.location, &name) } } naga::ShaderStage::Compute => {} - naga::ShaderStage::Task | naga::ShaderStage::Mesh => unreachable!(), + naga::ShaderStage::Task | naga::ShaderStage::Mesh | naga::ShaderStage::RayGeneration | naga::ShaderStage::AnyHit | naga::ShaderStage::ClosestHit | naga::ShaderStage::Miss => unreachable!(), } } @@ -175,7 +175,7 @@ impl super::Device { naga::ShaderStage::Vertex => glow::VERTEX_SHADER, naga::ShaderStage::Fragment => glow::FRAGMENT_SHADER, naga::ShaderStage::Compute => glow::COMPUTE_SHADER, - naga::ShaderStage::Task | naga::ShaderStage::Mesh => unreachable!(), + naga::ShaderStage::Task | naga::ShaderStage::Mesh | naga::ShaderStage::RayGeneration | naga::ShaderStage::AnyHit | naga::ShaderStage::ClosestHit | naga::ShaderStage::Miss => unreachable!(), }; let raw = unsafe { gl.create_shader(target) }.unwrap(); diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs index 6780bf3701d..05de5ded5d6 100644 --- a/wgpu-types/src/lib.rs +++ b/wgpu-types/src/lib.rs @@ -1568,6 +1568,14 @@ bitflags::bitflags! { const TASK = 1 << 3; /// Binding is visible from the mesh shader of a mesh pipeline. const MESH = 1 << 4; + /// Binding is visible from the ray generation shader of a ray tracing pipeline. + const RAY_GENERATION = 1 << 5; + /// Binding is visible from the ray any hit shader of a ray tracing pipeline. + const ANY_HIT = 1 << 5; + /// Binding is visible from the ray closest hit shader of a ray tracing pipeline. + const CLOSEST_HIT = 1 << 6; + /// Binding is visible from the ray miss shader of a ray tracing pipeline. + const MISS = 1 << 7; } } From ec60bfa555a54cb03292e35b25619288e9a4cc19 Mon Sep 17 00:00:00 2001 From: Vecvec Date: Sun, 23 Nov 2025 18:03:33 +1300 Subject: [PATCH 02/15] WGSL in --- naga/src/front/wgsl/error.rs | 9 ++ naga/src/front/wgsl/lower/mod.rs | 39 +++++++- naga/src/front/wgsl/parse/ast.rs | 1 + naga/src/front/wgsl/parse/conv.rs | 34 +++++++ .../wgsl/parse/directive/enable_extension.rs | 13 +++ naga/src/front/wgsl/parse/mod.rs | 98 ++++++++++++++++++- naga/tests/naga/snapshots.rs | 7 +- 7 files changed, 194 insertions(+), 7 deletions(-) diff --git a/naga/src/front/wgsl/error.rs b/naga/src/front/wgsl/error.rs index 0cd7e11c737..3db52fc10ee 100644 --- a/naga/src/front/wgsl/error.rs +++ b/naga/src/front/wgsl/error.rs @@ -415,6 +415,7 @@ pub(crate) enum Error<'a> { TypeTooLarge { span: Span, }, + MissingIncomingPayload(Span), } impl From for Error<'_> { @@ -1394,6 +1395,14 @@ impl<'a> Error<'a> { crate::valid::MAX_TYPE_SIZE )], }, + Error::MissingIncomingPayload(span) => ParseError { + message: "incoming payload is missing on ray hit or miss shader entry point".to_string(), + labels: vec![( + span, + "must be paired with a `@incoming_payload` attribute".into(), + )], + notes: vec![], + }, } } } diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index a9d34646b9a..8433dae9f2f 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -1629,6 +1629,21 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { None }; + let ray_incoming_payload = + if let Some((var_name, var_span)) = entry.ray_incoming_payload { + Some(match ctx.globals.get(var_name) { + Some(&LoweredGlobalDecl::Var(handle)) => handle, + Some(_) => { + return Err(Box::new(Error::ExpectedGlobalVariable { + name_span: var_span, + })) + } + None => return Err(Box::new(Error::UnknownIdent(var_span, var_name))), + }) + } else { + None + }; + ctx.module.entry_points.push(ir::EntryPoint { name: f.name.name.to_string(), stage: entry.stage, @@ -1638,7 +1653,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { function, mesh_info, task_payload, - ray_incoming_payload: None, + ray_incoming_payload, }); Ok(LoweredGlobalDecl::EntryPoint( ctx.module.entry_points.len() - 1, @@ -3383,6 +3398,28 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { ); return Ok(Some(result)); } + "traceRay" => { + let mut args = ctx.prepare_args(arguments, 3, span); + let acceleration_structure = self.expression(args.next()?, ctx)?; + let descriptor = self.expression(args.next()?, ctx)?; + let payload = self.expression(args.next()?, ctx)?; + args.finish()?; + + let _ = ctx.module.generate_ray_desc_type(); + let fun = ir::RayPipelineFunction::TraceRay { + acceleration_structure, + descriptor, + payload, + }; + + let rctx = ctx.runtime_expression_ctx(span)?; + rctx.block + .extend(rctx.emitter.finish(&rctx.function.expressions)); + rctx.emitter.start(&rctx.function.expressions); + rctx.block + .push(ir::Statement::RayPipelineFunction(fun), span); + return Ok(None); + } _ => { return Err(Box::new(Error::UnknownIdent(function.span, function.name))) } diff --git a/naga/src/front/wgsl/parse/ast.rs b/naga/src/front/wgsl/parse/ast.rs index 04964e7ba5f..2d6267d89c6 100644 --- a/naga/src/front/wgsl/parse/ast.rs +++ b/naga/src/front/wgsl/parse/ast.rs @@ -130,6 +130,7 @@ pub struct EntryPoint<'a> { pub workgroup_size: Option<[Option>>; 3]>, pub mesh_output_variable: Option<(&'a str, Span)>, pub task_payload: Option<(&'a str, Span)>, + pub ray_incoming_payload: Option<(&'a str, Span)>, } #[cfg(doc)] diff --git a/naga/src/front/wgsl/parse/conv.rs b/naga/src/front/wgsl/parse/conv.rs index 5431f0d2525..fc314806133 100644 --- a/naga/src/front/wgsl/parse/conv.rs +++ b/naga/src/front/wgsl/parse/conv.rs @@ -30,6 +30,26 @@ pub fn map_address_space<'a>( })) } } + "ray_payload" => { + if enable_extensions.contains(ImplementedEnableExtension::WgpuRayTracingPipeline) { + Ok(crate::AddressSpace::RayPayload) + } else { + Err(Box::new(Error::EnableExtensionNotEnabled { + span, + kind: ImplementedEnableExtension::WgpuRayTracingPipeline.into(), + })) + } + } + "incoming_ray_payload" => { + if enable_extensions.contains(ImplementedEnableExtension::WgpuRayTracingPipeline) { + Ok(crate::AddressSpace::IncomingRayPayload) + } else { + Err(Box::new(Error::EnableExtensionNotEnabled { + span, + kind: ImplementedEnableExtension::WgpuRayTracingPipeline.into(), + })) + } + } _ => Err(Box::new(Error::UnknownAddressSpace(span))), } } @@ -75,6 +95,20 @@ pub fn map_built_in( "vertices" => crate::BuiltIn::Vertices, "primitive_count" => crate::BuiltIn::PrimitiveCount, "primitives" => crate::BuiltIn::Primitives, + // ray tracing pipeline + "ray_invocation_id" => crate::BuiltIn::RayInvocationId, + "num_ray_invocations" => crate::BuiltIn::NumRayInvocations, + "instance_custom_data" => crate::BuiltIn::InstanceCustomData, + "geometry_index" => crate::BuiltIn::GeometryIndex, + "world_ray_origin" => crate::BuiltIn::WorldRayOrigin, + "world_ray_direction" => crate::BuiltIn::WorldRayDirection, + "object_ray_origin" => crate::BuiltIn::ObjectRayOrigin, + "object_ray_direction" => crate::BuiltIn::ObjectRayDirection, + "ray_t_min" => crate::BuiltIn::RayTmin, + "ray_t_current_max" => crate::BuiltIn::RayTCurrentMax, + "object_to_world" => crate::BuiltIn::ObjectToWorld, + "world_to_object" => crate::BuiltIn::WorldToObject, + "hit_kind" => crate::BuiltIn::HitKind, _ => return Err(Box::new(Error::UnknownBuiltin(span))), }; match built_in { diff --git a/naga/src/front/wgsl/parse/directive/enable_extension.rs b/naga/src/front/wgsl/parse/directive/enable_extension.rs index a4ca097df8f..13c163a6b8d 100644 --- a/naga/src/front/wgsl/parse/directive/enable_extension.rs +++ b/naga/src/front/wgsl/parse/directive/enable_extension.rs @@ -13,6 +13,7 @@ pub struct EnableExtensions { wgpu_mesh_shader: bool, wgpu_ray_query: bool, wgpu_ray_query_vertex_return: bool, + wgpu_ray_tracing_pipelines: bool, dual_source_blending: bool, /// Whether `enable f16;` was written earlier in the shader module. f16: bool, @@ -25,6 +26,7 @@ impl EnableExtensions { wgpu_mesh_shader: false, wgpu_ray_query: false, wgpu_ray_query_vertex_return: false, + wgpu_ray_tracing_pipelines: false, f16: false, dual_source_blending: false, clip_distances: false, @@ -39,6 +41,9 @@ impl EnableExtensions { ImplementedEnableExtension::WgpuRayQueryVertexReturn => { &mut self.wgpu_ray_query_vertex_return } + ImplementedEnableExtension::WgpuRayTracingPipeline => { + &mut self.wgpu_ray_tracing_pipelines + } ImplementedEnableExtension::DualSourceBlending => &mut self.dual_source_blending, ImplementedEnableExtension::F16 => &mut self.f16, ImplementedEnableExtension::ClipDistances => &mut self.clip_distances, @@ -54,6 +59,7 @@ impl EnableExtensions { ImplementedEnableExtension::WgpuRayQueryVertexReturn => { self.wgpu_ray_query_vertex_return } + ImplementedEnableExtension::WgpuRayTracingPipeline => self.wgpu_ray_tracing_pipelines, ImplementedEnableExtension::DualSourceBlending => self.dual_source_blending, ImplementedEnableExtension::F16 => self.f16, ImplementedEnableExtension::ClipDistances => self.clip_distances, @@ -89,6 +95,7 @@ impl EnableExtension { const MESH_SHADER: &'static str = "wgpu_mesh_shader"; const RAY_QUERY: &'static str = "wgpu_ray_query"; const RAY_QUERY_VERTEX_RETURN: &'static str = "wgpu_ray_query_vertex_return"; + const RAY_TRACING_PIPELINE: &'static str = "wgpu_ray_tracing_pipeline"; const SUBGROUPS: &'static str = "subgroups"; const PRIMITIVE_INDEX: &'static str = "primitive_index"; @@ -105,6 +112,9 @@ impl EnableExtension { Self::RAY_QUERY_VERTEX_RETURN => { Self::Implemented(ImplementedEnableExtension::WgpuRayQueryVertexReturn) } + Self::RAY_TRACING_PIPELINE => { + Self::Implemented(ImplementedEnableExtension::WgpuRayTracingPipeline) + } Self::SUBGROUPS => Self::Unimplemented(UnimplementedEnableExtension::Subgroups), Self::PRIMITIVE_INDEX => { Self::Unimplemented(UnimplementedEnableExtension::PrimitiveIndex) @@ -125,6 +135,7 @@ impl EnableExtension { ImplementedEnableExtension::DualSourceBlending => Self::DUAL_SOURCE_BLENDING, ImplementedEnableExtension::F16 => Self::F16, ImplementedEnableExtension::ClipDistances => Self::CLIP_DISTANCES, + ImplementedEnableExtension::WgpuRayTracingPipeline => Self::RAY_TRACING_PIPELINE, }, Self::Unimplemented(kind) => match kind { UnimplementedEnableExtension::Subgroups => Self::SUBGROUPS, @@ -161,6 +172,8 @@ pub enum ImplementedEnableExtension { WgpuRayQuery, /// Enables the `wgpu_ray_query_vertex_return` extension, native only. WgpuRayQueryVertexReturn, + /// Enables the `wgpu_ray_tracing_pipeline` extension, native only. + WgpuRayTracingPipeline, } /// A variant of [`EnableExtension::Unimplemented`]. diff --git a/naga/src/front/wgsl/parse/mod.rs b/naga/src/front/wgsl/parse/mod.rs index 142606b83e1..248b828f350 100644 --- a/naga/src/front/wgsl/parse/mod.rs +++ b/naga/src/front/wgsl/parse/mod.rs @@ -1933,7 +1933,11 @@ impl Parser { if !lexer .enable_extensions .contains(ImplementedEnableExtension::WgpuRayQuery) + && !lexer + .enable_extensions + .contains(ImplementedEnableExtension::WgpuRayTracingPipeline) { + // maybe we want a multi enable extension error? return Err(Box::new(Error::EnableExtensionNotEnabled { kind: EnableExtension::Implemented( ImplementedEnableExtension::WgpuRayQuery, @@ -2877,13 +2881,18 @@ impl Parser { // read attributes let mut binding = None; let mut stage = ParsedAttribute::default(); - let mut compute_like_span = Span::new(0, 0); + // Span in case we need to report an error for a shader stage missing something (e.g. its workgroup size). + // Doesn't need to be set in the vertex and fragment stages because they don't have errors like that. + let mut shader_stage_error_span = Span::new(0, 0); let mut workgroup_size = ParsedAttribute::default(); let mut early_depth_test = ParsedAttribute::default(); let (mut bind_index, mut bind_group) = (ParsedAttribute::default(), ParsedAttribute::default()); let mut id = ParsedAttribute::default(); + // the payload variable for a mesh shader let mut payload = ParsedAttribute::default(); + // the incoming payload from a traceRay call + let mut incoming_payload = ParsedAttribute::default(); let mut mesh_output = ParsedAttribute::default(); let mut must_use: ParsedAttribute = ParsedAttribute::default(); @@ -2943,7 +2952,7 @@ impl Parser { } "compute" => { stage.set(ShaderStage::Compute, name_span)?; - compute_like_span = name_span; + shader_stage_error_span = name_span; } "task" => { if !lexer @@ -2956,7 +2965,7 @@ impl Parser { })); } stage.set(ShaderStage::Task, name_span)?; - compute_like_span = name_span; + shader_stage_error_span = name_span; } "mesh" => { if !lexer @@ -2969,12 +2978,63 @@ impl Parser { })); } stage.set(ShaderStage::Mesh, name_span)?; - compute_like_span = name_span; + shader_stage_error_span = name_span; lexer.expect(Token::Paren('('))?; mesh_output.set(lexer.next_ident_with_span()?, name_span)?; lexer.expect(Token::Paren(')'))?; } + "ray_generation" => { + if !lexer + .enable_extensions + .contains(ImplementedEnableExtension::WgpuRayTracingPipeline) + { + return Err(Box::new(Error::EnableExtensionNotEnabled { + span: name_span, + kind: ImplementedEnableExtension::WgpuRayTracingPipeline.into(), + })); + } + stage.set(ShaderStage::RayGeneration, name_span)?; + } + "any_hit" => { + if !lexer + .enable_extensions + .contains(ImplementedEnableExtension::WgpuRayTracingPipeline) + { + return Err(Box::new(Error::EnableExtensionNotEnabled { + span: name_span, + kind: ImplementedEnableExtension::WgpuRayTracingPipeline.into(), + })); + } + stage.set(ShaderStage::AnyHit, name_span)?; + shader_stage_error_span = name_span; + } + "closest_hit" => { + if !lexer + .enable_extensions + .contains(ImplementedEnableExtension::WgpuRayTracingPipeline) + { + return Err(Box::new(Error::EnableExtensionNotEnabled { + span: name_span, + kind: ImplementedEnableExtension::WgpuRayTracingPipeline.into(), + })); + } + stage.set(ShaderStage::ClosestHit, name_span)?; + shader_stage_error_span = name_span; + } + "miss" => { + if !lexer + .enable_extensions + .contains(ImplementedEnableExtension::WgpuRayTracingPipeline) + { + return Err(Box::new(Error::EnableExtensionNotEnabled { + span: name_span, + kind: ImplementedEnableExtension::WgpuRayTracingPipeline.into(), + })); + } + stage.set(ShaderStage::Miss, name_span)?; + shader_stage_error_span = name_span; + } "payload" => { if !lexer .enable_extensions @@ -2989,6 +3049,20 @@ impl Parser { payload.set(lexer.next_ident_with_span()?, name_span)?; lexer.expect(Token::Paren(')'))?; } + "incoming_payload" => { + if !lexer + .enable_extensions + .contains(ImplementedEnableExtension::WgpuRayTracingPipeline) + { + return Err(Box::new(Error::EnableExtensionNotEnabled { + span: name_span, + kind: ImplementedEnableExtension::WgpuRayTracingPipeline.into(), + })); + } + lexer.expect(Token::Paren('('))?; + incoming_payload.set(lexer.next_ident_with_span()?, name_span)?; + lexer.expect(Token::Paren(')'))?; + } "workgroup_size" => { lexer.expect(Token::Paren('('))?; let mut new_workgroup_size = [None; 3]; @@ -3154,7 +3228,20 @@ impl Parser { Some(ast::GlobalDeclKind::Fn(ast::Function { entry_point: if let Some(stage) = stage.value { if stage.compute_like() && workgroup_size.value.is_none() { - return Err(Box::new(Error::MissingWorkgroupSize(compute_like_span))); + return Err(Box::new(Error::MissingWorkgroupSize( + shader_stage_error_span, + ))); + } + + match stage { + ShaderStage::AnyHit | ShaderStage::ClosestHit | ShaderStage::Miss => { + if incoming_payload.value.is_none() { + return Err(Box::new(Error::MissingIncomingPayload( + shader_stage_error_span, + ))); + } + } + _ => {} } Some(ast::EntryPoint { @@ -3163,6 +3250,7 @@ impl Parser { workgroup_size: workgroup_size.value, mesh_output_variable: mesh_output.value, task_payload: payload.value, + ray_incoming_payload: incoming_payload.value, }) } else { None diff --git a/naga/tests/naga/snapshots.rs b/naga/tests/naga/snapshots.rs index f324a5029ed..52062a70b3f 100644 --- a/naga/tests/naga/snapshots.rs +++ b/naga/tests/naga/snapshots.rs @@ -351,7 +351,12 @@ fn write_output_hlsl( naga::ShaderStage::Vertex => &mut config.vertex, naga::ShaderStage::Fragment => &mut config.fragment, naga::ShaderStage::Compute => &mut config.compute, - naga::ShaderStage::Task | naga::ShaderStage::Mesh | naga::ShaderStage::RayGeneration | naga::ShaderStage::AnyHit | naga::ShaderStage::ClosestHit | naga::ShaderStage::Miss => unreachable!(), + naga::ShaderStage::Task + | naga::ShaderStage::Mesh + | naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unreachable!(), } .push(hlsl_snapshots::ConfigItem { entry_point: name.clone(), From 08925f2a274bee2609743d282292db303d7621e1 Mon Sep 17 00:00:00 2001 From: Vecvec Date: Mon, 24 Nov 2025 12:19:22 +1300 Subject: [PATCH 03/15] add and rerun snapshots --- naga/src/valid/interface.rs | 2 +- naga/tests/in/wgsl/ray-tracing-pipeline.toml | 2 + naga/tests/in/wgsl/ray-tracing-pipeline.wgsl | 34 ++ naga/tests/out/analysis/spv-shadow.info.ron | 6 +- naga/tests/out/analysis/wgsl-access.info.ron | 38 +- naga/tests/out/analysis/wgsl-collatz.info.ron | 4 +- .../analysis/wgsl-mesh-shader-empty.info.ron | 4 +- .../analysis/wgsl-mesh-shader-lines.info.ron | 4 +- .../analysis/wgsl-mesh-shader-points.info.ron | 4 +- .../out/analysis/wgsl-mesh-shader.info.ron | 6 +- .../out/analysis/wgsl-overrides.info.ron | 2 +- .../analysis/wgsl-storage-textures.info.ron | 4 +- naga/tests/out/ir/spv-fetch_depth.compact.ron | 1 + naga/tests/out/ir/spv-fetch_depth.ron | 1 + naga/tests/out/ir/spv-shadow.compact.ron | 1 + naga/tests/out/ir/spv-shadow.ron | 1 + .../out/ir/spv-spec-constants.compact.ron | 1 + naga/tests/out/ir/spv-spec-constants.ron | 1 + naga/tests/out/ir/wgsl-access.compact.ron | 3 + naga/tests/out/ir/wgsl-access.ron | 3 + naga/tests/out/ir/wgsl-collatz.compact.ron | 1 + naga/tests/out/ir/wgsl-collatz.ron | 1 + .../out/ir/wgsl-const_assert.compact.ron | 1 + naga/tests/out/ir/wgsl-const_assert.ron | 1 + .../out/ir/wgsl-diagnostic-filter.compact.ron | 1 + naga/tests/out/ir/wgsl-diagnostic-filter.ron | 1 + .../out/ir/wgsl-index-by-value.compact.ron | 1 + naga/tests/out/ir/wgsl-index-by-value.ron | 1 + .../tests/out/ir/wgsl-local-const.compact.ron | 1 + naga/tests/out/ir/wgsl-local-const.ron | 1 + .../out/ir/wgsl-mesh-shader-empty.compact.ron | 2 + naga/tests/out/ir/wgsl-mesh-shader-empty.ron | 2 + .../out/ir/wgsl-mesh-shader-lines.compact.ron | 2 + naga/tests/out/ir/wgsl-mesh-shader-lines.ron | 2 + .../ir/wgsl-mesh-shader-points.compact.ron | 2 + naga/tests/out/ir/wgsl-mesh-shader-points.ron | 2 + .../tests/out/ir/wgsl-mesh-shader.compact.ron | 3 + naga/tests/out/ir/wgsl-mesh-shader.ron | 3 + naga/tests/out/ir/wgsl-must-use.compact.ron | 1 + naga/tests/out/ir/wgsl-must-use.ron | 1 + ...ides-atomicCompareExchangeWeak.compact.ron | 1 + ...sl-overrides-atomicCompareExchangeWeak.ron | 1 + .../ir/wgsl-overrides-ray-query.compact.ron | 1 + .../tests/out/ir/wgsl-overrides-ray-query.ron | 1 + naga/tests/out/ir/wgsl-overrides.compact.ron | 1 + naga/tests/out/ir/wgsl-overrides.ron | 1 + .../ir/wgsl-ray-tracing-pipeline.compact.ron | 350 ++++++++++++++++++ .../out/ir/wgsl-ray-tracing-pipeline.ron | 350 ++++++++++++++++++ .../out/ir/wgsl-storage-textures.compact.ron | 2 + naga/tests/out/ir/wgsl-storage-textures.ron | 2 + ...l-template-list-trailing-comma.compact.ron | 1 + .../ir/wgsl-template-list-trailing-comma.ron | 1 + .../out/ir/wgsl-texture-external.compact.ron | 3 + naga/tests/out/ir/wgsl-texture-external.ron | 3 + .../ir/wgsl-types_with_comments.compact.ron | 1 + .../tests/out/ir/wgsl-types_with_comments.ron | 1 + 56 files changed, 835 insertions(+), 37 deletions(-) create mode 100644 naga/tests/in/wgsl/ray-tracing-pipeline.toml create mode 100644 naga/tests/in/wgsl/ray-tracing-pipeline.wgsl create mode 100644 naga/tests/out/ir/wgsl-ray-tracing-pipeline.compact.ron create mode 100644 naga/tests/out/ir/wgsl-ray-tracing-pipeline.ron diff --git a/naga/src/valid/interface.rs b/naga/src/valid/interface.rs index 4ca46be6c05..a975ebf5bc4 100644 --- a/naga/src/valid/interface.rs +++ b/naga/src/valid/interface.rs @@ -1315,7 +1315,7 @@ impl super::Validator { } crate::AddressSpace::PushConstant => GlobalUse::READ, crate::AddressSpace::RayPayload => { - if matches!(ep.stage, crate::ShaderStage::RayGeneration | crate::ShaderStage::ClosestHit | crate::ShaderStage::Miss) { + if !matches!(ep.stage, crate::ShaderStage::RayGeneration | crate::ShaderStage::ClosestHit | crate::ShaderStage::Miss) { return Err(EntryPointError::RayPayloadInInvalidStage.with_span_handle(var_handle, &module.global_variables)); } GlobalUse::READ diff --git a/naga/tests/in/wgsl/ray-tracing-pipeline.toml b/naga/tests/in/wgsl/ray-tracing-pipeline.toml new file mode 100644 index 00000000000..f4611d2d304 --- /dev/null +++ b/naga/tests/in/wgsl/ray-tracing-pipeline.toml @@ -0,0 +1,2 @@ +god_mode = true +targets = "IR" \ No newline at end of file diff --git a/naga/tests/in/wgsl/ray-tracing-pipeline.wgsl b/naga/tests/in/wgsl/ray-tracing-pipeline.wgsl new file mode 100644 index 00000000000..02a40b05c5d --- /dev/null +++ b/naga/tests/in/wgsl/ray-tracing-pipeline.wgsl @@ -0,0 +1,34 @@ +enable wgpu_ray_tracing_pipeline; + +struct HitCounters { + hit_num: u32, + selected_hit: u32, +} + +var hit_num: HitCounters; + +@group(0) @binding(0) +var acc_struct: acceleration_structure; + +@ray_generation +fn ray_gen_main() { + hit_num = HitCounters(); + traceRay(acc_struct, RayDesc(RAY_FLAG_NONE, 0xff, 0.01, 100.0, vec3(0.0), vec3(0.0, 1.0, 0.0)), &hit_num); +} + +var incoming_hit_num: HitCounters; + +@miss +@incoming_payload(incoming_hit_num) +fn miss() {} + +@any_hit +@incoming_payload(incoming_hit_num) +fn any_hit_main() { + incoming_hit_num.hit_num++; + incoming_hit_num.selected_hit = incoming_hit_num.hit_num; +} + +@closest_hit +@incoming_payload(incoming_hit_num) +fn closest_hit_main() {} \ No newline at end of file diff --git a/naga/tests/out/analysis/spv-shadow.info.ron b/naga/tests/out/analysis/spv-shadow.info.ron index 381f841d5d9..c2adf7d22dd 100644 --- a/naga/tests/out/analysis/spv-shadow.info.ron +++ b/naga/tests/out/analysis/spv-shadow.info.ron @@ -18,7 +18,7 @@ functions: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(1), requirements: (""), @@ -416,7 +416,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(1), requirements: (""), @@ -1596,7 +1596,7 @@ entry_points: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(1), requirements: (""), diff --git a/naga/tests/out/analysis/wgsl-access.info.ron b/naga/tests/out/analysis/wgsl-access.info.ron index c22cd768f2e..32ba1b0b291 100644 --- a/naga/tests/out/analysis/wgsl-access.info.ron +++ b/naga/tests/out/analysis/wgsl-access.info.ron @@ -42,7 +42,7 @@ functions: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -1200,7 +1200,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -2526,7 +2526,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(0), requirements: (""), @@ -2566,7 +2566,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(0), requirements: (""), @@ -2615,7 +2615,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -2658,7 +2658,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -2752,7 +2752,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -2873,7 +2873,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(0), requirements: (""), @@ -2925,7 +2925,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -2980,7 +2980,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(0), requirements: (""), @@ -3032,7 +3032,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -3087,7 +3087,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(0), requirements: (""), @@ -3151,7 +3151,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(2), requirements: (""), @@ -3224,7 +3224,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(2), requirements: (""), @@ -3300,7 +3300,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -3400,7 +3400,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(1), requirements: (""), @@ -3598,7 +3598,7 @@ entry_points: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(0), requirements: (""), @@ -4293,7 +4293,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -4745,7 +4745,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(0), requirements: (""), diff --git a/naga/tests/out/analysis/wgsl-collatz.info.ron b/naga/tests/out/analysis/wgsl-collatz.info.ron index 219e016f8d7..9647887a0db 100644 --- a/naga/tests/out/analysis/wgsl-collatz.info.ron +++ b/naga/tests/out/analysis/wgsl-collatz.info.ron @@ -8,7 +8,7 @@ functions: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(3), requirements: (""), @@ -280,7 +280,7 @@ entry_points: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(3), requirements: (""), diff --git a/naga/tests/out/analysis/wgsl-mesh-shader-empty.info.ron b/naga/tests/out/analysis/wgsl-mesh-shader-empty.info.ron index 9f9feac9c09..b148e4d7bcd 100644 --- a/naga/tests/out/analysis/wgsl-mesh-shader-empty.info.ron +++ b/naga/tests/out/analysis/wgsl-mesh-shader-empty.info.ron @@ -14,7 +14,7 @@ entry_points: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -78,7 +78,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), diff --git a/naga/tests/out/analysis/wgsl-mesh-shader-lines.info.ron b/naga/tests/out/analysis/wgsl-mesh-shader-lines.info.ron index b4c8790508d..efd02461074 100644 --- a/naga/tests/out/analysis/wgsl-mesh-shader-lines.info.ron +++ b/naga/tests/out/analysis/wgsl-mesh-shader-lines.info.ron @@ -15,7 +15,7 @@ entry_points: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -79,7 +79,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), diff --git a/naga/tests/out/analysis/wgsl-mesh-shader-points.info.ron b/naga/tests/out/analysis/wgsl-mesh-shader-points.info.ron index fa62867576b..479da9b9c9e 100644 --- a/naga/tests/out/analysis/wgsl-mesh-shader-points.info.ron +++ b/naga/tests/out/analysis/wgsl-mesh-shader-points.info.ron @@ -14,7 +14,7 @@ entry_points: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -78,7 +78,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), diff --git a/naga/tests/out/analysis/wgsl-mesh-shader.info.ron b/naga/tests/out/analysis/wgsl-mesh-shader.info.ron index eacd33ad0f1..4018c5de449 100644 --- a/naga/tests/out/analysis/wgsl-mesh-shader.info.ron +++ b/naga/tests/out/analysis/wgsl-mesh-shader.info.ron @@ -17,7 +17,7 @@ entry_points: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -223,7 +223,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -1401,7 +1401,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: Some(0), requirements: (""), diff --git a/naga/tests/out/analysis/wgsl-overrides.info.ron b/naga/tests/out/analysis/wgsl-overrides.info.ron index 92e99112e53..d6fcee4aab3 100644 --- a/naga/tests/out/analysis/wgsl-overrides.info.ron +++ b/naga/tests/out/analysis/wgsl-overrides.info.ron @@ -8,7 +8,7 @@ entry_points: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), diff --git a/naga/tests/out/analysis/wgsl-storage-textures.info.ron b/naga/tests/out/analysis/wgsl-storage-textures.info.ron index 8bb298a6450..95b9484bd49 100644 --- a/naga/tests/out/analysis/wgsl-storage-textures.info.ron +++ b/naga/tests/out/analysis/wgsl-storage-textures.info.ron @@ -11,7 +11,7 @@ entry_points: [ ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), @@ -187,7 +187,7 @@ ), ( flags: ("EXPRESSIONS | BLOCKS | CONTROL_FLOW_UNIFORMITY | STRUCT_LAYOUTS | CONSTANTS | BINDINGS"), - available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK"), + available_stages: ("VERTEX | FRAGMENT | COMPUTE | MESH | TASK | RAY_GENERATION | ANY_HIT | CLOSEST_HIT | MISS"), uniformity: ( non_uniform_result: None, requirements: (""), diff --git a/naga/tests/out/ir/spv-fetch_depth.compact.ron b/naga/tests/out/ir/spv-fetch_depth.compact.ron index 98f4426c3eb..6a91852a1cf 100644 --- a/naga/tests/out/ir/spv-fetch_depth.compact.ron +++ b/naga/tests/out/ir/spv-fetch_depth.compact.ron @@ -198,6 +198,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/spv-fetch_depth.ron b/naga/tests/out/ir/spv-fetch_depth.ron index 104de852c17..32eafce671e 100644 --- a/naga/tests/out/ir/spv-fetch_depth.ron +++ b/naga/tests/out/ir/spv-fetch_depth.ron @@ -268,6 +268,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/spv-shadow.compact.ron b/naga/tests/out/ir/spv-shadow.compact.ron index bed86a5334d..e558039862b 100644 --- a/naga/tests/out/ir/spv-shadow.compact.ron +++ b/naga/tests/out/ir/spv-shadow.compact.ron @@ -1037,6 +1037,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/spv-shadow.ron b/naga/tests/out/ir/spv-shadow.ron index bdda1d18566..010ed83eba1 100644 --- a/naga/tests/out/ir/spv-shadow.ron +++ b/naga/tests/out/ir/spv-shadow.ron @@ -1315,6 +1315,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/spv-spec-constants.compact.ron b/naga/tests/out/ir/spv-spec-constants.compact.ron index 67eb29c2475..748f599b370 100644 --- a/naga/tests/out/ir/spv-spec-constants.compact.ron +++ b/naga/tests/out/ir/spv-spec-constants.compact.ron @@ -619,6 +619,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/spv-spec-constants.ron b/naga/tests/out/ir/spv-spec-constants.ron index 51686aa20eb..f5a97c51159 100644 --- a/naga/tests/out/ir/spv-spec-constants.ron +++ b/naga/tests/out/ir/spv-spec-constants.ron @@ -725,6 +725,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-access.compact.ron b/naga/tests/out/ir/wgsl-access.compact.ron index c3df0c8c500..a29a0777678 100644 --- a/naga/tests/out/ir/wgsl-access.compact.ron +++ b/naga/tests/out/ir/wgsl-access.compact.ron @@ -2657,6 +2657,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ( name: "foo_frag", @@ -2853,6 +2854,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ( name: "foo_compute", @@ -2914,6 +2916,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-access.ron b/naga/tests/out/ir/wgsl-access.ron index c3df0c8c500..a29a0777678 100644 --- a/naga/tests/out/ir/wgsl-access.ron +++ b/naga/tests/out/ir/wgsl-access.ron @@ -2657,6 +2657,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ( name: "foo_frag", @@ -2853,6 +2854,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ( name: "foo_compute", @@ -2914,6 +2916,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-collatz.compact.ron b/naga/tests/out/ir/wgsl-collatz.compact.ron index fc4daaa1296..2f55d55f7a6 100644 --- a/naga/tests/out/ir/wgsl-collatz.compact.ron +++ b/naga/tests/out/ir/wgsl-collatz.compact.ron @@ -336,6 +336,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-collatz.ron b/naga/tests/out/ir/wgsl-collatz.ron index fc4daaa1296..2f55d55f7a6 100644 --- a/naga/tests/out/ir/wgsl-collatz.ron +++ b/naga/tests/out/ir/wgsl-collatz.ron @@ -336,6 +336,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-const_assert.compact.ron b/naga/tests/out/ir/wgsl-const_assert.compact.ron index 648f4ff9bc9..5b312fcc773 100644 --- a/naga/tests/out/ir/wgsl-const_assert.compact.ron +++ b/naga/tests/out/ir/wgsl-const_assert.compact.ron @@ -36,6 +36,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-const_assert.ron b/naga/tests/out/ir/wgsl-const_assert.ron index 648f4ff9bc9..5b312fcc773 100644 --- a/naga/tests/out/ir/wgsl-const_assert.ron +++ b/naga/tests/out/ir/wgsl-const_assert.ron @@ -36,6 +36,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-diagnostic-filter.compact.ron b/naga/tests/out/ir/wgsl-diagnostic-filter.compact.ron index 9a2bf193f30..c1285a3abaf 100644 --- a/naga/tests/out/ir/wgsl-diagnostic-filter.compact.ron +++ b/naga/tests/out/ir/wgsl-diagnostic-filter.compact.ron @@ -75,6 +75,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [ diff --git a/naga/tests/out/ir/wgsl-diagnostic-filter.ron b/naga/tests/out/ir/wgsl-diagnostic-filter.ron index 9a2bf193f30..c1285a3abaf 100644 --- a/naga/tests/out/ir/wgsl-diagnostic-filter.ron +++ b/naga/tests/out/ir/wgsl-diagnostic-filter.ron @@ -75,6 +75,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [ diff --git a/naga/tests/out/ir/wgsl-index-by-value.compact.ron b/naga/tests/out/ir/wgsl-index-by-value.compact.ron index addd0e5871c..66b88c70476 100644 --- a/naga/tests/out/ir/wgsl-index-by-value.compact.ron +++ b/naga/tests/out/ir/wgsl-index-by-value.compact.ron @@ -467,6 +467,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-index-by-value.ron b/naga/tests/out/ir/wgsl-index-by-value.ron index addd0e5871c..66b88c70476 100644 --- a/naga/tests/out/ir/wgsl-index-by-value.ron +++ b/naga/tests/out/ir/wgsl-index-by-value.ron @@ -467,6 +467,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-local-const.compact.ron b/naga/tests/out/ir/wgsl-local-const.compact.ron index 0e4e2e4d40e..6a401b17b7c 100644 --- a/naga/tests/out/ir/wgsl-local-const.compact.ron +++ b/naga/tests/out/ir/wgsl-local-const.compact.ron @@ -102,6 +102,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-local-const.ron b/naga/tests/out/ir/wgsl-local-const.ron index 0e4e2e4d40e..6a401b17b7c 100644 --- a/naga/tests/out/ir/wgsl-local-const.ron +++ b/naga/tests/out/ir/wgsl-local-const.ron @@ -102,6 +102,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-mesh-shader-empty.compact.ron b/naga/tests/out/ir/wgsl-mesh-shader-empty.compact.ron index 8bdaa72da69..72f20f638c6 100644 --- a/naga/tests/out/ir/wgsl-mesh-shader-empty.compact.ron +++ b/naga/tests/out/ir/wgsl-mesh-shader-empty.compact.ron @@ -194,6 +194,7 @@ ), mesh_info: None, task_payload: Some(0), + ray_incoming_payload: None, ), ( name: "ms_main", @@ -226,6 +227,7 @@ output_variable: 1, )), task_payload: Some(0), + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-mesh-shader-empty.ron b/naga/tests/out/ir/wgsl-mesh-shader-empty.ron index 8bdaa72da69..72f20f638c6 100644 --- a/naga/tests/out/ir/wgsl-mesh-shader-empty.ron +++ b/naga/tests/out/ir/wgsl-mesh-shader-empty.ron @@ -194,6 +194,7 @@ ), mesh_info: None, task_payload: Some(0), + ray_incoming_payload: None, ), ( name: "ms_main", @@ -226,6 +227,7 @@ output_variable: 1, )), task_payload: Some(0), + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-mesh-shader-lines.compact.ron b/naga/tests/out/ir/wgsl-mesh-shader-lines.compact.ron index b298f296985..1759974b286 100644 --- a/naga/tests/out/ir/wgsl-mesh-shader-lines.compact.ron +++ b/naga/tests/out/ir/wgsl-mesh-shader-lines.compact.ron @@ -204,6 +204,7 @@ ), mesh_info: None, task_payload: Some(0), + ray_incoming_payload: None, ), ( name: "ms_main", @@ -236,6 +237,7 @@ output_variable: 1, )), task_payload: Some(0), + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-mesh-shader-lines.ron b/naga/tests/out/ir/wgsl-mesh-shader-lines.ron index b298f296985..1759974b286 100644 --- a/naga/tests/out/ir/wgsl-mesh-shader-lines.ron +++ b/naga/tests/out/ir/wgsl-mesh-shader-lines.ron @@ -204,6 +204,7 @@ ), mesh_info: None, task_payload: Some(0), + ray_incoming_payload: None, ), ( name: "ms_main", @@ -236,6 +237,7 @@ output_variable: 1, )), task_payload: Some(0), + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-mesh-shader-points.compact.ron b/naga/tests/out/ir/wgsl-mesh-shader-points.compact.ron index 558a88e28d8..05c0bf7dfe1 100644 --- a/naga/tests/out/ir/wgsl-mesh-shader-points.compact.ron +++ b/naga/tests/out/ir/wgsl-mesh-shader-points.compact.ron @@ -194,6 +194,7 @@ ), mesh_info: None, task_payload: Some(0), + ray_incoming_payload: None, ), ( name: "ms_main", @@ -226,6 +227,7 @@ output_variable: 1, )), task_payload: Some(0), + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-mesh-shader-points.ron b/naga/tests/out/ir/wgsl-mesh-shader-points.ron index 558a88e28d8..05c0bf7dfe1 100644 --- a/naga/tests/out/ir/wgsl-mesh-shader-points.ron +++ b/naga/tests/out/ir/wgsl-mesh-shader-points.ron @@ -194,6 +194,7 @@ ), mesh_info: None, task_payload: Some(0), + ray_incoming_payload: None, ), ( name: "ms_main", @@ -226,6 +227,7 @@ output_variable: 1, )), task_payload: Some(0), + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-mesh-shader.compact.ron b/naga/tests/out/ir/wgsl-mesh-shader.compact.ron index 9cb3c6479c3..55bb0ea5d01 100644 --- a/naga/tests/out/ir/wgsl-mesh-shader.compact.ron +++ b/naga/tests/out/ir/wgsl-mesh-shader.compact.ron @@ -325,6 +325,7 @@ ), mesh_info: None, task_payload: Some(0), + ray_incoming_payload: None, ), ( name: "ms_main", @@ -906,6 +907,7 @@ output_variable: 2, )), task_payload: Some(0), + ray_incoming_payload: None, ), ( name: "fs_main", @@ -972,6 +974,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-mesh-shader.ron b/naga/tests/out/ir/wgsl-mesh-shader.ron index 9cb3c6479c3..55bb0ea5d01 100644 --- a/naga/tests/out/ir/wgsl-mesh-shader.ron +++ b/naga/tests/out/ir/wgsl-mesh-shader.ron @@ -325,6 +325,7 @@ ), mesh_info: None, task_payload: Some(0), + ray_incoming_payload: None, ), ( name: "ms_main", @@ -906,6 +907,7 @@ output_variable: 2, )), task_payload: Some(0), + ray_incoming_payload: None, ), ( name: "fs_main", @@ -972,6 +974,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-must-use.compact.ron b/naga/tests/out/ir/wgsl-must-use.compact.ron index 16e925f2fb8..753db1f579d 100644 --- a/naga/tests/out/ir/wgsl-must-use.compact.ron +++ b/naga/tests/out/ir/wgsl-must-use.compact.ron @@ -203,6 +203,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-must-use.ron b/naga/tests/out/ir/wgsl-must-use.ron index 16e925f2fb8..753db1f579d 100644 --- a/naga/tests/out/ir/wgsl-must-use.ron +++ b/naga/tests/out/ir/wgsl-must-use.ron @@ -203,6 +203,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.compact.ron b/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.compact.ron index 28a824bb035..e0e5f68bd95 100644 --- a/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.compact.ron +++ b/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.compact.ron @@ -130,6 +130,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.ron b/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.ron index 28a824bb035..e0e5f68bd95 100644 --- a/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.ron +++ b/naga/tests/out/ir/wgsl-overrides-atomicCompareExchangeWeak.ron @@ -130,6 +130,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-overrides-ray-query.compact.ron b/naga/tests/out/ir/wgsl-overrides-ray-query.compact.ron index 152a45008c5..a3e0e5384f8 100644 --- a/naga/tests/out/ir/wgsl-overrides-ray-query.compact.ron +++ b/naga/tests/out/ir/wgsl-overrides-ray-query.compact.ron @@ -265,6 +265,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-overrides-ray-query.ron b/naga/tests/out/ir/wgsl-overrides-ray-query.ron index 152a45008c5..a3e0e5384f8 100644 --- a/naga/tests/out/ir/wgsl-overrides-ray-query.ron +++ b/naga/tests/out/ir/wgsl-overrides-ray-query.ron @@ -265,6 +265,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-overrides.compact.ron b/naga/tests/out/ir/wgsl-overrides.compact.ron index fe136e71e4d..95db50b89a6 100644 --- a/naga/tests/out/ir/wgsl-overrides.compact.ron +++ b/naga/tests/out/ir/wgsl-overrides.compact.ron @@ -223,6 +223,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-overrides.ron b/naga/tests/out/ir/wgsl-overrides.ron index fe136e71e4d..95db50b89a6 100644 --- a/naga/tests/out/ir/wgsl-overrides.ron +++ b/naga/tests/out/ir/wgsl-overrides.ron @@ -223,6 +223,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-ray-tracing-pipeline.compact.ron b/naga/tests/out/ir/wgsl-ray-tracing-pipeline.compact.ron new file mode 100644 index 00000000000..61148d02682 --- /dev/null +++ b/naga/tests/out/ir/wgsl-ray-tracing-pipeline.compact.ron @@ -0,0 +1,350 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: Some("HitCounters"), + inner: Struct( + members: [ + ( + name: Some("hit_num"), + ty: 0, + binding: None, + offset: 0, + ), + ( + name: Some("selected_hit"), + ty: 0, + binding: None, + offset: 4, + ), + ], + span: 8, + ), + ), + ( + name: None, + inner: AccelerationStructure( + vertex_return: false, + ), + ), + ( + name: None, + inner: Scalar(( + kind: Float, + width: 4, + )), + ), + ( + name: None, + inner: Vector( + size: Tri, + scalar: ( + kind: Float, + width: 4, + ), + ), + ), + ( + name: Some("RayDesc"), + inner: Struct( + members: [ + ( + name: Some("flags"), + ty: 0, + binding: None, + offset: 0, + ), + ( + name: Some("cull_mask"), + ty: 0, + binding: None, + offset: 4, + ), + ( + name: Some("tmin"), + ty: 3, + binding: None, + offset: 8, + ), + ( + name: Some("tmax"), + ty: 3, + binding: None, + offset: 12, + ), + ( + name: Some("origin"), + ty: 4, + binding: None, + offset: 16, + ), + ( + name: Some("dir"), + ty: 4, + binding: None, + offset: 32, + ), + ], + span: 48, + ), + ), + ], + special_types: ( + ray_desc: Some(5), + ray_intersection: None, + ray_vertex_return: None, + external_texture_params: None, + external_texture_transfer_function: None, + predeclared_types: {}, + ), + constants: [], + overrides: [], + global_variables: [ + ( + name: Some("hit_num"), + space: RayPayload, + binding: None, + ty: 1, + init: None, + ), + ( + name: Some("acc_struct"), + space: Handle, + binding: Some(( + group: 0, + binding: 0, + )), + ty: 2, + init: None, + ), + ( + name: Some("incoming_hit_num"), + space: IncomingRayPayload, + binding: None, + ty: 1, + init: None, + ), + ], + global_expressions: [], + functions: [], + entry_points: [ + ( + name: "ray_gen_main", + stage: RayGeneration, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("ray_gen_main"), + arguments: [], + result: None, + local_variables: [], + expressions: [ + GlobalVariable(0), + ZeroValue(1), + GlobalVariable(1), + Literal(U32(0)), + Literal(U32(255)), + Literal(F32(0.01)), + Literal(F32(100.0)), + Literal(F32(0.0)), + Splat( + size: Tri, + value: 7, + ), + Literal(F32(0.0)), + Literal(F32(1.0)), + Literal(F32(0.0)), + Compose( + ty: 4, + components: [ + 9, + 10, + 11, + ], + ), + Compose( + ty: 5, + components: [ + 3, + 4, + 5, + 6, + 8, + 12, + ], + ), + GlobalVariable(0), + ], + named_expressions: {}, + body: [ + Store( + pointer: 0, + value: 1, + ), + Emit(( + start: 0, + end: 0, + )), + Emit(( + start: 0, + end: 0, + )), + Emit(( + start: 8, + end: 9, + )), + Emit(( + start: 12, + end: 14, + )), + RayPipelineFunction(TraceRay( + acceleration_structure: 2, + descriptor: 13, + payload: 14, + )), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + mesh_info: None, + task_payload: None, + ray_incoming_payload: None, + ), + ( + name: "miss", + stage: Miss, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("miss"), + arguments: [], + result: None, + local_variables: [], + expressions: [], + named_expressions: {}, + body: [ + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + mesh_info: None, + task_payload: None, + ray_incoming_payload: Some(2), + ), + ( + name: "any_hit_main", + stage: AnyHit, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("any_hit_main"), + arguments: [], + result: None, + local_variables: [], + expressions: [ + GlobalVariable(2), + AccessIndex( + base: 0, + index: 0, + ), + Literal(U32(1)), + Load( + pointer: 1, + ), + Binary( + op: Add, + left: 3, + right: 2, + ), + GlobalVariable(2), + AccessIndex( + base: 5, + index: 1, + ), + GlobalVariable(2), + AccessIndex( + base: 7, + index: 0, + ), + Load( + pointer: 8, + ), + ], + named_expressions: {}, + body: [ + Emit(( + start: 1, + end: 2, + )), + Emit(( + start: 3, + end: 5, + )), + Store( + pointer: 1, + value: 4, + ), + Emit(( + start: 6, + end: 7, + )), + Emit(( + start: 8, + end: 10, + )), + Store( + pointer: 6, + value: 9, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + mesh_info: None, + task_payload: None, + ray_incoming_payload: Some(2), + ), + ( + name: "closest_hit_main", + stage: ClosestHit, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("closest_hit_main"), + arguments: [], + result: None, + local_variables: [], + expressions: [], + named_expressions: {}, + body: [ + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + mesh_info: None, + task_payload: None, + ray_incoming_payload: Some(2), + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/ir/wgsl-ray-tracing-pipeline.ron b/naga/tests/out/ir/wgsl-ray-tracing-pipeline.ron new file mode 100644 index 00000000000..61148d02682 --- /dev/null +++ b/naga/tests/out/ir/wgsl-ray-tracing-pipeline.ron @@ -0,0 +1,350 @@ +( + types: [ + ( + name: None, + inner: Scalar(( + kind: Uint, + width: 4, + )), + ), + ( + name: Some("HitCounters"), + inner: Struct( + members: [ + ( + name: Some("hit_num"), + ty: 0, + binding: None, + offset: 0, + ), + ( + name: Some("selected_hit"), + ty: 0, + binding: None, + offset: 4, + ), + ], + span: 8, + ), + ), + ( + name: None, + inner: AccelerationStructure( + vertex_return: false, + ), + ), + ( + name: None, + inner: Scalar(( + kind: Float, + width: 4, + )), + ), + ( + name: None, + inner: Vector( + size: Tri, + scalar: ( + kind: Float, + width: 4, + ), + ), + ), + ( + name: Some("RayDesc"), + inner: Struct( + members: [ + ( + name: Some("flags"), + ty: 0, + binding: None, + offset: 0, + ), + ( + name: Some("cull_mask"), + ty: 0, + binding: None, + offset: 4, + ), + ( + name: Some("tmin"), + ty: 3, + binding: None, + offset: 8, + ), + ( + name: Some("tmax"), + ty: 3, + binding: None, + offset: 12, + ), + ( + name: Some("origin"), + ty: 4, + binding: None, + offset: 16, + ), + ( + name: Some("dir"), + ty: 4, + binding: None, + offset: 32, + ), + ], + span: 48, + ), + ), + ], + special_types: ( + ray_desc: Some(5), + ray_intersection: None, + ray_vertex_return: None, + external_texture_params: None, + external_texture_transfer_function: None, + predeclared_types: {}, + ), + constants: [], + overrides: [], + global_variables: [ + ( + name: Some("hit_num"), + space: RayPayload, + binding: None, + ty: 1, + init: None, + ), + ( + name: Some("acc_struct"), + space: Handle, + binding: Some(( + group: 0, + binding: 0, + )), + ty: 2, + init: None, + ), + ( + name: Some("incoming_hit_num"), + space: IncomingRayPayload, + binding: None, + ty: 1, + init: None, + ), + ], + global_expressions: [], + functions: [], + entry_points: [ + ( + name: "ray_gen_main", + stage: RayGeneration, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("ray_gen_main"), + arguments: [], + result: None, + local_variables: [], + expressions: [ + GlobalVariable(0), + ZeroValue(1), + GlobalVariable(1), + Literal(U32(0)), + Literal(U32(255)), + Literal(F32(0.01)), + Literal(F32(100.0)), + Literal(F32(0.0)), + Splat( + size: Tri, + value: 7, + ), + Literal(F32(0.0)), + Literal(F32(1.0)), + Literal(F32(0.0)), + Compose( + ty: 4, + components: [ + 9, + 10, + 11, + ], + ), + Compose( + ty: 5, + components: [ + 3, + 4, + 5, + 6, + 8, + 12, + ], + ), + GlobalVariable(0), + ], + named_expressions: {}, + body: [ + Store( + pointer: 0, + value: 1, + ), + Emit(( + start: 0, + end: 0, + )), + Emit(( + start: 0, + end: 0, + )), + Emit(( + start: 8, + end: 9, + )), + Emit(( + start: 12, + end: 14, + )), + RayPipelineFunction(TraceRay( + acceleration_structure: 2, + descriptor: 13, + payload: 14, + )), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + mesh_info: None, + task_payload: None, + ray_incoming_payload: None, + ), + ( + name: "miss", + stage: Miss, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("miss"), + arguments: [], + result: None, + local_variables: [], + expressions: [], + named_expressions: {}, + body: [ + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + mesh_info: None, + task_payload: None, + ray_incoming_payload: Some(2), + ), + ( + name: "any_hit_main", + stage: AnyHit, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("any_hit_main"), + arguments: [], + result: None, + local_variables: [], + expressions: [ + GlobalVariable(2), + AccessIndex( + base: 0, + index: 0, + ), + Literal(U32(1)), + Load( + pointer: 1, + ), + Binary( + op: Add, + left: 3, + right: 2, + ), + GlobalVariable(2), + AccessIndex( + base: 5, + index: 1, + ), + GlobalVariable(2), + AccessIndex( + base: 7, + index: 0, + ), + Load( + pointer: 8, + ), + ], + named_expressions: {}, + body: [ + Emit(( + start: 1, + end: 2, + )), + Emit(( + start: 3, + end: 5, + )), + Store( + pointer: 1, + value: 4, + ), + Emit(( + start: 6, + end: 7, + )), + Emit(( + start: 8, + end: 10, + )), + Store( + pointer: 6, + value: 9, + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + mesh_info: None, + task_payload: None, + ray_incoming_payload: Some(2), + ), + ( + name: "closest_hit_main", + stage: ClosestHit, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("closest_hit_main"), + arguments: [], + result: None, + local_variables: [], + expressions: [], + named_expressions: {}, + body: [ + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + mesh_info: None, + task_payload: None, + ray_incoming_payload: Some(2), + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/ir/wgsl-storage-textures.compact.ron b/naga/tests/out/ir/wgsl-storage-textures.compact.ron index 68c867a19e2..0bf1bf3ee2a 100644 --- a/naga/tests/out/ir/wgsl-storage-textures.compact.ron +++ b/naga/tests/out/ir/wgsl-storage-textures.compact.ron @@ -220,6 +220,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ( name: "csStore", @@ -319,6 +320,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-storage-textures.ron b/naga/tests/out/ir/wgsl-storage-textures.ron index 68c867a19e2..0bf1bf3ee2a 100644 --- a/naga/tests/out/ir/wgsl-storage-textures.ron +++ b/naga/tests/out/ir/wgsl-storage-textures.ron @@ -220,6 +220,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ( name: "csStore", @@ -319,6 +320,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-template-list-trailing-comma.compact.ron b/naga/tests/out/ir/wgsl-template-list-trailing-comma.compact.ron index db619dff836..8d5cacc2889 100644 --- a/naga/tests/out/ir/wgsl-template-list-trailing-comma.compact.ron +++ b/naga/tests/out/ir/wgsl-template-list-trailing-comma.compact.ron @@ -192,6 +192,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-template-list-trailing-comma.ron b/naga/tests/out/ir/wgsl-template-list-trailing-comma.ron index db619dff836..8d5cacc2889 100644 --- a/naga/tests/out/ir/wgsl-template-list-trailing-comma.ron +++ b/naga/tests/out/ir/wgsl-template-list-trailing-comma.ron @@ -192,6 +192,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-texture-external.compact.ron b/naga/tests/out/ir/wgsl-texture-external.compact.ron index 689fe215e36..1fd051f2518 100644 --- a/naga/tests/out/ir/wgsl-texture-external.compact.ron +++ b/naga/tests/out/ir/wgsl-texture-external.compact.ron @@ -419,6 +419,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ( name: "vertex_main", @@ -457,6 +458,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ( name: "compute_main", @@ -490,6 +492,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-texture-external.ron b/naga/tests/out/ir/wgsl-texture-external.ron index 689fe215e36..1fd051f2518 100644 --- a/naga/tests/out/ir/wgsl-texture-external.ron +++ b/naga/tests/out/ir/wgsl-texture-external.ron @@ -419,6 +419,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ( name: "vertex_main", @@ -457,6 +458,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ( name: "compute_main", @@ -490,6 +492,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-types_with_comments.compact.ron b/naga/tests/out/ir/wgsl-types_with_comments.compact.ron index 7c0d856946f..6aa657e0d86 100644 --- a/naga/tests/out/ir/wgsl-types_with_comments.compact.ron +++ b/naga/tests/out/ir/wgsl-types_with_comments.compact.ron @@ -118,6 +118,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], diff --git a/naga/tests/out/ir/wgsl-types_with_comments.ron b/naga/tests/out/ir/wgsl-types_with_comments.ron index 34e44cb9653..2ea01ca40ec 100644 --- a/naga/tests/out/ir/wgsl-types_with_comments.ron +++ b/naga/tests/out/ir/wgsl-types_with_comments.ron @@ -174,6 +174,7 @@ ), mesh_info: None, task_payload: None, + ray_incoming_payload: None, ), ], diagnostic_filters: [], From 780f1a9b31fa6b43484d4d327830124618ee45ee Mon Sep 17 00:00:00 2001 From: Vecvec Date: Mon, 24 Nov 2025 15:19:20 +1300 Subject: [PATCH 04/15] add to rt spec --- docs/api-specs/ray_tracing.md | 40 +++++++++++++++++++++++++++++++++-- 1 file changed, 38 insertions(+), 2 deletions(-) diff --git a/docs/api-specs/ray_tracing.md b/docs/api-specs/ray_tracing.md index a4038547d46..c2ece2a6b75 100644 --- a/docs/api-specs/ray_tracing.md +++ b/docs/api-specs/ray_tracing.md @@ -93,9 +93,11 @@ fn render(/*whatever args you need to render*/) { ## `naga`'s raytracing API: -`naga` supports ray queries (also known as inline raytracing) only. Ray tracing pipelines are currently unsupported. +`naga` supports ray queries (also known as inline raytracing). Ray tracing pipelines are currently in development. Naming is mostly taken from vulkan. +### Ray Queries + ```wgsl // - Initializes the `ray_query` to check where (if anywhere) the ray defined by `ray_desc` hits in `acceleration_structure rayQueryInitialize(rq: ptr, acceleration_structure: acceleration_structure, ray_desc: RayDesc) @@ -145,7 +147,7 @@ getCandidateHitVertexPositions(rq: ptr>) -> a > [!CAUTION] > -> ### ⚠️Undefined behavior ⚠️: +> #### ⚠️Undefined behavior ⚠️: > - Calling `rayQueryGetCommittedIntersection` or `rayQueryGetCandidateIntersection` when `rayQueryProceed` has not been > called on this ray query since it was initialized (or if the ray query has not been previously initialized). > - Calling `rayQueryGetCommittedIntersection` when `rayQueryProceed`'s latest return on this ray query is considered @@ -268,3 +270,37 @@ const RAY_QUERY_INTERSECTION_GENERATED = 2; // if the ray intersects the bounding box for a custom object. const RAY_QUERY_INTERSECTION_AABB = 3; ``` + +### Ray Tracing Pipelines + +Functions +```wgsl +// Begins to check where (if anywhere) the ray defined by `ray_desc` hits in `acceleration_structure` calling through the `any_hit` shaders and `closest_hit` shader if something was hit or the `miss` shader if no hit was found +traceRay(acceleration_structure: acceleration_structure, ray_desc: RayDesc, payload: ptr) +``` + +> [!CAUTION] +> +> #### ⚠️Undefined behavior ⚠️: +> Calling `traceRay` inside another `traceRay` more than `max_recursion_depth` times +> +> *this is only known undefined behaviour, and will be worked around in the future. + +New shader stages +```wgsl +// First stage to be called, allowed to call `traceRay` +@ray_generation +fn rg() {} + +// Stage called on any hit that is not opaque, not allowed to call `traceRay` +@any_hit +fn ah() {} + +// Stage called on the closest hit, allowed to call `traceRay` +@closest_hit +fn ch() {} + +// Stage call if there was never a hit, allowed to call `traceRay` +@miss +fn miss() {} +``` From d8c62a6fc7cab4539df26c024ebf922a8faaaed9 Mon Sep 17 00:00:00 2001 From: Vecvec Date: Tue, 25 Nov 2025 18:08:26 +1300 Subject: [PATCH 05/15] Attempt to fix ci --- naga/src/back/pipeline_constants.rs | 2 +- naga/src/compact/statements.rs | 2 +- naga/src/valid/analyzer.rs | 8 ++--- naga/src/valid/function.rs | 6 ++-- naga/src/valid/interface.rs | 45 +++++++++++++++++++---------- wgpu-core/src/validation.rs | 7 ++++- wgpu-hal/src/gles/device.rs | 14 +++++++-- wgpu-hal/src/metal/command.rs | 32 ++++++++++++++++++++ wgpu-hal/src/metal/mod.rs | 4 +++ 9 files changed, 91 insertions(+), 29 deletions(-) diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs index 8bdc1349d56..8e596b6cc3f 100644 --- a/naga/src/back/pipeline_constants.rs +++ b/naga/src/back/pipeline_constants.rs @@ -860,7 +860,7 @@ fn adjust_stmt(new_pos: &HandleVec>, stmt: &mut S crate::RayQueryFunction::Terminate => {} } } - Statement::RayPipelineFunction(ref mut func) => match func { + Statement::RayPipelineFunction(ref mut func) => match *func { crate::RayPipelineFunction::TraceRay { ref mut acceleration_structure, ref mut descriptor, diff --git a/naga/src/compact/statements.rs b/naga/src/compact/statements.rs index 18e477c97ab..b490616fce0 100644 --- a/naga/src/compact/statements.rs +++ b/naga/src/compact/statements.rs @@ -382,7 +382,7 @@ impl FunctionMap { adjust(argument); adjust(result); } - St::RayPipelineFunction(ref mut func) => match func { + St::RayPipelineFunction(ref mut func) => match *func { crate::RayPipelineFunction::TraceRay { ref mut acceleration_structure, ref mut descriptor, diff --git a/naga/src/valid/analyzer.rs b/naga/src/valid/analyzer.rs index d5f416092f5..bfc70c5203c 100644 --- a/naga/src/valid/analyzer.rs +++ b/naga/src/valid/analyzer.rs @@ -1169,15 +1169,15 @@ impl FunctionInfo { FunctionUniformity::new() } S::RayPipelineFunction(ref fun) => { - match fun { + match *fun { crate::RayPipelineFunction::TraceRay { acceleration_structure, descriptor, payload, } => { - let _ = self.add_ref(*acceleration_structure); - let _ = self.add_ref(*descriptor); - let _ = self.add_ref(*payload); + let _ = self.add_ref(acceleration_structure); + let _ = self.add_ref(descriptor); + let _ = self.add_ref(payload); } } FunctionUniformity::new() diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index 6f00849291b..ef9e306b016 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -1645,10 +1645,8 @@ impl super::Validator { crate::TypeInner::AccelerationStructure { vertex_return } => { if !vertex_return { self.trace_rays_no_vertex_return = Some(Some(span)); - } else { - if let None = self.trace_rays_no_vertex_return { - self.trace_rays_no_vertex_return = Some(None); - } + } else if self.trace_rays_no_vertex_return.is_none() { + self.trace_rays_no_vertex_return = Some(None); } } _ => { diff --git a/naga/src/valid/interface.rs b/naga/src/valid/interface.rs index a975ebf5bc4..8e53dded306 100644 --- a/naga/src/valid/interface.rs +++ b/naga/src/valid/interface.rs @@ -735,7 +735,12 @@ impl VaryingContext<'_> { let needs_interpolation = match self.stage { crate::ShaderStage::Vertex => self.output, crate::ShaderStage::Fragment => !self.output && !per_primitive, - crate::ShaderStage::Compute | crate::ShaderStage::Task | crate::ShaderStage::RayGeneration | crate::ShaderStage::AnyHit | crate::ShaderStage::ClosestHit | crate::ShaderStage::Miss => false, + crate::ShaderStage::Compute + | crate::ShaderStage::Task + | crate::ShaderStage::RayGeneration + | crate::ShaderStage::AnyHit + | crate::ShaderStage::ClosestHit + | crate::ShaderStage::Miss => false, crate::ShaderStage::Mesh => self.output, }; @@ -962,15 +967,15 @@ impl super::Validator { ) } crate::AddressSpace::RayPayload | crate::AddressSpace::IncomingRayPayload => { - if !self.capabilities.contains(Capabilities::RAY_TRACING_PIPELINE) { + if !self + .capabilities + .contains(Capabilities::RAY_TRACING_PIPELINE) + { return Err(GlobalVariableError::UnsupportedCapability( Capabilities::PUSH_CONSTANT, )); } - ( - TypeFlags::DATA | TypeFlags::SIZED, - false, - ) + (TypeFlags::DATA | TypeFlags::SIZED, false) } }; @@ -1315,20 +1320,28 @@ impl super::Validator { } crate::AddressSpace::PushConstant => GlobalUse::READ, crate::AddressSpace::RayPayload => { - if !matches!(ep.stage, crate::ShaderStage::RayGeneration | crate::ShaderStage::ClosestHit | crate::ShaderStage::Miss) { - return Err(EntryPointError::RayPayloadInInvalidStage.with_span_handle(var_handle, &module.global_variables)); + if !matches!( + ep.stage, + crate::ShaderStage::RayGeneration + | crate::ShaderStage::ClosestHit + | crate::ShaderStage::Miss + ) { + return Err(EntryPointError::RayPayloadInInvalidStage + .with_span_handle(var_handle, &module.global_variables)); } - GlobalUse::READ - | GlobalUse::QUERY - | GlobalUse::WRITE + GlobalUse::READ | GlobalUse::QUERY | GlobalUse::WRITE } crate::AddressSpace::IncomingRayPayload => { - if !matches!(ep.stage, crate::ShaderStage::AnyHit | crate::ShaderStage::ClosestHit | crate::ShaderStage::Miss) { - return Err(EntryPointError::IncomingRayPayloadInInvalidStage.with_span_handle(var_handle, &module.global_variables)); + if !matches!( + ep.stage, + crate::ShaderStage::AnyHit + | crate::ShaderStage::ClosestHit + | crate::ShaderStage::Miss + ) { + return Err(EntryPointError::IncomingRayPayloadInInvalidStage + .with_span_handle(var_handle, &module.global_variables)); } - GlobalUse::READ - | GlobalUse::QUERY - | GlobalUse::WRITE + GlobalUse::READ | GlobalUse::QUERY | GlobalUse::WRITE } }; if !allowed_usage.contains(usage) { diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs index 5c499af749d..bc16dc18213 100644 --- a/wgpu-core/src/validation.rs +++ b/wgpu-core/src/validation.rs @@ -1301,7 +1301,12 @@ impl Interface { } naga::ShaderStage::Compute => (false, 0), // TODO: add validation for these, see https://github.com/gfx-rs/wgpu/issues/8003 - naga::ShaderStage::Task | naga::ShaderStage::Mesh | naga::ShaderStage::RayGeneration | naga::ShaderStage::AnyHit | naga::ShaderStage::ClosestHit | naga::ShaderStage::Miss => { + naga::ShaderStage::Task + | naga::ShaderStage::Mesh + | naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => { unreachable!() } }; diff --git a/wgpu-hal/src/gles/device.rs b/wgpu-hal/src/gles/device.rs index dbdd8d3f0bc..f024542d255 100644 --- a/wgpu-hal/src/gles/device.rs +++ b/wgpu-hal/src/gles/device.rs @@ -99,7 +99,12 @@ impl CompilationContext<'_> { unsafe { gl.bind_frag_data_location(program, location.location, &name) } } naga::ShaderStage::Compute => {} - naga::ShaderStage::Task | naga::ShaderStage::Mesh | naga::ShaderStage::RayGeneration | naga::ShaderStage::AnyHit | naga::ShaderStage::ClosestHit | naga::ShaderStage::Miss => unreachable!(), + naga::ShaderStage::Task + | naga::ShaderStage::Mesh + | naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unreachable!(), } } @@ -175,7 +180,12 @@ impl super::Device { naga::ShaderStage::Vertex => glow::VERTEX_SHADER, naga::ShaderStage::Fragment => glow::FRAGMENT_SHADER, naga::ShaderStage::Compute => glow::COMPUTE_SHADER, - naga::ShaderStage::Task | naga::ShaderStage::Mesh | naga::ShaderStage::RayGeneration | naga::ShaderStage::AnyHit | naga::ShaderStage::ClosestHit | naga::ShaderStage::Miss => unreachable!(), + naga::ShaderStage::Task + | naga::ShaderStage::Mesh + | naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unreachable!(), }; let raw = unsafe { gl.create_shader(target) }.unwrap(); diff --git a/wgpu-hal/src/metal/command.rs b/wgpu-hal/src/metal/command.rs index 86be90427d7..a1d4b3f5516 100644 --- a/wgpu-hal/src/metal/command.rs +++ b/wgpu-hal/src/metal/command.rs @@ -164,6 +164,10 @@ impl super::CommandEncoder { naga::ShaderStage::Task => &bg_info.base_resource_indices.ts, naga::ShaderStage::Mesh => &bg_info.base_resource_indices.ms, naga::ShaderStage::Compute => &bg_info.base_resource_indices.cs, + naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unimplemented!(), }; let buffers = match stage { naga::ShaderStage::Vertex => group.counters.vs.buffers, @@ -171,6 +175,10 @@ impl super::CommandEncoder { naga::ShaderStage::Task => group.counters.ts.buffers, naga::ShaderStage::Mesh => group.counters.ms.buffers, naga::ShaderStage::Compute => group.counters.cs.buffers, + naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unimplemented!(), }; let mut changes_sizes_buffer = false; for index in 0..buffers { @@ -190,6 +198,10 @@ impl super::CommandEncoder { naga::ShaderStage::Task => render_encoder.unwrap().set_object_buffer(a1, a2, a3), naga::ShaderStage::Mesh => render_encoder.unwrap().set_mesh_buffer(a1, a2, a3), naga::ShaderStage::Compute => compute_encoder.unwrap().set_buffer(a1, a2, a3), + naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unimplemented!(), } if let Some(size) = buf.binding_size { let br = naga::ResourceBinding { @@ -218,6 +230,10 @@ impl super::CommandEncoder { naga::ShaderStage::Task => render_encoder.unwrap().set_object_bytes(a1, a2, a3), naga::ShaderStage::Mesh => render_encoder.unwrap().set_mesh_bytes(a1, a2, a3), naga::ShaderStage::Compute => compute_encoder.unwrap().set_bytes(a1, a2, a3), + naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unimplemented!(), } } } @@ -227,6 +243,10 @@ impl super::CommandEncoder { naga::ShaderStage::Task => group.counters.ts.samplers, naga::ShaderStage::Mesh => group.counters.ms.samplers, naga::ShaderStage::Compute => group.counters.cs.samplers, + naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unimplemented!(), }; for index in 0..samplers { let res = group.samplers[(index_base.samplers + index) as usize]; @@ -242,6 +262,10 @@ impl super::CommandEncoder { naga::ShaderStage::Task => render_encoder.unwrap().set_object_sampler_state(a1, a2), naga::ShaderStage::Mesh => render_encoder.unwrap().set_mesh_sampler_state(a1, a2), naga::ShaderStage::Compute => compute_encoder.unwrap().set_sampler_state(a1, a2), + naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unimplemented!(), } } @@ -251,6 +275,10 @@ impl super::CommandEncoder { naga::ShaderStage::Task => group.counters.ts.textures, naga::ShaderStage::Mesh => group.counters.ms.textures, naga::ShaderStage::Compute => group.counters.cs.textures, + naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unimplemented!(), }; for index in 0..textures { let res = group.textures[(index_base.textures + index) as usize]; @@ -262,6 +290,10 @@ impl super::CommandEncoder { naga::ShaderStage::Task => render_encoder.unwrap().set_object_texture(a1, a2), naga::ShaderStage::Mesh => render_encoder.unwrap().set_mesh_texture(a1, a2), naga::ShaderStage::Compute => compute_encoder.unwrap().set_texture(a1, a2), + naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unimplemented!(), } } } diff --git a/wgpu-hal/src/metal/mod.rs b/wgpu-hal/src/metal/mod.rs index 7258a885f25..7c4d5d0a34c 100644 --- a/wgpu-hal/src/metal/mod.rs +++ b/wgpu-hal/src/metal/mod.rs @@ -631,6 +631,10 @@ impl ops::Index for MultiStageData { naga::ShaderStage::Compute => &self.cs, naga::ShaderStage::Task => &self.ts, naga::ShaderStage::Mesh => &self.ms, + naga::ShaderStage::RayGeneration + | naga::ShaderStage::AnyHit + | naga::ShaderStage::ClosestHit + | naga::ShaderStage::Miss => unimplemented!(), } } } From 50e99369b3a4137434f102efbbbcd4cdc4f5e5d4 Mon Sep 17 00:00:00 2001 From: Vecvec Date: Tue, 25 Nov 2025 20:26:17 +1300 Subject: [PATCH 06/15] satisfy taplo --- naga/tests/in/wgsl/ray-tracing-pipeline.toml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/naga/tests/in/wgsl/ray-tracing-pipeline.toml b/naga/tests/in/wgsl/ray-tracing-pipeline.toml index f4611d2d304..9845f3bb7e9 100644 --- a/naga/tests/in/wgsl/ray-tracing-pipeline.toml +++ b/naga/tests/in/wgsl/ray-tracing-pipeline.toml @@ -1,2 +1,2 @@ god_mode = true -targets = "IR" \ No newline at end of file +targets = "IR" From 263c0f217faf720ab6e56b0af964dddd8487c518 Mon Sep 17 00:00:00 2001 From: Vecvec Date: Tue, 25 Nov 2025 20:32:14 +1300 Subject: [PATCH 07/15] Fix typo --- naga/src/valid/function.rs | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index ef9e306b016..a246a55be1c 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -227,8 +227,8 @@ pub enum FunctionError { InvalidMeshShaderOutputType(Handle), #[error("The payload type passed to `traceRay` must be a pointer")] InvalidPayloadType, - #[error("The payload type passed to `traceRay` must be a pointer with an adress space of `ray_payload` or `incoming_ray_payload`, instead got {0:?}")] - InvalidPayloadAdressSpace(crate::AddressSpace), + #[error("The payload type passed to `traceRay` must be a pointer with an address space of `ray_payload` or `incoming_ray_payload`, instead got {0:?}")] + InvalidPayloadAddressSpace(crate::AddressSpace), #[error("The payload type ({0:?}) passed to `traceRay` does not match the previous one {1:?}")] MismatchedPayloadType(Handle, Handle), } @@ -1665,7 +1665,7 @@ impl super::Validator { AddressSpace::RayPayload | AddressSpace::IncomingRayPayload => { } space => { - return Err(FunctionError::InvalidPayloadAdressSpace(space) + return Err(FunctionError::InvalidPayloadAddressSpace(space) .with_span_handle(payload, context.expressions)) } } From 9139d153e0494fbeb28e50e9ad46c6fef40222f8 Mon Sep 17 00:00:00 2001 From: Vecvec Date: Tue, 25 Nov 2025 20:34:01 +1300 Subject: [PATCH 08/15] 2nd typo --- naga/src/ir/mod.rs | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/naga/src/ir/mod.rs b/naga/src/ir/mod.rs index 01c6057fb89..fd48e0891f6 100644 --- a/naga/src/ir/mod.rs +++ b/naga/src/ir/mod.rs @@ -2678,7 +2678,7 @@ pub struct MeshStageInfo { pub output_variable: Handle, } -/// Ray tracing pipeine intrinsics +/// Ray tracing pipeline intrinsics #[derive(Debug, Clone, Copy)] #[cfg_attr(feature = "serialize", derive(Serialize))] #[cfg_attr(feature = "deserialize", derive(Deserialize))] From e6b000a918fbfe7a873c064c9f4dde089305b47f Mon Sep 17 00:00:00 2001 From: Vecvec Date: Tue, 25 Nov 2025 20:48:48 +1300 Subject: [PATCH 09/15] Format from typo fixes --- naga/src/valid/function.rs | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index a246a55be1c..76ae7c286f7 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -1665,8 +1665,10 @@ impl super::Validator { AddressSpace::RayPayload | AddressSpace::IncomingRayPayload => { } space => { - return Err(FunctionError::InvalidPayloadAddressSpace(space) - .with_span_handle(payload, context.expressions)) + return Err(FunctionError::InvalidPayloadAddressSpace( + space, + ) + .with_span_handle(payload, context.expressions)) } } base From bee83af5860522b6a9bc121eca353ff26a64aa48 Mon Sep 17 00:00:00 2001 From: Vecvec Date: Thu, 27 Nov 2025 09:43:37 +1300 Subject: [PATCH 10/15] Add tests (plus some spots that had incorrect errors) --- naga/src/valid/interface.rs | 30 ++++-- naga/src/valid/mod.rs | 1 + naga/src/valid/type.rs | 5 +- naga/tests/naga/wgsl_errors.rs | 163 +++++++++++++++++++++++++++++++++ 4 files changed, 190 insertions(+), 9 deletions(-) diff --git a/naga/src/valid/interface.rs b/naga/src/valid/interface.rs index 8e53dded306..64ead99dcba 100644 --- a/naga/src/valid/interface.rs +++ b/naga/src/valid/interface.rs @@ -972,7 +972,7 @@ impl super::Validator { .contains(Capabilities::RAY_TRACING_PIPELINE) { return Err(GlobalVariableError::UnsupportedCapability( - Capabilities::PUSH_CONSTANT, + Capabilities::RAY_TRACING_PIPELINE, )); } (TypeFlags::DATA | TypeFlags::SIZED, false) @@ -1082,14 +1082,28 @@ impl super::Validator { module: &crate::Module, mod_info: &ModuleInfo, ) -> Result> { - if matches!( - ep.stage, + match ep.stage { crate::ShaderStage::Task | crate::ShaderStage::Mesh - ) && !self.capabilities.contains(Capabilities::MESH_SHADER) - { - return Err( - EntryPointError::UnsupportedCapability(Capabilities::MESH_SHADER).with_span(), - ); + if !self.capabilities.contains(Capabilities::MESH_SHADER) => + { + return Err( + EntryPointError::UnsupportedCapability(Capabilities::MESH_SHADER).with_span(), + ); + } + crate::ShaderStage::RayGeneration + | crate::ShaderStage::AnyHit + | crate::ShaderStage::ClosestHit + | crate::ShaderStage::Miss + if !self + .capabilities + .contains(Capabilities::RAY_TRACING_PIPELINE) => + { + return Err(EntryPointError::UnsupportedCapability( + Capabilities::RAY_TRACING_PIPELINE, + ) + .with_span()); + } + _ => {} } if ep.early_depth_test.is_some() { let required = Capabilities::EARLY_DEPTH_TEST; diff --git a/naga/src/valid/mod.rs b/naga/src/valid/mod.rs index 49fe6617871..fe362e0090a 100644 --- a/naga/src/valid/mod.rs +++ b/naga/src/valid/mod.rs @@ -212,6 +212,7 @@ impl Capabilities { Self::CLIP_DISTANCE => Some(Ext::ClipDistances), Self::RAY_QUERY => Some(Ext::WgpuRayQuery), Self::RAY_HIT_VERTEX_POSITION => Some(Ext::WgpuRayQueryVertexReturn), + Self::RAY_TRACING_PIPELINE => Some(Ext::WgpuRayTracingPipeline), _ => None, } } diff --git a/naga/src/valid/type.rs b/naga/src/valid/type.rs index 55885eff1eb..f0ebe1ebe17 100644 --- a/naga/src/valid/type.rs +++ b/naga/src/valid/type.rs @@ -766,7 +766,10 @@ impl super::Validator { Alignment::ONE, ), Ti::AccelerationStructure { vertex_return } => { - self.require_type_capability(Capabilities::RAY_QUERY)?; + self.require_type_capability(Capabilities::RAY_QUERY) + .or_else(|_| { + self.require_type_capability(Capabilities::RAY_TRACING_PIPELINE) + })?; if vertex_return { self.require_type_capability(Capabilities::RAY_HIT_VERTEX_POSITION)?; } diff --git a/naga/tests/naga/wgsl_errors.rs b/naga/tests/naga/wgsl_errors.rs index c9c41328408..2419b3d7c67 100644 --- a/naga/tests/naga/wgsl_errors.rs +++ b/naga/tests/naga/wgsl_errors.rs @@ -4379,3 +4379,166 @@ fn ray_query_vertex_return_enable_extension() { }) ); } + +/// Checks that every ray tracing pipeline binding in naga is invalid in other stages. +#[test] +fn check_ray_tracing_pipeline_bindings() { + for (builtin, ty) in [ + ("ray_invocation_id", "vec3"), + ("num_ray_invocations", "vec3"), + ("instance_custom_data", "u32"), + ("geometry_index", "u32"), + ("world_ray_origin", "vec3"), + ("world_ray_direction", "vec3"), + ("object_ray_origin", "vec3"), + ("object_ray_direction", "vec3"), + ("ray_t_min", "f32"), + ("ray_t_current_max", "f32"), + ("object_to_world", "mat4x3"), + ("world_to_object", "mat4x3"), + ("hit_kind", "u32"), + ] { + check_one_validation!( + &format!( + "@compute + @workgroup_size(1) + fn main(@builtin({builtin}) v: {ty}) {{}} + " + ), + Err(naga::valid::ValidationError::EntryPoint { + source: naga::valid::EntryPointError::Argument( + 0, + naga::valid::VaryingError::InvalidBuiltInStage(_), + ), + .. + },) + ); + check_one_validation!( + &format!( + "@vertex + fn main(@builtin({builtin}) v: {ty}) {{}} + " + ), + Err(naga::valid::ValidationError::EntryPoint { + source: naga::valid::EntryPointError::Argument( + 0, + naga::valid::VaryingError::InvalidBuiltInStage(_), + ), + .. + },) + ); + check_one_validation!( + &format!( + "@fragment + fn main(@builtin({builtin}) v: {ty}) {{}} + " + ), + Err(naga::valid::ValidationError::EntryPoint { + source: naga::valid::EntryPointError::Argument( + 0, + naga::valid::VaryingError::InvalidBuiltInStage(_), + ), + .. + },) + ); + } +} + +/// Checks ray generation stage is invalid without enable extension (other stages require `@incoming_payload` which forces a ray payload which is checked in [`check_ray_tracing_pipeline_payload`]) +#[test] +fn check_ray_tracing_pipeline_ray_generation() { + check_extension_validation!( + Capabilities::RAY_TRACING_PIPELINE, + "@ray_generation + fn main() {{}}", + "error: the `wgpu_ray_tracing_pipeline` enable extension is not enabled + ┌─ wgsl:1:2 + │ +1 │ @ray_generation + │ ^^^^^^^^^^^^^^ the `wgpu_ray_tracing_pipeline` \"Enable Extension\" is needed for this functionality, but it is not currently enabled. + │ + = note: You can enable this extension by adding `enable wgpu_ray_tracing_pipeline;` at the top of the shader, before any other items. + +", + Err(naga::valid::ValidationError::EntryPoint { + source: naga::valid::EntryPointError::UnsupportedCapability(naga::valid::Capabilities::RAY_TRACING_PIPELINE), + .. + },) + ); +} + +#[test] +fn check_ray_tracing_pipeline_payload() { + for space in ["ray_payload", "incoming_ray_payload"] { + // ascii is a byte per char so length is fine + let space_arrows = "^".to_string().repeat(space.len()); + check_extension_validation!( + Capabilities::RAY_TRACING_PIPELINE, + &format!("var<{space}> payload: u32;"), + &format!("error: the `wgpu_ray_tracing_pipeline` enable extension is not enabled + ┌─ wgsl:1:5 + │ +1 │ var<{space}> payload: u32; + │ {space_arrows} the `wgpu_ray_tracing_pipeline` \"Enable Extension\" is needed for this functionality, but it is not currently enabled. + │ + = note: You can enable this extension by adding `enable wgpu_ray_tracing_pipeline;` at the top of the shader, before any other items. + +"), + Err(naga::valid::ValidationError::GlobalVariable { + source: naga::valid::GlobalVariableError::UnsupportedCapability(naga::valid::Capabilities::RAY_TRACING_PIPELINE), + .. + },) + ); + } +} + +#[test] +fn check_ray_tracing_pipeline_incoming_payload_required() { + for stage in ["any_hit", "closest_hit", "miss"] { + // ascii is a byte per char so length is fine + let stage_arrows = "^".to_string().repeat(stage.len()); + check( + &format!("enable wgpu_ray_tracing_pipeline; @{stage} fn main() {{}}"), + &format!("error: incoming payload is missing on ray hit or miss shader entry point + ┌─ wgsl:1:36 + │ +1 │ enable wgpu_ray_tracing_pipeline; @{stage} fn main() {{}} + │ {stage_arrows} must be paired with a `@incoming_payload` attribute + +"), + ); + } +} + +#[test] +fn check_ray_tracing_pipeline_payload_disallowed() { + for (stage, output, stmt) in [ + ( + "var incoming: u32; @any_hit @incoming_payload(incoming)", + "", + "", + ), + ("@compute @workgroup_size(1)", "", ""), + ( + "@vertex", + " -> @builtin(position) vec4", + "return vec4();", + ), + ("@fragment", "", ""), + ] { + check_one_validation!( + &format!( + "enable wgpu_ray_tracing_pipeline; + @group(0) @binding(0) var acc_struct: acceleration_structure; + var payload: u32; + + {stage} fn main() {output} {{_ = payload; {stmt}}}" + ), + Err(naga::valid::ValidationError::EntryPoint { + source: naga::valid::EntryPointError::RayPayloadInInvalidStage, + .. + },), + Capabilities::RAY_TRACING_PIPELINE + ); + } +} From ee8883c17a393d5e37cde33325aa8e68c27fe6d4 Mon Sep 17 00:00:00 2001 From: Vecvec Date: Thu, 27 Nov 2025 10:29:14 +1300 Subject: [PATCH 11/15] Fix test broken by allowing 2 caps --- naga/src/valid/type.rs | 6 ++---- naga/tests/naga/wgsl_errors.rs | 35 ++++++++++++++++++++++++++++++---- 2 files changed, 33 insertions(+), 8 deletions(-) diff --git a/naga/src/valid/type.rs b/naga/src/valid/type.rs index f0ebe1ebe17..d2bba084df4 100644 --- a/naga/src/valid/type.rs +++ b/naga/src/valid/type.rs @@ -766,10 +766,8 @@ impl super::Validator { Alignment::ONE, ), Ti::AccelerationStructure { vertex_return } => { - self.require_type_capability(Capabilities::RAY_QUERY) - .or_else(|_| { - self.require_type_capability(Capabilities::RAY_TRACING_PIPELINE) - })?; + self.require_type_capability(Capabilities::RAY_TRACING_PIPELINE) + .or_else(|_| self.require_type_capability(Capabilities::RAY_QUERY))?; if vertex_return { self.require_type_capability(Capabilities::RAY_HIT_VERTEX_POSITION)?; } diff --git a/naga/tests/naga/wgsl_errors.rs b/naga/tests/naga/wgsl_errors.rs index 2419b3d7c67..59537cd962b 100644 --- a/naga/tests/naga/wgsl_errors.rs +++ b/naga/tests/naga/wgsl_errors.rs @@ -1029,7 +1029,10 @@ macro_rules! check_one_validation { /// NOTE: The only reason we don't use a function for this is because we need to syntactically /// re-use `$val_err_pat`. macro_rules! check_extension_validation { - ( $caps:expr, $source:expr, $parse_err:expr, $val_err_pat:pat ) => { + ( $caps:expr, $source:expr, $parse_err:expr, $val_err_pat:pat $(, $other_caps:expr)? ) => { + #[allow(unused_mut, unused_assignments)] + let mut other_caps = naga::valid::Capabilities::empty(); + $(other_caps = $other_caps;)? let caps = $caps; let source = $source; let mut ext = None; @@ -1081,7 +1084,9 @@ macro_rules! check_extension_validation { }; // Second check, for the expected validation error when the capability is not present - let error = naga::valid::Validator::new(naga::valid::ValidationFlags::all(), !caps) + // Don't check with explicitly allowed caps, as certain things (currently just + // `acceleration_structure`s) can be enabled by multiple extensions + let error = naga::valid::Validator::new(naga::valid::ValidationFlags::all(), !(caps | other_caps)) .validate(&module) .map_err(|e| e.into_inner()); // TODO(https://github.com/gfx-rs/wgpu/issues/8153): Add tests for spans #[allow(clippy::redundant_pattern_matching)] @@ -4285,7 +4290,7 @@ fn source_with_control_char() { } #[test] -fn ray_query_enable_extension() { +fn ray_types_enable_extension() { check_extension_validation!( Capabilities::RAY_QUERY, r#"fn foo() { @@ -4307,6 +4312,7 @@ fn ray_query_enable_extension() { }) ); + // can be enabled by either of these extensions check_extension_validation!( Capabilities::RAY_QUERY, r#"@group(0) @binding(0) @@ -4324,7 +4330,28 @@ fn ray_query_enable_extension() { Err(naga::valid::ValidationError::Type { source: naga::valid::TypeError::MissingCapability(Capabilities::RAY_QUERY), .. - }) + }), + Capabilities::RAY_TRACING_PIPELINE + ); + check_extension_validation!( + Capabilities::RAY_TRACING_PIPELINE, + r#"@group(0) @binding(0) + var acc_struct: acceleration_structure; + "#, + r#"error: the `wgpu_ray_query` enable extension is not enabled + ┌─ wgsl:2:25 + │ +2 │ var acc_struct: acceleration_structure; + │ ^^^^^^^^^^^^^^^^^^^^^^ the `wgpu_ray_query` "Enable Extension" is needed for this functionality, but it is not currently enabled. + │ + = note: You can enable this extension by adding `enable wgpu_ray_query;` at the top of the shader, before any other items. + +"#, + Err(naga::valid::ValidationError::Type { + source: naga::valid::TypeError::MissingCapability(Capabilities::RAY_QUERY), + .. + }), + Capabilities::RAY_QUERY ); } From d4574d548e70f107338cd7d66c49a88f08585166 Mon Sep 17 00:00:00 2001 From: Vecvec Date: Thu, 27 Nov 2025 13:50:17 +1300 Subject: [PATCH 12/15] Apply suggestions --- naga/src/back/glsl/mod.rs | 9 +++++---- naga/src/back/msl/writer.rs | 5 ++--- naga/src/ir/mod.rs | 1 + naga/src/valid/interface.rs | 8 ++++---- naga/tests/naga/wgsl_errors.rs | 2 +- 5 files changed, 13 insertions(+), 12 deletions(-) diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index b1e95d15bb1..2acac244aaa 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -140,10 +140,11 @@ impl crate::AddressSpace { | crate::AddressSpace::Storage { .. } | crate::AddressSpace::Handle | crate::AddressSpace::PushConstant - | crate::AddressSpace::TaskPayload - // just a default impl, not really supported - | crate::AddressSpace::RayPayload - | crate::AddressSpace::IncomingRayPayload => false, + | crate::AddressSpace::TaskPayload => false, + + crate::AddressSpace::RayPayload | crate::AddressSpace::IncomingRayPayload => { + unreachable!() + } } } } diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index fd6e27c7ace..4bb9dca0220 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -596,10 +596,9 @@ impl crate::AddressSpace { | Self::WorkGroup | Self::PushConstant | Self::Handle - | Self::TaskPayload - | Self::RayPayload - | Self::IncomingRayPayload => true, + | Self::TaskPayload => true, Self::Function => false, + Self::RayPayload | Self::IncomingRayPayload => unreachable!(), } } diff --git a/naga/src/ir/mod.rs b/naga/src/ir/mod.rs index fd48e0891f6..20484ada987 100644 --- a/naga/src/ir/mod.rs +++ b/naga/src/ir/mod.rs @@ -344,6 +344,7 @@ pub enum ShaderStage { /// A any hit shader, in a ray tracing pipeline. AnyHit, + /// A closest hit shader, in a ray tracing pipeline. ClosestHit, } diff --git a/naga/src/valid/interface.rs b/naga/src/valid/interface.rs index 64ead99dcba..ac125e08c8d 100644 --- a/naga/src/valid/interface.rs +++ b/naga/src/valid/interface.rs @@ -165,9 +165,9 @@ pub enum EntryPointError { #[error("Task payload must be at least 4 bytes, but is {0} bytes")] TaskPayloadTooSmall(u32), #[error("Only the `ray_generation`, `closest_hit`, and `any_hit` shader stages can access a global variable in the `ray_payload` address space")] - RayPayloadInInvalidStage, + RayPayloadInInvalidStage(crate::ShaderStage), #[error("Only the `closest_hit`, `any_hit`, and `miss` shader stages can access a global variable in the `incoming_ray_payload` address space")] - IncomingRayPayloadInInvalidStage, + IncomingRayPayloadInInvalidStage(crate::ShaderStage), } fn storage_usage(access: crate::StorageAccess) -> GlobalUse { @@ -1340,7 +1340,7 @@ impl super::Validator { | crate::ShaderStage::ClosestHit | crate::ShaderStage::Miss ) { - return Err(EntryPointError::RayPayloadInInvalidStage + return Err(EntryPointError::RayPayloadInInvalidStage(ep.stage) .with_span_handle(var_handle, &module.global_variables)); } GlobalUse::READ | GlobalUse::QUERY | GlobalUse::WRITE @@ -1352,7 +1352,7 @@ impl super::Validator { | crate::ShaderStage::ClosestHit | crate::ShaderStage::Miss ) { - return Err(EntryPointError::IncomingRayPayloadInInvalidStage + return Err(EntryPointError::IncomingRayPayloadInInvalidStage(ep.stage) .with_span_handle(var_handle, &module.global_variables)); } GlobalUse::READ | GlobalUse::QUERY | GlobalUse::WRITE diff --git a/naga/tests/naga/wgsl_errors.rs b/naga/tests/naga/wgsl_errors.rs index 59537cd962b..b2691a4d28b 100644 --- a/naga/tests/naga/wgsl_errors.rs +++ b/naga/tests/naga/wgsl_errors.rs @@ -4562,7 +4562,7 @@ fn check_ray_tracing_pipeline_payload_disallowed() { {stage} fn main() {output} {{_ = payload; {stmt}}}" ), Err(naga::valid::ValidationError::EntryPoint { - source: naga::valid::EntryPointError::RayPayloadInInvalidStage, + source: naga::valid::EntryPointError::RayPayloadInInvalidStage(_), .. },), Capabilities::RAY_TRACING_PIPELINE From fc50128562776f1ccaf7dbc1517f47a150ef10f5 Mon Sep 17 00:00:00 2001 From: Vecvec Date: Thu, 27 Nov 2025 14:45:23 +1300 Subject: [PATCH 13/15] Allow some previously available builtins --- naga/src/ir/mod.rs | 4 ++-- naga/src/valid/interface.rs | 15 ++++++++++++--- 2 files changed, 14 insertions(+), 5 deletions(-) diff --git a/naga/src/ir/mod.rs b/naga/src/ir/mod.rs index 20484ada987..de0215b57c5 100644 --- a/naga/src/ir/mod.rs +++ b/naga/src/ir/mod.rs @@ -414,7 +414,7 @@ pub enum BuiltIn { ClipDistance, /// Written in vertex & mesh shaders CullDistance, - /// Read in vertex shaders + /// Read in vertex, any- and closest-hit shaders InstanceIndex, /// Written in vertex & mesh shaders PointSize, @@ -429,7 +429,7 @@ pub enum BuiltIn { PointCoord, /// Read in fragment shaders FrontFacing, - /// Read in fragment shaders, written in mesh shaders + /// Read in fragment shaders, written in mesh shaders, read in any and closest hit shaders. PrimitiveIndex, /// Read in fragment shaders Barycentric, diff --git a/naga/src/valid/interface.rs b/naga/src/valid/interface.rs index ac125e08c8d..cf90addab66 100644 --- a/naga/src/valid/interface.rs +++ b/naga/src/valid/interface.rs @@ -233,7 +233,10 @@ impl VaryingContext<'_> { let required = match built_in { Bi::ClipDistance => Capabilities::CLIP_DISTANCE, Bi::CullDistance => Capabilities::CULL_DISTANCE, - Bi::PrimitiveIndex => Capabilities::PRIMITIVE_INDEX, + // Primitive index is allowed w/o any other extensions in any- and closest-hit shaders + Bi::PrimitiveIndex if !matches!(ep.stage, St::AnyHit | St::ClosestHit) => { + Capabilities::PRIMITIVE_INDEX + } Bi::Barycentric => Capabilities::SHADER_BARYCENTRICS, Bi::ViewIndex => Capabilities::MULTIVIEW, Bi::SampleIndex => Capabilities::MULTISAMPLED_SHADING, @@ -256,10 +259,15 @@ impl VaryingContext<'_> { } let (visible, type_good) = match built_in { - Bi::BaseInstance | Bi::BaseVertex | Bi::InstanceIndex | Bi::VertexIndex => ( + Bi::BaseInstance | Bi::BaseVertex | Bi::VertexIndex => ( self.stage == St::Vertex && !self.output, *ty_inner == Ti::Scalar(crate::Scalar::U32), ), + Bi::InstanceIndex => ( + matches!(self.stage, St::Vertex | St::AnyHit | St::ClosestHit) + && !self.output, + *ty_inner == Ti::Scalar(crate::Scalar::U32), + ), Bi::DrawID => ( // Always allowed in task/vertex stage. Allowed in mesh stage if there is no task stage in the pipeline. (self.stage == St::Vertex @@ -326,7 +334,8 @@ impl VaryingContext<'_> { *ty_inner == Ti::Scalar(crate::Scalar::BOOL), ), Bi::PrimitiveIndex => ( - (self.stage == St::Fragment && !self.output) + (matches!(self.stage, St::Fragment | St::AnyHit | St::ClosestHit) + && !self.output) || (self.stage == St::Mesh && self.output && self.mesh_output_type == MeshOutputType::PrimitiveOutput), From ec1c25bc8d541c36e456d0b5cc86d21c25e705c1 Mon Sep 17 00:00:00 2001 From: Vecvec Date: Fri, 28 Nov 2025 15:57:55 +1300 Subject: [PATCH 14/15] Changelog --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index b440074bd95..f7e4bc71716 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -135,6 +135,7 @@ By @cwfitzgerald in [#8579](https://github.com/gfx-rs/wgpu/pull/8579). - Added support for transient textures on Vulkan and Metal. By @opstic in [#8247](https://github.com/gfx-rs/wgpu/pull/8247) - Implement shader triangle barycentric coordinate builtins. By @atlv24 in [#8320](https://github.com/gfx-rs/wgpu/pull/8320). - Added support for binding arrays of storage textures on Metal. By @msvbg in [#8464](https://github.com/gfx-rs/wgpu/pull/8464) +- Initial wgsl-in ray tracing pipelines. By @Vecvec in [#8570](https://github.com/gfx-rs/wgpu/pull/8570). ### Changes From 7ecef079eb8f055e8ee7740cee42545159040a9f Mon Sep 17 00:00:00 2001 From: Vecvec Date: Sun, 30 Nov 2025 17:08:59 +1300 Subject: [PATCH 15/15] Fix merge --- naga/src/valid/mod.rs | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/naga/src/valid/mod.rs b/naga/src/valid/mod.rs index c1c34c39ad0..da414ab135e 100644 --- a/naga/src/valid/mod.rs +++ b/naga/src/valid/mod.rs @@ -83,7 +83,7 @@ bitflags::bitflags! { #[cfg_attr(feature = "serialize", derive(serde::Serialize))] #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))] #[derive(Clone, Copy, Debug, Eq, PartialEq)] - pub struct Capabilities: u32 { + pub struct Capabilities: u64 { /// Support for [`AddressSpace::PushConstant`][1]. /// /// [1]: crate::AddressSpace::PushConstant