Dawn: Add the r8unorm-storage feature.
This is a feature for use in Graphite to make compute shaders render
coverage masks. It will be superseeded by some storage texture tier in
the future.
Fixed: dawn:2101
Change-Id: I6fa127bd1d2db59816a2e7013c7ca11cc3ff0b7f
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/177543
Commit-Queue: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/docs/dawn/features/r8unorm_storage.md b/docs/dawn/features/r8unorm_storage.md
new file mode 100644
index 0000000..11123fe
--- /dev/null
+++ b/docs/dawn/features/r8unorm_storage.md
@@ -0,0 +1,7 @@
+# R8UnormStorage
+
+Adds the ability to read or write to r8unorm textures as storage for Graphite (but not both at the same time).
+
+Usage in WGSL requires the use of the `chromium_internal_graphite` enable.
+
+Note that in the future this will be merged in some form of storage texture support tier.
diff --git a/src/dawn/dawn.json b/src/dawn/dawn.json
index 5864f39..fe7c864 100644
--- a/src/dawn/dawn.json
+++ b/src/dawn/dawn.json
@@ -2189,6 +2189,7 @@
{"value": 1026, "name": "buffer map extended usages", "tags": ["dawn"]},
{"value": 1027, "name": "adapter properties memory heaps", "tags": ["dawn"]},
{"value": 1028, "name": "adapter properties D3D", "tags": ["dawn"]},
+ {"value": 1029, "name": "r8 unorm storage", "tags": ["dawn"]},
{"value": 1100, "name": "shared texture memory vk dedicated allocation", "tags": ["dawn", "native"]},
{"value": 1101, "name": "shared texture memory a hardware buffer", "tags": ["dawn", "native"]},
diff --git a/src/dawn/native/Features.cpp b/src/dawn/native/Features.cpp
index d2b9b06..74a7216 100644
--- a/src/dawn/native/Features.cpp
+++ b/src/dawn/native/Features.cpp
@@ -299,6 +299,11 @@
{"Support importing ID3D12Resource as shared buffer memory.",
"https://dawn.googlesource.com/dawn/+/refs/heads/main/docs/dawn/features/shared_buffer.md",
FeatureInfo::FeatureState::Experimental}},
+ {Feature::R8UnormStorage,
+ {"Supports using r8unorm texture as storage texture.",
+ "https://dawn.googlesource.com/dawn/+/refs/heads/main/docs/dawn/features/"
+ "r8unorm_storage.md",
+ FeatureInfo::FeatureState::Experimental}},
};
} // anonymous namespace
diff --git a/src/dawn/native/Format.cpp b/src/dawn/native/Format.cpp
index 6ce2515..91705d5 100644
--- a/src/dawn/native/Format.cpp
+++ b/src/dawn/native/Format.cpp
@@ -427,7 +427,8 @@
// clang-format off
// 1 byte color formats
- AddColorFormat(wgpu::TextureFormat::R8Unorm, Cap::Renderable | Cap::Multisample | Cap::Resolve, ByteSize(1), kAnyFloat, ComponentCount(1), RenderTargetPixelByteCost(1), RenderTargetComponentAlignment(1));
+ auto r8unormSupportsStorage = device->HasFeature(Feature::R8UnormStorage) ? Cap::StorageW : Cap::None;
+ AddColorFormat(wgpu::TextureFormat::R8Unorm, Cap::Renderable | Cap::Multisample | Cap::Resolve | r8unormSupportsStorage, ByteSize(1), kAnyFloat, ComponentCount(1), RenderTargetPixelByteCost(1), RenderTargetComponentAlignment(1));
AddColorFormat(wgpu::TextureFormat::R8Snorm, Cap::None, ByteSize(1), kAnyFloat, ComponentCount(1));
AddColorFormat(wgpu::TextureFormat::R8Uint, Cap::Renderable | Cap::Multisample, ByteSize(1), SampleTypeBit::Uint, ComponentCount(1), RenderTargetPixelByteCost(1), RenderTargetComponentAlignment(1));
AddColorFormat(wgpu::TextureFormat::R8Sint, Cap::Renderable | Cap::Multisample, ByteSize(1), SampleTypeBit::Sint, ComponentCount(1), RenderTargetPixelByteCost(1), RenderTargetComponentAlignment(1));
diff --git a/src/dawn/native/ShaderModule.cpp b/src/dawn/native/ShaderModule.cpp
index b71516c..d0857a5 100644
--- a/src/dawn/native/ShaderModule.cpp
+++ b/src/dawn/native/ShaderModule.cpp
@@ -132,6 +132,8 @@
return wgpu::TextureFormat::RGBA32Sint;
case tint::inspector::ResourceBinding::TexelFormat::kRgba32Float:
return wgpu::TextureFormat::RGBA32Float;
+ case tint::inspector::ResourceBinding::TexelFormat::kR8Unorm:
+ return wgpu::TextureFormat::R8Unorm;
case tint::inspector::ResourceBinding::TexelFormat::kNone:
return wgpu::TextureFormat::Undefined;
diff --git a/src/dawn/native/d3d11/PhysicalDeviceD3D11.cpp b/src/dawn/native/d3d11/PhysicalDeviceD3D11.cpp
index 8480dd7..f6f0863 100644
--- a/src/dawn/native/d3d11/PhysicalDeviceD3D11.cpp
+++ b/src/dawn/native/d3d11/PhysicalDeviceD3D11.cpp
@@ -191,6 +191,7 @@
EnableFeature(Feature::Norm16TextureFormats);
EnableFeature(Feature::AdapterPropertiesMemoryHeaps);
EnableFeature(Feature::AdapterPropertiesD3D);
+ EnableFeature(Feature::R8UnormStorage);
// Multi planar formats are always supported since Feature Level 11.0
// https://learn.microsoft.com/en-us/windows/win32/direct3ddxgi/format-support-for-direct3d-11-0-feature-level-hardware
diff --git a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
index f17def9..fb6cb21 100644
--- a/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
+++ b/src/dawn/native/d3d12/PhysicalDeviceD3D12.cpp
@@ -146,6 +146,7 @@
EnableFeature(Feature::AdapterPropertiesMemoryHeaps);
EnableFeature(Feature::AdapterPropertiesD3D);
EnableFeature(Feature::MultiPlanarRenderTargets);
+ EnableFeature(Feature::R8UnormStorage);
if (AreTimestampQueriesSupported()) {
EnableFeature(Feature::TimestampQuery);
diff --git a/src/dawn/native/metal/BackendMTL.mm b/src/dawn/native/metal/BackendMTL.mm
index 398c40a..c4c8314 100644
--- a/src/dawn/native/metal/BackendMTL.mm
+++ b/src/dawn/native/metal/BackendMTL.mm
@@ -616,6 +616,7 @@
EnableFeature(Feature::SurfaceCapabilities);
EnableFeature(Feature::MSAARenderToSingleSampled);
EnableFeature(Feature::DualSourceBlending);
+ EnableFeature(Feature::R8UnormStorage);
// SIMD-scoped permute operations is supported by GPU family Metal3, Apple6, Apple7, Apple8,
// and Mac2.
diff --git a/src/dawn/native/vulkan/DeviceVk.cpp b/src/dawn/native/vulkan/DeviceVk.cpp
index e404a84..9d3218b 100644
--- a/src/dawn/native/vulkan/DeviceVk.cpp
+++ b/src/dawn/native/vulkan/DeviceVk.cpp
@@ -474,6 +474,10 @@
usedKnobs.features.dualSrcBlend = VK_TRUE;
}
+ if (HasFeature(Feature::R8UnormStorage)) {
+ usedKnobs.features.shaderStorageImageExtendedFormats = VK_TRUE;
+ }
+
if (IsRobustnessEnabled() && mDeviceInfo.HasExt(DeviceExt::Robustness2)) {
DAWN_ASSERT(usedKnobs.HasExt(DeviceExt::Robustness2));
diff --git a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
index 5e4fb35..1d59695 100644
--- a/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
+++ b/src/dawn/native/vulkan/PhysicalDeviceVk.cpp
@@ -240,6 +240,10 @@
EnableFeature(Feature::DualSourceBlending);
}
+ if (mDeviceInfo.features.shaderStorageImageExtendedFormats == VK_TRUE) {
+ EnableFeature(Feature::R8UnormStorage);
+ }
+
if (mDeviceInfo.HasExt(DeviceExt::ShaderFloat16Int8) &&
mDeviceInfo.HasExt(DeviceExt::_16BitStorage) &&
mDeviceInfo.shaderFloat16Int8Features.shaderFloat16 == VK_TRUE &&
diff --git a/src/dawn/node/binding/Converter.cpp b/src/dawn/node/binding/Converter.cpp
index b91a63c..4ef2d00 100644
--- a/src/dawn/node/binding/Converter.cpp
+++ b/src/dawn/node/binding/Converter.cpp
@@ -1529,6 +1529,8 @@
case wgpu::FeatureName::FramebufferFetch:
case wgpu::FeatureName::BufferMapExtendedUsages:
case wgpu::FeatureName::AdapterPropertiesMemoryHeaps:
+ case wgpu::FeatureName::SharedBufferMemoryD3D12Resource:
+ case wgpu::FeatureName::R8UnormStorage:
case wgpu::FeatureName::Undefined:
return false;
}
diff --git a/src/dawn/tests/end2end/StorageTextureTests.cpp b/src/dawn/tests/end2end/StorageTextureTests.cpp
index 2ca3f9b..2660dc2 100644
--- a/src/dawn/tests/end2end/StorageTextureTests.cpp
+++ b/src/dawn/tests/end2end/StorageTextureTests.cpp
@@ -173,6 +173,12 @@
break;
}
+ case wgpu::TextureFormat::R8Unorm: {
+ uint8_t* valuePtr = static_cast<uint8_t*>(pixelValuePtr);
+ *valuePtr = pixelValue;
+ break;
+ }
+
default:
DAWN_UNREACHABLE();
break;
@@ -260,6 +266,9 @@
return "vec4f(f32(value) / 127.0, -f32(value) / 127.0, "
"f32(value) * 2.0 / 127.0, -f32(value) * 2.0 / 127.0)";
+ case wgpu::TextureFormat::R8Unorm:
+ return "vec4f(f32(value) / 255.0, 0.0, 0.0, 1.0)";
+
default:
DAWN_UNREACHABLE();
break;
@@ -303,6 +312,7 @@
// normalized signed/unsigned integer formats
case wgpu::TextureFormat::RGBA8Unorm:
case wgpu::TextureFormat::RGBA8Snorm:
+ case wgpu::TextureFormat::R8Unorm:
// On Windows Intel drivers the tests will fail if tolerance <= 0.00000001f.
return R"(
fn IsEqualTo(pixel : vec4f, expected : vec4f) -> bool {
@@ -318,6 +328,13 @@
return "";
}
+ const char* GetEnable(wgpu::TextureFormat format) {
+ if (format == wgpu::TextureFormat::R8Unorm) {
+ return "enable chromium_internal_graphite;";
+ }
+ return "";
+ }
+
std::string CommonWriteOnlyTestCode(
const char* stage,
wgpu::TextureFormat format,
@@ -353,6 +370,7 @@
const bool isFragment = strcmp(stage, "fragment") == 0;
std::ostringstream ostream;
+ ostream << GetEnable(format) << "\n";
ostream << GetImageDeclaration(format, "write", dimension, 0) << "\n";
ostream << "@" << stage << workgroupSize << "\n";
ostream << "fn main() ";
@@ -948,6 +966,71 @@
OpenGLESBackend(),
VulkanBackend());
+class R8UnormStorageTextureTests : public StorageTextureTests {
+ public:
+ std::vector<wgpu::FeatureName> GetRequiredFeatures() override {
+ if (SupportsFeatures({wgpu::FeatureName::R8UnormStorage})) {
+ mIsR8UnormStorageSupported = true;
+ return {wgpu::FeatureName::R8UnormStorage};
+ } else {
+ mIsR8UnormStorageSupported = false;
+ return {};
+ }
+ }
+
+ bool IsR8UnormStorageSupported() { return mIsR8UnormStorageSupported; }
+
+ private:
+ bool mIsR8UnormStorageSupported = false;
+};
+
+// Test that R8Unorm is supported to be used as storage texture in compute shaders when the
+// optional feature 'r8unorm-storage' is supported.
+TEST_P(R8UnormStorageTextureTests, WriteonlyStorageTextureInComputeShader) {
+ DAWN_TEST_UNSUPPORTED_IF(!IsR8UnormStorageSupported());
+
+ constexpr wgpu::TextureFormat kFormat = wgpu::TextureFormat::R8Unorm;
+ wgpu::Texture writeonlyStorageTexture =
+ CreateTexture(kFormat, wgpu::TextureUsage::StorageBinding | wgpu::TextureUsage::CopySrc,
+ {kWidth, kHeight});
+
+ // Write the expected pixel values into the write-only storage texture.
+ const std::string computeShader = CommonWriteOnlyTestCode("compute", kFormat);
+ WriteIntoStorageTextureInComputePass(writeonlyStorageTexture, computeShader.c_str());
+
+ // Verify the pixel data in the write-only storage texture is expected.
+ CheckOutputStorageTexture(writeonlyStorageTexture, kFormat, {kWidth, kHeight});
+}
+
+// Test that R8Unorm is supported to be used as storage texture in fragment shaders when the
+// optional feature 'r8unorm-storage' is supported.
+TEST_P(R8UnormStorageTextureTests, WriteonlyStorageTextureInFragmentShader) {
+ DAWN_TEST_UNSUPPORTED_IF(!IsR8UnormStorageSupported());
+
+ constexpr wgpu::TextureFormat kFormat = wgpu::TextureFormat::R8Unorm;
+
+ // Prepare the write-only storage texture.
+ wgpu::Texture writeonlyStorageTexture =
+ CreateTexture(kFormat, wgpu::TextureUsage::StorageBinding | wgpu::TextureUsage::CopySrc,
+ {kWidth, kHeight});
+
+ // Write the expected pixel values into the write-only storage texture.
+ const std::string fragmentShader = CommonWriteOnlyTestCode("fragment", kFormat);
+ WriteIntoStorageTextureInRenderPass(writeonlyStorageTexture, kSimpleVertexShader,
+ fragmentShader.c_str());
+
+ // Verify the pixel data in the write-only storage texture is expected.
+ CheckOutputStorageTexture(writeonlyStorageTexture, kFormat, {kWidth, kHeight});
+}
+
+DAWN_INSTANTIATE_TEST(R8UnormStorageTextureTests,
+ D3D11Backend(),
+ D3D12Backend(),
+ MetalBackend(),
+ OpenGLBackend(),
+ OpenGLESBackend(),
+ VulkanBackend());
+
class StorageTextureZeroInitTests : public StorageTextureTests {
public:
static std::vector<uint8_t> GetExpectedData() {
diff --git a/src/dawn/tests/unittests/validation/StorageTextureValidationTests.cpp b/src/dawn/tests/unittests/validation/StorageTextureValidationTests.cpp
index cfd9bcd..97e8423 100644
--- a/src/dawn/tests/unittests/validation/StorageTextureValidationTests.cpp
+++ b/src/dawn/tests/unittests/validation/StorageTextureValidationTests.cpp
@@ -1216,5 +1216,89 @@
}
}
+class R8UnormStorageValidationTests : public StorageTextureValidationTests {
+ WGPUDevice CreateTestDevice(native::Adapter dawnAdapter,
+ wgpu::DeviceDescriptor descriptor) override {
+ wgpu::FeatureName requiredFeatures[1] = {wgpu::FeatureName::R8UnormStorage};
+ descriptor.requiredFeatures = requiredFeatures;
+ descriptor.requiredFeatureCount = 1;
+ return dawnAdapter.CreateDevice(&descriptor);
+ }
+};
+
+// Check that it is allowed to create an R8Unorm texture with the storage usage.
+TEST_F(R8UnormStorageValidationTests, TextureCreation) {
+ wgpu::TextureDescriptor desc;
+ desc.format = wgpu::TextureFormat::R8Unorm;
+ desc.usage = wgpu::TextureUsage::StorageBinding;
+ desc.size = {1, 1};
+ device.CreateTexture(&desc);
+}
+
+// Check that it is allowed to create a BGL with a read-only or write-only R8unorm storage texture
+// entry.
+TEST_F(R8UnormStorageValidationTests, BGLEntry) {
+ // Control case: read-only or write-only are allowed.
+ utils::MakeBindGroupLayout(
+ device, {{0, wgpu::ShaderStage::Fragment, wgpu::StorageTextureAccess::ReadOnly,
+ wgpu::TextureFormat::R8Unorm}});
+ utils::MakeBindGroupLayout(
+ device, {{0, wgpu::ShaderStage::Fragment, wgpu::StorageTextureAccess::WriteOnly,
+ wgpu::TextureFormat::R8Unorm}});
+
+ // Error cases: read-write is disallowed.
+ ASSERT_DEVICE_ERROR(utils::MakeBindGroupLayout(
+ device, {{0, wgpu::ShaderStage::Fragment, wgpu::StorageTextureAccess::ReadWrite,
+ wgpu::TextureFormat::R8Unorm}}));
+}
+
+// Check that using the `r8unorm` to create a WGSL compute shader is allowed.
+TEST_F(R8UnormStorageValidationTests, ShaderModule) {
+ utils::CreateShaderModule(device, R"(
+ enable chromium_internal_graphite;
+ @group(0) @binding(0) var t : texture_storage_2d<r8unorm, write>;
+ )");
+}
+
+// Check that using an r8unorm storage texture read-only or write-only with implicit layout is
+// valid.
+TEST_F(R8UnormStorageValidationTests, End2endUsage) {
+ wgpu::ComputePipelineDescriptor cDesc;
+ cDesc.compute.module = utils::CreateShaderModule(device, R"(
+ enable chromium_internal_graphite;
+ @group(0) @binding(0) var input : texture_storage_2d<r8unorm, read>;
+ @group(0) @binding(1) var output : texture_storage_2d<r8unorm, write>;
+
+ @workgroup_size(4, 4) @compute fn main(@builtin(local_invocation_id) id : vec3<u32>) {
+ textureStore(output, id.xy, 2 * textureLoad(input, id.xy));
+ }
+ )");
+ wgpu::ComputePipeline pipeline = device.CreateComputePipeline(&cDesc);
+
+ wgpu::TextureDescriptor tDesc;
+ tDesc.format = wgpu::TextureFormat::R8Unorm;
+ tDesc.usage = wgpu::TextureUsage::StorageBinding | wgpu::TextureUsage::CopySrc |
+ wgpu::TextureUsage::CopyDst;
+ tDesc.size = {4, 4};
+ wgpu::Texture input = device.CreateTexture(&tDesc);
+ wgpu::Texture output = device.CreateTexture(&tDesc);
+
+ wgpu::BindGroup bg = utils::MakeBindGroup(device, pipeline.GetBindGroupLayout(0),
+ {
+ {0, input.CreateView()},
+ {1, output.CreateView()},
+ });
+
+ wgpu::CommandEncoder encoder = device.CreateCommandEncoder();
+ wgpu::ComputePassEncoder pass = encoder.BeginComputePass();
+ pass.SetPipeline(pipeline);
+ pass.SetBindGroup(0, bg);
+ pass.DispatchWorkgroups(1);
+ pass.End();
+
+ wgpu::CommandBuffer commands = encoder.Finish();
+ device.GetQueue().Submit(1, &commands);
+}
+
} // anonymous namespace
} // namespace dawn
diff --git a/src/dawn/wire/SupportedFeatures.cpp b/src/dawn/wire/SupportedFeatures.cpp
index 6f72263..5447227 100644
--- a/src/dawn/wire/SupportedFeatures.cpp
+++ b/src/dawn/wire/SupportedFeatures.cpp
@@ -95,6 +95,7 @@
case WGPUFeatureName_FramebufferFetch:
case WGPUFeatureName_AdapterPropertiesMemoryHeaps:
case WGPUFeatureName_AdapterPropertiesD3D:
+ case WGPUFeatureName_R8UnormStorage:
return true;
}
