Add support for bgra8unorm-storage (#4228)

* Add `BGRA8UNORM_STORAGE` extension * Leave a comment in the backends that don't support bgra8unorm-storage * Pass the appropriate storage format to naga * Check for bgra8unorm storage support in the vulkan backend * Add a test Co-authored-by: Jinlei Li <jinleili0@outlook.com> Co-authored-by: Teodor Tanasoaia <28601907+teoxoy@users.noreply.github.com>
2025-02-16 08:53:20 +00:00 · 2023-10-13 12:07:11 +02:00 · 2023-10-13 12:07:11 +02:00 · ff306d20e0
commit ff306d20e0
parent 6c8ccdeaa9
14 changed files with 261 additions and 5 deletions
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@ -115,6 +115,7 @@ By @teoxoy in [#4185](https://github.com/gfx-rs/wgpu/pull/4185)
 - Add `Rgb10a2Uint` format. By @teoxoy in [4199](https://github.com/gfx-rs/wgpu/pull/4199)
 - Validate that resources are used on the right device. By @nical in [4207](https://github.com/gfx-rs/wgpu/pull/4207)
 - Expose instance flags. By @nical in [4230](https://github.com/gfx-rs/wgpu/pull/4230)
+- Add support for the bgra8unorm-storage feature. By @jinleili and @nical in [https://github.com/gfx-rs/wgpu/pull/4228](4228)

 #### Vulkan

--- a/deno_webgpu/01_webgpu.js
+++ b/deno_webgpu/01_webgpu.js
@ -5163,6 +5163,7 @@ webidl.converters["GPUFeatureName"] = webidl.createEnumConverter(
    "texture-compression-etc2",
    "texture-compression-astc",
    "rg11b10ufloat-renderable",
+    "bgra8unorm-storage",

    // extended from spec

--- a/deno_webgpu/lib.rs
+++ b/deno_webgpu/lib.rs
@ -263,6 +263,9 @@ fn deserialize_features(features: &wgpu_types::Features) -> Vec<&'static str> {
    if features.contains(wgpu_types::Features::RG11B10UFLOAT_RENDERABLE) {
        return_features.push("rg11b10ufloat-renderable");
    }
+    if features.contains(wgpu_types::Features::BGRA8UNORM_STORAGE) {
+        return_features.push("bgra8unorm-storage");
+    }

    // extended from spec

@ -491,6 +494,10 @@ impl From<GpuRequiredFeatures> for wgpu_types::Features {
            wgpu_types::Features::RG11B10UFLOAT_RENDERABLE,
            required_features.0.contains("rg11b10ufloat-renderable"),
        );
+        features.set(
+            wgpu_types::Features::BGRA8UNORM_STORAGE,
+            required_features.0.contains("bgra8unorm-storage"),
+        );

        // extended from spec

--- a/deno_webgpu/webgpu.idl
+++ b/deno_webgpu/webgpu.idl
@ -102,6 +102,7 @@ enum GPUFeatureName {
    // shader
    "shader-f16",
    "rg11b10ufloat-renderable",
+    "bgra8unorm-storage",

    // extended from spec

--- a/tests/tests/bgra8unorm_storage.rs
+++ b/tests/tests/bgra8unorm_storage.rs
@ -0,0 +1,158 @@
+//! Tests for texture copy bounds checks.
+
+use std::borrow::Cow;
+
+use wasm_bindgen_test::*;
+use wgpu_test::{initialize_test, TestParameters};
+
+const SHADER_SRC: &str = "
+@group(0) @binding(0) var tex: texture_storage_2d<bgra8unorm, write>;
+@compute @workgroup_size(256)
+fn main(@builtin(workgroup_id) wgid: vec3<u32>) {
+    var texel = vec4f(0.0, 0.0, 1.0, 1.0);
+    textureStore(tex, wgid.xy, texel);
+}
+";
+
+#[test]
+#[wasm_bindgen_test]
+fn bgra8unorm_storage() {
+    let parameters = TestParameters::default()
+        .limits(wgpu::Limits {
+            max_storage_textures_per_shader_stage: 1,
+            ..Default::default()
+        })
+        .features(wgpu::Features::BGRA8UNORM_STORAGE);
+
+    initialize_test(parameters, |ctx| {
+        let device = &ctx.device;
+        let texture = ctx.device.create_texture(&wgpu::TextureDescriptor {
+            label: None,
+            size: wgpu::Extent3d {
+                width: 256,
+                height: 256,
+                depth_or_array_layers: 1,
+            },
+            mip_level_count: 1,
+            sample_count: 1,
+            dimension: wgpu::TextureDimension::D2,
+            format: wgpu::TextureFormat::Bgra8Unorm,
+            usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::COPY_SRC,
+            view_formats: &[],
+        });
+
+        let view = texture.create_view(&wgpu::TextureViewDescriptor {
+            label: None,
+            format: None,
+            dimension: None,
+            aspect: wgpu::TextureAspect::All,
+            base_mip_level: 0,
+            base_array_layer: 0,
+            mip_level_count: Some(1),
+            array_layer_count: Some(1),
+        });
+
+        let readback_buffer = device.create_buffer(&wgpu::BufferDescriptor {
+            label: None,
+            size: 256 * 256 * 4,
+            usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
+            mapped_at_creation: false,
+        });
+
+        let bgl = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: None,
+            entries: &[wgpu::BindGroupLayoutEntry {
+                binding: 0,
+                visibility: wgpu::ShaderStages::COMPUTE,
+                ty: wgpu::BindingType::StorageTexture {
+                    access: wgpu::StorageTextureAccess::WriteOnly,
+                    format: wgpu::TextureFormat::Bgra8Unorm,
+                    view_dimension: wgpu::TextureViewDimension::D2,
+                },
+                count: None,
+            }],
+        });
+
+        let bg = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: None,
+            layout: &bgl,
+            entries: &[wgpu::BindGroupEntry {
+                binding: 0,
+                resource: wgpu::BindingResource::TextureView(&view),
+            }],
+        });
+
+        let pl = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: None,
+            bind_group_layouts: &[&bgl],
+            push_constant_ranges: &[],
+        });
+
+        let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: None,
+            source: wgpu::ShaderSource::Wgsl(Cow::Borrowed(SHADER_SRC)),
+        });
+
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: None,
+            layout: Some(&pl),
+            entry_point: "main",
+            module: &module,
+        });
+
+        let mut encoder =
+            device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+
+        {
+            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: None,
+                timestamp_writes: None,
+            });
+
+            pass.set_bind_group(0, &bg, &[]);
+            pass.set_pipeline(&pipeline);
+            pass.dispatch_workgroups(256, 256, 1);
+        }
+
+        encoder.copy_texture_to_buffer(
+            wgpu::ImageCopyTexture {
+                texture: &texture,
+                mip_level: 0,
+                origin: wgpu::Origin3d { x: 0, y: 0, z: 0 },
+                aspect: wgpu::TextureAspect::All,
+            },
+            wgpu::ImageCopyBuffer {
+                buffer: &readback_buffer,
+                layout: wgpu::ImageDataLayout {
+                    offset: 0,
+                    bytes_per_row: Some(256 * 4),
+                    rows_per_image: Some(256),
+                },
+            },
+            wgpu::Extent3d {
+                width: 256,
+                height: 256,
+                depth_or_array_layers: 1,
+            },
+        );
+
+        ctx.queue.submit(Some(encoder.finish()));
+
+        let buffer_slice = readback_buffer.slice(..);
+        buffer_slice.map_async(wgpu::MapMode::Read, Result::unwrap);
+        device.poll(wgpu::Maintain::Wait);
+
+        {
+            let texels = buffer_slice.get_mapped_range();
+            assert_eq!(texels.len(), 256 * 256 * 4);
+            for texel in texels.chunks(4) {
+                assert_eq!(texel[0], 255); // b
+                assert_eq!(texel[1], 0); // g
+                assert_eq!(texel[2], 0); // r
+                assert_eq!(texel[3], 255); // a
+            }
+        }
+
+        readback_buffer.unmap();
+    });
+}
--- a/tests/tests/root.rs
+++ b/tests/tests/root.rs
@ -6,6 +6,7 @@ mod regression {
    mod issue_4122;
 }

