Repository: gfx-rs/wgpu (language: rust)

Issue:
WGSL: Support the `readonly_and_readwrite_storage_textures` language extension
**Is your feature request related to a problem? Please describe.**

The WGSL spec. defines a language extension called `readonly_and_readwrite_storage_textures`, which we can and should support.

**Describe the solution you'd like**

Implement this feature from the spec. It's not required for adherence to the spec., but there is interesting code in the wild (i.e., by ThreeJS) that already wants to use this language extension.

**Describe alternatives you've considered**

1. Don't implement this. If we encounter an implementation blocker, or find other concerns with this feature, we might reasonably do this, but that's the only situation in which we'd be justified, I think.

**Additional context**

Related to #6192 in that it's expected that the `requires` directive is supported with the `readonly_and_readwrite_storage_textures` identifier specifically for this feature.



Discussion / hints:
Firefox tracking for this issue: [bug 1916464](https://bugzilla.mozilla.org/show_bug.cgi?id=1916464)
We actually already support this, we're just missing some validation.
There is one new built-in function that we are missing: `textureBarrier()`. Brought up by @mikialex in https://github.com/gfx-rs/wgpu/issues/6260.

Fix patch:
diff --git a/CHANGELOG.md b/CHANGELOG.md
index 634f037f7a..aff32fcf59 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -151,6 +151,7 @@ By @jamienicol in [#6929](https://github.com/gfx-rs/wgpu/pull/6929) and [#7080](
 - Support @must_use attribute on function declarations. By @turbocrime in [#6801](https://github.com/gfx-rs/wgpu/pull/6801).
 - Support for generating the candidate intersections from AABB geometry, and confirming the hits. By @kvark in [#7047](https://github.com/gfx-rs/wgpu/pull/7047).
 - Make naga::back::spv::Function::to_words write the OpFunctionEnd instruction in itself, instead of making another call after it. By @junjunjd in [#7156](https://github.com/gfx-rs/wgpu/pull/7156).
+- Add support for texture memory barriers. By @Devon7925 in [#7173](https://github.com/gfx-rs/wgpu/pull/7173).
 
 ### Changes
 
diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs
index a9c32d7043..42d88d0d43 100644
--- a/naga/src/back/glsl/mod.rs
+++ b/naga/src/back/glsl/mod.rs
@@ -4628,6 +4628,9 @@ impl<'a, W: Write> Writer<'a, W> {
         if flags.contains(crate::Barrier::SUB_GROUP) {
             writeln!(self.out, "{level}subgroupMemoryBarrier();")?;
         }
+        if flags.contains(crate::Barrier::TEXTURE) {
+            writeln!(self.out, "{level}memoryBarrierImage();")?;
+        }
         writeln!(self.out, "{level}barrier();")?;
         Ok(())
     }
diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs
index 55ffad2d8b..89f0ee6d02 100644
--- a/naga/src/back/hlsl/writer.rs
+++ b/naga/src/back/hlsl/writer.rs
@@ -4130,6 +4130,9 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
         if barrier.contains(crate::Barrier::SUB_GROUP) {
             // Does not exist in DirectX
         }
+        if barrier.contains(crate::Barrier::TEXTURE) {
+            writeln!(self.out, "{level}DeviceMemoryBarrierWithGroupSync();")?;
+        }
         Ok(())
     }
 }
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index 6fe3e9f0c8..70609a2f43 100644
--- a/naga/src/back/msl/writer.rs
+++ b/naga/src/back/msl/writer.rs
@@ -6513,6 +6513,12 @@ template <typename A>
                 "{level}{NAMESPACE}::simdgroup_barrier({NAMESPACE}::mem_flags::mem_threadgroup);",
             )?;
         }
+        if flags.contains(crate::Barrier::TEXTURE) {
+            writeln!(
+                self.out,
+                "{level}{NAMESPACE}::threadgroup_barrier({NAMESPACE}::mem_flags::mem_texture);",
+            )?;
+        }
         Ok(())
     }
 }
diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs
index f1ba07a386..6d4cedc57f 100644
--- a/naga/src/back/spv/writer.rs
+++ b/naga/src/back/spv/writer.rs
@@ -1571,6 +1571,10 @@ impl Writer {
             spirv::MemorySemantics::WORKGROUP_MEMORY,
             flags.contains(crate::Barrier::WORK_GROUP),
         );
