[naga wgsl-in wgsl-out] WGSL support for texture_external texture type (#7822)

jamienicol · jimblandy · web-flow · commit 43a4d53107e7 · 2025-07-22T14:38:32.000-07:00
* [naga wgsl-in wgsl-out] WGSL support for texture_external texture type

Make wgsl-in correctly parse `texture_external` texture declarations,
and allow such textures to be used in `textureDimensions()`,
`textureSampleBaseClampToEdge()`, and `textureLoad()` function
calls. In IR these are represented by the `ImageClass::External` image
class, which is a 2D, non-multisampled, non-mipmapped, float-sampled
image.

Adds a new Capability `TEXTURE_EXTERNAL` and ensure validation rejects
shaders containing external textures if this capability flag is not
set. This capability is enabled for validation by wgpu devices which
support the `TEXTURE_EXTERNAL` feature (currently only when using the
noop backend), and by the Naga CLI when validating-only or when
outputting WGSL.

The WGSL backend can of course emit `ImageClass::External` images
directly as `texture_external` textures. Other backends are, for now,
unimplemented.

Lastly, we add a snapshot test covering all the valid uses of a
texture_external texture. These are:
  - As a global variable declaration
  - As an argument to the built-in functions `textureDimensions()`,
    `textureSampleBaseClampToEdge()`, and `textureLoad()`
  - As an argument to user-defined function declarations and calls.

We keep these in their own test so that we can control which targets
to run them against (currently WGSL and IR). When external textures
are supported by all Naga backends we can, if so inclined, integrate
these with existing texture tests.

* fixup! [naga wgsl-in wgsl-out] WGSL support for texture_external texture type

* fixup! [naga wgsl-in wgsl-out] WGSL support for texture_external texture type

---------

Co-authored-by: Jim Blandy &lt;jimb@red-bean.com&gt;
diff --git a/naga-cli/src/bin/naga.rs b/naga-cli/src/bin/naga.rs
@@ -538,8 +538,8 @@ fn run() -> anyhow::Result<()> {
                 use naga::valid::Capabilities as C;
                 let missing = match Path::new(path).extension().and_then(|ex| ex.to_str()) {
                     Some("wgsl") => C::CLIP_DISTANCE | C::CULL_DISTANCE,
-                    Some("metal") => C::CULL_DISTANCE,
-                    _ => C::empty(),
+                    Some("metal") => C::CULL_DISTANCE | C::TEXTURE_EXTERNAL,
+                    _ => C::TEXTURE_EXTERNAL,
                 };
                 caps & !missing
             });
diff --git a/naga/src/back/glsl/features.rs b/naga/src/back/glsl/features.rs
@@ -421,7 +421,8 @@ impl<W> Writer<'_, W> {
                             _ => {}
                         },
                         ImageClass::Sampled { multi: false, .. }
-                        | ImageClass::Depth { multi: false } => {}
+                        | ImageClass::Depth { multi: false }
+                        | ImageClass::External => {}
                     }
                 }
                 _ => {}
diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs
@@ -1176,6 +1176,7 @@ impl<'a, W: Write> Writer<'a, W> {
             Ic::Depth { multi: true } => ("sampler", float, "MS", ""),
             Ic::Depth { multi: false } => ("sampler", float, "", "Shadow"),
             Ic::Storage { format, .. } => ("image", format.into(), "", ""),
+            Ic::External => unimplemented!(),
         };
 
         let precision = if self.options.version.is_es() {
@@ -3302,6 +3303,7 @@ impl<'a, W: Write> Writer<'a, W> {
                                 write!(self.out, "imageSize(")?;
                                 self.write_expr(image, ctx)?;
                             }
+                            ImageClass::External => unimplemented!(),
                         }
                         write!(self.out, ")")?;
                         if components != 1 || self.options.version.is_es() {
@@ -3317,6 +3319,7 @@ impl<'a, W: Write> Writer<'a, W> {
                         let fun_name = match class {
                             ImageClass::Sampled { .. } | ImageClass::Depth { .. } => "textureSize",
                             ImageClass::Storage { .. } => "imageSize",
+                            ImageClass::External => unimplemented!(),
                         };
                         write!(self.out, "{fun_name}(")?;
                         self.write_expr(image, ctx)?;
@@ -3336,6 +3339,7 @@ impl<'a, W: Write> Writer<'a, W> {
                                 "textureSamples"
                             }
                             ImageClass::Storage { .. } => "imageSamples",
+                            ImageClass::External => unimplemented!(),
                         };
                         write!(self.out, "{fun_name}(")?;
                         self.write_expr(image, ctx)?;
@@ -4618,6 +4622,7 @@ impl<'a, W: Write> Writer<'a, W> {
                     "WGSL `textureLoad` from depth textures is not supported in GLSL".to_string(),
                 ))
             }
