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 <noreply+kokoro@google.com>
Reviewed-by: Austin Eng <enga@chromium.org>
Commit-Queue: Shrek Shao <shrekshao@google.com>
diff --git a/src/dawn/common/Numeric.h b/src/dawn/common/Numeric.h
index 6387f48..7be33c8 100644
--- a/src/dawn/common/Numeric.h
+++ b/src/dawn/common/Numeric.h
@@ -61,7 +61,23 @@
constexpr double kMax = static_cast<double>(std::numeric_limits<T>::max());
return kLowest <= value && value <= kMax;
} else {
- static_assert(sizeof(T) != sizeof(T), "Unsupported type");
+ static_assert(std::is_same_v<T, float> || std::is_integral_v<T>, "Unsupported type");
+ }
+}
+
+// Returns if two inclusive integral ranges [x0, x1] and [y0, y1] have overlap.
+template <typename T>
+bool RangesOverlap(T x0, T x1, T y0, T y1) {
+ ASSERT(x0 <= x1 && y0 <= y1);
+ if constexpr (std::is_integral_v<T>) {
+ // 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<T>, "Unsupported type");
}
}
diff --git a/src/dawn/native/CommandBufferStateTracker.cpp b/src/dawn/native/CommandBufferStateTracker.cpp
index 245f3a4..fc6cbc9 100644
--- a/src/dawn/native/CommandBufferStateTracker.cpp
+++ b/src/dawn/native/CommandBufferStateTracker.cpp
@@ -17,6 +17,7 @@
#include <optional>
#include <type_traits>
#include <utility>
+#include <variant>
#include "dawn/common/Assert.h"
#include "dawn/common/BitSetIterator.h"
@@ -53,7 +54,7 @@
return std::nullopt;
}
-struct BufferBindingAliasingResult {
+struct BufferAliasing {
struct Entry {
BindGroupIndex bindGroupIndex;
BindingIndex bindingIndex;
@@ -66,18 +67,42 @@
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<std::monostate, BufferAliasing, TextureAliasing>;
+
template <typename Return>
Return FindStorageBufferBindingAliasing(
const PipelineLayoutBase* pipelineLayout,
const ityp::array<BindGroupIndex, BindGroupBase*, kMaxBindGroups>& bindGroups,
- const ityp::array<BindGroupIndex, std::vector<uint32_t>, kMaxBindGroups> dynamicOffsets) {
+ const ityp::array<BindGroupIndex, std::vector<uint32_t>, 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<Return, WritableBindingAliasingResult>;
+
// 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<BufferBinding, 8> bindingsToCheck;
+ // There can at most be 8 storage buffer bindings (in default limits) per shader stage.
+ StackVector<BufferBinding, 8> storageBufferBindingsToCheck;
+ StackVector<std::pair<BindGroupIndex, BindingIndex>, 8> bufferBindingIndices;
- StackVector<std::pair<BindGroupIndex, BindingIndex>, 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<const TextureViewBase*, 8> storageTextureViewsToCheck;
+ StackVector<std::pair<BindGroupIndex, BindingIndex>, 8> textureBindingIndices;
for (BindGroupIndex groupIndex : IterateBitSet(pipelineLayout->GetBindGroupLayoutsMask())) {
BindGroupLayoutBase* bgl = bindGroups[groupIndex]->GetLayout();
@@ -107,55 +132,127 @@
adjustedOffset += dynamicOffsets[groupIndex][static_cast<uint32_t>(bindingIndex)];
}
- bindingsToCheck->push_back(BufferBinding{
+ storageBufferBindingsToCheck->push_back(BufferBinding{
bufferBinding.buffer,
adjustedOffset,
bufferBinding.size,
});
- if constexpr (std::is_same_v<Return, std::optional<BufferBindingAliasingResult>>) {
- 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<Return, bool>) {
+ 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,
- std::optional<BufferBindingAliasingResult>>) {
- 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, bool>) {
- return false;
- } else if constexpr (std::is_same_v<Return, std::optional<BufferBindingAliasingResult>>) {
- 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;
+ }
+ }
+ }
}
- UNREACHABLE();
+
+ if constexpr (kProduceDetails) {
+ return WritableBindingAliasingResult();
+ } else {
+ return false;
+ }
}
} // namespace
@@ -396,7 +493,7 @@
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 @@
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 @@
}
}
- auto result = FindStorageBufferBindingAliasing<std::optional<BufferBindingAliasingResult>>(
+ auto result = FindStorageBufferBindingAliasing<WritableBindingAliasingResult>(
mLastPipelineLayout, mBindgroups, mDynamicOffsets);
- if (result) {
+ if (std::holds_alternative<BufferAliasing>(result)) {
+ const auto& a = std::get<BufferAliasing>(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<uint32_t>(result->e0.bindGroupIndex),
- static_cast<uint32_t>(result->e0.bindingIndex),
- static_cast<uint32_t>(result->e1.bindGroupIndex),
- static_cast<uint32_t>(result->e1.bindingIndex), result->e0.offset, result->e0.size,
- result->e1.offset, result->e1.size);
+ static_cast<uint32_t>(a.e0.bindGroupIndex),
+ static_cast<uint32_t>(a.e0.bindingIndex),
+ static_cast<uint32_t>(a.e1.bindGroupIndex),
+ static_cast<uint32_t>(a.e1.bindingIndex), a.e0.offset, a.e0.size, a.e1.offset,
+ a.e1.size);
+ } else {
+ ASSERT(std::holds_alternative<TextureAliasing>(result));
+ const auto& a = std::get<TextureAliasing>(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<uint32_t>(a.e0.bindGroupIndex),
+ static_cast<uint32_t>(a.e0.bindingIndex),
+ static_cast<uint32_t>(a.e1.bindGroupIndex),
+ static_cast<uint32_t>(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 e7874e6..d328f82 100644
--- a/src/dawn/native/CommandValidation.cpp
+++ b/src/dawn/native/CommandValidation.cpp
@@ -21,6 +21,7 @@
#include <utility>
#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 @@
}
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<uint64_t>(minStart) + static_cast<uint64_t>(length) >
- static_cast<uint64_t>(maxStart);
+ if (length < 1) {
+ return false;
+ }
+ return RangesOverlap<uint64_t>(
+ static_cast<uint64_t>(startA),
+ static_cast<uint64_t>(startA) + static_cast<uint64_t>(length) - 1,
+ static_cast<uint64_t>(startB),
+ static_cast<uint64_t>(startB) + static_cast<uint64_t>(length) - 1);
}
ResultOrError<uint64_t> ComputeRequiredBytesInCopy(const TexelBlockInfo& blockInfo,
diff --git a/src/dawn/native/Subresource.h b/src/dawn/native/Subresource.h
index 473631a..4d4e387 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 a87fc1e..bfeb26a 100644
--- a/src/dawn/tests/BUILD.gn
+++ b/src/dawn/tests/BUILD.gn
@@ -292,6 +292,7 @@
"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 @@
"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 0000000..d94344a
--- /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 24d8e04..846b637 100644
--- a/src/dawn/tests/unittests/validation/ResourceUsageTrackingTests.cpp
+++ b/src/dawn/tests/unittests/validation/ResourceUsageTrackingTests.cpp
@@ -973,9 +973,11 @@
// 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 @@
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 63ec2f6..6b9366a 100644
--- a/src/dawn/tests/unittests/validation/WritableBufferBindingAliasingValidationTests.cpp
+++ b/src/dawn/tests/unittests/validation/WritableBufferBindingAliasingValidationTests.cpp
@@ -89,7 +89,6 @@
}
// Creates a compute shader with given bindings
-// std::string CreateComputeShaderWithBindings(const std::vector<BindingDescriptor>& 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 @@
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<wgpu::BindGroupLayout>& layouts,
const std::string& vertexShader,
@@ -165,12 +159,6 @@
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<BindingDescriptor>& bindings) {
std::vector<wgpu::BindGroupLayoutEntry> entries;
@@ -413,18 +401,10 @@
{{0, bufferStorage, 0, 16}, wgpu::BufferBindingType::Storage},
{{1, bufferStorage, 0, 8}, wgpu::BufferBindingType::Storage},
};
- // no overlap, but has dynamic offset
- std::vector<BindingDescriptor> 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 0000000..3bf6736
--- /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 <string>
+#include <vector>
+
+#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<std::vector<utils::BindingInitializationHelper>>;
+
+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<rgba8unorm, write>;\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<f32> {
+ return vec4<f32>();
+}
+)";
+
+} // 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<wgpu::BindGroupLayout>& 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<wgpu::BindGroupLayout>& 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<utils::BindingInitializationHelper>& bindings) {
+ std::vector<wgpu::BindGroupLayoutEntry> 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<uint32_t>(entries.size());
+ descriptor.entries = entries.data();
+ return device.CreateBindGroupLayout(&descriptor);
+ }
+
+ std::vector<wgpu::BindGroup> CreateBindGroups(const std::vector<wgpu::BindGroupLayout>& layouts,
+ const BindingDescriptorGroups& bindingsGroups) {
+ std::vector<wgpu::BindGroup> bindGroups;
+
+ ASSERT(layouts.size() == bindingsGroups.size());
+ for (size_t groupIdx = 0; groupIdx < layouts.size(); groupIdx++) {
+ const auto& bindings = bindingsGroups[groupIdx];
+
+ std::vector<wgpu::BindGroupEntry> 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<uint32_t>(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<wgpu::BindGroup>& 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<wgpu::BindGroup>& 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<wgpu::BindGroupLayout>& layouts,
+ const TestSet& test) {
+ std::vector<wgpu::BindGroup> 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> 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<wgpu::BindGroupLayout> layouts;
+ for (const std::vector<utils::BindingInitializationHelper>& 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<utils::BindingInitializationHelper> bindingDescriptor0 = {{
+ {0, view0},
+ {1, view1},
+ }};
+ // subresources intersect, create invalid bindGroups
+ std::vector<utils::BindingInitializationHelper> 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<wgpu::BindGroup> 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 8b06b49..8b2b244 100644
--- a/webgpu-cts/expectations.txt
+++ b/webgpu-cts/expectations.txt
@@ -259,6 +259,11 @@
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
################################################################################
