diff --git a/CHANGELOG.md b/CHANGELOG.md index d5096cddf5..a28d6f3c4a 100644 --- a/CHANGELOG.md +++ b/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 diff --git a/deno_webgpu/01_webgpu.js b/deno_webgpu/01_webgpu.js index 5d38ede30f..0e38aee70d 100644 --- 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 diff --git a/deno_webgpu/lib.rs b/deno_webgpu/lib.rs index 3109164ba3..d3b924bf7f 100644 --- 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 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 diff --git a/deno_webgpu/webgpu.idl b/deno_webgpu/webgpu.idl index 59b0547db4..0b6b04eb4e 100644 --- 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 diff --git a/tests/tests/bgra8unorm_storage.rs b/tests/tests/bgra8unorm_storage.rs new file mode 100644 index 0000000000..6acd543f59 --- /dev/null +++ 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; +@compute @workgroup_size(256) +fn main(@builtin(workgroup_id) wgid: vec3) { + 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(); + }); +} diff --git a/tests/tests/root.rs b/tests/tests/root.rs index ff910ef877..a97a160a2a 100644 --- 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; diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs index d512dd3245..1bf318bbaa 100644 --- 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 Sf::Rgba8Snorm, Tf::Rgba8Uint => Sf::Rgba8Uint, Tf::Rgba8Sint => Sf::Rgba8Sint, + Tf::Bgra8Unorm => Sf::Bgra8Unorm, Tf::Rgb10a2Uint => Sf::Rgb10a2Uint, Tf::Rgb10a2Unorm => Sf::Rgb10a2Unorm, diff --git a/wgpu-hal/src/dx11/adapter.rs b/wgpu-hal/src/dx11/adapter.rs index 290a9ade22..e2beb06571 100644 --- 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 // diff --git a/wgpu-hal/src/dx12/adapter.rs b/wgpu-hal/src/dx12/adapter.rs index e59de68781..f291e0808d 100644 --- 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::() 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(); diff --git a/wgpu-hal/src/gles/adapter.rs b/wgpu-hal/src/gles/adapter.rs index 4c81fca65b..e2f288d51f 100644 --- 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, diff --git a/wgpu-hal/src/metal/adapter.rs b/wgpu-hal/src/metal/adapter.rs index 538ccd64bf..2f48712b9b 100644 --- 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, diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs index 2224dbcc23..ef2f342084 100644 --- 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) + } +} diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs index c525d73de3..57f67f9ef1 100644 --- 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), diff --git a/wgpu/src/backend/web.rs b/wgpu/src/backend/web.rs index 5e3fc01767..a76fe5b142 100644 --- 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 {