[tint] Add textureBarrier() builtin function Must only be called from compute stages and only if the chromium_experimental_read_write_storage_texture extension is enabled. Bug: tint:2007 Change-Id: I1790f65b5a4caebd8c30daff9efe3bf92d97cd12 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/146243 Reviewed-by: Ben Clayton <bclayton@google.com> Auto-Submit: James Price <jrprice@google.com> Kokoro: Kokoro <noreply+kokoro@google.com>

diff --git a/src/tint/lang/core/core.def b/src/tint/lang/core/core.def
index eabb9dc..f08b591 100644
--- a/src/tint/lang/core/core.def
+++ b/src/tint/lang/core/core.def

@@ -735,6 +735,7 @@
 @stage("compute") fn workgroupBarrier()
 @must_use @stage("compute") fn workgroupUniformLoad<T>(ptr<workgroup, T, read_write>) -> T
 
+@stage("compute") fn textureBarrier()
 @must_use fn textureDimensions<T: fiu32>(texture: texture_1d<T>) -> u32
 @must_use fn textureDimensions<T: fiu32, L: iu32>(texture: texture_1d<T>, level: L) -> u32
 @must_use fn textureDimensions<T: fiu32>(texture: texture_2d<T>) -> vec2<u32>

diff --git a/src/tint/lang/core/function.cc b/src/tint/lang/core/function.cc
index 959b6fd..3d34627 100644
--- a/src/tint/lang/core/function.cc
+++ b/src/tint/lang/core/function.cc

@@ -286,6 +286,9 @@
     if (name == "workgroupUniformLoad") {
         return Function::kWorkgroupUniformLoad;
     }
+    if (name == "textureBarrier") {
+        return Function::kTextureBarrier;
+    }
     if (name == "textureDimensions") {
         return Function::kTextureDimensions;
     }
@@ -551,6 +554,8 @@
             return "workgroupBarrier";
         case Function::kWorkgroupUniformLoad:
             return "workgroupUniformLoad";
+        case Function::kTextureBarrier:
+            return "textureBarrier";
         case Function::kTextureDimensions:
             return "textureDimensions";
         case Function::kTextureGather:
