first successful raytracing shader compile (hit shader with rayPayloa…

…dInEXT)

cli/src/bin/naga.rs

@@ -287,7 +287,8 @@ fn run() -> Result<(), Box<dyn std::error::Error>> {
                 }
             }
         }
-        ext @ ("vert" | "frag" | "comp") => {
+        ext @ ("vert" | "frag" | "comp" | "rgen" | "rmiss" | "rcall" | "rchit" | "rahit"
+        | "rint") => {
             let input = String::from_utf8(input)?;
             let mut parser = naga::front::glsl::Frontend::default();
 
@@ -299,6 +300,12 @@ fn run() -> Result<(), Box<dyn std::error::Error>> {
                                 "vert" => naga::ShaderStage::Vertex,
                                 "frag" => naga::ShaderStage::Fragment,
                                 "comp" => naga::ShaderStage::Compute,
+                                "rgen" => naga::ShaderStage::RayGen,
+                                "rmiss" => naga::ShaderStage::Miss,
+                                "rcall" => naga::ShaderStage::Callable,
+                                "rchit" => naga::ShaderStage::ClosestHit,
+                                "rahit" => naga::ShaderStage::AnyHit,
+                                "rint" => naga::ShaderStage::Intersection,
                                 _ => unreachable!(),
                             },
                             defines: Default::default(),

src/back/glsl/mod.rs

@@ -107,6 +107,7 @@ impl crate::AddressSpace {
             | crate::AddressSpace::Storage { .. }
             | crate::AddressSpace::Handle
             | crate::AddressSpace::PushConstant => false,
+            crate::AddressSpace::IncomingRayPayload => unimplemented!(),
         }
     }
 }
@@ -335,6 +336,12 @@ impl fmt::Display for VaryingName<'_> {
             crate::Binding::Location { location, .. } => {
                 let prefix = match (self.stage, self.output) {
                     (ShaderStage::Compute, _) => unreachable!(),
+                    (ShaderStage::RayGen, _) => unreachable!(),
+                    (ShaderStage::Miss, _) => unreachable!(),
+                    (ShaderStage::Callable, _) => unreachable!(),
+                    (ShaderStage::ClosestHit, _) => unreachable!(),
+                    (ShaderStage::AnyHit, _) => unreachable!(),
+                    (ShaderStage::Intersection, _) => unreachable!(),
                     // pipeline to vertex
                     (ShaderStage::Vertex, false) => "p2vs",
                     // vertex to fragment
@@ -361,6 +368,12 @@ impl ShaderStage {
             ShaderStage::Compute => "cs",
             ShaderStage::Fragment => "fs",
             ShaderStage::Vertex => "vs",
+            ShaderStage::RayGen
+            | ShaderStage::Miss
+            | ShaderStage::Callable
+            | ShaderStage::ClosestHit
+            | ShaderStage::AnyHit
+            | ShaderStage::Intersection => unimplemented!(),
         }
     }
 }
@@ -1030,6 +1043,7 @@ impl<'a, W: Write> Writer<'a, W> {
             crate::AddressSpace::Function => unreachable!(),
             // Textures and samplers are handled directly in `Writer::write`.
             crate::AddressSpace::Handle => unreachable!(),
+            crate::AddressSpace::IncomingRayPayload => unimplemented!(),
         }
 
         Ok(())
@@ -1302,6 +1316,12 @@ impl<'a, W: Write> Writer<'a, W> {
             ShaderStage::Vertex => output,
             ShaderStage::Fragment => !output,
             ShaderStage::Compute => false,
+            ShaderStage::RayGen
+            | ShaderStage::Miss
+            | ShaderStage::Callable
+            | ShaderStage::ClosestHit
+            | ShaderStage::AnyHit
+            | ShaderStage::Intersection => unimplemented!(),
         };
 
         // Write the I/O locations, if allowed
@@ -4010,6 +4030,7 @@ const fn glsl_storage_qualifier(space: crate::AddressSpace) -> Option<&'static s
         As::Handle => Some("uniform"),
         As::WorkGroup => Some("shared"),
         As::PushConstant => Some("uniform"),
