Add support for bgra8unorm-storage

CHANGELOG.md

@@ -114,6 +114,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
 

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
 

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
 

deno_webgpu/webgpu.idl

@@ -102,6 +102,7 @@ enum GPUFeatureName {
     // shader
     "shader-f16",
     "rg11b10ufloat-renderable",
+    "bgra8unorm-storage",
 
     // extended from spec
 

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();
+    });
+}

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;

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,

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
         //

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();
 

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,

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,

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)
+    }
+}

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),