diff --git a/src/dawn_native/Pipeline.cpp b/src/dawn_native/Pipeline.cpp index 96258bc63c..ab232c3566 100644 --- a/src/dawn_native/Pipeline.cpp +++ b/src/dawn_native/Pipeline.cpp @@ -53,10 +53,43 @@ namespace dawn_native { // 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 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 7a7bbd7133..058f886e48 100644 --- a/src/dawn_native/ShaderModule.cpp +++ b/src/dawn_native/ShaderModule.cpp @@ -639,12 +639,19 @@ namespace dawn_native { 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 070d001650..09bede54cd 100644 --- a/src/dawn_native/ShaderModule.h +++ b/src/dawn_native/ShaderModule.h @@ -32,6 +32,7 @@ #include #include #include +#include #include namespace tint { @@ -200,11 +201,21 @@ namespace dawn_native { // 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 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 uninitializedOverridableConstants; + bool usesNumWorkgroups = false; }; diff --git a/src/dawn_native/vulkan/UtilsVulkan.cpp b/src/dawn_native/vulkan/UtilsVulkan.cpp index 5d9029073f..87d896d4aa 100644 --- a/src/dawn_native/vulkan/UtilsVulkan.cpp +++ b/src/dawn_native/vulkan/UtilsVulkan.cpp @@ -215,11 +215,11 @@ namespace dawn_native { namespace vulkan { 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 @@ namespace dawn_native { namespace vulkan { SpecializationDataEntry entry{}; switch (moduleConstant.type) { case EntryPointMetadata::OverridableConstant::Type::Boolean: - entry.b = static_cast(value); + entry.b = static_cast(value); break; case EntryPointMetadata::OverridableConstant::Type::Float32: entry.f32 = static_cast(value); diff --git a/src/dawn_native/vulkan/UtilsVulkan.h b/src/dawn_native/vulkan/UtilsVulkan.h index 8bd2a84544..c3859172d9 100644 --- a/src/dawn_native/vulkan/UtilsVulkan.h +++ b/src/dawn_native/vulkan/UtilsVulkan.h @@ -114,7 +114,8 @@ namespace dawn_native { namespace vulkan { // 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 7c32ad6d4b..aa40fe85f2 100644 --- a/src/tests/end2end/ShaderTests.cpp +++ b/src/tests/end2end/ShaderTests.cpp @@ -393,19 +393,13 @@ fn main([[location(0)]] pos : vec4) -> [[builtin(position)]] vec4 { } // 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 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 @@ TEST_P(ShaderTests, DISABLED_OverridableConstants) { [[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; + data : array; }; [[group(0), binding(0)]] var buf : Buf; @@ -443,10 +433,6 @@ TEST_P(ShaderTests, DISABLED_OverridableConstants) { 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 constants; @@ -461,10 +447,6 @@ TEST_P(ShaderTests, DISABLED_OverridableConstants) { 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 @@ TEST_P(ShaderTests, OverridableConstantsNumericIdentifiers) { 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 af96c394e9..af43f4cf86 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 @@ class ComputePipelineOverridableConstantsValidationTest : public ValidationTest [[override]] let c9: u32 = 0u; // default override [[override(1000)]] let c10: u32 = 10u; // default -[[block]] struct Buf { - data : array; -}; - -[[group(0), binding(0)]] var 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 @@ class ComputePipelineOverridableConstantsValidationTest : public ValidationTest // Basic constants lookup tests TEST_F(ComputePipelineOverridableConstantsValidationTest, ConstantsIdentifierLookUp) { + SetUpShadersWithDefaultValueConstants(); { // Valid: no constants specified std::vector 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 constants; + ASSERT_DEVICE_ERROR(TestCreatePipeline(constants)); + } + { + // Error: uninitialized constants exist + std::vector constants{ + {nullptr, "c0", false}, + {nullptr, "c2", 1}, + // c5 is missing + {nullptr, "c8", 1}, + }; + ASSERT_DEVICE_ERROR(TestCreatePipeline(constants)); + } + { + // Valid: all constants initialized + std::vector constants{ + {nullptr, "c0", false}, + {nullptr, "c2", 1}, + {nullptr, "c5", 1}, + {nullptr, "c8", 1}, + }; + TestCreatePipeline(constants); + } + { + // Valid: all constants initialized (with duplicate initializations) + std::vector 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_F(ComputePipelineOverridableConstantsValidationTest, } // 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 constants{{nullptr, "c0", 0}}; + TestCreatePipeline(constants); + } { // Error: constant with numeric id cannot be referenced with variable name std::vector constants{{nullptr, "c10", 0}}; diff --git a/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp b/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp index a7e915a3bd..624305b015 100644 --- a/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp +++ b/src/tests/unittests/validation/UnsafeAPIValidationTests.cpp @@ -129,8 +129,8 @@ TEST_F(UnsafeAPIValidationTest, PipelineOverridableConstants) { [[override(1000)]] let c1: u32; [[stage(compute), workgroup_size(1)]] fn main() { - ignore(c0); - ignore(c1); + _ = c0; + _ = c1; })")); }