tint: Add basic support for chromium_experimental_push_constant.
This extension adds support for the push_constant storage class such
that it can be tested with WGSL test files. The real goal is to allow
future transforms that will add push constants that the SPIRV writer
will output.
The extension:
- Adds the `chromium_experimental_push_constant` enable.
- Allows the push_constant storage class for global variables.
- Adds validation that the types are host-shareable for push_constant
variables, and that they don't contain f16 (must be 32bit types
only).
- Validates that at most one push_constant variable is statically used
per entry-point.
- Skips validation that the extension has been enabled if
kIgnoreStorageClass is used.
Tests are added:
- For parsing of var<push_constant>
- Caught a missing conversion.
- For each of the validation rules.
- For the wrapping of push constants in structs if needed by
AddSpirvBlockAttribute.
- For the layout and type rules of the storage class.
- For a shader with multiple entry-points using various push constants.
- Caught a missing reset of the previous push constant variable in
the validation check that at most one is used.
- Caught the missing wrapping in structs that had to be added to
AddSpirvBlockAttribute.
- Caught incorrect logic when adding diagnostics about the call
graph leading to the reference to push constants.
Bug: tint:1620
Change-Id: I04a5d8e5188c0dcef077f2233ba1359d1575bf51
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/96682
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/ast/extension.cc b/src/tint/ast/extension.cc
index 4283df1..e57b248 100644
--- a/src/tint/ast/extension.cc
+++ b/src/tint/ast/extension.cc
@@ -37,6 +37,9 @@
if (str == "chromium_disable_uniformity_analysis") {
return Extension::kChromiumDisableUniformityAnalysis;
}
+ if (str == "chromium_experimental_push_constant") {
+ return Extension::kChromiumExperimentalPushConstant;
+ }
return Extension::kInvalid;
}
@@ -50,6 +53,8 @@
return out << "chromium_experimental_dp4a";
case Extension::kChromiumDisableUniformityAnalysis:
return out << "chromium_disable_uniformity_analysis";
+ case Extension::kChromiumExperimentalPushConstant:
+ return out << "chromium_experimental_push_constant";
}
return out << "<unknown>";
}
diff --git a/src/tint/ast/extension.h b/src/tint/ast/extension.h
index 32bf507..ce3d48a 100644
--- a/src/tint/ast/extension.h
+++ b/src/tint/ast/extension.h
@@ -36,6 +36,7 @@
kF16,
kChromiumExperimentalDp4A,
kChromiumDisableUniformityAnalysis,
+ kChromiumExperimentalPushConstant,
};
/// @param out the std::ostream to write to
diff --git a/src/tint/ast/extension_bench.cc b/src/tint/ast/extension_bench.cc
index 47787e3..8fc9d1c 100644
--- a/src/tint/ast/extension_bench.cc
+++ b/src/tint/ast/extension_bench.cc
@@ -52,6 +52,13 @@
"chromiuE_disable_uniformity_analysis",
"chromium_disable_uniTTormity_aPPalsis",
"ddhromium_disabexxuniformity_analysis",
+ "c44romium_experimental_push_constant",
+ "chromium_experimental_pSSsVV_constant",
+ "chrom22Rm_experimental_pushRonstant",
+ "chromium_experimental_push_constant",
+ "chromium_exp9rimFntal_ush_constant",
+ "chrmium_experimental_push_constant",
+ "cOOromium_experiVeHtal_puh_conRRtant",
};
for (auto _ : state) {
for (auto& str : kStrings) {
diff --git a/src/tint/ast/extension_test.cc b/src/tint/ast/extension_test.cc
index 8c75613..283088d 100644
--- a/src/tint/ast/extension_test.cc
+++ b/src/tint/ast/extension_test.cc
@@ -45,6 +45,7 @@
{"f16", Extension::kF16},
{"chromium_experimental_dp4a", Extension::kChromiumExperimentalDp4A},
{"chromium_disable_uniformity_analysis", Extension::kChromiumDisableUniformityAnalysis},
+ {"chromium_experimental_push_constant", Extension::kChromiumExperimentalPushConstant},
};
static constexpr Case kInvalidCases[] = {
@@ -57,6 +58,9 @@
{"chromiumppdisableqquniformity_aalysHHs", Extension::kInvalid},
{"chromiu_disable_unifovmitc_analyi", Extension::kInvalid},
{"chromium_diable_uGbformity_analysis", Extension::kInvalid},
+ {"chvomium_experimental_push_constiint", Extension::kInvalid},
+ {"chromiu8WWexperimental_push_constant", Extension::kInvalid},
+ {"chromium_experiMental_push_costanxx", Extension::kInvalid},
};
using ExtensionParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/ast/storage_class.cc b/src/tint/ast/storage_class.cc
index 303c04e..450067b 100644
--- a/src/tint/ast/storage_class.cc
+++ b/src/tint/ast/storage_class.cc
@@ -43,6 +43,9 @@
if (str == "storage") {
return StorageClass::kStorage;
}
+ if (str == "push_constant") {
+ return StorageClass::kPushConstant;
+ }
return StorageClass::kInvalid;
}
@@ -62,6 +65,8 @@
return out << "uniform";
case StorageClass::kStorage:
return out << "storage";
+ case StorageClass::kPushConstant:
+ return out << "push_constant";
case StorageClass::kHandle:
return out << "handle";
case StorageClass::kIn:
diff --git a/src/tint/ast/storage_class.h b/src/tint/ast/storage_class.h
index 4da3db0..43e41e1 100644
--- a/src/tint/ast/storage_class.h
+++ b/src/tint/ast/storage_class.h
@@ -36,6 +36,7 @@
kWorkgroup,
kUniform,
kStorage,
+ kPushConstant,
kHandle, // Tint-internal enum entry - not parsed
kIn, // Tint-internal enum entry - not parsed
kOut, // Tint-internal enum entry - not parsed
@@ -55,7 +56,8 @@
/// @param sc the StorageClass
/// @see https://gpuweb.github.io/gpuweb/wgsl.html#host-shareable
inline bool IsHostShareable(StorageClass sc) {
- return sc == ast::StorageClass::kUniform || sc == ast::StorageClass::kStorage;
+ return sc == ast::StorageClass::kUniform || sc == ast::StorageClass::kStorage ||
+ sc == ast::StorageClass::kPushConstant;
}
} // namespace tint::ast
diff --git a/src/tint/ast/storage_class.h.tmpl b/src/tint/ast/storage_class.h.tmpl
index d885c72..c21d6f6 100644
--- a/src/tint/ast/storage_class.h.tmpl
+++ b/src/tint/ast/storage_class.h.tmpl
@@ -28,7 +28,8 @@
/// @param sc the StorageClass
/// @see https://gpuweb.github.io/gpuweb/wgsl.html#host-shareable
inline bool IsHostShareable(StorageClass sc) {
- return sc == ast::StorageClass::kUniform || sc == ast::StorageClass::kStorage;
+ return sc == ast::StorageClass::kUniform || sc == ast::StorageClass::kStorage ||
+ sc == ast::StorageClass::kPushConstant;
}
} // namespace tint::ast
diff --git a/src/tint/ast/storage_class_bench.cc b/src/tint/ast/storage_class_bench.cc
index d8e9ae9..00095c2 100644
--- a/src/tint/ast/storage_class_bench.cc
+++ b/src/tint/ast/storage_class_bench.cc
@@ -31,12 +31,48 @@
void StorageClassParser(::benchmark::State& state) {
std::array kStrings{
- "fccnctin", "ucti3", "functVon", "function", "1unction", "unJtqqon",
- "llun77tion", "ppqqivtHH", "prcv", "bivaGe", "private", "priviive",
- "8WWivate", "pxxvate", "wXkgrggup", "worXVup", "3orkgroup", "workgroup",
- "workgroEp", "woTTPkroup", "ddorkroxxp", "u44iform", "unSSfoVVm", "RniR22m",
- "uniform", "uFfo9m", "uniorm", "VOORRHrm", "straye", "llntrrr77ge",
- "stor4g00", "storage", "trooe", "zzrage", "siioppa1",
+ "fccnctin",
+ "ucti3",
+ "functVon",
+ "function",
+ "1unction",
+ "unJtqqon",
+ "llun77tion",
+ "ppqqivtHH",
+ "prcv",
+ "bivaGe",
+ "private",
+ "priviive",
+ "8WWivate",
+ "pxxvate",
+ "wXkgrggup",
+ "worXVup",
+ "3orkgroup",
+ "workgroup",
+ "workgroEp",
+ "woTTPkroup",
+ "ddorkroxxp",
+ "u44iform",
+ "unSSfoVVm",
+ "RniR22m",
+ "uniform",
+ "uFfo9m",
+ "uniorm",
+ "VOORRHrm",
+ "straye",
+ "llntrrr77ge",
+ "stor4g00",
+ "storage",
+ "trooe",
+ "zzrage",
+ "siioppa1",
+ "puXXh_constant",
+ "pusII9_nn55nstant",
+ "YusHH_coaastSSrnt",
+ "push_constant",
+ "pushonkkHan",
+ "jush_consgRt",
+ "puh_cobsant",
};
for (auto _ : state) {
for (auto& str : kStrings) {
diff --git a/src/tint/ast/storage_class_test.cc b/src/tint/ast/storage_class_test.cc
index 9688236..1d12d45 100644
--- a/src/tint/ast/storage_class_test.cc
+++ b/src/tint/ast/storage_class_test.cc
@@ -44,18 +44,19 @@
static constexpr Case kValidCases[] = {
{"function", StorageClass::kFunction}, {"private", StorageClass::kPrivate},
{"workgroup", StorageClass::kWorkgroup}, {"uniform", StorageClass::kUniform},
- {"storage", StorageClass::kStorage},
+ {"storage", StorageClass::kStorage}, {"push_constant", StorageClass::kPushConstant},
};
static constexpr Case kInvalidCases[] = {
- {"fccnctin", StorageClass::kInvalid}, {"ucti3", StorageClass::kInvalid},
- {"functVon", StorageClass::kInvalid}, {"priv1te", StorageClass::kInvalid},
- {"pqiJate", StorageClass::kInvalid}, {"privat7ll", StorageClass::kInvalid},
- {"workroppqHH", StorageClass::kInvalid}, {"workru", StorageClass::kInvalid},
- {"wbkgGoup", StorageClass::kInvalid}, {"unifiivm", StorageClass::kInvalid},
- {"8WWiform", StorageClass::kInvalid}, {"uxxform", StorageClass::kInvalid},
- {"sXraggg", StorageClass::kInvalid}, {"traXe", StorageClass::kInvalid},
- {"stor3ge", StorageClass::kInvalid},
+ {"fccnctin", StorageClass::kInvalid}, {"ucti3", StorageClass::kInvalid},
+ {"functVon", StorageClass::kInvalid}, {"priv1te", StorageClass::kInvalid},
+ {"pqiJate", StorageClass::kInvalid}, {"privat7ll", StorageClass::kInvalid},
+ {"workroppqHH", StorageClass::kInvalid}, {"workru", StorageClass::kInvalid},
+ {"wbkgGoup", StorageClass::kInvalid}, {"unifiivm", StorageClass::kInvalid},
+ {"8WWiform", StorageClass::kInvalid}, {"uxxform", StorageClass::kInvalid},
+ {"sXraggg", StorageClass::kInvalid}, {"traXe", StorageClass::kInvalid},
+ {"stor3ge", StorageClass::kInvalid}, {"push_constanE", StorageClass::kInvalid},
+ {"push_TTPnstant", StorageClass::kInvalid}, {"puxxdh_constan", StorageClass::kInvalid},
};
using StorageClassParseTest = testing::TestWithParam<Case>;
diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index 6e1d8fb..953000b 100644
--- a/src/tint/intrinsics.def
+++ b/src/tint/intrinsics.def
@@ -49,6 +49,8 @@
chromium_experimental_dp4a
// A Chromium-specific extension for disabling uniformity analysis.
chromium_disable_uniformity_analysis
+ // A Chromium-specific extension for push constants
+ chromium_experimental_push_constant
}
// https://gpuweb.github.io/gpuweb/wgsl/#storage-class
@@ -59,6 +61,7 @@
workgroup
uniform
storage
+ push_constant
@internal handle
@internal in
@internal out
diff --git a/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc b/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
index 3ed0948..2fee83c 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
@@ -88,6 +88,17 @@
EXPECT_EQ(v->source.range.end.column, 20u);
}
+TEST_F(ParserImplTest, VariableDecl_WithPushConstant) {
+ auto p = parser("var<push_constant> my_var : f32");
+ auto v = p->variable_decl();
+ EXPECT_TRUE(v.matched);
+ EXPECT_FALSE(v.errored);
+ EXPECT_FALSE(p->has_error());
+ EXPECT_EQ(v->name, "my_var");
+ EXPECT_TRUE(v->type->Is<ast::F32>());
+ EXPECT_EQ(v->storage_class, ast::StorageClass::kPushConstant);
+}
+
TEST_F(ParserImplTest, VariableDecl_InvalidStorageClass) {
auto p = parser("var<unknown> my_var : f32");
auto v = p->variable_decl();
diff --git a/src/tint/resolver/entry_point_validation_test.cc b/src/tint/resolver/entry_point_validation_test.cc
index 8df0425..5437828 100644
--- a/src/tint/resolver/entry_point_validation_test.cc
+++ b/src/tint/resolver/entry_point_validation_test.cc
@@ -306,6 +306,129 @@
"in its return type");
}
+TEST_F(ResolverEntryPointValidationTest, PushConstantAllowedWithEnable) {
+ // enable chromium_experimental_push_constant;
+ // var<push_constant> a : u32;
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar("a", ty.u32(), ast::StorageClass::kPushConstant);
+
+ EXPECT_TRUE(r()->Resolve());
+}
+
+TEST_F(ResolverEntryPointValidationTest, PushConstantDisallowedWithoutEnable) {
+ // var<push_constant> a : u32;
+ GlobalVar(Source{{1, 2}}, "a", ty.u32(), ast::StorageClass::kPushConstant);
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(),
+ "1:2 error: use of variable storage class 'push_constant' requires enabling "
+ "extension 'chromium_experimental_push_constant'");
+}
+
+TEST_F(ResolverEntryPointValidationTest, PushConstantAllowedWithIgnoreStorageClassAttribute) {
+ // var<push_constant> a : u32; // With ast::DisabledValidation::kIgnoreStorageClass
+ GlobalVar("a", ty.u32(), ast::StorageClass::kPushConstant,
+ ast::AttributeList{Disable(ast::DisabledValidation::kIgnoreStorageClass)});
+
+ EXPECT_TRUE(r()->Resolve());
+}
+
+TEST_F(ResolverEntryPointValidationTest, PushConstantOneVariableUsedInEntryPoint) {
+ // enable chromium_experimental_push_constant;
+ // var<push_constant> a : u32;
+ // @compute @workgroup_size(1) fn main() {
+ // _ = a;
+ // }
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar("a", ty.u32(), ast::StorageClass::kPushConstant);
+
+ Func("main", {}, ty.void_(), {Assign(Phony(), "a")},
+ {Stage(ast::PipelineStage::kCompute), create<ast::WorkgroupAttribute>(Expr(1_i))});
+
+ EXPECT_TRUE(r()->Resolve());
+}
+
+TEST_F(ResolverEntryPointValidationTest, PushConstantTwoVariablesUsedInEntryPoint) {
+ // enable chromium_experimental_push_constant;
+ // var<push_constant> a : u32;
+ // var<push_constant> b : u32;
+ // @compute @workgroup_size(1) fn main() {
+ // _ = a;
+ // _ = b;
+ // }
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar(Source{{1, 2}}, "a", ty.u32(), ast::StorageClass::kPushConstant);
+ GlobalVar(Source{{3, 4}}, "b", ty.u32(), ast::StorageClass::kPushConstant);
+
+ Func(Source{{5, 6}}, "main", {}, ty.void_(), {Assign(Phony(), "a"), Assign(Phony(), "b")},
+ {Stage(ast::PipelineStage::kCompute), create<ast::WorkgroupAttribute>(Expr(1_i))});
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(),
+ R"(5:6 error: entry point 'main' uses two different 'push_constant' variables.
+3:4 note: first 'push_constant' variable declaration is here
+1:2 note: second 'push_constant' variable declaration is here)");
+}
+
+TEST_F(ResolverEntryPointValidationTest,
+ PushConstantTwoVariablesUsedInEntryPointWithFunctionGraph) {
+ // enable chromium_experimental_push_constant;
+ // var<push_constant> a : u32;
+ // var<push_constant> b : u32;
+ // fn uses_a() {
+ // _ = a;
+ // }
+ // fn uses_b() {
+ // _ = b;
+ // }
+ // @compute @workgroup_size(1) fn main() {
+ // uses_a();
+ // uses_b();
+ // }
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar(Source{{1, 2}}, "a", ty.u32(), ast::StorageClass::kPushConstant);
+ GlobalVar(Source{{3, 4}}, "b", ty.u32(), ast::StorageClass::kPushConstant);
+
+ Func(Source{{5, 6}}, "uses_a", {}, ty.void_(), {Assign(Phony(), "a")});
+ Func(Source{{7, 8}}, "uses_b", {}, ty.void_(), {Assign(Phony(), "b")});
+
+ Func(Source{{9, 10}}, "main", {}, ty.void_(),
+ {CallStmt(Call("uses_a")), CallStmt(Call("uses_b"))},
+ {Stage(ast::PipelineStage::kCompute), create<ast::WorkgroupAttribute>(Expr(1_i))});
+
+ EXPECT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(),
+ R"(9:10 error: entry point 'main' uses two different 'push_constant' variables.
+3:4 note: first 'push_constant' variable declaration is here
+7:8 note: called by function 'uses_b'
+9:10 note: called by entry point 'main'
+1:2 note: second 'push_constant' variable declaration is here
+5:6 note: called by function 'uses_a'
+9:10 note: called by entry point 'main')");
+}
+
+TEST_F(ResolverEntryPointValidationTest, PushConstantTwoVariablesUsedInDifferentEntryPoint) {
+ // enable chromium_experimental_push_constant;
+ // var<push_constant> a : u32;
+ // var<push_constant> b : u32;
+ // @compute @workgroup_size(1) fn uses_a() {
+ // _ = a;
+ // }
+ // @compute @workgroup_size(1) fn uses_b() {
+ // _ = a;
+ // }
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar("a", ty.u32(), ast::StorageClass::kPushConstant);
+ GlobalVar("b", ty.u32(), ast::StorageClass::kPushConstant);
+
+ Func("uses_a", {}, ty.void_(), {Assign(Phony(), "a")},
+ {Stage(ast::PipelineStage::kCompute), create<ast::WorkgroupAttribute>(Expr(1_i))});
+ Func("uses_b", {}, ty.void_(), {Assign(Phony(), "b")},
+ {Stage(ast::PipelineStage::kCompute), create<ast::WorkgroupAttribute>(Expr(1_i))});
+
+ EXPECT_TRUE(r()->Resolve());
+}
+
namespace TypeValidationTests {
struct Params {
builder::ast_type_func_ptr create_ast_type;
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index 097d591..e5d4c65 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -159,6 +159,10 @@
return false;
}
+ if (!validator_.PushConstants(entry_points_)) {
+ return false;
+ }
+
if (!enabled_extensions_.contains(ast::Extension::kChromiumDisableUniformityAnalysis)) {
if (!AnalyzeUniformity(builder_, dependencies_)) {
// TODO(jrprice): Reject programs that fail uniformity analysis.
@@ -726,7 +730,7 @@
// TODO(bclayton): Call this at the end of resolve on all uniform and storage
// referenced structs
- if (!validator_.StorageClassLayout(sem, valid_type_storage_layouts_)) {
+ if (!validator_.StorageClassLayout(sem, enabled_extensions_, valid_type_storage_layouts_)) {
return nullptr;
}
diff --git a/src/tint/resolver/storage_class_layout_validation_test.cc b/src/tint/resolver/storage_class_layout_validation_test.cc
index 770361a..6952cb5 100644
--- a/src/tint/resolver/storage_class_layout_validation_test.cc
+++ b/src/tint/resolver/storage_class_layout_validation_test.cc
@@ -523,5 +523,47 @@
ASSERT_TRUE(r()->Resolve()) << r()->error();
}
+// Detect unaligned member for push constants buffers
+TEST_F(ResolverStorageClassLayoutValidationTest, PushConstant_UnalignedMember) {
+ // enable chromium_experimental_push_constant;
+ // struct S {
+ // @size(5) a : f32;
+ // @align(1) b : f32;
+ // };
+ // var<push_constant> a : S;
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ Structure(Source{{12, 34}}, "S",
+ {Member("a", ty.f32(), {MemberSize(5)}),
+ Member(Source{{34, 56}}, "b", ty.f32(), {MemberAlign(1)})});
+ GlobalVar(Source{{78, 90}}, "a", ty.type_name("S"), ast::StorageClass::kPushConstant);
+
+ ASSERT_FALSE(r()->Resolve());
+ EXPECT_EQ(
+ r()->error(),
+ R"(34:56 error: the offset of a struct member of type 'f32' in storage class 'push_constant' must be a multiple of 4 bytes, but 'b' is currently at offset 5. Consider setting @align(4) on this member
+12:34 note: see layout of struct:
+/* align(4) size(12) */ struct S {
+/* offset(0) align(4) size( 5) */ a : f32;
+/* offset(5) align(1) size( 4) */ b : f32;
+/* offset(9) align(1) size( 3) */ // -- implicit struct size padding --;
+/* */ };
+78:90 note: see declaration of variable)");
+}
+
+TEST_F(ResolverStorageClassLayoutValidationTest, PushConstant_Aligned) {
+ // enable chromium_experimental_push_constant;
+ // struct S {
+ // @size(5) a : f32;
+ // @align(4) b : f32;
+ // };
+ // var<push_constant> a : S;
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ Structure("S",
+ {Member("a", ty.f32(), {MemberSize(5)}), Member("b", ty.f32(), {MemberAlign(4)})});
+ GlobalVar("a", ty.type_name("S"), ast::StorageClass::kPushConstant);
+
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
} // namespace
} // namespace tint::resolver
diff --git a/src/tint/resolver/storage_class_validation_test.cc b/src/tint/resolver/storage_class_validation_test.cc
index f3e6f32..5e6e9ce 100644
--- a/src/tint/resolver/storage_class_validation_test.cc
+++ b/src/tint/resolver/storage_class_validation_test.cc
@@ -598,5 +598,89 @@
ASSERT_TRUE(r()->Resolve()) << r()->error();
}
+TEST_F(ResolverStorageClassValidationTest, PushConstantBool) {
+ // enable chromium_experimental_push_constant;
+ // var<push_constant> g : bool;
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar(Source{{56, 78}}, "g", ty.bool_(), ast::StorageClass::kPushConstant);
+
+ ASSERT_FALSE(r()->Resolve());
+ EXPECT_EQ(
+ r()->error(),
+ R"(56:78 error: Type 'bool' cannot be used in storage class 'push_constant' as it is non-host-shareable
+56:78 note: while instantiating 'var' g)");
+}
+
+TEST_F(ResolverStorageClassValidationTest, PushConstantF16) {
+ // enable chromium_experimental_push_constant;
+ // enable f16;
+ // var<push_constant> g : f16;
+ Enable(ast::Extension::kF16);
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar("g", ty.f16(Source{{56, 78}}), ast::StorageClass::kPushConstant);
+
+ ASSERT_FALSE(r()->Resolve());
+ EXPECT_EQ(r()->error(),
+ "56:78 error: using f16 types in 'push_constant' storage class is not "
+ "implemented yet");
+}
+
+TEST_F(ResolverStorageClassValidationTest, PushConstantPointer) {
+ // enable chromium_experimental_push_constant;
+ // var<push_constant> g : ptr<private, f32>;
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar(Source{{56, 78}}, "g", ty.pointer(ty.f32(), ast::StorageClass::kPrivate),
+ ast::StorageClass::kPushConstant);
+
+ ASSERT_FALSE(r()->Resolve());
+ EXPECT_EQ(
+ r()->error(),
+ R"(56:78 error: Type 'ptr<private, f32, read_write>' cannot be used in storage class 'push_constant' as it is non-host-shareable
+56:78 note: while instantiating 'var' g)");
+}
+
+TEST_F(ResolverStorageClassValidationTest, PushConstantIntScalar) {
+ // enable chromium_experimental_push_constant;
+ // var<push_constant> g : i32;
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar("g", ty.i32(), ast::StorageClass::kPushConstant);
+
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverStorageClassValidationTest, PushConstantVectorF32) {
+ // enable chromium_experimental_push_constant;
+ // var<push_constant> g : vec4<f32>;
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar("g", ty.vec4<f32>(), ast::StorageClass::kPushConstant);
+
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverStorageClassValidationTest, PushConstantArrayF32) {
+ // enable chromium_experimental_push_constant;
+ // struct S { a : f32}
+ // var<push_constant> g : array<S, 3u>;
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ auto* s = Structure("S", {Member("a", ty.f32())});
+ auto* a = ty.array(ty.Of(s), 3_u);
+ GlobalVar("g", a, ast::StorageClass::kPushConstant);
+
+ ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverStorageClassValidationTest, PushConstantWithInitializer) {
+ // enable chromium_experimental_push_constant;
+ // var<push_constant> a : u32 = 0u;
+ Enable(ast::Extension::kChromiumExperimentalPushConstant);
+ GlobalVar(Source{{1u, 2u}}, "a", ty.u32(), ast::StorageClass::kPushConstant,
+ Expr(Source{{3u, 4u}}, u32(0)));
+
+ ASSERT_FALSE(r()->Resolve());
+ EXPECT_EQ(
+ r()->error(),
+ R"(1:2 error: var of storage class 'push_constant' cannot have an initializer. var initializers are only supported for the storage classes 'private' and 'function')");
+}
+
} // namespace
} // namespace tint::resolver
diff --git a/src/tint/resolver/validator.cc b/src/tint/resolver/validator.cc
index ec858c4..61f13cb 100644
--- a/src/tint/resolver/validator.cc
+++ b/src/tint/resolver/validator.cc
@@ -395,7 +395,7 @@
// Temporally forbid using f16 types in "uniform" and "storage" storage class.
// TODO(tint:1473, tint:1502): Remove this error after f16 is supported in "uniform" and
- // "storage" storage class.
+ // "storage" storage class but keep for "push_constant" storage class.
if (Is<sem::F16>(sem::Type::DeepestElementOf(store_ty))) {
AddError(
"using f16 types in '" + utils::ToString(sc) + "' storage class is not implemented yet",
@@ -516,7 +516,19 @@
}
bool Validator::StorageClassLayout(const sem::Variable* var,
+ const ast::Extensions& enabled_extensions,
ValidTypeStorageLayouts& layouts) const {
+ if (var->StorageClass() == ast::StorageClass::kPushConstant &&
+ !enabled_extensions.contains(ast::Extension::kChromiumExperimentalPushConstant) &&
+ IsValidationEnabled(var->Declaration()->attributes,
+ ast::DisabledValidation::kIgnoreStorageClass)) {
+ AddError(
+ "use of variable storage class 'push_constant' requires enabling extension "
+ "'chromium_experimental_push_constant'",
+ var->Declaration()->source);
+ return false;
+ }
+
if (auto* str = var->Type()->UnwrapRef()->As<sem::Struct>()) {
if (!StorageClassLayout(str, var->StorageClass(), str->Declaration()->source, layouts)) {
AddNote("see declaration of variable", var->Declaration()->source);
@@ -2000,6 +2012,73 @@
return true;
}
+bool Validator::PushConstants(const std::vector<sem::Function*>& entry_points) const {
+ for (auto* entry_point : entry_points) {
+ // State checked and modified by check_push_constant so that it remembers previously seen
+ // push_constant variables for an entry-point.
+ const sem::Variable* push_constant_var = nullptr;
+ const sem::Function* push_constant_func = nullptr;
+
+ auto check_push_constant = [&](const sem::Function* func, const sem::Function* ep) {
+ for (auto* var : func->DirectlyReferencedGlobals()) {
+ if (var->StorageClass() != ast::StorageClass::kPushConstant ||
+ var == push_constant_var) {
+ continue;
+ }
+
+ if (push_constant_var == nullptr) {
+ push_constant_var = var;
+ push_constant_func = func;
+ continue;
+ }
+
+ AddError("entry point '" + symbols_.NameFor(ep->Declaration()->symbol) +
+ "' uses two different 'push_constant' variables.",
+ ep->Declaration()->source);
+ AddNote("first 'push_constant' variable declaration is here",
+ var->Declaration()->source);
+ if (func != ep) {
+ TraverseCallChain(diagnostics_, ep, func, [&](const sem::Function* f) {
+ AddNote("called by function '" +
+ symbols_.NameFor(f->Declaration()->symbol) + "'",
+ f->Declaration()->source);
+ });
+ AddNote("called by entry point '" +
+ symbols_.NameFor(ep->Declaration()->symbol) + "'",
+ ep->Declaration()->source);
+ }
+ AddNote("second 'push_constant' variable declaration is here",
+ push_constant_var->Declaration()->source);
+ if (push_constant_func != ep) {
+ TraverseCallChain(
+ diagnostics_, ep, push_constant_func, [&](const sem::Function* f) {
+ AddNote("called by function '" +
+ symbols_.NameFor(f->Declaration()->symbol) + "'",
+ f->Declaration()->source);
+ });
+ AddNote("called by entry point '" +
+ symbols_.NameFor(ep->Declaration()->symbol) + "'",
+ ep->Declaration()->source);
+ }
+ return false;
+ }
+
+ return true;
+ };
+
+ if (!check_push_constant(entry_point, entry_point)) {
+ return false;
+ }
+ for (auto* func : entry_point->TransitivelyCalledFunctions()) {
+ if (!check_push_constant(func, entry_point)) {
+ return false;
+ }
+ }
+ }
+
+ return true;
+}
+
bool Validator::Array(const sem::Array* arr, const Source& source) const {
auto* el_ty = arr->ElemType();
diff --git a/src/tint/resolver/validator.h b/src/tint/resolver/validator.h
index 385e020..935195e 100644
--- a/src/tint/resolver/validator.h
+++ b/src/tint/resolver/validator.h
@@ -116,6 +116,11 @@
/// @returns true on success, false otherwise.
bool PipelineStages(const std::vector<sem::Function*>& entry_points) const;
+ /// Validates push_constant variables
+ /// @param entry_points the entry points to the module
+ /// @returns true on success, false otherwise.
+ bool PushConstants(const std::vector<sem::Function*>& entry_points) const;
+
/// Validates aliases
/// @param alias the alias to validate
/// @returns true on success, false otherwise.
@@ -433,8 +438,11 @@
/// Validates a storage class layout
/// @param var the variable to validate
/// @param layouts previously validated storage layouts
+ /// @param enabled_extensions all the extensions declared in current module
/// @returns true on success, false otherwise.
- bool StorageClassLayout(const sem::Variable* var, ValidTypeStorageLayouts& layouts) const;
+ bool StorageClassLayout(const sem::Variable* var,
+ const ast::Extensions& enabled_extensions,
+ ValidTypeStorageLayouts& layouts) const;
/// @returns true if the attribute list contains a
/// ast::DisableValidationAttribute with the validation mode equal to
diff --git a/src/tint/transform/add_spirv_block_attribute.cc b/src/tint/transform/add_spirv_block_attribute.cc
index 85c20c3..06a199d 100644
--- a/src/tint/transform/add_spirv_block_attribute.cc
+++ b/src/tint/transform/add_spirv_block_attribute.cc
@@ -58,7 +58,8 @@
for (auto* var : ctx.src->AST().Globals<ast::Var>()) {
auto* sem_var = sem.Get<sem::GlobalVariable>(var);
if (var->declared_storage_class != ast::StorageClass::kStorage &&
- var->declared_storage_class != ast::StorageClass::kUniform) {
+ var->declared_storage_class != ast::StorageClass::kUniform &&
+ var->declared_storage_class != ast::StorageClass::kPushConstant) {
continue;
}
diff --git a/src/tint/transform/add_spirv_block_attribute_test.cc b/src/tint/transform/add_spirv_block_attribute_test.cc
index 455be60..90f9219 100644
--- a/src/tint/transform/add_spirv_block_attribute_test.cc
+++ b/src/tint/transform/add_spirv_block_attribute_test.cc
@@ -196,6 +196,71 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(AddSpirvBlockAttributeTest, BasicScalar_PushConstant) {
+ auto* src = R"(
+enable chromium_experimental_push_constant;
+var<push_constant> u : f32;
+
+@fragment
+fn main() {
+ let f = u;
+}
+)";
+ auto* expect = R"(
+enable chromium_experimental_push_constant;
+
+@internal(spirv_block)
+struct u_block {
+ inner : f32,
+}
+
+var<push_constant> u : u_block;
+
+@fragment
+fn main() {
+ let f = u.inner;
+}
+)";
+
+ auto got = Run<AddSpirvBlockAttribute>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AddSpirvBlockAttributeTest, BasicStruct_PushConstant) {
+ auto* src = R"(
+enable chromium_experimental_push_constant;
+struct S {
+ f : f32,
+};
+var<push_constant> u : S;
+
+@fragment
+fn main() {
+ let f = u.f;
+}
+)";
+ auto* expect = R"(
+enable chromium_experimental_push_constant;
+
+@internal(spirv_block)
+struct S {
+ f : f32,
+}
+
+var<push_constant> u : S;
+
+@fragment
+fn main() {
+ let f = u.f;
+}
+)";
+
+ auto got = Run<AddSpirvBlockAttribute>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(AddSpirvBlockAttributeTest, Nested_OuterBuffer_InnerNotBuffer) {
auto* src = R"(
struct Inner {
diff --git a/src/tint/transform/module_scope_var_to_entry_point_param.cc b/src/tint/transform/module_scope_var_to_entry_point_param.cc
index 7bd3991..6d313d2 100644
--- a/src/tint/transform/module_scope_var_to_entry_point_param.cc
+++ b/src/tint/transform/module_scope_var_to_entry_point_param.cc
@@ -26,6 +26,7 @@
#include "src/tint/sem/module.h"
#include "src/tint/sem/statement.h"
#include "src/tint/sem/variable.h"
+#include "src/tint/utils/string.h"
TINT_INSTANTIATE_TYPEINFO(tint::transform::ModuleScopeVarToEntryPointParam);
@@ -191,9 +192,16 @@
break;
}
+ case ast::StorageClass::kPushConstant: {
+ ctx.dst->Diagnostics().add_error(
+ diag::System::Transform,
+ "unhandled module-scope storage class (" + utils::ToString(sc) + ")");
+ break;
+ }
default: {
TINT_ICE(Transform, ctx.dst->Diagnostics())
<< "unhandled module-scope storage class (" << sc << ")";
+ break;
}
}
}
@@ -219,6 +227,12 @@
case ast::StorageClass::kHandle:
case ast::StorageClass::kWorkgroup:
break;
+ case ast::StorageClass::kPushConstant: {
+ ctx.dst->Diagnostics().add_error(
+ diag::System::Transform,
+ "unhandled module-scope storage class (" + utils::ToString(sc) + ")");
+ break;
+ }
default: {
TINT_ICE(Transform, ctx.dst->Diagnostics())
<< "unhandled module-scope storage class (" << sc << ")";
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index 4c513be..5403136 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -70,6 +70,7 @@
#include "src/tint/utils/defer.h"
#include "src/tint/utils/map.h"
#include "src/tint/utils/scoped_assignment.h"
+#include "src/tint/utils/string.h"
#include "src/tint/writer/append_vector.h"
#include "src/tint/writer/float_to_string.h"
#include "src/tint/writer/generate_external_texture_bindings.h"
@@ -1936,6 +1937,11 @@
case ast::StorageClass::kIn:
case ast::StorageClass::kOut:
return EmitIOVariable(sem);
+ case ast::StorageClass::kPushConstant:
+ diagnostics_.add_error(
+ diag::System::Writer,
+ "unhandled storage class " + utils::ToString(sem->StorageClass()));
+ return false;
default: {
TINT_ICE(Writer, diagnostics_)
<< "unhandled storage class " << sem->StorageClass();
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index d106804..be089c3 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -72,6 +72,7 @@
#include "src/tint/utils/defer.h"
#include "src/tint/utils/map.h"
#include "src/tint/utils/scoped_assignment.h"
+#include "src/tint/utils/string.h"
#include "src/tint/writer/append_vector.h"
#include "src/tint/writer/float_to_string.h"
#include "src/tint/writer/generate_external_texture_bindings.h"
@@ -2849,6 +2850,11 @@
return EmitPrivateVariable(sem);
case ast::StorageClass::kWorkgroup:
return EmitWorkgroupVariable(sem);
+ case ast::StorageClass::kPushConstant:
+ diagnostics_.add_error(
+ diag::System::Writer,
+ "unhandled storage class " + utils::ToString(sem->StorageClass()));
+ return false;
default: {
TINT_ICE(Writer, diagnostics_)
<< "unhandled storage class " << sem->StorageClass();
@@ -2863,6 +2869,7 @@
[&](Default) {
TINT_ICE(Writer, diagnostics_)
<< "unhandled global variable type " << global->TypeInfo().name;
+
return false;
});
}
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index 488cbf8..ccbdeae 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -4121,6 +4121,8 @@
return SpvStorageClassUniform;
case ast::StorageClass::kWorkgroup:
return SpvStorageClassWorkgroup;
+ case ast::StorageClass::kPushConstant:
+ return SpvStorageClassPushConstant;
case ast::StorageClass::kHandle:
return SpvStorageClassUniformConstant;
case ast::StorageClass::kStorage:
