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>
diff --git a/src/transform/decompose_memory_access.cc b/src/transform/decompose_memory_access.cc
index 09aed03..aba80cb 100644
--- a/src/transform/decompose_memory_access.cc
+++ b/src/transform/decompose_memory_access.cc
@@ -917,8 +917,15 @@
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.
- state.TakeAccess(call_expr->params()[0]);
+ // If X is an memory access, don't transform it into a load, as it
+ // 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) {
diff --git a/test/intrinsics/ignore.wgsl b/test/intrinsics/ignore/call.wgsl
similarity index 100%
rename from test/intrinsics/ignore.wgsl
rename to test/intrinsics/ignore/call.wgsl
diff --git a/test/intrinsics/ignore.wgsl.expected.hlsl b/test/intrinsics/ignore/call.wgsl.expected.hlsl
similarity index 100%
rename from test/intrinsics/ignore.wgsl.expected.hlsl
rename to test/intrinsics/ignore/call.wgsl.expected.hlsl
diff --git a/test/intrinsics/ignore.wgsl.expected.msl b/test/intrinsics/ignore/call.wgsl.expected.msl
similarity index 100%
rename from test/intrinsics/ignore.wgsl.expected.msl
rename to test/intrinsics/ignore/call.wgsl.expected.msl
diff --git a/test/intrinsics/ignore.wgsl.expected.spvasm b/test/intrinsics/ignore/call.wgsl.expected.spvasm
similarity index 100%
rename from test/intrinsics/ignore.wgsl.expected.spvasm
rename to test/intrinsics/ignore/call.wgsl.expected.spvasm
diff --git a/test/intrinsics/ignore.wgsl.expected.wgsl b/test/intrinsics/ignore/call.wgsl.expected.wgsl
similarity index 100%
rename from test/intrinsics/ignore.wgsl.expected.wgsl
rename to test/intrinsics/ignore/call.wgsl.expected.wgsl
diff --git a/test/intrinsics/ignore/runtime_array.wgsl b/test/intrinsics/ignore/runtime_array.wgsl
new file mode 100644
index 0000000..d934291
--- /dev/null
+++ b/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);
+}
diff --git a/test/intrinsics/ignore/runtime_array.wgsl.expected.hlsl b/test/intrinsics/ignore/runtime_array.wgsl.expected.hlsl
new file mode 100644
index 0000000..3f323c7
--- /dev/null
+++ b/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;
+}
diff --git a/test/intrinsics/ignore/runtime_array.wgsl.expected.msl b/test/intrinsics/ignore/runtime_array.wgsl.expected.msl
new file mode 100644
index 0000000..760ed42
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/intrinsics/ignore/runtime_array.wgsl.expected.spvasm b/test/intrinsics/ignore/runtime_array.wgsl.expected.spvasm
new file mode 100644
index 0000000..4539b59
--- /dev/null
+++ b/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
diff --git a/test/intrinsics/ignore/runtime_array.wgsl.expected.wgsl b/test/intrinsics/ignore/runtime_array.wgsl.expected.wgsl
new file mode 100644
index 0000000..77835c2
--- /dev/null
+++ b/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);
+}
diff --git a/test/intrinsics/ignore/storage_buffer.wgsl b/test/intrinsics/ignore/storage_buffer.wgsl
new file mode 100644
index 0000000..5e64740
--- /dev/null
+++ b/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);
+}
diff --git a/test/intrinsics/ignore/storage_buffer.wgsl.expected.hlsl b/test/intrinsics/ignore/storage_buffer.wgsl.expected.hlsl
new file mode 100644
index 0000000..179f697
--- /dev/null
+++ b/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;
+}
diff --git a/test/intrinsics/ignore/storage_buffer.wgsl.expected.msl b/test/intrinsics/ignore/storage_buffer.wgsl.expected.msl
new file mode 100644
index 0000000..968ab70
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/intrinsics/ignore/storage_buffer.wgsl.expected.spvasm b/test/intrinsics/ignore/storage_buffer.wgsl.expected.spvasm
new file mode 100644
index 0000000..bd5b51f
--- /dev/null
+++ b/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
diff --git a/test/intrinsics/ignore/storage_buffer.wgsl.expected.wgsl b/test/intrinsics/ignore/storage_buffer.wgsl.expected.wgsl
new file mode 100644
index 0000000..526bc3b
--- /dev/null
+++ b/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);
+}
diff --git a/test/intrinsics/ignore/uniform_buffer.wgsl b/test/intrinsics/ignore/uniform_buffer.wgsl
new file mode 100644
index 0000000..b96eb83
--- /dev/null
+++ b/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);
+}
diff --git a/test/intrinsics/ignore/uniform_buffer.wgsl.expected.hlsl b/test/intrinsics/ignore/uniform_buffer.wgsl.expected.hlsl
new file mode 100644
index 0000000..0b0aa2d
--- /dev/null
+++ b/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;
+}
diff --git a/test/intrinsics/ignore/uniform_buffer.wgsl.expected.msl b/test/intrinsics/ignore/uniform_buffer.wgsl.expected.msl
new file mode 100644
index 0000000..9220247
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/intrinsics/ignore/uniform_buffer.wgsl.expected.spvasm b/test/intrinsics/ignore/uniform_buffer.wgsl.expected.spvasm
new file mode 100644
index 0000000..3b232a3
--- /dev/null
+++ b/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
diff --git a/test/intrinsics/ignore/uniform_buffer.wgsl.expected.wgsl b/test/intrinsics/ignore/uniform_buffer.wgsl.expected.wgsl
new file mode 100644
index 0000000..f0aacfd
--- /dev/null
+++ b/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);
+}
