[spirv-reader][ir] Support `ArrayStride` annotation
Add support for processing the `ArrayStride` annotation. This creates an
`ExplicitLayoutArray` which will need to be cleaned up with a transform.
Bug: 391482902
Change-Id: I61ee0e3fcbfe3d074cec7cfa658ce9ab52923369
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/247774
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/reader/parser/composite_test.cc b/src/tint/lang/spirv/reader/parser/composite_test.cc
index 9327adf..b68ee7f 100644
--- a/src/tint/lang/spirv/reader/parser/composite_test.cc
+++ b/src/tint/lang/spirv/reader/parser/composite_test.cc
@@ -143,6 +143,155 @@
)");
}
+TEST_F(SpirvParserTest, CompositeConstruct_Array_ArrayStride_EqualsElementSize) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpDecorate %arr_ty ArrayStride 4
+ %void = OpTypeVoid
+ %u32 = OpTypeInt 32 0
+ %ep_type = OpTypeFunction %void
+ %u32_1 = OpConstant %u32 1
+ %u32_2 = OpConstant %u32 2
+ %u32_3 = OpConstant %u32 3
+ %u32_4 = OpConstant %u32 4
+ %arr_ty = OpTypeArray %u32 %u32_4
+ %fn_type = OpTypeFunction %arr_ty
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+
+ %foo = OpFunction %arr_ty None %fn_type
+ %foo_start = OpLabel
+ %arr = OpCompositeConstruct %arr_ty %u32_1 %u32_2 %u32_3 %u32_4
+ OpReturnValue %arr
+ OpFunctionEnd
+)",
+ R"(
+%2 = func():array<u32, 4> {
+ $B2: {
+ %3:array<u32, 4> = construct 1u, 2u, 3u, 4u
+ ret %3
+ }
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeConstruct_Array_ArrayStride_EqualsElementSize_ArrayVec3) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpDecorate %arr_ty ArrayStride 12
+ %void = OpTypeVoid
+ %u32 = OpTypeInt 32 0
+ %ep_type = OpTypeFunction %void
+ %u32_1 = OpConstant %u32 1
+ %u32_2 = OpConstant %u32 2
+ %u32_3 = OpConstant %u32 3
+ %u32_4 = OpConstant %u32 4
+ %vec3u = OpTypeVector %u32 3
+ %arr_ty = OpTypeArray %vec3u %u32_4
+ %ptr = OpTypePointer Private %arr_ty
+ %vs = OpVariable %ptr Private
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %2 = OpCopyObject %ptr %vs
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:ptr<private, spirv.explicit_layout_array<vec3<u32>, 4>, read_write> = var undef
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %3:ptr<private, spirv.explicit_layout_array<vec3<u32>, 4>, read_write> = let %1
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest,
+ CompositeConstruct_Array_ArrayStride_EqualsElementSize_ArrayVec3_MatchTint) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpDecorate %arr_ty ArrayStride 16
+ %void = OpTypeVoid
+ %u32 = OpTypeInt 32 0
+ %ep_type = OpTypeFunction %void
+ %u32_1 = OpConstant %u32 1
+ %u32_2 = OpConstant %u32 2
+ %u32_3 = OpConstant %u32 3
+ %u32_4 = OpConstant %u32 4
+ %vec3u = OpTypeVector %u32 3
+ %arr_ty = OpTypeArray %vec3u %u32_4
+ %ptr = OpTypePointer Private %arr_ty
+ %vs = OpVariable %ptr Private
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %2 = OpCopyObject %ptr %vs
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:ptr<private, array<vec3<u32>, 4>, read_write> = var undef
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %3:ptr<private, array<vec3<u32>, 4>, read_write> = let %1
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, CompositeConstruct_Array_ArrayStride) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpDecorate %arr_ty ArrayStride 16
+ %void = OpTypeVoid
+ %u32 = OpTypeInt 32 0
+ %ep_type = OpTypeFunction %void
+ %u32_1 = OpConstant %u32 1
+ %u32_2 = OpConstant %u32 2
+ %u32_3 = OpConstant %u32 3
+ %u32_4 = OpConstant %u32 4
+ %arr_ty = OpTypeArray %u32 %u32_4
+ %fn_type = OpTypeFunction %arr_ty
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+
+ %foo = OpFunction %arr_ty None %fn_type
+ %foo_start = OpLabel
+ %arr = OpCompositeConstruct %arr_ty %u32_1 %u32_2 %u32_3 %u32_4
+ OpReturnValue %arr
+ OpFunctionEnd
+)",
+ R"(
+%2 = func():spirv.explicit_layout_array<u32, 4> {
+ $B2: {
+ %3:spirv.explicit_layout_array<u32, 4> = construct 1u, 2u, 3u, 4u
+ ret %3
+ }
+)");
+}
+
TEST_F(SpirvParserTest, CompositeConstruct_ArrayOfVec) {
EXPECT_IR(R"(
OpCapability Shader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 5aeaee9..6890743 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -56,6 +56,7 @@
#include "src/tint/lang/core/type/builtin_structs.h"
#include "src/tint/lang/spirv/builtin_fn.h"
#include "src/tint/lang/spirv/ir/builtin_call.h"
+#include "src/tint/lang/spirv/type/explicit_layout_array.h"
#include "src/tint/lang/spirv/type/image.h"
#include "src/tint/lang/spirv/type/sampled_image.h"
#include "src/tint/lang/spirv/validate/validate.h"
@@ -478,10 +479,11 @@
}
// TODO(crbug.com/1907): Handle decorations that affect the type
+ uint32_t array_stride = 0;
for (auto& deco : type->decorations()) {
switch (spv::Decoration(deco[0])) {
- case spv::Decoration::SpecId: {
- // TODO(dsinclair): Trick clang, remove when actual decorations added
+ case spv::Decoration::ArrayStride: {
+ array_stride = deco[1];
break;
}
default: {
@@ -530,7 +532,7 @@
mat_ty->element_count());
}
case spvtools::opt::analysis::Type::kArray: {
- return EmitArray(type->AsArray());
+ return EmitArray(type->AsArray(), array_stride);
}
case spvtools::opt::analysis::Type::kRuntimeArray: {
auto* arr_ty = type->AsRuntimeArray();
@@ -650,7 +652,8 @@
/// @param arr_ty a SPIR-V array object
/// @returns a Tint array object
- const core::type::Type* EmitArray(const spvtools::opt::analysis::Array* arr_ty) {
+ const core::type::Type* EmitArray(const spvtools::opt::analysis::Array* arr_ty,
+ uint32_t array_stride) {
const auto& length = arr_ty->length_info();
TINT_ASSERT(!length.words.empty());
if (length.words[0] != spvtools::opt::analysis::Array::LengthInfo::kConstant) {
@@ -664,7 +667,15 @@
const uint64_t count_val = count_const->GetZeroExtendedValue();
TINT_ASSERT(count_val <= UINT32_MAX);
- return ty_.array(Type(arr_ty->element_type()), static_cast<uint32_t>(count_val));
+ auto* elem_ty = Type(arr_ty->element_type());
+ uint32_t implicit_stride = tint::RoundUp(elem_ty->Align(), elem_ty->Size());
+ if (array_stride == 0 || array_stride == implicit_stride) {
+ return ty_.array(elem_ty, static_cast<uint32_t>(count_val));
+ }
+
+ return ty_.Get<spirv::type::ExplicitLayoutArray>(
+ elem_ty, ty_.Get<core::type::ConstantArrayCount>(static_cast<uint32_t>(count_val)),
+ elem_ty->Align(), static_cast<uint32_t>(array_stride * count_val), array_stride);
}
/// @param struct_ty a SPIR-V struct object
diff --git a/src/tint/lang/spirv/reader/reader_test.cc b/src/tint/lang/spirv/reader/reader_test.cc
index 5b297f2..e406e90 100644
--- a/src/tint/lang/spirv/reader/reader_test.cc
+++ b/src/tint/lang/spirv/reader/reader_test.cc
@@ -264,8 +264,7 @@
)");
}
-// TODO(dsinclair): Requires `ArrayStride` decoration support
-TEST_F(SpirvReaderTest, DISABLED_ClipDistances) {
+TEST_F(SpirvReaderTest, ClipDistances) {
auto got = Run(R"(
OpCapability Shader
OpCapability ClipDistance
