[msl] Fix read-write textureStore in entry points

When emitting calls to the `fence()` builtin member function, wrap the object in a `const_cast<>` to remove the `const` qualifier. We do this to work around an MSL bug that prevents us from being able to use texture fence intrinsics when texture handles are stored inside const-qualified structures. Fixed: 365570202 Change-Id: I8571804f063e60f544e2150f1454a6d7c63a4af3 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/206194 Auto-Submit: James Price <[email protected]> Reviewed-by: Antonio Maiorano <[email protected]> Commit-Queue: James Price <[email protected]>
google · Sep 10, 2024 · bc56248 · bc56248
1 parent 369919f
commit bc56248
Show file tree

Hide file tree

Showing 373 changed files with 490 additions and 362 deletions.
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
@@ -902,7 +902,20 @@ class Printer : public tint::TextGenerator {
     }
 
     void EmitMslMemberBuiltinCall(StringStream& out, const msl::ir::MemberBuiltinCall* c) {
-        EmitValue(out, c->Object());
+        if (c->Func() == BuiltinFn::kFence) {
+            // If this is a fence builtin, we need to `const_cast<>` the object to remove the
+            // `const` qualifier. We do this to work around an MSL bug that prevents us from being
+            // able to use texture fence intrinsics when texture handles are stored inside
+            // const-qualified structures (see crbug.com/365570202).
+            out << "const_cast<";
+            EmitType(out, c->Object()->Type());
+            out << "thread &>(";
+            EmitValue(out, c->Object());
+            out << ")";
+        } else {
+            EmitValue(out, c->Object());
+        }
+
         out << "." << c->Func() << "(";
         bool needs_comma = false;
         for (const auto* arg : c->Args()) {

diff --git a/test/tint/bug/tint/365570202.wgsl b/test/tint/bug/tint/365570202.wgsl
@@ -0,0 +1,7 @@
+@group(0) @binding(0)
+var tex : texture_storage_2d<rgba8unorm, read_write>;
+
+@compute @workgroup_size(1)
+fn main() {
+  textureStore(tex, vec2(), vec4());
+}
diff --git a/test/tint/bug/tint/365570202.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/365570202.wgsl.expected.dxc.hlsl
@@ -0,0 +1,7 @@
+RWTexture2D<float4> tex : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tex[(0).xx] = (0.0f).xxxx;
+  return;
+}
diff --git a/test/tint/bug/tint/365570202.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/365570202.wgsl.expected.fxc.hlsl
@@ -0,0 +1,7 @@
+RWTexture2D<float4> tex : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tex[(0).xx] = (0.0f).xxxx;
+  return;
+}
diff --git a/test/tint/bug/tint/365570202.wgsl.expected.glsl b/test/tint/bug/tint/365570202.wgsl.expected.glsl
@@ -0,0 +1,12 @@
+#version 310 es
+
+layout(binding = 0, rgba8) uniform highp writeonly image2D tex;
+void tint_symbol() {
+  imageStore(tex, ivec2(0), vec4(0.0f));
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol();
+  return;
+}
diff --git a/test/tint/bug/tint/365570202.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/tint/365570202.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,7 @@
+
+RWTexture2D<float4> tex : register(u0);
+[numthreads(1, 1, 1)]
+void main() {
+  tex[(0).xx] = (0.0f).xxxx;
+}
+
diff --git a/test/tint/bug/tint/365570202.wgsl.expected.ir.fxc.hlsl b/test/tint/bug/tint/365570202.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,7 @@
+
+RWTexture2D<float4> tex : register(u0);
+[numthreads(1, 1, 1)]
+void main() {
+  tex[(0).xx] = (0.0f).xxxx;
+}
+
diff --git a/test/tint/bug/tint/365570202.wgsl.expected.ir.glsl b/test/tint/bug/tint/365570202.wgsl.expected.ir.glsl
@@ -0,0 +1,11 @@
+SKIP: FAILED
+
+../../src/tint/lang/glsl/writer/printer/printer.cc:729 internal compiler error: TINT_UNREACHABLE 
+********************************************************************
+*  The tint shader compiler has encountered an unexpected error.   *
+*                                                                  *
+*  Please help us fix this issue by submitting a bug report at     *
+*  crbug.com/tint with the source program that triggered the bug.  *
+********************************************************************
+
+tint executable returned error: signal: trace/BPT trap
diff --git a/test/tint/bug/tint/365570202.wgsl.expected.ir.msl b/test/tint/bug/tint/365570202.wgsl.expected.ir.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+  texture2d<float, access::read_write> tex;
+};
+
+kernel void tint_symbol(texture2d<float, access::read_write> tex [[texture(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.tex=tex};
+  tint_module_vars.tex.write(float4(0.0f), uint2(int2(0)));
+  const_cast<texture2d<float, access::read_write>thread &>(tint_module_vars.tex).fence();
+}
diff --git a/test/tint/bug/tint/365570202.wgsl.expected.msl b/test/tint/bug/tint/365570202.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(texture2d<float, access::read_write> tint_symbol_1 [[texture(0)]]) {
+  tint_symbol_1.write(float4(0.0f), uint2(int2(0))); tint_symbol_1.fence();
+  return;
+}
+
diff --git a/test/tint/bug/tint/365570202.wgsl.expected.spvasm b/test/tint/bug/tint/365570202.wgsl.expected.spvasm
@@ -0,0 +1,31 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %tex "tex"
+               OpName %main "main"
+               OpDecorate %tex DescriptorSet 0
+               OpDecorate %tex Binding 0
+               OpDecorate %tex Coherent
+      %float = OpTypeFloat 32
+          %3 = OpTypeImage %float 2D 0 0 0 2 Rgba8
+%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
+        %tex = OpVariable %_ptr_UniformConstant_3 UniformConstant
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+         %11 = OpConstantNull %v2int
+    %v4float = OpTypeVector %float 4
+         %14 = OpConstantNull %v4float
+       %main = OpFunction %void None %7
+          %8 = OpLabel
+          %9 = OpLoad %3 %tex None
+               OpImageWrite %9 %11 %14 None
+               OpReturn
+               OpFunctionEnd
diff --git a/test/tint/bug/tint/365570202.wgsl.expected.wgsl b/test/tint/bug/tint/365570202.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+@group(0) @binding(0) var tex : texture_storage_2d<rgba8unorm, read_write>;
+
+@compute @workgroup_size(1)
+fn main() {
+  textureStore(tex, vec2(), vec4());
+}
diff --git a/test/tint/builtins/gen/literal/textureStore/031506.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/031506.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_031506(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint2(int2(1)));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/036d0e.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/036d0e.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_036d0e(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint3(int3(1)));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture3d<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture3d<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/03e7a0.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/03e7a0.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_03e7a0(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(uint4(1u), uint2(int2(1)));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d<uint, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d<uint, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/042b06.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/042b06.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_042b06(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint2(1u), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d_array<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d_array<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/052a4e.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/052a4e.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_052a4e(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture1d<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture1d<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/053664.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/053664.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_053664(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint2(int2(1)), 1);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d_array<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d_array<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/09e4d5.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/09e4d5.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_09e4d5(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture1d<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture1d<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/0ad124.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/0ad124.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_0ad124(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint(1));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture1d<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture1d<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/0ade9a.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/0ade9a.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_0ade9a(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(uint4(1u), uint2(1u), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d_array<uint, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d_array<uint, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/0ff97a.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/0ff97a.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_0ff97a(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint2(1u), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d_array<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d_array<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/101325.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/101325.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_101325(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(uint4(1u), uint3(1u));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture3d<uint, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture3d<uint, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/145061.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/145061.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_145061(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(uint4(1u), uint2(1u));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d<uint, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d<uint, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/178e69.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/178e69.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_178e69(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint(1));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture1d<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture1d<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/195d1b.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/195d1b.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_195d1b(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint3(1u));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture3d<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture3d<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/197637.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/197637.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_197637(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint(1));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture1d<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture1d<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/1a6c0b.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/1a6c0b.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_1a6c0b(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(int4(1), uint2(1u), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d_array<int, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d_array<int, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/1af236.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/1af236.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_1af236(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint2(int2(1)), 1);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d_array<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d_array<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/1e9fbd.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/1e9fbd.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_1e9fbd(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(int4(1), uint2(1u), 1);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d_array<int, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d_array<int, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/1fef04.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/1fef04.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_1fef04(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(int4(1), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture1d<int, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture1d<int, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/2046db.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/2046db.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_2046db(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint2(1u), 1);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d_array<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d_array<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/2173fd.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/2173fd.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_2173fd(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(uint4(1u), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture1d<uint, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture1d<uint, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/24e6b7.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/24e6b7.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_24e6b7(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(int4(1), uint3(int3(1)));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture3d<int, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture3d<int, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/26a26d.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/26a26d.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_26a26d(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture1d<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture1d<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/272f5a.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/272f5a.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_272f5a(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint2(int2(1)), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d_array<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d_array<float, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/28e109.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/28e109.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_28e109(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(int4(1), uint2(int2(1)));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d<int, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d<int, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/2a60c9.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/2a60c9.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_2a60c9(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(int4(1), uint2(1u), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d_array<int, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d_array<int, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/2addd6.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/2addd6.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_2addd6(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(uint4(1u), uint2(1u), 1u);
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d_array<uint, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d_array<uint, access::read_write> arg_0 [[texture(0)]]) {

diff --git a/test/tint/builtins/gen/literal/textureStore/2c76db.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureStore/2c76db.wgsl.expected.ir.msl
@@ -7,7 +7,7 @@ struct tint_module_vars_struct {
 
 void textureStore_2c76db(tint_module_vars_struct tint_module_vars) {
   tint_module_vars.arg_0.write(float4(1.0f), uint2(1u));
-  tint_module_vars.arg_0.fence();
+  const_cast<texture2d<float, access::read_write>thread &>(tint_module_vars.arg_0).fence();
 }
 
 fragment void fragment_main(texture2d<float, access::read_write> arg_0 [[texture(0)]]) {