Test for TEXTURE_BINDING vs STORAGE_BINDING issue
The GL backend has an issue if you used a texture as a TEXTURE_BINDING
with a limited subview and then use the same texture has a
STORAGE_BINDING with a different subview. In the backend the first
usage sets TEXTURE_BASE_LEVEL and TEXTURE_MAX_LEVEL on the texture
and this affected the 2nd usage.
This is just a test for this issue. It's suppressed because it fails
on various backends.
Bug: 392121637,392121643,392121648
Change-Id: I2d048ee3d6a065ffb1e1f397f439c335ce96f4b7
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/223594
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Auto-Submit: Gregg Tavares <gman@chromium.org>
Commit-Queue: Gregg Tavares <gman@chromium.org>
diff --git a/src/dawn/tests/end2end/StorageTextureTests.cpp b/src/dawn/tests/end2end/StorageTextureTests.cpp
index 3aa99cb..f6868bb 100644
--- a/src/dawn/tests/end2end/StorageTextureTests.cpp
+++ b/src/dawn/tests/end2end/StorageTextureTests.cpp
@@ -646,19 +646,28 @@
queue.Submit(1, &commandBuffer);
}
- void CheckOutputStorageTexture(wgpu::Texture writeonlyStorageTexture,
+ void CheckOutputStorageTexture(wgpu::Texture storageTexture,
wgpu::TextureFormat format,
const wgpu::Extent3D& size) {
const std::vector<uint8_t>& expectedData = GetExpectedData(format, size.depthOrArrayLayers);
- CheckOutputStorageTexture(writeonlyStorageTexture, format, size, expectedData.data(),
+ CheckOutputStorageTexture(storageTexture, format, size, expectedData.data(),
expectedData.size());
}
- void CheckOutputStorageTexture(wgpu::Texture writeonlyStorageTexture,
+ void CheckOutputStorageTexture(wgpu::Texture storageTexture,
wgpu::TextureFormat format,
const wgpu::Extent3D& size,
const uint8_t* expectedData,
size_t expectedDataSize) {
+ CheckOutputStorageTexture(storageTexture, format, 0, size, expectedData, expectedDataSize);
+ }
+
+ void CheckOutputStorageTexture(wgpu::Texture storageTexture,
+ wgpu::TextureFormat format,
+ uint32_t mipLevel,
+ const wgpu::Extent3D& size,
+ 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 =
@@ -669,7 +678,7 @@
wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
{
wgpu::ImageCopyTexture imageCopyTexture =
- utils::CreateImageCopyTexture(writeonlyStorageTexture, 0, {0, 0, 0});
+ utils::CreateImageCopyTexture(storageTexture, mipLevel, {0, 0, 0});
wgpu::ImageCopyBuffer imageCopyBuffer = utils::CreateImageCopyBuffer(
resultBuffer, 0, kTextureBytesPerRowAlignment, size.height);
encoder.CopyTextureToBuffer(&imageCopyTexture, &imageCopyBuffer, &size);
@@ -1420,6 +1429,296 @@
expectedData.size() * sizeof(uint32_t));
}
+// Tests reading from mip level 0 of a mipLevelCount = 3 texture using a TEXTURE_BINDING
+// and then writing to mip level 1 as STORAGE_BINDING. This surfaced an issue in the GL
+// backend where the first usage sets TEXTURE_BASE_LEVEL and TEXTURE_MAX_LEVEL on the texture
+// and they end up affecting the 2nd usage.
+TEST_P(ReadWriteStorageTextureTests, ReadMipLevel0WriteMipLevel1) {
+ // https://crbug.com/392121637
+ DAWN_SUPPRESS_TEST_IF(IsOpenGL() || IsOpenGLES());
+
+ wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"(
+ @binding(0) @group(0) var<storage, read_write> buf : array<f32>;
+ @binding(1) @group(0) var t_in: texture_2d<f32>;
+ @binding(2) @group(0) var t_out: texture_storage_2d<rgba8unorm, write>;
+
+ @compute @workgroup_size(1) fn csLoad() {
+ // just make sure we actually read (don't want this optimized out)
+ buf[0] = textureLoad(t_in, vec2u(0), 0).w;
+ }
+ @compute @workgroup_size(1) fn csStore() {
+ textureStore(t_out, vec2u(0), vec4f(64, 128, 192, 255) / 255);
+ }
+ )");
+
+ wgpu::ComputePipelineDescriptor pipelineDescriptor;
+ pipelineDescriptor.layout = nullptr;
+ pipelineDescriptor.compute.module = csModule;
+
+ pipelineDescriptor.compute.entryPoint = "csLoad";
+ wgpu::ComputePipeline loadPipeline = device.CreateComputePipeline(&pipelineDescriptor);
+
+ pipelineDescriptor.compute.entryPoint = "csStore";
+ wgpu::ComputePipeline storePipeline = device.CreateComputePipeline(&pipelineDescriptor);
+
+ wgpu::BufferDescriptor bufferDesc;
+ bufferDesc.size = 16;
+ bufferDesc.usage = wgpu::BufferUsage::Storage;
+ wgpu::Buffer storageBuffer = device.CreateBuffer(&bufferDesc);
+
+ bufferDesc.size = 4;
+ bufferDesc.usage = wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::CopySrc;
+ wgpu::Buffer resultBuffer0 = device.CreateBuffer(&bufferDesc);
+ wgpu::Buffer resultBuffer1 = device.CreateBuffer(&bufferDesc);
+
+ wgpu::TextureDescriptor textureDesc;
+ textureDesc.format = wgpu::TextureFormat::RGBA8Unorm;
+ textureDesc.size = {2, 1};
+ textureDesc.mipLevelCount = 2;
+ textureDesc.usage = wgpu::TextureUsage::StorageBinding | wgpu::TextureUsage::TextureBinding |
+ wgpu::TextureUsage::CopySrc;
+ wgpu::Texture texture = device.CreateTexture(&textureDesc);
+
+ wgpu::TextureViewDescriptor textureViewDesc1;
+ textureViewDesc1.baseMipLevel = 0;
+ textureViewDesc1.mipLevelCount = 1;
+ wgpu::BindGroup loadBindGroup =
+ utils::MakeBindGroup(device, loadPipeline.GetBindGroupLayout(0),
+ {{1, texture.CreateView(&textureViewDesc1)}, {0, storageBuffer}});
+
+ wgpu::TextureViewDescriptor textureViewDesc2;
+ textureViewDesc2.baseMipLevel = 1;
+ textureViewDesc2.mipLevelCount = 1;
+ wgpu::BindGroup storeBindGroup = utils::MakeBindGroup(
+ device, storePipeline.GetBindGroupLayout(0), {{2, texture.CreateView(&textureViewDesc2)}});
+
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder computeEncoder = encoder.BeginComputePass();
+
+ computeEncoder.SetBindGroup(0, loadBindGroup);
+ computeEncoder.SetPipeline(loadPipeline);
+ computeEncoder.DispatchWorkgroups(1);
+
+ computeEncoder.SetBindGroup(0, storeBindGroup);
+ computeEncoder.SetPipeline(storePipeline);
+ computeEncoder.DispatchWorkgroups(1);
+
+ computeEncoder.End();
+
+ {
+ wgpu::ImageCopyTexture imageCopyTexture =
+ utils::CreateImageCopyTexture(texture, 0, {0, 0, 0});
+ wgpu::ImageCopyBuffer imageCopyBuffer =
+ utils::CreateImageCopyBuffer(resultBuffer0, 0, 256, 1);
+ wgpu::Extent3D size({1, 1, 1});
+ encoder.CopyTextureToBuffer(&imageCopyTexture, &imageCopyBuffer, &size);
+ }
+
+ {
+ wgpu::ImageCopyTexture imageCopyTexture =
+ utils::CreateImageCopyTexture(texture, 1, {0, 0, 0});
+ wgpu::ImageCopyBuffer imageCopyBuffer =
+ utils::CreateImageCopyBuffer(resultBuffer1, 0, 256, 1);
+ wgpu::Extent3D size({1, 1, 1});
+ encoder.CopyTextureToBuffer(&imageCopyTexture, &imageCopyBuffer, &size);
+ }
+
+ wgpu::CommandBuffer commandBuffer = encoder.Finish();
+ queue.Submit(1, &commandBuffer);
+
+ static uint8_t expectedData0[]{0, 0, 0, 0};
+ EXPECT_BUFFER_U8_RANGE_EQ(expectedData0, resultBuffer0, 0, 4);
+
+ static uint8_t expectedData1[]{64, 128, 192, 255};
+ EXPECT_BUFFER_U8_RANGE_EQ(expectedData1, resultBuffer1, 0, 4);
+}
+
+// Tests reading from both a TEXTURE_BINDING and a STORAGE_BINDING from the same
+// texture at the same time. This test is to double check on a workaround for
+// fixing the previous test above where in the GL backend we try to reset the
+// TEXTURE_BASE_LEVEL and TEXTURE_MAX_LEVEL. If we mistakenly apply the workaround
+// to read only textures then this test will fail.
+TEST_P(ReadWriteStorageTextureTests, ReadMipLevel2AsBothTextureBindingAndStorageBinding) {
+ // This asserts in TextureVK.cpp, see https://crbug.com/392121643
+ DAWN_SUPPRESS_TEST_IF(IsVulkan());
+ // https://crbug.com/392121648
+ DAWN_SUPPRESS_TEST_IF(IsANGLED3D11());
+
+ wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"(
+ @binding(0) @group(0) var<storage, read_write> buf : array<vec4u>;
+ @binding(1) @group(0) var t_in: texture_2d<f32>;
+ @binding(2) @group(0) var s_in: texture_storage_2d<rgba8unorm, read>;
+
+ @compute @workgroup_size(1) fn cs() {
+ buf[0] = vec4u(
+ u32(textureLoad(t_in, vec2u(0), 0).r * 255),
+ u32(textureLoad(s_in, vec2u(0)).r * 255),
+ 123,
+ 456,
+ );
+ }
+ )");
+
+ wgpu::ComputePipelineDescriptor pipelineDescriptor;
+ pipelineDescriptor.layout = nullptr;
+ pipelineDescriptor.compute.module = csModule;
+ wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDescriptor);
+
+ wgpu::BufferDescriptor bufferDesc;
+ bufferDesc.size = 16;
+ bufferDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc;
+ wgpu::Buffer storageBuffer = device.CreateBuffer(&bufferDesc);
+
+ // make a 3 mip level texture
+ wgpu::TextureDescriptor textureDesc;
+ textureDesc.format = wgpu::TextureFormat::RGBA8Unorm;
+ textureDesc.size = {4, 1};
+ textureDesc.mipLevelCount = 3;
+ textureDesc.usage = wgpu::TextureUsage::StorageBinding | wgpu::TextureUsage::TextureBinding |
+ wgpu::TextureUsage::CopySrc | wgpu::TextureUsage::CopyDst;
+ wgpu::Texture texture = device.CreateTexture(&textureDesc);
+
+ // put 1 in first mip, 2 in 2nd, 3 in 3rd.
+ for (uint32_t mipLevel = 0; mipLevel < 3; ++mipLevel) {
+ uint32_t width = 4 >> mipLevel;
+ uint32_t bytesPerRow = width * 4;
+ wgpu::Extent3D copySize({width, 1, 1});
+ wgpu::ImageCopyTexture imageCopyTexture =
+ utils::CreateImageCopyTexture(texture, mipLevel, {0, 0, 0});
+ wgpu::TextureDataLayout textureDataLayout = utils::CreateTextureDataLayout(0, bytesPerRow);
+ std::vector<uint8_t> data(bytesPerRow, mipLevel + 1);
+ queue.WriteTexture(&imageCopyTexture, data.data(), bytesPerRow, &textureDataLayout,
+ ©Size);
+ }
+
+ // View mip level 2
+ wgpu::TextureViewDescriptor textureViewDesc;
+ textureViewDesc.baseMipLevel = 2;
+ textureViewDesc.mipLevelCount = 1;
+ wgpu::TextureView view = texture.CreateView(&textureViewDesc);
+ wgpu::BindGroup bindGroup = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
+ {{0, storageBuffer}, {1, view}, {2, view}});
+
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder computeEncoder = encoder.BeginComputePass();
+
+ computeEncoder.SetBindGroup(0, bindGroup);
+ computeEncoder.SetPipeline(pipeline);
+ computeEncoder.DispatchWorkgroups(1);
+
+ computeEncoder.End();
+
+ wgpu::CommandBuffer commandBuffer = encoder.Finish();
+ queue.Submit(1, &commandBuffer);
+
+ // expect 3 from reading through the texture binding and
+ // also 3 from reading through the storage binding.
+ static uint32_t expectedData[]{3, 3, 123, 456};
+ EXPECT_BUFFER_U32_RANGE_EQ(expectedData, storageBuffer, 0, 4);
+}
+
+// Tests reading from mip level 1 via TEXTURE_BINDING and write to mip level 2 via
+// STORAGE_BINDING at the same time.
+TEST_P(ReadWriteStorageTextureTests, ReadMipLevel1AndWriteLevel2AtTheSameTime) {
+ // Compat mode doesn't support different views of the same texture
+ DAWN_SUPPRESS_TEST_IF(IsCompatibilityMode());
+
+ wgpu::ShaderModule csModule = utils::CreateShaderModule(device, R"(
+ @binding(0) @group(0) var<storage, read_write> buf : array<vec4u>;
+ @binding(1) @group(0) var t_in: texture_2d<f32>;
+ @binding(2) @group(0) var s_out: texture_storage_2d<rgba8unorm, write>;
+
+ @compute @workgroup_size(1) fn cs() {
+ buf[0] = vec4u(textureLoad(t_in, vec2u(0), 0) * 255);
+ textureStore(s_out, vec2u(0), vec4f(64, 128, 192, 255) / 255);
+ }
+ )");
+
+ wgpu::ComputePipelineDescriptor pipelineDescriptor;
+ pipelineDescriptor.layout = nullptr;
+ pipelineDescriptor.compute.module = csModule;
+ wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&pipelineDescriptor);
+
+ wgpu::BufferDescriptor bufferDesc;
+ bufferDesc.size = 16;
+ bufferDesc.usage = wgpu::BufferUsage::Storage | wgpu::BufferUsage::CopySrc;
+ wgpu::Buffer storageBuffer = device.CreateBuffer(&bufferDesc);
+
+ bufferDesc.size = 16;
+ bufferDesc.usage = wgpu::BufferUsage::CopyDst | wgpu::BufferUsage::CopySrc;
+ wgpu::Buffer resultBuffer = device.CreateBuffer(&bufferDesc);
+
+ // make a 3 mip level texture
+ wgpu::TextureDescriptor textureDesc;
+ textureDesc.format = wgpu::TextureFormat::RGBA8Unorm;
+ textureDesc.size = {4, 1};
+ textureDesc.mipLevelCount = 3;
+ textureDesc.usage = wgpu::TextureUsage::StorageBinding | wgpu::TextureUsage::TextureBinding |
+ wgpu::TextureUsage::CopySrc | wgpu::TextureUsage::CopyDst;
+ wgpu::Texture texture = device.CreateTexture(&textureDesc);
+
+ // put 1 in first mip, 2 in 2nd, 3 in 3rd.
+ for (uint32_t mipLevel = 0; mipLevel < 3; ++mipLevel) {
+ uint32_t width = 4 >> mipLevel;
+ uint32_t bytesPerRow = width * 4;
+ wgpu::Extent3D copySize({width, 1, 1});
+ wgpu::ImageCopyTexture imageCopyTexture =
+ utils::CreateImageCopyTexture(texture, mipLevel, {0, 0, 0});
+ wgpu::TextureDataLayout textureDataLayout = utils::CreateTextureDataLayout(0, bytesPerRow);
+ std::vector<uint8_t> data(bytesPerRow, mipLevel + 1);
+ queue.WriteTexture(&imageCopyTexture, data.data(), bytesPerRow, &textureDataLayout,
+ ©Size);
+ }
+
+ // View mip level 1
+ wgpu::TextureViewDescriptor textureViewDesc;
+ textureViewDesc.baseMipLevel = 1;
+ textureViewDesc.mipLevelCount = 1;
+ wgpu::TextureView viewL1 = texture.CreateView(&textureViewDesc);
+
+ // View mip level 2
+ textureViewDesc.baseMipLevel = 2;
+ textureViewDesc.mipLevelCount = 1;
+ wgpu::TextureView viewL2 = texture.CreateView(&textureViewDesc);
+
+ wgpu::BindGroup bindGroup = utils::MakeBindGroup(
+ device, pipeline.GetBindGroupLayout(0), {{0, storageBuffer}, {1, viewL1}, {2, viewL2}});
+
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder computeEncoder = encoder.BeginComputePass();
+
+ computeEncoder.SetBindGroup(0, bindGroup);
+ computeEncoder.SetPipeline(pipeline);
+ computeEncoder.DispatchWorkgroups(1);
+
+ computeEncoder.End();
+
+ // copy a texel from mip level 2
+ {
+ wgpu::ImageCopyTexture imageCopyTexture =
+ utils::CreateImageCopyTexture(texture, 2, {0, 0, 0});
+ wgpu::ImageCopyBuffer imageCopyBuffer =
+ utils::CreateImageCopyBuffer(resultBuffer, 0, 256, 1);
+ wgpu::Extent3D size({1, 1, 1});
+ encoder.CopyTextureToBuffer(&imageCopyTexture, &imageCopyBuffer, &size);
+ }
+
+ wgpu::CommandBuffer commandBuffer = encoder.Finish();
+ queue.Submit(1, &commandBuffer);
+
+ // expect 2 from reading mip level 1.
+ {
+ static uint32_t expectedData[]{2, 2, 2, 2};
+ EXPECT_BUFFER_U32_RANGE_EQ(expectedData, storageBuffer, 0, 4);
+ }
+
+ // expect 0.25, 0.5, 0.75, 1 in mip level 2
+ {
+ static uint8_t expectedData[]{64, 128, 192, 255};
+ EXPECT_BUFFER_U8_RANGE_EQ(expectedData, resultBuffer, 0, 4);
+ }
+}
+
DAWN_INSTANTIATE_TEST(ReadWriteStorageTextureTests,
D3D11Backend(),
D3D12Backend(),