+        As::IncomingRayPayload => unimplemented!(),
     }
 }
 

src/back/hlsl/mod.rs

@@ -149,6 +149,12 @@ impl crate::ShaderStage {
             Self::Vertex => "vs",
             Self::Fragment => "ps",
             Self::Compute => "cs",
+            Self::RayGen
+            | Self::Miss
+            | Self::Callable
+            | Self::ClosestHit
+            | Self::AnyHit
+            | Self::Intersection => unimplemented!(),
         }
     }
 }

src/back/hlsl/writer.rs

@@ -676,6 +676,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                 write!(self.out, "ConstantBuffer<")?;
                 "b"
             }
+            crate::AddressSpace::IncomingRayPayload => unimplemented!(),
         };
 
         // If the global is a push constant write the type now because it will be a

src/back/msl/writer.rs

@@ -438,6 +438,7 @@ impl crate::AddressSpace {
             | Self::PushConstant
             | Self::Handle => true,
             Self::Function => false,
+            Self::IncomingRayPayload => unimplemented!(),
         }
     }
 
@@ -455,6 +456,7 @@ impl crate::AddressSpace {
             Self::Uniform | Self::PushConstant => false,
             // Not applicable.
             Self::Handle | Self::Function => false,
+            Self::IncomingRayPayload => unimplemented!(),
         }
     }
 
@@ -465,6 +467,7 @@ impl crate::AddressSpace {
             Self::Storage { .. } => Some("device"),
             Self::Private | Self::Function => Some("thread"),
             Self::WorkGroup => Some("threadgroup"),
+            Self::IncomingRayPayload => unimplemented!(),
         }
     }
 }
@@ -3432,6 +3435,7 @@ impl<W: Write> Writer<W> {
                         crate::AddressSpace::Function
                         | crate::AddressSpace::Private
                         | crate::AddressSpace::WorkGroup => {}
+                        crate::AddressSpace::IncomingRayPayload => unimplemented!(),
                     }
                 }
                 if supports_array_length {
@@ -3464,6 +3468,12 @@ impl<W: Write> Writer<W> {
                 crate::ShaderStage::Compute { .. } => {
                     ("kernel", LocationMode::Uniform, LocationMode::Uniform)
                 }
+                crate::ShaderStage::RayGen
+                | crate::ShaderStage::Miss
+                | crate::ShaderStage::Callable
+                | crate::ShaderStage::ClosestHit
+                | crate::ShaderStage::AnyHit
+                | crate::ShaderStage::Intersection => unimplemented!(),
             };
 
             // List all the Naga `EntryPoint`'s `Function`'s arguments,

src/back/spv/helpers.rs

@@ -29,6 +29,7 @@ pub(super) const fn map_storage_class(space: crate::AddressSpace) -> spirv::Stor
         crate::AddressSpace::Uniform => spirv::StorageClass::Uniform,
         crate::AddressSpace::WorkGroup => spirv::StorageClass::Workgroup,
         crate::AddressSpace::PushConstant => spirv::StorageClass::PushConstant,
+        crate::AddressSpace::IncomingRayPayload => spirv::StorageClass::IncomingRayPayloadKHR,
     }
 }
 

src/back/spv/writer.rs

@@ -766,6 +766,13 @@ impl Writer {
                 .to_words(&mut self.logical_layout.execution_modes);
                 spirv::ExecutionModel::GLCompute
             }
+            // TODO
+            crate::ShaderStage::RayGen => spirv::ExecutionModel::RayGenerationKHR,
+            crate::ShaderStage::Miss => spirv::ExecutionModel::MissKHR,
+            crate::ShaderStage::Callable => spirv::ExecutionModel::CallableKHR,
+            crate::ShaderStage::ClosestHit => spirv::ExecutionModel::ClosestHitKHR,
+            crate::ShaderStage::AnyHit => spirv::ExecutionModel::AnyHitKHR,
+            crate::ShaderStage::Intersection => spirv::ExecutionModel::IntersectionKHR,
         };
         //self.check(exec_model.required_capabilities())?;
 
