diff --git a/dawn.json b/dawn.json index 569dde198d..1e92c03ed2 100644 --- a/dawn.json +++ b/dawn.json @@ -438,7 +438,7 @@ "category": "structure", "extensible": "in", "members": [ - {"name": "key", "type": "char", "annotation": "const*"}, + {"name": "key", "type": "char", "annotation": "const*", "length": "strlen"}, {"name": "value", "type": "double"} ] }, @@ -1348,8 +1348,8 @@ "members": [ {"name": "module", "type": "shader module"}, {"name": "entry point", "type": "char", "annotation": "const*", "length": "strlen"}, - {"name": "constant count", "type": "uint32_t", "tags": ["upstream"]}, - {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count", "tags": ["upstream"]} + {"name": "constant count", "type": "uint32_t", "default": 0}, + {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count"} ] }, "primitive topology": { @@ -1817,8 +1817,8 @@ "members": [ {"name": "module", "type": "shader module"}, {"name": "entry point", "type": "char", "annotation": "const*", "length": "strlen"}, - {"name": "constant count", "type": "uint32_t", "tags": ["upstream"]}, - {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count", "tags": ["upstream"]}, + {"name": "constant count", "type": "uint32_t", "default": 0}, + {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count"}, {"name": "buffer count", "type": "uint32_t", "default": 0}, {"name": "buffers", "type": "vertex buffer layout", "annotation": "const*", "length": "buffer count"} ] @@ -1876,8 +1876,8 @@ "members": [ {"name": "module", "type": "shader module"}, {"name": "entry point", "type": "char", "annotation": "const*", "length": "strlen"}, - {"name": "constant count", "type": "uint32_t", "tags": ["upstream"]}, - {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count", "tags": ["upstream"]}, + {"name": "constant count", "type": "uint32_t", "default": 0}, + {"name": "constants", "type": "constant entry", "annotation": "const*", "length": "constant count"}, {"name": "target count", "type": "uint32_t"}, {"name": "targets", "type": "color target state", "annotation": "const*", "length": "target count"} ] diff --git a/src/dawn_native/ComputePipeline.cpp b/src/dawn_native/ComputePipeline.cpp index 8360b5608a..f789235b02 100644 --- a/src/dawn_native/ComputePipeline.cpp +++ b/src/dawn_native/ComputePipeline.cpp @@ -30,9 +30,10 @@ namespace dawn_native { DAWN_TRY(device->ValidateObject(descriptor->layout)); } - return ValidateProgrammableStage(device, descriptor->compute.module, - descriptor->compute.entryPoint, descriptor->layout, - SingleShaderStage::Compute); + return ValidateProgrammableStage( + device, descriptor->compute.module, descriptor->compute.entryPoint, + descriptor->compute.constantCount, descriptor->compute.constants, descriptor->layout, + SingleShaderStage::Compute); } // ComputePipelineBase @@ -43,7 +44,8 @@ namespace dawn_native { descriptor->layout, descriptor->label, {{SingleShaderStage::Compute, descriptor->compute.module, - descriptor->compute.entryPoint}}) { + descriptor->compute.entryPoint, descriptor->compute.constantCount, + descriptor->compute.constants}}) { } ComputePipelineBase::ComputePipelineBase(DeviceBase* device, ObjectBase::ErrorTag tag) diff --git a/src/dawn_native/Device.cpp b/src/dawn_native/Device.cpp index e1c87acb09..d557fb492c 100644 --- a/src/dawn_native/Device.cpp +++ b/src/dawn_native/Device.cpp @@ -135,9 +135,13 @@ namespace dawn_native { if (outDescriptor->layout == nullptr) { DAWN_TRY_ASSIGN(layoutRef, PipelineLayoutBase::CreateDefault( - device, {{SingleShaderStage::Compute, - outDescriptor->compute.module, - outDescriptor->compute.entryPoint}})); + device, {{ + SingleShaderStage::Compute, + outDescriptor->compute.module, + outDescriptor->compute.entryPoint, + outDescriptor->compute.constantCount, + outDescriptor->compute.constants, + }})); outDescriptor->layout = layoutRef.Get(); } diff --git a/src/dawn_native/Pipeline.cpp b/src/dawn_native/Pipeline.cpp index 4f658c5b5a..b4e822378b 100644 --- a/src/dawn_native/Pipeline.cpp +++ b/src/dawn_native/Pipeline.cpp @@ -26,6 +26,8 @@ namespace dawn_native { MaybeError ValidateProgrammableStage(DeviceBase* device, const ShaderModuleBase* module, const std::string& entryPoint, + uint32_t constantCount, + const ConstantEntry* constants, const PipelineLayoutBase* layout, SingleShaderStage stage) { DAWN_TRY(device->ValidateObject(module)); @@ -44,6 +46,14 @@ namespace dawn_native { DAWN_TRY(ValidateCompatibilityWithPipelineLayout(device, metadata, layout)); } + // Validate if overridable constants exist in shader module + // pipelineBase is not yet constructed at this moment so iterate constants from descriptor + 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); + } + return {}; } @@ -68,7 +78,12 @@ namespace dawn_native { // Record them internally. bool isFirstStage = mStageMask == wgpu::ShaderStage::None; mStageMask |= StageBit(shaderStage); - mStages[shaderStage] = {module, entryPointName, &metadata}; + mStages[shaderStage] = {module, entryPointName, &metadata, + std::vector()}; + auto& constants = mStages[shaderStage].constants; + for (uint32_t i = 0; i < stage.constantCount; i++) { + constants.emplace_back(stage.constants[i].key, stage.constants[i].value); + } // Compute the max() of all minBufferSizes across all stages. RequiredBufferSizes stageMinBufferSizes = diff --git a/src/dawn_native/Pipeline.h b/src/dawn_native/Pipeline.h index c4a7098ceb..c73d38968e 100644 --- a/src/dawn_native/Pipeline.h +++ b/src/dawn_native/Pipeline.h @@ -32,15 +32,20 @@ namespace dawn_native { MaybeError ValidateProgrammableStage(DeviceBase* device, const ShaderModuleBase* module, const std::string& entryPoint, + uint32_t constantCount, + const ConstantEntry* constants, const PipelineLayoutBase* layout, SingleShaderStage stage); + using PipelineConstantEntry = std::pair; struct ProgrammableStage { Ref module; std::string entryPoint; // The metadata lives as long as module, that's ref-ed in the same structure. const EntryPointMetadata* metadata = nullptr; + + std::vector constants; }; class PipelineBase : public ApiObjectBase, public CachedObject { diff --git a/src/dawn_native/PipelineLayout.h b/src/dawn_native/PipelineLayout.h index d130720bac..7371dab46d 100644 --- a/src/dawn_native/PipelineLayout.h +++ b/src/dawn_native/PipelineLayout.h @@ -44,6 +44,8 @@ namespace dawn_native { SingleShaderStage shaderStage; ShaderModuleBase* module; std::string entryPoint; + uint32_t constantCount = 0u; + ConstantEntry const* constants = nullptr; }; class PipelineLayoutBase : public ApiObjectBase, public CachedObject { diff --git a/src/dawn_native/RenderPipeline.cpp b/src/dawn_native/RenderPipeline.cpp index e4f70d589e..a73a401e9f 100644 --- a/src/dawn_native/RenderPipeline.cpp +++ b/src/dawn_native/RenderPipeline.cpp @@ -114,6 +114,7 @@ namespace dawn_native { } DAWN_TRY(ValidateProgrammableStage(device, descriptor->module, descriptor->entryPoint, + descriptor->constantCount, descriptor->constants, layout, SingleShaderStage::Vertex)); const EntryPointMetadata& vertexMetadata = descriptor->module->GetEntryPoint(descriptor->entryPoint); @@ -302,6 +303,7 @@ namespace dawn_native { } DAWN_TRY(ValidateProgrammableStage(device, descriptor->module, descriptor->entryPoint, + descriptor->constantCount, descriptor->constants, layout, SingleShaderStage::Fragment)); if (descriptor->targetCount > kMaxColorAttachments) { @@ -422,17 +424,20 @@ namespace dawn_native { DeviceBase* device, const RenderPipelineDescriptor* descriptor) { std::vector stages; - stages.push_back( - {SingleShaderStage::Vertex, descriptor->vertex.module, descriptor->vertex.entryPoint}); + stages.push_back({SingleShaderStage::Vertex, descriptor->vertex.module, + descriptor->vertex.entryPoint, descriptor->vertex.constantCount, + descriptor->vertex.constants}); if (descriptor->fragment != nullptr) { stages.push_back({SingleShaderStage::Fragment, descriptor->fragment->module, - descriptor->fragment->entryPoint}); + descriptor->fragment->entryPoint, descriptor->fragment->constantCount, + descriptor->fragment->constants}); } else if (device->IsToggleEnabled(Toggle::UseDummyFragmentInVertexOnlyPipeline)) { InternalPipelineStore* store = device->GetInternalPipelineStore(); // The dummy fragment shader module should already be initialized DAWN_ASSERT(store->dummyFragmentShader != nullptr); ShaderModuleBase* dummyFragmentShader = store->dummyFragmentShader.Get(); - stages.push_back({SingleShaderStage::Fragment, dummyFragmentShader, "fs_empty_main"}); + stages.push_back( + {SingleShaderStage::Fragment, dummyFragmentShader, "fs_empty_main", 0, nullptr}); } return stages; } diff --git a/src/dawn_native/ShaderModule.cpp b/src/dawn_native/ShaderModule.cpp index 9650f357c7..47badae871 100644 --- a/src/dawn_native/ShaderModule.cpp +++ b/src/dawn_native/ShaderModule.cpp @@ -33,6 +33,22 @@ namespace dawn_native { + EntryPointMetadata::OverridableConstant::Type GetDawnOverridableConstantType( + tint::inspector::OverridableConstant::Type type) { + switch (type) { + case tint::inspector::OverridableConstant::Type::kBool: + return EntryPointMetadata::OverridableConstant::Type::Boolean; + case tint::inspector::OverridableConstant::Type::kFloat32: + return EntryPointMetadata::OverridableConstant::Type::Float32; + case tint::inspector::OverridableConstant::Type::kInt32: + return EntryPointMetadata::OverridableConstant::Type::Int32; + case tint::inspector::OverridableConstant::Type::kUint32: + return EntryPointMetadata::OverridableConstant::Type::Uint32; + default: + UNREACHABLE(); + } + } + namespace { std::string GetShaderDeclarationString(BindGroupIndex group, BindingNumber binding) { @@ -629,13 +645,22 @@ namespace dawn_native { for (auto& entryPoint : entryPoints) { ASSERT(result.count(entryPoint.name) == 0); - if (!entryPoint.overridable_constants.empty()) { - return DAWN_VALIDATION_ERROR( - "Pipeline overridable constants are not implemented yet"); - } - auto metadata = std::make_unique(); + if (!entryPoint.overridable_constants.empty()) { + const auto& name2Id = inspector.GetConstantNameToIdMap(); + + for (auto& c : entryPoint.overridable_constants) { + EntryPointMetadata::OverridableConstant constant = { + name2Id.at(c.name), GetDawnOverridableConstantType(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; + } + } + DAWN_TRY_ASSIGN(metadata->stage, TintPipelineStageToShaderStage(entryPoint.stage)); if (metadata->stage == SingleShaderStage::Compute) { diff --git a/src/dawn_native/ShaderModule.h b/src/dawn_native/ShaderModule.h index 82737848c2..8f8081cedb 100644 --- a/src/dawn_native/ShaderModule.h +++ b/src/dawn_native/ShaderModule.h @@ -194,6 +194,16 @@ namespace dawn_native { // The shader stage for this binding. SingleShaderStage stage; + + struct OverridableConstant { + uint32_t id; + // 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; + }; + + // Store overridableConstants from tint program + std::unordered_map overridableConstants; }; class ShaderModuleBase : public ApiObjectBase, public CachedObject { diff --git a/src/dawn_native/vulkan/ComputePipelineVk.cpp b/src/dawn_native/vulkan/ComputePipelineVk.cpp index e289845df8..649ab5478f 100644 --- a/src/dawn_native/vulkan/ComputePipelineVk.cpp +++ b/src/dawn_native/vulkan/ComputePipelineVk.cpp @@ -52,8 +52,15 @@ namespace dawn_native { namespace vulkan { ToBackend(computeStage.module.Get()) ->GetTransformedModuleHandle(computeStage.entryPoint.c_str(), ToBackend(GetLayout()))); + createInfo.stage.pName = computeStage.entryPoint.c_str(); - createInfo.stage.pSpecializationInfo = nullptr; + + std::vector specializationDataEntries; + std::vector specializationMapEntries; + VkSpecializationInfo specializationInfo{}; + createInfo.stage.pSpecializationInfo = + GetVkSpecializationInfo(computeStage, &specializationInfo, &specializationDataEntries, + &specializationMapEntries); Device* device = ToBackend(GetDevice()); diff --git a/src/dawn_native/vulkan/RenderPipelineVk.cpp b/src/dawn_native/vulkan/RenderPipelineVk.cpp index 9fa582d870..5b196eb400 100644 --- a/src/dawn_native/vulkan/RenderPipelineVk.cpp +++ b/src/dawn_native/vulkan/RenderPipelineVk.cpp @@ -341,21 +341,25 @@ namespace dawn_native { namespace vulkan { // There are at most 2 shader stages in render pipeline, i.e. vertex and fragment std::array shaderStages; + std::array, 2> specializationDataEntriesPerStages; + std::array, 2> specializationMapEntriesPerStages; + std::array specializationInfoPerStages; uint32_t stageCount = 0; for (auto stage : IterateStages(this->GetStageMask())) { VkPipelineShaderStageCreateInfo shaderStage; + const ProgrammableStage& programmableStage = GetStage(stage); DAWN_TRY_ASSIGN(shaderStage.module, - ToBackend(GetStage(stage).module) - ->GetTransformedModuleHandle(GetStage(stage).entryPoint.c_str(), + ToBackend(programmableStage.module) + ->GetTransformedModuleHandle(programmableStage.entryPoint.c_str(), ToBackend(GetLayout()))); shaderStage.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO; shaderStage.pNext = nullptr; shaderStage.flags = 0; shaderStage.pSpecializationInfo = nullptr; - shaderStage.pName = GetStage(stage).entryPoint.c_str(); + shaderStage.pName = programmableStage.entryPoint.c_str(); switch (stage) { case dawn_native::SingleShaderStage::Vertex: { @@ -373,6 +377,11 @@ namespace dawn_native { namespace vulkan { } } + shaderStage.pSpecializationInfo = + GetVkSpecializationInfo(programmableStage, &specializationInfoPerStages[stageCount], + &specializationDataEntriesPerStages[stageCount], + &specializationMapEntriesPerStages[stageCount]); + DAWN_ASSERT(stageCount < 2); shaderStages[stageCount] = shaderStage; stageCount++; diff --git a/src/dawn_native/vulkan/UtilsVulkan.cpp b/src/dawn_native/vulkan/UtilsVulkan.cpp index b04c0a4fe1..6d73eec289 100644 --- a/src/dawn_native/vulkan/UtilsVulkan.cpp +++ b/src/dawn_native/vulkan/UtilsVulkan.cpp @@ -17,6 +17,8 @@ #include "common/Assert.h" #include "dawn_native/EnumMaskIterator.h" #include "dawn_native/Format.h" +#include "dawn_native/Pipeline.h" +#include "dawn_native/ShaderModule.h" #include "dawn_native/vulkan/DeviceVk.h" #include "dawn_native/vulkan/Forward.h" #include "dawn_native/vulkan/TextureVk.h" @@ -195,4 +197,63 @@ namespace dawn_native { namespace vulkan { device->fn.SetDebugUtilsObjectNameEXT(device->GetVkDevice(), &objectNameInfo); } } + + VkSpecializationInfo* GetVkSpecializationInfo( + const ProgrammableStage& programmableStage, + VkSpecializationInfo* specializationInfo, + std::vector* specializationDataEntries, + std::vector* specializationMapEntries) { + ASSERT(specializationInfo); + ASSERT(specializationDataEntries); + ASSERT(specializationMapEntries); + + if (programmableStage.constants.size() == 0) { + return nullptr; + } + + const EntryPointMetadata& entryPointMetaData = + programmableStage.module->GetEntryPoint(programmableStage.entryPoint); + + for (const auto& pipelineConstant : programmableStage.constants) { + const std::string& name = pipelineConstant.first; + double value = pipelineConstant.second; + + // This is already validated so `name` must exist + const auto& moduleConstant = entryPointMetaData.overridableConstants.at(name); + + specializationMapEntries->push_back( + VkSpecializationMapEntry{moduleConstant.id, + static_cast(specializationDataEntries->size() * + sizeof(SpecializationDataEntry)), + sizeof(SpecializationDataEntry)}); + + SpecializationDataEntry entry; + switch (moduleConstant.type) { + case EntryPointMetadata::OverridableConstant::Type::Boolean: + entry.b = static_cast(value); + break; + case EntryPointMetadata::OverridableConstant::Type::Float32: + entry.f32 = static_cast(value); + break; + case EntryPointMetadata::OverridableConstant::Type::Int32: + entry.i32 = static_cast(value); + break; + case EntryPointMetadata::OverridableConstant::Type::Uint32: + entry.u32 = static_cast(value); + break; + default: + UNREACHABLE(); + } + specializationDataEntries->push_back(entry); + } + + specializationInfo->mapEntryCount = static_cast(specializationMapEntries->size()); + specializationInfo->pMapEntries = specializationMapEntries->data(); + specializationInfo->dataSize = + specializationDataEntries->size() * sizeof(SpecializationDataEntry); + specializationInfo->pData = specializationDataEntries->data(); + + return specializationInfo; + } + }} // namespace dawn_native::vulkan diff --git a/src/dawn_native/vulkan/UtilsVulkan.h b/src/dawn_native/vulkan/UtilsVulkan.h index 23c36a9f93..8bd2a84544 100644 --- a/src/dawn_native/vulkan/UtilsVulkan.h +++ b/src/dawn_native/vulkan/UtilsVulkan.h @@ -19,6 +19,10 @@ #include "dawn_native/Commands.h" #include "dawn_native/dawn_platform.h" +namespace dawn_native { + struct ProgrammableStage; +} // namespace dawn_native + namespace dawn_native { namespace vulkan { class Device; @@ -107,6 +111,24 @@ namespace dawn_native { namespace vulkan { const char* prefix, std::string label = ""); + // Helpers for creating VkSpecializationInfo + // The WebGPU overridable constants only support these scalar types + union SpecializationDataEntry { + bool b; + float f32; + int32_t i32; + uint32_t u32; + }; + + // Returns nullptr or &specializationInfo + // specializationInfo, specializationDataEntries, specializationMapEntries needs to + // be alive at least until VkSpecializationInfo is passed into Vulkan Create*Pipelines + VkSpecializationInfo* GetVkSpecializationInfo( + const ProgrammableStage& programmableStage, + VkSpecializationInfo* specializationInfo, + std::vector* specializationDataEntries, + std::vector* specializationMapEntries); + }} // namespace dawn_native::vulkan #endif // DAWNNATIVE_VULKAN_UTILSVULKAN_H_ diff --git a/src/tests/BUILD.gn b/src/tests/BUILD.gn index ce271362e9..b3aaa208b2 100644 --- a/src/tests/BUILD.gn +++ b/src/tests/BUILD.gn @@ -207,6 +207,7 @@ test("dawn_unittests") { "unittests/validation/LabelTests.cpp", "unittests/validation/MinimumBufferSizeValidationTests.cpp", "unittests/validation/MultipleDeviceTests.cpp", + "unittests/validation/OverridableConstantsValidationTests.cpp", "unittests/validation/QueryValidationTests.cpp", "unittests/validation/QueueOnSubmittedWorkDoneValidationTests.cpp", "unittests/validation/QueueSubmitValidationTests.cpp", diff --git a/src/tests/end2end/ShaderTests.cpp b/src/tests/end2end/ShaderTests.cpp index c635048867..bbb4791f84 100644 --- a/src/tests/end2end/ShaderTests.cpp +++ b/src/tests/end2end/ShaderTests.cpp @@ -17,18 +17,37 @@ #include "utils/ComboRenderPipelineDescriptor.h" #include "utils/WGPUHelpers.h" +#include #include -class ShaderTests : public DawnTest {}; +class ShaderTests : public DawnTest { + public: + wgpu::Buffer CreateBuffer(const uint32_t count) { + std::vector data(count, 0); + uint64_t bufferSize = static_cast(data.size() * sizeof(uint32_t)); + return utils::CreateBufferFromData(device, data.data(), bufferSize, + wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc); + } + wgpu::ComputePipeline CreateComputePipeline( + const std::string& shader, + const char* entryPoint, + const std::vector* constants = nullptr) { + wgpu::ComputePipelineDescriptor csDesc; + csDesc.compute.module = utils::CreateShaderModule(device, shader.c_str()); + csDesc.compute.entryPoint = entryPoint; + if (constants) { + csDesc.compute.constants = constants->data(); + csDesc.compute.constantCount = constants->size(); + } + return device.CreateComputePipeline(&csDesc); + } +}; // Test that log2 is being properly calculated, base on crbug.com/1046622 TEST_P(ShaderTests, ComputeLog2) { uint32_t const kSteps = 19; - std::vector data(kSteps, 0); std::vector expected{0, 1, 1, 2, 2, 3, 3, 4, 4, 5, 5, 6, 6, 7, 7, 8, 8, 9, 32}; - uint64_t bufferSize = static_cast(data.size() * sizeof(uint32_t)); - wgpu::Buffer buffer = utils::CreateBufferFromData( - device, data.data(), bufferSize, wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc); + wgpu::Buffer buffer = CreateBuffer(kSteps); std::string shader = R"( [[block]] struct Buf { @@ -61,10 +80,7 @@ TEST_P(ShaderTests, ComputeLog2) { buf.data[18] = u32(log2(4294967295.0 * factor)); })"; - wgpu::ComputePipelineDescriptor csDesc; - csDesc.compute.module = utils::CreateShaderModule(device, shader.c_str()); - csDesc.compute.entryPoint = "main"; - wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc); + wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main"); wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}}); @@ -305,23 +321,6 @@ fn fragmentMain(input : VertexOut) -> [[location(0)]] vec4 { wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&rpDesc); } -// Feature currently not implemented in Tint, so should fail validation. -TEST_P(ShaderTests, PipelineOverridableUsed) { - DAWN_TEST_UNSUPPORTED_IF(HasToggleEnabled("skip_validation")); - DAWN_TEST_UNSUPPORTED_IF(!HasToggleEnabled("use_tint_generator")); - - std::string shader = R"( -[[override]] let foo : f32; - -[[stage(compute), workgroup_size(1)]] -fn ep_func() { - var local_foo : f32; - local_foo = foo; - return; -})"; - ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, shader.c_str())); -} - // This is a regression test for an issue caused by the FirstIndexOffset transfrom being done before // the BindingRemapper, causing an intermediate AST to be invalid (and fail the overall // compilation). @@ -393,6 +392,299 @@ fn main([[location(0)]] pos : vec4) -> [[builtin(position)]] vec4 { device.CreateRenderPipeline(&descriptor); } +// Test overridable constants without numeric identifiers +// TODO(tint:1155): Implicit numeric ID is undetermined in tint +TEST_P(ShaderTests, DISABLED_OverridableConstants) { + // TODO(dawn:1041): Only Vulkan backend is implemented + DAWN_TEST_UNSUPPORTED_IF(!IsVulkan()); + + uint32_t const kCount = 15; + 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"( +[[override]] let c0: bool; // type: bool +[[override]] let c1: bool = false; // default override +[[override]] let c2: f32; // type: float32 +[[override]] let c3: f32 = 0.0; // default override +[[override]] let c4: f32 = 4.0; // default +[[override]] let c5: i32; // type: int32 +[[override]] let c6: i32 = 0; // default override +[[override]] let c7: i32 = 7; // default +[[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; +}; + +[[group(0), binding(0)]] var buf : Buf; + +[[stage(compute), workgroup_size(1)]] fn main() { + 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); + buf.data[11] = u32(c11); + buf.data[12] = u32(c12); + buf.data[13] = u32(c13); + buf.data[14] = u32(c14); +})"; + + std::vector constants; + constants.push_back({nullptr, "c0", 0}); + constants.push_back({nullptr, "c1", 1}); + constants.push_back({nullptr, "c2", 2}); + constants.push_back({nullptr, "c3", 3}); + // c4 is not assigned, testing default value + constants.push_back({nullptr, "c5", 5}); + constants.push_back({nullptr, "c6", 6}); + // c7 is not assigned, testing default value + 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); + + wgpu::BindGroup bindGroup = + utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}}); + + wgpu::CommandBuffer commands; + { + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetPipeline(pipeline); + pass.SetBindGroup(0, bindGroup); + pass.Dispatch(1); + pass.EndPass(); + + commands = encoder.Finish(); + } + + queue.Submit(1, &commands); + + EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), buffer, 0, kCount); +} + +// Test overridable constants with numeric identifiers +TEST_P(ShaderTests, OverridableConstantsNumericIdentifiers) { + // TODO(dawn:1041): Only Vulkan backend is implemented + DAWN_TEST_UNSUPPORTED_IF(!IsVulkan()); + + uint32_t const kCount = 4; + std::vector expected{1u, 2u, 3u, 0u}; + wgpu::Buffer buffer = CreateBuffer(kCount); + + std::string shader = R"( +[[override(1001)]] let c1: u32; // some big numeric id +[[override(1)]] let c2: u32 = 0u; // id == 1 might collide with some generated constant id +[[override(1003)]] let c3: u32 = 3u; // default +[[override(1004)]] let c4: u32; // default unspecified + +[[block]] struct Buf { + data : array; +}; + +[[group(0), binding(0)]] var buf : Buf; + +[[stage(compute), workgroup_size(1)]] fn main() { + buf.data[0] = c1; + buf.data[1] = c2; + buf.data[2] = c3; + buf.data[3] = c4; +})"; + + std::vector constants; + 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 + + wgpu::ComputePipeline pipeline = CreateComputePipeline(shader, "main", &constants); + + wgpu::BindGroup bindGroup = + utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer}}); + + wgpu::CommandBuffer commands; + { + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetPipeline(pipeline); + pass.SetBindGroup(0, bindGroup); + pass.Dispatch(1); + pass.EndPass(); + + commands = encoder.Finish(); + } + + queue.Submit(1, &commands); + + EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), buffer, 0, kCount); +} + +// Test overridable constants for different entry points +TEST_P(ShaderTests, OverridableConstantsMultipleEntryPoints) { + // TODO(dawn:1041): Only Vulkan backend is implemented + DAWN_TEST_UNSUPPORTED_IF(!IsVulkan()); + + uint32_t const kCount = 1; + std::vector expected1{1u}; + std::vector expected2{2u}; + + wgpu::Buffer buffer1 = CreateBuffer(kCount); + wgpu::Buffer buffer2 = CreateBuffer(kCount); + + std::string shader = R"( +[[override(1001)]] let c1: u32; +[[override(1002)]] let c2: u32; + +[[block]] struct Buf { + data : array; +}; + +[[group(0), binding(0)]] var buf : Buf; + +[[stage(compute), workgroup_size(1)]] fn main1() { + buf.data[0] = c1; +} + +[[stage(compute), workgroup_size(1)]] fn main2() { + buf.data[0] = c2; +} +)"; + + std::vector constants1; + constants1.push_back({nullptr, "1001", 1}); + std::vector constants2; + constants2.push_back({nullptr, "1002", 2}); + + wgpu::ShaderModule shaderModule = utils::CreateShaderModule(device, shader.c_str()); + + wgpu::ComputePipelineDescriptor csDesc1; + csDesc1.compute.module = shaderModule; + csDesc1.compute.entryPoint = "main1"; + csDesc1.compute.constants = constants1.data(); + csDesc1.compute.constantCount = constants1.size(); + wgpu::ComputePipeline pipeline1 = device.CreateComputePipeline(&csDesc1); + + wgpu::ComputePipelineDescriptor csDesc2; + csDesc2.compute.module = shaderModule; + csDesc2.compute.entryPoint = "main2"; + csDesc2.compute.constants = constants2.data(); + csDesc2.compute.constantCount = constants2.size(); + wgpu::ComputePipeline pipeline2 = device.CreateComputePipeline(&csDesc2); + + wgpu::BindGroup bindGroup1 = + utils::MakeBindGroup(device, pipeline1.GetBindGroupLayout(0), {{0, buffer1}}); + wgpu::BindGroup bindGroup2 = + utils::MakeBindGroup(device, pipeline2.GetBindGroupLayout(0), {{0, buffer2}}); + + wgpu::CommandBuffer commands; + { + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); + pass.SetPipeline(pipeline1); + pass.SetBindGroup(0, bindGroup1); + pass.Dispatch(1); + + pass.SetPipeline(pipeline2); + pass.SetBindGroup(0, bindGroup2); + pass.Dispatch(1); + + pass.EndPass(); + + commands = encoder.Finish(); + } + + queue.Submit(1, &commands); + + EXPECT_BUFFER_U32_RANGE_EQ(expected1.data(), buffer1, 0, kCount); + EXPECT_BUFFER_U32_RANGE_EQ(expected2.data(), buffer2, 0, kCount); +} + +// Test overridable constants with render pipeline +// Draw a triangle covering the render target, with vertex position and color values from +// overridable constants +TEST_P(ShaderTests, OverridableConstantsRenderPipeline) { + // TODO(dawn:1041): Only Vulkan backend is implemented + DAWN_TEST_UNSUPPORTED_IF(!IsVulkan()); + + wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"( +[[override(1111)]] let xright: f32; +[[override(2222)]] let ytop: f32; +[[stage(vertex)]] +fn main([[builtin(vertex_index)]] VertexIndex : u32) + -> [[builtin(position)]] vec4 { + var pos = array, 3>( + vec2(-1.0, ytop), + vec2(-1.0, -ytop), + vec2(xright, 0.0)); + + return vec4(pos[VertexIndex], 0.0, 1.0); +})"); + + wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"( +[[override(1000)]] let intensity: f32 = 0.0; +[[stage(fragment)]] fn main() + -> [[location(0)]] vec4 { + return vec4(intensity, intensity, intensity, 1.0); +})"); + + utils::BasicRenderPass renderPass = utils::CreateBasicRenderPass(device, 1, 1); + + utils::ComboRenderPipelineDescriptor descriptor; + descriptor.vertex.module = vsModule; + descriptor.cFragment.module = fsModule; + descriptor.primitive.topology = wgpu::PrimitiveTopology::TriangleList; + descriptor.cTargets[0].format = renderPass.colorFormat; + + std::vector vertexConstants; + vertexConstants.push_back({nullptr, "1111", 3.0}); // x right + vertexConstants.push_back({nullptr, "2222", 3.0}); // y top + descriptor.vertex.constants = vertexConstants.data(); + descriptor.vertex.constantCount = vertexConstants.size(); + std::vector fragmentConstants; + fragmentConstants.push_back({nullptr, "1000", 1.0}); // color intensity + descriptor.cFragment.constants = fragmentConstants.data(); + descriptor.cFragment.constantCount = fragmentConstants.size(); + + wgpu::RenderPipeline pipeline = device.CreateRenderPipeline(&descriptor); + + wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); + wgpu::RenderPassEncoder pass = encoder.BeginRenderPass(&renderPass.renderPassInfo); + pass.SetPipeline(pipeline); + pass.Draw(3); + pass.EndPass(); + wgpu::CommandBuffer commands = encoder.Finish(); + queue.Submit(1, &commands); + + EXPECT_PIXEL_RGBA8_EQ(RGBA8(255, 255, 255, 255), renderPass.color, 0, 0); +} + +// TODO(tint:1155): Test overridable constants used for workgroup size + DAWN_INSTANTIATE_TEST(ShaderTests, D3D12Backend(), MetalBackend(), diff --git a/src/tests/unittests/validation/OverridableConstantsValidationTests.cpp b/src/tests/unittests/validation/OverridableConstantsValidationTests.cpp new file mode 100644 index 0000000000..af96c394e9 --- /dev/null +++ b/src/tests/unittests/validation/OverridableConstantsValidationTests.cpp @@ -0,0 +1,147 @@ +// Copyright 2021 The Dawn Authors +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "common/Constants.h" +#include "tests/unittests/validation/ValidationTest.h" +#include "utils/WGPUHelpers.h" + +class ComputePipelineOverridableConstantsValidationTest : public ValidationTest { + protected: + void SetUp() override { + ValidationTest::SetUp(); + + computeModule = utils::CreateShaderModule(device, R"( +[[override]] let c0: bool; // type: bool +[[override]] let c1: bool = false; // default override +[[override]] let c2: f32; // type: float32 +[[override]] let c3: f32 = 0.0; // default override +[[override]] let c4: f32 = 4.0; // default +[[override]] let c5: i32; // type: int32 +[[override]] let c6: i32 = 0; // default override +[[override]] let c7: i32 = 7; // default +[[override]] let c8: u32; // type: uint32 +[[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); +})"); + } + + void TestCreatePipeline(const std::vector& constants) { + wgpu::ComputePipelineDescriptor csDesc; + csDesc.compute.module = computeModule; + csDesc.compute.entryPoint = "main"; + csDesc.compute.constants = constants.data(); + csDesc.compute.constantCount = constants.size(); + wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&csDesc); + } + + wgpu::ShaderModule computeModule; + wgpu::Buffer buffer; +}; + +// Basic constants lookup tests +TEST_F(ComputePipelineOverridableConstantsValidationTest, ConstantsIdentifierLookUp) { + { + // Valid: no constants specified + std::vector constants; + TestCreatePipeline(constants); + } + { + // Valid: find by constant name + std::vector constants{{nullptr, "c0", 0}}; + TestCreatePipeline(constants); + } + { + // Valid: set the same constant twice + std::vector constants{ + {nullptr, "c0", 0}, + {nullptr, "c0", 1}, + }; + TestCreatePipeline(constants); + } + { + // Valid: find by constant numeric id + std::vector constants{{nullptr, "1000", 0}}; + TestCreatePipeline(constants); + } + { + // Error: constant numeric id not specified + std::vector constants{{nullptr, "9999", 0}}; + ASSERT_DEVICE_ERROR(TestCreatePipeline(constants)); + } + { + // Error: constant name doesn't exit + std::vector constants{{nullptr, "c99", 0}}; + ASSERT_DEVICE_ERROR(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) { + { + // Error: constant numeric id not explicitly specified + // But could be impliciltly assigned to one of the constants + std::vector constants{{nullptr, "0", 0}}; + ASSERT_DEVICE_ERROR(TestCreatePipeline(constants)); + } + { + // Error: constant numeric id not explicitly specified + // But could be impliciltly assigned to one of the constants + std::vector constants{{nullptr, "1", 0}}; + ASSERT_DEVICE_ERROR(TestCreatePipeline(constants)); + } + { + // Error: constant numeric id not explicitly specified + // But could be impliciltly assigned to one of the constants + std::vector constants{{nullptr, "2", 0}}; + ASSERT_DEVICE_ERROR(TestCreatePipeline(constants)); + } + { + // Error: constant numeric id not explicitly specified + // But could be impliciltly assigned to one of the constants + std::vector constants{{nullptr, "3", 0}}; + ASSERT_DEVICE_ERROR(TestCreatePipeline(constants)); + } +} + +// Test that identifiers are unique +// TODO(tint:1155): missing feature in tint to differentiate explicitly specified numeric ID +TEST_F(ComputePipelineOverridableConstantsValidationTest, DISABLED_ConstantsIdentifierUnique) { + { + // Error: constant with numeric id cannot be referenced with variable name + std::vector constants{{nullptr, "c10", 0}}; + ASSERT_DEVICE_ERROR(TestCreatePipeline(constants)); + } +} \ No newline at end of file diff --git a/src/tests/unittests/validation/ShaderModuleValidationTests.cpp b/src/tests/unittests/validation/ShaderModuleValidationTests.cpp index 2c495babc5..da78826207 100644 --- a/src/tests/unittests/validation/ShaderModuleValidationTests.cpp +++ b/src/tests/unittests/validation/ShaderModuleValidationTests.cpp @@ -493,3 +493,22 @@ TEST_F(ShaderModuleValidationTest, ComputeWorkgroupStorageSizeLimits) { ASSERT_DEVICE_ERROR(MakeShaderWithWorkgroupStorage(0, kMaxMat4Count + 1)); ASSERT_DEVICE_ERROR(MakeShaderWithWorkgroupStorage(4, kMaxMat4Count)); } + +// Test that numeric ID must be unique +TEST_F(ShaderModuleValidationTest, OverridableConstantsNumericIDConflicts) { + ASSERT_DEVICE_ERROR(utils::CreateShaderModule(device, R"( +[[override(1234)]] let c0: u32; +[[override(1234)]] let c1: u32; + +[[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] = c0; + buf.data[1] = c1; +})")); +} \ No newline at end of file