+mod bgra8unorm_storage;
 mod bind_group_layout_dedup;
 mod buffer;
 mod buffer_copy;
--- a/wgpu-core/src/validation.rs
+++ b/wgpu-core/src/validation.rs
@ -295,6 +295,7 @@ fn map_storage_format_to_naga(format: wgt::TextureFormat) -> Option<naga::Storag
        Tf::Rgba8Snorm => Sf::Rgba8Snorm,
        Tf::Rgba8Uint => Sf::Rgba8Uint,
        Tf::Rgba8Sint => Sf::Rgba8Sint,
+        Tf::Bgra8Unorm => Sf::Bgra8Unorm,

        Tf::Rgb10a2Uint => Sf::Rgb10a2Uint,
        Tf::Rgb10a2Unorm => Sf::Rgb10a2Unorm,
--- a/wgpu-hal/src/dx11/adapter.rs
+++ b/wgpu-hal/src/dx11/adapter.rs
@ -152,6 +152,9 @@ impl super::Adapter {
            features |= wgt::Features::VERTEX_WRITABLE_STORAGE;
        }

+        // bgra8unorm-storage is never supported on dx11 according to:
+        // https://learn.microsoft.com/en-us/windows/win32/direct3ddxgi/format-support-for-direct3d-11-0-feature-level-hardware#dxgi_format_b8g8r8a8_unormfcs-87
+
        //
        // Fill out limits and alignments
        //
