From c6c4588036c11e33efade087a60c22b048625d4e Mon Sep 17 00:00:00 2001 From: shrekshao Date: Thu, 18 Nov 2021 17:47:45 +0000 Subject: [PATCH] Pipeline overridable constants: D3D12 backend D3D12 doesn't have native pipeline constants feature. This is done by using #define. Add some new tests to make sure these define approaches work as expected. Also makes duplicate pipeline constant entries an invalid case. Bug: dawn:1137 Change-Id: Iefed44a749625b535bbafbb39f42699f0b42e06a Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/68860 Reviewed-by: Austin Eng Commit-Queue: Shrek Shao --- src/dawn_native/Pipeline.cpp | 20 +- src/dawn_native/Pipeline.h | 6 +- src/dawn_native/ShaderModule.h | 6 +- .../d3d12/ComputePipelineD3D12.cpp | 5 +- src/dawn_native/d3d12/RenderPipelineD3D12.cpp | 8 +- src/dawn_native/d3d12/ShaderModuleD3D12.cpp | 204 +++++++++++++++++- src/dawn_native/d3d12/ShaderModuleD3D12.h | 6 +- src/tests/end2end/ShaderTests.cpp | 86 +++++++- .../OverridableConstantsValidationTests.cpp | 8 +- 9 files changed, 309 insertions(+), 40 deletions(-) diff --git a/src/dawn_native/Pipeline.cpp b/src/dawn_native/Pipeline.cpp index 2a5e31696e..7b24460451 100644 --- a/src/dawn_native/Pipeline.cpp +++ b/src/dawn_native/Pipeline.cpp @@ -55,17 +55,22 @@ namespace dawn_native { // pipelineBase is not yet constructed at this moment so iterate constants from descriptor size_t numUninitializedConstants = metadata.uninitializedOverridableConstants.size(); // Keep an 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.", + "Pipeline overridable constant \"%s\" not found in %s.", constants[i].key, module); - if (metadata.uninitializedOverridableConstants.count(constants[i].key) > 0 && - stageInitializedConstantIdentifiers.count(constants[i].key) == 0) { - numUninitializedConstants--; + if (stageInitializedConstantIdentifiers.count(constants[i].key) == 0) { + if (metadata.uninitializedOverridableConstants.count(constants[i].key) > 0) { + numUninitializedConstants--; + } stageInitializedConstantIdentifiers.insert(constants[i].key); + } else { + // There are duplicate initializations + return DAWN_FORMAT_VALIDATION_ERROR( + "Pipeline overridable constants \"%s\" is set more than once in %s", + constants[i].key, module); } } @@ -116,11 +121,10 @@ namespace dawn_native { // Record them internally. bool isFirstStage = mStageMask == wgpu::ShaderStage::None; mStageMask |= StageBit(shaderStage); - mStages[shaderStage] = {module, entryPointName, &metadata, - std::vector()}; + mStages[shaderStage] = {module, entryPointName, &metadata, {}}; 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); + constants.emplace(stage.constants[i].key, stage.constants[i].value); } // Compute the max() of all minBufferSizes across all stages. diff --git a/src/dawn_native/Pipeline.h b/src/dawn_native/Pipeline.h index 7a1e09a595..10aec50ca7 100644 --- a/src/dawn_native/Pipeline.h +++ b/src/dawn_native/Pipeline.h @@ -37,7 +37,9 @@ namespace dawn_native { const PipelineLayoutBase* layout, SingleShaderStage stage); - using PipelineConstantEntry = std::pair; + // Use map to make sure constant keys are sorted for creating shader cache keys + using PipelineConstantEntries = std::map; + struct ProgrammableStage { Ref module; std::string entryPoint; @@ -45,7 +47,7 @@ namespace dawn_native { // The metadata lives as long as module, that's ref-ed in the same structure. const EntryPointMetadata* metadata = nullptr; - std::vector constants; + PipelineConstantEntries constants; }; class PipelineBase : public ApiObjectBase, public CachedObject { diff --git a/src/dawn_native/ShaderModule.h b/src/dawn_native/ShaderModule.h index 4c69343955..660f925ffa 100644 --- a/src/dawn_native/ShaderModule.h +++ b/src/dawn_native/ShaderModule.h @@ -222,9 +222,11 @@ namespace dawn_native { OverridableConstantScalar defaultValue; }; + using OverridableConstantsMap = std::unordered_map; + // Map identifier to overridable constant // Identifier is unique: either the variable name or the numeric ID if specified - std::unordered_map overridableConstants; + OverridableConstantsMap overridableConstants; // Overridable constants that are not initialized in shaders // They need value initialization from pipeline stage or it is a validation error @@ -254,7 +256,7 @@ namespace dawn_native { // Return true iff the program has an entrypoint called `entryPoint`. bool HasEntryPoint(const std::string& entryPoint) const; - // Returns the metadata for the given `entryPoint`. HasEntryPoint with the same argument + // Return the metadata for the given `entryPoint`. HasEntryPoint with the same argument // must be true. const EntryPointMetadata& GetEntryPoint(const std::string& entryPoint) const; diff --git a/src/dawn_native/d3d12/ComputePipelineD3D12.cpp b/src/dawn_native/d3d12/ComputePipelineD3D12.cpp index 744cb68b93..dd8d225400 100644 --- a/src/dawn_native/d3d12/ComputePipelineD3D12.cpp +++ b/src/dawn_native/d3d12/ComputePipelineD3D12.cpp @@ -48,9 +48,8 @@ namespace dawn_native { namespace d3d12 { d3dDesc.pRootSignature = ToBackend(GetLayout())->GetRootSignature(); CompiledShader compiledShader; - DAWN_TRY_ASSIGN(compiledShader, - module->Compile(computeStage.entryPoint.c_str(), SingleShaderStage::Compute, - ToBackend(GetLayout()), compileFlags)); + DAWN_TRY_ASSIGN(compiledShader, module->Compile(computeStage, SingleShaderStage::Compute, + ToBackend(GetLayout()), compileFlags)); d3dDesc.CS = compiledShader.GetD3D12ShaderBytecode(); auto* d3d12Device = device->GetD3D12Device(); DAWN_TRY(CheckHRESULT( diff --git a/src/dawn_native/d3d12/RenderPipelineD3D12.cpp b/src/dawn_native/d3d12/RenderPipelineD3D12.cpp index 44fd0c4772..dbe62522d4 100644 --- a/src/dawn_native/d3d12/RenderPipelineD3D12.cpp +++ b/src/dawn_native/d3d12/RenderPipelineD3D12.cpp @@ -348,10 +348,10 @@ namespace dawn_native { namespace d3d12 { PerStage compiledShader; for (auto stage : IterateStages(GetStageMask())) { - DAWN_TRY_ASSIGN(compiledShader[stage], - ToBackend(pipelineStages[stage].module) - ->Compile(pipelineStages[stage].entryPoint.c_str(), stage, - ToBackend(GetLayout()), compileFlags)); + DAWN_TRY_ASSIGN( + compiledShader[stage], + ToBackend(pipelineStages[stage].module) + ->Compile(pipelineStages[stage], stage, ToBackend(GetLayout()), compileFlags)); *shaders[stage] = compiledShader[stage].GetD3D12ShaderBytecode(); } diff --git a/src/dawn_native/d3d12/ShaderModuleD3D12.cpp b/src/dawn_native/d3d12/ShaderModuleD3D12.cpp index 34ea0d50aa..ec5b3b3863 100644 --- a/src/dawn_native/d3d12/ShaderModuleD3D12.cpp +++ b/src/dawn_native/d3d12/ShaderModuleD3D12.cpp @@ -17,6 +17,8 @@ #include "common/Assert.h" #include "common/BitSetIterator.h" #include "common/Log.h" +#include "common/WindowsUtils.h" +#include "dawn_native/Pipeline.h" #include "dawn_native/TintUtils.h" #include "dawn_native/d3d12/BindGroupLayoutD3D12.h" #include "dawn_native/d3d12/D3D12Error.h" @@ -32,6 +34,94 @@ #include #include +namespace dawn_native { + template + struct NumberToString { + static StringType ToStringAsValue(T v); + static StringType ToStringAsId(T v); + }; + + template + struct NumberToString { + static constexpr char kSpecConstantPrefix[] = "WGSL_SPEC_CONSTANT_"; + static std::string ToStringAsValue(T v) { + return std::to_string(v); + } + static std::string ToStringAsId(T v) { + return std::to_string(v); + } + }; + + template + struct NumberToString { + static constexpr WCHAR kSpecConstantPrefix[] = L"WGSL_SPEC_CONSTANT_"; + static std::wstring ToStringAsValue(T v) { + return std::to_wstring(v); + } + static std::wstring ToStringAsId(T v) { + return std::to_wstring(v); + } + }; + + template <> + struct NumberToString { + static std::string ToStringAsValue(float v) { + std::ostringstream out; + // 32 bit float has 7 decimal digits of precision so setting n to 8 should be enough + out.precision(8); + out << std::fixed << v; + return out.str(); + } + }; + + template <> + struct NumberToString { + static std::wstring ToStringAsValue(float v) { + std::basic_ostringstream out; + // 32 bit float has 7 decimal digits of precision so setting n to 8 should be enough + out.precision(8); + out << std::fixed << v; + return out.str(); + } + }; + + template <> + struct NumberToString { + static std::string ToStringAsValue(uint32_t v) { + return std::to_string(v) + "u"; + } + }; + + template <> + struct NumberToString { + static std::wstring ToStringAsValue(uint32_t v) { + return std::to_wstring(v) + L"u"; + } + }; + + template + StringType GetHLSLValueString(EntryPointMetadata::OverridableConstant::Type dawnType, + const OverridableConstantScalar* entry, + double value = 0) { + switch (dawnType) { + case EntryPointMetadata::OverridableConstant::Type::Boolean: + return NumberToString::ToStringAsValue( + entry ? entry->b : static_cast(value)); + case EntryPointMetadata::OverridableConstant::Type::Float32: + return NumberToString::ToStringAsValue( + entry ? entry->f32 : static_cast(value)); + case EntryPointMetadata::OverridableConstant::Type::Int32: + return NumberToString::ToStringAsValue( + entry ? entry->i32 : static_cast(value)); + case EntryPointMetadata::OverridableConstant::Type::Uint32: + return NumberToString::ToStringAsValue( + entry ? entry->u32 : static_cast(value)); + default: + UNREACHABLE(); + } + } +} // namespace dawn_native + namespace dawn_native { namespace d3d12 { namespace { @@ -109,6 +199,8 @@ namespace dawn_native { namespace d3d12 { bool usesNumWorkgroups; uint32_t numWorkgroupsRegisterSpace; uint32_t numWorkgroupsShaderRegister; + const PipelineConstantEntries* pipelineConstantEntries; + const EntryPointMetadata::OverridableConstantsMap* shaderEntryPointConstants; // FXC/DXC common inputs bool disableWorkgroupInit; @@ -128,7 +220,8 @@ namespace dawn_native { namespace d3d12 { uint32_t compileFlags, const Device* device, const tint::Program* program, - const EntryPointMetadata& entryPoint) { + const EntryPointMetadata& entryPoint, + const ProgrammableStage& programmableStage) { Compiler compiler; uint64_t dxcVersion = 0; if (device->IsToggleEnabled(Toggle::UseDXC)) { @@ -200,6 +293,10 @@ namespace dawn_native { namespace d3d12 { request.dxcVersion = compiler == Compiler::DXC ? dxcVersion : 0; request.deviceInfo = &device->GetDeviceInfo(); request.hasShaderFloat16Feature = device->IsFeatureEnabled(Feature::ShaderFloat16); + request.pipelineConstantEntries = &programmableStage.constants; + request.shaderEntryPointConstants = + &programmableStage.module->GetEntryPoint(programmableStage.entryPoint) + .overridableConstants; return std::move(request); } @@ -251,6 +348,21 @@ namespace dawn_native { namespace d3d12 { stream << " fxcVersion=" << fxcVersion; stream << " dxcVersion=" << dxcVersion; stream << " hasShaderFloat16Feature=" << hasShaderFloat16Feature; + + stream << " overridableConstants={"; + for (const auto& pipelineConstant : *pipelineConstantEntries) { + const std::string& name = pipelineConstant.first; + double value = pipelineConstant.second; + + // This is already validated so `name` must exist + const auto& moduleConstant = shaderEntryPointConstants->at(name); + + stream << " <" << name << "," + << GetHLSLValueString(moduleConstant.type, nullptr, value) + << ">"; + } + stream << " }"; + stream << ")"; stream << "\n"; @@ -310,6 +422,53 @@ namespace dawn_native { namespace d3d12 { return arguments; } + template + const std::vector> GetOverridableConstantsDefines( + const PipelineConstantEntries* pipelineConstantEntries, + const EntryPointMetadata::OverridableConstantsMap* shaderEntryPointConstants) { + std::vector> defineStrings; + + std::unordered_set overriddenConstants; + + // Set pipeline overridden values + for (const auto& pipelineConstant : *pipelineConstantEntries) { + const std::string& name = pipelineConstant.first; + double value = pipelineConstant.second; + + overriddenConstants.insert(name); + + // This is already validated so `name` must exist + const auto& moduleConstant = shaderEntryPointConstants->at(name); + + defineStrings.emplace_back( + NumberToString::kSpecConstantPrefix + + NumberToString::ToStringAsId( + static_cast(moduleConstant.id)), + GetHLSLValueString(moduleConstant.type, nullptr, value)); + } + + // Set shader initialized default values + for (const auto& iter : *shaderEntryPointConstants) { + const std::string& name = iter.first; + if (overriddenConstants.count(name) != 0) { + // This constant already has overridden value + continue; + } + + const auto& moduleConstant = shaderEntryPointConstants->at(name); + + // Uninitialized default values are okay since they are only defined to pass + // compilation but not used + defineStrings.emplace_back(NumberToString::kSpecConstantPrefix + + NumberToString::ToStringAsId( + static_cast(moduleConstant.id)), + GetHLSLValueString( + moduleConstant.type, &moduleConstant.defaultValue)); + } + + return defineStrings; + } + ResultOrError> CompileShaderDXC(IDxcLibrary* dxcLibrary, IDxcCompiler* dxcCompiler, const ShaderCompilationRequest& request, @@ -326,12 +485,21 @@ namespace dawn_native { namespace d3d12 { std::vector arguments = GetDXCArguments(request.compileFlags, request.hasShaderFloat16Feature); + // Build defines for overridable constants + const auto& defineStrings = GetOverridableConstantsDefines( + request.pipelineConstantEntries, request.shaderEntryPointConstants); + std::vector dxcDefines; + dxcDefines.reserve(defineStrings.size()); + for (const auto& d : defineStrings) { + dxcDefines.push_back({d.first.c_str(), d.second.c_str()}); + } + ComPtr result; DAWN_TRY(CheckHRESULT( dxcCompiler->Compile(sourceBlob.Get(), nullptr, entryPointW.c_str(), request.deviceInfo->shaderProfiles[request.stage].c_str(), - arguments.data(), arguments.size(), nullptr, 0, nullptr, - &result), + arguments.data(), arguments.size(), dxcDefines.data(), + dxcDefines.size(), nullptr, &result), "DXC compile")); HRESULT hr; @@ -369,8 +537,24 @@ namespace dawn_native { namespace d3d12 { ComPtr compiledShader; ComPtr errors; + // Build defines for overridable constants + const auto& defineStrings = GetOverridableConstantsDefines( + request.pipelineConstantEntries, request.shaderEntryPointConstants); + + const D3D_SHADER_MACRO* pDefines = nullptr; + std::vector fxcDefines; + if (defineStrings.size() > 0) { + fxcDefines.reserve(defineStrings.size() + 1); + for (const auto& d : defineStrings) { + fxcDefines.push_back({d.first.c_str(), d.second.c_str()}); + } + // d3dCompile D3D_SHADER_MACRO* pDefines is a nullptr terminated array + fxcDefines.push_back({nullptr, nullptr}); + pDefines = fxcDefines.data(); + } + DAWN_INVALID_IF(FAILED(functions->d3dCompile( - hlslSource.c_str(), hlslSource.length(), nullptr, nullptr, nullptr, + hlslSource.c_str(), hlslSource.length(), nullptr, pDefines, nullptr, request.entryPointName, targetProfile, request.compileFlags, 0, &compiledShader, &errors)), "D3D compile failed with: %s", @@ -392,6 +576,9 @@ namespace dawn_native { namespace d3d12 { } transformManager.Add(); + transformManager.Add(); + transformInputs.Add(request.entryPointName); + transformManager.Add(); if (request.disableSymbolRenaming) { @@ -508,7 +695,7 @@ namespace dawn_native { namespace d3d12 { return InitializeBase(parseResult); } - ResultOrError ShaderModule::Compile(const char* entryPointName, + ResultOrError ShaderModule::Compile(const ProgrammableStage& programmableStage, SingleShaderStage stage, PipelineLayout* layout, uint32_t compileFlags) { @@ -555,9 +742,10 @@ namespace dawn_native { namespace d3d12 { } ShaderCompilationRequest request; - DAWN_TRY_ASSIGN(request, ShaderCompilationRequest::Create(entryPointName, stage, layout, - compileFlags, device, program, - GetEntryPoint(entryPointName))); + DAWN_TRY_ASSIGN( + request, ShaderCompilationRequest::Create( + programmableStage.entryPoint.c_str(), stage, layout, compileFlags, device, + program, GetEntryPoint(programmableStage.entryPoint), programmableStage)); PersistentCacheKey shaderCacheKey; DAWN_TRY_ASSIGN(shaderCacheKey, request.CreateCacheKey()); diff --git a/src/dawn_native/d3d12/ShaderModuleD3D12.h b/src/dawn_native/d3d12/ShaderModuleD3D12.h index 880a35cf45..02c1f8cea0 100644 --- a/src/dawn_native/d3d12/ShaderModuleD3D12.h +++ b/src/dawn_native/d3d12/ShaderModuleD3D12.h @@ -20,6 +20,10 @@ #include "dawn_native/d3d12/d3d12_platform.h" +namespace dawn_native { + struct ProgrammableStage; +} // namespace dawn_native + namespace dawn_native { namespace d3d12 { class Device; @@ -49,7 +53,7 @@ namespace dawn_native { namespace d3d12 { const ShaderModuleDescriptor* descriptor, ShaderModuleParseResult* parseResult); - ResultOrError Compile(const char* entryPointName, + ResultOrError Compile(const ProgrammableStage& programmableStage, SingleShaderStage stage, PipelineLayout* layout, uint32_t compileFlags); diff --git a/src/tests/end2end/ShaderTests.cpp b/src/tests/end2end/ShaderTests.cpp index 19865e01dd..2902a35799 100644 --- a/src/tests/end2end/ShaderTests.cpp +++ b/src/tests/end2end/ShaderTests.cpp @@ -391,8 +391,8 @@ fn main([[location(0)]] pos : vec4) -> [[builtin(position)]] vec4 { // Test overridable constants without numeric identifiers TEST_P(ShaderTests, OverridableConstants) { - // TODO(dawn:1137): D3D12 backend is unimplemented - DAWN_TEST_UNSUPPORTED_IF(!IsVulkan() && !IsMetal()); + DAWN_TEST_UNSUPPORTED_IF(IsOpenGL()); + DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES()); uint32_t const kCount = 11; std::vector expected(kCount); @@ -469,8 +469,8 @@ TEST_P(ShaderTests, OverridableConstants) { // Test overridable constants with numeric identifiers TEST_P(ShaderTests, OverridableConstantsNumericIdentifiers) { - // TODO(dawn:1137): D3D12 backend is unimplemented - DAWN_TEST_UNSUPPORTED_IF(!IsVulkan() && !IsMetal()); + DAWN_TEST_UNSUPPORTED_IF(IsOpenGL()); + DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES()); uint32_t const kCount = 4; std::vector expected{1u, 2u, 3u, 0u}; @@ -523,17 +523,71 @@ TEST_P(ShaderTests, OverridableConstantsNumericIdentifiers) { EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), buffer, 0, kCount); } +// Test overridable constants precision +// D3D12 HLSL shader uses defines so we want float number to have enough precision +TEST_P(ShaderTests, OverridableConstantsPrecision) { + DAWN_TEST_UNSUPPORTED_IF(IsOpenGL()); + DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES()); + + uint32_t const kCount = 2; + float const kValue1 = 3.14159; + float const kValue2 = 3.141592653589793238; + std::vector expected{kValue1, kValue2}; + wgpu::Buffer buffer = CreateBuffer(kCount); + + std::string shader = R"( +[[override(1001)]] let c1: f32; +[[override(1002)]] let c2: f32; + +[[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; +})"; + + std::vector constants; + constants.push_back({nullptr, "1001", kValue1}); + constants.push_back({nullptr, "1002", kValue2}); + 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_FLOAT_RANGE_EQ(expected.data(), buffer, 0, kCount); +} + // Test overridable constants for different entry points TEST_P(ShaderTests, OverridableConstantsMultipleEntryPoints) { - // TODO(dawn:1137): D3D12 backend is unimplemented - DAWN_TEST_UNSUPPORTED_IF(!IsVulkan() && !IsMetal()); + DAWN_TEST_UNSUPPORTED_IF(IsOpenGL()); + DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES()); uint32_t const kCount = 1; std::vector expected1{1u}; std::vector expected2{2u}; + std::vector expected3{3u}; wgpu::Buffer buffer1 = CreateBuffer(kCount); wgpu::Buffer buffer2 = CreateBuffer(kCount); + wgpu::Buffer buffer3 = CreateBuffer(kCount); std::string shader = R"( [[override(1001)]] let c1: u32; @@ -552,6 +606,10 @@ TEST_P(ShaderTests, OverridableConstantsMultipleEntryPoints) { [[stage(compute), workgroup_size(1)]] fn main2() { buf.data[0] = c2; } + +[[stage(compute), workgroup_size(1)]] fn main3() { + buf.data[0] = 3u; +} )"; std::vector constants1; @@ -575,10 +633,17 @@ TEST_P(ShaderTests, OverridableConstantsMultipleEntryPoints) { csDesc2.compute.constantCount = constants2.size(); wgpu::ComputePipeline pipeline2 = device.CreateComputePipeline(&csDesc2); + wgpu::ComputePipelineDescriptor csDesc3; + csDesc3.compute.module = shaderModule; + csDesc3.compute.entryPoint = "main3"; + wgpu::ComputePipeline pipeline3 = device.CreateComputePipeline(&csDesc3); + wgpu::BindGroup bindGroup1 = utils::MakeBindGroup(device, pipeline1.GetBindGroupLayout(0), {{0, buffer1}}); wgpu::BindGroup bindGroup2 = utils::MakeBindGroup(device, pipeline2.GetBindGroupLayout(0), {{0, buffer2}}); + wgpu::BindGroup bindGroup3 = + utils::MakeBindGroup(device, pipeline3.GetBindGroupLayout(0), {{0, buffer3}}); wgpu::CommandBuffer commands; { @@ -592,6 +657,10 @@ TEST_P(ShaderTests, OverridableConstantsMultipleEntryPoints) { pass.SetBindGroup(0, bindGroup2); pass.Dispatch(1); + pass.SetPipeline(pipeline3); + pass.SetBindGroup(0, bindGroup3); + pass.Dispatch(1); + pass.EndPass(); commands = encoder.Finish(); @@ -601,14 +670,15 @@ TEST_P(ShaderTests, OverridableConstantsMultipleEntryPoints) { EXPECT_BUFFER_U32_RANGE_EQ(expected1.data(), buffer1, 0, kCount); EXPECT_BUFFER_U32_RANGE_EQ(expected2.data(), buffer2, 0, kCount); + EXPECT_BUFFER_U32_RANGE_EQ(expected3.data(), buffer3, 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:1137): D3D12 backend is unimplemented - DAWN_TEST_UNSUPPORTED_IF(!IsVulkan() && !IsMetal()); + DAWN_TEST_UNSUPPORTED_IF(IsOpenGL()); + DAWN_TEST_UNSUPPORTED_IF(IsOpenGLES()); wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"( [[override(1111)]] let xright: f32; diff --git a/src/tests/unittests/validation/OverridableConstantsValidationTests.cpp b/src/tests/unittests/validation/OverridableConstantsValidationTests.cpp index af43f4cf86..f8c5337be5 100644 --- a/src/tests/unittests/validation/OverridableConstantsValidationTests.cpp +++ b/src/tests/unittests/validation/OverridableConstantsValidationTests.cpp @@ -105,12 +105,12 @@ TEST_F(ComputePipelineOverridableConstantsValidationTest, ConstantsIdentifierLoo TestCreatePipeline(constants); } { - // Valid: set the same constant twice + // Error: set the same constant twice std::vector constants{ {nullptr, "c0", 0}, {nullptr, "c0", 1}, }; - TestCreatePipeline(constants); + ASSERT_DEVICE_ERROR(TestCreatePipeline(constants)); } { // Valid: find by constant numeric id @@ -158,12 +158,12 @@ TEST_F(ComputePipelineOverridableConstantsValidationTest, UninitializedConstants TestCreatePipeline(constants); } { - // Valid: all constants initialized (with duplicate initializations) + // Error: duplicate initializations std::vector constants{ {nullptr, "c0", false}, {nullptr, "c2", 1}, {nullptr, "c5", 1}, {nullptr, "c8", 1}, {nullptr, "c2", 2}, }; - TestCreatePipeline(constants); + ASSERT_DEVICE_ERROR(TestCreatePipeline(constants)); } }