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>
diff --git a/src/dawn_native/Pipeline.cpp b/src/dawn_native/Pipeline.cpp
index 96258bc..ab232c3 100644
--- a/src/dawn_native/Pipeline.cpp
+++ b/src/dawn_native/Pipeline.cpp
@@ -53,10 +53,43 @@
// Validate if overridable constants exist in shader module
// 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++) {
DAWN_INVALID_IF(metadata.overridableConstants.count(constants[i].key) == 0,
"Pipeline overridable constant \"%s\" not found in shader module %s.",
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 {};
diff --git a/src/dawn_native/ShaderModule.cpp b/src/dawn_native/ShaderModule.cpp
index 7a7bbd7..058f886e 100644
--- a/src/dawn_native/ShaderModule.cpp
+++ b/src/dawn_native/ShaderModule.cpp
@@ -639,12 +639,19 @@
for (auto& c : entryPoint.overridable_constants) {
EntryPointMetadata::OverridableConstant constant = {
- name2Id.at(c.name), FromTintOverridableConstantType(c.type)};
- metadata->overridableConstants[c.name] = constant;
- // TODO(tint:1155) tint needs ways to differentiate whether a pipeline
- // constant id is specified explicitly. Now we just store numeric id and
- // variable name in the index at the same time
- metadata->overridableConstants[std::to_string(constant.id)] = constant;
+ name2Id.at(c.name), FromTintOverridableConstantType(c.type),
+ c.is_initialized};
+
+ std::string identifier =
+ c.is_numeric_id_specified ? std::to_string(constant.id) : c.name;
+ 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);
+ }
}
}
diff --git a/src/dawn_native/ShaderModule.h b/src/dawn_native/ShaderModule.h
index 070d001..09bede5 100644
--- a/src/dawn_native/ShaderModule.h
+++ b/src/dawn_native/ShaderModule.h
@@ -32,6 +32,7 @@
#include <bitset>
#include <map>
#include <unordered_map>
+#include <unordered_set>
#include <vector>
namespace tint {
@@ -200,11 +201,21 @@
// Match tint::inspector::OverridableConstant::Type
// Bool is defined as a macro on linux X11 and cannot compile
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;
+ // 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;
};
diff --git a/src/dawn_native/vulkan/UtilsVulkan.cpp b/src/dawn_native/vulkan/UtilsVulkan.cpp
index 5d90290..87d896d 100644
--- a/src/dawn_native/vulkan/UtilsVulkan.cpp
+++ b/src/dawn_native/vulkan/UtilsVulkan.cpp
@@ -215,11 +215,11 @@
programmableStage.module->GetEntryPoint(programmableStage.entryPoint);
for (const auto& pipelineConstant : programmableStage.constants) {
- const std::string& name = pipelineConstant.first;
+ const std::string& identifier = pipelineConstant.first;
double value = pipelineConstant.second;
- // This is already validated so `name` must exist
- const auto& moduleConstant = entryPointMetaData.overridableConstants.at(name);
+ // This is already validated so `identifier` must exist
+ const auto& moduleConstant = entryPointMetaData.overridableConstants.at(identifier);
specializationMapEntries->push_back(
VkSpecializationMapEntry{moduleConstant.id,
@@ -230,7 +230,7 @@
SpecializationDataEntry entry{};
switch (moduleConstant.type) {
case EntryPointMetadata::OverridableConstant::Type::Boolean:
- entry.b = static_cast<bool>(value);
+ entry.b = static_cast<int32_t>(value);
break;
case EntryPointMetadata::OverridableConstant::Type::Float32:
entry.f32 = static_cast<float>(value);
diff --git a/src/dawn_native/vulkan/UtilsVulkan.h b/src/dawn_native/vulkan/UtilsVulkan.h
index 8bd2a84..c385917 100644
--- a/src/dawn_native/vulkan/UtilsVulkan.h
+++ b/src/dawn_native/vulkan/UtilsVulkan.h
@@ -114,7 +114,8 @@
// Helpers for creating VkSpecializationInfo
// The WebGPU overridable constants only support these scalar types
union SpecializationDataEntry {
- bool b;
+ // Use int32_t for boolean to initialize the full 32bit
+ int32_t b;
float f32;
int32_t i32;
uint32_t u32;
diff --git a/src/tests/end2end/ShaderTests.cpp b/src/tests/end2end/ShaderTests.cpp
index 7c32ad6..aa40fe8 100644
--- a/src/tests/end2end/ShaderTests.cpp
+++ b/src/tests/end2end/ShaderTests.cpp
@@ -393,19 +393,13 @@
}
// Test overridable constants without numeric identifiers
-// TODO(tint:1155): Implicit numeric ID is undetermined in tint
-TEST_P(ShaderTests, DISABLED_OverridableConstants) {
+TEST_P(ShaderTests, OverridableConstants) {
// TODO(dawn:1041): Only Vulkan backend is implemented
DAWN_TEST_UNSUPPORTED_IF(!IsVulkan());
- uint32_t const kCount = 15;
+ uint32_t const kCount = 11;
std::vector<uint32_t> expected(kCount);
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);
std::string shader = R"(
@@ -420,13 +414,9 @@
[[override]] let c8: u32; // type: uint32
[[override]] let c9: u32 = 0u; // default override
[[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 {
- data : array<u32, 15>;
+ data : array<u32, 11>;
};
[[group(0), binding(0)]] var<storage, read_write> buf : Buf;
@@ -443,10 +433,6 @@
buf.data[8] = u32(c8);
buf.data[9] = u32(c9);
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;
@@ -461,10 +447,6 @@
constants.push_back({nullptr, "c8", 8});
constants.push_back({nullptr, "c9", 9});
// 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);
@@ -520,7 +502,7 @@
constants.push_back({nullptr, "1001", 1});
constants.push_back({nullptr, "1", 2});
// 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);
diff --git a/src/tests/unittests/validation/OverridableConstantsValidationTests.cpp b/src/tests/unittests/validation/OverridableConstantsValidationTests.cpp
index af96c39..af43f4c 100644
--- a/src/tests/unittests/validation/OverridableConstantsValidationTests.cpp
+++ b/src/tests/unittests/validation/OverridableConstantsValidationTests.cpp
@@ -18,9 +18,37 @@
class ComputePipelineOverridableConstantsValidationTest : public ValidationTest {
protected:
- void SetUp() override {
- ValidationTest::SetUp();
+ void SetUpShadersWithDefaultValueConstants() {
+ 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"(
[[override]] let c0: bool; // type: bool
[[override]] let c1: bool = false; // default override
@@ -34,25 +62,19 @@
[[override]] let c9: u32 = 0u; // default override
[[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() {
// make sure the overridable constants are not optimized out
- buf.data[0] = u32(c0);
- buf.data[1] = u32(c1);
- buf.data[2] = u32(c2);
- buf.data[3] = u32(c3);
- buf.data[4] = u32(c4);
- buf.data[5] = u32(c5);
- buf.data[6] = u32(c6);
- buf.data[7] = u32(c7);
- buf.data[8] = u32(c8);
- buf.data[9] = u32(c9);
- buf.data[10] = u32(c10);
+ _ = u32(c0);
+ _ = u32(c1);
+ _ = u32(c2);
+ _ = u32(c3);
+ _ = u32(c4);
+ _ = u32(c5);
+ _ = u32(c6);
+ _ = u32(c7);
+ _ = u32(c8);
+ _ = u32(c9);
+ _ = u32(c10);
})");
}
@@ -71,6 +93,7 @@
// Basic constants lookup tests
TEST_F(ComputePipelineOverridableConstantsValidationTest, ConstantsIdentifierLookUp) {
+ SetUpShadersWithDefaultValueConstants();
{
// Valid: no constants specified
std::vector<wgpu::ConstantEntry> constants;
@@ -106,10 +129,47 @@
}
}
+// 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
-// TODO(tint:1155): missing feature in tint to differentiate explicitly specified numeric ID
-TEST_F(ComputePipelineOverridableConstantsValidationTest,
- DISABLED_ConstantsIdentifierExplicitNumericID) {
+TEST_F(ComputePipelineOverridableConstantsValidationTest, ConstantsIdentifierExplicitNumericID) {
+ SetUpShadersWithDefaultValueConstants();
{
// Error: constant numeric id not explicitly specified
// But could be impliciltly assigned to one of the constants
@@ -137,8 +197,13 @@
}
// Test that identifiers are unique
-// TODO(tint:1155): missing feature in tint to differentiate explicitly specified numeric ID
-TEST_F(ComputePipelineOverridableConstantsValidationTest, DISABLED_ConstantsIdentifierUnique) {
+TEST_F(ComputePipelineOverridableConstantsValidationTest, 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
std::vector<wgpu::ConstantEntry> constants{{nullptr, "c10", 0}};
diff --git a/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp b/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp
index a7e915a..624305b 100644
--- a/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp
+++ b/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp
@@ -129,8 +129,8 @@
[[override(1000)]] let c1: u32;
[[stage(compute), workgroup_size(1)]] fn main() {
- ignore(c0);
- ignore(c1);
+ _ = c0;
+ _ = c1;
})"));
}