@@ -657,7 +662,8 @@
 }
 
 bool IsBarrierBuiltin(Function f) {
-    return f == Function::kWorkgroupBarrier || f == Function::kStorageBarrier;
+    return f == Function::kWorkgroupBarrier || f == Function::kStorageBarrier ||
+           f == Function::kTextureBarrier;
 }
 
 bool IsAtomicBuiltin(Function f) {

diff --git a/src/tint/lang/core/function.cc.tmpl b/src/tint/lang/core/function.cc.tmpl
index 5bceedc..32ed851 100644
--- a/src/tint/lang/core/function.cc.tmpl
+++ b/src/tint/lang/core/function.cc.tmpl

@@ -87,7 +87,8 @@
 }
 
 bool IsBarrierBuiltin(Function f) {
-    return f == Function::kWorkgroupBarrier || f == Function::kStorageBarrier;
+    return f == Function::kWorkgroupBarrier || f == Function::kStorageBarrier ||
+           f == Function::kTextureBarrier;
 }
 
 bool IsAtomicBuiltin(Function f) {

diff --git a/src/tint/lang/core/function.h b/src/tint/lang/core/function.h
index c7a3263..0daf5b0 100644
--- a/src/tint/lang/core/function.h
+++ b/src/tint/lang/core/function.h

@@ -119,6 +119,7 @@
     kUnpack4X8Unorm,
     kWorkgroupBarrier,
     kWorkgroupUniformLoad,
+    kTextureBarrier,
     kTextureDimensions,
     kTextureGather,
     kTextureGatherCompare,
@@ -256,6 +257,7 @@
     Function::kUnpack4X8Unorm,
     Function::kWorkgroupBarrier,
     Function::kWorkgroupUniformLoad,
+    Function::kTextureBarrier,
     Function::kTextureDimensions,
     Function::kTextureGather,
     Function::kTextureGatherCompare,
@@ -375,6 +377,7 @@
     "unpack4x8unorm",
     "workgroupBarrier",
     "workgroupUniformLoad",
+    "textureBarrier",
     "textureDimensions",
     "textureGather",
     "textureGatherCompare",

diff --git a/src/tint/lang/core/intrinsic/data/data.cc b/src/tint/lang/core/intrinsic/data/data.cc
index 547c83d..cbc2a3f 100644
--- a/src/tint/lang/core/intrinsic/data/data.cc
+++ b/src/tint/lang/core/intrinsic/data/data.cc

@@ -11478,6 +11478,12 @@
   },
   {
     /* [87] */
+    /* fn textureBarrier() */
+    /* num overloads */ 1,
+    /* overloads */ OverloadIndex(447),
+  },
+  {
+    /* [88] */
     /* fn textureDimensions<T : fiu32>(texture: texture_1d<T>) -> u32 */
     /* fn textureDimensions<T : fiu32, L : iu32>(texture: texture_1d<T>, level: L) -> u32 */
     /* fn textureDimensions<T : fiu32>(texture: texture_2d<T>) -> vec2<u32> */
@@ -11509,7 +11515,7 @@
     /* overloads */ OverloadIndex(0),
   },
   {
-    /* [88] */
+    /* [89] */
     /* fn textureGather<T : fiu32, C : iu32>(@const component: C, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>) -> vec4<T> */
     /* fn textureGather<T : fiu32, C : iu32>(@const component: C, texture: texture_2d<T>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<T> */
     /* fn textureGather<T : fiu32, C : iu32, A : iu32>(@const component: C, texture: texture_2d_array<T>, sampler: sampler, coords: vec2<f32>, array_index: A) -> vec4<T> */
@@ -11526,7 +11532,7 @@
     /* overloads */ OverloadIndex(93),
   },
   {
-    /* [89] */
+    /* [90] */
     /* fn textureGatherCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> vec4<f32> */
     /* fn textureGatherCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, @const offset: vec2<i32>) -> vec4<f32> */
     /* fn textureGatherCompare<A : iu32>(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: A, depth_ref: f32) -> vec4<f32> */
@@ -11537,7 +11543,7 @@
     /* overloads */ OverloadIndex(174),
   },
   {
-    /* [90] */
+    /* [91] */
     /* fn textureNumLayers<T : fiu32>(texture: texture_2d_array<T>) -> u32 */
     /* fn textureNumLayers<T : fiu32>(texture: texture_cube_array<T>) -> u32 */
     /* fn textureNumLayers(texture: texture_depth_2d_array) -> u32 */
@@ -11547,7 +11553,7 @@
     /* overloads */ OverloadIndex(246),
   },
   {
-    /* [91] */
+    /* [92] */
     /* fn textureNumLevels<T : fiu32>(texture: texture_1d<T>) -> u32 */
     /* fn textureNumLevels<T : fiu32>(texture: texture_2d<T>) -> u32 */
     /* fn textureNumLevels<T : fiu32>(texture: texture_2d_array<T>) -> u32 */
@@ -11562,14 +11568,14 @@
     /* overloads */ OverloadIndex(129),
   },
   {
-    /* [92] */
+    /* [93] */
     /* fn textureNumSamples<T : fiu32>(texture: texture_multisampled_2d<T>) -> u32 */
     /* fn textureNumSamples(texture: texture_depth_multisampled_2d) -> u32 */
     /* num overloads */ 2,
     /* overloads */ OverloadIndex(404),
   },
   {
-    /* [93] */
+    /* [94] */
     /* fn textureSample(texture: texture_1d<f32>, sampler: sampler, coords: f32) -> vec4<f32> */
     /* fn textureSample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32> */
     /* fn textureSample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32> */
@@ -11589,7 +11595,7 @@
     /* overloads */ OverloadIndex(64),
   },
   {
-    /* [94] */
+    /* [95] */
     /* fn textureSampleBias(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, bias: f32) -> vec4<f32> */
     /* fn textureSampleBias(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, bias: f32, @const offset: vec2<i32>) -> vec4<f32> */
     /* fn textureSampleBias<A : iu32>(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, bias: f32) -> vec4<f32> */
@@ -11602,7 +11608,7 @@
     /* overloads */ OverloadIndex(158),
   },
   {
-    /* [95] */
+    /* [96] */
     /* fn textureSampleCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> f32 */
     /* fn textureSampleCompare(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, @const offset: vec2<i32>) -> f32 */
     /* fn textureSampleCompare<A : iu32>(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: A, depth_ref: f32) -> f32 */
@@ -11613,7 +11619,7 @@
     /* overloads */ OverloadIndex(180),
   },
   {
-    /* [96] */
+    /* [97] */
     /* fn textureSampleCompareLevel(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32) -> f32 */
     /* fn textureSampleCompareLevel(texture: texture_depth_2d, sampler: sampler_comparison, coords: vec2<f32>, depth_ref: f32, @const offset: vec2<i32>) -> f32 */
     /* fn textureSampleCompareLevel<A : iu32>(texture: texture_depth_2d_array, sampler: sampler_comparison, coords: vec2<f32>, array_index: A, depth_ref: f32) -> f32 */
@@ -11624,7 +11630,7 @@
     /* overloads */ OverloadIndex(186),
   },
   {
-    /* [97] */
+    /* [98] */
     /* fn textureSampleGrad(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, ddx: vec2<f32>, ddy: vec2<f32>) -> vec4<f32> */
     /* fn textureSampleGrad(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, ddx: vec2<f32>, ddy: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32> */
     /* fn textureSampleGrad<A : iu32>(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, ddx: vec2<f32>, ddy: vec2<f32>) -> vec4<f32> */
@@ -11637,7 +11643,7 @@
     /* overloads */ OverloadIndex(166),
   },
   {
-    /* [98] */
+    /* [99] */
     /* fn textureSampleLevel(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, level: f32) -> vec4<f32> */
     /* fn textureSampleLevel(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, level: f32, @const offset: vec2<i32>) -> vec4<f32> */
     /* fn textureSampleLevel<A : iu32>(texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, level: f32) -> vec4<f32> */
@@ -11656,14 +11662,14 @@
     /* overloads */ OverloadIndex(79),
   },
   {
-    /* [99] */
+    /* [100] */
     /* fn textureSampleBaseClampToEdge(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32> */
     /* fn textureSampleBaseClampToEdge(texture: texture_external, sampler: sampler, coords: vec2<f32>) -> vec4<f32> */
     /* num overloads */ 2,
     /* overloads */ OverloadIndex(406),
   },
   {
-    /* [100] */
+    /* [101] */
     /* fn textureStore<C : iu32>(texture: texture_storage_1d<f32_texel_format, writable>, coords: C, value: vec4<f32>) */
     /* fn textureStore<C : iu32>(texture: texture_storage_2d<f32_texel_format, writable>, coords: vec2<C>, value: vec4<f32>) */
     /* fn textureStore<C : iu32, A : iu32>(texture: texture_storage_2d_array<f32_texel_format, writable>, coords: vec2<C>, array_index: A, value: vec4<f32>) */
@@ -11680,7 +11686,7 @@
     /* overloads */ OverloadIndex(105),
   },
   {
-    /* [101] */
+    /* [102] */
     /* fn textureLoad<T : fiu32, C : iu32, L : iu32>(texture: texture_1d<T>, coords: C, level: L) -> vec4<T> */
     /* fn textureLoad<T : fiu32, C : iu32, L : iu32>(texture: texture_2d<T>, coords: vec2<C>, level: L) -> vec4<T> */
     /* fn textureLoad<T : fiu32, C : iu32, A : iu32, L : iu32>(texture: texture_2d_array<T>, coords: vec2<C>, array_index: A, level: L) -> vec4<T> */
@@ -11706,79 +11712,79 @@
     /* overloads */ OverloadIndex(27),
   },
   {
-    /* [102] */
+    /* [103] */
     /* fn atomicLoad<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>) -> T */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(455),
   },
   {
-    /* [103] */
+    /* [104] */
     /* fn atomicStore<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(456),
   },
   {
-    /* [104] */
+    /* [105] */
     /* fn atomicAdd<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(457),
   },
   {
-    /* [105] */
+    /* [106] */
     /* fn atomicSub<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(457),
   },
   {
-    /* [106] */
+    /* [107] */
     /* fn atomicMax<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(457),
   },
   {
-    /* [107] */
+    /* [108] */
     /* fn atomicMin<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(457),
   },
   {
-    /* [108] */
+    /* [109] */
     /* fn atomicAnd<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(457),
   },
   {
-    /* [109] */
+    /* [110] */
     /* fn atomicOr<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(457),
   },
   {
-    /* [110] */
+    /* [111] */
     /* fn atomicXor<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(457),
   },
   {
-    /* [111] */
+    /* [112] */
     /* fn atomicExchange<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(457),
   },
   {
-    /* [112] */
+    /* [113] */
     /* fn atomicCompareExchangeWeak<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T> */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(458),
   },
   {
-    /* [113] */
+    /* [114] */
     /* fn subgroupBallot() -> vec4<u32> */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(459),
   },
   {
-    /* [114] */
+    /* [115] */
     /* fn _tint_materialize<T>(T) -> T */
     /* num overloads */ 1,
     /* overloads */ OverloadIndex(460),

diff --git a/src/tint/lang/wgsl/resolver/builtin_validation_test.cc b/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
index bcbd5f4..038be27 100644
--- a/src/tint/lang/wgsl/resolver/builtin_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/builtin_validation_test.cc

@@ -680,5 +680,31 @@
     EXPECT_TRUE(r()->Resolve());
 }
 