--- a/wgpu-hal/src/dx12/adapter.rs
+++ b/wgpu-hal/src/dx12/adapter.rs
@ -4,7 +4,9 @@ use crate::{
 };
 use std::{mem, ptr, sync::Arc, thread};
 use winapi::{
-    shared::{dxgi, dxgi1_2, minwindef::DWORD, windef, winerror},
+    shared::{
+        dxgi, dxgi1_2, dxgiformat::DXGI_FORMAT_B8G8R8A8_UNORM, minwindef::DWORD, windef, winerror,
+    },
    um::{d3d12 as d3d12_ty, d3d12sdklayers, winuser},
 };

@ -276,6 +278,25 @@ impl super::Adapter {
            shader_model_support.HighestShaderModel >= d3d12_ty::D3D_SHADER_MODEL_5_1,
        );

+        let bgra8unorm_storage_supported = {
+            let mut bgra8unorm_info: d3d12_ty::D3D12_FEATURE_DATA_FORMAT_SUPPORT =
+                unsafe { mem::zeroed() };
+            bgra8unorm_info.Format = DXGI_FORMAT_B8G8R8A8_UNORM;
+            let hr = unsafe {
+                device.CheckFeatureSupport(
+                    d3d12_ty::D3D12_FEATURE_FORMAT_SUPPORT,
+                    &mut bgra8unorm_info as *mut _ as *mut _,
+                    mem::size_of::<d3d12_ty::D3D12_FEATURE_DATA_FORMAT_SUPPORT>() as _,
+                )
+            };
+            hr == 0
+                && (bgra8unorm_info.Support2 & d3d12_ty::D3D12_FORMAT_SUPPORT2_UAV_TYPED_STORE != 0)
+        };
+        features.set(
+            wgt::Features::BGRA8UNORM_STORAGE,
+            bgra8unorm_storage_supported,
+        );
+
        // TODO: Determine if IPresentationManager is supported
        let presentation_timer = auxil::dxgi::time::PresentationTimer::new_dxgi();

--- a/wgpu-hal/src/gles/adapter.rs
+++ b/wgpu-hal/src/gles/adapter.rs
@ -434,6 +434,8 @@ impl super::Adapter {
            );
        }

