encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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>
This commit is contained in:
dan sinclair 2022-08-02 15:55:35 +00:00 committed by Dawn LUCI CQ
parent 24d36b0227
commit 4abf28e29b
31 changed files with 979 additions and 21 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

28
docs/tint/extensions/chromium_experimental_push_constant.md Normal file
View File

@ -0,0 +1,28 @@
# Chromium Experimental Push Constant
The `chromium_experimental_push_constant` extension adds support for push constant global variables to WGSL.
Push constants are small amounts of data that are passed to the shader and are expected to be more lightweight to set / modify than uniform buffer bindings.
The concept of push constant comes from Vulkan but D3D12 has similar "root constants".
Metal doesn't have the same concept but push constants can be efficiently implemented with the `setBytes` family of command encoder methods.
## Status
Push constant support in Tint is highly experimental and only meant to be used in internal transforms at this stage.
Specification work in the WebGPU group hasn't started.
## Pseudo-specification
This extension adds a new `push_constant` storage class that's only allowed on global variable declarations.
Push constant variables must only contain 32bit data types (or aggregates of such types).
Push constant variable declarations must not have an initializer.
It is an error for a entry point to statically use more than one `push_constant` variable.
## Example usage
```
var<push_constant> draw_id : u32;
@fragment fn main() -> u32 {
return draw_id;
}
```

5
src/tint/ast/extension.cc
View File

