Revert "[dawn] Switch indirect dispatch validation to use immediates" This reverts commit 104068a1a6c0bacad1e9896f18a7c558f82f5444. Reason for revert: Breaks without unsafe webgpu flags Bug: 493976342 Original change's description: > [dawn] Switch indirect dispatch validation to use immediates > > Change the indirect dispatch validation shader to use immediates instead > of a uniform buffer. > > Fix: 488346117 > Change-Id: Ib0804f399bd387bcd93b5540490b566d57ff88d4 > Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/294716 > Commit-Queue: Corentin Wallez <cwallez@chromium.org> > Reviewed-by: Corentin Wallez <cwallez@chromium.org> # Not skipping CQ checks because original CL landed > 1 day ago. Change-Id: If6b86bb60eb681f1cdfb61ef3bf97c3c021ffb7a Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/298038 Commit-Queue: Alan Baker <alanbaker@google.com> Auto-Submit: Alan Baker <alanbaker@google.com> Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/dawn/native/ComputePassEncoder.cpp b/src/dawn/native/ComputePassEncoder.cpp index 76b60a1..cc3a2e7 100644 --- a/src/dawn/native/ComputePassEncoder.cpp +++ b/src/dawn/native/ComputePassEncoder.cpp
@@ -51,18 +51,6 @@ namespace { -// Neither 'enableValidation' nor 'duplicateNumWorkgroups' can be declared as 'bool' as -// currently in WGSL type 'bool' cannot be used in address space 'uniform' as 'it is -// non-host-shareable'. -struct IndirectDispatchParams { - uint32_t maxComputeWorkgroupsPerDimension; - uint32_t clientOffsetInU32; - uint32_t enableValidation; - uint32_t duplicateNumWorkgroups; - uint32_t linearIndexing; - uint32_t overflowValue; -}; - ResultOrError<ComputePipelineBase*> GetOrCreateIndirectDispatchValidationPipeline( DeviceBase* device) { InternalPipelineStore* store = device->GetInternalPipelineStore(); @@ -71,12 +59,13 @@ return store->dispatchIndirectValidationPipeline.Get(); } + // TODO(https://crbug.com/dawn/488346117): Use immediates instead of uniform. // TODO(https://crbug.com/dawn/1108): Propagate validation feedback from this // shader in various failure modes. // Type 'bool' cannot be used in address space 'uniform' as it is non-host-shareable. Ref<ShaderModuleBase> shaderModule; DAWN_TRY_ASSIGN(shaderModule, utils::CreateShaderModule(device, DAWN_MULTILINE( - struct Params { + struct UniformParams { maxComputeWorkgroupsPerDimension: u32, clientOffsetInU32: u32, enableValidation: u32, @@ -93,23 +82,23 @@ data: array<u32> } - var<immediate> params: Params; - @group(0) @binding(0) var<storage, read_write> clientParams: IndirectParams; - @group(0) @binding(1) var<storage, read_write> validatedParams: ValidatedParams; + @group(0) @binding(0) var<uniform> uniformParams: UniformParams; + @group(0) @binding(1) var<storage, read_write> clientParams: IndirectParams; + @group(0) @binding(2) var<storage, read_write> validatedParams: ValidatedParams; @compute @workgroup_size(1, 1, 1) fn main() { - var workgroups = vec3u(clientParams.data[params.clientOffsetInU32 + 0], - clientParams.data[params.clientOffsetInU32 + 1], - clientParams.data[params.clientOffsetInU32 + 2]); - if (params.enableValidation > 0u) { + var workgroups = vec3u(clientParams.data[uniformParams.clientOffsetInU32 + 0], + clientParams.data[uniformParams.clientOffsetInU32 + 1], + clientParams.data[uniformParams.clientOffsetInU32 + 2]); + if (uniformParams.enableValidation > 0u) { var invalid = false; - if (max(workgroups.x, max(workgroups.y, workgroups.z)) > params.maxComputeWorkgroupsPerDimension) { + if (max(workgroups.x, max(workgroups.y, workgroups.z)) > uniformParams.maxComputeWorkgroupsPerDimension) { invalid = true; - } else if (params.linearIndexing > 0u) { - invalid |= workgroups.x > (params.overflowValue / workgroups.y); + } else if (uniformParams.linearIndexing > 0u) { + invalid |= workgroups.x > (uniformParams.overflowValue / workgroups.y); let xy = workgroups.x * workgroups.y; - invalid |= xy > (params.overflowValue / workgroups.z); + invalid |= xy > (uniformParams.overflowValue / workgroups.z); } if (invalid) { @@ -119,7 +108,7 @@ validatedParams.data[0] = workgroups.x; validatedParams.data[1] = workgroups.y; validatedParams.data[2] = workgroups.z; - if (params.duplicateNumWorkgroups > 0u) { + if (uniformParams.duplicateNumWorkgroups > 0u) { validatedParams.data[3] = workgroups.x; validatedParams.data[4] = workgroups.y; validatedParams.data[5] = workgroups.z; @@ -132,14 +121,14 @@ utils::MakeBindGroupLayout( device, { - {0, wgpu::ShaderStage::Compute, kInternalStorageBufferBinding}, - {1, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage}, + {0, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Uniform}, + {1, wgpu::ShaderStage::Compute, kInternalStorageBufferBinding}, + {2, wgpu::ShaderStage::Compute, wgpu::BufferBindingType::Storage}, }, /* allowInternalBinding */ true)); Ref<PipelineLayoutBase> pipelineLayout; - DAWN_TRY_ASSIGN(pipelineLayout, utils::MakeBasicPipelineLayout(device, bindGroupLayout, - sizeof(IndirectDispatchParams))); + DAWN_TRY_ASSIGN(pipelineLayout, utils::MakeBasicPipelineLayout(device, bindGroupLayout)); ComputePipelineDescriptor computePipelineDescriptor = {}; computePipelineDescriptor.layout = pipelineLayout.Get(); @@ -363,15 +352,33 @@ const uint64_t clientIndirectBindingSize = kDispatchIndirectSize + clientOffsetFromAlignedBoundary; - // Set the immediate params. - IndirectDispatchParams params; - params.maxComputeWorkgroupsPerDimension = - device->GetLimits().v1.maxComputeWorkgroupsPerDimension; - params.clientOffsetInU32 = clientOffsetFromAlignedBoundary / sizeof(uint32_t); - params.enableValidation = static_cast<uint32_t>(IsValidationEnabled()); - params.duplicateNumWorkgroups = static_cast<uint32_t>(shouldDuplicateNumWorkgroups); - params.linearIndexing = static_cast<uint32_t>(usesLinearIndexing); - params.overflowValue = overflowValue; + // Neither 'enableValidation' nor 'duplicateNumWorkgroups' can be declared as 'bool' as + // currently in WGSL type 'bool' cannot be used in address space 'uniform' as 'it is + // non-host-shareable'. + struct UniformParams { + uint32_t maxComputeWorkgroupsPerDimension; + uint32_t clientOffsetInU32; + uint32_t enableValidation; + uint32_t duplicateNumWorkgroups; + uint32_t linearIndexing; + uint32_t overflowValue; + }; + + // Create a uniform buffer to hold parameters for the shader. + Ref<BufferBase> uniformBuffer; + { + UniformParams params; + params.maxComputeWorkgroupsPerDimension = + device->GetLimits().v1.maxComputeWorkgroupsPerDimension; + params.clientOffsetInU32 = clientOffsetFromAlignedBoundary / sizeof(uint32_t); + params.enableValidation = static_cast<uint32_t>(IsValidationEnabled()); + params.duplicateNumWorkgroups = static_cast<uint32_t>(shouldDuplicateNumWorkgroups); + params.linearIndexing = static_cast<uint32_t>(usesLinearIndexing); + params.overflowValue = overflowValue; + + DAWN_TRY_ASSIGN(uniformBuffer, + utils::CreateBufferFromData(device, wgpu::BufferUsage::Uniform, {params})); + } // Reserve space in the scratch buffer to hold the validated indirect params. ScratchBuffer& scratchBuffer = store->scratchIndirectStorage; @@ -385,15 +392,15 @@ DAWN_TRY_ASSIGN(validationBindGroup, utils::MakeBindGroup(device, layout, { - {0, indirectBuffer, clientIndirectBindingOffset, + {0, uniformBuffer}, + {1, indirectBuffer, clientIndirectBindingOffset, clientIndirectBindingSize}, - {1, validatedIndirectBuffer, 0, scratchBufferSize}, + {2, validatedIndirectBuffer, 0, scratchBufferSize}, }, UsageValidationMode::Internal)); // Issue commands to validate the indirect buffer. APISetPipeline(validationPipeline.Get()); - APISetImmediates(0, ¶ms, sizeof(IndirectDispatchParams)); APISetBindGroup(0, validationBindGroup.Get()); APIDispatchWorkgroups(1);
diff --git a/src/dawn/native/utils/WGPUHelpers.cpp b/src/dawn/native/utils/WGPUHelpers.cpp index bb47ec9..4592036 100644 --- a/src/dawn/native/utils/WGPUHelpers.cpp +++ b/src/dawn/native/utils/WGPUHelpers.cpp
@@ -78,13 +78,11 @@ ResultOrError<Ref<PipelineLayoutBase>> MakeBasicPipelineLayout( DeviceBase* device, - const Ref<BindGroupLayoutBase>& bindGroupLayout, - uint32_t immediateSize) { + const Ref<BindGroupLayoutBase>& bindGroupLayout) { PipelineLayoutDescriptor descriptor; descriptor.bindGroupLayoutCount = 1; BindGroupLayoutBase* bgl = bindGroupLayout.Get(); descriptor.bindGroupLayouts = &bgl; - descriptor.immediateSize = immediateSize; return device->CreatePipelineLayout(&descriptor); }
diff --git a/src/dawn/native/utils/WGPUHelpers.h b/src/dawn/native/utils/WGPUHelpers.h index 2536348..d734467 100644 --- a/src/dawn/native/utils/WGPUHelpers.h +++ b/src/dawn/native/utils/WGPUHelpers.h
@@ -73,8 +73,7 @@ ResultOrError<Ref<PipelineLayoutBase>> MakeBasicPipelineLayout( DeviceBase* device, - const Ref<BindGroupLayoutBase>& bindGroupLayout, - uint32_t immediateSize = 0); + const Ref<BindGroupLayoutBase>& bindGroupLayout); // Helpers to make creating bind group layouts look nicer: //
diff --git a/src/dawn/tests/unittests/native/CommandBufferEncodingTests.cpp b/src/dawn/tests/unittests/native/CommandBufferEncodingTests.cpp index 2ccd330..1e68eaf 100644 --- a/src/dawn/tests/unittests/native/CommandBufferEncodingTests.cpp +++ b/src/dawn/tests/unittests/native/CommandBufferEncodingTests.cpp
@@ -213,13 +213,6 @@ } }; - auto ExpectSetValidationImmediates = [&](CommandIterator* commands) { - auto* cmd = commands->NextCommand<SetImmediatesCmd>(); - commands->NextData<uint8_t>(cmd->size); - ASSERT_EQ(cmd->offset, 0u); - ASSERT_EQ(cmd->size, 6 * sizeof(uint32_t)); - }; - auto ExpectSetValidationBindGroup = [&](CommandIterator* commands) { auto* cmd = commands->NextCommand<SetBindGroupCmd>(); ASSERT_EQ(cmd->index, BindGroupIndex(0)); @@ -246,7 +239,6 @@ // Expect the validation. {Command::SetComputePipeline, ExpectSetValidationPipeline}, - {Command::SetImmediates, ExpectSetValidationImmediates}, {Command::SetBindGroup, ExpectSetValidationBindGroup}, {Command::Dispatch, ExpectSetValidationDispatch}, @@ -260,7 +252,6 @@ // Expect the validation. {Command::SetComputePipeline, ExpectSetValidationPipeline}, - {Command::SetImmediates, ExpectSetValidationImmediates}, {Command::SetBindGroup, ExpectSetValidationBindGroup}, {Command::Dispatch, ExpectSetValidationDispatch}, @@ -279,7 +270,6 @@ // Expect the validation. {Command::SetComputePipeline, ExpectSetValidationPipeline}, - {Command::SetImmediates, ExpectSetValidationImmediates}, {Command::SetBindGroup, ExpectSetValidationBindGroup}, {Command::Dispatch, ExpectSetValidationDispatch},