+TEST_F(ResolverBuiltinValidationTest, TextureBarrierWithoutExtension) {
+    // fn func { textureBarrier(); }
+    Func("func", tint::Empty, ty.void_(),
+         Vector{
+             CallStmt(Call(Source{Source::Location{12, 34}}, "textureBarrier")),
+         });
+
+    EXPECT_FALSE(r()->Resolve());
+    EXPECT_EQ(
+        r()->error(),
+        R"(12:34 error: cannot call built-in function 'textureBarrier' without extension chromium_experimental_read_write_storage_texture)");
+}
+
+TEST_F(ResolverBuiltinValidationTest, TextureBarrierWithExtension) {
+    // enable chromium_experimental_read_write_storage_texture;
+    // fn func { textureBarrier(); }
+    Enable(core::Extension::kChromiumExperimentalReadWriteStorageTexture);
+
+    Func("func", tint::Empty, ty.void_(),
+         Vector{
+             CallStmt(Call(Source{Source::Location{12, 34}}, "textureBarrier")),
+         });
+
+    EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
 }  // namespace
 }  // namespace tint::resolver

diff --git a/src/tint/lang/wgsl/resolver/uniformity_test.cc b/src/tint/lang/wgsl/resolver/uniformity_test.cc
index fc48e7f..0cba7df 100644
--- a/src/tint/lang/wgsl/resolver/uniformity_test.cc
+++ b/src/tint/lang/wgsl/resolver/uniformity_test.cc

