Enable ReadWrite storage texture usage on D3D12, Metal and Vulkan
This patch adds the support of ReadWrite storage texture usage on
D3D12, Metal and Vulkan.
Bug: dawn:1972
Test: dawn_end2end_tests
Change-Id: I4c2ac4daaeec33e5d3ea67e85ad0ee5f97a8063d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/147420
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Jiawei Shao <jiawei.shao@intel.com>
Reviewed-by: Austin Eng <enga@chromium.org>
diff --git a/src/dawn/native/CommandBufferStateTracker.cpp b/src/dawn/native/CommandBufferStateTracker.cpp
index 003b6eb..6768fba 100644
--- a/src/dawn/native/CommandBufferStateTracker.cpp
+++ b/src/dawn/native/CommandBufferStateTracker.cpp
@@ -156,11 +156,11 @@
switch (bindingInfo.storageTexture.access) {
case wgpu::StorageTextureAccess::WriteOnly:
+ case wgpu::StorageTextureAccess::ReadWrite:
break;
- // TODO(dawn:1972): Implement ReadOnly and ReadWrite storage texture
+ // TODO(dawn:1972): Implement ReadOnly storage texture
case wgpu::StorageTextureAccess::ReadOnly:
- case wgpu::StorageTextureAccess::ReadWrite:
case wgpu::StorageTextureAccess::Undefined:
default:
UNREACHABLE();
diff --git a/src/dawn/native/PassResourceUsageTracker.cpp b/src/dawn/native/PassResourceUsageTracker.cpp
index 91e2a7f..34e4bd0 100644
--- a/src/dawn/native/PassResourceUsageTracker.cpp
+++ b/src/dawn/native/PassResourceUsageTracker.cpp
@@ -130,12 +130,12 @@
TextureViewBase* view = group->GetBindingAsTextureView(bindingIndex);
switch (bindingInfo.storageTexture.access) {
case wgpu::StorageTextureAccess::WriteOnly:
+ case wgpu::StorageTextureAccess::ReadWrite:
TextureViewUsedAs(view, wgpu::TextureUsage::StorageBinding);
break;
- // TODO(dawn:1972): Implement ReadOnly and ReadWrite storage texture
+ // TODO(dawn:1972): Implement ReadOnly storage texture
case wgpu::StorageTextureAccess::ReadOnly:
- case wgpu::StorageTextureAccess::ReadWrite:
case wgpu::StorageTextureAccess::Undefined:
UNREACHABLE();
}
diff --git a/src/dawn/native/d3d12/BindGroupD3D12.cpp b/src/dawn/native/d3d12/BindGroupD3D12.cpp
index 0bede63..10312a0 100644
--- a/src/dawn/native/d3d12/BindGroupD3D12.cpp
+++ b/src/dawn/native/d3d12/BindGroupD3D12.cpp
@@ -162,7 +162,8 @@
}
switch (bindingInfo.storageTexture.access) {
- case wgpu::StorageTextureAccess::WriteOnly: {
+ case wgpu::StorageTextureAccess::WriteOnly:
+ case wgpu::StorageTextureAccess::ReadWrite: {
D3D12_UNORDERED_ACCESS_VIEW_DESC uav = view->GetUAVDescriptor();
d3d12Device->CreateUnorderedAccessView(
resource, nullptr, &uav,
@@ -171,9 +172,8 @@
break;
}
- // TODO(dawn:1972): Implement ReadOnly and ReadWrite storage texture
+ // TODO(dawn:1972): Implement ReadWrite storage texture
case wgpu::StorageTextureAccess::ReadOnly:
- case wgpu::StorageTextureAccess::ReadWrite:
case wgpu::StorageTextureAccess::Undefined:
UNREACHABLE();
}
diff --git a/src/dawn/native/d3d12/BindGroupLayoutD3D12.cpp b/src/dawn/native/d3d12/BindGroupLayoutD3D12.cpp
index 71fdd11..5b3bab7 100644
--- a/src/dawn/native/d3d12/BindGroupLayoutD3D12.cpp
+++ b/src/dawn/native/d3d12/BindGroupLayoutD3D12.cpp
@@ -48,11 +48,11 @@
case BindingInfoType::StorageTexture:
switch (bindingInfo.storageTexture.access) {
case wgpu::StorageTextureAccess::WriteOnly:
+ case wgpu::StorageTextureAccess::ReadWrite:
return D3D12_DESCRIPTOR_RANGE_TYPE_UAV;
- // TODO(dawn:1972): Implement ReadOnly and ReadWrite storage texture
+ // TODO(dawn:1972): Implement ReadOnly storage texture
case wgpu::StorageTextureAccess::ReadOnly:
- case wgpu::StorageTextureAccess::ReadWrite:
case wgpu::StorageTextureAccess::Undefined:
UNREACHABLE();
}
diff --git a/src/dawn/tests/end2end/StorageTextureTests.cpp b/src/dawn/tests/end2end/StorageTextureTests.cpp
index 6ebc829..faa5b03 100644
--- a/src/dawn/tests/end2end/StorageTextureTests.cpp
+++ b/src/dawn/tests/end2end/StorageTextureTests.cpp
@@ -397,7 +397,8 @@
}
wgpu::Texture CreateTextureWithTestData(
- const std::vector<uint8_t>& initialTextureData,
+ const uint8_t* initialTextureData,
+ size_t initialTextureDataSize,
wgpu::TextureFormat format,
wgpu::TextureViewDimension dimension = wgpu::TextureViewDimension::e2D) {
uint32_t texelSize = utils::GetTexelBlockSizeInBytes(format);
@@ -405,7 +406,7 @@
const uint32_t bytesPerTextureRow = texelSize * kWidth;
const uint32_t sliceCount =
- static_cast<uint32_t>(initialTextureData.size() / texelSize / (kWidth * kHeight));
+ static_cast<uint32_t>(initialTextureDataSize / texelSize / (kWidth * kHeight));
const size_t uploadBufferSize =
kTextureBytesPerRowAlignment * (kHeight * sliceCount - 1) + kWidth * bytesPerTextureRow;
@@ -427,7 +428,9 @@
wgpu::BufferUsage::CopySrc | wgpu::BufferUsage::CopyDst);
wgpu::Texture outputTexture = CreateTexture(
- format, wgpu::TextureUsage::StorageBinding | wgpu::TextureUsage::CopyDst,
+ format,
+ wgpu::TextureUsage::StorageBinding | wgpu::TextureUsage::CopySrc |
+ wgpu::TextureUsage::CopyDst,
{kWidth, kHeight, sliceCount}, utils::ViewDimensionToTextureDimension(dimension));
wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
@@ -605,13 +608,15 @@
wgpu::TextureFormat format,
const wgpu::Extent3D& size) {
const std::vector<uint8_t>& expectedData = GetExpectedData(format, size.depthOrArrayLayers);
- CheckOutputStorageTexture(writeonlyStorageTexture, format, size, expectedData);
+ CheckOutputStorageTexture(writeonlyStorageTexture, format, size, expectedData.data(),
+ expectedData.size());
}
void CheckOutputStorageTexture(wgpu::Texture writeonlyStorageTexture,
wgpu::TextureFormat format,
const wgpu::Extent3D& size,
- const std::vector<uint8_t>& expectedData) {
+ const uint8_t* expectedData,
+ size_t expectedDataSize) {
// Copy the content from the write-only storage texture to the result buffer.
wgpu::BufferDescriptor descriptor;
descriptor.size =
@@ -640,7 +645,7 @@
kTextureBytesPerRowAlignment * (size.height * z + y);
const size_t expectedDataOffset = texelSize * size.width * (size.height * z + y);
EXPECT_BUFFER_U32_RANGE_EQ(
- reinterpret_cast<const uint32_t*>(expectedData.data() + expectedDataOffset),
+ reinterpret_cast<const uint32_t*>(expectedData + expectedDataOffset),
resultBuffer, resultBufferOffset, texelSize);
}
}
@@ -985,8 +990,9 @@
WriteIntoStorageTextureInRenderPass(writeonlyStorageTexture, kSimpleVertexShader,
kCommonWriteOnlyZeroInitTestCodeFragment);
+ std::vector<uint8_t> expectedData = GetExpectedData();
CheckOutputStorageTexture(writeonlyStorageTexture, wgpu::TextureFormat::R32Uint,
- {kWidth, kHeight}, GetExpectedData());
+ {kWidth, kHeight}, expectedData.data(), expectedData.size());
}
// Verify that the texture is correctly cleared to 0 before its first usage as a write-only storage
@@ -999,8 +1005,9 @@
WriteIntoStorageTextureInComputePass(writeonlyStorageTexture,
kCommonWriteOnlyZeroInitTestCodeCompute);
+ std::vector<uint8_t> expectedData = GetExpectedData();
CheckOutputStorageTexture(writeonlyStorageTexture, wgpu::TextureFormat::R32Uint,
- {kWidth, kHeight}, GetExpectedData());
+ {kWidth, kHeight}, expectedData.data(), expectedData.size());
}
DAWN_INSTANTIATE_TEST(StorageTextureZeroInitTests,
@@ -1011,5 +1018,145 @@
MetalBackend({"nonzero_clear_resources_on_creation_for_testing"}),
VulkanBackend({"nonzero_clear_resources_on_creation_for_testing"}));
+class ReadWriteStorageTextureTests : public StorageTextureTests {
+ public:
+ std::vector<wgpu::FeatureName> GetRequiredFeatures() override {
+ if (SupportsFeatures({wgpu::FeatureName::ChromiumExperimentalReadWriteStorageTexture})) {
+ mIsReadWriteStorageTextureSupported = true;
+ return {wgpu::FeatureName::ChromiumExperimentalReadWriteStorageTexture};
+ } else {
+ mIsReadWriteStorageTextureSupported = false;
+ return {};
+ }
+ }
+
+ bool IsReadWriteStorageTextureSupported() { return mIsReadWriteStorageTextureSupported; }
+
+ private:
+ bool mIsReadWriteStorageTextureSupported = false;
+};
+
+// Verify read-write storage texture can work correctly in compute shaders.
+TEST_P(ReadWriteStorageTextureTests, ReadWriteStorageTextureInComputeShader) {
+ DAWN_TEST_UNSUPPORTED_IF(!IsReadWriteStorageTextureSupported());
+
+ std::array<uint32_t, kWidth * kHeight> inputData;
+ std::array<uint32_t, kWidth * kHeight> expectedData;
+ for (size_t i = 0; i < inputData.size(); ++i) {
+ inputData[i] = i + 1;
+ expectedData[i] = inputData[i] * 2;
+ }
+
+ wgpu::Texture readWriteStorageTexture = CreateTextureWithTestData(
+ reinterpret_cast<const uint8_t*>(inputData.data()), inputData.size() * sizeof(uint32_t),
+ wgpu::TextureFormat::R32Uint);
+
+ std::ostringstream sstream;
+ sstream << R"(
+enable chromium_experimental_read_write_storage_texture;
+@group(0) @binding(0) var rwImage : texture_storage_2d<r32uint, read_write>;
+
+@compute @workgroup_size()"
+ << kWidth << ", " << kHeight << R"()
+fn main(@builtin(local_invocation_id) local_id: vec3<u32>,) {
+ var data1 = textureLoad(rwImage, vec2i(local_id.xy));
+ data1.x = data1.x * 2;
+ textureStore(rwImage, vec2i(local_id.xy), data1);
+})";
+
+ wgpu::ComputePipeline pipeline = CreateComputePipeline(sstream.str().c_str());
+ wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
+ {{0, readWriteStorageTexture.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(readWriteStorageTexture, wgpu::TextureFormat::R32Uint,
+ {kWidth, kHeight},
+ reinterpret_cast<const uint8_t*>(expectedData.data()),
+ expectedData.size() * sizeof(uint32_t));
+}
+
+// Verify read-write storage texture can work correctly in fragment shaders.
+TEST_P(ReadWriteStorageTextureTests, ReadWriteStorageTextureInFragmentShader) {
+ DAWN_TEST_UNSUPPORTED_IF(!IsReadWriteStorageTextureSupported());
+
+ std::array<uint32_t, kWidth * kHeight> inputData;
+ std::array<uint32_t, kWidth * kHeight> expectedData;
+ for (size_t i = 0; i < inputData.size(); ++i) {
+ inputData[i] = i + 1;
+ expectedData[i] = inputData[i] * 2;
+ }
+
+ wgpu::Texture readWriteStorageTexture = CreateTextureWithTestData(
+ reinterpret_cast<const uint8_t*>(inputData.data()), inputData.size() * sizeof(uint32_t),
+ wgpu::TextureFormat::R32Uint);
+
+ wgpu::TextureDescriptor colorTextureDescriptor;
+ colorTextureDescriptor.format = wgpu::TextureFormat::RGBA8Unorm;
+ colorTextureDescriptor.size = {kWidth, kHeight, 1};
+ colorTextureDescriptor.usage = wgpu::TextureUsage::RenderAttachment;
+ wgpu::Texture dummyColorTexture = device.CreateTexture(&colorTextureDescriptor);
+
+ wgpu::ShaderModule vsModule = utils::CreateShaderModule(device, R"(
+ @vertex fn main(@builtin(vertex_index) VertexIndex : u32) -> @builtin(position) vec4f {
+ var pos = array(
+ vec2f(-2.0, -2.0),
+ vec2f(-2.0, 2.0),
+ vec2f( 2.0, -2.0),
+ vec2f(-2.0, 2.0),
+ vec2f( 2.0, -2.0),
+ vec2f( 2.0, 2.0));
+ return vec4f(pos[VertexIndex], 0.0, 1.0);
+})");
+
+ wgpu::ShaderModule fsModule = utils::CreateShaderModule(device, R"(
+enable chromium_experimental_read_write_storage_texture;
+@group(0) @binding(0) var rwImage : texture_storage_2d<r32uint, read_write>;
+@fragment fn main(@builtin(position) fragcoord: vec4f) -> @location(0) vec4f {
+ var data1 = textureLoad(rwImage, vec2i(fragcoord.xy));
+ data1.x = data1.x * 2;
+ textureStore(rwImage, vec2i(fragcoord.xy), data1);
+ return vec4f(0.0, 1.0, 0.0, 1.0);
+})");
+
+ utils::ComboRenderPipelineDescriptor pipelineDescriptor;
+ pipelineDescriptor.vertex.module = vsModule;
+ pipelineDescriptor.cFragment.module = fsModule;
+ pipelineDescriptor.cTargets[0].format = colorTextureDescriptor.format;
+ wgpu::RenderPipeline renderPipeline = device.CreateRenderPipeline(&pipelineDescriptor);
+
+ wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, renderPipeline.GetBindGroupLayout(0),
+ {{0, readWriteStorageTexture.CreateView()}});
+
+ utils::ComboRenderPassDescriptor renderPassDescriptor({dummyColorTexture.CreateView()});
+
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::RenderPassEncoder renderPassEncoder = encoder.BeginRenderPass(&renderPassDescriptor);
+ renderPassEncoder.SetBindGroup(0, bindGroup);
+ renderPassEncoder.SetPipeline(renderPipeline);
+ renderPassEncoder.Draw(6);
+ renderPassEncoder.End();
+ wgpu::CommandBuffer commandBuffer = encoder.Finish();
+ queue.Submit(1, &commandBuffer);
+
+ CheckOutputStorageTexture(readWriteStorageTexture, wgpu::TextureFormat::R32Uint,
+ {kWidth, kHeight},
+ reinterpret_cast<const uint8_t*>(expectedData.data()),
+ expectedData.size() * sizeof(uint32_t));
+}
+
+// TODO(dawn:1972): Support ReadWrite storage texture access on D3D11 and OpenGL backends
+DAWN_INSTANTIATE_TEST(ReadWriteStorageTextureTests,
+ D3D12Backend(),
+ MetalBackend(),
+ VulkanBackend());
+
} // anonymous namespace
} // namespace dawn