+            crate::ImageClass::External => unimplemented!(),
         };
 
         // openGL es doesn't have 1D images so we need workaround it
diff --git a/naga/src/back/hlsl/help.rs b/naga/src/back/hlsl/help.rs
@@ -195,6 +195,7 @@ impl<W: Write> super::Writer<'_, W> {
                 let storage_format_str = format.to_hlsl_str();
                 write!(self.out, "<{storage_format_str}>")?
             }
+            crate::ImageClass::External => unimplemented!(),
         }
         Ok(())
     }
@@ -290,6 +291,7 @@ impl<W: Write> super::Writer<'_, W> {
             crate::ImageClass::Depth { multi: false } => "Depth",
             crate::ImageClass::Sampled { multi: false, .. } => "",
             crate::ImageClass::Storage { .. } => "RW",
+            crate::ImageClass::External => unimplemented!(),
         };
         let arrayed_str = if query.arrayed { "Array" } else { "" };
         let query_str = match query.query {
@@ -349,6 +351,7 @@ impl<W: Write> super::Writer<'_, W> {
         let extra_coords = match wiq.class {
             crate::ImageClass::Storage { .. } => 0,
             crate::ImageClass::Sampled { .. } | crate::ImageClass::Depth { .. } => 1,
+            crate::ImageClass::External => unimplemented!(),
         };
 
         // GetDimensions Overloaded Methods
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
@@ -321,6 +321,7 @@ impl Display for TypeContext<'_> {
                         };
                         ("texture", "", format.into(), access)
                     }
+                    crate::ImageClass::External => unimplemented!(),
                 };
                 let base_name = scalar.to_msl_name();
                 let array_str = if arrayed { "_array" } else { "" };
@@ -6637,6 +6638,7 @@ template <typename A>
                                             "read-write textures".to_string(),
                                         ));
                                     }
+                                    crate::ImageClass::External => unimplemented!(),
                                 },
                                 _ => {
                                     return Err(Error::UnsupportedArrayOfType(base));
diff --git a/naga/src/back/spv/image.rs b/naga/src/back/spv/image.rs
@@ -118,6 +118,7 @@ impl Load {
             crate::ImageClass::Depth { .. } | crate::ImageClass::Sampled { .. } => {
                 spirv::Op::ImageFetch
             }
+            crate::ImageClass::External => unimplemented!(),
         };
 
         // `OpImageRead` and `OpImageFetch` instructions produce vec4<f32>
diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs
@@ -276,6 +276,7 @@ impl LocalImageType {
                 flags: make_flags(false, ImageTypeFlags::empty()),
                 image_format: format.into(),
             },
+            crate::ImageClass::External => unimplemented!(),
         }
     }
 }
diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs
@@ -1246,6 +1246,7 @@ impl Writer {
                         self.request_image_format_capabilities(format.into())?;
                         false
                     }
+                    crate::ImageClass::External => unimplemented!(),
                 };
 
                 match dim {
diff --git a/naga/src/common/wgsl/types.rs b/naga/src/common/wgsl/types.rs
@@ -250,6 +250,9 @@ where
                         "texture_storage_{dim_str}{arrayed_str}<{format_str}{access_str}>"
                     )?;
                 }
+                Ic::External => {
+                    write!(out, "texture_external")?;
+                }
             }
         }
         TypeInner::Scalar(scalar) => {
diff --git a/naga/src/front/glsl/builtins.rs b/naga/src/front/glsl/builtins.rs
@@ -2138,6 +2138,7 @@ impl Frontend {
                 ImageClass::Depth { .. } => (true, false),
                 ImageClass::Storage { .. } => (false, true),
                 ImageClass::Sampled { .. } => (false, false),
+                ImageClass::External => unreachable!(),
             };
 
             let coordinate = match (image_size, coord_size) {
@@ -2259,6 +2260,7 @@ pub fn sampled_to_depth(
                 kind: ErrorKind::SemanticError("Not a texture".into()),
                 meta,
             }),
+            ImageClass::External => unreachable!(),
         },
         _ => errors.push(Error {
             kind: ErrorKind::SemanticError("Not a texture".into()),
diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs
@@ -3587,9 +3587,12 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                         self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?
                     }
 
-                    // Sampling `Storage` textures isn't allowed at all. Let the
-                    // validator report the error.
-                    ir::ImageClass::Storage { .. } => self.expression(args.next()?, ctx)?,
+                    // Sampling `External` textures with a specified level isn't
+                    // allowed, and sampling `Storage` textures isn't allowed at
+                    // all. Let the validator report the error.
+                    ir::ImageClass::Storage { .. } | ir::ImageClass::External => {
+                        self.expression(args.next()?, ctx)?
+                    }
                 };
                 level = ir::SampleLevel::Exact(exact);
                 depth_ref = None;