+        semantics.set(
+            spirv::MemorySemantics::IMAGE_MEMORY,
+            flags.contains(crate::Barrier::TEXTURE),
+        );
         let exec_scope_id = if flags.contains(crate::Barrier::SUB_GROUP) {
             self.get_index_constant(spirv::Scope::Subgroup as u32)
         } else {
diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs
index a7cd8f95c9..f6aaaffa4f 100644
--- a/naga/src/back/wgsl/writer.rs
+++ b/naga/src/back/wgsl/writer.rs
@@ -973,6 +973,10 @@ impl<W: Write> Writer<W> {
                 if barrier.contains(crate::Barrier::SUB_GROUP) {
                     writeln!(self.out, "{level}subgroupBarrier();")?;
                 }
+
+                if barrier.contains(crate::Barrier::TEXTURE) {
+                    writeln!(self.out, "{level}textureBarrier();")?;
+                }
             }
             Statement::RayQuery { .. } => unreachable!(),
             Statement::SubgroupBallot { result, predicate } => {
diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs
index d34edf367e..fdb81379af 100644
--- a/naga/src/front/spv/mod.rs
+++ b/naga/src/front/spv/mod.rs
@@ -3846,6 +3846,10 @@ impl<I: Iterator<Item = u32>> Frontend<I> {
                                     .bits()
                                 != 0,
                         );
+                        flags.set(
+                            crate::Barrier::TEXTURE,
+                            semantics & spirv::MemorySemantics::IMAGE_MEMORY.bits() != 0,
+                        );
                         block.push(crate::Statement::Barrier(flags), span);
                     } else {
                         log::warn!("Unsupported barrier execution scope: {}", exec_scope);
diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs
index 7d8fbb7c80..60070284ec 100644
--- a/naga/src/front/wgsl/lower/mod.rs
+++ b/naga/src/front/wgsl/lower/mod.rs
@@ -2534,6 +2534,14 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
                                 .push(crate::Statement::Barrier(crate::Barrier::SUB_GROUP), span);
                             return Ok(None);
                         }
+                        "textureBarrier" => {
+                            ctx.prepare_args(arguments, 0, span).finish()?;
+
+                            let rctx = ctx.runtime_expression_ctx(span)?;
+                            rctx.block
+                                .push(crate::Statement::Barrier(crate::Barrier::TEXTURE), span);
+                            return Ok(None);
+                        }
                         "workgroupUniformLoad" => {
                             let mut args = ctx.prepare_args(arguments, 1, span);
                             let expr = args.next()?;
diff --git a/naga/src/lib.rs b/naga/src/lib.rs
index 2e917d34e0..b4d4c05f2b 100644
--- a/naga/src/lib.rs
+++ b/naga/src/lib.rs
@@ -1371,6 +1371,8 @@ bitflags::bitflags! {
         const WORK_GROUP = 1 << 1;
         /// Barrier synchronizes execution across all invocations within a subgroup that execute this instruction.
         const SUB_GROUP = 1 << 2;
+        /// Barrier synchronizes texture memory accesses in a workgroup.
+        const TEXTURE = 1 << 3;
     }
 }
 


Test patch:
diff --git a/naga/tests/in/wgsl/control-flow.toml b/naga/tests/in/wgsl/control-flow.toml
index c9811b10d1..02263c96e4 100644
--- a/naga/tests/in/wgsl/control-flow.toml
+++ b/naga/tests/in/wgsl/control-flow.toml
@@ -1,1 +1,4 @@
 targets = "SPIRV | METAL | GLSL | HLSL | WGSL"
+
+[msl]
+lang_version = [1, 2]
diff --git a/naga/tests/in/wgsl/control-flow.wgsl b/naga/tests/in/wgsl/control-flow.wgsl
index a25c899a44..564f1ab2a5 100644
--- a/naga/tests/in/wgsl/control-flow.wgsl
+++ b/naga/tests/in/wgsl/control-flow.wgsl
@@ -3,6 +3,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
     //TODO: execution-only barrier?
     storageBarrier();
     workgroupBarrier();
+    textureBarrier();
 
     var pos: i32;
     // switch without cases
diff --git a/naga/tests/in/wgsl/separate-entry-points.wgsl b/naga/tests/in/wgsl/separate-entry-points.wgsl
index a7ec3b083a..68a2adeae3 100644
--- a/naga/tests/in/wgsl/separate-entry-points.wgsl
+++ b/naga/tests/in/wgsl/separate-entry-points.wgsl
@@ -9,6 +9,7 @@ fn derivatives() {
 fn barriers() {
     storageBarrier();
     workgroupBarrier();
+    textureBarrier();
 }
 
 @fragment
diff --git a/naga/tests/out/glsl/control-flow.main.Compute.glsl b/naga/tests/out/glsl/control-flow.main.Compute.glsl
index 391fca84f4..48f1bffa30 100644
--- a/naga/tests/out/glsl/control-flow.main.Compute.glsl
+++ b/naga/tests/out/glsl/control-flow.main.Compute.glsl
@@ -149,6 +149,8 @@ void main() {
     barrier();
     memoryBarrierShared();
     barrier();
+    memoryBarrierImage();
+    barrier();
     do {
         pos = 1;
     } while(false);
diff --git a/naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl b/naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl
index 869b7ca418..2214046a4a 100644
--- a/naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl
+++ b/naga/tests/out/glsl/separate-entry-points.compute.Compute.glsl
@@ -11,6 +11,8 @@ void barriers() {
     barrier();
     memoryBarrierShared();
     barrier();
+    memoryBarrierImage();
+    barrier();
     return;
 }
 
diff --git a/naga/tests/out/hlsl/control-flow.hlsl b/naga/tests/out/hlsl/control-flow.hlsl
index ad36cc6201..19e4c8d2a2 100644
--- a/naga/tests/out/hlsl/control-flow.hlsl
+++ b/naga/tests/out/hlsl/control-flow.hlsl
@@ -193,6 +193,7 @@ void main(uint3 global_id : SV_DispatchThreadID)
 
     DeviceMemoryBarrierWithGroupSync();
     GroupMemoryBarrierWithGroupSync();
+    DeviceMemoryBarrierWithGroupSync();
     do {
         pos = int(1);
     } while(false);
diff --git a/naga/tests/out/msl/control-flow.msl b/naga/tests/out/msl/control-flow.msl
index 3c73656af7..eecf31fb9b 100644
--- a/naga/tests/out/msl/control-flow.msl
+++ b/naga/tests/out/msl/control-flow.msl
@@ -1,4 +1,4 @@
-// language: metal1.0
+// language: metal1.2
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -178,6 +178,7 @@ kernel void main_(
     int pos = {};
     metal::threadgroup_barrier(metal::mem_flags::mem_device);
     metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
+    metal::threadgroup_barrier(metal::mem_flags::mem_texture);
     switch(1) {
         default: {
             pos = 1;
diff --git a/naga/tests/out/spv/control-flow.spvasm b/naga/tests/out/spv/control-flow.spvasm
index 1fb3b52b1a..a2d6dc070e 100644
--- a/naga/tests/out/spv/control-flow.spvasm
+++ b/naga/tests/out/spv/control-flow.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 207
+; Bound: 208
 OpCapability Shader
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
@@ -37,6 +37,7 @@ OpDecorate %176 BuiltIn GlobalInvocationId
 %185 = OpConstant  %4  2
 %186 = OpConstant  %4  72
 %187 = OpConstant  %4  264
+%188 = OpConstant  %4  2056
 %8 = OpFunction  %2  None %9
 %7 = OpFunctionParameter  %5
 %6 = OpLabel
@@ -318,55 +319,56 @@ OpBranch %184
 %184 = OpLabel
 OpControlBarrier %185 %35 %186
 OpControlBarrier %185 %185 %187
-OpSelectionMerge %188 None
-OpSwitch %127 %189
-%189 = OpLabel
+OpControlBarrier %185 %185 %188
+OpSelectionMerge %189 None
+OpSwitch %127 %190
+%190 = OpLabel
 OpStore %182 %127
-OpBranch %188
-%188 = OpLabel
-%190 = OpLoad  %5  %182
-OpSelectionMerge %191 None
-OpSwitch %190 %196 1 %192 2 %193 3 %194 4 %194 5 %195 6 %196
-%192 = OpLabel
-OpStore %182 %16
-OpBranch %191
+OpBranch %189
+%189 = OpLabel
+%191 = OpLoad  %5  %182
+OpSelectionMerge %192 None
+OpSwitch %191 %197 1 %193 2 %194 3 %195 4 %195 5 %196 6 %197
 %193 = OpLabel
-OpStore %182 %127
-OpBranch %191
+OpStore %182 %16
+OpBranch %192
 %194 = OpLabel
-OpStore %182 %128
-OpBranch %191
+OpStore %182 %127
+OpBranch %192
 %195 = OpLabel
-OpStore %182 %180
-OpBranch %191
+OpStore %182 %128
+OpBranch %192
 %196 = OpLabel
+OpStore %182 %180
+OpBranch %192
+%197 = OpLabel
 OpStore %182 %181
-OpBranch %191
-%191 = OpLabel
-OpSelectionMerge %197 None
-OpSwitch %33 %199 0 %198
-%198 = OpLabel
-OpBranch %197
+OpBranch %192
+%192 = OpLabel
+OpSelectionMerge %198 None
+OpSwitch %33 %200 0 %199
 %199 = OpLabel
-OpBranch %197
-%197 = OpLabel
-%200 = OpLoad  %5  %182
-OpSelectionMerge %201 None
-OpSwitch %200 %206 1 %202 2 %203 3 %204 4 %205
-%202 = OpLabel
-OpStore %182 %16
-OpBranch %201
+OpBranch %198
+%200 = OpLabel
+OpBranch %198
+%198 = OpLabel
+%201 = OpLoad  %5  %182
+OpSelectionMerge %202 None
+OpSwitch %201 %207 1 %203 2 %204 3 %205 4 %206
 %203 = OpLabel
-OpStore %182 %127
-OpReturn
+OpStore %182 %16
+OpBranch %202
 %204 = OpLabel
-OpStore %182 %128
+OpStore %182 %127
 OpReturn
 %205 = OpLabel
+OpStore %182 %128
 OpReturn
 %206 = OpLabel
+OpReturn
+%207 = OpLabel
 OpStore %182 %180
 OpReturn
-%201 = OpLabel
+%202 = OpLabel
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/spv/separate-entry-points.compute.spvasm b/naga/tests/out/spv/separate-entry-points.compute.spvasm
index 38b7ea417e..a8faea48ea 100644
--- a/naga/tests/out/spv/separate-entry-points.compute.spvasm
+++ b/naga/tests/out/spv/separate-entry-points.compute.spvasm
@@ -1,12 +1,12 @@
 ; SPIR-V
 ; Version: 1.0
 ; Generator: rspirv
-; Bound: 18
+; Bound: 19
 OpCapability Shader
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint GLCompute %15 "compute"
-OpExecutionMode %15 LocalSize 1 1 1
+OpEntryPoint GLCompute %16 "compute"
+OpExecutionMode %16 LocalSize 1 1 1
 %2 = OpTypeVoid
 %4 = OpTypeFloat 32
 %3 = OpTypeVector %4 4
@@ -16,18 +16,20 @@ OpExecutionMode %15 LocalSize 1 1 1
 %11 = OpConstant  %10  1
 %12 = OpConstant  %10  72
 %13 = OpConstant  %10  264
+%14 = OpConstant  %10  2056
 %6 = OpFunction  %2  None %7
 %5 = OpLabel
 OpBranch %8
 %8 = OpLabel
 OpControlBarrier %9 %11 %12
 OpControlBarrier %9 %9 %13
+OpControlBarrier %9 %9 %14
 OpReturn
 OpFunctionEnd
-%15 = OpFunction  %2  None %7
-%14 = OpLabel
-OpBranch %16
-%16 = OpLabel
-%17 = OpFunctionCall  %2  %6
+%16 = OpFunction  %2  None %7
+%15 = OpLabel
+OpBranch %17
+%17 = OpLabel
+%18 = OpFunctionCall  %2  %6
 OpReturn
 OpFunctionEnd
\ No newline at end of file
diff --git a/naga/tests/out/wgsl/control-flow.wgsl b/naga/tests/out/wgsl/control-flow.wgsl
index ad071af58a..eaee5aa214 100644
--- a/naga/tests/out/wgsl/control-flow.wgsl
+++ b/naga/tests/out/wgsl/control-flow.wgsl
@@ -122,6 +122,7 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
 
     storageBarrier();
     workgroupBarrier();
+    textureBarrier();
     switch 1i {
         default: {
             pos = 1i;