@@ -1587,6 +1594,24 @@ impl Writer {
             }
         };
 
+        match global_variable.space {
+            crate::AddressSpace::Function
+            | crate::AddressSpace::Private
+            | crate::AddressSpace::WorkGroup
+            | crate::AddressSpace::Uniform
+            | crate::AddressSpace::Storage { .. }
+            | crate::AddressSpace::Handle
+            | crate::AddressSpace::PushConstant => {}
+            crate::AddressSpace::IncomingRayPayload => {
+                self.require_any("Incoming Ray Payload", &[spirv::Capability::RayTracingKHR])?;
+                self.use_extension("SPV_KHR_ray_tracing");
+            }
+        };
+
+        if let Some(location) = global_variable.location {
+            self.decorate(id, Decoration::Location, &[location]);
+        };
+
         let init_word = global_variable
             .init
             .map(|constant| self.constant_ids[constant.index()]);

src/back/wgsl/writer.rs

@@ -157,6 +157,12 @@ impl<W: Write> Writer<W> {
                     Attribute::Stage(ShaderStage::Compute),
                     Attribute::WorkGroupSize(ep.workgroup_size),
                 ],
+                ShaderStage::RayGen
+                | ShaderStage::Miss
+                | ShaderStage::Callable
+                | ShaderStage::ClosestHit
+                | ShaderStage::AnyHit
+                | ShaderStage::Intersection => unimplemented!(),
             };
 
             self.write_attributes(&attributes)?;
@@ -209,6 +215,12 @@ impl<W: Write> Writer<W> {
                     ShaderStage::Compute => "ComputeOutput",
                     ShaderStage::Fragment => "FragmentOutput",
                     ShaderStage::Vertex => "VertexOutput",
+                    ShaderStage::RayGen
+                    | ShaderStage::Miss
+                    | ShaderStage::Callable
+                    | ShaderStage::ClosestHit
+                    | ShaderStage::AnyHit
+                    | ShaderStage::Intersection => unimplemented!(),
                 };
 
                 write!(self.out, "{name}")?;
@@ -343,6 +355,12 @@ impl<W: Write> Writer<W> {
                         ShaderStage::Vertex => "vertex",
                         ShaderStage::Fragment => "fragment",
                         ShaderStage::Compute => "compute",
+                        ShaderStage::RayGen
+                        | ShaderStage::Miss
+                        | ShaderStage::Callable
+                        | ShaderStage::ClosestHit
+                        | ShaderStage::AnyHit
+                        | ShaderStage::Intersection => unimplemented!(),
                     };
                     write!(self.out, "@{stage_str} ")?;
                 }
@@ -1931,6 +1949,7 @@ const fn address_space_str(
             As::WorkGroup => "workgroup",
             As::Handle => return (None, None),
             As::Function => "function",
+            As::IncomingRayPayload => unimplemented!(),
         }),
         None,
     )

src/front/glsl/lex.rs

@@ -108,6 +108,7 @@ impl<'a> Iterator for Lexer<'a> {
                     // types
                     "void" => TokenValue::Void,
                     "struct" => TokenValue::Struct,
+                    "rayPayloadInEXT" => TokenValue::RayPayloadInEXT,
                     word => match parse_type(word) {
                         Some(t) => TokenValue::TypeName(t),
                         None => TokenValue::Identifier(String::from(word)),

src/front/glsl/parser/types.rs

@@ -156,7 +156,8 @@ impl<'source> ParsingContext<'source> {
             | TokenValue::Buffer
             | TokenValue::Restrict
             | TokenValue::MemoryQualifier(_)
-            | TokenValue::Layout => true,
+            | TokenValue::Layout
+            | TokenValue::RayPayloadInEXT => true,
             _ => false,
         })
     }