@@ -113,6 +113,7 @@
         kUserRequiredToBeUniform,
         kWorkgroupBarrier,
         kStorageBarrier,
+        kTextureBarrier,
         kWorkgroupUniformLoad,
         kTextureSample,
         kTextureSampleBias,
@@ -182,6 +183,8 @@
                 return "workgroupBarrier()";
             case kStorageBarrier:
                 return "storageBarrier()";
+            case kTextureBarrier:
+                return "textureBarrier()";
             case kWorkgroupUniformLoad:
                 return "_ = workgroupUniformLoad(&w)";
             case kTextureSample:
@@ -257,6 +260,7 @@
             CASE(kUserRequiredToBeUniform);
             CASE(kWorkgroupBarrier);
             CASE(kStorageBarrier);
+            CASE(kTextureBarrier);
             CASE(kWorkgroupUniformLoad);
             CASE(kTextureSample);
             CASE(kTextureSampleBias);
@@ -284,6 +288,8 @@
     auto condition = static_cast<Condition>(std::get<0>(GetParam()));
     auto function = static_cast<Function>(std::get<1>(GetParam()));
     std::string src = R"(
+enable chromium_experimental_read_write_storage_texture;
+
 var<private> p : i32;
 var<workgroup> w : i32;
 @group(0) @binding(0) var<uniform> u : i32;

