diff --git a/naga/src/front/glsl/builtins.rs b/naga/src/front/glsl/builtins.rs index d07ba9c1281..b421d49ff68 100644 --- a/naga/src/front/glsl/builtins.rs +++ b/naga/src/front/glsl/builtins.rs @@ -464,89 +464,10 @@ fn inject_standard_builtins( module: &mut Module, name: &str, ) { - // Some samplers (sampler1D, etc...) can be float, int, or uint - let anykind_sampler = if name.starts_with("sampler") { - Some((name, Sk::Float)) - } else if name.starts_with("usampler") { - Some((&name[1..], Sk::Uint)) - } else if name.starts_with("isampler") { - Some((&name[1..], Sk::Sint)) - } else { - None - }; - if let Some((sampler, kind)) = anykind_sampler { - match sampler { - "sampler1D" | "sampler1DArray" | "sampler2D" | "sampler2DArray" | "sampler2DMS" - | "sampler2DMSArray" | "sampler3D" | "samplerCube" | "samplerCubeArray" => { - declaration.overloads.push(module.add_builtin( - vec![ - TypeInner::Image { - dim: match sampler { - "sampler1D" | "sampler1DArray" => Dim::D1, - "sampler2D" | "sampler2DArray" | "sampler2DMS" - | "sampler2DMSArray" => Dim::D2, - "sampler3D" => Dim::D3, - _ => Dim::Cube, - }, - arrayed: matches!( - sampler, - "sampler1DArray" - | "sampler2DArray" - | "sampler2DMSArray" - | "samplerCubeArray" - ), - class: ImageClass::Sampled { - kind, - multi: matches!(sampler, "sampler2DMS" | "sampler2DMSArray"), - }, - }, - TypeInner::Sampler { comparison: false }, - ], - MacroCall::Sampler, - )); - return; - } - _ => (), - } - } + // `sampler2D(tex, samp)` constructors are handled in constructor_many() — the lexer + // emits CombinedSamplerTypeName, so they never arrive here as function identifiers. match name { - // Shadow sampler can only be of kind `Sk::Float` - "sampler1DShadow" - | "sampler1DArrayShadow" - | "sampler2DShadow" - | "sampler2DArrayShadow" - | "samplerCubeShadow" - | "samplerCubeArrayShadow" => { - let dim = match name { - "sampler1DShadow" | "sampler1DArrayShadow" => Dim::D1, - "sampler2DShadow" | "sampler2DArrayShadow" => Dim::D2, - _ => Dim::Cube, - }; - let arrayed = matches!( - name, - "sampler1DArrayShadow" | "sampler2DArrayShadow" | "samplerCubeArrayShadow" - ); - - for i in 0..2 { - let ty = TypeInner::Image { - dim, - arrayed, - class: match i { - 0 => ImageClass::Sampled { - kind: Sk::Float, - multi: false, - }, - _ => ImageClass::Depth { multi: false }, - }, - }; - - declaration.overloads.push(module.add_builtin( - vec![ty, TypeInner::Sampler { comparison: true }], - MacroCall::SamplerShadow, - )) - } - } "sin" | "exp" | "exp2" | "sinh" | "cos" | "cosh" | "tan" | "tanh" | "acos" | "asin" | "log" | "log2" | "radians" | "degrees" | "asinh" | "acosh" | "atanh" | "floatBitsToInt" | "floatBitsToUint" | "dFdx" | "dFdxFine" | "dFdxCoarse" | "dFdy" @@ -1543,8 +1464,6 @@ pub enum TextureLevelType { /// A compiler defined builtin function #[derive(Clone, Copy, PartialEq, Debug)] pub enum MacroCall { - Sampler, - SamplerShadow, Texture { proj: bool, offset: bool, @@ -1593,16 +1512,6 @@ impl MacroCall { meta: Span, ) -> Result>> { Ok(Some(match *self { - MacroCall::Sampler => { - ctx.samplers.insert(args[0], args[1]); - args[0] - } - MacroCall::SamplerShadow => { - sampled_to_depth(ctx, args[0], meta, &mut frontend.errors); - ctx.invalidate_expression(args[0], meta)?; - ctx.samplers.insert(args[0], args[1]); - args[0] - } MacroCall::Texture { proj, offset, diff --git a/naga/src/front/glsl/context.rs b/naga/src/front/glsl/context.rs index 1b681068d3a..4ecde71f4cc 100644 --- a/naga/src/front/glsl/context.rs +++ b/naga/src/front/glsl/context.rs @@ -121,6 +121,35 @@ impl<'a> Context<'a> { for &(ref name, lookup) in frontend.global_variables.iter() { this.add_global(name, lookup)? } + + // For each combined image-sampler uniform (e.g. `sampler2D u_tex`), add an + // entry to `ctx.samplers` mapping the image expression to its implicit + // companion sampler expression. This enables `texture(u_tex, uv)` to be + // lowered to `textureSample(u_tex, u_tex_sampler, uv)` without requiring the + // caller to write `texture(sampler2D(u_tex, u_samp), uv)`. + // + // The image global's expression was appended by the `add_global` loop above. + // The implicit sampler global was NOT added to `frontend.global_variables` (it + // is anonymous from the GLSL source's perspective), so we add its expression + // here and immediately map it. + for &(image_handle, sampler_handle) in frontend.combined_samplers.iter() { + // Locate the expression that was already created for the image global. + let maybe_image_expr = this + .expressions + .iter() + .find(|&(_, expr)| *expr == Expression::GlobalVariable(image_handle)) + .map(|(h, _)| h); + + if let Some(image_expr) = maybe_image_expr { + let span = this.module.global_variables.get_span(sampler_handle); + // The implicit sampler global was not added to frontend.global_variables, + // so we add its expression here and populate the samplers map. + let sampler_expr = + this.add_expression(Expression::GlobalVariable(sampler_handle), span)?; + this.samplers.insert(image_expr, sampler_expr); + } + } + this.is_const = is_const; Ok(this) diff --git a/naga/src/front/glsl/functions.rs b/naga/src/front/glsl/functions.rs index b26061cecbd..69e8818cb58 100644 --- a/naga/src/front/glsl/functions.rs +++ b/naga/src/front/glsl/functions.rs @@ -489,6 +489,26 @@ impl Frontend { .map(|member| scalar_components(&ctx.module.types[member.ty].inner)) .collect::>(), ), + TypeInner::Image { + class: crate::ImageClass::Sampled { .. }, + .. + } if args.len() == 2 => { + let (image, _) = args[0]; + let (sampler, _) = args[1]; + ctx.samplers.insert(image, sampler); + return Ok(image); + } + TypeInner::Image { + class: crate::ImageClass::Depth { .. }, + .. + } if args.len() == 2 => { + let (image, image_meta) = args[0]; + let (sampler, _) = args[1]; + sampled_to_depth(ctx, image, image_meta, &mut self.errors); + ctx.invalidate_expression(image, image_meta)?; + ctx.samplers.insert(image, sampler); + return Ok(image); + } _ => { return Err(Error { kind: ErrorKind::SemanticError("Constructor: Too many arguments".into()), diff --git a/naga/src/front/glsl/lex.rs b/naga/src/front/glsl/lex.rs index 9337491b5da..61ef3d410ee 100644 --- a/naga/src/front/glsl/lex.rs +++ b/naga/src/front/glsl/lex.rs @@ -8,7 +8,7 @@ use pp_rs::{ use super::{ ast::Precision, token::{Directive, DirectiveKind, Token, TokenValue}, - types::parse_type, + types::{is_combined_sampler, parse_type}, }; use crate::{FastHashMap, Span, StorageAccess}; @@ -112,6 +112,9 @@ impl Iterator for Lexer<'_> { "void" => TokenValue::Void, "struct" => TokenValue::Struct, word => match parse_type(word) { + Some(t) if is_combined_sampler(word) => { + TokenValue::CombinedSamplerTypeName(t) + } Some(t) => TokenValue::TypeName(t), None => TokenValue::Identifier(String::from(word)), }, diff --git a/naga/src/front/glsl/mod.rs b/naga/src/front/glsl/mod.rs index e5eda6b3ad9..4ae3e3b60cc 100644 --- a/naga/src/front/glsl/mod.rs +++ b/naga/src/front/glsl/mod.rs @@ -18,7 +18,10 @@ pub use token::TokenValue; use alloc::{string::String, vec::Vec}; -use crate::{proc::Layouter, FastHashMap, FastHashSet, Handle, Module, ShaderStage, Span, Type}; +use crate::{ + proc::Layouter, FastHashMap, FastHashSet, GlobalVariable, Handle, Module, ShaderStage, Span, + Type, +}; use ast::{EntryArg, FunctionDeclaration, GlobalLookup}; use parser::ParsingContext; @@ -174,6 +177,14 @@ pub struct Frontend { entry_args: Vec, + /// Maps each combined-sampler image global (e.g. from `sampler2D` / `sampler2DShadow` + /// uniforms) to an implicit paired sampler global that was synthesized for it. + /// + /// When `texture(u_tex, uv)` is called with a combined-sampler uniform, the texture + /// call machinery looks up `ctx.samplers[image_expr]` to find the WGSL sampler. + /// These pairs get translated into `ctx.samplers` entries in `Context::new`. + pub(crate) combined_samplers: Vec<(Handle, Handle)>, + layouter: Layouter, errors: Vec, @@ -187,6 +198,7 @@ impl Frontend { self.lookup_type.clear(); self.global_variables.clear(); self.entry_args.clear(); + self.combined_samplers.clear(); self.layouter.clear(); } diff --git a/naga/src/front/glsl/parser.rs b/naga/src/front/glsl/parser.rs index 959a9b8e60d..24513f87780 100644 --- a/naga/src/front/glsl/parser.rs +++ b/naga/src/front/glsl/parser.rs @@ -181,6 +181,66 @@ impl<'source> ParsingContext<'source> { self.parse_external_declaration(frontend, &mut ctx)?; } + // Fix up bindings for all implicit samplers that were synthesised for combined + // image-sampler uniforms (`sampler2D`, `sampler2DShadow`, …). The sampler + // globals were initially created with the same binding as their paired image, + // which would cause `BindingCollision` validation errors. Now that all + // declarations have been processed we can see the full binding landscape and + // assign each implicit sampler the next free slot on its descriptor set. + // + // We do this in a single pass: for each set we collect the highest binding + // currently allocated, then hand out sequential slots for the samplers. + if !frontend.combined_samplers.is_empty() { + // Collect the set of implicit sampler handles so we can exclude them + // from the "already allocated" scan. + let implicit_sampler_handles: alloc::collections::BTreeSet<_> = + frontend.combined_samplers.iter().map(|&(_, s)| s).collect(); + + // Build a map: group -> (one past the highest binding currently in use), + // counting only the non-implicit globals. + let mut group_next_binding: crate::FastHashMap = + crate::FastHashMap::default(); + for (handle, var) in ctx.module.global_variables.iter() { + if implicit_sampler_handles.contains(&handle) { + continue; + } + if let Some(ref rb) = var.binding { + let entry = group_next_binding.entry(rb.group).or_insert(0); + if rb.binding + 1 > *entry { + *entry = rb.binding + 1; + } + } + } + + // Assign the implicit samplers non-conflicting bindings in the same group + // as their paired image, after all the declared bindings. + for &(image_handle, sampler_handle) in frontend.combined_samplers.iter() { + let group = ctx + .module + .global_variables + .get_mut(image_handle) + .binding + .as_ref() + .map(|rb| rb.group) + .unwrap_or(0); + + let next = group_next_binding.entry(group).or_insert(0); + if let Some(ref mut rb) = + ctx.module.global_variables.get_mut(sampler_handle).binding + { + rb.group = group; + rb.binding = *next; + } else { + ctx.module.global_variables.get_mut(sampler_handle).binding = + Some(crate::ResourceBinding { + group, + binding: *next, + }); + } + *next += 1; + } + } + // Add an `EntryPoint` to `parser.module` for `main`, if a // suitable overload exists. Error out if we can't find one. if let Some(declaration) = frontend.lookup_function.get("main") { @@ -411,6 +471,9 @@ pub struct DeclarationContext<'ctx, 'qualifiers, 'a> { external: bool, is_inside_loop: bool, ctx: &'ctx mut Context<'a>, + /// `true` when the current declaration's type was a combined image-sampler type + /// name in GLSL source (`sampler2D`, `isampler3D`, `sampler2DShadow`, …). + is_combined_sampler: bool, } impl DeclarationContext<'_, '_, '_> { @@ -428,6 +491,7 @@ impl DeclarationContext<'_, '_, '_> { name: Some(name), init, meta, + is_combined_sampler: self.is_combined_sampler, }; match self.external { diff --git a/naga/src/front/glsl/parser/declarations.rs b/naga/src/front/glsl/parser/declarations.rs index e3cb321c944..92cfb6f0140 100644 --- a/naga/src/front/glsl/parser/declarations.rs +++ b/naga/src/front/glsl/parser/declarations.rs @@ -315,7 +315,15 @@ impl ParsingContext<'_> { if self.peek_type_name(frontend) { // This branch handles variables and function prototypes and if - // external is true also function definitions + // external is true also function definitions. + // + // Check whether the upcoming token is a combined image-sampler type + // name (`sampler2D`, `isampler3D`, `sampler2DShadow`, …) before + // consuming it — after `parse_type` the token has been consumed and + // the flag cannot be recovered from the produced `Handle` alone. + let is_combined_sampler = self + .peek(frontend) + .is_some_and(|t| matches!(t.value, TokenValue::CombinedSamplerTypeName(_))); let (ty, mut meta) = self.parse_type(frontend, ctx)?; let token = self.bump(frontend)?; @@ -403,6 +411,7 @@ impl ParsingContext<'_> { external, is_inside_loop, ctx, + is_combined_sampler, }; self.backtrack(token_fallthrough)?; @@ -597,6 +606,7 @@ impl ParsingContext<'_> { name, init: None, meta, + is_combined_sampler: false, }, )?; diff --git a/naga/src/front/glsl/parser/functions.rs b/naga/src/front/glsl/parser/functions.rs index 441e99130a0..7776a85b09d 100644 --- a/naga/src/front/glsl/parser/functions.rs +++ b/naga/src/front/glsl/parser/functions.rs @@ -469,6 +469,7 @@ impl ParsingContext<'_> { name: Some(name), init: None, meta, + is_combined_sampler: false, }; let pointer = frontend.add_local_var(ctx, decl)?; diff --git a/naga/src/front/glsl/parser/types.rs b/naga/src/front/glsl/parser/types.rs index 3eef1873c63..3683612c060 100644 --- a/naga/src/front/glsl/parser/types.rs +++ b/naga/src/front/glsl/parser/types.rs @@ -80,7 +80,15 @@ impl ParsingContext<'_> { let token = self.bump(frontend)?; let mut handle = match token.value { TokenValue::Void => return Ok((None, token.meta)), - TokenValue::TypeName(ty) => ctx.module.types.insert(ty, token.meta), + // Combined sampler types (sampler2D, isampler3D, sampler2DShadow, …) carry + // the same `Image` inner type as their `textureXX` counterparts but are + // tokenised separately so variable declarations can synthesise a paired + // implicit sampler. Here they are inserted into the type arena the same way + // as any other type name — the `CombinedSamplerTypeName` distinction is only + // relevant at global-variable-declaration time (see `add_global_var`). + TokenValue::TypeName(ty) | TokenValue::CombinedSamplerTypeName(ty) => { + ctx.module.types.insert(ty, token.meta) + } TokenValue::Struct => { let mut meta = token.meta; let ty_name = self.expect_ident(frontend)?.0; @@ -390,7 +398,9 @@ impl ParsingContext<'_> { pub fn peek_type_name(&mut self, frontend: &mut Frontend) -> bool { self.peek(frontend).is_some_and(|t| match t.value { - TokenValue::TypeName(_) | TokenValue::Void => true, + TokenValue::TypeName(_) | TokenValue::CombinedSamplerTypeName(_) | TokenValue::Void => { + true + } TokenValue::Struct => true, TokenValue::Identifier(ref ident) => frontend.lookup_type.contains_key(ident), _ => false, diff --git a/naga/src/front/glsl/token.rs b/naga/src/front/glsl/token.rs index 78c23c56c29..89aa429fa0e 100644 --- a/naga/src/front/glsl/token.rs +++ b/naga/src/front/glsl/token.rs @@ -69,6 +69,15 @@ pub enum TokenValue { Void, Struct, TypeName(Type), + /// A combined image-sampler type name (`sampler2D`, `isampler3D`, `sampler2DShadow`, …). + /// + /// Emitted for GLSL `samplerXX` types (excluding the bare `sampler` and `samplerShadow` + /// tokens, which carry an explicit `TypeInner::Sampler` and are tokenised as plain + /// [`TypeName`](TokenValue::TypeName)). The inner `Type` is the same `Image` type that + /// `TypeName` would carry — the separate variant exists so that the variable-declaration + /// code can distinguish "this was a combined sampler" from "this was a plain texture type" + /// at the point where a paired implicit sampler needs to be synthesised. + CombinedSamplerTypeName(Type), Assign, AddAssign, diff --git a/naga/src/front/glsl/types.rs b/naga/src/front/glsl/types.rs index dffdf96e6dd..a5953617261 100644 --- a/naga/src/front/glsl/types.rs +++ b/naga/src/front/glsl/types.rs @@ -188,14 +188,101 @@ pub fn parse_type(type_name: &str) -> Option { }) }; + let sampler_parse = |word: &str| { + let sampler_kind = |prefix: &str| { + Some(match prefix { + "" => ScalarKind::Float, + "i" => ScalarKind::Sint, + "u" => ScalarKind::Uint, + _ => return None, + }) + }; + + let shadow_parse = |suffix: &str| { + let suffix = suffix.strip_prefix("sampler")?; + let (dim, arrayed) = match suffix { + "1DShadow" => (ImageDimension::D1, false), + "1DArrayShadow" => (ImageDimension::D1, true), + "2DShadow" => (ImageDimension::D2, false), + "2DArrayShadow" => (ImageDimension::D2, true), + "CubeShadow" => (ImageDimension::Cube, false), + "CubeArrayShadow" => (ImageDimension::Cube, true), + _ => return None, + }; + Some(Type { + name: None, + inner: TypeInner::Image { + dim, + arrayed, + class: ImageClass::Depth { multi: false }, + }, + }) + }; + + if let Some(ty) = shadow_parse(word) { + return Some(ty); + } + + let (kind_prefix, rest) = if let Some(rest) = word.strip_prefix("sampler") { + ("", rest) + } else if let Some(rest) = word.strip_prefix("isampler") { + ("i", rest) + } else if let Some(rest) = word.strip_prefix("usampler") { + ("u", rest) + } else { + return None; + }; + + let kind = sampler_kind(kind_prefix)?; + let sampled = |multi| ImageClass::Sampled { kind, multi }; + + let (dim, arrayed, class) = match rest { + "1D" => (ImageDimension::D1, false, sampled(false)), + "1DArray" => (ImageDimension::D1, true, sampled(false)), + "2D" => (ImageDimension::D2, false, sampled(false)), + "2DArray" => (ImageDimension::D2, true, sampled(false)), + "2DMS" => (ImageDimension::D2, false, sampled(true)), + "2DMSArray" => (ImageDimension::D2, true, sampled(true)), + "3D" => (ImageDimension::D3, false, sampled(false)), + "Cube" => (ImageDimension::Cube, false, sampled(false)), + "CubeArray" => (ImageDimension::Cube, true, sampled(false)), + _ => return None, + }; + + Some(Type { + name: None, + inner: TypeInner::Image { + dim, + arrayed, + class, + }, + }) + }; + vec_parse(word) .or_else(|| mat_parse(word)) .or_else(|| texture_parse(word)) .or_else(|| image_parse(word)) + .or_else(|| sampler_parse(word)) } } } +/// Returns `true` for GLSL combined image-sampler type names (`sampler2D`, +/// `isampler3D`, `sampler2DShadow`, …). Excludes the plain-sampler types +/// `sampler` and `samplerShadow`, which map to [`TypeInner::Sampler`], not +/// [`TypeInner::Image`]. +/// +/// Used by the lexer to emit [`CombinedSamplerTypeName`] tokens. +/// +/// [`CombinedSamplerTypeName`]: super::token::TokenValue::CombinedSamplerTypeName +pub fn is_combined_sampler(word: &str) -> bool { + if word == "sampler" || word == "samplerShadow" { + return false; + } + word.starts_with("sampler") || word.starts_with("isampler") || word.starts_with("usampler") +} + pub const fn scalar_components(ty: &TypeInner) -> Option { match *ty { TypeInner::Scalar(scalar) diff --git a/naga/src/front/glsl/variables.rs b/naga/src/front/glsl/variables.rs index 31d6411b785..fdec5d42d77 100644 --- a/naga/src/front/glsl/variables.rs +++ b/naga/src/front/glsl/variables.rs @@ -17,6 +17,11 @@ pub struct VarDeclaration<'a, 'key> { pub name: Option, pub init: Option>, pub meta: Span, + /// `true` when the GLSL source type name is a combined image-sampler type + /// (`sampler2D`, `isampler3D`, `sampler2DShadow`, …). When set, `add_global_var` + /// will synthesise a paired implicit sampler global at the same (set, binding) so + /// that `texture(u_tex, uv)` can be lowered to `textureSample(u_tex, u_tex_sampler, uv)`. + pub is_combined_sampler: bool, } /// Information about a builtin used in [`add_builtin`](Frontend::add_builtin). @@ -421,6 +426,7 @@ impl Frontend { name, init, meta, + is_combined_sampler, }: VarDeclaration, ) -> Result { let storage = qualifiers.storage.0; @@ -641,6 +647,72 @@ impl Frontend { meta, ); + // If this is a combined image-sampler uniform (`sampler2D`, `sampler2DShadow`, + // etc.), synthesise an implicit companion `sampler` global. The companion is + // recorded in `Frontend::combined_samplers` so that every `Context` built + // afterwards can populate its `samplers` map, which is what `texture_call` uses + // to emit `textureSample`. + // + // WGSL and naga validation require each (set, binding) pair to be unique even + // across different resource types. Therefore the implicit sampler is placed at a + // binding that is guaranteed not to collide: the next unused binding slot on the + // same set. + if is_combined_sampler && space == AddressSpace::Handle { + let comparison = matches!( + ctx.module.types[ty].inner, + TypeInner::Image { + class: crate::ImageClass::Depth { .. }, + .. + } + ); + + let sampler_ty = ctx.module.types.insert( + Type { + name: None, + inner: TypeInner::Sampler { comparison }, + }, + meta, + ); + + // Name the implicit sampler `_sampler` for debuggability. + let sampler_name = name.as_deref().map(|n| alloc::format!("{n}_sampler")); + + // Compute a unique sampler binding: one past the highest binding already + // in use on the same descriptor set (or binding 0 if the set is empty). + let sampler_binding = if let Some(ref image_rb) = binding { + let group = image_rb.group; + let next_binding = ctx + .module + .global_variables + .iter() + .filter_map(|(_, var)| var.binding.as_ref()) + .filter(|rb| rb.group == group) + .map(|rb| rb.binding + 1) + .max() + .unwrap_or(0); + Some(ResourceBinding { + group, + binding: next_binding, + }) + } else { + None + }; + + let sampler_handle = ctx.module.global_variables.append( + GlobalVariable { + name: sampler_name, + space: AddressSpace::Handle, + binding: sampler_binding, + ty: sampler_ty, + init: None, + memory_decorations: crate::MemoryDecorations::empty(), + }, + meta, + ); + + self.combined_samplers.push((handle, sampler_handle)); + } + let lookup = GlobalLookup { kind: GlobalLookupKind::Variable(handle), entry_arg: None, diff --git a/naga/src/front/mod.rs b/naga/src/front/mod.rs index b2b1fbc93b3..2afa6f1be2a 100644 --- a/naga/src/front/mod.rs +++ b/naga/src/front/mod.rs @@ -2,7 +2,7 @@ Frontend parsers that consume binary and text shaders and load them into [`Module`](super::Module)s. */ -#[cfg(any(feature = "spv-in", feature = "wgsl-in"))] +#[cfg(any(feature = "glsl-in", feature = "spv-in", feature = "wgsl-in"))] pub(crate) mod interpolator; mod type_gen; diff --git a/naga/tests/in/glsl/image-compute.comp b/naga/tests/in/glsl/image-compute.comp new file mode 100644 index 00000000000..b8ed9bc4203 --- /dev/null +++ b/naga/tests/in/glsl/image-compute.comp @@ -0,0 +1,21 @@ +// Exercises imageLoad and imageStore in a compute shader context. +// Covers the readonly/writeonly/read-write access qualifiers on image2D uniforms. +#version 460 core + +layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in; + +layout(rgba8, set = 0, binding = 0) uniform readonly image2D src; +layout(rgba8, set = 0, binding = 1) uniform writeonly image2D dst; +layout(rgba8, set = 0, binding = 2) uniform image2D rw; + +void main() { + ivec2 pos = ivec2(gl_GlobalInvocationID.xy); + + // Read from a readonly image and write to a writeonly image. + vec4 val = imageLoad(src, pos); + imageStore(dst, pos, val); + + // Read-modify-write via a read-write image. + vec4 prev = imageLoad(rw, pos); + imageStore(rw, pos, prev * 0.5); +} diff --git a/naga/tests/in/glsl/sampler-combined-texture.frag b/naga/tests/in/glsl/sampler-combined-texture.frag new file mode 100644 index 00000000000..fd9d4561813 --- /dev/null +++ b/naga/tests/in/glsl/sampler-combined-texture.frag @@ -0,0 +1,20 @@ +#version 450 core + +// Exercises calling `texture()` directly on combined image-sampler uniforms +// (`sampler2D`, `sampler2DShadow`). Naga synthesises a paired implicit sampler +// for each such uniform so the call can be lowered to `textureSample` / +// `textureSampleCompare` without requiring the explicit constructor syntax +// `texture(sampler2D(u_tex, u_samp), uv)`. + +layout(set = 0, binding = 0) uniform sampler2D u_tex; +layout(set = 0, binding = 1) uniform sampler2DShadow u_shadow; + +layout(location = 0) in vec2 v_uv; +layout(location = 1) in vec3 v_shadow_coord; +layout(location = 0) out vec4 o_color; +layout(location = 1) out float o_shadow; + +void main() { + o_color = texture(u_tex, v_uv); + o_shadow = texture(u_shadow, v_shadow_coord); +} diff --git a/naga/tests/in/glsl/sampler-combined-types.frag b/naga/tests/in/glsl/sampler-combined-types.frag new file mode 100644 index 00000000000..0831be08d0b --- /dev/null +++ b/naga/tests/in/glsl/sampler-combined-types.frag @@ -0,0 +1,27 @@ +#version 450 core + +// Exercises combined image-sampler type names (`sampler2D` etc.) as direct +// uniform variable types with explicit descriptor-set / binding layout qualifiers. +// This regressed as NotImplemented("variable qualifier") before the fix. + +layout(set = 0, binding = 0) uniform sampler2D u_sampled; +layout(set = 0, binding = 1) uniform isampler2D u_isampled; +layout(set = 0, binding = 2) uniform usampler3D u_usampled_3d; +layout(set = 0, binding = 3) uniform samplerCube u_cube; +layout(set = 0, binding = 4) uniform sampler2DArray u_array; +layout(set = 0, binding = 5) uniform sampler2DShadow u_shadow; + +layout(set = 1, binding = 0) uniform sampler samp; + +void main() { + ivec2 coord2 = ivec2(gl_FragCoord.xy); + vec3 coord3 = vec3(gl_FragCoord.xy, 0.0); + + vec4 c0 = texelFetch(u_sampled, coord2, 0); + ivec4 c1 = texelFetch(u_isampled, coord2, 0); + uvec4 c2 = texelFetch(u_usampled_3d, ivec3(coord2, 0), 0); + vec4 c3 = texture(samplerCube(u_cube, samp), coord3); + vec4 c4 = texture(sampler2DArray(u_array, samp), coord3); + + gl_FragDepth = c0.r + float(c1.r) + float(c2.r) + c3.r + c4.r; +} diff --git a/naga/tests/naga/glsl_errors.rs b/naga/tests/naga/glsl_errors.rs new file mode 100644 index 00000000000..67500961395 --- /dev/null +++ b/naga/tests/naga/glsl_errors.rs @@ -0,0 +1,100 @@ +//! Tests for known GLSL front-end limitations and error cases. +//! +//! Complements the snapshot tests in `snapshots.rs` by documenting parse/compile +//! failures that are expected given the current state of the GLSL frontend. + +#![cfg(feature = "glsl-in")] + +use naga::front::glsl::{Frontend, Options}; + +fn parse_frag(source: &str) -> Result { + Frontend::default().parse( + &Options { + stage: naga::ShaderStage::Fragment, + defines: Default::default(), + }, + source, + ) +} + +fn parse_comp(source: &str) -> Result { + Frontend::default().parse( + &Options { + stage: naga::ShaderStage::Compute, + defines: Default::default(), + }, + source, + ) +} + +/// `texture(sampler2D_uniform, uv)` — calling `texture()` directly on a `sampler2D` +/// uniform that was declared as a combined image+sampler type — is now supported. +/// +/// When a `sampler2D` uniform is declared, naga synthesises a paired implicit sampler +/// global at the same binding. `texture(u_tex, uv)` is then lowered to +/// `textureSample(u_tex, u_tex_sampler, uv)` using that implicit sampler. +/// +/// The explicit constructor syntax `texture(sampler2D(u_tex, u_samp), uv)` with +/// separate `texture2D` + `sampler` uniforms (as in `tests/in/glsl/samplers.frag`) +/// continues to work unchanged. +/// +/// The snapshot for this shader lives in `tests/in/glsl/sampler-combined-texture.frag`. +#[test] +fn texture_call_on_combined_sampler_uniform_is_unsupported() { + let src = r#" + #version 450 + layout(set = 0, binding = 0) uniform sampler2D u_tex; + layout(location = 0) in vec2 v_uv; + layout(location = 0) out vec4 o_color; + void main() { + o_color = texture(u_tex, v_uv); + } + "#; + + assert!( + parse_frag(src).is_ok(), + "expected texture(sampler2D_uniform, uv) to compile successfully" + ); +} + +/// Shadow sampler variant: `texture(sampler2DShadow_uniform, coord)` is now supported. +/// +/// When a `sampler2DShadow` uniform is declared, naga synthesises a paired implicit +/// comparison sampler at the same binding. `texture(u_shadow, coord)` is then lowered +/// to `textureSampleCompare(u_shadow, u_shadow_sampler, coord.xy, coord.z)`. +#[test] +fn texture_call_on_shadow_sampler_uniform_is_unsupported() { + let src = r#" + #version 450 + layout(set = 0, binding = 0) uniform sampler2DShadow u_shadow; + layout(location = 0) in vec3 v_coord; + layout(location = 0) out float o_depth; + void main() { + o_depth = texture(u_shadow, v_coord); + } + "#; + + assert!( + parse_frag(src).is_ok(), + "expected texture(sampler2DShadow_uniform, coord) to compile successfully" + ); +} + +/// imageLoad and imageStore in a compute shader work correctly. +/// This is a compile-only smoke test; the full snapshot lives in +/// `tests/in/glsl/image-compute.comp`. +#[test] +fn image_load_store_compute_compiles() { + let src = r#" + #version 460 + layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in; + layout(rgba8, set = 0, binding = 0) uniform readonly image2D src; + layout(rgba8, set = 0, binding = 1) uniform writeonly image2D dst; + void main() { + ivec2 pos = ivec2(gl_GlobalInvocationID.xy); + imageStore(dst, pos, imageLoad(src, pos)); + } + "#; + + parse_comp(src).expect("imageLoad/imageStore in compute should compile"); +} diff --git a/naga/tests/naga/main.rs b/naga/tests/naga/main.rs index 208de772d14..77450e3c571 100644 --- a/naga/tests/naga/main.rs +++ b/naga/tests/naga/main.rs @@ -1,4 +1,5 @@ mod example_wgsl; +mod glsl_errors; mod snapshots; mod spirv_capabilities; mod validation; diff --git a/naga/tests/out/wgsl/glsl-image-compute.comp.wgsl b/naga/tests/out/wgsl/glsl-image-compute.comp.wgsl new file mode 100644 index 00000000000..c2f80b08fc3 --- /dev/null +++ b/naga/tests/out/wgsl/glsl-image-compute.comp.wgsl @@ -0,0 +1,36 @@ +@group(0) @binding(0) +var src: texture_storage_2d; +@group(0) @binding(1) +var dst: texture_storage_2d; +@group(0) @binding(2) +var rw: texture_storage_2d; +var gl_GlobalInvocationID_1: vec3; + +fn main_1() { + var pos: vec2; + var val: vec4; + var prev: vec4; + + let _e4 = gl_GlobalInvocationID_1; + pos = vec2(_e4.xy); + let _e8 = pos; + let _e9 = textureLoad(src, _e8); + val = _e9; + let _e11 = pos; + let _e12 = val; + textureStore(dst, _e11, _e12); + let _e13 = pos; + let _e14 = textureLoad(rw, _e13); + prev = _e14; + let _e16 = pos; + let _e17 = prev; + textureStore(rw, _e16, (_e17 * 0.5f)); + return; +} + +@compute @workgroup_size(8, 8, 1) +fn main(@builtin(global_invocation_id) gl_GlobalInvocationID: vec3) { + gl_GlobalInvocationID_1 = gl_GlobalInvocationID; + main_1(); + return; +} diff --git a/naga/tests/out/wgsl/glsl-sampler-combined-texture.frag.wgsl b/naga/tests/out/wgsl/glsl-sampler-combined-texture.frag.wgsl new file mode 100644 index 00000000000..9da705b341f --- /dev/null +++ b/naga/tests/out/wgsl/glsl-sampler-combined-texture.frag.wgsl @@ -0,0 +1,37 @@ +struct FragmentOutput { + @location(0) o_color: vec4, + @location(1) o_shadow: f32, +} + +@group(0) @binding(0) +var u_tex: texture_2d; +@group(0) @binding(2) +var u_tex_sampler: sampler; +@group(0) @binding(1) +var u_shadow: texture_depth_2d; +@group(0) @binding(3) +var u_shadow_sampler: sampler_comparison; +var v_uv_1: vec2; +var v_shadow_coord_1: vec3; +var o_color: vec4; +var o_shadow: f32; + +fn main_1() { + let _e8 = v_uv_1; + let _e9 = textureSample(u_tex, u_tex_sampler, _e8); + o_color = _e9; + let _e10 = v_shadow_coord_1; + let _e13 = textureSampleCompare(u_shadow, u_shadow_sampler, _e10.xy, _e10.z); + o_shadow = _e13; + return; +} + +@fragment +fn main(@location(0) v_uv: vec2, @location(1) v_shadow_coord: vec3) -> FragmentOutput { + v_uv_1 = v_uv; + v_shadow_coord_1 = v_shadow_coord; + main_1(); + let _e5 = o_color; + let _e7 = o_shadow; + return FragmentOutput(_e5, _e7); +} diff --git a/naga/tests/out/wgsl/glsl-sampler-combined-types.frag.wgsl b/naga/tests/out/wgsl/glsl-sampler-combined-types.frag.wgsl new file mode 100644 index 00000000000..7de464ec1fd --- /dev/null +++ b/naga/tests/out/wgsl/glsl-sampler-combined-types.frag.wgsl @@ -0,0 +1,64 @@ +struct FragmentOutput { + @builtin(frag_depth) gl_FragDepth: f32, +} + +@group(0) @binding(0) +var u_sampled: texture_2d; +@group(0) @binding(1) +var u_isampled: texture_2d; +@group(0) @binding(2) +var u_usampled_3d: texture_3d; +@group(0) @binding(3) +var u_cube: texture_cube; +@group(0) @binding(4) +var u_array: texture_2d_array; +@group(1) @binding(0) +var samp: sampler; +var gl_FragCoord_1: vec4; +var gl_FragDepth: f32; + +fn main_1() { + var coord2_: vec2; + var coord3_: vec3; + var c0_: vec4; + var c1_: vec4; + var c2_: vec4; + var c3_: vec4; + var c4_: vec4; + + let _e7 = gl_FragCoord_1; + coord2_ = vec2(_e7.xy); + let _e11 = gl_FragCoord_1; + let _e12 = _e11.xy; + coord3_ = vec3(_e12.x, _e12.y, 0f); + let _e18 = coord2_; + let _e20 = textureLoad(u_sampled, _e18, 0i); + c0_ = _e20; + let _e22 = coord2_; + let _e24 = textureLoad(u_isampled, _e22, 0i); + c1_ = _e24; + let _e26 = coord2_; + let _e32 = textureLoad(u_usampled_3d, vec3(_e26.x, _e26.y, 0i), 0i); + c2_ = _e32; + let _e34 = coord3_; + let _e35 = textureSample(u_cube, samp, _e34); + c3_ = _e35; + let _e37 = coord3_; + let _e41 = textureSample(u_array, samp, _e37.xy, i32(_e37.z)); + c4_ = _e41; + let _e44 = c0_; + let _e46 = c1_; + let _e50 = c2_; + let _e54 = c3_; + let _e57 = c4_; + gl_FragDepth = ((((_e44.x + f32(_e46.x)) + f32(_e50.x)) + _e54.x) + _e57.x); + return; +} + +@fragment +fn main(@builtin(position) gl_FragCoord: vec4) -> FragmentOutput { + gl_FragCoord_1 = gl_FragCoord; + main_1(); + let _e3 = gl_FragDepth; + return FragmentOutput(_e3); +}