@@ -210,7 +211,8 @@ impl<'source> ParsingContext<'source> {
                 | TokenValue::Out
                 | TokenValue::Uniform
                 | TokenValue::Shared
-                | TokenValue::Buffer => {
+                | TokenValue::Buffer
+                | TokenValue::RayPayloadInEXT => {
                     let storage = match token.value {
                         TokenValue::Const => StorageQualifier::Const,
                         TokenValue::In => StorageQualifier::Input,
@@ -226,6 +228,9 @@ impl<'source> ParsingContext<'source> {
                                 access: crate::StorageAccess::all(),
                             })
                         }
+                        TokenValue::RayPayloadInEXT => {
+                            StorageQualifier::AddressSpace(AddressSpace::IncomingRayPayload)
+                        }
                         _ => unreachable!(),
                     };
 

src/front/glsl/token.rs

@@ -36,6 +36,7 @@ pub enum TokenValue {
     Buffer,
     Const,
     Shared,
+    RayPayloadInEXT,
 
     Restrict,
     /// A `glsl` memory qualifier such as `writeonly`

src/front/glsl/variables.rs

@@ -58,6 +58,7 @@ impl Frontend {
                 name: Some(name.into()),
                 space: AddressSpace::Private,
                 binding: None,
+                location: None,
                 ty,
                 init: None,
             },
@@ -462,6 +463,7 @@ impl Frontend {
                         name: name.clone(),
                         space: AddressSpace::Private,
                         binding: None,
+                        location: None,
                         ty,
                         init,
                     },
@@ -600,11 +602,24 @@ impl Frontend {
                     _ => None,
                 };
 
+                let location = match space {
+                    AddressSpace::IncomingRayPayload => {
+                        // TODO: glslang seems to use a counter for variables without
+                        // explicit location (even if that causes collisions)
+                        let location = qualifiers
+                            .uint_layout_qualifier("location", &mut self.errors)
+                            .unwrap_or(0);
+                        Some(location)
+                    }
+                    _ => None,
+                };
+
                 let handle = self.module.global_variables.append(
                     GlobalVariable {
                         name: name.clone(),
                         space,
                         binding,
+                        location,
                         ty,
                         init,
                     },

src/front/spv/mod.rs

@@ -5003,6 +5003,7 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                 }
                 let var = crate::GlobalVariable {
                     binding: dec.resource_binding(),
+                    location: None,
                     name: dec.name,
                     space,
                     ty,
@@ -5048,6 +5049,7 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                     name: dec.name.clone(),
                     space: crate::AddressSpace::Private,
                     binding: None,
+                    location: None,
                     ty,
                     init: None,
                 };
@@ -5121,6 +5123,7 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                     name: dec.name,
                     space: crate::AddressSpace::Private,
                     binding: None,
+                    location: None,
                     ty,
                     init,
                 };

src/front/wgsl/lower/mod.rs

@@ -666,6 +666,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                             name: Some(v.name.name.to_string()),
                             space: v.space,
                             binding: v.binding.clone(),
+                            location: None,
                             ty,
                             init,
                         },

src/lib.rs

@@ -302,6 +302,12 @@ pub enum ShaderStage {
     Vertex,
     Fragment,
     Compute,
+    RayGen,
+    Miss,
+    Callable,
+    ClosestHit,
+    AnyHit,
+    Intersection,
 }
 
 /// Addressing space of variables.
@@ -324,6 +330,8 @@ pub enum AddressSpace {
     Handle,
     /// Push constants.
     PushConstant,
+    /// Incoming ray payload for SPV_KHR_ray_tracing
+    IncomingRayPayload,
 }
 
 /// Built-in inputs and outputs.
@@ -861,6 +869,8 @@ pub struct GlobalVariable {
     pub space: AddressSpace,
     /// For resources, defines the binding point.
     pub binding: Option<ResourceBinding>,
+    /// For some ray-tracing storage qualifiers, define the location
+    pub location: Option<u32>,
     /// The type of this variable.
     pub ty: Handle<Type>,
     /// Initial value for this variable.