diff --git a/src/tint/lang/wgsl/sem/builtin.cc b/src/tint/lang/wgsl/sem/builtin.cc
index ece6ded..3845778 100644
--- a/src/tint/lang/wgsl/sem/builtin.cc
+++ b/src/tint/lang/wgsl/sem/builtin.cc

@@ -99,6 +99,9 @@
     if (IsSubgroup()) {
         return core::Extension::kChromiumExperimentalSubgroups;
     }
+    if (type_ == core::Function::kTextureBarrier) {
+        return core::Extension::kChromiumExperimentalReadWriteStorageTexture;
+    }
     return core::Extension::kUndefined;
 }
 

diff --git a/test/tint/builtins/gen/gen.wgsl.tmpl b/test/tint/builtins/gen/gen.wgsl.tmpl
index d74ecf0..a9d5718 100644
--- a/test/tint/builtins/gen/gen.wgsl.tmpl
+++ b/test/tint/builtins/gen/gen.wgsl.tmpl

@@ -208,7 +208,7 @@
 {{    end -}}
 
 {{- /* Check and emit chromium_experimental_read_write_storage_texture */ -}}
-{{-   if OverloadUsesReadWriteStorageTexture $overload }}
+{{-   if or (OverloadUsesReadWriteStorageTexture $overload) (eq "textureBarrier" $builtin_name)}}
 enable chromium_experimental_read_write_storage_texture;
 {{    end -}}
 

diff --git a/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl
new file mode 100644
index 0000000..0e66018
--- /dev/null
+++ b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl

@@ -0,0 +1,33 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+// fn textureBarrier()
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}

diff --git a/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..07414b9
--- /dev/null
+++ b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.dxc.hlsl

@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}
+
+Failed to generate: builtins/gen/literal/textureBarrier/3d0f7e.wgsl:24:8 error: HLSL backend does not support extension 'chromium_experimental_read_write_storage_texture'
+enable chromium_experimental_read_write_storage_texture;
+       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+

diff --git a/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..07414b9
--- /dev/null
+++ b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.fxc.hlsl

@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}
+
+Failed to generate: builtins/gen/literal/textureBarrier/3d0f7e.wgsl:24:8 error: HLSL backend does not support extension 'chromium_experimental_read_write_storage_texture'
+enable chromium_experimental_read_write_storage_texture;
+       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+

diff --git a/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.glsl b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.glsl
new file mode 100644
index 0000000..22638ea
--- /dev/null
+++ b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.glsl

@@ -0,0 +1,15 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}
+
+Failed to generate: error: Unknown builtin method: textureBarrier

diff --git a/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.msl b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.msl
new file mode 100644
index 0000000..6ab09ee
--- /dev/null
+++ b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.msl

