Tint: Resolve @input_attachment_index.
Bug: 341117913
Change-Id: I120d1c5ea7cb9fe365a6396e5564edf730753b80
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/189480
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: Quyen Le <lehoangquyen@chromium.org>
diff --git a/src/tint/lang/wgsl/resolver/BUILD.bazel b/src/tint/lang/wgsl/resolver/BUILD.bazel
index 9893abb..e5ea130 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.bazel
+++ b/src/tint/lang/wgsl/resolver/BUILD.bazel
@@ -125,6 +125,7 @@
"host_shareable_validation_test.cc",
"increment_decrement_validation_test.cc",
"inferred_type_test.cc",
+ "input_attachments_extension_test.cc",
"is_host_shareable_test.cc",
"is_storeable_test.cc",
"language_features_test.cc",
diff --git a/src/tint/lang/wgsl/resolver/BUILD.cmake b/src/tint/lang/wgsl/resolver/BUILD.cmake
index 429b9c0..326af91 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.cmake
+++ b/src/tint/lang/wgsl/resolver/BUILD.cmake
@@ -123,6 +123,7 @@
lang/wgsl/resolver/host_shareable_validation_test.cc
lang/wgsl/resolver/increment_decrement_validation_test.cc
lang/wgsl/resolver/inferred_type_test.cc
+ lang/wgsl/resolver/input_attachments_extension_test.cc
lang/wgsl/resolver/is_host_shareable_test.cc
lang/wgsl/resolver/is_storeable_test.cc
lang/wgsl/resolver/language_features_test.cc
diff --git a/src/tint/lang/wgsl/resolver/BUILD.gn b/src/tint/lang/wgsl/resolver/BUILD.gn
index 9456dbb..6fb89d0 100644
--- a/src/tint/lang/wgsl/resolver/BUILD.gn
+++ b/src/tint/lang/wgsl/resolver/BUILD.gn
@@ -125,6 +125,7 @@
"host_shareable_validation_test.cc",
"increment_decrement_validation_test.cc",
"inferred_type_test.cc",
+ "input_attachments_extension_test.cc",
"is_host_shareable_test.cc",
"is_storeable_test.cc",
"language_features_test.cc",
diff --git a/src/tint/lang/wgsl/resolver/attribute_validation_test.cc b/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
index e94ebbe..809c8e3 100644
--- a/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
+++ b/src/tint/lang/wgsl/resolver/attribute_validation_test.cc
@@ -65,6 +65,7 @@
kDiagnostic,
kGroup,
kId,
+ kInputAttachmentIndex,
kInterpolate,
kInvariant,
kLocation,
@@ -93,6 +94,8 @@
return o << "@group";
case AttributeKind::kId:
return o << "@id";
+ case AttributeKind::kInputAttachmentIndex:
+ return o << "@input_attachment_index";
case AttributeKind::kInterpolate:
return o << "@interpolate";
case AttributeKind::kInvariant:
@@ -166,6 +169,10 @@
"1:2 error: '@id' is not valid for " + thing,
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ "1:2 error: '@input_attachment_index' is not valid for " + thing,
+ },
+ TestParams{
{AttributeKind::kInterpolate},
"1:2 error: '@interpolate' is not valid for " + thing,
},
@@ -230,6 +237,8 @@
return builder.Group(source, 1_a);
case AttributeKind::kId:
return builder.Id(source, 0_a);
+ case AttributeKind::kInputAttachmentIndex:
+ return builder.InputAttachmentIndex(source, 2_a);
case AttributeKind::kBlendSrc:
return builder.BlendSrc(source, 0_a);
case AttributeKind::kInterpolate:
@@ -264,6 +273,9 @@
case AttributeKind::kBlendSrc:
Enable(wgsl::Extension::kChromiumInternalDualSourceBlending);
break;
+ case AttributeKind::kInputAttachmentIndex:
+ Enable(wgsl::Extension::kChromiumInternalInputAttachments);
+ break;
default:
break;
}
@@ -345,6 +357,10 @@
R"(1:2 error: '@id' is not valid for functions)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for functions)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(1:2 error: '@interpolate' is not valid for functions)",
},
@@ -393,81 +409,86 @@
CHECK();
}
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
- NonVoidFunctionAttributeTest,
- testing::Values(
- TestParams{
- {AttributeKind::kAlign},
- R"(1:2 error: '@align' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kBinding},
- R"(1:2 error: '@binding' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kBlendSrc},
- R"(1:2 error: '@blend_src' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kBuiltinPosition},
- R"(1:2 error: '@builtin' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kColor},
- R"(1:2 error: '@color' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kDiagnostic},
- Pass,
- },
- TestParams{
- {AttributeKind::kGroup},
- R"(1:2 error: '@group' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kId},
- R"(1:2 error: '@id' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kInterpolate},
- R"(1:2 error: '@interpolate' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kInvariant},
- R"(1:2 error: '@invariant' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kLocation},
- R"(1:2 error: '@location' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kMustUse},
- Pass,
- },
- TestParams{
- {AttributeKind::kOffset},
- R"(1:2 error: '@offset' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kSize},
- R"(1:2 error: '@size' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kStageCompute},
- R"(9:9 error: missing entry point IO attribute on return type)",
- },
- TestParams{
- {AttributeKind::kStageCompute, AttributeKind::kWorkgroupSize},
- R"(9:9 error: missing entry point IO attribute on return type)",
- },
- TestParams{
- {AttributeKind::kStride},
- R"(1:2 error: '@stride' is not valid for functions)",
- },
- TestParams{
- {AttributeKind::kWorkgroupSize},
- R"(1:2 error: '@workgroup_size' is only valid for compute stages)",
- }));
+INSTANTIATE_TEST_SUITE_P(
+ ResolverAttributeValidationTest,
+ NonVoidFunctionAttributeTest,
+ testing::Values(
+ TestParams{
+ {AttributeKind::kAlign},
+ R"(1:2 error: '@align' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kBinding},
+ R"(1:2 error: '@binding' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kBlendSrc},
+ R"(1:2 error: '@blend_src' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kBuiltinPosition},
+ R"(1:2 error: '@builtin' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kColor},
+ R"(1:2 error: '@color' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kDiagnostic},
+ Pass,
+ },
+ TestParams{
+ {AttributeKind::kGroup},
+ R"(1:2 error: '@group' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kId},
+ R"(1:2 error: '@id' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kInterpolate},
+ R"(1:2 error: '@interpolate' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kInvariant},
+ R"(1:2 error: '@invariant' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kLocation},
+ R"(1:2 error: '@location' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kMustUse},
+ Pass,
+ },
+ TestParams{
+ {AttributeKind::kOffset},
+ R"(1:2 error: '@offset' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kSize},
+ R"(1:2 error: '@size' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kStageCompute},
+ R"(9:9 error: missing entry point IO attribute on return type)",
+ },
+ TestParams{
+ {AttributeKind::kStageCompute, AttributeKind::kWorkgroupSize},
+ R"(9:9 error: missing entry point IO attribute on return type)",
+ },
+ TestParams{
+ {AttributeKind::kStride},
+ R"(1:2 error: '@stride' is not valid for functions)",
+ },
+ TestParams{
+ {AttributeKind::kWorkgroupSize},
+ R"(1:2 error: '@workgroup_size' is only valid for compute stages)",
+ }));
} // namespace FunctionTests
namespace FunctionInputAndOutputTests {
@@ -520,6 +541,10 @@
R"(1:2 error: '@id' is not valid for function parameters)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(1:2 error: '@interpolate' is not valid for non-entry point function parameters)",
},
@@ -605,6 +630,10 @@
R"(1:2 error: '@id' is not valid for non-entry point function return types)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for non-entry point function return types)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(1:2 error: '@interpolate' is not valid for non-entry point function return types)",
},
@@ -695,6 +724,10 @@
R"(1:2 error: '@id' is not valid for function parameters)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(1:2 error: '@interpolate' cannot be used by compute shaders)",
},
@@ -784,6 +817,10 @@
R"(1:2 error: '@id' is not valid for function parameters)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(9:9 error: missing entry point IO attribute on parameter)",
},
@@ -887,6 +924,10 @@
R"(1:2 error: '@id' is not valid for function parameters)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for function parameters)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(9:9 error: missing entry point IO attribute on parameter)",
},
@@ -992,6 +1033,10 @@
R"(1:2 error: '@id' is not valid for entry point return types)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for entry point return types)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(1:2 error: '@interpolate' cannot be used by compute shaders)",
},
@@ -1082,6 +1127,10 @@
R"(1:2 error: '@id' is not valid for entry point return types)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for entry point return types)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(9:9 error: missing entry point IO attribute on return type)",
},
@@ -1190,6 +1239,10 @@
R"(1:2 error: '@id' is not valid for entry point return types)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for entry point return types)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(1:2 error: '@interpolate' can only be used with '@location')",
},
@@ -1315,6 +1368,10 @@
R"(1:2 error: '@id' is not valid for 'struct' declarations)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for 'struct' declarations)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(1:2 error: '@interpolate' is not valid for 'struct' declarations)",
},
@@ -1399,6 +1456,10 @@
R"(1:2 error: '@id' is not valid for 'struct' members)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for 'struct' members)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(1:2 error: '@interpolate' can only be used with '@location')",
},
@@ -1635,82 +1696,87 @@
CHECK();
}
-INSTANTIATE_TEST_SUITE_P(ResolverAttributeValidationTest,
- ArrayAttributeTest,
- testing::Values(
- TestParams{
- {AttributeKind::kAlign},
- R"(1:2 error: '@align' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kBinding},
- R"(1:2 error: '@binding' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kBlendSrc},
- R"(1:2 error: '@blend_src' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kBuiltinPosition},
- R"(1:2 error: '@builtin' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kDiagnostic},
- R"(1:2 error: '@diagnostic' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kGroup},
- R"(1:2 error: '@group' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kId},
- R"(1:2 error: '@id' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kInterpolate},
- R"(1:2 error: '@interpolate' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kInvariant},
- R"(1:2 error: '@invariant' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kLocation},
- R"(1:2 error: '@location' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kMustUse},
- R"(1:2 error: '@must_use' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kOffset},
- R"(1:2 error: '@offset' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kSize},
- R"(1:2 error: '@size' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kStageCompute},
- R"(1:2 error: '@compute' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kStride},
- Pass,
- },
- TestParams{
- {AttributeKind::kWorkgroupSize},
- R"(1:2 error: '@workgroup_size' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kBinding, AttributeKind::kGroup},
- R"(1:2 error: '@binding' is not valid for 'array' types)",
- },
- TestParams{
- {AttributeKind::kStride, AttributeKind::kStride},
- R"(3:4 error: duplicate stride attribute
+INSTANTIATE_TEST_SUITE_P(
+ ResolverAttributeValidationTest,
+ ArrayAttributeTest,
+ testing::Values(
+ TestParams{
+ {AttributeKind::kAlign},
+ R"(1:2 error: '@align' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kBinding},
+ R"(1:2 error: '@binding' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kBlendSrc},
+ R"(1:2 error: '@blend_src' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kBuiltinPosition},
+ R"(1:2 error: '@builtin' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kDiagnostic},
+ R"(1:2 error: '@diagnostic' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kGroup},
+ R"(1:2 error: '@group' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kId},
+ R"(1:2 error: '@id' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kInterpolate},
+ R"(1:2 error: '@interpolate' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kInvariant},
+ R"(1:2 error: '@invariant' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kLocation},
+ R"(1:2 error: '@location' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kMustUse},
+ R"(1:2 error: '@must_use' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kOffset},
+ R"(1:2 error: '@offset' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kSize},
+ R"(1:2 error: '@size' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kStageCompute},
+ R"(1:2 error: '@compute' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kStride},
+ Pass,
+ },
+ TestParams{
+ {AttributeKind::kWorkgroupSize},
+ R"(1:2 error: '@workgroup_size' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kBinding, AttributeKind::kGroup},
+ R"(1:2 error: '@binding' is not valid for 'array' types)",
+ },
+ TestParams{
+ {AttributeKind::kStride, AttributeKind::kStride},
+ R"(3:4 error: duplicate stride attribute
1:2 note: first attribute declared here)",
- }));
+ }));
using VariableAttributeTest = TestWithParams;
TEST_P(VariableAttributeTest, IsValid) {
@@ -1862,6 +1928,10 @@
R"(1:2 error: '@id' is not valid for 'const' declaration)",
},
TestParams{
+ {AttributeKind::kInputAttachmentIndex},
+ R"(1:2 error: '@input_attachment_index' is not valid for 'const' declaration)",
+ },
+ TestParams{
{AttributeKind::kInterpolate},
R"(1:2 error: '@interpolate' is not valid for 'const' declaration)",
},
diff --git a/src/tint/lang/wgsl/resolver/dependency_graph.cc b/src/tint/lang/wgsl/resolver/dependency_graph.cc
index 4838cde..bffa92e 100644
--- a/src/tint/lang/wgsl/resolver/dependency_graph.cc
+++ b/src/tint/lang/wgsl/resolver/dependency_graph.cc
@@ -52,6 +52,7 @@
#include "src/tint/lang/wgsl/ast/identifier.h"
#include "src/tint/lang/wgsl/ast/if_statement.h"
#include "src/tint/lang/wgsl/ast/increment_decrement_statement.h"
+#include "src/tint/lang/wgsl/ast/input_attachment_index_attribute.h"
#include "src/tint/lang/wgsl/ast/internal_attribute.h"
#include "src/tint/lang/wgsl/ast/interpolate_attribute.h"
#include "src/tint/lang/wgsl/ast/invariant_attribute.h"
@@ -379,6 +380,7 @@
[&](const ast::ColorAttribute* color) { TraverseExpression(color->expr); },
[&](const ast::GroupAttribute* group) { TraverseExpression(group->expr); },
[&](const ast::IdAttribute* id) { TraverseExpression(id->expr); },
+ [&](const ast::InputAttachmentIndexAttribute* idx) { TraverseExpression(idx->expr); },
[&](const ast::BlendSrcAttribute* index) { TraverseExpression(index->expr); },
[&](const ast::InterpolateAttribute* interpolate) {
TraverseExpression(interpolate->type);
diff --git a/src/tint/lang/wgsl/resolver/dependency_graph_test.cc b/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
index e435301..f64306a 100644
--- a/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
+++ b/src/tint/lang/wgsl/resolver/dependency_graph_test.cc
@@ -1761,6 +1761,7 @@
GlobalVar(Sym(), ty.sampler(core::type::SamplerKind::kSampler));
GlobalVar(Sym(), ty.i32(), Vector{Binding(V), Group(V)});
+ GlobalVar(Sym(), ty.input_attachment(T), Vector{Binding(V), Group(V), InputAttachmentIndex(V)});
GlobalVar(Sym(), ty.i32(), Vector{Location(V)});
Override(Sym(), ty.i32(), Vector{Id(V)});
diff --git a/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc b/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc
new file mode 100644
index 0000000..2b0449d
--- /dev/null
+++ b/src/tint/lang/wgsl/resolver/input_attachments_extension_test.cc
@@ -0,0 +1,166 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/wgsl/resolver/resolver.h"
+#include "src/tint/lang/wgsl/resolver/resolver_helper_test.h"
+
+namespace tint::resolver {
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+using InputAttachmenExtensionTest = ResolverTest;
+
+// Test that input_attachment cannot be used without extension.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentWithoutExtension) {
+ // @group(0) @binding(0) @input_attachment_index(3)
+ // var input_tex : input_attachment<f32>;
+
+ GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+ Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(
+ r()->error(),
+ R"(error: use of 'input_attachment' requires enabling extension 'chromium_internal_input_attachments')");
+}
+
+// Test that input_attachment cannot be declared locally.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentLocalDecl) {
+ // enable chromium_internal_input_attachments;
+ // @fragment fn f() {
+ // var input_tex : input_attachment<f32>;
+ // }
+
+ Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+ Func("f", Empty, ty.void_(),
+ Vector{
+ Decl(Var("input_tex", ty.input_attachment(ty.Of<f32>()))),
+ },
+ Vector{Stage(ast::PipelineStage::kFragment)});
+ EXPECT_FALSE(r()->Resolve());
+}
+
+// Test that input_attachment cannot be declared without index.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentWithoutIndex) {
+ // enable chromium_internal_input_attachments;
+ // @group(0) @binding(0)
+ // var input_tex : input_attachment<f32>;
+
+ Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+ GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()), Vector{Binding(0_u), Group(0_u)});
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(),
+ R"(error: 'input_attachment' variables require '@input_attachment_index' attribute)");
+}
+
+// Test that Resolver can get input_attachment_index value.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexValue) {
+ // enable chromium_internal_input_attachments;
+ // @group(0) @binding(0) @input_attachment_index(3)
+ // var input_tex : input_attachment<f32>;
+
+ Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+ auto* ast_var = GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+ Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+ EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+ auto* sem_var = Sem().Get<sem::GlobalVariable>(ast_var);
+ ASSERT_NE(sem_var, nullptr);
+ EXPECT_EQ(sem_var->Attributes().input_attachment_index, 3u);
+}
+
+// Test that @input_attachment_index cannot be used without extension.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexWithoutExtension) {
+ // @group(0) @binding(0) @input_attachment_index(3)
+ // var input_tex : texture_2d<f32>;
+
+ GlobalVar("input_tex", ty.sampled_texture(core::type::TextureDimension::k2d, ty.Of<f32>()),
+ Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(
+ r()->error(),
+ R"(error: use of '@input_attachment_index' requires enabling extension 'chromium_internal_input_attachments')");
+}
+
+// Test that input_attachment_index's value cannot be float.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexInvalidValueType) {
+ // enable chromium_internal_input_attachments;
+ // @group(0) @binding(0) @input_attachment_index(3.0)
+ // var input_tex : input_attachment<f32>;
+
+ Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+ GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+ Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_f)});
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(), R"(error: '@input_attachment_index' must be an 'i32' or 'u32' value)");
+}
+
+// Test that input_attachment_index's value cannot be negative.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexNegative) {
+ // enable chromium_internal_input_attachments;
+ // @group(0) @binding(0) @input_attachment_index(-2)
+ // var input_tex : input_attachment<f32>;
+
+ Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+ GlobalVar("input_tex", ty.input_attachment(ty.Of<f32>()),
+ Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(core::i32(-2))});
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(), R"(error: '@input_attachment_index' value must be non-negative)");
+}
+
+// Test that input_attachment_index cannot be used on non input_attachment variable.
+TEST_F(InputAttachmenExtensionTest, InputAttachmentIndexInvalidType) {
+ // enable chromium_internal_input_attachments;
+ // @group(0) @binding(0) @input_attachment_index(3)
+ // var input_tex : texture_2d<f32>;
+
+ Enable(Source{{12, 34}}, wgsl::Extension::kChromiumInternalInputAttachments);
+
+ GlobalVar("input_tex", ty.sampled_texture(core::type::TextureDimension::k2d, ty.Of<f32>()),
+ Vector{Binding(0_u), Group(0_u), InputAttachmentIndex(3_u)});
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(
+ r()->error(),
+ R"(error: cannot apply '@input_attachment_index' to declaration of type 'texture_2d<f32>'
+note: '@input_attachment_index' must only be applied to declarations of 'input_attachment' type)");
+}
+
+} // namespace
+} // namespace tint::resolver
diff --git a/src/tint/lang/wgsl/resolver/resolver.cc b/src/tint/lang/wgsl/resolver/resolver.cc
index 0a17841..fe24ffc 100644
--- a/src/tint/lang/wgsl/resolver/resolver.cc
+++ b/src/tint/lang/wgsl/resolver/resolver.cc
@@ -65,6 +65,7 @@
#include "src/tint/lang/wgsl/ast/for_loop_statement.h"
#include "src/tint/lang/wgsl/ast/id_attribute.h"
#include "src/tint/lang/wgsl/ast/if_statement.h"
+#include "src/tint/lang/wgsl/ast/input_attachment_index_attribute.h"
#include "src/tint/lang/wgsl/ast/internal_attribute.h"
#include "src/tint/lang/wgsl/ast/interpolate_attribute.h"
#include "src/tint/lang/wgsl/ast/loop_statement.h"
@@ -614,7 +615,7 @@
bool has_io_address_space = sem->AddressSpace() == core::AddressSpace::kIn ||
sem->AddressSpace() == core::AddressSpace::kOut;
- std::optional<uint32_t> group, binding;
+ std::optional<uint32_t> group, binding, input_attachment_index;
for (auto* attribute : var->attributes) {
Mark(attribute);
enum Status { kSuccess, kErrored, kInvalid };
@@ -636,6 +637,14 @@
group = value.Get();
return kSuccess;
},
+ [&](const ast::InputAttachmentIndexAttribute* attr) {
+ auto value = InputAttachmentIndexAttribute(attr);
+ if (value != Success) {
+ return kErrored;
+ }
+ input_attachment_index = value.Get();
+ return kSuccess;
+ },
[&](const ast::LocationAttribute* attr) {
if (!has_io_address_space) {
return kInvalid;
@@ -708,6 +717,10 @@
global->Attributes().binding_point = BindingPoint{group.value(), binding.value()};
}
+ if (input_attachment_index) {
+ global->Attributes().input_attachment_index = input_attachment_index;
+ }
+
} else {
for (auto* attribute : var->attributes) {
Mark(attribute);
@@ -3876,6 +3889,31 @@
return static_cast<uint32_t>(value);
}
+tint::Result<uint32_t> Resolver::InputAttachmentIndexAttribute(
+ const ast::InputAttachmentIndexAttribute* attr) {
+ ExprEvalStageConstraint constraint{core::EvaluationStage::kConstant, "@input_attachment_index"};
+ TINT_SCOPED_ASSIGNMENT(expr_eval_stage_constraint_, constraint);
+
+ auto* materialized = Materialize(ValueExpression(attr->expr));
+ if (!materialized) {
+ return Failure{};
+ }
+ if (!materialized->Type()->IsAnyOf<core::type::I32, core::type::U32>()) {
+ AddError(attr->source) << style::Attribute("@input_attachment_index") << " must be an "
+ << style::Type("i32") << " or " << style::Type("u32") << " value";
+ return Failure{};
+ }
+
+ auto const_value = materialized->ConstantValue();
+ auto value = const_value->ValueAs<AInt>();
+ if (value < 0) {
+ AddError(attr->source) << style::Attribute("@input_attachment_index")
+ << " value must be non-negative";
+ return Failure{};
+ }
+ return static_cast<uint32_t>(value);
+}
+
tint::Result<sem::WorkgroupSize> Resolver::WorkgroupAttribute(const ast::WorkgroupAttribute* attr) {
// Set work-group size defaults.
sem::WorkgroupSize ws;
diff --git a/src/tint/lang/wgsl/resolver/resolver.h b/src/tint/lang/wgsl/resolver/resolver.h
index df18f91..1781493 100644
--- a/src/tint/lang/wgsl/resolver/resolver.h
+++ b/src/tint/lang/wgsl/resolver/resolver.h
@@ -442,6 +442,11 @@
/// @returns the group value on success.
tint::Result<uint32_t> GroupAttribute(const ast::GroupAttribute* attr);
+ /// Resolves the `@input_attachment_index` attribute @p attr
+ /// @returns the index value on success.
+ tint::Result<uint32_t> InputAttachmentIndexAttribute(
+ const ast::InputAttachmentIndexAttribute* attr);
+
/// Resolves the `@workgroup_size` attribute @p attr
/// @returns the workgroup size on success.
tint::Result<sem::WorkgroupSize> WorkgroupAttribute(const ast::WorkgroupAttribute* attr);
diff --git a/src/tint/lang/wgsl/resolver/validator.cc b/src/tint/lang/wgsl/resolver/validator.cc
index f482497..559aa50 100644
--- a/src/tint/lang/wgsl/resolver/validator.cc
+++ b/src/tint/lang/wgsl/resolver/validator.cc
@@ -442,6 +442,12 @@
}
bool Validator::InputAttachment(const core::type::InputAttachment* t, const Source& source) const {
+ if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumInternalInputAttachments)) {
+ AddError(source) << "use of " << style::Type("input_attachment")
+ << " requires enabling extension "
+ << style::Code("chromium_internal_input_attachments");
+ return false;
+ }
if (!t->type()->UnwrapRef()->IsAnyOf<core::type::F32, core::type::I32, core::type::U32>()) {
AddError(source) << "input_attachment<type>: type must be f32, i32 or u32";
return false;
@@ -450,6 +456,29 @@
return true;
}
+bool Validator::InputAttachmentIndexAttribute(const ast::InputAttachmentIndexAttribute* attr,
+ const core::type::Type* type,
+ const Source& source) const {
+ if (!enabled_extensions_.Contains(wgsl::Extension::kChromiumInternalInputAttachments)) {
+ AddError(source) << "use of " << style::Attribute("@input_attachment_index")
+ << " requires enabling extension "
+ << style::Code("chromium_internal_input_attachments");
+ return false;
+ }
+
+ if (!type->Is<core::type::InputAttachment>()) {
+ std::string invalid_type = sem_.TypeNameOf(type);
+ AddError(source) << "cannot apply " << style::Attribute("@input_attachment_index")
+ << " to declaration of type " << style::Type(invalid_type);
+ AddNote(attr->source) << style::Attribute("@input_attachment_index")
+ << " must only be applied to declarations of "
+ << style::Type("input_attachment") << " type";
+ return false;
+ }
+
+ return true;
+}
+
bool Validator::Materialize(const core::type::Type* to,
const core::type::Type* from,
const Source& source) const {
@@ -708,6 +737,13 @@
return false;
}
+ auto* input_attachment_index_attr =
+ ast::GetAttribute<ast::InputAttachmentIndexAttribute>(decl->attributes);
+ if (input_attachment_index_attr &&
+ !InputAttachmentIndexAttribute(input_attachment_index_attr, global->Type()->UnwrapRef(),
+ decl->source)) {
+ return false;
+ }
switch (global->AddressSpace()) {
case core::AddressSpace::kUniform:
case core::AddressSpace::kStorage:
@@ -720,6 +756,13 @@
<< style::Attribute("@binding") << " attributes";
return false;
}
+ if (global->Type()->UnwrapRef()->Is<core::type::InputAttachment>() &&
+ !input_attachment_index_attr) {
+ AddError(decl->source)
+ << style::Type("input_attachment") << " variables require "
+ << style::Attribute("@input_attachment_index") << " attribute";
+ return false;
+ }
break;
}
default: {
diff --git a/src/tint/lang/wgsl/resolver/validator.h b/src/tint/lang/wgsl/resolver/validator.h
index 6cdbfaf..fb54048 100644
--- a/src/tint/lang/wgsl/resolver/validator.h
+++ b/src/tint/lang/wgsl/resolver/validator.h
@@ -434,6 +434,15 @@
/// @returns true on success, false otherwise
bool InputAttachment(const core::type::InputAttachment* t, const Source& source) const;
+ /// Validates a input attachment index attribute
+ /// @param attr the input attachment index attribute to validate
+ /// @param type the variable type
+ /// @param source the source of declaration using the attribute
+ /// @returns true on success, false otherwise.
+ bool InputAttachmentIndexAttribute(const ast::InputAttachmentIndexAttribute* attr,
+ const core::type::Type* type,
+ const Source& source) const;
+
/// Validates a structure
/// @param str the structure to validate
/// @param stage the current pipeline stage
diff --git a/src/tint/lang/wgsl/sem/variable.h b/src/tint/lang/wgsl/sem/variable.h
index bf97929..51cac2a 100644
--- a/src/tint/lang/wgsl/sem/variable.h
+++ b/src/tint/lang/wgsl/sem/variable.h
@@ -169,6 +169,8 @@
/// @note a GlobalVariable generally doesn't have a `color` in WGSL, as it isn't allowed by
/// the spec. The location maybe attached by transforms such as CanonicalizeEntryPointIO.
std::optional<uint32_t> color;
+ /// The `input_attachment_index` attribute value for the variable, if set
+ std::optional<uint32_t> input_attachment_index;
};
/// GlobalVariable is a module-scope variable
