Pipeline constants: validate if any uninitialized overridable constants exist

Now an overridable constant need to be initialized, either via value specified
in shader, or via constant entry from pipeline stage. Otherwise it is invalid.

Together fix the bool 32bit initialize problem on vulkan and re-enable the
overridable constants shader test, and use assigning to _ instead of
ignore() in tests.

Bug: dawn:1041
Change-Id: I49e7885c8d6134647b09926ceb15234ac21ee35d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/67560
Commit-Queue: Shrek Shao <shrekshao@google.com>
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
This commit is contained in:
Shrek Shao 2021-10-28 00:04:07 +00:00 committed by Dawn LUCI CQ
parent 729cc97f87
commit 5e2e2d863e
8 changed files with 159 additions and 60 deletions

View File

@ -53,10 +53,43 @@ namespace dawn_native {
// Validate if overridable constants exist in shader module // Validate if overridable constants exist in shader module
// pipelineBase is not yet constructed at this moment so iterate constants from descriptor // pipelineBase is not yet constructed at this moment so iterate constants from descriptor
size_t numUninitializedConstants = metadata.uninitializedOverridableConstants.size();
// Keep a initialized constants sets to handle duplicate initialization cases
// Only storing that of uninialized constants is needed
std::unordered_set<std::string> stageInitializedConstantIdentifiers;
for (uint32_t i = 0; i < constantCount; i++) { for (uint32_t i = 0; i < constantCount; i++) {
DAWN_INVALID_IF(metadata.overridableConstants.count(constants[i].key) == 0, DAWN_INVALID_IF(metadata.overridableConstants.count(constants[i].key) == 0,
"Pipeline overridable constant \"%s\" not found in shader module %s.", "Pipeline overridable constant \"%s\" not found in shader module %s.",
constants[i].key, module); constants[i].key, module);
if (metadata.uninitializedOverridableConstants.count(constants[i].key) > 0 &&
stageInitializedConstantIdentifiers.count(constants[i].key) == 0) {
numUninitializedConstants--;
stageInitializedConstantIdentifiers.insert(constants[i].key);
}
}
// Validate if any overridable constant is left uninitialized
if (DAWN_UNLIKELY(numUninitializedConstants > 0)) {
std::string uninitializedConstantsArray;
bool isFirst = true;
for (std::string identifier : metadata.uninitializedOverridableConstants) {
if (stageInitializedConstantIdentifiers.count(identifier) > 0) {
continue;
}
if (isFirst) {
isFirst = false;
} else {
uninitializedConstantsArray.append(", ");
}
uninitializedConstantsArray.append(identifier);
}
return DAWN_FORMAT_VALIDATION_ERROR(
"There are uninitialized pipeline overridable constants in shader module %s, their "
"identifiers:[%s]",
module, uninitializedConstantsArray);
} }
return {}; return {};

View File

@ -639,12 +639,19 @@ namespace dawn_native {
for (auto& c : entryPoint.overridable_constants) { for (auto& c : entryPoint.overridable_constants) {
EntryPointMetadata::OverridableConstant constant = { EntryPointMetadata::OverridableConstant constant = {
name2Id.at(c.name), FromTintOverridableConstantType(c.type)}; name2Id.at(c.name), FromTintOverridableConstantType(c.type),
metadata->overridableConstants[c.name] = constant; c.is_initialized};
// TODO(tint:1155) tint needs ways to differentiate whether a pipeline
// constant id is specified explicitly. Now we just store numeric id and std::string identifier =
// variable name in the index at the same time c.is_numeric_id_specified ? std::to_string(constant.id) : c.name;
metadata->overridableConstants[std::to_string(constant.id)] = constant; metadata->overridableConstants[identifier] = constant;
if (!c.is_initialized) {
auto it = metadata->uninitializedOverridableConstants.emplace(
std::move(identifier));
// The insertion should have taken place
ASSERT(it.second);
}
} }
} }

View File

@ -32,6 +32,7 @@
#include <bitset> #include <bitset>
#include <map> #include <map>
#include <unordered_map> #include <unordered_map>
#include <unordered_set>
#include <vector> #include <vector>
namespace tint { namespace tint {
@ -200,11 +201,21 @@ namespace dawn_native {
// Match tint::inspector::OverridableConstant::Type // Match tint::inspector::OverridableConstant::Type
// Bool is defined as a macro on linux X11 and cannot compile // Bool is defined as a macro on linux X11 and cannot compile
enum class Type { Boolean, Float32, Uint32, Int32 } type; enum class Type { Boolean, Float32, Uint32, Int32 } type;
// If the constant doesn't not have an initializer in the shader
// Then it is required for the pipeline stage to have a constant record to initialize a
// value
bool isInitialized;
}; };
// Store overridableConstants from tint program // Map identifier to overridable constant
// Identifier is unique: either the variable name or the numeric ID if specified
std::unordered_map<std::string, OverridableConstant> overridableConstants; std::unordered_map<std::string, OverridableConstant> overridableConstants;
// Overridable constants that are not initialized in shaders
// They need value initialization from pipeline stage or it is a validation error
std::unordered_set<std::string> uninitializedOverridableConstants;
bool usesNumWorkgroups = false; bool usesNumWorkgroups = false;
}; };