@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}
+
+Failed to generate: builtins/gen/literal/textureBarrier/3d0f7e.wgsl:24:8 error: MSL backend does not support extension 'chromium_experimental_read_write_storage_texture'
+enable chromium_experimental_read_write_storage_texture;
+       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+

diff --git a/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.spvasm b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.spvasm
new file mode 100644
index 0000000..101d0e9
--- /dev/null
+++ b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.spvasm

@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}
+
+Failed to generate: builtins/gen/literal/textureBarrier/3d0f7e.wgsl:24:8 error: SPIR-V backend does not support extension 'chromium_experimental_read_write_storage_texture'
+enable chromium_experimental_read_write_storage_texture;
+       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+

diff --git a/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.wgsl b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.wgsl
new file mode 100644
index 0000000..6dd62df
--- /dev/null
+++ b/test/tint/builtins/gen/literal/textureBarrier/3d0f7e.wgsl.expected.wgsl

@@ -0,0 +1,10 @@
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}

diff --git a/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl
new file mode 100644
index 0000000..0e66018
--- /dev/null
+++ b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl

@@ -0,0 +1,33 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+////////////////////////////////////////////////////////////////////////////////
+// File generated by tools/src/cmd/gen
+// using the template:
+//   test/tint/builtins/gen/gen.wgsl.tmpl
+//
+// Do not modify this file directly
+////////////////////////////////////////////////////////////////////////////////
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+// fn textureBarrier()
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}

diff --git a/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.dxc.hlsl b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..321204a
--- /dev/null
+++ b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.dxc.hlsl

@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}
+
+Failed to generate: builtins/gen/var/textureBarrier/3d0f7e.wgsl:24:8 error: HLSL backend does not support extension 'chromium_experimental_read_write_storage_texture'
+enable chromium_experimental_read_write_storage_texture;
+       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+

diff --git a/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.fxc.hlsl b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..321204a
--- /dev/null
+++ b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.fxc.hlsl

@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}
+
+Failed to generate: builtins/gen/var/textureBarrier/3d0f7e.wgsl:24:8 error: HLSL backend does not support extension 'chromium_experimental_read_write_storage_texture'
+enable chromium_experimental_read_write_storage_texture;
+       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+

diff --git a/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.glsl b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.glsl
new file mode 100644
index 0000000..22638ea
--- /dev/null
+++ b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.glsl

@@ -0,0 +1,15 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}
+
+Failed to generate: error: Unknown builtin method: textureBarrier

diff --git a/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.msl b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.msl
new file mode 100644
index 0000000..192df8b
--- /dev/null
+++ b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.msl

@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}
+
+Failed to generate: builtins/gen/var/textureBarrier/3d0f7e.wgsl:24:8 error: MSL backend does not support extension 'chromium_experimental_read_write_storage_texture'
+enable chromium_experimental_read_write_storage_texture;
+       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+

diff --git a/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.spvasm b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.spvasm
new file mode 100644
index 0000000..bc66566
--- /dev/null
+++ b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.spvasm

@@ -0,0 +1,18 @@
+SKIP: FAILED
+
+
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}
+
+Failed to generate: builtins/gen/var/textureBarrier/3d0f7e.wgsl:24:8 error: SPIR-V backend does not support extension 'chromium_experimental_read_write_storage_texture'
+enable chromium_experimental_read_write_storage_texture;
+       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+

diff --git a/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.wgsl b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.wgsl
new file mode 100644
index 0000000..6dd62df
--- /dev/null
+++ b/test/tint/builtins/gen/var/textureBarrier/3d0f7e.wgsl.expected.wgsl

@@ -0,0 +1,10 @@
+enable chromium_experimental_read_write_storage_texture;
+
+fn textureBarrier_3d0f7e() {
+  textureBarrier();
+}
+
+@compute @workgroup_size(1)
+fn compute_main() {
+  textureBarrier_3d0f7e();
+}