Remove support for the @stride attribute
We still use the StrideAttribute AST node in the SPIR-V reader for
strided arrays and matrices, which are then removed by transforms.
The WGSL parser no longer has to handle attributes on types.
Bug: tint:1381
Change-Id: Ifa39575ce207d3fdfcbef7125fe6a3686fad5f20
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/83963
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/docs/tint/origin-trial-changes.md b/docs/tint/origin-trial-changes.md
index 2371a83..ff3c4f6 100644
--- a/docs/tint/origin-trial-changes.md
+++ b/docs/tint/origin-trial-changes.md
@@ -5,6 +5,7 @@
### Breaking changes
* The `@block` attribute has been removed. [tint:1324](crbug.com/tint/1324)
+* The `@stride` attribute has been removed. [tint:1381](crbug.com/tint/1381)
* Attributes using `[[attribute]]` syntax are no longer supported. [tint:1382](crbug.com/tint/1382)
## Changes for M101
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 54578e7..91bee00 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -467,8 +467,6 @@
"transform/multiplanar_external_texture.h",
"transform/num_workgroups_from_uniform.cc",
"transform/num_workgroups_from_uniform.h",
- "transform/pad_array_elements.cc",
- "transform/pad_array_elements.h",
"transform/promote_initializers_to_const_var.cc",
"transform/promote_initializers_to_const_var.h",
"transform/promote_side_effects_to_decl.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 011fc39..5e077e6 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -345,8 +345,6 @@
transform/multiplanar_external_texture.h
transform/num_workgroups_from_uniform.cc
transform/num_workgroups_from_uniform.h
- transform/pad_array_elements.cc
- transform/pad_array_elements.h
transform/promote_initializers_to_const_var.cc
transform/promote_initializers_to_const_var.h
transform/promote_side_effects_to_decl.cc
@@ -1023,7 +1021,6 @@
transform/module_scope_var_to_entry_point_param_test.cc
transform/multiplanar_external_texture_test.cc
transform/num_workgroups_from_uniform_test.cc
- transform/pad_array_elements_test.cc
transform/promote_initializers_to_const_var_test.cc
transform/promote_side_effects_to_decl_test.cc
transform/remove_phonies_test.cc
diff --git a/src/tint/ast/module_clone_test.cc b/src/tint/ast/module_clone_test.cc
index 12853ba..983d1b0 100644
--- a/src/tint/ast/module_clone_test.cc
+++ b/src/tint/ast/module_clone_test.cc
@@ -41,7 +41,7 @@
let c0 : i32 = 10;
let c1 : bool = true;
-type t0 = @stride(16) array<vec4<f32>>;
+type t0 = array<vec4<f32>>;
type t1 = array<vec4<f32>>;
var<private> g0 : u32 = 20u;
diff --git a/src/tint/ast/stride_attribute.h b/src/tint/ast/stride_attribute.h
index 4edc54c..0aa3baf 100644
--- a/src/tint/ast/stride_attribute.h
+++ b/src/tint/ast/stride_attribute.h
@@ -18,11 +18,13 @@
#include <string>
#include "src/tint/ast/attribute.h"
+#include "src/tint/ast/internal_attribute.h"
namespace tint {
namespace ast {
-/// A stride attribute
+/// A stride attribute used by the SPIR-V reader for strided arrays and
+/// matrices.
class StrideAttribute final : public Castable<StrideAttribute, Attribute> {
public:
/// constructor
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index eaf5fb7..0f62a2f 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -119,7 +119,6 @@
const char kSizeAttribute[] = "size";
const char kAlignAttribute[] = "align";
const char kStageAttribute[] = "stage";
-const char kStrideAttribute[] = "stride";
const char kWorkgroupSizeAttribute[] = "workgroup_size";
// https://gpuweb.github.io/gpuweb/wgsl.html#reserved-keywords
@@ -799,7 +798,7 @@
}
// variable_ident_decl
-// : IDENT COLON variable_attribute_list* type_decl
+// : IDENT COLON type_decl
Expect<ParserImpl::TypedIdentifier> ParserImpl::expect_variable_ident_decl(
std::string_view use,
bool allow_inferred) {
@@ -814,20 +813,13 @@
if (!expect(use, Token::Type::kColon))
return Failure::kErrored;
- auto attrs = attribute_list();
- if (attrs.errored)
- return Failure::kErrored;
-
auto t = peek();
- auto type = type_decl(attrs.value);
+ auto type = type_decl();
if (type.errored)
return Failure::kErrored;
if (!type.matched)
return add_error(t.source(), "invalid type", use);
- if (!expect_attributes_consumed(attrs.value))
- return Failure::kErrored;
-
return TypedIdentifier{type.value, ident.value, ident.source};
}
@@ -929,25 +921,6 @@
// | MAT4x4 LESS_THAN type_decl GREATER_THAN
// | texture_sampler_types
Maybe<const ast::Type*> ParserImpl::type_decl() {
- auto attrs = attribute_list();
- if (attrs.errored)
- return Failure::kErrored;
-
- auto type = type_decl(attrs.value);
- if (type.errored) {
- return Failure::kErrored;
- }
- if (!expect_attributes_consumed(attrs.value)) {
- return Failure::kErrored;
- }
- if (!type.matched) {
- return Failure::kNoMatch;
- }
-
- return type;
-}
-
-Maybe<const ast::Type*> ParserImpl::type_decl(ast::AttributeList& attrs) {
auto t = peek();
Source source;
if (match(Token::Type::kIdentifier, &source)) {
@@ -981,7 +954,7 @@
}
if (match(Token::Type::kArray, &source)) {
- return expect_type_decl_array(t, std::move(attrs));
+ return expect_type_decl_array(t);
}
if (t.IsMatrix()) {
@@ -1080,9 +1053,7 @@
return builder_.ty.vec(make_source_range_from(t.source()), subtype, count);
}
-Expect<const ast::Type*> ParserImpl::expect_type_decl_array(
- Token t,
- ast::AttributeList attrs) {
+Expect<const ast::Type*> ParserImpl::expect_type_decl_array(Token t) {
const char* use = "array declaration";
const ast::Expression* size = nullptr;
@@ -1111,7 +1082,7 @@
}
return builder_.ty.array(make_source_range_from(t.source()), subtype.value,
- size, std::move(attrs));
+ size);
}
Expect<const ast::Type*> ParserImpl::expect_type_decl_matrix(Token t) {
@@ -1321,19 +1292,7 @@
}
return_attributes = attrs.value;
- // Apply stride attributes to the type node instead of the function.
- ast::AttributeList type_attributes;
- auto itr =
- std::find_if(return_attributes.begin(), return_attributes.end(),
- [](auto* attr) { return Is<ast::StrideAttribute>(attr); });
- if (itr != return_attributes.end()) {
- type_attributes.emplace_back(*itr);
- return_attributes.erase(itr);
- }
-
- auto tok = peek();
-
- auto type = type_decl(type_attributes);
+ auto type = type_decl();
if (type.errored) {
errored = true;
} else if (!type.matched) {
@@ -2974,19 +2933,6 @@
});
}
- if (t == kStrideAttribute) {
- const char* use = "stride attribute";
- return expect_paren_block(use, [&]() -> Result {
- auto val = expect_nonzero_positive_sint(use);
- if (val.errored)
- return Failure::kErrored;
- deprecated(t.source(),
- "the @stride attribute is deprecated; use a larger type if "
- "necessary");
- return create<ast::StrideAttribute>(t.source(), val.value);
- });
- }
-
if (t == kSizeAttribute) {
const char* use = "size attribute";
return expect_paren_block(use, [&]() -> Result {
diff --git a/src/tint/reader/wgsl/parser_impl.h b/src/tint/reader/wgsl/parser_impl.h
index 50344e7..7cce8fd 100644
--- a/src/tint/reader/wgsl/parser_impl.h
+++ b/src/tint/reader/wgsl/parser_impl.h
@@ -419,11 +419,6 @@
/// Parses a `type_decl` grammar element
/// @returns the parsed Type or nullptr if none matched.
Maybe<const ast::Type*> type_decl();
- /// Parses a `type_decl` grammar element with the given pre-parsed
- /// attributes.
- /// @param attrs the list of attributes for the type.
- /// @returns the parsed Type or nullptr if none matched.
- Maybe<const ast::Type*> type_decl(ast::AttributeList& attrs);
/// Parses a `storage_class` grammar element, erroring on parse failure.
/// @param use a description of what was being parsed if an error was raised.
/// @returns the storage class or StorageClass::kNone if none matched
@@ -849,8 +844,7 @@
Expect<const ast::Type*> expect_type_decl_pointer(Token t);
Expect<const ast::Type*> expect_type_decl_atomic(Token t);
Expect<const ast::Type*> expect_type_decl_vector(Token t);
- Expect<const ast::Type*> expect_type_decl_array(Token t,
- ast::AttributeList attrs);
+ Expect<const ast::Type*> expect_type_decl_array(Token t);
Expect<const ast::Type*> expect_type_decl_matrix(Token t);
Expect<const ast::Type*> expect_type(std::string_view use);
diff --git a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
index 7c21634..f68c973 100644
--- a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
@@ -445,9 +445,9 @@
TEST_F(ParserImplErrorTest, FunctionScopeUnusedDecl) {
EXPECT("fn f(a:i32)->i32{return a;@size(1)}",
- R"(test.wgsl:1:28 error: unexpected attributes
+ R"(test.wgsl:1:27 error: expected '}'
fn f(a:i32)->i32{return a;@size(1)}
- ^^^^
+ ^
)");
}
@@ -774,31 +774,6 @@
)");
}
-TEST_F(ParserImplErrorTest, GlobalDeclVarArrayAttrNotArray) {
- EXPECT("var i : @location(1) i32;",
- R"(test.wgsl:1:10 error: unexpected attributes
-var i : @location(1) i32;
- ^^^^^^^^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclVarArrayStrideInvalid) {
- EXPECT(
- "var i : @stride(x) array<i32>;",
- R"(test.wgsl:1:17 error: expected signed integer literal for stride attribute
-var i : @stride(x) array<i32>;
- ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclVarArrayStrideNegative) {
- EXPECT("var i : @stride(-1) array<i32>;",
- R"(test.wgsl:1:17 error: stride attribute must be greater than 0
-var i : @stride(-1) array<i32>;
- ^^
-)");
-}
-
TEST_F(ParserImplErrorTest, GlobalDeclVarArrayMissingType) {
EXPECT("var i : array<1, 3>;",
R"(test.wgsl:1:15 error: invalid type for array declaration
diff --git a/src/tint/reader/wgsl/parser_impl_function_header_test.cc b/src/tint/reader/wgsl/parser_impl_function_header_test.cc
index 45e4d5d..eff2184 100644
--- a/src/tint/reader/wgsl/parser_impl_function_header_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_header_test.cc
@@ -75,27 +75,6 @@
EXPECT_TRUE(f->return_type_attributes[0]->Is<ast::InvariantAttribute>());
}
-TEST_F(ParserImplTest, FunctionHeader_AttributeReturnType_WithArrayStride) {
- auto p = parser("fn main() -> @location(1) @stride(16) array<f32, 4>");
- auto f = p->function_header();
- ASSERT_FALSE(p->has_error()) << p->error();
- EXPECT_TRUE(f.matched);
- EXPECT_FALSE(f.errored);
-
- EXPECT_EQ(f->name, "main");
- EXPECT_EQ(f->params.size(), 0u);
- ASSERT_EQ(f->return_type_attributes.size(), 1u);
- auto* loc = f->return_type_attributes[0]->As<ast::LocationAttribute>();
- ASSERT_TRUE(loc != nullptr);
- EXPECT_EQ(loc->value, 1u);
-
- auto* array_type = f->return_type->As<ast::Array>();
- ASSERT_EQ(array_type->attributes.size(), 1u);
- auto* stride = array_type->attributes[0]->As<ast::StrideAttribute>();
- ASSERT_TRUE(stride != nullptr);
- EXPECT_EQ(stride->stride, 16u);
-}
-
TEST_F(ParserImplTest, FunctionHeader_MissingIdent) {
auto p = parser("fn ()");
auto f = p->function_header();
diff --git a/src/tint/reader/wgsl/parser_impl_global_decl_test.cc b/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
index 964f498..29d29a6 100644
--- a/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
@@ -166,33 +166,6 @@
EXPECT_EQ(str->members.size(), 2u);
}
-TEST_F(ParserImplTest, GlobalDecl_Struct_WithStride) {
- auto p = parser("struct A { data: @stride(4) array<f32>; }");
-
- p->expect_global_decl();
- ASSERT_FALSE(p->has_error()) << p->error();
-
- auto program = p->program();
- ASSERT_EQ(program.AST().TypeDecls().size(), 1u);
-
- auto* t = program.AST().TypeDecls()[0];
- ASSERT_NE(t, nullptr);
- ASSERT_TRUE(t->Is<ast::Struct>());
-
- auto* str = t->As<ast::Struct>();
- EXPECT_EQ(str->name, program.Symbols().Get("A"));
- EXPECT_EQ(str->members.size(), 1u);
-
- const auto* ty = str->members[0]->type;
- ASSERT_TRUE(ty->Is<ast::Array>());
- const auto* arr = ty->As<ast::Array>();
-
- ASSERT_EQ(arr->attributes.size(), 1u);
- auto* stride = arr->attributes[0];
- ASSERT_TRUE(stride->Is<ast::StrideAttribute>());
- ASSERT_EQ(stride->As<ast::StrideAttribute>()->stride, 4u);
-}
-
TEST_F(ParserImplTest, GlobalDecl_Struct_Invalid) {
auto p = parser("A {}");
p->expect_global_decl();
diff --git a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
index d74ddf6..5845c40 100644
--- a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
@@ -439,172 +439,6 @@
EXPECT_EQ(p->builder().Symbols().NameFor(count_expr->symbol), "size");
}
-TEST_F(ParserImplTest, TypeDecl_Array_Stride) {
- auto p = parser("@stride(16) array<f32, 5>");
- auto t = p->type_decl();
- EXPECT_TRUE(t.matched);
- EXPECT_FALSE(t.errored);
- ASSERT_NE(t.value, nullptr) << p->error();
- ASSERT_FALSE(p->has_error());
- ASSERT_TRUE(t.value->Is<ast::Array>());
-
- auto* a = t.value->As<ast::Array>();
- ASSERT_FALSE(a->IsRuntimeArray());
- ASSERT_TRUE(a->type->Is<ast::F32>());
-
- auto* size = a->count->As<ast::SintLiteralExpression>();
- ASSERT_NE(size, nullptr);
- EXPECT_EQ(size->ValueAsI32(), 5);
-
- ASSERT_EQ(a->attributes.size(), 1u);
- auto* stride = a->attributes[0];
- ASSERT_TRUE(stride->Is<ast::StrideAttribute>());
- ASSERT_EQ(stride->As<ast::StrideAttribute>()->stride, 16u);
- EXPECT_EQ(t.value->source.range, (Source::Range{{1u, 13u}, {1u, 26u}}));
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Runtime_Stride) {
- auto p = parser("@stride(16) array<f32>");
- auto t = p->type_decl();
- EXPECT_TRUE(t.matched);
- EXPECT_FALSE(t.errored);
- ASSERT_NE(t.value, nullptr) << p->error();
- ASSERT_FALSE(p->has_error());
- ASSERT_TRUE(t.value->Is<ast::Array>());
-
- auto* a = t.value->As<ast::Array>();
- ASSERT_TRUE(a->IsRuntimeArray());
- ASSERT_TRUE(a->type->Is<ast::F32>());
-
- ASSERT_EQ(a->attributes.size(), 1u);
- auto* stride = a->attributes[0];
- ASSERT_TRUE(stride->Is<ast::StrideAttribute>());
- ASSERT_EQ(stride->As<ast::StrideAttribute>()->stride, 16u);
- EXPECT_EQ(t.value->source.range, (Source::Range{{1u, 13u}, {1u, 23u}}));
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_MultipleAttributes_OneBlock) {
- auto p = parser("@stride(16) @stride(32) array<f32>");
- auto t = p->type_decl();
- EXPECT_TRUE(t.matched);
- EXPECT_FALSE(t.errored);
- ASSERT_NE(t.value, nullptr) << p->error();
- ASSERT_FALSE(p->has_error());
- ASSERT_TRUE(t.value->Is<ast::Array>());
-
- auto* a = t.value->As<ast::Array>();
- ASSERT_TRUE(a->IsRuntimeArray());
- ASSERT_TRUE(a->type->Is<ast::F32>());
-
- auto& attrs = a->attributes;
- ASSERT_EQ(attrs.size(), 2u);
- EXPECT_TRUE(attrs[0]->Is<ast::StrideAttribute>());
- EXPECT_EQ(attrs[0]->As<ast::StrideAttribute>()->stride, 16u);
- EXPECT_TRUE(attrs[1]->Is<ast::StrideAttribute>());
- EXPECT_EQ(attrs[1]->As<ast::StrideAttribute>()->stride, 32u);
- EXPECT_EQ(t.value->source.range, (Source::Range{{1u, 25u}, {1u, 35u}}));
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_MultipleAttributes_MultipleBlocks) {
- auto p = parser("@stride(16) @stride(32) array<f32>");
- auto t = p->type_decl();
- EXPECT_TRUE(t.matched);
- EXPECT_FALSE(t.errored);
- ASSERT_NE(t.value, nullptr) << p->error();
- ASSERT_FALSE(p->has_error());
- ASSERT_TRUE(t.value->Is<ast::Array>());
-
- auto* a = t.value->As<ast::Array>();
- ASSERT_TRUE(a->IsRuntimeArray());
- ASSERT_TRUE(a->type->Is<ast::F32>());
-
- auto& attrs = a->attributes;
- ASSERT_EQ(attrs.size(), 2u);
- EXPECT_TRUE(attrs[0]->Is<ast::StrideAttribute>());
- EXPECT_EQ(attrs[0]->As<ast::StrideAttribute>()->stride, 16u);
- EXPECT_TRUE(attrs[1]->Is<ast::StrideAttribute>());
- EXPECT_EQ(attrs[1]->As<ast::StrideAttribute>()->stride, 32u);
- EXPECT_EQ(t.value->source.range, (Source::Range{{1u, 25u}, {1u, 35u}}));
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Attribute_MissingArray) {
- auto p = parser("@stride(16) f32");
- auto t = p->type_decl();
- EXPECT_TRUE(t.errored);
- EXPECT_FALSE(t.matched);
- ASSERT_EQ(t.value, nullptr);
- ASSERT_TRUE(p->has_error());
- EXPECT_EQ(
- p->error(),
- R"(1:2: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
-1:2: unexpected attributes)");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Attribute_UnknownAttribute) {
- auto p = parser("@unknown(16) array<f32, 5>");
- auto t = p->type_decl();
- EXPECT_TRUE(t.errored);
- EXPECT_FALSE(t.matched);
- ASSERT_EQ(t.value, nullptr);
- ASSERT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), R"(1:2: expected attribute)");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingLeftParen) {
- auto p = parser("@stride 4) array<f32, 5>");
- auto t = p->type_decl();
- EXPECT_TRUE(t.errored);
- EXPECT_FALSE(t.matched);
- ASSERT_EQ(t.value, nullptr);
- ASSERT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), R"(1:9: expected '(' for stride attribute)");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingRightParen) {
- auto p = parser("@stride(4 array<f32, 5>");
- auto t = p->type_decl();
- EXPECT_TRUE(t.errored);
- EXPECT_FALSE(t.matched);
- ASSERT_EQ(t.value, nullptr);
- ASSERT_TRUE(p->has_error());
- EXPECT_EQ(
- p->error(),
- R"(1:2: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
-1:11: expected ')' for stride attribute)");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingValue) {
- auto p = parser("@stride() array<f32, 5>");
- auto t = p->type_decl();
- EXPECT_TRUE(t.errored);
- EXPECT_FALSE(t.matched);
- ASSERT_EQ(t.value, nullptr);
- ASSERT_TRUE(p->has_error());
- EXPECT_EQ(p->error(),
- "1:9: expected signed integer literal for stride attribute");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Stride_InvalidValue) {
- auto p = parser("@stride(invalid) array<f32, 5>");
- auto t = p->type_decl();
- EXPECT_TRUE(t.errored);
- EXPECT_FALSE(t.matched);
- ASSERT_EQ(t.value, nullptr);
- ASSERT_TRUE(p->has_error());
- EXPECT_EQ(p->error(),
- "1:9: expected signed integer literal for stride attribute");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Stride_InvalidValue_Negative) {
- auto p = parser("@stride(-1) array<f32, 5>");
- auto t = p->type_decl();
- EXPECT_TRUE(t.errored);
- EXPECT_FALSE(t.matched);
- ASSERT_EQ(t.value, nullptr);
- ASSERT_TRUE(p->has_error());
- EXPECT_EQ(p->error(), "1:9: stride attribute must be greater than 0");
-}
-
TEST_F(ParserImplTest, TypeDecl_Array_Runtime) {
auto p = parser("array<u32>");
auto t = p->type_decl();
diff --git a/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc b/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
index 976284e..6f07375 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
@@ -75,35 +75,6 @@
ASSERT_EQ(p->error(), "1:1: expected identifier for test");
}
-TEST_F(ParserImplTest, VariableIdentDecl_NonAccessAttrFail) {
- auto p = parser("my_var : @location(1) S");
-
- auto* mem = Member("a", ty.i32(), ast::AttributeList{});
- ast::StructMemberList members;
- members.push_back(mem);
-
- auto decl = p->expect_variable_ident_decl("test");
- ASSERT_TRUE(p->has_error());
- ASSERT_TRUE(decl.errored);
- ASSERT_EQ(p->error(), "1:11: unexpected attributes");
-}
-
-TEST_F(ParserImplTest, VariableIdentDecl_AttributeMissingRightParen) {
- auto p = parser("my_var : @location(4 S");
- auto decl = p->expect_variable_ident_decl("test");
- ASSERT_TRUE(p->has_error());
- ASSERT_TRUE(decl.errored);
- ASSERT_EQ(p->error(), "1:22: expected ')' for location attribute");
-}
-
-TEST_F(ParserImplTest, VariableIdentDecl_AttributeMissingLeftParen) {
- auto p = parser("my_var : @stride 4) S");
- auto decl = p->expect_variable_ident_decl("test");
- ASSERT_TRUE(p->has_error());
- ASSERT_TRUE(decl.errored);
- ASSERT_EQ(p->error(), "1:18: expected '(' for stride attribute");
-}
-
} // namespace
} // namespace wgsl
} // namespace reader
diff --git a/src/tint/transform/array_length_from_uniform_test.cc b/src/tint/transform/array_length_from_uniform_test.cc
index ffb63af..8bd6659 100644
--- a/src/tint/transform/array_length_from_uniform_test.cc
+++ b/src/tint/transform/array_length_from_uniform_test.cc
@@ -177,94 +177,6 @@
got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
}
-TEST_F(ArrayLengthFromUniformTest, WithStride) {
- auto* src = R"(
-@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
- var len : u32 = arrayLength(&sb);
-}
-)";
-
- auto* expect = R"(
-struct tint_symbol {
- buffer_size : array<vec4<u32>, 1u>;
-}
-
-@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
-
-@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
- var len : u32 = (tint_symbol_1.buffer_size[0u][0u] / 64u);
-}
-)";
-
- ArrayLengthFromUniform::Config cfg({0, 30u});
- cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
-
- DataMap data;
- data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
-
- auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
-
- EXPECT_EQ(expect, str(got));
- EXPECT_EQ(std::unordered_set<uint32_t>({0}),
- got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
-}
-
-TEST_F(ArrayLengthFromUniformTest, WithStride_InStruct) {
- auto* src = R"(
-struct SB {
- x : i32;
- y : f32;
- arr : @stride(64) array<i32>;
-};
-
-@group(0) @binding(0) var<storage, read> sb : SB;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
- var len : u32 = arrayLength(&sb.arr);
-}
-)";
-
- auto* expect = R"(
-struct tint_symbol {
- buffer_size : array<vec4<u32>, 1u>;
-}
-
-@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
-
-struct SB {
- x : i32;
- y : f32;
- arr : @stride(64) array<i32>;
-}
-
-@group(0) @binding(0) var<storage, read> sb : SB;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
- var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 8u) / 64u);
-}
-)";
-
- ArrayLengthFromUniform::Config cfg({0, 30u});
- cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
-
- DataMap data;
- data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
-
- auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
-
- EXPECT_EQ(expect, str(got));
- EXPECT_EQ(std::unordered_set<uint32_t>({0}),
- got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
-}
-
TEST_F(ArrayLengthFromUniformTest, MultipleStorageBuffers) {
auto* src = R"(
struct SB1 {
diff --git a/src/tint/transform/calculate_array_length_test.cc b/src/tint/transform/calculate_array_length_test.cc
index 64589a0..c2ecbc9 100644
--- a/src/tint/transform/calculate_array_length_test.cc
+++ b/src/tint/transform/calculate_array_length_test.cc
@@ -287,78 +287,6 @@
EXPECT_EQ(expect, str(got));
}
-TEST_F(CalculateArrayLengthTest, WithStride) {
- auto* src = R"(
-@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
- var len : u32 = arrayLength(&sb);
-}
-)";
-
- auto* expect = R"(
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : @stride(64) array<i32>, result : ptr<function, u32>)
-
-@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
- var tint_symbol_1 : u32 = 0u;
- tint_symbol(sb, &(tint_symbol_1));
- let tint_symbol_2 : u32 = (tint_symbol_1 / 64u);
- var len : u32 = tint_symbol_2;
-}
-)";
-
- auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(CalculateArrayLengthTest, WithStride_InStruct) {
- auto* src = R"(
-struct SB {
- x : i32;
- y : f32;
- arr : @stride(64) array<i32>;
-};
-
-@group(0) @binding(0) var<storage, read> sb : SB;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
- var len : u32 = arrayLength(&sb.arr);
-}
-)";
-
- auto* expect = R"(
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
-struct SB {
- x : i32;
- y : f32;
- arr : @stride(64) array<i32>;
-}
-
-@group(0) @binding(0) var<storage, read> sb : SB;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
- var tint_symbol_1 : u32 = 0u;
- tint_symbol(sb, &(tint_symbol_1));
- let tint_symbol_2 : u32 = ((tint_symbol_1 - 8u) / 64u);
- var len : u32 = tint_symbol_2;
-}
-)";
-
- auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
TEST_F(CalculateArrayLengthTest, Nested) {
auto* src = R"(
struct SB {
diff --git a/src/tint/transform/decompose_memory_access_test.cc b/src/tint/transform/decompose_memory_access_test.cc
index 8ca3c2b..9d1956b 100644
--- a/src/tint/transform/decompose_memory_access_test.cc
+++ b/src/tint/transform/decompose_memory_access_test.cc
@@ -1970,22 +1970,26 @@
TEST_F(DecomposeMemoryAccessTest, ComplexStaticAccessChain) {
auto* src = R"(
+// sizeof(S1) == 32
+// alignof(S1) == 16
struct S1 {
a : i32;
b : vec3<f32>;
c : i32;
};
+// sizeof(S2) == 116
+// alignof(S2) == 16
struct S2 {
a : i32;
- b : @stride(32) array<S1, 3>;
+ b : array<S1, 3>;
c : i32;
};
struct SB {
@size(128)
a : i32;
- b : @stride(256) array<S2>;
+ b : array<S2>;
};
@group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -1999,9 +2003,9 @@
// sb.b[4].b[1].b.z
// ^ ^ ^ ^ ^ ^
// | | | | | |
- // 128 | |1200| 1224
+ // 128 | |688 | 712
// | | |
- // 1152 1168 1216
+ // 640 656 704
auto* expect = R"(
struct S1 {
@@ -2012,14 +2016,14 @@
struct S2 {
a : i32;
- b : @stride(32) array<S1, 3>;
+ b : array<S1, 3>;
c : i32;
}
struct SB {
@size(128)
a : i32;
- b : @stride(256) array<S2>;
+ b : array<S2>;
}
@group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2029,7 +2033,7 @@
@stage(compute) @workgroup_size(1)
fn main() {
- var x : f32 = tint_symbol(sb, 1224u);
+ var x : f32 = tint_symbol(sb, 712u);
}
)";
@@ -2050,12 +2054,12 @@
struct SB {
@size(128)
a : i32;
- b : @stride(256) array<S2>;
+ b : array<S2>;
};
struct S2 {
a : i32;
- b : @stride(32) array<S1, 3>;
+ b : array<S1, 3>;
c : i32;
};
@@ -2069,9 +2073,9 @@
// sb.b[4].b[1].b.z
// ^ ^ ^ ^ ^ ^
// | | | | | |
- // 128 | |1200| 1224
+ // 128 | |688 | 712
// | | |
- // 1152 1168 1216
+ // 640 656 704
auto* expect = R"(
@internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body)
@@ -2079,7 +2083,7 @@
@stage(compute) @workgroup_size(1)
fn main() {
- var x : f32 = tint_symbol(sb, 1224u);
+ var x : f32 = tint_symbol(sb, 712u);
}
@group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2087,12 +2091,12 @@
struct SB {
@size(128)
a : i32;
- b : @stride(256) array<S2>;
+ b : array<S2>;
}
struct S2 {
a : i32;
- b : @stride(32) array<S1, 3>;
+ b : array<S1, 3>;
c : i32;
}
@@ -2118,14 +2122,14 @@
struct S2 {
a : i32;
- b : @stride(32) array<S1, 3>;
+ b : array<S1, 3>;
c : i32;
};
struct SB {
@size(128)
a : i32;
- b : @stride(256) array<S2>;
+ b : array<S2>;
};
@group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2148,14 +2152,14 @@
struct S2 {
a : i32;
- b : @stride(32) array<S1, 3>;
+ b : array<S1, 3>;
c : i32;
}
struct SB {
@size(128)
a : i32;
- b : @stride(256) array<S2>;
+ b : array<S2>;
}
@group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2168,7 +2172,7 @@
var i : i32 = 4;
var j : u32 = 1u;
var k : i32 = 2;
- var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
+ var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
}
)";
@@ -2192,12 +2196,12 @@
struct SB {
@size(128)
a : i32;
- b : @stride(256) array<S2>;
+ b :array<S2>;
};
struct S2 {
a : i32;
- b : @stride(32) array<S1, 3>;
+ b : array<S1, 3>;
c : i32;
};
@@ -2217,7 +2221,7 @@
var i : i32 = 4;
var j : u32 = 1u;
var k : i32 = 2;
- var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
+ var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
}
@group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2225,12 +2229,12 @@
struct SB {
@size(128)
a : i32;
- b : @stride(256) array<S2>;
+ b : array<S2>;
}
struct S2 {
a : i32;
- b : @stride(32) array<S1, 3>;
+ b : array<S1, 3>;
c : i32;
}
@@ -2256,7 +2260,7 @@
type A1 = S1;
-type A1_Array = @stride(32) array<S1, 3>;
+type A1_Array = array<S1, 3>;
struct S2 {
a : i32;
@@ -2266,7 +2270,7 @@
type A2 = S2;
-type A2_Array = @stride(256) array<S2>;
+type A2_Array = array<S2>;
struct SB {
@size(128)
@@ -2294,7 +2298,7 @@
type A1 = S1;
-type A1_Array = @stride(32) array<S1, 3>;
+type A1_Array = array<S1, 3>;
struct S2 {
a : i32;
@@ -2304,7 +2308,7 @@
type A2 = S2;
-type A2_Array = @stride(256) array<S2>;
+type A2_Array = array<S2>;
struct SB {
@size(128)
@@ -2322,7 +2326,7 @@
var i : i32 = 4;
var j : u32 = 1u;
var k : i32 = 2;
- var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
+ var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
}
)";
@@ -2350,7 +2354,7 @@
b : A2_Array;
};
-type A2_Array = @stride(256) array<S2>;
+type A2_Array = array<S2>;
type A2 = S2;
@@ -2362,7 +2366,7 @@
type A1 = S1;
-type A1_Array = @stride(32) array<S1, 3>;
+type A1_Array = array<S1, 3>;
struct S1 {
a : i32;
@@ -2380,7 +2384,7 @@
var i : i32 = 4;
var j : u32 = 1u;
var k : i32 = 2;
- var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
+ var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
}
@group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2391,7 +2395,7 @@
b : A2_Array;
}
-type A2_Array = @stride(256) array<S2>;
+type A2_Array = array<S2>;
type A2 = S2;
@@ -2403,7 +2407,7 @@
type A1 = S1;
-type A1_Array = @stride(32) array<S1, 3>;
+type A1_Array = array<S1, 3>;
struct S1 {
a : i32;
diff --git a/src/tint/transform/glsl.cc b/src/tint/transform/glsl.cc
index 39f81e6..c7352c4 100644
--- a/src/tint/transform/glsl.cc
+++ b/src/tint/transform/glsl.cc
@@ -27,7 +27,6 @@
#include "src/tint/transform/fold_trivial_single_use_lets.h"
#include "src/tint/transform/loop_to_for_loop.h"
#include "src/tint/transform/manager.h"
-#include "src/tint/transform/pad_array_elements.h"
#include "src/tint/transform/promote_initializers_to_const_var.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
#include "src/tint/transform/remove_phonies.h"
@@ -105,7 +104,6 @@
}
manager.Add<PromoteInitializersToConstVar>();
- manager.Add<PadArrayElements>();
manager.Add<AddEmptyEntryPoint>();
manager.Add<AddSpirvBlockAttribute>();
diff --git a/src/tint/transform/pad_array_elements.cc b/src/tint/transform/pad_array_elements.cc
deleted file mode 100644
index bf406df..0000000
--- a/src/tint/transform/pad_array_elements.cc
+++ /dev/null
@@ -1,177 +0,0 @@
-// Copyright 2021 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "src/tint/transform/pad_array_elements.h"
-
-#include <unordered_map>
-#include <utility>
-
-#include "src/tint/program_builder.h"
-#include "src/tint/sem/array.h"
-#include "src/tint/sem/call.h"
-#include "src/tint/sem/expression.h"
-#include "src/tint/sem/type_constructor.h"
-#include "src/tint/utils/map.h"
-
-TINT_INSTANTIATE_TYPEINFO(tint::transform::PadArrayElements);
-
-namespace tint {
-namespace transform {
-namespace {
-
-using ArrayBuilder = std::function<const ast::Array*()>;
-
-/// PadArray returns a function that constructs a new array in `ctx.dst` with
-/// the element type padded to account for the explicit stride. PadArray will
-/// recursively pad arrays-of-arrays. The new array element type will be added
-/// to module-scope type declarations of `ctx.dst`.
-/// @param ctx the CloneContext
-/// @param create_ast_type_for Transform::CreateASTTypeFor()
-/// @param padded_arrays a map of src array type to the new array name
-/// @param array the array type
-/// @return the new AST array
-template <typename CREATE_AST_TYPE_FOR>
-ArrayBuilder PadArray(
- CloneContext& ctx,
- CREATE_AST_TYPE_FOR&& create_ast_type_for,
- std::unordered_map<const sem::Array*, ArrayBuilder>& padded_arrays,
- const sem::Array* array) {
- if (array->IsStrideImplicit()) {
- // We don't want to wrap arrays that have an implicit stride
- return nullptr;
- }
-
- return utils::GetOrCreate(padded_arrays, array, [&] {
- // Generate a unique name for the array element type
- auto name = ctx.dst->Symbols().New("tint_padded_array_element");
-
- // Examine the element type. Is it also an array?
- const ast::Type* el_ty = nullptr;
- if (auto* el_array = array->ElemType()->As<sem::Array>()) {
- // Array of array - call PadArray() on the element type
- if (auto p =
- PadArray(ctx, create_ast_type_for, padded_arrays, el_array)) {
- el_ty = p();
- }
- }
-
- // If the element wasn't a padded array, just create the typical AST type
- // for it
- if (el_ty == nullptr) {
- el_ty = create_ast_type_for(ctx, array->ElemType());
- }
-
- // Structure() will create and append the ast::Struct to the
- // global declarations of `ctx.dst`. As we haven't finished building the
- // current module-scope statement or function, this will be placed
- // immediately before the usage.
- ctx.dst->Structure(
- name,
- {ctx.dst->Member("el", el_ty, {ctx.dst->MemberSize(array->Stride())})});
-
- auto* dst = ctx.dst;
- return [=] {
- if (array->IsRuntimeSized()) {
- return dst->ty.array(dst->create<ast::TypeName>(name));
- } else {
- return dst->ty.array(dst->create<ast::TypeName>(name), array->Count());
- }
- };
- });
-}
-
-} // namespace
-
-PadArrayElements::PadArrayElements() = default;
-
-PadArrayElements::~PadArrayElements() = default;
-
-bool PadArrayElements::ShouldRun(const Program* program, const DataMap&) const {
- for (auto* node : program->ASTNodes().Objects()) {
- if (auto* var = node->As<ast::Type>()) {
- if (auto* arr = program->Sem().Get<sem::Array>(var)) {
- if (!arr->IsStrideImplicit()) {
- return true;
- }
- }
- }
- }
- return false;
-}
-
-void PadArrayElements::Run(CloneContext& ctx, const DataMap&, DataMap&) const {
- auto& sem = ctx.src->Sem();
-
- std::unordered_map<const sem::Array*, ArrayBuilder> padded_arrays;
- auto pad = [&](const sem::Array* array) {
- return PadArray(ctx, CreateASTTypeFor, padded_arrays, array);
- };
-
- // Replace all array types with their corresponding padded array type
- ctx.ReplaceAll([&](const ast::Type* ast_type) -> const ast::Type* {
- auto* type = ctx.src->TypeOf(ast_type);
- if (auto* array = type->UnwrapRef()->As<sem::Array>()) {
- if (auto p = pad(array)) {
- return p();
- }
- }
- return nullptr;
- });
-
- // Fix up index accessors so `a[1]` becomes `a[1].el`
- ctx.ReplaceAll([&](const ast::IndexAccessorExpression* accessor)
- -> const ast::Expression* {
- if (auto* array = tint::As<sem::Array>(
- sem.Get(accessor->object)->Type()->UnwrapRef())) {
- if (pad(array)) {
- // Array element is wrapped in a structure. Emit a member accessor
- // to get to the actual array element.
- auto* idx = ctx.CloneWithoutTransform(accessor);
- return ctx.dst->MemberAccessor(idx, "el");
- }
- }
- return nullptr;
- });
-
- // Fix up array constructors so `A(1,2)` becomes
- // `A(padded(1), padded(2))`
- ctx.ReplaceAll(
- [&](const ast::CallExpression* expr) -> const ast::Expression* {
- auto* call = sem.Get(expr);
- if (auto* ctor = call->Target()->As<sem::TypeConstructor>()) {
- if (auto* array = ctor->ReturnType()->As<sem::Array>()) {
- if (auto p = pad(array)) {
- auto* arr_ty = p();
- auto el_typename = arr_ty->type->As<ast::TypeName>()->name;
-
- ast::ExpressionList args;
- args.reserve(call->Arguments().size());
- for (auto* arg : call->Arguments()) {
- auto* val = ctx.Clone(arg->Declaration());
- args.emplace_back(ctx.dst->Construct(
- ctx.dst->create<ast::TypeName>(el_typename), val));
- }
-
- return ctx.dst->Construct(arr_ty, args);
- }
- }
- }
- return nullptr;
- });
-
- ctx.Clone();
-}
-
-} // namespace transform
-} // namespace tint
diff --git a/src/tint/transform/pad_array_elements.h b/src/tint/transform/pad_array_elements.h
deleted file mode 100644
index cf5bfee..0000000
--- a/src/tint/transform/pad_array_elements.h
+++ /dev/null
@@ -1,62 +0,0 @@
-// Copyright 2021 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef SRC_TINT_TRANSFORM_PAD_ARRAY_ELEMENTS_H_
-#define SRC_TINT_TRANSFORM_PAD_ARRAY_ELEMENTS_H_
-
-#include "src/tint/transform/transform.h"
-
-namespace tint {
-namespace transform {
-
-/// PadArrayElements is a transform that replaces array types with an explicit
-/// stride that is larger than the implicit stride, with an array of a new
-/// structure type. This structure holds with a single field of the element
-/// type, decorated with a `@size` attribute to pad the structure to the
-/// required array stride. The new array types have no explicit stride,
-/// structure size is equal to the desired stride.
-/// Array index expressions and constructors are also adjusted to deal with this
-/// structure element type.
-/// This transform helps with backends that cannot directly return arrays or use
-/// them as parameters.
-class PadArrayElements : public Castable<PadArrayElements, Transform> {
- public:
- /// Constructor
- PadArrayElements();
-
- /// Destructor
- ~PadArrayElements() override;
-
- /// @param program the program to inspect
- /// @param data optional extra transform-specific input data
- /// @returns true if this transform should be run for the given program
- bool ShouldRun(const Program* program,
- const DataMap& data = {}) const override;
-
- protected:
- /// Runs the transform using the CloneContext built for transforming a
- /// program. Run() is responsible for calling Clone() on the CloneContext.
- /// @param ctx the CloneContext primed with the input program and
- /// ProgramBuilder
- /// @param inputs optional extra transform-specific input data
- /// @param outputs optional extra transform-specific output data
- void Run(CloneContext& ctx,
- const DataMap& inputs,
- DataMap& outputs) const override;
-};
-
-} // namespace transform
-} // namespace tint
-
-#endif // SRC_TINT_TRANSFORM_PAD_ARRAY_ELEMENTS_H_
diff --git a/src/tint/transform/pad_array_elements_test.cc b/src/tint/transform/pad_array_elements_test.cc
deleted file mode 100644
index f5e921b..0000000
--- a/src/tint/transform/pad_array_elements_test.cc
+++ /dev/null
@@ -1,518 +0,0 @@
-// Copyright 2021 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-// http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "src/tint/transform/pad_array_elements.h"
-
-#include <utility>
-
-#include "src/tint/transform/test_helper.h"
-
-namespace tint {
-namespace transform {
-namespace {
-
-using PadArrayElementsTest = TransformTest;
-
-TEST_F(PadArrayElementsTest, ShouldRunEmptyModule) {
- auto* src = R"()";
-
- EXPECT_FALSE(ShouldRun<PadArrayElements>(src));
-}
-
-TEST_F(PadArrayElementsTest, ShouldRunHasImplicitArrayStride) {
- auto* src = R"(
-var<private> arr : array<i32, 4>;
-)";
-
- EXPECT_FALSE(ShouldRun<PadArrayElements>(src));
-}
-
-TEST_F(PadArrayElementsTest, ShouldRunHasExplicitArrayStride) {
- auto* src = R"(
-var<private> arr : @stride(8) array<i32, 4>;
-)";
-
- EXPECT_TRUE(ShouldRun<PadArrayElements>(src));
-}
-
-TEST_F(PadArrayElementsTest, EmptyModule) {
- auto* src = "";
- auto* expect = "";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ImplicitArrayStride) {
- auto* src = R"(
-var<private> arr : array<i32, 4>;
-)";
- auto* expect = src;
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArrayAsGlobal) {
- auto* src = R"(
-var<private> arr : @stride(8) array<i32, 4>;
-)";
- auto* expect = R"(
-struct tint_padded_array_element {
- @size(8)
- el : i32;
-}
-
-var<private> arr : array<tint_padded_array_element, 4u>;
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, RuntimeArray) {
- auto* src = R"(
-struct S {
- rta : @stride(8) array<i32>;
-};
-)";
- auto* expect = R"(
-struct tint_padded_array_element {
- @size(8)
- el : i32;
-}
-
-struct S {
- rta : array<tint_padded_array_element>;
-}
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArrayFunctionVar) {
- auto* src = R"(
-fn f() {
- var arr : @stride(16) array<i32, 4>;
- arr = @stride(16) array<i32, 4>();
- arr = @stride(16) array<i32, 4>(1, 2, 3, 4);
- let x = arr[3];
-}
-)";
- auto* expect = R"(
-struct tint_padded_array_element {
- @size(16)
- el : i32;
-}
-
-fn f() {
- var arr : array<tint_padded_array_element, 4u>;
- arr = array<tint_padded_array_element, 4u>();
- arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
- let x = arr[3].el;
-}
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArrayAsParam) {
- auto* src = R"(
-fn f(a : @stride(12) array<i32, 4>) -> i32 {
- return a[2];
-}
-)";
- auto* expect = R"(
-struct tint_padded_array_element {
- @size(12)
- el : i32;
-}
-
-fn f(a : array<tint_padded_array_element, 4u>) -> i32 {
- return a[2].el;
-}
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-// TODO(crbug.com/tint/781): Cannot parse the stride on the return array type.
-TEST_F(PadArrayElementsTest, DISABLED_ArrayAsReturn) {
- auto* src = R"(
-fn f() -> @stride(8) array<i32, 4> {
- return array<i32, 4>(1, 2, 3, 4);
-}
-)";
- auto* expect = R"(
-struct tint_padded_array_element {
- el : i32;
- @size(4)
- padding : u32;
-};
-
-fn f() -> array<tint_padded_array_element, 4> {
- return array<tint_padded_array_element, 4>(tint_padded_array_element(1, 0u), tint_padded_array_element(2, 0u), tint_padded_array_element(3, 0u), tint_padded_array_element(4, 0u));
-}
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArrayAlias) {
- auto* src = R"(
-type Array = @stride(16) array<i32, 4>;
-
-fn f() {
- var arr : Array;
- arr = Array();
- arr = Array(1, 2, 3, 4);
- let vals : Array = Array(1, 2, 3, 4);
- arr = vals;
- let x = arr[3];
-}
-)";
- auto* expect = R"(
-struct tint_padded_array_element {
- @size(16)
- el : i32;
-}
-
-type Array = array<tint_padded_array_element, 4u>;
-
-fn f() {
- var arr : array<tint_padded_array_element, 4u>;
- arr = array<tint_padded_array_element, 4u>();
- arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
- let vals : array<tint_padded_array_element, 4u> = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
- arr = vals;
- let x = arr[3].el;
-}
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArrayAlias_OutOfOrder) {
- auto* src = R"(
-fn f() {
- var arr : Array;
- arr = Array();
- arr = Array(1, 2, 3, 4);
- let vals : Array = Array(1, 2, 3, 4);
- arr = vals;
- let x = arr[3];
-}
-
-type Array = @stride(16) array<i32, 4>;
-)";
- auto* expect = R"(
-struct tint_padded_array_element {
- @size(16)
- el : i32;
-}
-
-fn f() {
- var arr : array<tint_padded_array_element, 4u>;
- arr = array<tint_padded_array_element, 4u>();
- arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
- let vals : array<tint_padded_array_element, 4u> = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
- arr = vals;
- let x = arr[3].el;
-}
-
-type Array = array<tint_padded_array_element, 4u>;
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArraysInStruct) {
- auto* src = R"(
-struct S {
- a : @stride(8) array<i32, 4>;
- b : @stride(8) array<i32, 8>;
- c : @stride(8) array<i32, 4>;
- d : @stride(12) array<i32, 8>;
-};
-)";
- auto* expect = R"(
-struct tint_padded_array_element {
- @size(8)
- el : i32;
-}
-
-struct tint_padded_array_element_1 {
- @size(8)
- el : i32;
-}
-
-struct tint_padded_array_element_2 {
- @size(12)
- el : i32;
-}
-
-struct S {
- a : array<tint_padded_array_element, 4u>;
- b : array<tint_padded_array_element_1, 8u>;
- c : array<tint_padded_array_element, 4u>;
- d : array<tint_padded_array_element_2, 8u>;
-}
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArraysOfArraysInStruct) {
- auto* src = R"(
-struct S {
- a : @stride(512) array<i32, 4>;
- b : @stride(512) array<@stride(32) array<i32, 4>, 4>;
- c : @stride(512) array<@stride(64) array<@stride(8) array<i32, 4>, 4>, 4>;
-};
-)";
- auto* expect = R"(
-struct tint_padded_array_element {
- @size(512)
- el : i32;
-}
-
-struct tint_padded_array_element_2 {
- @size(32)
- el : i32;
-}
-
-struct tint_padded_array_element_1 {
- @size(512)
- el : array<tint_padded_array_element_2, 4u>;
-}
-
-struct tint_padded_array_element_5 {
- @size(8)
- el : i32;
-}
-
-struct tint_padded_array_element_4 {
- @size(64)
- el : array<tint_padded_array_element_5, 4u>;
-}
-
-struct tint_padded_array_element_3 {
- @size(512)
- el : array<tint_padded_array_element_4, 4u>;
-}
-
-struct S {
- a : array<tint_padded_array_element, 4u>;
- b : array<tint_padded_array_element_1, 4u>;
- c : array<tint_padded_array_element_3, 4u>;
-}
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, AccessArraysOfArraysInStruct) {
- auto* src = R"(
-struct S {
- a : @stride(512) array<i32, 4>;
- b : @stride(512) array<@stride(32) array<i32, 4>, 4>;
- c : @stride(512) array<@stride(64) array<@stride(8) array<i32, 4>, 4>, 4>;
-};
-
-fn f(s : S) -> i32 {
- return s.a[2] + s.b[1][2] + s.c[3][1][2];
-}
-)";
- auto* expect = R"(
-struct tint_padded_array_element {
- @size(512)
- el : i32;
-}
-
-struct tint_padded_array_element_2 {
- @size(32)
- el : i32;
-}
-
-struct tint_padded_array_element_1 {
- @size(512)
- el : array<tint_padded_array_element_2, 4u>;
-}
-
-struct tint_padded_array_element_5 {
- @size(8)
- el : i32;
-}
-
-struct tint_padded_array_element_4 {
- @size(64)
- el : array<tint_padded_array_element_5, 4u>;
-}
-
-struct tint_padded_array_element_3 {
- @size(512)
- el : array<tint_padded_array_element_4, 4u>;
-}
-
-struct S {
- a : array<tint_padded_array_element, 4u>;
- b : array<tint_padded_array_element_1, 4u>;
- c : array<tint_padded_array_element_3, 4u>;
-}
-
-fn f(s : S) -> i32 {
- return ((s.a[2].el + s.b[1].el[2].el) + s.c[3].el[1].el[2].el);
-}
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, AccessArraysOfArraysInStruct_OutOfOrder) {
- auto* src = R"(
-fn f(s : S) -> i32 {
- return s.a[2] + s.b[1][2] + s.c[3][1][2];
-}
-
-struct S {
- a : @stride(512) array<i32, 4>;
- b : @stride(512) array<@stride(32) array<i32, 4>, 4>;
- c : @stride(512) array<@stride(64) array<@stride(8) array<i32, 4>, 4>, 4>;
-};
-)";
- auto* expect = R"(
-struct tint_padded_array_element {
- @size(512)
- el : i32;
-}
-
-struct tint_padded_array_element_1 {
- @size(32)
- el : i32;
-}
-
-struct tint_padded_array_element_2 {
- @size(512)
- el : array<tint_padded_array_element_1, 4u>;
-}
-
-struct tint_padded_array_element_3 {
- @size(8)
- el : i32;
-}
-
-struct tint_padded_array_element_4 {
- @size(64)
- el : array<tint_padded_array_element_3, 4u>;
-}
-
-struct tint_padded_array_element_5 {
- @size(512)
- el : array<tint_padded_array_element_4, 4u>;
-}
-
-fn f(s : S) -> i32 {
- return ((s.a[2].el + s.b[1].el[2].el) + s.c[3].el[1].el[2].el);
-}
-
-struct S {
- a : array<tint_padded_array_element, 4u>;
- b : array<tint_padded_array_element_2, 4u>;
- c : array<tint_padded_array_element_5, 4u>;
-}
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, DeclarationOrder) {
- auto* src = R"(
-type T0 = i32;
-
-type T1 = @stride(8) array<i32, 1>;
-
-type T2 = i32;
-
-fn f1(a : @stride(8) array<i32, 2>) {
-}
-
-type T3 = i32;
-
-fn f2() {
- var v : @stride(8) array<i32, 3>;
-}
-)";
- auto* expect = R"(
-type T0 = i32;
-
-struct tint_padded_array_element {
- @size(8)
- el : i32;
-}
-
-type T1 = array<tint_padded_array_element, 1u>;
-
-type T2 = i32;
-
-struct tint_padded_array_element_1 {
- @size(8)
- el : i32;
-}
-
-fn f1(a : array<tint_padded_array_element_1, 2u>) {
-}
-
-type T3 = i32;
-
-struct tint_padded_array_element_2 {
- @size(8)
- el : i32;
-}
-
-fn f2() {
- var v : array<tint_padded_array_element_2, 3u>;
-}
-)";
-
- auto got = Run<PadArrayElements>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-} // namespace
-} // namespace transform
-} // namespace tint
diff --git a/src/tint/transform/robustness_test.cc b/src/tint/transform/robustness_test.cc
index edc33f5..1d19987 100644
--- a/src/tint/transform/robustness_test.cc
+++ b/src/tint/transform/robustness_test.cc
@@ -1431,7 +1431,7 @@
};
@group(0) @binding(0) var<storage, read> s : S;
-type UArr = @stride(16) array<f32, 4>;
+type UArr = array<vec4<f32>, 4>;
struct U {
a : UArr;
};
@@ -1451,11 +1451,11 @@
var i32_sb4 : f32 = s.b[-1];
var i32_sb5 : f32 = s.b[-4];
- var i32_ua1 : f32 = u.a[4];
- var i32_ua2 : f32 = u.a[1];
- var i32_ua3 : f32 = u.a[0];
- var i32_ua4 : f32 = u.a[-1];
- var i32_ua5 : f32 = u.a[-4];
+ var i32_ua1 : f32 = u.a[4].x;
+ var i32_ua2 : f32 = u.a[1].x;
+ var i32_ua3 : f32 = u.a[0].x;
+ var i32_ua4 : f32 = u.a[-1].x;
+ var i32_ua5 : f32 = u.a[-4].x;
// Unsigned
var u32_sa1 : f32 = s.a[0u];
@@ -1472,12 +1472,12 @@
var u32_sb5 : f32 = s.b[10u];
var u32_sb6 : f32 = s.b[100u];
- var u32_ua1 : f32 = u.a[0u];
- var u32_ua2 : f32 = u.a[1u];
- var u32_ua3 : f32 = u.a[3u];
- var u32_ua4 : f32 = u.a[4u];
- var u32_ua5 : f32 = u.a[10u];
- var u32_ua6 : f32 = u.a[100u];
+ var u32_ua1 : f32 = u.a[0u].x;
+ var u32_ua2 : f32 = u.a[1u].x;
+ var u32_ua3 : f32 = u.a[3u].x;
+ var u32_ua4 : f32 = u.a[4u].x;
+ var u32_ua5 : f32 = u.a[10u].x;
+ var u32_ua6 : f32 = u.a[100u].x;
}
)";
@@ -1490,7 +1490,7 @@
@group(0) @binding(0) var<storage, read> s : S;
-type UArr = @stride(16) array<f32, 4>;
+type UArr = array<vec4<f32>, 4>;
struct U {
a : UArr;
@@ -1509,11 +1509,11 @@
var i32_sb3 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
var i32_sb4 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
var i32_sb5 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
- var i32_ua1 : f32 = u.a[3];
- var i32_ua2 : f32 = u.a[1];
- var i32_ua3 : f32 = u.a[0];
- var i32_ua4 : f32 = u.a[0];
- var i32_ua5 : f32 = u.a[0];
+ var i32_ua1 : f32 = u.a[3].x;
+ var i32_ua2 : f32 = u.a[1].x;
+ var i32_ua3 : f32 = u.a[0].x;
+ var i32_ua4 : f32 = u.a[0].x;
+ var i32_ua5 : f32 = u.a[0].x;
var u32_sa1 : f32 = s.a[0u];
var u32_sa2 : f32 = s.a[1u];
var u32_sa3 : f32 = s.a[3u];
@@ -1526,12 +1526,12 @@
var u32_sb4 : f32 = s.b[min(4u, (arrayLength(&(s.b)) - 1u))];
var u32_sb5 : f32 = s.b[min(10u, (arrayLength(&(s.b)) - 1u))];
var u32_sb6 : f32 = s.b[min(100u, (arrayLength(&(s.b)) - 1u))];
- var u32_ua1 : f32 = u.a[0u];
- var u32_ua2 : f32 = u.a[1u];
- var u32_ua3 : f32 = u.a[3u];
- var u32_ua4 : f32 = u.a[3u];
- var u32_ua5 : f32 = u.a[3u];
- var u32_ua6 : f32 = u.a[3u];
+ var u32_ua1 : f32 = u.a[0u].x;
+ var u32_ua2 : f32 = u.a[1u].x;
+ var u32_ua3 : f32 = u.a[3u].x;
+ var u32_ua4 : f32 = u.a[3u].x;
+ var u32_ua5 : f32 = u.a[3u].x;
+ var u32_ua6 : f32 = u.a[3u].x;
}
)";
@@ -1553,7 +1553,7 @@
@group(0) @binding(0) var<storage, read> s : S;
-type UArr = @stride(16) array<f32, 4>;
+type UArr = array<vec4<f32>, 4>;
struct U {
a : UArr;
@@ -1572,11 +1572,11 @@
var i32_sb3 : f32 = s.b[0];
var i32_sb4 : f32 = s.b[-1];
var i32_sb5 : f32 = s.b[-4];
- var i32_ua1 : f32 = u.a[3];
- var i32_ua2 : f32 = u.a[1];
- var i32_ua3 : f32 = u.a[0];
- var i32_ua4 : f32 = u.a[0];
- var i32_ua5 : f32 = u.a[0];
+ var i32_ua1 : f32 = u.a[3].x;
+ var i32_ua2 : f32 = u.a[1].x;
+ var i32_ua3 : f32 = u.a[0].x;
+ var i32_ua4 : f32 = u.a[0].x;
+ var i32_ua5 : f32 = u.a[0].x;
var u32_sa1 : f32 = s.a[0u];
var u32_sa2 : f32 = s.a[1u];
var u32_sa3 : f32 = s.a[3u];
@@ -1589,12 +1589,12 @@
var u32_sb4 : f32 = s.b[4u];
var u32_sb5 : f32 = s.b[10u];
var u32_sb6 : f32 = s.b[100u];
- var u32_ua1 : f32 = u.a[0u];
- var u32_ua2 : f32 = u.a[1u];
- var u32_ua3 : f32 = u.a[3u];
- var u32_ua4 : f32 = u.a[3u];
- var u32_ua5 : f32 = u.a[3u];
- var u32_ua6 : f32 = u.a[3u];
+ var u32_ua1 : f32 = u.a[0u].x;
+ var u32_ua2 : f32 = u.a[1u].x;
+ var u32_ua3 : f32 = u.a[3u].x;
+ var u32_ua4 : f32 = u.a[3u].x;
+ var u32_ua5 : f32 = u.a[3u].x;
+ var u32_ua6 : f32 = u.a[3u].x;
}
)";
@@ -1618,7 +1618,7 @@
@group(0) @binding(0) var<storage, read> s : S;
-type UArr = @stride(16) array<f32, 4>;
+type UArr = array<vec4<f32>, 4>;
struct U {
a : UArr;
@@ -1637,11 +1637,11 @@
var i32_sb3 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
var i32_sb4 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
var i32_sb5 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
- var i32_ua1 : f32 = u.a[4];
- var i32_ua2 : f32 = u.a[1];
- var i32_ua3 : f32 = u.a[0];
- var i32_ua4 : f32 = u.a[-1];
- var i32_ua5 : f32 = u.a[-4];
+ var i32_ua1 : f32 = u.a[4].x;
+ var i32_ua2 : f32 = u.a[1].x;
+ var i32_ua3 : f32 = u.a[0].x;
+ var i32_ua4 : f32 = u.a[-1].x;
+ var i32_ua5 : f32 = u.a[-4].x;
var u32_sa1 : f32 = s.a[0u];
var u32_sa2 : f32 = s.a[1u];
var u32_sa3 : f32 = s.a[3u];
@@ -1654,12 +1654,12 @@
var u32_sb4 : f32 = s.b[min(4u, (arrayLength(&(s.b)) - 1u))];
var u32_sb5 : f32 = s.b[min(10u, (arrayLength(&(s.b)) - 1u))];
var u32_sb6 : f32 = s.b[min(100u, (arrayLength(&(s.b)) - 1u))];
- var u32_ua1 : f32 = u.a[0u];
- var u32_ua2 : f32 = u.a[1u];
- var u32_ua3 : f32 = u.a[3u];
- var u32_ua4 : f32 = u.a[4u];
- var u32_ua5 : f32 = u.a[10u];
- var u32_ua6 : f32 = u.a[100u];
+ var u32_ua1 : f32 = u.a[0u].x;
+ var u32_ua2 : f32 = u.a[1u].x;
+ var u32_ua3 : f32 = u.a[3u].x;
+ var u32_ua4 : f32 = u.a[4u].x;
+ var u32_ua5 : f32 = u.a[10u].x;
+ var u32_ua6 : f32 = u.a[100u].x;
}
)";
@@ -1683,7 +1683,7 @@
@group(0) @binding(0) var<storage, read> s : S;
-type UArr = @stride(16) array<f32, 4>;
+type UArr = array<vec4<f32>, 4>;
struct U {
a : UArr;
@@ -1702,11 +1702,11 @@
var i32_sb3 : f32 = s.b[0];
var i32_sb4 : f32 = s.b[-1];
var i32_sb5 : f32 = s.b[-4];
- var i32_ua1 : f32 = u.a[4];
- var i32_ua2 : f32 = u.a[1];
- var i32_ua3 : f32 = u.a[0];
- var i32_ua4 : f32 = u.a[-1];
- var i32_ua5 : f32 = u.a[-4];
+ var i32_ua1 : f32 = u.a[4].x;
+ var i32_ua2 : f32 = u.a[1].x;
+ var i32_ua3 : f32 = u.a[0].x;
+ var i32_ua4 : f32 = u.a[-1].x;
+ var i32_ua5 : f32 = u.a[-4].x;
var u32_sa1 : f32 = s.a[0u];
var u32_sa2 : f32 = s.a[1u];
var u32_sa3 : f32 = s.a[3u];
@@ -1719,12 +1719,12 @@
var u32_sb4 : f32 = s.b[4u];
var u32_sb5 : f32 = s.b[10u];
var u32_sb6 : f32 = s.b[100u];
- var u32_ua1 : f32 = u.a[0u];
- var u32_ua2 : f32 = u.a[1u];
- var u32_ua3 : f32 = u.a[3u];
- var u32_ua4 : f32 = u.a[4u];
- var u32_ua5 : f32 = u.a[10u];
- var u32_ua6 : f32 = u.a[100u];
+ var u32_ua1 : f32 = u.a[0u].x;
+ var u32_ua2 : f32 = u.a[1u].x;
+ var u32_ua3 : f32 = u.a[3u].x;
+ var u32_ua4 : f32 = u.a[4u].x;
+ var u32_ua5 : f32 = u.a[10u].x;
+ var u32_ua6 : f32 = u.a[100u].x;
}
)";
diff --git a/src/tint/transform/vertex_pulling.cc b/src/tint/transform/vertex_pulling.cc
index e515f04..e0d3128 100644
--- a/src/tint/transform/vertex_pulling.cc
+++ b/src/tint/transform/vertex_pulling.cc
@@ -258,7 +258,7 @@
ctx.dst->Symbols().New(kStructName),
{
ctx.dst->Member(GetStructBufferName(),
- ctx.dst->ty.array<ProgramBuilder::u32>(4)),
+ ctx.dst->ty.array<ProgramBuilder::u32>()),
});
for (uint32_t i = 0; i < cfg.vertex_state.size(); ++i) {
// The decorated variable with struct type
diff --git a/src/tint/transform/vertex_pulling_test.cc b/src/tint/transform/vertex_pulling_test.cc
index c1e63a9..725b7de 100644
--- a/src/tint/transform/vertex_pulling_test.cc
+++ b/src/tint/transform/vertex_pulling_test.cc
@@ -108,7 +108,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@stage(vertex)
@@ -137,7 +137,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -175,7 +175,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -213,7 +213,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(5) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -256,7 +256,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -304,7 +304,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -364,7 +364,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -444,7 +444,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -526,7 +526,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -603,7 +603,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -670,7 +670,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -715,7 +715,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -771,7 +771,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data_1 : @stride(4) array<u32>;
+ tint_vertex_data_1 : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0_1 : TintVertexData;
@@ -848,7 +848,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -1008,7 +1008,7 @@
auto* expect =
R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -1167,7 +1167,7 @@
auto* expect = R"(
struct TintVertexData {
- tint_vertex_data : @stride(4) array<u32>;
+ tint_vertex_data : array<u32>;
}
@binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
diff --git a/src/tint/writer/glsl/generator_impl_member_accessor_test.cc b/src/tint/writer/glsl/generator_impl_member_accessor_test.cc
index 4bf817f..e297050 100644
--- a/src/tint/writer/glsl/generator_impl_member_accessor_test.cc
+++ b/src/tint/writer/glsl/generator_impl_member_accessor_test.cc
@@ -363,7 +363,7 @@
TEST_F(GlslGeneratorImplTest_MemberAccessor,
EmitExpression_IndexAccessor_StorageBuffer_Load_Int_FromArray) {
// struct Data {
- // a : @stride(4) array<i32, 5>;
+ // a : array<i32, 5>;
// };
// var<storage> data : Data;
// data.a[2];
@@ -409,7 +409,7 @@
TEST_F(GlslGeneratorImplTest_MemberAccessor,
EmitExpression_IndexAccessor_StorageBuffer_Load_Int_FromArray_ExprIdx) {
// struct Data {
- // a : @stride(4) array<i32, 5>;
+ // a : array<i32, 5>;
// };
// var<storage> data : Data;
// data.a[(2 + 4) - 3];
@@ -455,7 +455,7 @@
TEST_F(GlslGeneratorImplTest_MemberAccessor, StorageBuffer_Store_ToArray) {
// struct Data {
- // a : @stride(4) array<i32, 5>;
+ // a : array<i32, 5>;
// };
// var<storage> data : Data;
// data.a[2] = 2;
@@ -503,7 +503,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
@@ -562,7 +562,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
@@ -623,7 +623,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
@@ -684,7 +684,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
@@ -744,7 +744,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
@@ -802,7 +802,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index cf0a320..909fcff 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -57,7 +57,6 @@
#include "src/tint/transform/loop_to_for_loop.h"
#include "src/tint/transform/manager.h"
#include "src/tint/transform/num_workgroups_from_uniform.h"
-#include "src/tint/transform/pad_array_elements.h"
#include "src/tint/transform/promote_initializers_to_const_var.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
#include "src/tint/transform/remove_phonies.h"
@@ -209,7 +208,6 @@
manager.Add<transform::CalculateArrayLength>();
manager.Add<transform::PromoteInitializersToConstVar>();
- manager.Add<transform::PadArrayElements>();
manager.Add<transform::AddEmptyEntryPoint>();
data.Add<transform::CanonicalizeEntryPointIO::Config>(
diff --git a/src/tint/writer/hlsl/generator_impl_member_accessor_test.cc b/src/tint/writer/hlsl/generator_impl_member_accessor_test.cc
index 7942ba4..79be125 100644
--- a/src/tint/writer/hlsl/generator_impl_member_accessor_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_member_accessor_test.cc
@@ -391,7 +391,7 @@
TEST_F(HlslGeneratorImplTest_MemberAccessor,
EmitExpression_IndexAccessor_StorageBuffer_Load_Int_FromArray) {
// struct Data {
- // a : @stride(4) array<i32, 5>;
+ // a : array<i32, 5>;
// };
// var<storage> data : Data;
// data.a[2];
@@ -423,7 +423,7 @@
TEST_F(HlslGeneratorImplTest_MemberAccessor,
EmitExpression_IndexAccessor_StorageBuffer_Load_Int_FromArray_ExprIdx) {
// struct Data {
- // a : @stride(4) array<i32, 5>;
+ // a : array<i32, 5>;
// };
// var<storage> data : Data;
// data.a[(2 + 4) - 3];
@@ -455,7 +455,7 @@
TEST_F(HlslGeneratorImplTest_MemberAccessor, StorageBuffer_Store_ToArray) {
// struct Data {
- // a : @stride(4) array<i32, 5>;
+ // a : array<i32, 5>;
// };
// var<storage> data : Data;
// data.a[2] = 2;
@@ -489,7 +489,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
@@ -531,7 +531,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
@@ -575,7 +575,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
@@ -619,7 +619,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
@@ -662,7 +662,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
@@ -703,7 +703,7 @@
// b : vec3<f32>;
// };
// struct Data {
- // var c : @stride(32) array<Inner, 4>;
+ // var c : array<Inner, 4>;
// };
//
// var<storage> data : Pre;
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index fad7968..78388dd 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -63,7 +63,6 @@
#include "src/tint/transform/canonicalize_entry_point_io.h"
#include "src/tint/transform/manager.h"
#include "src/tint/transform/module_scope_var_to_entry_point_param.h"
-#include "src/tint/transform/pad_array_elements.h"
#include "src/tint/transform/promote_initializers_to_const_var.h"
#include "src/tint/transform/promote_side_effects_to_decl.h"
#include "src/tint/transform/remove_phonies.h"
@@ -179,7 +178,6 @@
manager.Add<transform::VectorizeScalarMatrixConstructors>();
manager.Add<transform::WrapArraysInStructs>();
- manager.Add<transform::PadArrayElements>();
manager.Add<transform::RemovePhonies>();
manager.Add<transform::SimplifyPointers>();
// ArrayLengthFromUniform must come after SimplifyPointers, as
@@ -2945,9 +2943,8 @@
[&](const sem::Array* arr) {
if (!arr->IsStrideImplicit()) {
- TINT_ICE(Writer, diagnostics_)
- << "arrays with explicit strides should have "
- "removed with the PadArrayElements transform";
+ TINT_ICE(Writer, diagnostics_) << "arrays with explicit strides not "
+ "exist past the SPIR-V reader";
return SizeAndAlign{};
}
auto num_els = std::max<uint32_t>(arr->Count(), 1);
diff --git a/src/tint/writer/msl/generator_impl_type_test.cc b/src/tint/writer/msl/generator_impl_type_test.cc
index 08f49a2..b59ca32 100644
--- a/src/tint/writer/msl/generator_impl_type_test.cc
+++ b/src/tint/writer/msl/generator_impl_type_test.cc
@@ -119,30 +119,6 @@
EXPECT_EQ(out.str(), "bool ary[1]");
}
-TEST_F(MslGeneratorImplTest, EmitType_ArrayWithStride) {
- auto* s = Structure("s", {Member("arr", ty.array<f32, 4>(64))});
- auto* ubo = Global("ubo", ty.Of(s), ast::StorageClass::kUniform,
- ast::AttributeList{
- create<ast::GroupAttribute>(0),
- create<ast::BindingAttribute>(1),
- });
- WrapInFunction(MemberAccessor(ubo, "arr"));
-
- GeneratorImpl& gen = SanitizeAndBuild();
-
- ASSERT_TRUE(gen.Generate()) << gen.error();
- EXPECT_THAT(gen.result(), HasSubstr(R"(struct tint_padded_array_element {
- /* 0x0000 */ float el;
- /* 0x0004 */ int8_t tint_pad[60];
-};)"));
- EXPECT_THAT(gen.result(), HasSubstr(R"(struct tint_array_wrapper {
- /* 0x0000 */ tint_padded_array_element arr[4];
-};)"));
- EXPECT_THAT(gen.result(), HasSubstr(R"(struct s {
- /* 0x0000 */ tint_array_wrapper arr;
-};)"));
-}
-
TEST_F(MslGeneratorImplTest, EmitType_Bool) {
auto* bool_ = create<sem::Bool>();
diff --git a/src/tint/writer/spirv/builder_builtin_test.cc b/src/tint/writer/spirv/builder_builtin_test.cc
index b7f6a6a..6738546 100644
--- a/src/tint/writer/spirv/builder_builtin_test.cc
+++ b/src/tint/writer/spirv/builder_builtin_test.cc
@@ -1622,7 +1622,7 @@
TEST_F(BuiltinBuilderTest, Call_ArrayLength_ViaLets_WithPtrNoise) {
// struct my_struct {
- // a : @stride(4) array<f32>;
+ // a : array<f32>;
// };
// @binding(1) @group(2) var<storage, read> b : my_struct;
//
diff --git a/test/tint/BUILD.gn b/test/tint/BUILD.gn
index 302fb65..834d948 100644
--- a/test/tint/BUILD.gn
+++ b/test/tint/BUILD.gn
@@ -327,7 +327,6 @@
"../../src/tint/transform/module_scope_var_to_entry_point_param_test.cc",
"../../src/tint/transform/multiplanar_external_texture_test.cc",
"../../src/tint/transform/num_workgroups_from_uniform_test.cc",
- "../../src/tint/transform/pad_array_elements_test.cc",
"../../src/tint/transform/promote_initializers_to_const_var_test.cc",
"../../src/tint/transform/promote_side_effects_to_decl_test.cc",
"../../src/tint/transform/remove_phonies_test.cc",
diff --git a/test/tint/benchmark/animometer.wgsl.expected.glsl b/test/tint/benchmark/animometer.wgsl.expected.glsl
deleted file mode 100644
index 2a44a68..0000000
--- a/test/tint/benchmark/animometer.wgsl.expected.glsl
+++ /dev/null
@@ -1,98 +0,0 @@
-#version 310 es
-
-float tint_float_modulo(float lhs, float rhs) {
- return (lhs - rhs * trunc(lhs / rhs));
-}
-
-
-layout(location = 0) in vec4 position_1;
-layout(location = 1) in vec4 color_1;
-layout(location = 0) out vec4 v_color_1;
-struct Time {
- float value;
-};
-
-struct Uniforms {
- float scale;
- float offsetX;
- float offsetY;
- float scalar;
- float scalarOffset;
-};
-
-layout(binding = 0) uniform Time_1 {
- float value;
-} time;
-
-layout(binding = 1) uniform Uniforms_1 {
- float scale;
- float offsetX;
- float offsetY;
- float scalar;
- float scalarOffset;
-} uniforms;
-
-struct VertexOutput {
- vec4 Position;
- vec4 v_color;
-};
-
-VertexOutput vert_main(vec4 position, vec4 color) {
- float fade = tint_float_modulo((uniforms.scalarOffset + ((time.value * uniforms.scalar) / 10.0f)), 1.0f);
- if ((fade < 0.5f)) {
- fade = (fade * 2.0f);
- } else {
- fade = ((1.0f - fade) * 2.0f);
- }
- float xpos = (position.x * uniforms.scale);
- float ypos = (position.y * uniforms.scale);
- float angle = ((3.141590118f * 2.0f) * fade);
- float xrot = ((xpos * cos(angle)) - (ypos * sin(angle)));
- float yrot = ((xpos * sin(angle)) + (ypos * cos(angle)));
- xpos = (xrot + uniforms.offsetX);
- ypos = (yrot + uniforms.offsetY);
- VertexOutput tint_symbol = VertexOutput(vec4(0.0f, 0.0f, 0.0f, 0.0f), vec4(0.0f, 0.0f, 0.0f, 0.0f));
- tint_symbol.v_color = (vec4(fade, (1.0f - fade), 0.0f, 1.0f) + color);
- tint_symbol.Position = vec4(xpos, ypos, 0.0f, 1.0f);
- return tint_symbol;
-}
-
-void main() {
- VertexOutput inner_result = vert_main(position_1, color_1);
- gl_Position = inner_result.Position;
- v_color_1 = inner_result.v_color;
- gl_Position.y = -(gl_Position.y);
- gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
- return;
-}
-#version 310 es
-precision mediump float;
-
-layout(location = 0) in vec4 v_color_1;
-layout(location = 0) out vec4 value_1;
-struct Time {
- float value;
-};
-
-struct Uniforms {
- float scale;
- float offsetX;
- float offsetY;
- float scalar;
- float scalarOffset;
-};
-
-struct VertexOutput {
- vec4 Position;
- vec4 v_color;
-};
-
-vec4 frag_main(vec4 v_color) {
- return v_color;
-}
-
-void main() {
- vec4 inner_result = frag_main(v_color_1);
- value_1 = inner_result;
- return;
-}
diff --git a/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl b/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl
index 4021327..0397c89 100644
--- a/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl
+++ b/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl
@@ -48,7 +48,7 @@
dirIntensity : f32;
dirDirection : vec3<f32>;
lightCount : u32;
- lights : @stride(32) array<Light>;
+ lights : array<Light>;
}
@binding(2) @group(0) var<storage, read> globalLights : GlobalLights;
diff --git a/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl.expected.glsl b/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl.expected.glsl
deleted file mode 100644
index de73d59..0000000
--- a/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl.expected.glsl
+++ /dev/null
@@ -1,381 +0,0 @@
-benchmark/skinned-shadowed-pbr-fragment.wgsl:51:13 warning: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
- lights : @stride(32) array<Light>;
- ^^^^^^
-
-#version 310 es
-precision mediump float;
-
-layout(location = 0) in vec3 worldPos_1;
-layout(location = 1) in vec3 view_1;
-layout(location = 2) in vec2 texcoord_1;
-layout(location = 3) in vec2 texcoord2_1;
-layout(location = 4) in vec4 color_1;
-layout(location = 5) in vec4 instanceColor_1;
-layout(location = 6) in vec3 normal_1;
-layout(location = 7) in vec3 tangent_1;
-layout(location = 8) in vec3 bitangent_1;
-layout(location = 0) out vec4 color_2;
-layout(location = 1) out vec4 emissive_1;
-const float GAMMA = 2.200000048f;
-vec3 linearTosRGB(vec3 linear) {
- float INV_GAMMA = (1.0f / GAMMA);
- return pow(linear, vec3(INV_GAMMA));
-}
-
-struct Camera {
- mat4 projection;
- mat4 inverseProjection;
- mat4 view;
- vec3 position;
- float time;
- vec2 outputSize;
- float zNear;
- float zFar;
-};
-
-layout(binding = 0) uniform Camera_1 {
- mat4 projection;
- mat4 inverseProjection;
- mat4 view;
- vec3 position;
- float time;
- vec2 outputSize;
- float zNear;
- float zFar;
-} camera;
-
-struct ClusterLights {
- uint offset;
- uint count;
-};
-
-struct ClusterLightGroup {
- uint offset;
- ClusterLights lights[27648];
- uint indices[1769472];
-};
-
-layout(binding = 1, std430) buffer ClusterLightGroup_1 {
- uint offset;
- ClusterLights lights[27648];
- uint indices[1769472];
-} clusterLights;
-struct Light {
- vec3 position;
- float range;
- vec3 color;
- float intensity;
-};
-
-layout(binding = 2, std430) buffer GlobalLights_1 {
- vec3 ambient;
- vec3 dirColor;
- float dirIntensity;
- vec3 dirDirection;
- uint lightCount;
- Light lights[];
-} globalLights;
-const uvec3 tileCount = uvec3(32u, 18u, 48u);
-float linearDepth(float depthSample) {
- return ((camera.zFar * camera.zNear) / ((depthSample) * ((camera.zNear - camera.zFar)) + (camera.zFar)));
-}
-
-uvec3 getTile(vec4 fragCoord) {
- float sliceScale = (float(tileCount.z) / log2((camera.zFar / camera.zNear)));
- float sliceBias = -(((float(tileCount.z) * log2(camera.zNear)) / log2((camera.zFar / camera.zNear))));
- float tint_symbol_3 = linearDepth(fragCoord.z);
- float tint_symbol_4 = log2(tint_symbol_3);
- float tint_symbol_5 = max(((tint_symbol_4 * sliceScale) + sliceBias), 0.0f);
- uint zTile = uint(tint_symbol_5);
- return uvec3(uint((fragCoord.x / (camera.outputSize.x / float(tileCount.x)))), uint((fragCoord.y / (camera.outputSize.y / float(tileCount.y)))), zTile);
-}
-
-uint getClusterIndex(vec4 fragCoord) {
- uvec3 tile = getTile(fragCoord);
- return ((tile.x + (tile.y * tileCount.x)) + ((tile.z * tileCount.x) * tileCount.y));
-}
-
-layout(binding = 6, std430) buffer LightShadowTable_1 {
- int light[];
-} lightShadowTable;
-vec2 shadowSampleOffsets[16] = vec2[16](vec2(-1.5f, -1.5f), vec2(-1.5f, -0.5f), vec2(-1.5f, 0.5f), vec2(-1.5f, 1.5f), vec2(-0.5f, -1.5f), vec2(-0.5f, -0.5f), vec2(-0.5f, 0.5f), vec2(-0.5f, 1.5f), vec2(0.5f, -1.5f), vec2(0.5f, -0.5f), vec2(0.5f, 0.5f), vec2(0.5f, 1.5f), vec2(1.5f, -1.5f), vec2(1.5f, -0.5f), vec2(1.5f, 0.5f), vec2(1.5f, 1.5f));
-const uint shadowSampleCount = 16u;
-struct ShadowProperties {
- vec4 viewport;
- mat4 viewProj;
-};
-
-layout(binding = 7, std430) buffer LightShadows_1 {
- ShadowProperties properties[];
-} shadow;
-uniform highp sampler2D shadowTexture_1;
-uniform highp sampler2DShadow shadowTexture_shadowSampler;
-
-float dirLightVisibility(vec3 worldPos) {
- int shadowIndex = lightShadowTable.light[0u];
- if ((shadowIndex == -1)) {
- return 1.0f;
- }
- vec4 viewport = shadow.properties[shadowIndex].viewport;
- vec4 lightPos = (shadow.properties[shadowIndex].viewProj * vec4(worldPos, 1.0f));
- vec3 shadowPos = vec3((((lightPos.xy / lightPos.w) * vec2(0.5f, -0.5f)) + vec2(0.5f, 0.5f)), (lightPos.z / lightPos.w));
- vec2 viewportPos = vec2((viewport.xy + (shadowPos.xy * viewport.zw)));
- vec2 texelSize = (1.0f / vec2(textureSize(shadowTexture_1, 0)));
- vec4 clampRect = vec4((viewport.xy - texelSize), ((viewport.xy + viewport.zw) + texelSize));
- float visibility = 0.0f;
- {
- for(uint i = 0u; (i < shadowSampleCount); i = (i + 1u)) {
- visibility = (visibility + texture(shadowTexture_shadowSampler, vec3(clamp((viewportPos + (shadowSampleOffsets[i] * texelSize)), clampRect.xy, clampRect.zw), (shadowPos.z - 0.003f))));
- }
- }
- return (visibility / float(shadowSampleCount));
-}
-
-int getCubeFace(vec3 v) {
- vec3 vAbs = abs(v);
- bool tint_tmp = (vAbs.z >= vAbs.x);
- if (tint_tmp) {
- tint_tmp = (vAbs.z >= vAbs.y);
- }
- if ((tint_tmp)) {
- if ((v.z < 0.0f)) {
- return 5;
- }
- return 4;
- }
- if ((vAbs.y >= vAbs.x)) {
- if ((v.y < 0.0f)) {
- return 3;
- }
- return 2;
- }
- if ((v.x < 0.0f)) {
- return 1;
- }
- return 0;
-}
-
-float pointLightVisibility(uint lightIndex, vec3 worldPos, vec3 pointToLight) {
- int shadowIndex = lightShadowTable.light[(lightIndex + 1u)];
- if ((shadowIndex == -1)) {
- return 1.0f;
- }
- int tint_symbol_6 = shadowIndex;
- int tint_symbol_7 = getCubeFace((pointToLight * -1.0f));
- shadowIndex = (tint_symbol_6 + tint_symbol_7);
- vec4 viewport = shadow.properties[shadowIndex].viewport;
- vec4 lightPos = (shadow.properties[shadowIndex].viewProj * vec4(worldPos, 1.0f));
- vec3 shadowPos = vec3((((lightPos.xy / lightPos.w) * vec2(0.5f, -0.5f)) + vec2(0.5f, 0.5f)), (lightPos.z / lightPos.w));
- vec2 viewportPos = vec2((viewport.xy + (shadowPos.xy * viewport.zw)));
- vec2 texelSize = (1.0f / vec2(textureSize(shadowTexture_1, 0)));
- vec4 clampRect = vec4(viewport.xy, (viewport.xy + viewport.zw));
- float visibility = 0.0f;
- {
- for(uint i = 0u; (i < shadowSampleCount); i = (i + 1u)) {
- visibility = (visibility + texture(shadowTexture_shadowSampler, vec3(clamp((viewportPos + (shadowSampleOffsets[i] * texelSize)), clampRect.xy, clampRect.zw), (shadowPos.z - 0.01f))));
- }
- }
- return (visibility / float(shadowSampleCount));
-}
-
-struct VertexOutput {
- vec4 position;
- vec3 worldPos;
- vec3 view;
- vec2 texcoord;
- vec2 texcoord2;
- vec4 color;
- vec4 instanceColor;
- vec3 normal;
- vec3 tangent;
- vec3 bitangent;
-};
-
-struct Material {
- vec4 baseColorFactor;
- vec3 emissiveFactor;
- float occlusionStrength;
- vec2 metallicRoughnessFactor;
- float alphaCutoff;
-};
-
-layout(binding = 8) uniform Material_1 {
- vec4 baseColorFactor;
- vec3 emissiveFactor;
- float occlusionStrength;
- vec2 metallicRoughnessFactor;
- float alphaCutoff;
-} material;
-
-struct SurfaceInfo {
- vec4 baseColor;
- vec3 albedo;
- float metallic;
- float roughness;
- vec3 normal;
- vec3 f0;
- float ao;
- vec3 emissive;
- vec3 v;
-};
-
-uniform highp sampler2D normalTexture_normalSampler;
-uniform highp sampler2D baseColorTexture_baseColorSampler;
-uniform highp sampler2D metallicRoughnessTexture_metallicRoughnessSampler;
-uniform highp sampler2D occlusionTexture_occlusionSampler;
-uniform highp sampler2D emissiveTexture_emissiveSampler;
-
-SurfaceInfo GetSurfaceInfo(VertexOutput tint_symbol) {
- SurfaceInfo surface = SurfaceInfo(vec4(0.0f, 0.0f, 0.0f, 0.0f), vec3(0.0f, 0.0f, 0.0f), 0.0f, 0.0f, vec3(0.0f, 0.0f, 0.0f), vec3(0.0f, 0.0f, 0.0f), 0.0f, vec3(0.0f, 0.0f, 0.0f), vec3(0.0f, 0.0f, 0.0f));
- surface.v = normalize(tint_symbol.view);
- mat3 tbn = mat3(tint_symbol.tangent, tint_symbol.bitangent, tint_symbol.normal);
- vec3 normalMap = texture(normalTexture_normalSampler, tint_symbol.texcoord).rgb;
- surface.normal = normalize((tbn * ((2.0f * normalMap) - vec3(1.0f))));
- vec4 baseColorMap = texture(baseColorTexture_baseColorSampler, tint_symbol.texcoord);
- surface.baseColor = ((tint_symbol.color * material.baseColorFactor) * baseColorMap);
- if ((surface.baseColor.a < material.alphaCutoff)) {
- discard;
- }
- surface.albedo = surface.baseColor.rgb;
- vec4 metallicRoughnessMap = texture(metallicRoughnessTexture_metallicRoughnessSampler, tint_symbol.texcoord);
- surface.metallic = (material.metallicRoughnessFactor.x * metallicRoughnessMap.b);
- surface.roughness = (material.metallicRoughnessFactor.y * metallicRoughnessMap.g);
- vec3 dielectricSpec = vec3(0.039999999f);
- surface.f0 = mix(dielectricSpec, surface.albedo, vec3(surface.metallic));
- vec4 occlusionMap = texture(occlusionTexture_occlusionSampler, tint_symbol.texcoord);
- surface.ao = (material.occlusionStrength * occlusionMap.r);
- vec4 emissiveMap = texture(emissiveTexture_emissiveSampler, tint_symbol.texcoord);
- surface.emissive = (material.emissiveFactor * emissiveMap.rgb);
- if ((tint_symbol.instanceColor.a == 0.0f)) {
- surface.albedo = (surface.albedo + tint_symbol.instanceColor.rgb);
- } else {
- surface.albedo = (surface.albedo * tint_symbol.instanceColor.rgb);
- }
- return surface;
-}
-
-const float PI = 3.141592741f;
-const uint LightType_Point = 0u;
-const uint LightType_Directional = 2u;
-struct PuctualLight {
- uint lightType;
- vec3 pointToLight;
- float range;
- vec3 color;
- float intensity;
-};
-
-vec3 FresnelSchlick(float cosTheta, vec3 F0) {
- return (F0 + ((vec3(1.0f) - F0) * pow((1.0f - cosTheta), 5.0f)));
-}
-
-float DistributionGGX(vec3 N, vec3 H, float roughness) {
- float a_1 = (roughness * roughness);
- float a2 = (a_1 * a_1);
- float NdotH = max(dot(N, H), 0.0f);
- float NdotH2 = (NdotH * NdotH);
- float num = a2;
- float denom = ((NdotH2 * (a2 - 1.0f)) + 1.0f);
- return (num / ((PI * denom) * denom));
-}
-
-float GeometrySchlickGGX(float NdotV, float roughness) {
- float r_1 = (roughness + 1.0f);
- float k = ((r_1 * r_1) / 8.0f);
- float num = NdotV;
- float denom = ((NdotV * (1.0f - k)) + k);
- return (num / denom);
-}
-
-float GeometrySmith(vec3 N, vec3 V, vec3 L, float roughness) {
- float NdotV = max(dot(N, V), 0.0f);
- float NdotL = max(dot(N, L), 0.0f);
- float ggx2 = GeometrySchlickGGX(NdotV, roughness);
- float ggx1 = GeometrySchlickGGX(NdotL, roughness);
- return (ggx1 * ggx2);
-}
-
-float lightAttenuation(PuctualLight light) {
- if ((light.lightType == LightType_Directional)) {
- return 1.0f;
- }
- float tint_symbol_1 = length(light.pointToLight);
- if ((light.range <= 0.0f)) {
- return (1.0f / pow(tint_symbol_1, 2.0f));
- }
- return (clamp((1.0f - pow((tint_symbol_1 / light.range), 4.0f)), 0.0f, 1.0f) / pow(tint_symbol_1, 2.0f));
-}
-
-vec3 lightRadiance(PuctualLight light, SurfaceInfo surface) {
- vec3 L = normalize(light.pointToLight);
- vec3 H = normalize((surface.v + L));
- float NDF = DistributionGGX(surface.normal, H, surface.roughness);
- float G = GeometrySmith(surface.normal, surface.v, L, surface.roughness);
- vec3 F = FresnelSchlick(max(dot(H, surface.v), 0.0f), surface.f0);
- vec3 kD = ((vec3(1.0f) - F) * (1.0f - surface.metallic));
- float NdotL = max(dot(surface.normal, L), 0.0f);
- vec3 numerator = ((NDF * G) * F);
- float denominator = max(((4.0f * max(dot(surface.normal, surface.v), 0.0f)) * NdotL), 0.001f);
- vec3 specular = (numerator / vec3(denominator));
- vec3 tint_symbol_8 = (light.color * light.intensity);
- float tint_symbol_9 = lightAttenuation(light);
- vec3 radiance = (tint_symbol_8 * tint_symbol_9);
- return (((((kD * surface.albedo) / vec3(PI)) + specular) * radiance) * NdotL);
-}
-
-struct FragmentOutput {
- vec4 color;
- vec4 emissive;
-};
-
-uniform highp sampler2D ssaoTexture_1;
-uniform highp sampler2D ssaoTexture_defaultSampler;
-FragmentOutput fragmentMain(VertexOutput tint_symbol) {
- SurfaceInfo surface = GetSurfaceInfo(tint_symbol);
- vec3 Lo = vec3(0.0f, 0.0f, 0.0f);
- if ((globalLights.dirIntensity > 0.0f)) {
- PuctualLight light = PuctualLight(0u, vec3(0.0f, 0.0f, 0.0f), 0.0f, vec3(0.0f, 0.0f, 0.0f), 0.0f);
- light.lightType = LightType_Directional;
- light.pointToLight = globalLights.dirDirection;
- light.color = globalLights.dirColor;
- light.intensity = globalLights.dirIntensity;
- float lightVis = dirLightVisibility(tint_symbol.worldPos);
- vec3 tint_symbol_10 = Lo;
- vec3 tint_symbol_11 = lightRadiance(light, surface);
- Lo = (tint_symbol_10 + (tint_symbol_11 * lightVis));
- }
- uint clusterIndex = getClusterIndex(tint_symbol.position);
- uint lightOffset = clusterLights.lights[clusterIndex].offset;
- uint lightCount = clusterLights.lights[clusterIndex].count;
- {
- for(uint lightIndex = 0u; (lightIndex < lightCount); lightIndex = (lightIndex + 1u)) {
- uint i = clusterLights.indices[(lightOffset + lightIndex)];
- PuctualLight light = PuctualLight(0u, vec3(0.0f, 0.0f, 0.0f), 0.0f, vec3(0.0f, 0.0f, 0.0f), 0.0f);
- light.lightType = LightType_Point;
- light.pointToLight = (globalLights.lights[i].position.xyz - tint_symbol.worldPos);
- light.range = globalLights.lights[i].range;
- light.color = globalLights.lights[i].color;
- light.intensity = globalLights.lights[i].intensity;
- float lightVis = pointLightVisibility(i, tint_symbol.worldPos, light.pointToLight);
- vec3 tint_symbol_12 = Lo;
- vec3 tint_symbol_13 = lightRadiance(light, surface);
- Lo = (tint_symbol_12 + (tint_symbol_13 * lightVis));
- }
- }
- vec2 ssaoCoord = (tint_symbol.position.xy / vec2(textureSize(ssaoTexture_1, 0).xy));
- float ssaoFactor = texture(ssaoTexture_defaultSampler, ssaoCoord).r;
- vec3 ambient = (((globalLights.ambient * surface.albedo) * surface.ao) * ssaoFactor);
- vec3 color = linearTosRGB(((Lo + ambient) + surface.emissive));
- FragmentOutput tint_symbol_2 = FragmentOutput(vec4(0.0f, 0.0f, 0.0f, 0.0f), vec4(0.0f, 0.0f, 0.0f, 0.0f));
- tint_symbol_2.color = vec4(color, surface.baseColor.a);
- tint_symbol_2.emissive = vec4(surface.emissive, surface.baseColor.a);
- return tint_symbol_2;
-}
-
-void main() {
- VertexOutput tint_symbol_14 = VertexOutput(gl_FragCoord, worldPos_1, view_1, texcoord_1, texcoord2_1, color_1, instanceColor_1, normal_1, tangent_1, bitangent_1);
- FragmentOutput inner_result = fragmentMain(tint_symbol_14);
- color_2 = inner_result.color;
- emissive_1 = inner_result.emissive;
- return;
-}
