From 84532462f62cf2ae98d9a3080c8e8cb371fa36fb Mon Sep 17 00:00:00 2001 From: shrekshao Date: Wed, 1 Mar 2023 21:04:28 +0000 Subject: [PATCH] Validate writable storage texture bindings don't alias Followup of storage buffer bindings aliasing validation. Bug: dawn:1642 Change-Id: I84bf33895320053630ed80d3503ff53d1eaa83b9 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/121420 Kokoro: Kokoro Reviewed-by: Austin Eng Commit-Queue: Shrek Shao --- src/dawn/common/Numeric.h | 18 +- src/dawn/native/CommandBufferStateTracker.cpp | 203 +++++-- src/dawn/native/CommandValidation.cpp | 13 +- src/dawn/native/Subresource.h | 2 +- src/dawn/tests/BUILD.gn | 2 + src/dawn/tests/unittests/NumericTests.cpp | 59 ++ .../validation/ResourceUsageTrackingTests.cpp | 12 +- ...leBufferBindingAliasingValidationTests.cpp | 20 - ...eTextureBindingAliasingValidationTests.cpp | 532 ++++++++++++++++++ webgpu-cts/expectations.txt | 5 + 10 files changed, 793 insertions(+), 73 deletions(-) create mode 100644 src/dawn/tests/unittests/NumericTests.cpp create mode 100644 src/dawn/tests/unittests/validation/WritableTextureBindingAliasingValidationTests.cpp diff --git a/src/dawn/common/Numeric.h b/src/dawn/common/Numeric.h index 6387f48185..7be33c87f9 100644 --- a/src/dawn/common/Numeric.h +++ b/src/dawn/common/Numeric.h @@ -61,7 +61,23 @@ bool IsDoubleValueRepresentable(double value) { constexpr double kMax = static_cast(std::numeric_limits::max()); return kLowest <= value && value <= kMax; } else { - static_assert(sizeof(T) != sizeof(T), "Unsupported type"); + static_assert(std::is_same_v || std::is_integral_v, "Unsupported type"); + } +} + +// Returns if two inclusive integral ranges [x0, x1] and [y0, y1] have overlap. +template +bool RangesOverlap(T x0, T x1, T y0, T y1) { + ASSERT(x0 <= x1 && y0 <= y1); + if constexpr (std::is_integral_v) { + // Two ranges DON'T have overlap if and only if: + // 1. [x0, x1] [y0, y1], or + // 2. [y0, y1] [x0, x1] + // which is (x1 < y0 || y1 < x0) + // The inverse of which ends in the following statement. + return x0 <= y1 && y0 <= x1; + } else { + static_assert(std::is_integral_v, "Unsupported type"); } } diff --git a/src/dawn/native/CommandBufferStateTracker.cpp b/src/dawn/native/CommandBufferStateTracker.cpp index 245f3a46a4..fc6cbc93ea 100644 --- a/src/dawn/native/CommandBufferStateTracker.cpp +++ b/src/dawn/native/CommandBufferStateTracker.cpp @@ -17,6 +17,7 @@ #include #include #include +#include #include "dawn/common/Assert.h" #include "dawn/common/BitSetIterator.h" @@ -53,7 +54,7 @@ std::optional FindFirstUndersizedBuffer( return std::nullopt; } -struct BufferBindingAliasingResult { +struct BufferAliasing { struct Entry { BindGroupIndex bindGroupIndex; BindingIndex bindingIndex; @@ -66,18 +67,42 @@ struct BufferBindingAliasingResult { Entry e1; }; -// TODO(dawn:1642): Find storage texture binding aliasing as well. +struct TextureAliasing { + struct Entry { + BindGroupIndex bindGroupIndex; + BindingIndex bindingIndex; + + uint32_t baseMipLevel; + uint32_t mipLevelCount; + uint32_t baseArrayLayer; + uint32_t arrayLayerCount; + }; + Entry e0; + Entry e1; +}; + +using WritableBindingAliasingResult = std::variant; + template Return FindStorageBufferBindingAliasing( const PipelineLayoutBase* pipelineLayout, const ityp::array& bindGroups, - const ityp::array, kMaxBindGroups> dynamicOffsets) { + const ityp::array, kMaxBindGroups>& dynamicOffsets) { + // If true, returns detailed validation error info. Otherwise simply returns if any binding + // aliasing is found. + constexpr bool kProduceDetails = std::is_same_v; + // Reduce the bindings array first to only preserve storage buffer bindings that could // potentially have ranges overlap. - // There can at most be 8 storage buffer bindings per shader stage. - StackVector bindingsToCheck; + // There can at most be 8 storage buffer bindings (in default limits) per shader stage. + StackVector storageBufferBindingsToCheck; + StackVector, 8> bufferBindingIndices; - StackVector, 8> bindingIndices; + // Reduce the bindings array first to only preserve writable storage texture bindings that could + // potentially have ranges overlap. + // There can at most be 8 storage texture bindings (in default limits) per shader stage. + StackVector storageTextureViewsToCheck; + StackVector, 8> textureBindingIndices; for (BindGroupIndex groupIndex : IterateBitSet(pipelineLayout->GetBindGroupLayoutsMask())) { BindGroupLayoutBase* bgl = bindGroups[groupIndex]->GetLayout(); @@ -107,55 +132,127 @@ Return FindStorageBufferBindingAliasing( adjustedOffset += dynamicOffsets[groupIndex][static_cast(bindingIndex)]; } - bindingsToCheck->push_back(BufferBinding{ + storageBufferBindingsToCheck->push_back(BufferBinding{ bufferBinding.buffer, adjustedOffset, bufferBinding.size, }); - if constexpr (std::is_same_v>) { - bindingIndices->emplace_back(groupIndex, bindingIndex); + if constexpr (kProduceDetails) { + bufferBindingIndices->emplace_back(groupIndex, bindingIndex); + } + } + + // TODO(dawn:1642): optimize: precompute start/end range of storage textures bindings. + for (BindingIndex bindingIndex{bgl->GetBufferCount()}; + bindingIndex < bgl->GetBindingCount(); ++bindingIndex) { + const BindingInfo& bindingInfo = bgl->GetBindingInfo(bindingIndex); + + if (bindingInfo.bindingType != BindingInfoType::StorageTexture) { + continue; + } + + switch (bindingInfo.storageTexture.access) { + case wgpu::StorageTextureAccess::WriteOnly: + break; + // Continue for other StorageTextureAccess type when we have any. + default: + UNREACHABLE(); + } + + const TextureViewBase* textureView = + bindGroups[groupIndex]->GetBindingAsTextureView(bindingIndex); + + storageTextureViewsToCheck->push_back(textureView); + + if constexpr (kProduceDetails) { + textureBindingIndices->emplace_back(groupIndex, bindingIndex); } } } - // Iterate through each bindings to find if any writable storage bindings aliasing exists. - // Given that maxStorageBuffersPerShaderStage is 8, - // it doesn't seem too bad to do a nested loop check. + // Iterate through each buffer bindings to find if any writable storage bindings aliasing + // exists. Given that maxStorageBuffersPerShaderStage is 8, it doesn't seem too bad to do a + // nested loop check. // TODO(dawn:1642): Maybe do algorithm optimization from O(N^2) to O(N*logN). - for (size_t i = 0; i < bindingsToCheck->size(); i++) { - const auto& bufferBinding0 = bindingsToCheck[i]; + for (size_t i = 0; i < storageBufferBindingsToCheck->size(); i++) { + const auto& bufferBinding0 = storageBufferBindingsToCheck[i]; - for (size_t j = i + 1; j < bindingsToCheck->size(); j++) { - const auto& bufferBinding1 = bindingsToCheck[j]; + for (size_t j = i + 1; j < storageBufferBindingsToCheck->size(); j++) { + const auto& bufferBinding1 = storageBufferBindingsToCheck[j]; if (bufferBinding0.buffer != bufferBinding1.buffer) { continue; } - if (bufferBinding0.offset <= bufferBinding1.offset + bufferBinding1.size - 1 && - bufferBinding1.offset <= bufferBinding0.offset + bufferBinding0.size - 1) { - if constexpr (std::is_same_v) { + if (RangesOverlap( + bufferBinding0.offset, bufferBinding0.offset + bufferBinding0.size - 1, + bufferBinding1.offset, bufferBinding1.offset + bufferBinding1.size - 1)) { + if constexpr (kProduceDetails) { + return WritableBindingAliasingResult{BufferAliasing{ + {bufferBindingIndices[i].first, bufferBindingIndices[i].second, + bufferBinding0.offset, bufferBinding0.size}, + {bufferBindingIndices[j].first, bufferBindingIndices[j].second, + bufferBinding1.offset, bufferBinding1.size}, + }}; + } else { return true; - } else if constexpr (std::is_same_v>) { - return BufferBindingAliasingResult{ - {bindingIndices[i].first, bindingIndices[i].second, bufferBinding0.offset, - bufferBinding0.size}, - {bindingIndices[j].first, bindingIndices[j].second, bufferBinding1.offset, - bufferBinding1.size}, - }; } } } } - if constexpr (std::is_same_v) { - return false; - } else if constexpr (std::is_same_v>) { - return std::nullopt; + // Iterate through each texture views to find if any writable storage bindings aliasing exists. + // Given that maxStorageTexturesPerShaderStage is 8, + // it doesn't seem too bad to do a nested loop check. + // TODO(dawn:1642): Maybe do algorithm optimization from O(N^2) to O(N*logN). + for (size_t i = 0; i < storageTextureViewsToCheck->size(); i++) { + const TextureViewBase* textureView0 = storageTextureViewsToCheck[i]; + + ASSERT(textureView0->GetAspects() == Aspect::Color); + + uint32_t baseMipLevel0 = textureView0->GetBaseMipLevel(); + uint32_t mipLevelCount0 = textureView0->GetLevelCount(); + uint32_t baseArrayLayer0 = textureView0->GetBaseArrayLayer(); + uint32_t arrayLayerCount0 = textureView0->GetLayerCount(); + + for (size_t j = i + 1; j < storageTextureViewsToCheck->size(); j++) { + const TextureViewBase* textureView1 = storageTextureViewsToCheck[j]; + + if (textureView0->GetTexture() != textureView1->GetTexture()) { + continue; + } + + ASSERT(textureView1->GetAspects() == Aspect::Color); + + uint32_t baseMipLevel1 = textureView1->GetBaseMipLevel(); + uint32_t mipLevelCount1 = textureView1->GetLevelCount(); + uint32_t baseArrayLayer1 = textureView1->GetBaseArrayLayer(); + uint32_t arrayLayerCount1 = textureView1->GetLayerCount(); + + if (RangesOverlap(baseMipLevel0, baseMipLevel0 + mipLevelCount0 - 1, baseMipLevel1, + baseMipLevel1 + mipLevelCount1 - 1) && + RangesOverlap(baseArrayLayer0, baseArrayLayer0 + arrayLayerCount0 - 1, + baseArrayLayer1, baseArrayLayer1 + arrayLayerCount1 - 1)) { + if constexpr (kProduceDetails) { + return WritableBindingAliasingResult{TextureAliasing{ + {textureBindingIndices[i].first, textureBindingIndices[i].second, + baseMipLevel0, mipLevelCount0, baseArrayLayer0, arrayLayerCount0}, + {textureBindingIndices[j].first, textureBindingIndices[j].second, + baseMipLevel1, mipLevelCount1, baseArrayLayer1, arrayLayerCount1}, + }}; + } else { + return true; + } + } + } + } + + if constexpr (kProduceDetails) { + return WritableBindingAliasingResult(); + } else { + return false; } - UNREACHABLE(); } } // namespace @@ -396,7 +493,7 @@ MaybeError CommandBufferStateTracker::CheckMissingAspects(ValidationAspects aspe DAWN_INVALID_IF(aspects[VALIDATION_ASPECT_PIPELINE], "No pipeline set."); - if (DAWN_UNLIKELY(aspects[VALIDATION_ASPECT_INDEX_BUFFER])) { + if (aspects[VALIDATION_ASPECT_INDEX_BUFFER]) { DAWN_INVALID_IF(!mIndexBufferSet, "Index buffer was not set."); RenderPipelineBase* lastRenderPipeline = GetRenderPipeline(); @@ -436,7 +533,7 @@ MaybeError CommandBufferStateTracker::CheckMissingAspects(ValidationAspects aspe uint8_t(firstMissing), GetRenderPipeline()); } - if (DAWN_UNLIKELY(aspects[VALIDATION_ASPECT_BIND_GROUPS])) { + if (aspects[VALIDATION_ASPECT_BIND_GROUPS]) { for (BindGroupIndex i : IterateBitSet(mLastPipelineLayout->GetBindGroupLayoutsMask())) { ASSERT(HasPipeline()); @@ -495,19 +592,39 @@ MaybeError CommandBufferStateTracker::CheckMissingAspects(ValidationAspects aspe } } - auto result = FindStorageBufferBindingAliasing>( + auto result = FindStorageBufferBindingAliasing( mLastPipelineLayout, mBindgroups, mDynamicOffsets); - if (result) { + if (std::holds_alternative(result)) { + const auto& a = std::get(result); return DAWN_VALIDATION_ERROR( - "Writable storage buffer binding found between bind group index %u, binding index " - "%u, and bind group index %u, binding index %u, with overlapping ranges (offset: " + "Writable storage buffer binding aliasing found between bind group index %u, " + "binding index " + "%u, and bind group index %u, binding index %u, with overlapping ranges " + "(offset: " "%u, size: %u) and (offset: %u, size: %u).", - static_cast(result->e0.bindGroupIndex), - static_cast(result->e0.bindingIndex), - static_cast(result->e1.bindGroupIndex), - static_cast(result->e1.bindingIndex), result->e0.offset, result->e0.size, - result->e1.offset, result->e1.size); + static_cast(a.e0.bindGroupIndex), + static_cast(a.e0.bindingIndex), + static_cast(a.e1.bindGroupIndex), + static_cast(a.e1.bindingIndex), a.e0.offset, a.e0.size, a.e1.offset, + a.e1.size); + } else { + ASSERT(std::holds_alternative(result)); + const auto& a = std::get(result); + return DAWN_VALIDATION_ERROR( + "Writable storage texture binding aliasing found between bind group " + "index %u, binding index " + "%u, and bind group index %u, binding index %u, with subresources " + "(base mipmap level: " + "%u, mip level count: %u, base array layer: %u, array layer count: %u) and " + "(base mipmap level: %u, mip level count: " + "%u, base array layer: %u, array layer count: %u).", + static_cast(a.e0.bindGroupIndex), + static_cast(a.e0.bindingIndex), + static_cast(a.e1.bindGroupIndex), + static_cast(a.e1.bindingIndex), a.e0.baseMipLevel, a.e0.mipLevelCount, + a.e0.baseArrayLayer, a.e0.arrayLayerCount, a.e1.baseMipLevel, a.e1.mipLevelCount, + a.e1.baseArrayLayer, a.e1.arrayLayerCount); } // The chunk of code above should be similar to the one in |RecomputeLazyAspects|. diff --git a/src/dawn/native/CommandValidation.cpp b/src/dawn/native/CommandValidation.cpp index e7874e6635..d328f82cc1 100644 --- a/src/dawn/native/CommandValidation.cpp +++ b/src/dawn/native/CommandValidation.cpp @@ -21,6 +21,7 @@ #include #include "dawn/common/BitSetIterator.h" +#include "dawn/common/Numeric.h" #include "dawn/native/Adapter.h" #include "dawn/native/BindGroup.h" #include "dawn/native/Buffer.h" @@ -115,10 +116,14 @@ MaybeError ValidateWriteBuffer(const DeviceBase* device, } bool IsRangeOverlapped(uint32_t startA, uint32_t startB, uint32_t length) { - uint32_t maxStart = std::max(startA, startB); - uint32_t minStart = std::min(startA, startB); - return static_cast(minStart) + static_cast(length) > - static_cast(maxStart); + if (length < 1) { + return false; + } + return RangesOverlap( + static_cast(startA), + static_cast(startA) + static_cast(length) - 1, + static_cast(startB), + static_cast(startB) + static_cast(length) - 1); } ResultOrError ComputeRequiredBytesInCopy(const TexelBlockInfo& blockInfo, diff --git a/src/dawn/native/Subresource.h b/src/dawn/native/Subresource.h index 473631afc7..4d4e387787 100644 --- a/src/dawn/native/Subresource.h +++ b/src/dawn/native/Subresource.h @@ -21,7 +21,7 @@ namespace dawn::native { // Note: Subresource indices are computed by iterating the aspects in increasing order. -// D3D12 uses these directly, so the order much match D3D12's indices. +// D3D12 uses these directly, so the order must match D3D12's indices. // - Depth/Stencil textures have Depth as Plane 0, and Stencil as Plane 1. enum class Aspect : uint8_t { None = 0x0, diff --git a/src/dawn/tests/BUILD.gn b/src/dawn/tests/BUILD.gn index a87fc1e953..bfeb26a930 100644 --- a/src/dawn/tests/BUILD.gn +++ b/src/dawn/tests/BUILD.gn @@ -292,6 +292,7 @@ dawn_test("dawn_unittests") { "unittests/LimitsTests.cpp", "unittests/LinkedListTests.cpp", "unittests/MathTests.cpp", + "unittests/NumericTests.cpp", "unittests/ObjectBaseTests.cpp", "unittests/PerStageTests.cpp", "unittests/PerThreadProcTests.cpp", @@ -365,6 +366,7 @@ dawn_test("dawn_unittests") { "unittests/validation/VertexStateValidationTests.cpp", "unittests/validation/VideoViewsValidationTests.cpp", "unittests/validation/WritableBufferBindingAliasingValidationTests.cpp", + "unittests/validation/WritableTextureBindingAliasingValidationTests.cpp", "unittests/validation/WriteBufferTests.cpp", "unittests/wire/WireAdapterTests.cpp", "unittests/wire/WireArgumentTests.cpp", diff --git a/src/dawn/tests/unittests/NumericTests.cpp b/src/dawn/tests/unittests/NumericTests.cpp new file mode 100644 index 0000000000..d94344aa7e --- /dev/null +++ b/src/dawn/tests/unittests/NumericTests.cpp @@ -0,0 +1,59 @@ +// Copyright 2023 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 "dawn/common/Numeric.h" +#include "gtest/gtest.h" + +// Tests for RangesOverlap +TEST(Numeric, RangesOverlap) { + // Range contains only one number + ASSERT_EQ(true, RangesOverlap(0, 0, 0, 0)); + ASSERT_EQ(false, RangesOverlap(0, 0, 1, 1)); + + // [ ] + // [ ] + ASSERT_EQ(false, RangesOverlap(0, 8, 9, 16)); + + // [ ] + // [ ] + ASSERT_EQ(false, RangesOverlap(9, 16, 0, 8)); + + // [ ] + // [ ] + ASSERT_EQ(true, RangesOverlap(2, 3, 0, 8)); + + // [ ] + // [ ] + ASSERT_EQ(true, RangesOverlap(0, 8, 2, 3)); + + // [ ] + // [ ] + ASSERT_EQ(true, RangesOverlap(0, 8, 4, 12)); + + // [ ] + // [ ] + ASSERT_EQ(true, RangesOverlap(4, 12, 0, 8)); + + // [ ] + // [ ] + ASSERT_EQ(true, RangesOverlap(0, 8, 8, 12)); + + // [ ] + // [ ] + ASSERT_EQ(true, RangesOverlap(8, 12, 0, 8)); + + // Negative numbers + ASSERT_EQ(true, RangesOverlap(-9, 12, 4, 16)); + ASSERT_EQ(false, RangesOverlap(-9, -3, -2, 0)); +} diff --git a/src/dawn/tests/unittests/validation/ResourceUsageTrackingTests.cpp b/src/dawn/tests/unittests/validation/ResourceUsageTrackingTests.cpp index 24d8e04e26..846b637909 100644 --- a/src/dawn/tests/unittests/validation/ResourceUsageTrackingTests.cpp +++ b/src/dawn/tests/unittests/validation/ResourceUsageTrackingTests.cpp @@ -973,9 +973,11 @@ TEST_F(ResourceUsageTrackingTest, TextureWithMultipleWriteUsage) { // Create a bind group to use the texture as sampled and writeonly bindings wgpu::BindGroupLayout bgl = utils::MakeBindGroupLayout( device, - {{0, wgpu::ShaderStage::Compute, wgpu::StorageTextureAccess::WriteOnly, kFormat}, - {1, wgpu::ShaderStage::Compute, wgpu::StorageTextureAccess::WriteOnly, kFormat}}); - wgpu::BindGroup bg = utils::MakeBindGroup(device, bgl, {{0, view}, {1, view}}); + {{0, wgpu::ShaderStage::Compute, wgpu::StorageTextureAccess::WriteOnly, kFormat}}); + // Create 2 bind groups with same texture subresources and dispatch twice to avoid + // storage texture binding aliasing + wgpu::BindGroup bg0 = utils::MakeBindGroup(device, bgl, {{0, view}}); + wgpu::BindGroup bg1 = utils::MakeBindGroup(device, bgl, {{0, view}}); // Create a no-op compute pipeline wgpu::ComputePipeline cp = CreateNoOpComputePipeline({bgl}); @@ -985,7 +987,9 @@ TEST_F(ResourceUsageTrackingTest, TextureWithMultipleWriteUsage) { wgpu::CommandEncoder encoder = device.CreateCommandEncoder(); wgpu::ComputePassEncoder pass = encoder.BeginComputePass(); pass.SetPipeline(cp); - pass.SetBindGroup(0, bg); + pass.SetBindGroup(0, bg0); + pass.DispatchWorkgroups(1); + pass.SetBindGroup(0, bg1); pass.DispatchWorkgroups(1); pass.End(); encoder.Finish(); diff --git a/src/dawn/tests/unittests/validation/WritableBufferBindingAliasingValidationTests.cpp b/src/dawn/tests/unittests/validation/WritableBufferBindingAliasingValidationTests.cpp index 63ec2f6c07..6b9366a70c 100644 --- a/src/dawn/tests/unittests/validation/WritableBufferBindingAliasingValidationTests.cpp +++ b/src/dawn/tests/unittests/validation/WritableBufferBindingAliasingValidationTests.cpp @@ -89,7 +89,6 @@ std::string GenerateReferenceString(const BindingDescriptorGroups& bindingsGroup } // Creates a compute shader with given bindings -// std::string CreateComputeShaderWithBindings(const std::vector& bindings) { std::string CreateComputeShaderWithBindings(const BindingDescriptorGroups& bindingsGroups) { return GenerateBindingString(bindingsGroups) + "@compute @workgroup_size(1,1,1) fn main() {\n" + GenerateReferenceString(bindingsGroups, wgpu::ShaderStage::Compute) + "}"; @@ -137,11 +136,6 @@ class WritableBufferBindingAliasingValidationTests : public ValidationTest { return device.CreateComputePipeline(&csDesc); } - // Creates compute pipeline with default layout - wgpu::ComputePipeline CreateComputePipelineWithDefaultLayout(const std::string& shader) { - return CreateComputePipeline({}, shader); - } - // Creates render pipeline given layouts and shaders wgpu::RenderPipeline CreateRenderPipeline(const std::vector& layouts, const std::string& vertexShader, @@ -165,12 +159,6 @@ class WritableBufferBindingAliasingValidationTests : public ValidationTest { return device.CreateRenderPipeline(&pipelineDescriptor); } - // Creates render pipeline with default layout - wgpu::RenderPipeline CreateRenderPipelineWithDefaultLayout(const std::string& vertexShader, - const std::string& fragShader) { - return CreateRenderPipeline({}, vertexShader, fragShader); - } - // Creates bind group layout with given minimum sizes for each binding wgpu::BindGroupLayout CreateBindGroupLayout(const std::vector& bindings) { std::vector entries; @@ -413,18 +401,10 @@ TEST_F(WritableBufferBindingAliasingValidationTests, SetBindGroupLazyAspect) { {{0, bufferStorage, 0, 16}, wgpu::BufferBindingType::Storage}, {{1, bufferStorage, 0, 8}, wgpu::BufferBindingType::Storage}, }; - // no overlap, but has dynamic offset - std::vector bindingDescriptorDynamicOffset = { - {{0, bufferStorage, 256, 16}, wgpu::BufferBindingType::Storage, true}, - {{1, bufferStorage, 0, 8}, wgpu::BufferBindingType::Storage, true}, - }; // bindingDescriptor0 and 1 share the same bind group layout, shader and pipeline wgpu::BindGroupLayout layout = CreateBindGroupLayout(bindingDescriptor0); - wgpu::BindGroupLayout layoutHasDynamicOffset = - CreateBindGroupLayout(bindingDescriptorDynamicOffset); - std::string computeShader = CreateComputeShaderWithBindings({bindingDescriptor0}); wgpu::ComputePipeline computePipeline = CreateComputePipeline({layout}, computeShader); std::string vertexShader = CreateVertexShaderWithBindings({bindingDescriptor0}); diff --git a/src/dawn/tests/unittests/validation/WritableTextureBindingAliasingValidationTests.cpp b/src/dawn/tests/unittests/validation/WritableTextureBindingAliasingValidationTests.cpp new file mode 100644 index 0000000000..3bf6736248 --- /dev/null +++ b/src/dawn/tests/unittests/validation/WritableTextureBindingAliasingValidationTests.cpp @@ -0,0 +1,532 @@ +// Copyright 2023 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 +#include + +#include "dawn/common/Assert.h" +#include "dawn/common/Constants.h" +#include "dawn/common/Numeric.h" +#include "dawn/tests/unittests/validation/ValidationTest.h" +#include "dawn/utils/ComboRenderPipelineDescriptor.h" +#include "dawn/utils/WGPUHelpers.h" + +namespace { +using BindingDescriptorGroups = std::vector>; + +struct TestSet { + bool valid; + BindingDescriptorGroups bindingEntries; +}; + +constexpr wgpu::TextureFormat kTextureFormat = wgpu::TextureFormat::RGBA8Unorm; + +wgpu::TextureViewDescriptor GetTextureViewDescriptor( + uint32_t baseMipLevel, + uint32_t mipLevelcount, + uint32_t baseArrayLayer, + uint32_t arrayLayerCount, + wgpu::TextureAspect aspect = wgpu::TextureAspect::All) { + wgpu::TextureViewDescriptor descriptor; + descriptor.dimension = wgpu::TextureViewDimension::e2DArray; + descriptor.baseMipLevel = baseMipLevel; + descriptor.mipLevelCount = mipLevelcount; + descriptor.baseArrayLayer = baseArrayLayer; + descriptor.arrayLayerCount = arrayLayerCount; + descriptor.aspect = aspect; + return descriptor; +} + +// Creates a bind group with given bindings for shader text. +std::string GenerateBindingString(const BindingDescriptorGroups& descriptors) { + std::ostringstream ostream; + size_t index = 0; + uint32_t groupIndex = 0; + for (const auto& entries : descriptors) { + for (uint32_t bindingIndex = 0; bindingIndex < entries.size(); bindingIndex++) { + // All texture view binding format uses RGBA8Unorm in this test. + ostream << "@group(" << groupIndex << ") @binding(" << bindingIndex << ") " + << "var b" << index << " : texture_storage_2d_array;\n"; + + index++; + } + groupIndex++; + } + return ostream.str(); +} + +// Creates reference shader text to make sure variables don't get optimized out. +std::string GenerateReferenceString(const BindingDescriptorGroups& descriptors) { + std::ostringstream ostream; + size_t index = 0; + for (const auto& entries : descriptors) { + for (uint32_t bindingIndex = 0; bindingIndex < entries.size(); bindingIndex++) { + ostream << "_ = b" << index << ";\n"; + index++; + } + } + return ostream.str(); +} + +// Creates a compute shader with given bindings +std::string CreateComputeShaderWithBindings(const BindingDescriptorGroups& bindingsGroups) { + return GenerateBindingString(bindingsGroups) + "@compute @workgroup_size(1,1,1) fn main() {\n" + + GenerateReferenceString(bindingsGroups) + "}"; +} + +// Creates a fragment shader with given bindings +std::string CreateFragmentShaderWithBindings(const BindingDescriptorGroups& bindingsGroups) { + return GenerateBindingString(bindingsGroups) + "@fragment fn main() {\n" + + GenerateReferenceString(bindingsGroups) + "}"; +} + +const char* kVertexShader = R"( +@vertex fn main() -> @builtin(position) vec4 { + return vec4(); +} +)"; + +} // namespace + +class WritableTextureBindingAliasingValidationTests : public ValidationTest { + public: + wgpu::Texture CreateTexture(wgpu::TextureUsage usage, + wgpu::TextureFormat format, + uint32_t mipLevelCount, + uint32_t arrayLayerCount, + wgpu::TextureDimension dimension = wgpu::TextureDimension::e2D) { + wgpu::TextureDescriptor descriptor; + descriptor.dimension = dimension; + descriptor.size = {16, 16, arrayLayerCount}; + descriptor.sampleCount = 1; + descriptor.format = format; + descriptor.mipLevelCount = mipLevelCount; + descriptor.usage = usage; + return device.CreateTexture(&descriptor); + } + + // Creates compute pipeline given a layout and shader + wgpu::ComputePipeline CreateComputePipeline(const std::vector& layouts, + const std::string& shader) { + wgpu::ShaderModule csModule = utils::CreateShaderModule(device, shader.c_str()); + + wgpu::ComputePipelineDescriptor csDesc; + wgpu::PipelineLayoutDescriptor descriptor; + descriptor.bindGroupLayoutCount = layouts.size(); + descriptor.bindGroupLayouts = layouts.data(); + csDesc.layout = device.CreatePipelineLayout(&descriptor); + csDesc.compute.module = csModule; + csDesc.compute.entryPoint = "main"; + + return device.CreateComputePipeline(&csDesc); + } + + // Creates render pipeline given layouts and shaders + wgpu::RenderPipeline CreateRenderPipeline(const std::vector& layouts, + const std::string& vertexShader, + const std::string& fragShader) { + wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, vertexShader.c_str()); + + wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, fragShader.c_str()); + + utils::ComboRenderPipelineDescriptor pipelineDescriptor; + pipelineDescriptor.vertex.module = vsModule; + pipelineDescriptor.cFragment.module = fsModule; + pipelineDescriptor.cTargets[0].writeMask = wgpu::ColorWriteMask::None; + pipelineDescriptor.layout = nullptr; + + ASSERT(!layouts.empty()); + wgpu::PipelineLayoutDescriptor descriptor; + descriptor.bindGroupLayoutCount = layouts.size(); + descriptor.bindGroupLayouts = layouts.data(); + pipelineDescriptor.layout = device.CreatePipelineLayout(&descriptor); + + return device.CreateRenderPipeline(&pipelineDescriptor); + } + + // Creates bind group layout with given minimum sizes for each binding + wgpu::BindGroupLayout CreateBindGroupLayout( + const std::vector& bindings) { + std::vector entries; + + for (size_t i = 0; i < bindings.size(); ++i) { + const utils::BindingInitializationHelper& b = bindings[i]; + wgpu::BindGroupLayoutEntry e = {}; + e.binding = b.binding; + e.visibility = wgpu::ShaderStage::Compute | wgpu::ShaderStage::Fragment; + e.storageTexture.access = wgpu::StorageTextureAccess::WriteOnly; // only enum supported + e.storageTexture.format = kTextureFormat; + e.storageTexture.viewDimension = wgpu::TextureViewDimension::e2DArray; + + entries.push_back(e); + } + + wgpu::BindGroupLayoutDescriptor descriptor; + descriptor.entryCount = static_cast(entries.size()); + descriptor.entries = entries.data(); + return device.CreateBindGroupLayout(&descriptor); + } + + std::vector CreateBindGroups(const std::vector& layouts, + const BindingDescriptorGroups& bindingsGroups) { + std::vector bindGroups; + + ASSERT(layouts.size() == bindingsGroups.size()); + for (size_t groupIdx = 0; groupIdx < layouts.size(); groupIdx++) { + const auto& bindings = bindingsGroups[groupIdx]; + + std::vector entries; + entries.reserve(bindings.size()); + for (const auto& binding : bindings) { + entries.push_back(binding.GetAsBinding()); + } + + wgpu::BindGroupDescriptor descriptor; + descriptor.layout = layouts[groupIdx]; + descriptor.entryCount = static_cast(entries.size()); + descriptor.entries = entries.data(); + + bindGroups.push_back(device.CreateBindGroup(&descriptor)); + } + + return bindGroups; + } + + // Runs a single dispatch with given pipeline and bind group + void TestDispatch(const wgpu::ComputePipeline& computePipeline, + const std::vector& bindGroups, + const TestSet& test) { + wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder computePassEncoder = commandEncoder.BeginComputePass(); + computePassEncoder.SetPipeline(computePipeline); + + ASSERT(bindGroups.size() == test.bindingEntries.size()); + ASSERT(bindGroups.size() > 0); + for (size_t i = 0; i < bindGroups.size(); ++i) { + computePassEncoder.SetBindGroup(i, bindGroups[i]); + } + + computePassEncoder.DispatchWorkgroups(1); + computePassEncoder.End(); + if (!test.valid) { + ASSERT_DEVICE_ERROR(commandEncoder.Finish()); + } else { + commandEncoder.Finish(); + } + } + + // Runs a single draw with given pipeline and bind group + void TestDraw(const wgpu::RenderPipeline& renderPipeline, + const std::vector& bindGroups, + const TestSet& test) { + PlaceholderRenderPass renderPass(device); + + wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); + wgpu::RenderPassEncoder renderPassEncoder = commandEncoder.BeginRenderPass(&renderPass); + renderPassEncoder.SetPipeline(renderPipeline); + + ASSERT(bindGroups.size() == test.bindingEntries.size()); + ASSERT(bindGroups.size() > 0); + for (size_t i = 0; i < bindGroups.size(); ++i) { + renderPassEncoder.SetBindGroup(i, bindGroups[i]); + } + + renderPassEncoder.Draw(3); + renderPassEncoder.End(); + if (!test.valid) { + ASSERT_DEVICE_ERROR(commandEncoder.Finish()); + } else { + commandEncoder.Finish(); + } + } + + void TestBindings(const wgpu::ComputePipeline& computePipeline, + const wgpu::RenderPipeline& renderPipeline, + const std::vector& layouts, + const TestSet& test) { + std::vector bindGroups = CreateBindGroups(layouts, test.bindingEntries); + + TestDispatch(computePipeline, bindGroups, test); + TestDraw(renderPipeline, bindGroups, test); + } +}; + +// Test various combinations of texture mip levels, array layers, aspects, bind groups, etc. +// validating aliasing +TEST_F(WritableTextureBindingAliasingValidationTests, BasicTest) { + wgpu::Texture textureStorage = + CreateTexture(wgpu::TextureUsage::StorageBinding, kTextureFormat, 4, 4); + wgpu::Texture textureStorage2 = + CreateTexture(wgpu::TextureUsage::StorageBinding, kTextureFormat, 4, 4); + + // view0 and view1 don't intersect at all + wgpu::TextureViewDescriptor viewDescriptor0 = GetTextureViewDescriptor(0, 1, 0, 1); + wgpu::TextureView view0 = textureStorage.CreateView(&viewDescriptor0); + wgpu::TextureViewDescriptor viewDescriptor1 = GetTextureViewDescriptor(1, 1, 1, 1); + wgpu::TextureView view1 = textureStorage.CreateView(&viewDescriptor1); + + // view2 and view3 intersects in mip levels only + wgpu::TextureViewDescriptor viewDescriptor2 = GetTextureViewDescriptor(0, 1, 0, 1); + wgpu::TextureView view2 = textureStorage.CreateView(&viewDescriptor2); + wgpu::TextureViewDescriptor viewDescriptor3 = GetTextureViewDescriptor(0, 1, 1, 1); + wgpu::TextureView view3 = textureStorage.CreateView(&viewDescriptor3); + + // view4 and view5 intersects in array layers only + wgpu::TextureViewDescriptor viewDescriptor4 = GetTextureViewDescriptor(0, 1, 0, 3); + wgpu::TextureView view4 = textureStorage.CreateView(&viewDescriptor4); + wgpu::TextureViewDescriptor viewDescriptor5 = GetTextureViewDescriptor(1, 1, 1, 3); + wgpu::TextureView view5 = textureStorage.CreateView(&viewDescriptor5); + + // view6 and view7 intersects in both mip levels and array layers + wgpu::TextureViewDescriptor viewDescriptor6 = GetTextureViewDescriptor(0, 1, 0, 3); + wgpu::TextureView view6 = textureStorage.CreateView(&viewDescriptor6); + wgpu::TextureViewDescriptor viewDescriptor7 = GetTextureViewDescriptor(0, 1, 1, 3); + wgpu::TextureView view7 = textureStorage.CreateView(&viewDescriptor7); + + // view72 is created by another texture, so no aliasing at all. + wgpu::TextureView view72 = textureStorage2.CreateView(&viewDescriptor7); + + std::vector testSet = { + // same texture, subresources don't intersect + {true, + {{ + {0, view0}, + {1, view1}, + }}}, + // same texture, subresources don't intersect + {true, + {{ + {0, view2}, + {1, view3}, + }}}, + // same texture, subresources don't intersect, in different bind groups + {true, + {{ + {0, view0}, + }, + { + {0, view1}, + }}}, + // same texture, subresources intersect in array layers + {true, + {{ + {0, view4}, + {1, view5}, + }}}, + + // same texture, subresources intersect in both mip levels and array layers + {false, + {{ + {0, view6}, + {1, view7}, + }}}, + // reverse order to test range overlap logic + {false, + {{ + {0, view6}, + {1, view7}, + }}}, + // subreources intersect in different bind groups + {false, + {{ + {0, view6}, + }, + { + {0, view7}, + }}}, + // different texture, no aliasing at all + {true, + {{ + {0, view6}, + {1, view72}, + }}}, + // Altough spec says texture aspect could also affect whether two texture view intersects, + // It is not possible to create storage texture with depth stencil format, with different + // aspect values (all, depth only, stencil only) + // So we don't have tests for this case. + }; + + for (const auto& test : testSet) { + std::vector layouts; + for (const std::vector& bindings : + test.bindingEntries) { + layouts.push_back(CreateBindGroupLayout(bindings)); + } + + std::string computeShader = CreateComputeShaderWithBindings(test.bindingEntries); + wgpu::ComputePipeline computePipeline = CreateComputePipeline(layouts, computeShader); + std::string fragmentShader = CreateFragmentShaderWithBindings(test.bindingEntries); + wgpu::RenderPipeline renderPipeline = + CreateRenderPipeline(layouts, kVertexShader, fragmentShader); + + TestBindings(computePipeline, renderPipeline, layouts, test); + } +} + +// Test if validate bind group lazy aspect flag is set and checked properly +TEST_F(WritableTextureBindingAliasingValidationTests, SetBindGroupLazyAspect) { + wgpu::Texture textureStorage = + CreateTexture(wgpu::TextureUsage::StorageBinding, kTextureFormat, 4, 4); + + // view0 and view1 don't intersect + wgpu::TextureViewDescriptor viewDescriptor0 = GetTextureViewDescriptor(0, 1, 0, 1); + wgpu::TextureView view0 = textureStorage.CreateView(&viewDescriptor0); + wgpu::TextureViewDescriptor viewDescriptor1 = GetTextureViewDescriptor(1, 1, 1, 1); + wgpu::TextureView view1 = textureStorage.CreateView(&viewDescriptor1); + + // view2 and view3 intersects + wgpu::TextureViewDescriptor viewDescriptor2 = GetTextureViewDescriptor(0, 1, 0, 2); + wgpu::TextureView view2 = textureStorage.CreateView(&viewDescriptor2); + wgpu::TextureViewDescriptor viewDescriptor3 = GetTextureViewDescriptor(0, 1, 1, 2); + wgpu::TextureView view3 = textureStorage.CreateView(&viewDescriptor3); + + // subresources don't intersect, create valid bindGroups + std::vector bindingDescriptor0 = {{ + {0, view0}, + {1, view1}, + }}; + // subresources intersect, create invalid bindGroups + std::vector bindingDescriptor1 = {{ + {0, view2}, + {1, view3}, + }}; + + // bindingDescriptor0 and 1 share the same bind group layout, shader and pipeline + wgpu::BindGroupLayout layout = CreateBindGroupLayout(bindingDescriptor0); + + std::string computeShader = CreateComputeShaderWithBindings({bindingDescriptor0}); + wgpu::ComputePipeline computePipeline = CreateComputePipeline({layout}, computeShader); + std::string fragmentShader = CreateFragmentShaderWithBindings({bindingDescriptor0}); + wgpu::RenderPipeline renderPipeline = + CreateRenderPipeline({layout}, kVertexShader, fragmentShader); + + std::vector bindGroups = + CreateBindGroups({layout, layout}, {bindingDescriptor0, bindingDescriptor1}); + + // Test compute pass dispatch + + // bindGroups[0] is valid + { + wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder computePassEncoder = commandEncoder.BeginComputePass(); + computePassEncoder.SetPipeline(computePipeline); + + computePassEncoder.SetBindGroup(0, bindGroups[0]); + computePassEncoder.DispatchWorkgroups(1); + + computePassEncoder.End(); + commandEncoder.Finish(); + } + + // bindGroups[1] is invalid + { + wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder computePassEncoder = commandEncoder.BeginComputePass(); + computePassEncoder.SetPipeline(computePipeline); + + computePassEncoder.SetBindGroup(0, bindGroups[1]); + computePassEncoder.DispatchWorkgroups(1); + + computePassEncoder.End(); + ASSERT_DEVICE_ERROR(commandEncoder.Finish()); + } + + // setting bindGroups[1] first and then resetting to bindGroups[0] is valid + { + wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder computePassEncoder = commandEncoder.BeginComputePass(); + computePassEncoder.SetPipeline(computePipeline); + + computePassEncoder.SetBindGroup(0, bindGroups[1]); + computePassEncoder.SetBindGroup(0, bindGroups[0]); + computePassEncoder.DispatchWorkgroups(1); + + computePassEncoder.End(); + commandEncoder.Finish(); + } + + // bindGroups[0] is valid, bindGroups[1] is invalid but set to an unused slot, should still be + // valid + { + wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); + wgpu::ComputePassEncoder computePassEncoder = commandEncoder.BeginComputePass(); + computePassEncoder.SetPipeline(computePipeline); + + computePassEncoder.SetBindGroup(0, bindGroups[0]); + computePassEncoder.SetBindGroup(1, bindGroups[1]); + computePassEncoder.DispatchWorkgroups(1); + + computePassEncoder.End(); + commandEncoder.Finish(); + } + + // Test render pass draw + + PlaceholderRenderPass renderPass(device); + + // bindGroups[0] is valid + { + wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); + wgpu::RenderPassEncoder renderPassEncoder = commandEncoder.BeginRenderPass(&renderPass); + renderPassEncoder.SetPipeline(renderPipeline); + + renderPassEncoder.SetBindGroup(0, bindGroups[0]); + renderPassEncoder.Draw(3); + + renderPassEncoder.End(); + commandEncoder.Finish(); + } + + // bindGroups[1] is invalid + { + wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); + wgpu::RenderPassEncoder renderPassEncoder = commandEncoder.BeginRenderPass(&renderPass); + renderPassEncoder.SetPipeline(renderPipeline); + + renderPassEncoder.SetBindGroup(0, bindGroups[1]); + renderPassEncoder.Draw(3); + + renderPassEncoder.End(); + ASSERT_DEVICE_ERROR(commandEncoder.Finish()); + } + + // setting bindGroups[1] first and then resetting to bindGroups[0] is valid + { + wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); + wgpu::RenderPassEncoder renderPassEncoder = commandEncoder.BeginRenderPass(&renderPass); + renderPassEncoder.SetPipeline(renderPipeline); + + renderPassEncoder.SetBindGroup(0, bindGroups[1]); + renderPassEncoder.SetBindGroup(0, bindGroups[0]); + renderPassEncoder.Draw(3); + + renderPassEncoder.End(); + commandEncoder.Finish(); + } + + // bindGroups[0] is valid, bindGroups[1] is invalid but set to an unused slot, should still be + // valid + { + wgpu::CommandEncoder commandEncoder = device.CreateCommandEncoder(); + wgpu::RenderPassEncoder renderPassEncoder = commandEncoder.BeginRenderPass(&renderPass); + renderPassEncoder.SetPipeline(renderPipeline); + + renderPassEncoder.SetBindGroup(0, bindGroups[0]); + renderPassEncoder.SetBindGroup(1, bindGroups[1]); + renderPassEncoder.Draw(3); + + renderPassEncoder.End(); + commandEncoder.Finish(); + } +} diff --git a/webgpu-cts/expectations.txt b/webgpu-cts/expectations.txt index 8b06b49f9f..8b2b24455c 100644 --- a/webgpu-cts/expectations.txt +++ b/webgpu-cts/expectations.txt @@ -258,6 +258,11 @@ crbug.com/dawn/0000 [ win ] webgpu:shader,execution,memory_model,* [ Failure ] ################################################################################ crbug.com/tint/0000 webgpu:shader,validation,parse,blankspace:null_characters:contains_null=true;placement="comment" [ Failure ] +################################################################################ +# Storage texture binding validation failures +################################################################################ +crbug.com/dawn/1642 webgpu:api,validation,resource_usages,texture,in_pass_encoder:subresources_and_binding_types_combination_for_color:compute=true;type0="writeonly-storage-texture";type1="writeonly-storage-texture" [ Failure ] + ################################################################################ # Flaky on Intel Mac # KEEP