diff --git a/naga/src/front/wgsl/parse/mod.rs b/naga/src/front/wgsl/parse/mod.rs
@@ -676,6 +676,7 @@ impl Parser {
             | "texture_depth_cube"
             | "texture_depth_cube_array"
             | "texture_depth_multisampled_2d"
+            | "texture_external"
             | "texture_storage_1d"
             | "texture_storage_1d_array"
             | "texture_storage_2d"
@@ -1867,6 +1868,11 @@ impl Parser {
                 arrayed: false,
                 class: crate::ImageClass::Depth { multi: true },
             },
+            "texture_external" => ast::Type::Image {
+                dim: crate::ImageDimension::D2,
+                arrayed: false,
+                class: crate::ImageClass::External,
+            },
             "texture_storage_1d" => {
                 let (format, access) = lexer.next_format_generic()?;
                 ast::Type::Image {
diff --git a/naga/src/ir/mod.rs b/naga/src/ir/mod.rs
@@ -652,6 +652,8 @@ pub enum ImageClass {
         /// Multi-sampled depth image.
         multi: bool,
     },
+    /// External texture.
+    External,
     /// Storage image.
     Storage {
         format: StorageFormat,
diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs
@@ -383,13 +383,15 @@ impl super::ImageClass {
         match self {
             crate::ImageClass::Sampled { multi, .. } | crate::ImageClass::Depth { multi } => multi,
             crate::ImageClass::Storage { .. } => false,
+            crate::ImageClass::External => false,
         }
     }
 
     pub const fn is_mipmapped(self) -> bool {
         match self {
             crate::ImageClass::Sampled { multi, .. } | crate::ImageClass::Depth { multi } => !multi,
             crate::ImageClass::Storage { .. } => false,
+            crate::ImageClass::External => false,
         }
     }
 
diff --git a/naga/src/proc/typifier.rs b/naga/src/proc/typifier.rs
@@ -512,6 +512,10 @@ impl<'a> ResolveContext<'a> {
                         scalar: format.into(),
                         size: crate::VectorSize::Quad,
                     },
+                    crate::ImageClass::External => Ti::Vector {
+                        scalar: crate::Scalar::F32,
+                        size: crate::VectorSize::Quad,
+                    },
                 }),
                 ref other => {
                     log::error!("Image type {other:?}");
diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs
@@ -460,6 +460,7 @@ impl super::Validator {
                         kind: crate::ScalarKind::Uint | crate::ScalarKind::Sint,
                         multi: false,
                     } if gather.is_some() => false,
+                    crate::ImageClass::External => false,
                     crate::ImageClass::Depth { multi: false } => true,
                     _ => return Err(ExpressionError::InvalidImageClass(class)),
                 };
@@ -551,7 +552,7 @@ impl super::Validator {
                         crate::ImageClass::Sampled {
                             kind: crate::ScalarKind::Float,
                             multi: false
-                        }
+                        } | crate::ImageClass::External
                     ) {
                         return Err(ExpressionError::InvalidSampleClampCoordinateToEdge(
                             alloc::format!("image class `{class:?}`"),
diff --git a/naga/src/valid/mod.rs b/naga/src/valid/mod.rs
@@ -168,6 +168,8 @@ bitflags::bitflags! {
         const RAY_HIT_VERTEX_POSITION = 1 << 25;
         /// Support for 16-bit floating-point types.
         const SHADER_FLOAT16 = 1 << 26;
+        /// Support for [`ImageClass::External`]
+        const TEXTURE_EXTERNAL = 1 << 27;
     }
 }
 
diff --git a/naga/src/valid/type.rs b/naga/src/valid/type.rs
@@ -732,6 +732,16 @@ impl super::Validator {
                 if arrayed && matches!(dim, crate::ImageDimension::Cube) {
                     self.require_type_capability(Capabilities::CUBE_ARRAY_TEXTURES)?;
                 }
+                if matches!(class, crate::ImageClass::External) {
+                    if dim != crate::ImageDimension::D2 || arrayed {
+                        return Err(TypeError::UnsupportedImageType {
+                            dim,
+                            arrayed,
+                            class,
+                        });
+                    }
+                    self.require_type_capability(Capabilities::TEXTURE_EXTERNAL)?;
+                }
                 TypeInfo::new(
                     TypeFlags::ARGUMENT | TypeFlags::CREATION_RESOLVED,
                     Alignment::ONE,
diff --git a/naga/tests/in/wgsl/texture-external.toml b/naga/tests/in/wgsl/texture-external.toml
@@ -0,0 +1,2 @@
+god_mode = true
+targets = "IR | WGSL"
diff --git a/naga/tests/in/wgsl/texture-external.wgsl b/naga/tests/in/wgsl/texture-external.wgsl
@@ -0,0 +1,27 @@
+@group(0) @binding(0)
+var tex: texture_external;
+@group(0) @binding(1)
+var samp: sampler;
+
+fn test(t: texture_external) -> vec4<f32> {
+  var a = textureSampleBaseClampToEdge(t, samp, vec2(0.0f));
+  var b = textureLoad(t, vec2(0u));
+  var c = textureDimensions(t);
+
+  return a + b + vec2f(c).xyxy;
+}
+
+@fragment
+fn fragment_main() -> @location(0) vec4<f32> {
+  return test(tex);
+}
+
+@vertex
+fn vertex_main() -> @builtin(position) vec4<f32> {
+  return test(tex);
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  test(tex);
+}
diff --git a/naga/tests/out/ir/wgsl-texture-external.compact.ron b/naga/tests/out/ir/wgsl-texture-external.compact.ron
diff --git a/naga/tests/out/ir/wgsl-texture-external.ron b/naga/tests/out/ir/wgsl-texture-external.ron
diff --git a/naga/tests/out/wgsl/wgsl-texture-external.wgsl b/naga/tests/out/wgsl/wgsl-texture-external.wgsl
diff --git a/wgpu-core/src/device/mod.rs b/wgpu-core/src/device/mod.rs
diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs

Original file line number	Diff line number	Diff line change
`@@ -421,7 +421,8 @@ impl<W> Writer<'_, W> {`
`421`	`421`	`_ => {}`
`422`	`422`	`},`
`423`	`423`	`ImageClass::Sampled { multi: false, .. }`
`424`		`- \| ImageClass::Depth { multi: false } => {}`
	`424`	`+ \| ImageClass::Depth { multi: false }`
	`425`	`+ \| ImageClass::External => {}`
`425`	`426`	`}`
`426`	`427`	`}`
`427`	`428`	`_ => {}`
Original file line number	Diff line number	Diff line change
`@@ -118,6 +118,7 @@ impl Load {`
`118`	`118`	`crate::ImageClass::Depth { .. } \| crate::ImageClass::Sampled { .. } => {`
`119`	`119`	`spirv::Op::ImageFetch`
`120`	`120`	`}`
	`121`	`+ crate::ImageClass::External => unimplemented!(),`
`121`	`122`	`};`
`122`	`123`
`123`	`124`	// `OpImageRead` and `OpImageFetch` instructions produce vec4<f32>
Original file line number	Diff line number	Diff line change
`@@ -276,6 +276,7 @@ impl LocalImageType {`
`276`	`276`	`flags: make_flags(false, ImageTypeFlags::empty()),`
`277`	`277`	`image_format: format.into(),`
`278`	`278`	`},`
	`279`	`+ crate::ImageClass::External => unimplemented!(),`
`279`	`280`	`}`
`280`	`281`	`}`
`281`	`282`	`}`
Original file line number	Diff line number	Diff line change
`@@ -1246,6 +1246,7 @@ impl Writer {`
`1246`	`1246`	`self.request_image_format_capabilities(format.into())?;`
`1247`	`1247`	`false`
`1248`	`1248`	`}`
	`1249`	`+ crate::ImageClass::External => unimplemented!(),`
`1249`	`1250`	`};`
`1250`	`1251`
`1251`	`1252`	`match dim {`
Original file line number	Diff line number	Diff line change
`@@ -250,6 +250,9 @@ where`
`250`	`250`	`"texture_storage_{dim_str}{arrayed_str}<{format_str}{access_str}>"`
`251`	`251`	`)?;`
`252`	`252`	`}`
	`253`	`+ Ic::External => {`
	`254`	`+ write!(out, "texture_external")?;`
	`255`	`+ }`
`253`	`256`	`}`
`254`	`257`	`}`
`255`	`258`	`TypeInner::Scalar(scalar) => {`