@ -37,6 +37,9 @@ Extension ParseExtension(std::string_view str) {
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 @@ std::ostream& operator<<(std::ostream& out, Extension value) {
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>";
}

1
src/tint/ast/extension.h
View File

@ -36,6 +36,7 @@ enum class Extension {
kF16,
kChromiumExperimentalDp4A,
kChromiumDisableUniformityAnalysis,
kChromiumExperimentalPushConstant,
};
/// @param out the std::ostream to write to

7
src/tint/ast/extension_bench.cc
View File

@ -52,6 +52,13 @@ void ExtensionParser(::benchmark::State& state) {
"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) {

4
src/tint/ast/extension_test.cc
View File

@ -45,6 +45,7 @@ static constexpr Case kValidCases[] = {
{"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 @@ static constexpr Case kInvalidCases[] = {
{"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>;

5
src/tint/ast/storage_class.cc
View File

@ -43,6 +43,9 @@ StorageClass ParseStorageClass(std::string_view str) {
if (str == "storage") {
return StorageClass::kStorage;
}
if (str == "push_constant") {
return StorageClass::kPushConstant;
}
return StorageClass::kInvalid;
}
@ -62,6 +65,8 @@ std::ostream& operator<<(std::ostream& out, StorageClass value) {
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:

4
src/tint/ast/storage_class.h
View File

@ -36,6 +36,7 @@ enum class StorageClass {
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 @@ StorageClass ParseStorageClass(std::string_view str);
/// @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

3
src/tint/ast/storage_class.h.tmpl
View File

@ -28,7 +28,8 @@ namespace tint::ast {
/// @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

48
src/tint/ast/storage_class_bench.cc
View File

@ -31,12 +31,48 @@ namespace {
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) {

5
src/tint/ast/storage_class_test.cc
View File

@ -44,7 +44,7 @@ inline std::ostream& operator<<(std::ostream& out, Case c) {
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[] = {
@ -55,7 +55,8 @@ static constexpr Case kInvalidCases[] = {
{"wbkgGoup", StorageClass::kInvalid}, {"unifiivm", StorageClass::kInvalid},
{"8WWiform", StorageClass::kInvalid}, {"uxxform", StorageClass::kInvalid},
{"sXraggg", StorageClass::kInvalid}, {"traXe", StorageClass::kInvalid},
{"stor3ge", StorageClass::kInvalid},
{"stor3ge", StorageClass::kInvalid}, {"push_constanE", StorageClass::kInvalid},
{"push_TTPnstant", StorageClass::kInvalid}, {"puxxdh_constan", StorageClass::kInvalid},
};
using StorageClassParseTest = testing::TestWithParam<Case>;

3
src/tint/intrinsics.def
View File

@ -49,6 +49,8 @@ enum extension {
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 @@ enum storage_class {
workgroup
uniform
storage
push_constant
@internal handle
@internal in
@internal out

11
src/tint/reader/wgsl/parser_impl_variable_decl_test.cc
View File

@ -88,6 +88,17 @@ TEST_F(ParserImplTest, VariableDecl_WithStorageClass) {
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();

123
src/tint/resolver/entry_point_validation_test.cc
View File

@ -306,6 +306,129 @@ TEST_F(ResolverEntryPointValidationTest, VertexShaderMustReturnPosition) {
"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;

6
src/tint/resolver/resolver.cc
View File

@ -159,6 +159,10 @@ bool Resolver::ResolveInternal() {
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 @@ sem::GlobalVariable* Resolver::GlobalVariable(const ast::Variable* v) {
// 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;
}

42
src/tint/resolver/storage_class_layout_validation_test.cc
View File

@ -523,5 +523,47 @@ TEST_F(ResolverStorageClassLayoutValidationTest, UniformBuffer_InvalidArrayStrid
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

84
src/tint/resolver/storage_class_validation_test.cc
View File

@ -598,5 +598,89 @@ TEST_F(ResolverStorageClassValidationTest, UniformBufferStructI32Aliases) {
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

81
src/tint/resolver/validator.cc
View File

@ -395,7 +395,7 @@ bool Validator::StorageClassLayout(const sem::Type* store_ty,
// 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::Type* store_ty,
}
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 @@ bool Validator::PipelineStages(const std::vector<sem::Function*>& entry_points)
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();

10
src/tint/resolver/validator.h
View File

@ -116,6 +116,11 @@ class Validator {
/// @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 @@ class Validator {
/// 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

3
src/tint/transform/add_spirv_block_attribute.cc
View File

@ -58,7 +58,8 @@ void AddSpirvBlockAttribute::Run(CloneContext& ctx, const DataMap&, DataMap&) co
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;
}

65
src/tint/transform/add_spirv_block_attribute_test.cc
View File

@ -196,6 +196,71 @@ fn main() {
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 {

14
src/tint/transform/module_scope_var_to_entry_point_param.cc
View File

@ -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 @@ struct ModuleScopeVarToEntryPointParam::State {
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 @@ struct ModuleScopeVarToEntryPointParam::State {
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 << ")";

6
src/tint/writer/glsl/generator_impl.cc
View File

@ -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 @@ bool GeneratorImpl::EmitGlobalVariable(const ast::Variable* global) {
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();

7
src/tint/writer/hlsl/generator_impl.cc
View File

@ -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 @@ bool GeneratorImpl::EmitGlobalVariable(const ast::Variable* global) {
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 @@ bool GeneratorImpl::EmitGlobalVariable(const ast::Variable* global) {
[&](Default) {
TINT_ICE(Writer, diagnostics_)
<< "unhandled global variable type " << global->TypeInfo().name;
return false;
});
}

2
src/tint/writer/spirv/builder.cc
View File

@ -4121,6 +4121,8 @@ SpvStorageClass Builder::ConvertStorageClass(ast::StorageClass klass) const {
return SpvStorageClassUniform;
case ast::StorageClass::kWorkgroup:
return SpvStorageClassWorkgroup;
case ast::StorageClass::kPushConstant:
return SpvStorageClassPushConstant;
case ast::StorageClass::kHandle:
return SpvStorageClassUniformConstant;
case ast::StorageClass::kStorage:

36
test/tint/var/uses/push_constant.wgsl Normal file
View File

@ -0,0 +1,36 @@
enable chromium_experimental_push_constant;
var<push_constant> a : i32;
var<push_constant> b : i32;
var<push_constant> c : i32; // unused
fn uses_a() {
let foo = a;
}
fn uses_uses_a() {
uses_a();
}
fn uses_b() {
let foo = b;
}
@compute @workgroup_size(1)
fn main1() {
uses_a();
}
@compute @workgroup_size(1)
fn main2() {
uses_uses_a();
}
@compute @workgroup_size(1)
fn main3() {
uses_b();
}
@compute @workgroup_size(1)
fn main4() {
}

43
test/tint/var/uses/push_constant.wgsl.expected.dxc.hlsl Normal file
View File

@ -0,0 +1,43 @@
SKIP: FAILED
enable chromium_experimental_push_constant;
var<push_constant> a : i32;
var<push_constant> b : i32;
var<push_constant> c : i32;
fn uses_a() {
let foo = a;
}
fn uses_uses_a() {
uses_a();
}
fn uses_b() {
let foo = b;
}
@compute @workgroup_size(1)
fn main1() {
uses_a();
}
@compute @workgroup_size(1)
fn main2() {
uses_uses_a();
}
@compute @workgroup_size(1)
fn main3() {
uses_b();
}
@compute @workgroup_size(1)
fn main4() {
}
Failed to generate: error: unhandled storage class push_constant

43
test/tint/var/uses/push_constant.wgsl.expected.fxc.hlsl Normal file
View File

@ -0,0 +1,43 @@
SKIP: FAILED
enable chromium_experimental_push_constant;
var<push_constant> a : i32;
var<push_constant> b : i32;
var<push_constant> c : i32;
fn uses_a() {
let foo = a;
}
fn uses_uses_a() {
uses_a();
}
fn uses_b() {
let foo = b;
}
@compute @workgroup_size(1)
fn main1() {
uses_a();
}
@compute @workgroup_size(1)
fn main2() {
uses_uses_a();
}
@compute @workgroup_size(1)
fn main3() {
uses_b();
}
@compute @workgroup_size(1)
fn main4() {
}
Failed to generate: error: unhandled storage class push_constant

135
test/tint/var/uses/push_constant.wgsl.expected.glsl Normal file
View File

@ -0,0 +1,135 @@
SKIP: FAILED
enable chromium_experimental_push_constant;
var<push_constant> a : i32;
var<push_constant> b : i32;
var<push_constant> c : i32;
fn uses_a() {
let foo = a;
}
fn uses_uses_a() {
uses_a();
}
fn uses_b() {
let foo = b;
}
@compute @workgroup_size(1)
fn main1() {
uses_a();
}
@compute @workgroup_size(1)
fn main2() {
uses_uses_a();
}
@compute @workgroup_size(1)
fn main3() {
uses_b();
}
@compute @workgroup_size(1)
fn main4() {
}
Failed to generate: error: unhandled storage class push_constant
enable chromium_experimental_push_constant;
var<push_constant> a : i32;
var<push_constant> b : i32;
var<push_constant> c : i32;
fn uses_a() {
let foo = a;
}
fn uses_uses_a() {
uses_a();
}
fn uses_b() {
let foo = b;
}
@compute @workgroup_size(1)
fn main1() {
uses_a();
}
@compute @workgroup_size(1)
fn main2() {
uses_uses_a();
}
@compute @workgroup_size(1)
fn main3() {
uses_b();
}
@compute @workgroup_size(1)
fn main4() {
}
Failed to generate: error: unhandled storage class push_constant
enable chromium_experimental_push_constant;
var<push_constant> a : i32;
var<push_constant> b : i32;
var<push_constant> c : i32;
fn uses_a() {
let foo = a;
}
fn uses_uses_a() {
uses_a();
}
fn uses_b() {
let foo = b;
}
@compute @workgroup_size(1)
fn main1() {
uses_a();
}
@compute @workgroup_size(1)
fn main2() {
uses_uses_a();
}
@compute @workgroup_size(1)
fn main3() {
uses_b();
}
@compute @workgroup_size(1)
fn main4() {
}
Failed to generate: error: unhandled storage class push_constant
#version 310 es
void main4() {
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
main4();
return;
}

48
test/tint/var/uses/push_constant.wgsl.expected.msl Normal file
View File

@ -0,0 +1,48 @@
SKIP: FAILED
enable chromium_experimental_push_constant;
var<push_constant> a : i32;
var<push_constant> b : i32;
var<push_constant> c : i32;
fn uses_a() {
let foo = a;
}
fn uses_uses_a() {
uses_a();
}
fn uses_b() {
let foo = b;
}
@compute @workgroup_size(1)
fn main1() {
uses_a();
}
@compute @workgroup_size(1)
fn main2() {
uses_uses_a();
}
@compute @workgroup_size(1)
fn main3() {
uses_b();
}
@compute @workgroup_size(1)
fn main4() {
}
Failed to generate: error: unhandled module-scope storage class (push_constant)
error: unhandled module-scope storage class (push_constant)
error: unhandled module-scope storage class (push_constant)
error: unhandled module-scope storage class (push_constant)
error: unhandled module-scope storage class (push_constant)
error: unhandled module-scope storage class (push_constant)

76
test/tint/var/uses/push_constant.wgsl.expected.spvasm Normal file
View File

@ -0,0 +1,76 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 34
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main1 "main1"
OpEntryPoint GLCompute %main2 "main2"
OpEntryPoint GLCompute %main3 "main3"
OpEntryPoint GLCompute %main4 "main4"
OpExecutionMode %main1 LocalSize 1 1 1
OpExecutionMode %main2 LocalSize 1 1 1
OpExecutionMode %main3 LocalSize 1 1 1
OpExecutionMode %main4 LocalSize 1 1 1
OpName %a_block "a_block"
OpMemberName %a_block 0 "inner"
OpName %a "a"
OpName %b "b"
OpName %c "c"
OpName %uses_a "uses_a"
OpName %uses_uses_a "uses_uses_a"
OpName %uses_b "uses_b"
OpName %main1 "main1"
OpName %main2 "main2"
OpName %main3 "main3"
OpName %main4 "main4"
OpDecorate %a_block Block
OpMemberDecorate %a_block 0 Offset 0
%int = OpTypeInt 32 1
%a_block = OpTypeStruct %int
%_ptr_PushConstant_a_block = OpTypePointer PushConstant %a_block
%a = OpVariable %_ptr_PushConstant_a_block PushConstant
%b = OpVariable %_ptr_PushConstant_a_block PushConstant
%c = OpVariable %_ptr_PushConstant_a_block PushConstant
%void = OpTypeVoid
%7 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%uint_0 = OpConstant %uint 0
%_ptr_PushConstant_int = OpTypePointer PushConstant %int
%uses_a = OpFunction %void None %7
%10 = OpLabel
%14 = OpAccessChain %_ptr_PushConstant_int %a %uint_0
%15 = OpLoad %int %14
OpReturn
OpFunctionEnd
%uses_uses_a = OpFunction %void None %7
%17 = OpLabel
%18 = OpFunctionCall %void %uses_a
OpReturn
OpFunctionEnd
%uses_b = OpFunction %void None %7
%20 = OpLabel
%21 = OpAccessChain %_ptr_PushConstant_int %b %uint_0
%22 = OpLoad %int %21
OpReturn
OpFunctionEnd
%main1 = OpFunction %void None %7
%24 = OpLabel
%25 = OpFunctionCall %void %uses_a
OpReturn
OpFunctionEnd
%main2 = OpFunction %void None %7
%27 = OpLabel
%28 = OpFunctionCall %void %uses_uses_a
OpReturn
OpFunctionEnd
%main3 = OpFunction %void None %7
%30 = OpLabel
%31 = OpFunctionCall %void %uses_b
OpReturn
OpFunctionEnd
%main4 = OpFunction %void None %7
%33 = OpLabel
OpReturn
OpFunctionEnd

38
test/tint/var/uses/push_constant.wgsl.expected.wgsl Normal file
View File

@ -0,0 +1,38 @@
enable chromium_experimental_push_constant;
var<push_constant> a : i32;
var<push_constant> b : i32;
var<push_constant> c : i32;
fn uses_a() {
let foo = a;
}
fn uses_uses_a() {
uses_a();
}
fn uses_b() {
let foo = b;
}
@compute @workgroup_size(1)
fn main1() {
uses_a();
}
@compute @workgroup_size(1)
fn main2() {
uses_uses_a();
}
@compute @workgroup_size(1)
fn main3() {
uses_b();
}
@compute @workgroup_size(1)
fn main4() {
}