View File

@ -215,11 +215,11 @@ namespace dawn_native { namespace vulkan {
programmableStage.module->GetEntryPoint(programmableStage.entryPoint); programmableStage.module->GetEntryPoint(programmableStage.entryPoint);
for (const auto& pipelineConstant : programmableStage.constants) { for (const auto& pipelineConstant : programmableStage.constants) {
const std::string& name = pipelineConstant.first; const std::string& identifier = pipelineConstant.first;
double value = pipelineConstant.second; double value = pipelineConstant.second;
// This is already validated so `name` must exist // This is already validated so `identifier` must exist
const auto& moduleConstant = entryPointMetaData.overridableConstants.at(name); const auto& moduleConstant = entryPointMetaData.overridableConstants.at(identifier);
specializationMapEntries->push_back( specializationMapEntries->push_back(
VkSpecializationMapEntry{moduleConstant.id, VkSpecializationMapEntry{moduleConstant.id,
@ -230,7 +230,7 @@ namespace dawn_native { namespace vulkan {
SpecializationDataEntry entry{}; SpecializationDataEntry entry{};
switch (moduleConstant.type) { switch (moduleConstant.type) {
case EntryPointMetadata::OverridableConstant::Type::Boolean: case EntryPointMetadata::OverridableConstant::Type::Boolean:
entry.b = static_cast<bool>(value); entry.b = static_cast<int32_t>(value);
break; break;
case EntryPointMetadata::OverridableConstant::Type::Float32: case EntryPointMetadata::OverridableConstant::Type::Float32:
entry.f32 = static_cast<float>(value); entry.f32 = static_cast<float>(value);

View File

@ -114,7 +114,8 @@ namespace dawn_native { namespace vulkan {
// Helpers for creating VkSpecializationInfo // Helpers for creating VkSpecializationInfo
// The WebGPU overridable constants only support these scalar types // The WebGPU overridable constants only support these scalar types
union SpecializationDataEntry { union SpecializationDataEntry {
bool b; // Use int32_t for boolean to initialize the full 32bit
int32_t b;
float f32; float f32;
int32_t i32; int32_t i32;
uint32_t u32; uint32_t u32;

View File

@ -393,19 +393,13 @@ fn main([[location(0)]] pos : vec4<f32>) -> [[builtin(position)]] vec4<f32> {
} }
// Test overridable constants without numeric identifiers // Test overridable constants without numeric identifiers
// TODO(tint:1155): Implicit numeric ID is undetermined in tint TEST_P(ShaderTests, OverridableConstants) {
TEST_P(ShaderTests, DISABLED_OverridableConstants) {
// TODO(dawn:1041): Only Vulkan backend is implemented // TODO(dawn:1041): Only Vulkan backend is implemented
DAWN_TEST_UNSUPPORTED_IF(!IsVulkan()); DAWN_TEST_UNSUPPORTED_IF(!IsVulkan());
uint32_t const kCount = 15; uint32_t const kCount = 11;
std::vector<uint32_t> expected(kCount); std::vector<uint32_t> expected(kCount);
std::iota(expected.begin(), expected.end(), 0); std::iota(expected.begin(), expected.end(), 0);
// Test last entry with unspecified default value
expected[kCount - 1] = 0u;
expected[kCount - 2] = 0u;
expected[kCount - 3] = 0u;
expected[kCount - 4] = 0u;
wgpu::Buffer buffer = CreateBuffer(kCount); wgpu::Buffer buffer = CreateBuffer(kCount);
std::string shader = R"( std::string shader = R"(
@ -420,13 +414,9 @@ TEST_P(ShaderTests, DISABLED_OverridableConstants) {
[[override]] let c8: u32; // type: uint32 [[override]] let c8: u32; // type: uint32
[[override]] let c9: u32 = 0u; // default override [[override]] let c9: u32 = 0u; // default override
[[override]] let c10: u32 = 10u; // default [[override]] let c10: u32 = 10u; // default
[[override]] let c11: bool; // default unspecified
[[override]] let c12: f32; // default unspecified
[[override]] let c13: i32; // default unspecified
[[override]] let c14: u32; // default unspecified
[[block]] struct Buf { [[block]] struct Buf {
data : array<u32, 15>; data : array<u32, 11>;
}; };
[[group(0), binding(0)]] var<storage, read_write> buf : Buf; [[group(0), binding(0)]] var<storage, read_write> buf : Buf;
@ -443,10 +433,6 @@ TEST_P(ShaderTests, DISABLED_OverridableConstants) {
buf.data[8] = u32(c8); buf.data[8] = u32(c8);
buf.data[9] = u32(c9); buf.data[9] = u32(c9);
buf.data[10] = u32(c10); buf.data[10] = u32(c10);
buf.data[11] = u32(c11);
buf.data[12] = u32(c12);
buf.data[13] = u32(c13);
buf.data[14] = u32(c14);
})"; })";
std::vector<wgpu::ConstantEntry> constants; std::vector<wgpu::ConstantEntry> constants;
@ -461,10 +447,6 @@ TEST_P(ShaderTests, DISABLED_OverridableConstants) {
constants.push_back({nullptr, "c8", 8}); constants.push_back({nullptr, "c8", 8});
constants.push_back({nullptr, "c9", 9}); constants.push_back({nullptr, "c9", 9});
// c10 is not assigned, testing default value // c10 is not assigned, testing default value
// c11 is not assigned, testing unspecified default value
// c12 is not assigned, testing unspecified default value
// c13 is not assigned, testing unspecified default value
// c14 is not assigned, testing unspecified default value
wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main", &constants); wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main", &constants);
@ -520,7 +502,7 @@ TEST_P(ShaderTests, OverridableConstantsNumericIdentifiers) {
constants.push_back({nullptr, "1001", 1}); constants.push_back({nullptr, "1001", 1});
constants.push_back({nullptr, "1", 2}); constants.push_back({nullptr, "1", 2});
// c3 is not assigned, testing default value // c3 is not assigned, testing default value
// c4 is not assigned, testing unspecified default value constants.push_back({nullptr, "1004", 0});
wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main", &constants); wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main", &constants);

View File

@ -18,9 +18,37 @@
class ComputePipelineOverridableConstantsValidationTest : public ValidationTest { class ComputePipelineOverridableConstantsValidationTest : public ValidationTest {
protected: protected:
void SetUp() override { void SetUpShadersWithDefaultValueConstants() {
ValidationTest::SetUp(); computeModule = utils::CreateShaderModule(device, R"(
[[override]] let c0: bool = true; // type: bool
[[override]] let c1: bool = false; // default override
[[override]] let c2: f32 = 0.0; // type: float32
[[override]] let c3: f32 = 0.0; // default override
[[override]] let c4: f32 = 4.0; // default
[[override]] let c5: i32 = 0; // type: int32
[[override]] let c6: i32 = 0; // default override
[[override]] let c7: i32 = 7; // default
[[override]] let c8: u32 = 0u; // type: uint32
[[override]] let c9: u32 = 0u; // default override
[[override(1000)]] let c10: u32 = 10u; // default
[[stage(compute), workgroup_size(1)]] fn main() {
// make sure the overridable constants are not optimized out
_ = u32(c0);
_ = u32(c1);
_ = u32(c2);
_ = u32(c3);
_ = u32(c4);
_ = u32(c5);
_ = u32(c6);
_ = u32(c7);
_ = u32(c8);
_ = u32(c9);
_ = u32(c10);
})");
}
void SetUpShadersWithUninitializedConstants() {
computeModule = utils::CreateShaderModule(device, R"( computeModule = utils::CreateShaderModule(device, R"(
[[override]] let c0: bool; // type: bool [[override]] let c0: bool; // type: bool
[[override]] let c1: bool = false; // default override [[override]] let c1: bool = false; // default override
@ -34,25 +62,19 @@ class ComputePipelineOverridableConstantsValidationTest : public ValidationTest
[[override]] let c9: u32 = 0u; // default override [[override]] let c9: u32 = 0u; // default override
[[override(1000)]] let c10: u32 = 10u; // default [[override(1000)]] let c10: u32 = 10u; // default
[[block]] struct Buf {
data : array<u32, 11>;
};
[[group(0), binding(0)]] var<storage, read_write> buf : Buf;
[[stage(compute), workgroup_size(1)]] fn main() { [[stage(compute), workgroup_size(1)]] fn main() {
// make sure the overridable constants are not optimized out // make sure the overridable constants are not optimized out
buf.data[0] = u32(c0); _ = u32(c0);
buf.data[1] = u32(c1); _ = u32(c1);
buf.data[2] = u32(c2); _ = u32(c2);
buf.data[3] = u32(c3); _ = u32(c3);
buf.data[4] = u32(c4); _ = u32(c4);
buf.data[5] = u32(c5); _ = u32(c5);
buf.data[6] = u32(c6); _ = u32(c6);
buf.data[7] = u32(c7); _ = u32(c7);
buf.data[8] = u32(c8); _ = u32(c8);
buf.data[9] = u32(c9); _ = u32(c9);
buf.data[10] = u32(c10); _ = u32(c10);
})"); })");
} }
@ -71,6 +93,7 @@ class ComputePipelineOverridableConstantsValidationTest : public ValidationTest
// Basic constants lookup tests // Basic constants lookup tests
TEST_F(ComputePipelineOverridableConstantsValidationTest, ConstantsIdentifierLookUp) { TEST_F(ComputePipelineOverridableConstantsValidationTest, ConstantsIdentifierLookUp) {
SetUpShadersWithDefaultValueConstants();
{ {
// Valid: no constants specified // Valid: no constants specified
std::vector<wgpu::ConstantEntry> constants; std::vector<wgpu::ConstantEntry> constants;
@ -106,10 +129,47 @@ TEST_F(ComputePipelineOverridableConstantsValidationTest, ConstantsIdentifierLoo
} }
} }
// Test that it is invalid to leave any constants uninitialized
TEST_F(ComputePipelineOverridableConstantsValidationTest, UninitializedConstants) {
SetUpShadersWithUninitializedConstants();
{
// Error: uninitialized constants exist
std::vector<wgpu::ConstantEntry> constants;
ASSERT_DEVICE_ERROR(TestCreatePipeline(constants));
}
{
// Error: uninitialized constants exist
std::vector<wgpu::ConstantEntry> constants{
{nullptr, "c0", false},
{nullptr, "c2", 1},
// c5 is missing
{nullptr, "c8", 1},
};
ASSERT_DEVICE_ERROR(TestCreatePipeline(constants));
}
{
// Valid: all constants initialized
std::vector<wgpu::ConstantEntry> constants{
{nullptr, "c0", false},
{nullptr, "c2", 1},
{nullptr, "c5", 1},
{nullptr, "c8", 1},
};
TestCreatePipeline(constants);
}
{
// Valid: all constants initialized (with duplicate initializations)
std::vector<wgpu::ConstantEntry> constants{
{nullptr, "c0", false}, {nullptr, "c2", 1}, {nullptr, "c5", 1},
{nullptr, "c8", 1}, {nullptr, "c2", 2},
};
TestCreatePipeline(constants);
}
}
// Test that only explicitly specified numeric ID can be referenced // Test that only explicitly specified numeric ID can be referenced
// TODO(tint:1155): missing feature in tint to differentiate explicitly specified numeric ID TEST_F(ComputePipelineOverridableConstantsValidationTest, ConstantsIdentifierExplicitNumericID) {
TEST_F(ComputePipelineOverridableConstantsValidationTest, SetUpShadersWithDefaultValueConstants();
DISABLED_ConstantsIdentifierExplicitNumericID) {
{ {
// Error: constant numeric id not explicitly specified // Error: constant numeric id not explicitly specified
// But could be impliciltly assigned to one of the constants // But could be impliciltly assigned to one of the constants
@ -137,8 +197,13 @@ TEST_F(ComputePipelineOverridableConstantsValidationTest,
} }
// Test that identifiers are unique // Test that identifiers are unique
// TODO(tint:1155): missing feature in tint to differentiate explicitly specified numeric ID TEST_F(ComputePipelineOverridableConstantsValidationTest, ConstantsIdentifierUnique) {
TEST_F(ComputePipelineOverridableConstantsValidationTest, DISABLED_ConstantsIdentifierUnique) { SetUpShadersWithDefaultValueConstants();
{
// Valid: constant without numeric id can be referenced with variable name
std::vector<wgpu::ConstantEntry> constants{{nullptr, "c0", 0}};
TestCreatePipeline(constants);
}
{ {
// Error: constant with numeric id cannot be referenced with variable name // Error: constant with numeric id cannot be referenced with variable name
std::vector<wgpu::ConstantEntry> constants{{nullptr, "c10", 0}}; std::vector<wgpu::ConstantEntry> constants{{nullptr, "c10", 0}};

View File

@ -129,8 +129,8 @@ TEST_F(UnsafeAPIValidationTest, PipelineOverridableConstants) {
[[override(1000)]] let c1: u32; [[override(1000)]] let c1: u32;
[[stage(compute), workgroup_size(1)]] fn main() { [[stage(compute), workgroup_size(1)]] fn main() {
ignore(c0); _ = c0;
ignore(c1); _ = c1;
})")); })"));
} }