Fix internal storage buffer usage

TimestampQueryTests.ResolveTwiceToSameBuffer fails on Intel latest
driver on Windows, because the kInternalStorageBuffer is not treated
in buffer usage when adding resource barrier.

Add missed kInternalStorageBuffer in buffer usage and remove
D3D12_RESOURCE_STATE_UNORDERED_ACCESS from QueryResolve, which will be
added by kInternalStorageBuffer.

Bug: dawn:797
Change-Id: I78607002179ba443b0db09c9c3bbc85fcc97a85b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/56523
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Hao Li <hao.x.li@intel.com>

src/tests/BUILD.gn

@@ -435,6 +435,7 @@ source_set("dawn_white_box_tests_sources") {
 
   sources += [
     "white_box/InternalResourceUsageTests.cpp",
+    "white_box/InternalStorageBufferBindingTests.cpp",
     "white_box/QueryInternalShaderTests.cpp",
   ]
 

src/tests/white_box/InternalResourceUsageTests.cpp

@@ -16,18 +16,25 @@
 
 #include "dawn_native/dawn_platform.h"
 
-class InternalResourceUsageTests : public DawnTest {};
+class InternalResourceUsageTests : public DawnTest {
+  protected:
+    wgpu::Buffer CreateBuffer(wgpu::BufferUsage usage) {
+        wgpu::BufferDescriptor descriptor;
+        descriptor.size = 4;
+        descriptor.usage = usage;
+
+        return device.CreateBuffer(&descriptor);
+    }
+};
 
 // Verify it is an error to create a buffer with a buffer usage that should only be used
 // internally.
 TEST_P(InternalResourceUsageTests, InternalBufferUsage) {
     DAWN_TEST_UNSUPPORTED_IF(HasToggleEnabled("skip_validation"));
 
-    wgpu::BufferDescriptor descriptor;
-    descriptor.size = 4;
-    descriptor.usage = dawn_native::kReadOnlyStorageBuffer;
+    ASSERT_DEVICE_ERROR(CreateBuffer(dawn_native::kReadOnlyStorageBuffer));
 
-    ASSERT_DEVICE_ERROR(device.CreateBuffer(&descriptor));
+    ASSERT_DEVICE_ERROR(CreateBuffer(dawn_native::kInternalStorageBuffer));
 }
 
 // Verify it is an error to create a texture with a texture usage that should only be used
@@ -43,3 +50,23 @@ TEST_P(InternalResourceUsageTests, InternalTextureUsage) {
 }
 
 DAWN_INSTANTIATE_TEST(InternalResourceUsageTests, NullBackend());
+
+class InternalBindingTypeTests : public DawnTest {};
+
+// Verify it is an error to create a bind group layout with a buffer binding type that should only
+// be used internally.
+TEST_P(InternalBindingTypeTests, InternalStorageBufferBindingType) {
+    DAWN_TEST_UNSUPPORTED_IF(HasToggleEnabled("skip_validation"));
+
+    wgpu::BindGroupLayoutEntry bglEntry;
+    bglEntry.binding = 0;
+    bglEntry.buffer.type = dawn_native::kInternalStorageBufferBinding;
+    bglEntry.visibility = wgpu::ShaderStage::Compute;
+
+    wgpu::BindGroupLayoutDescriptor bglDesc;
+    bglDesc.entryCount = 1;
+    bglDesc.entries = &bglEntry;
+    ASSERT_DEVICE_ERROR(device.CreateBindGroupLayout(&bglDesc));
+}
+
+DAWN_INSTANTIATE_TEST(InternalBindingTypeTests, NullBackend());

src/tests/white_box/InternalStorageBufferBindingTests.cpp

@@ -0,0 +1,113 @@
+// Copyright 2021 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 "tests/DawnTest.h"
+
+#include "dawn_native/BindGroupLayout.h"
+#include "dawn_native/Device.h"
+#include "dawn_native/dawn_platform.h"
+#include "utils/WGPUHelpers.h"
+
+class InternalStorageBufferBindingTests : public DawnTest {
+  protected:
+    static constexpr uint32_t kNumValues = 4;
+    static constexpr uint32_t kIterations = 4;
+
+    void SetUp() override {
+        DawnTest::SetUp();
+        DAWN_TEST_UNSUPPORTED_IF(UsesWire());
+    }
+
+    wgpu::ComputePipeline CreateComputePipelineWithInternalStorage() {
+        wgpu::ShaderModule module = utils::CreateShaderModule(device, R"(
+            [[block]] struct Buf {
+                data : array<u32, 4>;
+            };
+
+            [[group(0), binding(0)]] var<storage, read_write> buf : Buf;
+
+            [[stage(compute), workgroup_size(1)]]
+            fn main([[builtin(global_invocation_id)]] GlobalInvocationID : vec3<u32>) {
+                buf.data[GlobalInvocationID.x] = buf.data[GlobalInvocationID.x] + 0x1234u;
+            }
+        )");
+
+        // Create binding group layout with internal storage buffer binding type
+        dawn_native::BindGroupLayoutEntry bglEntry;
+        bglEntry.binding = 0;
+        bglEntry.buffer.type = dawn_native::kInternalStorageBufferBinding;
+        bglEntry.visibility = wgpu::ShaderStage::Compute;
+
+        dawn_native::BindGroupLayoutDescriptor bglDesc;
+        bglDesc.entryCount = 1;
+        bglDesc.entries = &bglEntry;
+
+        dawn_native::DeviceBase* nativeDevice =
+            reinterpret_cast<dawn_native::DeviceBase*>(device.Get());
+
+        Ref<dawn_native::BindGroupLayoutBase> bglRef =
+            nativeDevice->CreateBindGroupLayout(&bglDesc, true).AcquireSuccess();
+
+        wgpu::BindGroupLayout bgl =
+            wgpu::BindGroupLayout::Acquire(reinterpret_cast<WGPUBindGroupLayout>(bglRef.Detach()));
+
+        // Create pipeline layout
+        wgpu::PipelineLayoutDescriptor plDesc;
+        plDesc.bindGroupLayoutCount = 1;
+        plDesc.bindGroupLayouts = &bgl;
+        wgpu::PipelineLayout layout = device.CreatePipelineLayout(&plDesc);
+
+        wgpu::ComputePipelineDescriptor pipelineDesc = {};
+        pipelineDesc.layout = layout;
+        pipelineDesc.compute.module = module;
+        pipelineDesc.compute.entryPoint = "main";
+
+        return device.CreateComputePipeline(&pipelineDesc);
+    }
+};
+
+// Test that query resolve buffer can be bound as internal storage buffer, multiple dispatches to
+// increment values in the query resolve buffer are synchronized.
+TEST_P(InternalStorageBufferBindingTests, QueryResolveBufferBoundAsInternalStorageBuffer) {
+    std::vector<uint32_t> data(kNumValues, 0);
+    std::vector<uint32_t> expected(kNumValues, 0x1234u * kIterations);
+
+    uint64_t bufferSize = static_cast<uint64_t>(data.size() * sizeof(uint32_t));
+    wgpu::Buffer buffer =
+        utils::CreateBufferFromData(device, data.data(), bufferSize,
+                                    wgpu::BufferUsage::QueryResolve | wgpu::BufferUsage::CopySrc);
+
+    wgpu::ComputePipeline pipeline = CreateComputePipelineWithInternalStorage();
+
+    wgpu::BindGroup bindGroup =
+        utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0), {{0, buffer, 0, bufferSize}});
+
+    wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+    wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+    pass.SetPipeline(pipeline);
+    pass.SetBindGroup(0, bindGroup);
+    for (uint32_t i = 0; i < kIterations; ++i) {
+        pass.Dispatch(kNumValues);
+    }
+    pass.EndPass();
+    wgpu::CommandBuffer commands = encoder.Finish();
+    queue.Submit(1, &commands);
+
+    EXPECT_BUFFER_U32_RANGE_EQ(expected.data(), buffer, 0, kNumValues);
+}
+
+DAWN_INSTANTIATE_TEST(InternalStorageBufferBindingTests,
+                      D3D12Backend(),
+                      MetalBackend(),
+                      VulkanBackend());