tint: Don't override alignment with @offset
Overriding the alignment to 1 would cause nested structures to be
incorrectly laid out. The fix: Don't override the alignment.
All struct layout validation works on the sem offsets, so none of this
has to change.
Bug: tint:1776
Change-Id: Ic01d45fb2790cd823ed9a55e336860ebdc351aea
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/112603
Kokoro: Ben Clayton <bclayton@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index 3ad8661..60f8207 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -3173,7 +3173,6 @@
AddError("offsets must be in ascending order", o->source);
return false;
}
- align = 1;
has_offset_attr = true;
return true;
},
diff --git a/src/tint/resolver/struct_layout_test.cc b/src/tint/resolver/struct_layout_test.cc
index 1e5ec45..acfe691 100644
--- a/src/tint/resolver/struct_layout_test.cc
+++ b/src/tint/resolver/struct_layout_test.cc
@@ -555,5 +555,47 @@
}
}
+TEST_F(ResolverStructLayoutTest, OffsetAttributes) {
+ auto* inner = Structure("Inner", utils::Vector{
+ Member("a", ty.f32(), utils::Vector{MemberOffset(8_i)}),
+ Member("b", ty.f32(), utils::Vector{MemberOffset(16_i)}),
+ Member("c", ty.f32(), utils::Vector{MemberOffset(32_i)}),
+ });
+ auto* s = Structure("S", utils::Vector{
+ Member("a", ty.f32(), utils::Vector{MemberOffset(4_i)}),
+ Member("b", ty.u32(), utils::Vector{MemberOffset(8_i)}),
+ Member("c", ty.Of(inner), utils::Vector{MemberOffset(32_i)}),
+ Member("d", ty.i32()),
+ Member("e", ty.i32(), utils::Vector{MemberOffset(128_i)}),
+ });
+
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* sem = TypeOf(s)->As<sem::Struct>();
+ ASSERT_NE(sem, nullptr);
+ EXPECT_EQ(sem->Size(), 132u);
+ EXPECT_EQ(sem->SizeNoPadding(), 132u);
+ EXPECT_EQ(sem->Align(), 4u);
+ ASSERT_EQ(sem->Members().size(), 5u);
+ EXPECT_EQ(sem->Members()[0]->Offset(), 4u);
+ EXPECT_EQ(sem->Members()[0]->Align(), 4u);
+ EXPECT_EQ(sem->Members()[0]->Size(), 4u);
+ EXPECT_EQ(sem->Members()[1]->Offset(), 8u);
+ EXPECT_EQ(sem->Members()[1]->Align(), 4u);
+ EXPECT_EQ(sem->Members()[1]->Size(), 4u);
+ EXPECT_EQ(sem->Members()[2]->Offset(), 32u);
+ EXPECT_EQ(sem->Members()[2]->Align(), 4u);
+ EXPECT_EQ(sem->Members()[2]->Size(), 36u);
+ EXPECT_EQ(sem->Members()[3]->Offset(), 68u);
+ EXPECT_EQ(sem->Members()[3]->Align(), 4u);
+ EXPECT_EQ(sem->Members()[3]->Size(), 4u);
+ EXPECT_EQ(sem->Members()[4]->Offset(), 128u);
+ EXPECT_EQ(sem->Members()[4]->Align(), 4u);
+ EXPECT_EQ(sem->Members()[4]->Size(), 4u);
+ for (auto& m : sem->Members()) {
+ EXPECT_EQ(m->Struct()->Declaration(), s);
+ }
+}
+
} // namespace
} // namespace tint::resolver
diff --git a/src/tint/writer/wgsl/generator_impl.cc b/src/tint/writer/wgsl/generator_impl.cc
index f701146..cfd512b 100644
--- a/src/tint/writer/wgsl/generator_impl.cc
+++ b/src/tint/writer/wgsl/generator_impl.cc
@@ -606,8 +606,8 @@
increment_indent();
uint32_t offset = 0;
for (auto* mem : str->members) {
- // TODO(crbug.com/tint/798) move the @offset attribute handling to the
- // transform::Wgsl sanitizer.
+ // TODO(crbug.com/tint/798) move the @offset attribute handling to the transform::Wgsl
+ // sanitizer.
if (auto* mem_sem = program_->Sem().Get(mem)) {
offset = utils::RoundUp(mem_sem->Align(), offset);
if (uint32_t padding = mem_sem->Offset() - offset) {
diff --git a/test/tint/bug/tint/1088.spvasm.expected.wgsl b/test/tint/bug/tint/1088.spvasm.expected.wgsl
index 54ce4f2..5468399 100644
--- a/test/tint/bug/tint/1088.spvasm.expected.wgsl
+++ b/test/tint/bug/tint/1088.spvasm.expected.wgsl
@@ -12,8 +12,6 @@
worldViewProjection : mat4x4<f32>,
/* @offset(64) */
time : f32,
- @size(12)
- padding : u32,
/* @offset(80) */
test2 : Arr,
/* @offset(208) */
diff --git a/test/tint/bug/tint/1520.spvasm.expected.wgsl b/test/tint/bug/tint/1520.spvasm.expected.wgsl
index 931944d..812b115 100644
--- a/test/tint/bug/tint/1520.spvasm.expected.wgsl
+++ b/test/tint/bug/tint/1520.spvasm.expected.wgsl
@@ -3,8 +3,6 @@
padding : u32,
/* @offset(16) */
unknownInput_S1_c0 : f32,
- @size(12)
- padding_1 : u32,
/* @offset(32) */
ucolorRed_S1_c0 : vec4<f32>,
/* @offset(48) */
diff --git a/test/tint/bug/tint/1776.spvasm b/test/tint/bug/tint/1776.spvasm
new file mode 100644
index 0000000..3e85c9e
--- /dev/null
+++ b/test/tint/bug/tint/1776.spvasm
@@ -0,0 +1,44 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 19
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %sb_block "sb_block"
+ OpMemberName %sb_block 0 "inner"
+ OpName %S "S"
+ OpMemberName %S 0 "a"
+ OpMemberName %S 1 "b"
+ OpName %sb "sb"
+ OpName %main "main"
+ OpDecorate %sb_block Block
+ OpMemberDecorate %sb_block 0 Offset 0
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 16
+ OpDecorate %_runtimearr_S ArrayStride 32
+ OpDecorate %sb NonWritable
+ OpDecorate %sb DescriptorSet 0
+ OpDecorate %sb Binding 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %v4float %int
+%_runtimearr_S = OpTypeRuntimeArray %S
+ %sb_block = OpTypeStruct %_runtimearr_S
+%_ptr_StorageBuffer_sb_block = OpTypePointer StorageBuffer %sb_block
+ %sb = OpVariable %_ptr_StorageBuffer_sb_block StorageBuffer
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+ %int_1 = OpConstant %int 1
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %main = OpFunction %void None %9
+ %12 = OpLabel
+ %17 = OpAccessChain %_ptr_StorageBuffer_S %sb %uint_0 %int_1
+ %18 = OpLoad %S %17
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl b/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl
new file mode 100644
index 0000000..1cd2986
--- /dev/null
+++ b/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl
@@ -0,0 +1,22 @@
+struct S {
+ float4 a;
+ int b;
+};
+
+RWByteAddressBuffer sb : register(u0, space0);
+
+S tint_symbol(RWByteAddressBuffer buffer, uint offset) {
+ const S tint_symbol_3 = {asfloat(buffer.Load4((offset + 0u))), asint(buffer.Load((offset + 16u)))};
+ return tint_symbol_3;
+}
+
+void main_1() {
+ const S x_18 = tint_symbol(sb, 32u);
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ main_1();
+ return;
+}
diff --git a/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl b/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl
new file mode 100644
index 0000000..1cd2986
--- /dev/null
+++ b/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl
@@ -0,0 +1,22 @@
+struct S {
+ float4 a;
+ int b;
+};
+
+RWByteAddressBuffer sb : register(u0, space0);
+
+S tint_symbol(RWByteAddressBuffer buffer, uint offset) {
+ const S tint_symbol_3 = {asfloat(buffer.Load4((offset + 0u))), asint(buffer.Load((offset + 16u)))};
+ return tint_symbol_3;
+}
+
+void main_1() {
+ const S x_18 = tint_symbol(sb, 32u);
+ return;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ main_1();
+ return;
+}
diff --git a/test/tint/bug/tint/1776.spvasm.expected.glsl b/test/tint/bug/tint/1776.spvasm.expected.glsl
new file mode 100644
index 0000000..e73e547
--- /dev/null
+++ b/test/tint/bug/tint/1776.spvasm.expected.glsl
@@ -0,0 +1,28 @@
+#version 310 es
+
+struct S {
+ vec4 a;
+ int b;
+ uint pad;
+ uint pad_1;
+ uint pad_2;
+};
+
+layout(binding = 0, std430) buffer sb_block_ssbo {
+ S inner[];
+} sb;
+
+void main_1() {
+ S x_18 = sb.inner[1];
+ return;
+}
+
+void tint_symbol() {
+ main_1();
+}
+
+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/1776.spvasm.expected.msl b/test/tint/bug/tint/1776.spvasm.expected.msl
new file mode 100644
index 0000000..0933ac0
--- /dev/null
+++ b/test/tint/bug/tint/1776.spvasm.expected.msl
@@ -0,0 +1,36 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct S {
+ /* 0x0000 */ float4 a;
+ /* 0x0010 */ int b;
+ /* 0x0014 */ tint_array<int8_t, 12> tint_pad;
+};
+
+struct sb_block {
+ /* 0x0000 */ tint_array<S, 1> inner;
+};
+
+void main_1(device sb_block* const tint_symbol_1) {
+ S const x_18 = (*(tint_symbol_1)).inner[1];
+ return;
+}
+
+kernel void tint_symbol(device sb_block* tint_symbol_2 [[buffer(0)]]) {
+ main_1(tint_symbol_2);
+ return;
+}
+
diff --git a/test/tint/bug/tint/1776.spvasm.expected.spvasm b/test/tint/bug/tint/1776.spvasm.expected.spvasm
new file mode 100644
index 0000000..6e076d1
--- /dev/null
+++ b/test/tint/bug/tint/1776.spvasm.expected.spvasm
@@ -0,0 +1,49 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 22
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %sb_block "sb_block"
+ OpMemberName %sb_block 0 "inner"
+ OpName %S "S"
+ OpMemberName %S 0 "a"
+ OpMemberName %S 1 "b"
+ OpName %sb "sb"
+ OpName %main_1 "main_1"
+ OpName %main "main"
+ OpDecorate %sb_block Block
+ OpMemberDecorate %sb_block 0 Offset 0
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 16
+ OpDecorate %_runtimearr_S ArrayStride 32
+ OpDecorate %sb DescriptorSet 0
+ OpDecorate %sb Binding 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %v4float %int
+%_runtimearr_S = OpTypeRuntimeArray %S
+ %sb_block = OpTypeStruct %_runtimearr_S
+%_ptr_StorageBuffer_sb_block = OpTypePointer StorageBuffer %sb_block
+ %sb = OpVariable %_ptr_StorageBuffer_sb_block StorageBuffer
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+ %int_1 = OpConstant %int 1
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %main_1 = OpFunction %void None %9
+ %12 = OpLabel
+ %17 = OpAccessChain %_ptr_StorageBuffer_S %sb %uint_0 %int_1
+ %18 = OpLoad %S %17
+ OpReturn
+ OpFunctionEnd
+ %main = OpFunction %void None %9
+ %20 = OpLabel
+ %21 = OpFunctionCall %void %main_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/1776.spvasm.expected.wgsl b/test/tint/bug/tint/1776.spvasm.expected.wgsl
new file mode 100644
index 0000000..c5f6bc7
--- /dev/null
+++ b/test/tint/bug/tint/1776.spvasm.expected.wgsl
@@ -0,0 +1,25 @@
+struct S {
+ /* @offset(0) */
+ a : vec4<f32>,
+ /* @offset(16) */
+ b : i32,
+}
+
+type RTArr = array<S>;
+
+struct sb_block {
+ /* @offset(0) */
+ inner : RTArr,
+}
+
+@group(0) @binding(0) var<storage, read_write> sb : sb_block;
+
+fn main_1() {
+ let x_18 : S = sb.inner[1i];
+ return;
+}
+
+@compute @workgroup_size(1i, 1i, 1i)
+fn main() {
+ main_1();
+}
diff --git a/test/tint/bug/tint/1776.wgsl b/test/tint/bug/tint/1776.wgsl
new file mode 100644
index 0000000..893b64c
--- /dev/null
+++ b/test/tint/bug/tint/1776.wgsl
@@ -0,0 +1,11 @@
+struct S { /* size: 32 align: 16 */
+ a : vec4<f32>,
+ b : i32,
+}
+
+@group(0) @binding(0) var<storage> sb : array<S>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let x = sb[1];
+}
diff --git a/test/tint/bug/tint/1776.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1776.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..36f87a7
--- /dev/null
+++ b/test/tint/bug/tint/1776.wgsl.expected.dxc.hlsl
@@ -0,0 +1,17 @@
+struct S {
+ float4 a;
+ int b;
+};
+
+ByteAddressBuffer sb : register(t0, space0);
+
+S tint_symbol(ByteAddressBuffer buffer, uint offset) {
+ const S tint_symbol_3 = {asfloat(buffer.Load4((offset + 0u))), asint(buffer.Load((offset + 16u)))};
+ return tint_symbol_3;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ const S x = tint_symbol(sb, 32u);
+ return;
+}
diff --git a/test/tint/bug/tint/1776.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1776.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..36f87a7
--- /dev/null
+++ b/test/tint/bug/tint/1776.wgsl.expected.fxc.hlsl
@@ -0,0 +1,17 @@
+struct S {
+ float4 a;
+ int b;
+};
+
+ByteAddressBuffer sb : register(t0, space0);
+
+S tint_symbol(ByteAddressBuffer buffer, uint offset) {
+ const S tint_symbol_3 = {asfloat(buffer.Load4((offset + 0u))), asint(buffer.Load((offset + 16u)))};
+ return tint_symbol_3;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+ const S x = tint_symbol(sb, 32u);
+ return;
+}
diff --git a/test/tint/bug/tint/1776.wgsl.expected.glsl b/test/tint/bug/tint/1776.wgsl.expected.glsl
new file mode 100644
index 0000000..6a83ec2
--- /dev/null
+++ b/test/tint/bug/tint/1776.wgsl.expected.glsl
@@ -0,0 +1,23 @@
+#version 310 es
+
+struct S {
+ vec4 a;
+ int b;
+ uint pad;
+ uint pad_1;
+ uint pad_2;
+};
+
+layout(binding = 0, std430) buffer sb_block_ssbo {
+ S inner[];
+} sb;
+
+void tint_symbol() {
+ S x = sb.inner[1];
+}
+
+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/1776.wgsl.expected.msl b/test/tint/bug/tint/1776.wgsl.expected.msl
new file mode 100644
index 0000000..38097d0
--- /dev/null
+++ b/test/tint/bug/tint/1776.wgsl.expected.msl
@@ -0,0 +1,31 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+ const constant T& operator[](size_t i) const constant { return elements[i]; }
+ device T& operator[](size_t i) device { return elements[i]; }
+ const device T& operator[](size_t i) const device { return elements[i]; }
+ thread T& operator[](size_t i) thread { return elements[i]; }
+ const thread T& operator[](size_t i) const thread { return elements[i]; }
+ threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+ const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+ T elements[N];
+};
+
+struct S {
+ /* 0x0000 */ float4 a;
+ /* 0x0010 */ int b;
+ /* 0x0014 */ tint_array<int8_t, 12> tint_pad;
+};
+
+struct tint_symbol_2 {
+ /* 0x0000 */ tint_array<S, 1> arr;
+};
+
+kernel void tint_symbol(const device tint_symbol_2* tint_symbol_1 [[buffer(0)]]) {
+ S const x = (*(tint_symbol_1)).arr[1];
+ return;
+}
+
diff --git a/test/tint/bug/tint/1776.wgsl.expected.spvasm b/test/tint/bug/tint/1776.wgsl.expected.spvasm
new file mode 100644
index 0000000..3e85c9e
--- /dev/null
+++ b/test/tint/bug/tint/1776.wgsl.expected.spvasm
@@ -0,0 +1,44 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 19
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %sb_block "sb_block"
+ OpMemberName %sb_block 0 "inner"
+ OpName %S "S"
+ OpMemberName %S 0 "a"
+ OpMemberName %S 1 "b"
+ OpName %sb "sb"
+ OpName %main "main"
+ OpDecorate %sb_block Block
+ OpMemberDecorate %sb_block 0 Offset 0
+ OpMemberDecorate %S 0 Offset 0
+ OpMemberDecorate %S 1 Offset 16
+ OpDecorate %_runtimearr_S ArrayStride 32
+ OpDecorate %sb NonWritable
+ OpDecorate %sb DescriptorSet 0
+ OpDecorate %sb Binding 0
+ %float = OpTypeFloat 32
+ %v4float = OpTypeVector %float 4
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %v4float %int
+%_runtimearr_S = OpTypeRuntimeArray %S
+ %sb_block = OpTypeStruct %_runtimearr_S
+%_ptr_StorageBuffer_sb_block = OpTypePointer StorageBuffer %sb_block
+ %sb = OpVariable %_ptr_StorageBuffer_sb_block StorageBuffer
+ %void = OpTypeVoid
+ %9 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+ %int_1 = OpConstant %int 1
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %main = OpFunction %void None %9
+ %12 = OpLabel
+ %17 = OpAccessChain %_ptr_StorageBuffer_S %sb %uint_0 %int_1
+ %18 = OpLoad %S %17
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/tint/1776.wgsl.expected.wgsl b/test/tint/bug/tint/1776.wgsl.expected.wgsl
new file mode 100644
index 0000000..08733dc
--- /dev/null
+++ b/test/tint/bug/tint/1776.wgsl.expected.wgsl
@@ -0,0 +1,11 @@
+struct S {
+ a : vec4<f32>,
+ b : i32,
+}
+
+@group(0) @binding(0) var<storage> sb : array<S>;
+
+@compute @workgroup_size(1)
+fn main() {
+ let x = sb[1];
+}
diff --git a/test/tint/bug/tint/870.spvasm.expected.glsl b/test/tint/bug/tint/870.spvasm.expected.glsl
index 0cddbf5..f56b3de 100644
--- a/test/tint/bug/tint/870.spvasm.expected.glsl
+++ b/test/tint/bug/tint/870.spvasm.expected.glsl
@@ -6,6 +6,7 @@
vec4 offset;
int essence;
int orientation[6];
+ uint pad;
};
struct x_B4_BuildInformation {
diff --git a/test/tint/bug/tint/870.spvasm.expected.msl b/test/tint/bug/tint/870.spvasm.expected.msl
index 9a97441..408a9712 100644
--- a/test/tint/bug/tint/870.spvasm.expected.msl
+++ b/test/tint/bug/tint/870.spvasm.expected.msl
@@ -19,6 +19,7 @@
/* 0x0010 */ float4 offset;
/* 0x0020 */ int essence;
/* 0x0024 */ tint_array<int, 6> orientation;
+ /* 0x003c */ tint_array<int8_t, 4> tint_pad;
};
struct x_B4_BuildInformation {
