Allow readwrite storage texture in pipelinelayout matches writeonly in shader
This patch enables using read-write storage texture access in pipeline layout
and declaring the correspondent storage texture as write-only in shader.
Note that currently we still disallow read-write access in pipeline layout be
compatible with read-only in shader as on D3D12 read-write access in pipeline
layout will be implemented with an UAV, while read-only access in shader will
consume an SRV slot.
Bug: dawn:1972
Test: dawn_unittests, dawn_end2end_tests
Change-Id: I26387eb4c57cd850bc7672db9a5a3b68a4f7f8ef
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/150560
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Austin Eng <enga@chromium.org>
Commit-Queue: Jiawei Shao <jiawei.shao@intel.com>
diff --git a/src/dawn/native/ShaderModule.cpp b/src/dawn/native/ShaderModule.cpp
index 41b66a8..56d8a2c 100644
--- a/src/dawn/native/ShaderModule.cpp
+++ b/src/dawn/native/ShaderModule.cpp
@@ -371,6 +371,14 @@
return requiredBufferSizes;
}
+bool IsShaderCompatibleWithPipelineLayoutOnStorageTextureAccess(
+ const BindingInfo& bindingInfo,
+ const ShaderBindingInfo& shaderBindingInfo) {
+ return bindingInfo.storageTexture.access == shaderBindingInfo.storageTexture.access ||
+ (bindingInfo.storageTexture.access == wgpu::StorageTextureAccess::ReadWrite &&
+ shaderBindingInfo.storageTexture.access == wgpu::StorageTextureAccess::WriteOnly);
+}
+
MaybeError ValidateCompatibilityOfSingleBindingWithLayout(const DeviceBase* device,
const BindGroupLayoutInternalBase* layout,
SingleShaderStage entryPointStage,
@@ -451,10 +459,11 @@
ASSERT(layoutInfo.storageTexture.format != wgpu::TextureFormat::Undefined);
ASSERT(shaderInfo.storageTexture.format != wgpu::TextureFormat::Undefined);
- DAWN_INVALID_IF(layoutInfo.storageTexture.access != shaderInfo.storageTexture.access,
- "The layout's binding access (%s) isn't compatible with the shader's "
- "binding access (%s).",
- layoutInfo.storageTexture.access, shaderInfo.storageTexture.access);
+ DAWN_INVALID_IF(
+ !IsShaderCompatibleWithPipelineLayoutOnStorageTextureAccess(layoutInfo, shaderInfo),
+ "The layout's binding access (%s) isn't compatible with the shader's "
+ "binding access (%s).",
+ layoutInfo.storageTexture.access, shaderInfo.storageTexture.access);
DAWN_INVALID_IF(layoutInfo.storageTexture.format != shaderInfo.storageTexture.format,
"The layout's binding format (%s) doesn't match the shader's binding "
diff --git a/src/dawn/tests/end2end/StorageTextureTests.cpp b/src/dawn/tests/end2end/StorageTextureTests.cpp
index 54098c9..d57d921 100644
--- a/src/dawn/tests/end2end/StorageTextureTests.cpp
+++ b/src/dawn/tests/end2end/StorageTextureTests.cpp
@@ -1291,6 +1291,61 @@
CheckDrawsGreen(kSimpleVertexShader, fsstream.str().c_str(), readonlyStorageTexture);
}
+// Verify using read-write storage texture access in pipeline layout is compatible with write-only
+// storage texture access in shader.
+TEST_P(ReadWriteStorageTextureTests, ReadWriteInPipelineLayoutAndWriteOnlyInShader) {
+ DAWN_TEST_UNSUPPORTED_IF(!IsReadWriteStorageTextureSupported());
+
+ constexpr wgpu::TextureFormat kStorageTextureFormat = wgpu::TextureFormat::R32Uint;
+ std::array<uint32_t, kWidth * kHeight> expectedData;
+ for (size_t i = 0; i < expectedData.size(); ++i) {
+ expectedData[i] = i + 1;
+ }
+
+ wgpu::Texture storageTexture = CreateTexture(
+ wgpu::TextureFormat::R32Uint,
+ wgpu::TextureUsage::StorageBinding | wgpu::TextureUsage::CopySrc, {kWidth, kHeight, 1});
+
+ std::ostringstream sstream;
+ sstream << R"(
+enable chromium_experimental_read_write_storage_texture;
+@group(0) @binding(0) var rwImage : texture_storage_2d<r32uint, write>;
+
+@compute @workgroup_size()"
+ << kWidth << ", " << kHeight << R"()
+fn main(
+ @builtin(local_invocation_id) local_id: vec3u,
+ @builtin(local_invocation_index) local_index : u32) {
+ let data1 = vec4u(local_index + 1u, 0, 0, 1);
+ textureStore(rwImage, vec2i(local_id.xy), data1);
+})";
+
+ wgpu::BindGroupLayout bindGroupLayout = utils::MakeBindGroupLayout(
+ device, {{0, wgpu::ShaderStage::Compute, wgpu::StorageTextureAccess::ReadWrite,
+ kStorageTextureFormat, wgpu::TextureViewDimension::e2D}});
+ wgpu::ComputePipelineDescriptor computeDescriptor;
+ computeDescriptor.layout = utils::MakePipelineLayout(device, {bindGroupLayout});
+ computeDescriptor.compute.module = utils::CreateShaderModule(device, sstream.str().c_str());
+ computeDescriptor.compute.entryPoint = "main";
+ wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&computeDescriptor);
+
+ wgpu::BindGroup bindGroup =
+ utils::MakeBindGroup(device, bindGroupLayout, {{0, storageTexture.CreateView()}});
+
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder computePassEncoder = encoder.BeginComputePass();
+ computePassEncoder.SetBindGroup(0, bindGroup);
+ computePassEncoder.SetPipeline(pipeline);
+ computePassEncoder.DispatchWorkgroups(1);
+ computePassEncoder.End();
+ wgpu::CommandBuffer commandBuffer = encoder.Finish();
+ queue.Submit(1, &commandBuffer);
+
+ CheckOutputStorageTexture(storageTexture, wgpu::TextureFormat::R32Uint, {kWidth, kHeight},
+ reinterpret_cast<const uint8_t*>(expectedData.data()),
+ expectedData.size() * sizeof(uint32_t));
+}
+
DAWN_INSTANTIATE_TEST(ReadWriteStorageTextureTests,
D3D11Backend(),
D3D12Backend(),
diff --git a/src/dawn/tests/unittests/validation/StorageTextureValidationTests.cpp b/src/dawn/tests/unittests/validation/StorageTextureValidationTests.cpp
index f7f1f3b..2dee034 100644
--- a/src/dawn/tests/unittests/validation/StorageTextureValidationTests.cpp
+++ b/src/dawn/tests/unittests/validation/StorageTextureValidationTests.cpp
@@ -989,8 +989,9 @@
}
}
-// Test that storage texture access in shader must match the one in pipeline layout when we create
-// a pipeline with storage texture.
+// Test that storage texture access in shader must be compatible with the one in pipeline layout
+// when we create a pipeline with storage texture. Note that read-write storage texture access in
+// pipeline layout is compatible with write-only storage texture access in shader.
TEST_F(ReadWriteStorageTextureValidationTests, StorageTextureAccessInPipeline) {
constexpr std::array<wgpu::StorageTextureAccess, 3> kStorageTextureAccesses = {
{wgpu::StorageTextureAccess::ReadOnly, wgpu::StorageTextureAccess::WriteOnly,
@@ -1010,7 +1011,9 @@
computePipelineDescriptor.compute.entryPoint = "main";
computePipelineDescriptor.layout =
utils::MakePipelineLayout(device, {{bindGroupLayout}});
- if (accessInShader == accessInBindGroupLayout) {
+ if (accessInShader == accessInBindGroupLayout ||
+ (accessInShader == wgpu::StorageTextureAccess::WriteOnly &&
+ accessInBindGroupLayout == wgpu::StorageTextureAccess::ReadWrite)) {
device.CreateComputePipeline(&computePipelineDescriptor);
} else {
ASSERT_DEVICE_ERROR(device.CreateComputePipeline(&computePipelineDescriptor));
