diff --git a/CHANGELOG.md b/CHANGELOG.md index 55628e5cd2..7231ae3521 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -143,6 +143,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 diff --git a/docs/api-specs/ray_tracing.md b/docs/api-specs/ray_tracing.md index 4c69e476de..94eaafe158 100644 --- a/docs/api-specs/ray_tracing.md +++ b/docs/api-specs/ray_tracing.md @@ -93,10 +93,12 @@ fn render(/*whatever args you need to render*/) { ## `naga`'s raytracing API: -`naga` supports ray queries (also known as inline raytracing) only. To enable basic ray query functions you must add +`naga` supports ray queries (also known as inline raytracing). To enable basic ray query functions you must add `enable wgpu_ray_query` to the shader, ray queries and acceleration structures also support tags which require extra -`enable` extensions (see Acceleration structure tags for more info). Ray tracing pipelines -are currently unsupported. Naming is mostly taken from vulkan. +`enable` extensions (see Acceleration structure tags for more info). 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` @@ -147,7 +149,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 @@ -271,6 +273,39 @@ const RAY_QUERY_INTERSECTION_GENERATED = 2; 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() {} +``` ### Acceleration structure tags These are tags that can be added to a acceleration structure (`acceleration_structure` -> diff --git a/naga/src/back/dot/mod.rs b/naga/src/back/dot/mod.rs index 826dad1c21..6feb1ca73a 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/conv.rs b/naga/src/back/glsl/conv.rs index 65e332df0b..3cae94d061 100644 --- a/naga/src/back/glsl/conv.rs +++ b/naga/src/back/glsl/conv.rs @@ -125,7 +125,20 @@ pub(in crate::back::glsl) const fn glsl_built_in( | 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!() } } @@ -145,7 +158,7 @@ pub(in crate::back::glsl) const fn glsl_storage_qualifier( As::Handle => Some("uniform"), As::WorkGroup => Some("shared"), As::Immediate => Some("uniform"), - As::TaskPayload => unreachable!(), + As::TaskPayload | As::RayPayload | As::IncomingRayPayload => unreachable!(), } } diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index 9f9e504a44..e2ff8a60c2 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -147,6 +147,10 @@ impl crate::AddressSpace { | crate::AddressSpace::Handle | crate::AddressSpace::Immediate | crate::AddressSpace::TaskPayload => false, + + crate::AddressSpace::RayPayload | crate::AddressSpace::IncomingRayPayload => { + unreachable!() + } } } } @@ -510,7 +514,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}",) } @@ -527,7 +539,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!(), } } } diff --git a/naga/src/back/glsl/writer.rs b/naga/src/back/glsl/writer.rs index 4d23d2873a..014551ef88 100644 --- a/naga/src/back/glsl/writer.rs +++ b/naga/src/back/glsl/writer.rs @@ -727,6 +727,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(()) @@ -1093,7 +1097,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 @@ -2228,6 +2237,7 @@ impl<'a, W: Write> Writer<'a, W> { } writeln!(self.out, ");")?; } + Statement::RayPipelineFunction(_) => unimplemented!(), } Ok(()) diff --git a/naga/src/back/hlsl/conv.rs b/naga/src/back/hlsl/conv.rs index 6cd3679e81..ab534019dc 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 0230434a74..496221f7da 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 cc85e512f2..a4952d45b5 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 immediate data 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::Immediate - | 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 d764b17ebb..c37821a2ed 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 77d6814c60..28fc4fd311 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -598,6 +598,7 @@ impl crate::AddressSpace { | Self::Handle | Self::TaskPayload => true, Self::Function => false, + Self::RayPayload | Self::IncomingRayPayload => unreachable!(), } } @@ -609,7 +610,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 +625,13 @@ impl crate::AddressSpace { Self::Handle => None, Self::Uniform | Self::Immediate => 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 +4199,7 @@ impl Writer { } writeln!(self.out, ");")?; } + crate::Statement::RayPipelineFunction(_) => unreachable!(), } } @@ -6672,6 +6678,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 +6754,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 de643b82fa..8e596b6cc3 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 ace9565a1c..1a95b3eb03 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 16d5867593..5d302ec3b6 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::Immediate => 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 71b761c736..2c3bce65d6 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 daf32a7116..253e59f6d1 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 c11e8f9adf..0ebbfac3e7 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 39d6065f5f..b490616fce 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 ba096a82b3..5d1f930452 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 48b23e7c4c..a88bcf306b 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 e8b73fb7ac..c72278350a 100644 --- a/naga/src/front/spv/mod.rs +++ b/naga/src/front/spv/mod.rs @@ -1592,7 +1592,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/error.rs b/naga/src/front/wgsl/error.rs index 0cd7e11c73..3db52fc10e 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 c47941f0f8..8433dae9f2 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,6 +1653,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { function, mesh_info, task_payload, + ray_incoming_payload, }); Ok(LoweredGlobalDecl::EntryPoint( ctx.module.entry_points.len() - 1, @@ -3382,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 04964e7ba5..2d6267d89c 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 3fe4643ac5..9e0ebd3ec5 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 a4ca097df8..13c163a6b8 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 142606b83e..248b828f35 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/src/ir/mod.rs b/naga/src/ir/mod.rs index eb6e989e37..7b738b8dfb 100644 --- a/naga/src/ir/mod.rs +++ b/naga/src/ir/mod.rs @@ -335,6 +335,18 @@ 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 +385,14 @@ pub enum AddressSpace { Immediate, /// 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. @@ -394,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, @@ -409,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, @@ -459,6 +479,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 +2280,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 +2461,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 +2679,35 @@ pub struct MeshStageInfo { pub output_variable: Handle, } +/// Ray tracing pipeline 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 b800444d58..a8c69efdd6 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 127e346f3a..108b377c3d 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 b29ccb054a..f15e5ecc83 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 a25377c538..87c35da596 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 abf6bc430a..76ae7c286f 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 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), } 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 self.trace_rays_no_vertex_return.is_none() { + 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::InvalidPayloadAddressSpace( + 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 5b7fb3fab7..2e61a3b9ad 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 1e054700c0..a1e1918d54 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(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(crate::ShaderStage), } fn storage_usage(access: crate::StorageAccess) -> GlobalUse { @@ -229,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, @@ -252,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 @@ -294,6 +306,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 +317,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), ), @@ -317,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), @@ -361,7 +379,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 +423,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) @@ -531,7 +743,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 => 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, }; @@ -757,6 +974,17 @@ impl super::Validator { false, ) } + crate::AddressSpace::RayPayload | crate::AddressSpace::IncomingRayPayload => { + if !self + .capabilities + .contains(Capabilities::RAY_TRACING_PIPELINE) + { + return Err(GlobalVariableError::UnsupportedCapability( + Capabilities::RAY_TRACING_PIPELINE, + )); + } + (TypeFlags::DATA | TypeFlags::SIZED, false) + } }; if !type_info.flags.contains(required_type_flags) { @@ -862,14 +1090,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; @@ -958,6 +1200,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) { @@ -1095,6 +1341,30 @@ impl super::Validator { } } crate::AddressSpace::Immediate => GlobalUse::READ, + crate::AddressSpace::RayPayload => { + if !matches!( + ep.stage, + crate::ShaderStage::RayGeneration + | crate::ShaderStage::ClosestHit + | crate::ShaderStage::Miss + ) { + return Err(EntryPointError::RayPayloadInInvalidStage(ep.stage) + .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(ep.stage) + .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 787746ff56..08f95cde2d 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::Immediate`][1]. /// /// [1]: crate::AddressSpace::Immediate @@ -192,6 +192,8 @@ bitflags::bitflags! { const MESH_SHADER = 1 << 30; /// Support for mesh shaders which output points. const MESH_SHADER_POINT_TOPOLOGY = 1 << 31; + /// Support for ray generation, any hit, closest hit, and miss shaders. + const RAY_TRACING_PIPELINE = 1 << 32; } } @@ -210,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, } } @@ -282,12 +285,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 +369,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 +580,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 62aaf97b79..90c054ace6 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 @@ -765,7 +765,8 @@ impl super::Validator { Alignment::ONE, ), Ti::AccelerationStructure { vertex_return } => { - self.require_type_capability(Capabilities::RAY_QUERY)?; + 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/in/wgsl/ray-tracing-pipeline.toml b/naga/tests/in/wgsl/ray-tracing-pipeline.toml new file mode 100644 index 0000000000..9845f3bb7e --- /dev/null +++ b/naga/tests/in/wgsl/ray-tracing-pipeline.toml @@ -0,0 +1,2 @@ +god_mode = true +targets = "IR" 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 0000000000..02a40b05c5 --- /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/naga/snapshots.rs b/naga/tests/naga/snapshots.rs index e01610f29a..52062a70b3 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 => 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/naga/tests/naga/wgsl_errors.rs b/naga/tests/naga/wgsl_errors.rs index bc4386d015..7d371a5ce1 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 ); } @@ -4379,3 +4406,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 + ); + } +} diff --git a/naga/tests/out/analysis/spv-shadow.info.ron b/naga/tests/out/analysis/spv-shadow.info.ron index 381f841d5d..c2adf7d22d 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 c22cd768f2..32ba1b0b29 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 219e016f8d..9647887a0d 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 9f9feac9c0..b148e4d7bc 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 b4c8790508..efd0246107 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 fa62867576..479da9b9c9 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 eacd33ad0f..4018c5de44 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 92e99112e5..d6fcee4aab 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 8bb298a645..95b9484bd4 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 98f4426c3e..6a91852a1c 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 104de852c1..32eafce671 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 bed86a5334..e558039862 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 bdda1d1856..010ed83eba 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 67eb29c247..748f599b37 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 51686aa20e..f5a97c5115 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 c3df0c8c50..a29a077767 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 c3df0c8c50..a29a077767 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 fc4daaa129..2f55d55f7a 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 fc4daaa129..2f55d55f7a 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 648f4ff9bc..5b312fcc77 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 648f4ff9bc..5b312fcc77 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 9a2bf193f3..c1285a3aba 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 9a2bf193f3..c1285a3aba 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 addd0e5871..66b88c7047 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 addd0e5871..66b88c7047 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 0e4e2e4d40..6a401b17b7 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 0e4e2e4d40..6a401b17b7 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 8bdaa72da6..72f20f638c 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 8bdaa72da6..72f20f638c 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 b298f29698..1759974b28 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 b298f29698..1759974b28 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 558a88e28d..05c0bf7dfe 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 558a88e28d..05c0bf7dfe 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 9cb3c6479c..55bb0ea5d0 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 9cb3c6479c..55bb0ea5d0 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 16e925f2fb..753db1f579 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 16e925f2fb..753db1f579 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 28a824bb03..e0e5f68bd9 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 28a824bb03..e0e5f68bd9 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 152a45008c..a3e0e5384f 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 152a45008c..a3e0e5384f 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 fe136e71e4..95db50b89a 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 fe136e71e4..95db50b89a 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 0000000000..61148d0268 --- /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 0000000000..61148d0268 --- /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 68c867a19e..0bf1bf3ee2 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 68c867a19e..0bf1bf3ee2 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 db619dff83..8d5cacc288 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 db619dff83..8d5cacc288 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 689fe215e3..1fd051f251 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 689fe215e3..1fd051f251 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 7c0d856946..6aa657e0d8 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 34e44cb965..2ea01ca40e 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: [], diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs index 1a4ca444da..1fc16b5d6d 100644 --- a/wgpu-core/src/validation.rs +++ b/wgpu-core/src/validation.rs @@ -1302,7 +1302,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::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 3fef171d9b..c4d3c3fc86 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 9f41b77383..f2192fcd38 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 => 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 => 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 173fb2fa55..1241741c3e 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 c3f879627c..5c4aace444 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!(), } } } diff --git a/wgpu-types/src/binding.rs b/wgpu-types/src/binding.rs index edd3e55bd2..b75ecc0e34 100644 --- a/wgpu-types/src/binding.rs +++ b/wgpu-types/src/binding.rs @@ -221,6 +221,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; } }