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>

src/dawn/common/Numeric.h

@@ -61,7 +61,23 @@ bool IsDoubleValueRepresentable(double value) {
         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");
     }
 }
 

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 @@ std::optional<uint32_t> 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<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.
+    // There can at most be 8 storage buffer bindings (in default limits) per shader stage.
-    StackVector<BufferBinding, 8> bindingsToCheck;
+    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 @@ Return FindStorageBufferBindingAliasing(
                 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>>) {
+            if constexpr (kProduceDetails) {
-                bindingIndices->emplace_back(groupIndex, bindingIndex);
+                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.
+    // Iterate through each buffer bindings to find if any writable storage bindings aliasing
-    // Given that maxStorageBuffersPerShaderStage is 8,
+    // exists. Given that maxStorageBuffersPerShaderStage is 8, it doesn't seem too bad to do a
-    // it doesn't seem too bad to do a nested loop check.
+    // 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++) {
+    for (size_t i = 0; i < storageBufferBindingsToCheck->size(); i++) {
-        const auto& bufferBinding0 = bindingsToCheck[i];
+        const auto& bufferBinding0 = storageBufferBindingsToCheck[i];
 
-        for (size_t j = i + 1; j < bindingsToCheck->size(); j++) {
+        for (size_t j = i + 1; j < storageBufferBindingsToCheck->size(); j++) {
-            const auto& bufferBinding1 = bindingsToCheck[j];
+            const auto& bufferBinding1 = storageBufferBindingsToCheck[j];
 
             if (bufferBinding0.buffer != bufferBinding1.buffer) {
                 continue;
             }
 
-            if (bufferBinding0.offset <= bufferBinding1.offset + bufferBinding1.size - 1 &&
+            if (RangesOverlap(
-                bufferBinding1.offset <= bufferBinding0.offset + bufferBinding0.size - 1) {
+                    bufferBinding0.offset, bufferBinding0.offset + bufferBinding0.size - 1,
-                if constexpr (std::is_same_v<Return, bool>) {
+                    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>) {
+    // Iterate through each texture views to find if any writable storage bindings aliasing exists.
-        return false;
+    // Given that maxStorageTexturesPerShaderStage is 8,
-    } else if constexpr (std::is_same_v<Return, std::optional<BufferBindingAliasingResult>>) {
+    // it doesn't seem too bad to do a nested loop check.
-        return std::nullopt;
+    // 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<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 "
+                "Writable storage buffer binding aliasing found between bind group index %u, "
-                "%u, and bind group index %u, binding index %u, with overlapping ranges (offset: "
+                "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>(a.e0.bindGroupIndex),
-                static_cast<uint32_t>(result->e0.bindingIndex),
+                static_cast<uint32_t>(a.e0.bindingIndex),
-                static_cast<uint32_t>(result->e1.bindGroupIndex),
+                static_cast<uint32_t>(a.e1.bindGroupIndex),
-                static_cast<uint32_t>(result->e1.bindingIndex), result->e0.offset, result->e0.size,
+                static_cast<uint32_t>(a.e1.bindingIndex), a.e0.offset, a.e0.size, a.e1.offset,
-                result->e1.offset, result->e1.size);
+                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|.

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 @@ MaybeError ValidateWriteBuffer(const DeviceBase* device,
 }
 
 bool IsRangeOverlapped(uint32_t startA, uint32_t startB, uint32_t length) {
-    uint32_t maxStart = std::max(startA, startB);
+    if (length < 1) {
-    uint32_t minStart = std::min(startA, startB);
+        return false;
-    return static_cast<uint64_t>(minStart) + static_cast<uint64_t>(length) >
+    }
-           static_cast<uint64_t>(maxStart);
+    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,

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,

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",

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));
+}

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},
+            {{0, wgpu::ShaderStage::Compute, wgpu::StorageTextureAccess::WriteOnly, kFormat}});
-             {1, wgpu::ShaderStage::Compute, wgpu::StorageTextureAccess::WriteOnly, kFormat}});
+        // Create 2 bind groups with same texture subresources and dispatch twice to avoid
-        wgpu::BindGroup bg = utils::MakeBindGroup(device, bgl, {{0, view}, {1, view}});
+        // 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();

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<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 @@ 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<wgpu::BindGroupLayout>& 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<BindingDescriptor>& bindings) {
         std::vector<wgpu::BindGroupLayoutEntry> 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<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});

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();
+    }
+}

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