From 43a4d53107e7a2ea89707af31620264a58febdc0 Mon Sep 17 00:00:00 2001 From: Jamie Nicol Date: Tue, 22 Jul 2025 22:38:32 +0100 Subject: [PATCH] [naga wgsl-in wgsl-out] WGSL support for texture_external texture type (#7822) * [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 --- naga-cli/src/bin/naga.rs | 4 +- naga/src/back/glsl/features.rs | 3 +- naga/src/back/glsl/mod.rs | 5 + naga/src/back/hlsl/help.rs | 3 + naga/src/back/msl/writer.rs | 2 + naga/src/back/spv/image.rs | 1 + naga/src/back/spv/mod.rs | 1 + naga/src/back/spv/writer.rs | 1 + naga/src/common/wgsl/types.rs | 3 + naga/src/front/glsl/builtins.rs | 2 + naga/src/front/wgsl/lower/mod.rs | 9 +- naga/src/front/wgsl/parse/mod.rs | 6 + naga/src/ir/mod.rs | 2 + naga/src/proc/mod.rs | 2 + naga/src/proc/typifier.rs | 4 + naga/src/valid/expression.rs | 3 +- naga/src/valid/mod.rs | 2 + naga/src/valid/type.rs | 10 + naga/tests/in/wgsl/texture-external.toml | 2 + naga/tests/in/wgsl/texture-external.wgsl | 27 ++ .../out/ir/wgsl-texture-external.compact.ron | 320 ++++++++++++++++++ naga/tests/out/ir/wgsl-texture-external.ron | 320 ++++++++++++++++++ .../tests/out/wgsl/wgsl-texture-external.wgsl | 39 +++ wgpu-core/src/device/mod.rs | 4 + wgpu-core/src/validation.rs | 6 + 25 files changed, 774 insertions(+), 7 deletions(-) create mode 100644 naga/tests/in/wgsl/texture-external.toml create mode 100644 naga/tests/in/wgsl/texture-external.wgsl create mode 100644 naga/tests/out/ir/wgsl-texture-external.compact.ron create mode 100644 naga/tests/out/ir/wgsl-texture-external.ron create mode 100644 naga/tests/out/wgsl/wgsl-texture-external.wgsl diff --git a/naga-cli/src/bin/naga.rs b/naga-cli/src/bin/naga.rs index 591ea0d9d7..b018058613 100644 --- 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 index 1eb3df65f3..a6dfe4e310 100644 --- a/naga/src/back/glsl/features.rs +++ b/naga/src/back/glsl/features.rs @@ -421,7 +421,8 @@ impl 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 index 75917fb76f..e78af74c84 100644 --- 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 index f5c9d4b3b9..bea917547b 100644 --- a/naga/src/back/hlsl/help.rs +++ b/naga/src/back/hlsl/help.rs @@ -195,6 +195,7 @@ impl 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 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 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 index ce83b08dc4..2525855cd7 100644 --- 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 "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 index faac982872..3aec1333f0 100644 --- 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 diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index 2dcd95957d..21c0001547 100644 --- 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 index ff6a9c1813..0688eb6c97 100644 --- 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 index 1c20d96724..82b8eeaa67 100644 --- 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 index 6dcddda44e..3d7588d27b 100644 --- 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 index a59bdad1e1..c8545ca654 100644 --- 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 index 892930b430..cf4dd4d4bb 100644 --- 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 index 125ba818c3..0101954849 100644 --- 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 index 0843e709b5..f2584c64b3 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -383,6 +383,7 @@ impl super::ImageClass { match self { crate::ImageClass::Sampled { multi, .. } | crate::ImageClass::Depth { multi } => multi, crate::ImageClass::Storage { .. } => false, + crate::ImageClass::External => false, } } @@ -390,6 +391,7 @@ impl super::ImageClass { 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 index d675a898d9..79b4f95e10 100644 --- 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 index a2675d8c71..a34f3aa9ec 100644 --- 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 index cee5b7420d..6b4b514c6c 100644 --- 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 index b3ae13b7d4..095fd7f882 100644 --- 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 new file mode 100644 index 0000000000..f8f46d4223 --- /dev/null +++ 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 new file mode 100644 index 0000000000..086551ab44 --- /dev/null +++ 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 { + 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 { + return test(tex); +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + 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 new file mode 100644 index 0000000000..4c0876618d --- /dev/null +++ b/naga/tests/out/ir/wgsl-texture-external.compact.ron @@ -0,0 +1,320 @@ +( + types: [ + ( + name: None, + inner: Image( + dim: D2, + arrayed: false, + class: External, + ), + ), + ( + name: None, + inner: Sampler( + comparison: false, + ), + ), + ( + name: None, + inner: Vector( + size: Quad, + scalar: ( + kind: Float, + width: 4, + ), + ), + ), + ( + name: None, + inner: Vector( + size: Bi, + scalar: ( + kind: Uint, + width: 4, + ), + ), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + ray_vertex_return: None, + predeclared_types: {}, + ), + constants: [], + overrides: [], + global_variables: [ + ( + name: Some("tex"), + space: Handle, + binding: Some(( + group: 0, + binding: 0, + )), + ty: 0, + init: None, + ), + ( + name: Some("samp"), + space: Handle, + binding: Some(( + group: 0, + binding: 1, + )), + ty: 1, + init: None, + ), + ], + global_expressions: [], + functions: [ + ( + name: Some("test"), + arguments: [ + ( + name: Some("t"), + ty: 0, + binding: None, + ), + ], + result: Some(( + ty: 2, + binding: None, + )), + local_variables: [ + ( + name: Some("a"), + ty: 2, + init: None, + ), + ( + name: Some("b"), + ty: 2, + init: None, + ), + ( + name: Some("c"), + ty: 3, + init: None, + ), + ], + expressions: [ + FunctionArgument(0), + GlobalVariable(1), + Literal(F32(0.0)), + Splat( + size: Bi, + value: 2, + ), + ImageSample( + image: 0, + sampler: 1, + gather: None, + coordinate: 3, + array_index: None, + offset: None, + level: Zero, + depth_ref: None, + clamp_to_edge: true, + ), + LocalVariable(0), + Literal(U32(0)), + Splat( + size: Bi, + value: 6, + ), + ImageLoad( + image: 0, + coordinate: 7, + array_index: None, + sample: None, + level: None, + ), + LocalVariable(1), + ImageQuery( + image: 0, + query: Size( + level: None, + ), + ), + LocalVariable(2), + Load( + pointer: 5, + ), + Load( + pointer: 9, + ), + Binary( + op: Add, + left: 12, + right: 13, + ), + Load( + pointer: 11, + ), + As( + expr: 15, + kind: Float, + convert: Some(4), + ), + Swizzle( + size: Quad, + vector: 16, + pattern: (X, Y, X, Y), + ), + Binary( + op: Add, + left: 14, + right: 17, + ), + ], + named_expressions: { + 0: "t", + }, + body: [ + Emit(( + start: 3, + end: 5, + )), + Store( + pointer: 5, + value: 4, + ), + Emit(( + start: 7, + end: 9, + )), + Store( + pointer: 9, + value: 8, + ), + Emit(( + start: 10, + end: 11, + )), + Store( + pointer: 11, + value: 10, + ), + Emit(( + start: 12, + end: 19, + )), + Return( + value: Some(18), + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "fragment_main", + stage: Fragment, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("fragment_main"), + arguments: [], + result: Some(( + ty: 2, + binding: Some(Location( + location: 0, + interpolation: Some(Perspective), + sampling: Some(Center), + blend_src: None, + )), + )), + local_variables: [], + expressions: [ + GlobalVariable(0), + CallResult(0), + ], + named_expressions: {}, + body: [ + Call( + function: 0, + arguments: [ + 0, + ], + result: Some(1), + ), + Return( + value: Some(1), + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ( + name: "vertex_main", + stage: Vertex, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("vertex_main"), + arguments: [], + result: Some(( + ty: 2, + binding: Some(BuiltIn(Position( + invariant: false, + ))), + )), + local_variables: [], + expressions: [ + GlobalVariable(0), + CallResult(0), + ], + named_expressions: {}, + body: [ + Call( + function: 0, + arguments: [ + 0, + ], + result: Some(1), + ), + Return( + value: Some(1), + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ( + name: "compute_main", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("compute_main"), + arguments: [], + result: None, + local_variables: [], + expressions: [ + GlobalVariable(0), + CallResult(0), + ], + named_expressions: {}, + body: [ + Call( + function: 0, + arguments: [ + 0, + ], + result: Some(1), + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/ir/wgsl-texture-external.ron b/naga/tests/out/ir/wgsl-texture-external.ron new file mode 100644 index 0000000000..4c0876618d --- /dev/null +++ b/naga/tests/out/ir/wgsl-texture-external.ron @@ -0,0 +1,320 @@ +( + types: [ + ( + name: None, + inner: Image( + dim: D2, + arrayed: false, + class: External, + ), + ), + ( + name: None, + inner: Sampler( + comparison: false, + ), + ), + ( + name: None, + inner: Vector( + size: Quad, + scalar: ( + kind: Float, + width: 4, + ), + ), + ), + ( + name: None, + inner: Vector( + size: Bi, + scalar: ( + kind: Uint, + width: 4, + ), + ), + ), + ], + special_types: ( + ray_desc: None, + ray_intersection: None, + ray_vertex_return: None, + predeclared_types: {}, + ), + constants: [], + overrides: [], + global_variables: [ + ( + name: Some("tex"), + space: Handle, + binding: Some(( + group: 0, + binding: 0, + )), + ty: 0, + init: None, + ), + ( + name: Some("samp"), + space: Handle, + binding: Some(( + group: 0, + binding: 1, + )), + ty: 1, + init: None, + ), + ], + global_expressions: [], + functions: [ + ( + name: Some("test"), + arguments: [ + ( + name: Some("t"), + ty: 0, + binding: None, + ), + ], + result: Some(( + ty: 2, + binding: None, + )), + local_variables: [ + ( + name: Some("a"), + ty: 2, + init: None, + ), + ( + name: Some("b"), + ty: 2, + init: None, + ), + ( + name: Some("c"), + ty: 3, + init: None, + ), + ], + expressions: [ + FunctionArgument(0), + GlobalVariable(1), + Literal(F32(0.0)), + Splat( + size: Bi, + value: 2, + ), + ImageSample( + image: 0, + sampler: 1, + gather: None, + coordinate: 3, + array_index: None, + offset: None, + level: Zero, + depth_ref: None, + clamp_to_edge: true, + ), + LocalVariable(0), + Literal(U32(0)), + Splat( + size: Bi, + value: 6, + ), + ImageLoad( + image: 0, + coordinate: 7, + array_index: None, + sample: None, + level: None, + ), + LocalVariable(1), + ImageQuery( + image: 0, + query: Size( + level: None, + ), + ), + LocalVariable(2), + Load( + pointer: 5, + ), + Load( + pointer: 9, + ), + Binary( + op: Add, + left: 12, + right: 13, + ), + Load( + pointer: 11, + ), + As( + expr: 15, + kind: Float, + convert: Some(4), + ), + Swizzle( + size: Quad, + vector: 16, + pattern: (X, Y, X, Y), + ), + Binary( + op: Add, + left: 14, + right: 17, + ), + ], + named_expressions: { + 0: "t", + }, + body: [ + Emit(( + start: 3, + end: 5, + )), + Store( + pointer: 5, + value: 4, + ), + Emit(( + start: 7, + end: 9, + )), + Store( + pointer: 9, + value: 8, + ), + Emit(( + start: 10, + end: 11, + )), + Store( + pointer: 11, + value: 10, + ), + Emit(( + start: 12, + end: 19, + )), + Return( + value: Some(18), + ), + ], + diagnostic_filter_leaf: None, + ), + ], + entry_points: [ + ( + name: "fragment_main", + stage: Fragment, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("fragment_main"), + arguments: [], + result: Some(( + ty: 2, + binding: Some(Location( + location: 0, + interpolation: Some(Perspective), + sampling: Some(Center), + blend_src: None, + )), + )), + local_variables: [], + expressions: [ + GlobalVariable(0), + CallResult(0), + ], + named_expressions: {}, + body: [ + Call( + function: 0, + arguments: [ + 0, + ], + result: Some(1), + ), + Return( + value: Some(1), + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ( + name: "vertex_main", + stage: Vertex, + early_depth_test: None, + workgroup_size: (0, 0, 0), + workgroup_size_overrides: None, + function: ( + name: Some("vertex_main"), + arguments: [], + result: Some(( + ty: 2, + binding: Some(BuiltIn(Position( + invariant: false, + ))), + )), + local_variables: [], + expressions: [ + GlobalVariable(0), + CallResult(0), + ], + named_expressions: {}, + body: [ + Call( + function: 0, + arguments: [ + 0, + ], + result: Some(1), + ), + Return( + value: Some(1), + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ( + name: "compute_main", + stage: Compute, + early_depth_test: None, + workgroup_size: (1, 1, 1), + workgroup_size_overrides: None, + function: ( + name: Some("compute_main"), + arguments: [], + result: None, + local_variables: [], + expressions: [ + GlobalVariable(0), + CallResult(0), + ], + named_expressions: {}, + body: [ + Call( + function: 0, + arguments: [ + 0, + ], + result: Some(1), + ), + Return( + value: None, + ), + ], + diagnostic_filter_leaf: None, + ), + ), + ], + diagnostic_filters: [], + diagnostic_filter_leaf: None, + doc_comments: None, +) \ No newline at end of file diff --git a/naga/tests/out/wgsl/wgsl-texture-external.wgsl b/naga/tests/out/wgsl/wgsl-texture-external.wgsl new file mode 100644 index 0000000000..d165fa2590 --- /dev/null +++ b/naga/tests/out/wgsl/wgsl-texture-external.wgsl @@ -0,0 +1,39 @@ +@group(0) @binding(0) +var tex: texture_external; +@group(0) @binding(1) +var samp: sampler; + +fn test(t: texture_external) -> vec4 { + var a: vec4; + var b: vec4; + var c: vec2; + + let _e4 = textureSampleBaseClampToEdge(t, samp, vec2(0f)); + a = _e4; + let _e8 = textureLoad(t, vec2(0u)); + b = _e8; + let _e10 = textureDimensions(t); + c = _e10; + let _e12 = a; + let _e13 = b; + let _e15 = c; + return ((_e12 + _e13) + vec2(_e15).xyxy); +} + +@fragment +fn fragment_main() -> @location(0) vec4 { + let _e1 = test(tex); + return _e1; +} + +@vertex +fn vertex_main() -> @builtin(position) vec4 { + let _e1 = test(tex); + return _e1; +} + +@compute @workgroup_size(1, 1, 1) +fn compute_main() { + let _e1 = test(tex); + return; +} diff --git a/wgpu-core/src/device/mod.rs b/wgpu-core/src/device/mod.rs index ad60e106d6..a57714900b 100644 --- a/wgpu-core/src/device/mod.rs +++ b/wgpu-core/src/device/mod.rs @@ -494,6 +494,10 @@ pub fn create_validator( Caps::RAY_HIT_VERTEX_POSITION, features.intersects(wgt::Features::EXPERIMENTAL_RAY_HIT_VERTEX_RETURN), ); + caps.set( + Caps::TEXTURE_EXTERNAL, + features.intersects(wgt::Features::EXTERNAL_TEXTURE), + ); naga::valid::Validator::new(flags, caps) } diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs index f720ef68b4..5bddd8a5b0 100644 --- a/wgpu-core/src/validation.rs +++ b/wgpu-core/src/validation.rs @@ -46,6 +46,10 @@ impl From<&ResourceType> for BindingTypeName { fn from(ty: &ResourceType) -> BindingTypeName { match ty { ResourceType::Buffer { .. } => BindingTypeName::Buffer, + ResourceType::Texture { + class: naga::ImageClass::External, + .. + } => BindingTypeName::ExternalTexture, ResourceType::Texture { .. } => BindingTypeName::Texture, ResourceType::Sampler { .. } => BindingTypeName::Sampler, ResourceType::AccelerationStructure { .. } => BindingTypeName::AccelerationStructure, @@ -594,6 +598,7 @@ impl Resource { access: naga_access, } } + BindingType::ExternalTexture => naga::ImageClass::External, _ => { return Err(BindingError::WrongType { binding: (&entry.ty).into(), @@ -701,6 +706,7 @@ impl Resource { f }, }, + naga::ImageClass::External => BindingType::ExternalTexture, } } ResourceType::AccelerationStructure { vertex_return } => {