+        // We *might* be able to emulate bgra8unorm-storage but currently don't attempt to.
+
        let mut private_caps = super::PrivateCapabilities::empty();
        private_caps.set(
            super::PrivateCapabilities::BUFFER_ALLOCATION,
--- a/wgpu-hal/src/metal/adapter.rs
+++ b/wgpu-hal/src/metal/adapter.rs
@ -825,7 +825,8 @@ impl super::PrivateCapabilities {
            | F::TEXTURE_FORMAT_16BIT_NORM
            | F::SHADER_F16
            | F::DEPTH32FLOAT_STENCIL8
-            | F::MULTI_DRAW_INDIRECT;
+            | F::MULTI_DRAW_INDIRECT
+            | F::BGRA8UNORM_STORAGE;

        features.set(
            F::TIMESTAMP_QUERY,
--- a/wgpu-hal/src/vulkan/adapter.rs
+++ b/wgpu-hal/src/vulkan/adapter.rs
@ -524,6 +524,11 @@ impl PhysicalDeviceFeatures {
        features.set(F::RG11B10UFLOAT_RENDERABLE, rg11b10ufloat_renderable);
        features.set(F::SHADER_UNUSED_VERTEX_OUTPUT, true);

+        features.set(
+            F::BGRA8UNORM_STORAGE,
+            supports_bgra8unorm_storage(instance, phd, caps.effective_api_version),
+        );
+
        (features, dl_flags)
    }

@ -1263,6 +1268,9 @@ impl super::Adapter {
            ) {
                capabilities.push(spv::Capability::ShaderNonUniform);
            }
+            if features.contains(wgt::Features::BGRA8UNORM_STORAGE) {
+                capabilities.push(spv::Capability::StorageImageWriteWithoutFormat);
+            }

            let mut flags = spv::WriterFlags::empty();
            flags.set(
@ -1748,3 +1756,35 @@ fn supports_format(
        _ => false,
    }
 }
+
+fn supports_bgra8unorm_storage(
+    instance: &ash::Instance,
+    phd: vk::PhysicalDevice,
+    api_version: u32,
+) -> bool {
+    // See https://github.com/KhronosGroup/Vulkan-Docs/issues/2027#issuecomment-1380608011
+
+    // This check gates the function call and structures used below.
+    // TODO: check for (`VK_KHR_get_physical_device_properties2` or VK1.1) and (`VK_KHR_format_feature_flags2` or VK1.3).
+    // Right now we only check for VK1.3.
+    if api_version < vk::API_VERSION_1_3 {
+        return false;
+    }
+
+    unsafe {
+        let mut properties3 = vk::FormatProperties3::default();
+        let mut properties2 = vk::FormatProperties2::builder().push_next(&mut properties3);
+
+        instance.get_physical_device_format_properties2(
+            phd,
+            vk::Format::B8G8R8A8_UNORM,
+            &mut properties2,
+        );
+
+        let features2 = properties2.format_properties.optimal_tiling_features;
+        let features3 = properties3.optimal_tiling_features;
+
+        features2.contains(vk::FormatFeatureFlags::STORAGE_IMAGE)
+            && features3.contains(vk::FormatFeatureFlags2::STORAGE_WRITE_WITHOUT_FORMAT)
+    }
+}
--- a/wgpu-types/src/lib.rs
+++ b/wgpu-types/src/lib.rs
@ -309,7 +309,17 @@ bitflags::bitflags! {
        //
        // ? const FORMATS_TIER_1 = 1 << 14; (https://github.com/gpuweb/gpuweb/issues/3837)
        // ? const RW_STORAGE_TEXTURE_TIER_1 = 1 << 15; (https://github.com/gpuweb/gpuweb/issues/3838)
-        // TODO const BGRA8UNORM_STORAGE = 1 << 16;
+
+        /// Allows the [`wgpu::TextureUsages::STORAGE_BINDING`] usage on textures with format [`TextureFormat::Bgra8unorm`]
+        ///
+        /// Supported Platforms:
+        /// - Vulkan
+        /// - DX12
+        /// - Metal
+        ///
+        /// This is a web and native feature.
+        const BGRA8UNORM_STORAGE = 1 << 16;
+
        // ? const NORM16_FILTERABLE = 1 << 17; (https://github.com/gpuweb/gpuweb/issues/3839)
        // ? const NORM16_RESOLVE = 1 << 18; (https://github.com/gpuweb/gpuweb/issues/3839)
        // TODO const FLOAT32_FILTERABLE = 1 << 19;
@ -2977,6 +2987,11 @@ impl TextureFormat {
        } else {
            basic
        };
+        let bgra8unorm = if device_features.contains(Features::BGRA8UNORM_STORAGE) {
+            attachment | TextureUsages::STORAGE_BINDING
+        } else {
+            attachment
+        };

        #[rustfmt::skip] // lets make a nice table
        let (
@ -3005,7 +3020,7 @@ impl TextureFormat {
            Self::Rgba8Snorm =>           (        noaa,    storage),
            Self::Rgba8Uint =>            (        msaa,  all_flags),
            Self::Rgba8Sint =>            (        msaa,  all_flags),
-            Self::Bgra8Unorm =>           (msaa_resolve, attachment),
+            Self::Bgra8Unorm =>           (msaa_resolve, bgra8unorm),
            Self::Bgra8UnormSrgb =>       (msaa_resolve, attachment),
            Self::Rgb10a2Uint =>          (        msaa, attachment),
            Self::Rgb10a2Unorm =>         (msaa_resolve, attachment),
--- a/wgpu/src/backend/web.rs
+++ b/wgpu/src/backend/web.rs
@ -659,7 +659,7 @@ fn map_map_mode(mode: crate::MapMode) -> u32 {
    }
 }

-const FEATURES_MAPPING: [(wgt::Features, web_sys::GpuFeatureName); 9] = [
+const FEATURES_MAPPING: [(wgt::Features, web_sys::GpuFeatureName); 10] = [
    //TODO: update the name
    (
        wgt::Features::DEPTH_CLIP_CONTROL,
@ -697,6 +697,10 @@ const FEATURES_MAPPING: [(wgt::Features, web_sys::GpuFeatureName); 9] = [
        wgt::Features::RG11B10UFLOAT_RENDERABLE,
        web_sys::GpuFeatureName::Rg11b10ufloatRenderable,
    ),
+    (
+        wgt::Features::BGRA8UNORM_STORAGE,
+        web_sys::GpuFeatureName::Bgra8unormStorage,
+    ),
 ];

 fn map_wgt_features(supported_features: web_sys::GpuSupportedFeatures) -> wgt::Features {