transform/DMA: Fix ignore() for buffer members

https://dawn-review.googlesource.com/c/tint/+/60213 special cased ignore() to work around tint:1046.
This fix produced bad output for structures when they are fully decomposed into ByteAddressBuffers, as the final HLSL references a structure that no longer exists.

Fixes CTS tests, and tint->dawn roll.

Change-Id: If6eab083c5f0bcca4a90c582df255b77e97a8e9f
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60347
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>

src/transform/decompose_memory_access.cc

@@ -917,8 +917,15 @@ void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) {
       if (auto* intrinsic = call->Target()->As<sem::Intrinsic>()) {
         if (intrinsic->Type() == sem::IntrinsicType::kIgnore) {
           // ignore(X)
-          // Don't convert X into a load, this isn't actually used.
+          // If X is an memory access, don't transform it into a load, as it
-          state.TakeAccess(call_expr->params()[0]);
+          // may refer to a structure holding a runtime array, which cannot be
+          // loaded. Instead replace X with the underlying storage / uniform
+          // buffer variable.
+          if (auto access = state.TakeAccess(call_expr->params()[0])) {
+            ctx.Replace(call_expr->params()[0], [=, &ctx] {
+              return ctx.CloneWithoutTransform(access.var->Declaration());
+            });
+          }
           continue;
         }
         if (intrinsic->Type() == sem::IntrinsicType::kArrayLength) {

test/intrinsics/ignore/runtime_array.wgsl

@@ -0,0 +1,11 @@
+[[block]]
+struct S {
+    arr : array<i32>;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(s.arr);
+}

test/intrinsics/ignore/runtime_array.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+RWByteAddressBuffer s : register(u0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  s;
+  return;
+}

test/intrinsics/ignore/runtime_array.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  /* 0x0000 */ int arr[1];
+};
+
+kernel void tint_symbol(device S& s [[buffer(0)]]) {
+  (void) s.arr;
+  return;
+}
+

test/intrinsics/ignore/runtime_array.wgsl.expected.spvasm

@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "arr"
+               OpName %s "s"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+          %S = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer__runtimearr_int = OpTypePointer StorageBuffer %_runtimearr_int
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %14 = OpAccessChain %_ptr_StorageBuffer__runtimearr_int %s %uint_0
+         %15 = OpLoad %_runtimearr_int %14
+               OpReturn
+               OpFunctionEnd

test/intrinsics/ignore/runtime_array.wgsl.expected.wgsl

@@ -0,0 +1,11 @@
+[[block]]
+struct S {
+  arr : array<i32>;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(s.arr);
+}

test/intrinsics/ignore/storage_buffer.wgsl

@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+    i : i32;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(s);
+    ignore(s.i);
+}

test/intrinsics/ignore/storage_buffer.wgsl.expected.hlsl

@@ -0,0 +1,8 @@
+RWByteAddressBuffer s : register(u0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  s;
+  s;
+  return;
+}

test/intrinsics/ignore/storage_buffer.wgsl.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  /* 0x0000 */ int i;
+};
+
+kernel void tint_symbol(device S& s [[buffer(0)]]) {
+  (void) s;
+  (void) s.i;
+  return;
+}
+

test/intrinsics/ignore/storage_buffer.wgsl.expected.spvasm

@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "i"
+               OpName %s "s"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %10 = OpLoad %S %s
+         %15 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0
+         %16 = OpLoad %int %15
+               OpReturn
+               OpFunctionEnd

test/intrinsics/ignore/storage_buffer.wgsl.expected.wgsl

@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+  i : i32;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(s);
+  ignore(s.i);
+}

test/intrinsics/ignore/uniform_buffer.wgsl

@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+    i : i32;
+};
+
+[[binding(0), group(0)]] var<uniform> u : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(u);
+    ignore(u.i);
+}

test/intrinsics/ignore/uniform_buffer.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  u;
+  u;
+  return;
+}

test/intrinsics/ignore/uniform_buffer.wgsl.expected.msl

@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  /* 0x0000 */ int i;
+};
+
+kernel void tint_symbol(constant S& u [[buffer(0)]]) {
+  (void) u;
+  (void) u.i;
+  return;
+}
+

test/intrinsics/ignore/uniform_buffer.wgsl.expected.spvasm

@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "i"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %u NonWritable
+               OpDecorate %u Binding 0
+               OpDecorate %u DescriptorSet 0
+        %int = OpTypeInt 32 1
+          %S = OpTypeStruct %int
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+          %u = OpVariable %_ptr_Uniform_S Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %10 = OpLoad %S %u
+         %15 = OpAccessChain %_ptr_Uniform_int %u %uint_0
+         %16 = OpLoad %int %15
+               OpReturn
+               OpFunctionEnd

test/intrinsics/ignore/uniform_buffer.wgsl.expected.wgsl

@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+  i : i32;
+};
+
+[[binding(0), group(0)]] var<uniform> u : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(u);
+  ignore(u.i);
+}