From 79b2708a251397f6d4d1e08aa8761a337c8aa115 Mon Sep 17 00:00:00 2001 From: atlas Date: Mon, 2 Dec 2024 05:23:12 -0500 Subject: [PATCH] Align Storage Access enums to spec --- .../tests/data/zero-init-texture-binding.ron | 2 +- tests/tests/bgra8unorm_storage.rs | 5 +- wgpu-core/src/binding_model.rs | 4 ++ wgpu-core/src/command/compute.rs | 2 +- wgpu-core/src/conv.rs | 10 ++-- wgpu-core/src/device/resource.rs | 46 ++++++++++++++----- wgpu-core/src/instance.rs | 14 +++++- wgpu-hal/src/auxil/dxgi/conv.rs | 3 +- wgpu-hal/src/dx12/adapter.rs | 4 +- wgpu-hal/src/dx12/conv.rs | 6 +-- wgpu-hal/src/dx12/device.rs | 2 +- wgpu-hal/src/gles/adapter.rs | 3 +- wgpu-hal/src/gles/queue.rs | 8 +++- wgpu-hal/src/lib.rs | 22 +++++---- wgpu-hal/src/metal/adapter.rs | 41 ++++++++++------- wgpu-hal/src/metal/conv.rs | 4 +- wgpu-hal/src/vulkan/adapter.rs | 2 +- wgpu-hal/src/vulkan/conv.rs | 34 ++++++++++---- wgpu-types/src/lib.rs | 10 +++- 19 files changed, 154 insertions(+), 68 deletions(-) diff --git a/player/tests/data/zero-init-texture-binding.ron b/player/tests/data/zero-init-texture-binding.ron index 48415f43c57..2368a080237 100644 --- a/player/tests/data/zero-init-texture-binding.ron +++ b/player/tests/data/zero-init-texture-binding.ron @@ -1,5 +1,5 @@ ( - features: [], + features: ["TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES"], expectations: [ ( name: "Sampled Texture", diff --git a/tests/tests/bgra8unorm_storage.rs b/tests/tests/bgra8unorm_storage.rs index 30babe0d82b..b1c9da6abdf 100644 --- a/tests/tests/bgra8unorm_storage.rs +++ b/tests/tests/bgra8unorm_storage.rs @@ -21,7 +21,10 @@ static BGRA8_UNORM_STORAGE: GpuTestConfiguration = GpuTestConfiguration::new() max_storage_textures_per_shader_stage: 1, ..Default::default() }) - .features(wgpu::Features::BGRA8UNORM_STORAGE), + .features( + wgpu::Features::BGRA8UNORM_STORAGE + | wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES, + ), ) .run_async(|ctx| async move { let device = &ctx.device; diff --git a/wgpu-core/src/binding_model.rs b/wgpu-core/src/binding_model.rs index b0d9b5f68ca..c7867ab2103 100644 --- a/wgpu-core/src/binding_model.rs +++ b/wgpu-core/src/binding_model.rs @@ -185,6 +185,10 @@ pub enum CreateBindGroupError { DepthStencilAspect, #[error("The adapter does not support read access for storage textures of format {0:?}")] StorageReadNotSupported(wgt::TextureFormat), + #[error("The adapter does not support write access for storage textures of format {0:?}")] + StorageWriteNotSupported(wgt::TextureFormat), + #[error("The adapter does not support read-write access for storage textures of format {0:?}")] + StorageReadWriteNotSupported(wgt::TextureFormat), #[error(transparent)] ResourceUsageCompatibility(#[from] ResourceUsageCompatibilityError), #[error(transparent)] diff --git a/wgpu-core/src/command/compute.rs b/wgpu-core/src/command/compute.rs index 89681b8e271..7eaafea4691 100644 --- a/wgpu-core/src/command/compute.rs +++ b/wgpu-core/src/command/compute.rs @@ -944,7 +944,7 @@ fn dispatch_indirect( let src_transition = state .intermediate_trackers .buffers - .set_single(&buffer, hal::BufferUses::STORAGE_READ); + .set_single(&buffer, hal::BufferUses::STORAGE_READ_ONLY); let src_barrier = src_transition.map(|transition| transition.into_hal(&buffer, &state.snatch_guard)); unsafe { diff --git a/wgpu-core/src/conv.rs b/wgpu-core/src/conv.rs index a9f48ff6915..b114e9826e9 100644 --- a/wgpu-core/src/conv.rs +++ b/wgpu-core/src/conv.rs @@ -82,7 +82,7 @@ pub fn map_buffer_usage(usage: wgt::BufferUsages) -> hal::BufferUses { usage.contains(wgt::BufferUsages::UNIFORM), ); u.set( - hal::BufferUses::STORAGE_READ | hal::BufferUses::STORAGE_READ_WRITE, + hal::BufferUses::STORAGE_READ_WRITE, usage.contains(wgt::BufferUsages::STORAGE), ); u.set( @@ -122,7 +122,7 @@ pub fn map_texture_usage( usage.contains(wgt::TextureUsages::TEXTURE_BINDING), ); u.set( - hal::TextureUses::STORAGE_READ | hal::TextureUses::STORAGE_READ_WRITE, + hal::TextureUses::STORAGE_READ_WRITE, usage.contains(wgt::TextureUsages::STORAGE_BINDING), ); let is_color = aspect.contains(hal::FormatAspects::COLOR); @@ -179,7 +179,11 @@ pub fn map_texture_usage_from_hal(uses: hal::TextureUses) -> wgt::TextureUsages ); u.set( wgt::TextureUsages::STORAGE_BINDING, - uses.contains(hal::TextureUses::STORAGE_READ | hal::TextureUses::STORAGE_READ_WRITE), + uses.intersects( + hal::TextureUses::STORAGE_READ_ONLY + | hal::TextureUses::STORAGE_WRITE_ONLY + | hal::TextureUses::STORAGE_READ_WRITE, + ), ); u.set( wgt::TextureUsages::RENDER_ATTACHMENT, diff --git a/wgpu-core/src/device/resource.rs b/wgpu-core/src/device/resource.rs index a89bf0b6c24..97c616d8070 100644 --- a/wgpu-core/src/device/resource.rs +++ b/wgpu-core/src/device/resource.rs @@ -521,7 +521,16 @@ impl Device { self.require_downlevel_flags(wgt::DownlevelFlags::INDIRECT_EXECUTION)?; // We are going to be reading from it, internally; // when validating the content of the buffer - usage |= hal::BufferUses::STORAGE_READ | hal::BufferUses::STORAGE_READ_WRITE; + if !usage.intersects( + hal::BufferUses::STORAGE_READ_ONLY | hal::BufferUses::STORAGE_READ_WRITE, + ) { + if usage.contains(hal::BufferUses::STORAGE_WRITE_ONLY) { + usage |= hal::BufferUses::STORAGE_READ_WRITE; + usage &= !hal::BufferUses::STORAGE_WRITE_ONLY; + } else { + usage |= hal::BufferUses::STORAGE_READ_ONLY; + } + } } if desc.mapped_at_creation { @@ -1254,7 +1263,8 @@ impl Device { } TextureViewDimension::D3 => { hal::TextureUses::RESOURCE - | hal::TextureUses::STORAGE_READ + | hal::TextureUses::STORAGE_READ_ONLY + | hal::TextureUses::STORAGE_WRITE_ONLY | hal::TextureUses::STORAGE_READ_WRITE } _ => hal::TextureUses::all(), @@ -1916,7 +1926,7 @@ impl Device { wgt::BufferBindingType::Storage { read_only } => ( wgt::BufferUsages::STORAGE, if read_only { - hal::BufferUses::STORAGE_READ + hal::BufferUses::STORAGE_READ_ONLY } else { hal::BufferUses::STORAGE_READ_WRITE }, @@ -2489,24 +2499,36 @@ impl Device { } let internal_use = match access { - wgt::StorageTextureAccess::WriteOnly => hal::TextureUses::STORAGE_READ_WRITE, + wgt::StorageTextureAccess::WriteOnly => { + if !view.format_features.flags.intersects( + wgt::TextureFormatFeatureFlags::STORAGE_WRITE_ONLY + | wgt::TextureFormatFeatureFlags::STORAGE_READ_WRITE, + ) { + println!( + "{:?} {:?}", + view.desc.format, + view.format_features.flags.iter_names().collect::>() + ); + return Err(Error::StorageWriteNotSupported(view.desc.format)); + } + hal::TextureUses::STORAGE_WRITE_ONLY + } wgt::StorageTextureAccess::ReadOnly => { - if !view - .format_features - .flags - .contains(wgt::TextureFormatFeatureFlags::STORAGE_WRITE) - { + if !view.format_features.flags.intersects( + wgt::TextureFormatFeatureFlags::STORAGE_READ_ONLY + | wgt::TextureFormatFeatureFlags::STORAGE_READ_WRITE, + ) { return Err(Error::StorageReadNotSupported(view.desc.format)); } - hal::TextureUses::STORAGE_READ + hal::TextureUses::STORAGE_READ_ONLY } wgt::StorageTextureAccess::ReadWrite => { if !view .format_features .flags - .contains(wgt::TextureFormatFeatureFlags::STORAGE_WRITE) + .contains(wgt::TextureFormatFeatureFlags::STORAGE_READ_WRITE) { - return Err(Error::StorageReadNotSupported(view.desc.format)); + return Err(Error::StorageReadWriteNotSupported(view.desc.format)); } hal::TextureUses::STORAGE_READ_WRITE diff --git a/wgpu-core/src/instance.rs b/wgpu-core/src/instance.rs index 16278e16f80..2f4365e9270 100644 --- a/wgpu-core/src/instance.rs +++ b/wgpu-core/src/instance.rs @@ -512,7 +512,9 @@ impl Adapter { ); allowed_usages.set( wgt::TextureUsages::STORAGE_BINDING, - caps.contains(Tfc::STORAGE_WRITE), + caps.intersects( + Tfc::STORAGE_WRITE_ONLY | Tfc::STORAGE_READ_ONLY | Tfc::STORAGE_READ_WRITE, + ), ); allowed_usages.set( wgt::TextureUsages::RENDER_ATTACHMENT, @@ -521,7 +523,15 @@ impl Adapter { let mut flags = wgt::TextureFormatFeatureFlags::empty(); flags.set( - wgt::TextureFormatFeatureFlags::STORAGE_WRITE, + wgt::TextureFormatFeatureFlags::STORAGE_READ_ONLY, + caps.contains(Tfc::STORAGE_READ_ONLY), + ); + flags.set( + wgt::TextureFormatFeatureFlags::STORAGE_WRITE_ONLY, + caps.contains(Tfc::STORAGE_WRITE_ONLY), + ); + flags.set( + wgt::TextureFormatFeatureFlags::STORAGE_READ_WRITE, caps.contains(Tfc::STORAGE_READ_WRITE), ); diff --git a/wgpu-hal/src/auxil/dxgi/conv.rs b/wgpu-hal/src/auxil/dxgi/conv.rs index 878dab39e93..ad64f044cc7 100644 --- a/wgpu-hal/src/auxil/dxgi/conv.rs +++ b/wgpu-hal/src/auxil/dxgi/conv.rs @@ -206,7 +206,8 @@ pub fn map_texture_format_for_resource( } else if format.is_depth_stencil_format() && usage.intersects( crate::TextureUses::RESOURCE - | crate::TextureUses::STORAGE_READ + | crate::TextureUses::STORAGE_READ_ONLY + | crate::TextureUses::STORAGE_WRITE_ONLY | crate::TextureUses::STORAGE_READ_WRITE, ) { diff --git a/wgpu-hal/src/dx12/adapter.rs b/wgpu-hal/src/dx12/adapter.rs index b88941c81e0..a8940846d6b 100644 --- a/wgpu-hal/src/dx12/adapter.rs +++ b/wgpu-hal/src/dx12/adapter.rs @@ -670,13 +670,13 @@ impl crate::Adapter for super::Adapter { ); // UAVs use srv_uav_format caps.set( - Tfc::STORAGE_WRITE, + Tfc::STORAGE_WRITE_ONLY, data_srv_uav .Support1 .contains(Direct3D12::D3D12_FORMAT_SUPPORT1_TYPED_UNORDERED_ACCESS_VIEW), ); caps.set( - Tfc::STORAGE_READ_WRITE, + Tfc::STORAGE_READ_WRITE | Tfc::STORAGE_READ_ONLY | Tfc::STORAGE_WRITE_ONLY, data_srv_uav .Support2 .contains(Direct3D12::D3D12_FORMAT_SUPPORT2_UAV_TYPED_LOAD), diff --git a/wgpu-hal/src/dx12/conv.rs b/wgpu-hal/src/dx12/conv.rs index 8e60f6e0645..70ba83fa429 100644 --- a/wgpu-hal/src/dx12/conv.rs +++ b/wgpu-hal/src/dx12/conv.rs @@ -128,9 +128,9 @@ pub fn map_buffer_usage_to_state(usage: crate::BufferUses) -> Direct3D12::D3D12_ if usage.intersects(Bu::VERTEX | Bu::UNIFORM) { state |= Direct3D12::D3D12_RESOURCE_STATE_VERTEX_AND_CONSTANT_BUFFER; } - if usage.intersects(Bu::STORAGE_READ_WRITE) { + if usage.intersects(Bu::STORAGE_READ_WRITE | Bu::STORAGE_WRITE_ONLY) { state |= Direct3D12::D3D12_RESOURCE_STATE_UNORDERED_ACCESS; - } else if usage.intersects(Bu::STORAGE_READ) { + } else if usage.intersects(Bu::STORAGE_READ_ONLY) { state |= Direct3D12::D3D12_RESOURCE_STATE_PIXEL_SHADER_RESOURCE | Direct3D12::D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE; } @@ -168,7 +168,7 @@ pub fn map_texture_usage_to_state(usage: crate::TextureUses) -> Direct3D12::D3D1 if usage.intersects(Tu::DEPTH_STENCIL_WRITE) { state |= Direct3D12::D3D12_RESOURCE_STATE_DEPTH_WRITE; } - if usage.intersects(Tu::STORAGE_READ | Tu::STORAGE_READ_WRITE) { + if usage.intersects(Tu::STORAGE_READ_ONLY | Tu::STORAGE_READ_WRITE) { state |= Direct3D12::D3D12_RESOURCE_STATE_UNORDERED_ACCESS; } state diff --git a/wgpu-hal/src/dx12/device.rs b/wgpu-hal/src/dx12/device.rs index 9e089bbd123..2a7d88592e0 100644 --- a/wgpu-hal/src/dx12/device.rs +++ b/wgpu-hal/src/dx12/device.rs @@ -577,7 +577,7 @@ impl crate::Device for super::Device { None }, handle_uav: if desc.usage.intersects( - crate::TextureUses::STORAGE_READ | crate::TextureUses::STORAGE_READ_WRITE, + crate::TextureUses::STORAGE_READ_ONLY | crate::TextureUses::STORAGE_READ_WRITE, ) { match unsafe { view_desc.to_uav() } { Some(raw_desc) => { diff --git a/wgpu-hal/src/gles/adapter.rs b/wgpu-hal/src/gles/adapter.rs index 7d2ba2d23d8..a95511908e2 100644 --- a/wgpu-hal/src/gles/adapter.rs +++ b/wgpu-hal/src/gles/adapter.rs @@ -1038,7 +1038,8 @@ impl crate::Adapter for super::Adapter { let renderable = unfilterable | Tfc::COLOR_ATTACHMENT | sample_count | Tfc::MULTISAMPLE_RESOLVE; let filterable_renderable = filterable | renderable | Tfc::COLOR_ATTACHMENT_BLEND; - let storage = base | Tfc::STORAGE_WRITE | Tfc::STORAGE_READ_WRITE; + let storage = + base | Tfc::STORAGE_READ_WRITE | Tfc::STORAGE_READ_ONLY | Tfc::STORAGE_WRITE_ONLY; let feature_fn = |f, caps| { if self.shared.features.contains(f) { diff --git a/wgpu-hal/src/gles/queue.rs b/wgpu-hal/src/gles/queue.rs index 03edf30831f..892e854b526 100644 --- a/wgpu-hal/src/gles/queue.rs +++ b/wgpu-hal/src/gles/queue.rs @@ -1225,7 +1225,9 @@ impl super::Queue { flags |= glow::BUFFER_UPDATE_BARRIER_BIT; } if usage.intersects( - crate::BufferUses::STORAGE_READ | crate::BufferUses::STORAGE_READ_WRITE, + crate::BufferUses::STORAGE_READ_ONLY + | crate::BufferUses::STORAGE_WRITE_ONLY + | crate::BufferUses::STORAGE_READ_WRITE, ) { flags |= glow::SHADER_STORAGE_BARRIER_BIT; } @@ -1237,7 +1239,9 @@ impl super::Queue { flags |= glow::TEXTURE_FETCH_BARRIER_BIT; } if usage.intersects( - crate::TextureUses::STORAGE_READ | crate::TextureUses::STORAGE_READ_WRITE, + crate::TextureUses::STORAGE_READ_ONLY + | crate::TextureUses::STORAGE_WRITE_ONLY + | crate::TextureUses::STORAGE_READ_WRITE, ) { flags |= glow::SHADER_IMAGE_ACCESS_BARRIER_BIT; } diff --git a/wgpu-hal/src/lib.rs b/wgpu-hal/src/lib.rs index 44ee9fb9aa1..0a263104e52 100644 --- a/wgpu-hal/src/lib.rs +++ b/wgpu-hal/src/lib.rs @@ -1543,8 +1543,10 @@ bitflags!( /// Format can be sampled with a min/max reduction sampler. const SAMPLED_MINMAX = 1 << 2; + /// Format can be used as storage with read-only access. + const STORAGE_READ_ONLY = 1 << 16; /// Format can be used as storage with write-only access. - const STORAGE_WRITE = 1 << 3; + const STORAGE_WRITE_ONLY = 1 << 3; /// Format can be used as storage with read and read/write access. const STORAGE_READ_WRITE = 1 << 4; /// Format can be used as storage with atomics. @@ -1675,8 +1677,10 @@ bitflags::bitflags! { /// A uniform buffer bound in a bind group. const UNIFORM = 1 << 6; /// A read-only storage buffer used in a bind group. - const STORAGE_READ = 1 << 7; - /// A read-write or write-only buffer used in a bind group. + const STORAGE_READ_ONLY = 1 << 7; + /// A write-only storage buffer used in a bind group. + const STORAGE_WRITE_ONLY = 1 << 8; + /// A read-write buffer used in a bind group. const STORAGE_READ_WRITE = 1 << 8; /// The indirect or count buffer in a indirect draw or dispatch. const INDIRECT = 1 << 9; @@ -1688,7 +1692,7 @@ bitflags::bitflags! { /// The combination of states that a buffer may be in _at the same time_. const INCLUSIVE = Self::MAP_READ.bits() | Self::COPY_SRC.bits() | Self::INDEX.bits() | Self::VERTEX.bits() | Self::UNIFORM.bits() | - Self::STORAGE_READ.bits() | Self::INDIRECT.bits() | Self::BOTTOM_LEVEL_ACCELERATION_STRUCTURE_INPUT.bits() | Self::TOP_LEVEL_ACCELERATION_STRUCTURE_INPUT.bits(); + Self::STORAGE_READ_ONLY.bits() | Self::INDIRECT.bits() | Self::BOTTOM_LEVEL_ACCELERATION_STRUCTURE_INPUT.bits() | Self::TOP_LEVEL_ACCELERATION_STRUCTURE_INPUT.bits(); /// The combination of states that a buffer must exclusively be in. const EXCLUSIVE = Self::MAP_WRITE.bits() | Self::COPY_DST.bits() | Self::STORAGE_READ_WRITE.bits() | Self::ACCELERATION_STRUCTURE_SCRATCH.bits(); /// The combination of all usages that the are guaranteed to be be ordered by the hardware. @@ -1719,17 +1723,19 @@ bitflags::bitflags! { /// Read-write depth stencil usage const DEPTH_STENCIL_WRITE = 1 << 7; /// Read-only storage buffer usage. Corresponds to a UAV in d3d, so is exclusive, despite being read only. - const STORAGE_READ = 1 << 8; - /// Read-write or write-only storage buffer usage. + const STORAGE_READ_ONLY = 1 << 8; + /// Write-only storage buffer usage. + const STORAGE_WRITE_ONLY = 1 << 9; + /// Read-write storage buffer usage. const STORAGE_READ_WRITE = 1 << 9; /// The combination of states that a texture may be in _at the same time_. const INCLUSIVE = Self::COPY_SRC.bits() | Self::RESOURCE.bits() | Self::DEPTH_STENCIL_READ.bits(); /// The combination of states that a texture must exclusively be in. - const EXCLUSIVE = Self::COPY_DST.bits() | Self::COLOR_TARGET.bits() | Self::DEPTH_STENCIL_WRITE.bits() | Self::STORAGE_READ.bits() | Self::STORAGE_READ_WRITE.bits() | Self::PRESENT.bits(); + const EXCLUSIVE = Self::COPY_DST.bits() | Self::COLOR_TARGET.bits() | Self::DEPTH_STENCIL_WRITE.bits() | Self::STORAGE_READ_ONLY.bits() | Self::STORAGE_READ_WRITE.bits() | Self::PRESENT.bits(); /// The combination of all usages that the are guaranteed to be be ordered by the hardware. /// If a usage is ordered, then if the texture state doesn't change between draw calls, there /// are no barriers needed for synchronization. - const ORDERED = Self::INCLUSIVE.bits() | Self::COLOR_TARGET.bits() | Self::DEPTH_STENCIL_WRITE.bits() | Self::STORAGE_READ.bits(); + const ORDERED = Self::INCLUSIVE.bits() | Self::COLOR_TARGET.bits() | Self::DEPTH_STENCIL_WRITE.bits() | Self::STORAGE_READ_ONLY.bits(); /// Flag used by the wgpu-core texture tracker to say a texture is in different states for every sub-resource const COMPLEX = 1 << 10; diff --git a/wgpu-hal/src/metal/adapter.rs b/wgpu-hal/src/metal/adapter.rs index bb7b9926ac4..0c00b964a12 100644 --- a/wgpu-hal/src/metal/adapter.rs +++ b/wgpu-hal/src/metal/adapter.rs @@ -111,7 +111,9 @@ impl crate::Adapter for super::Adapter { // Metal defined pixel format capabilities let all_caps = Tfc::SAMPLED_LINEAR - | Tfc::STORAGE_WRITE + | Tfc::STORAGE_READ_ONLY + | Tfc::STORAGE_WRITE_ONLY + | Tfc::STORAGE_READ_WRITE | Tfc::COLOR_ATTACHMENT | Tfc::COLOR_ATTACHMENT_BLEND | msaa_count @@ -134,7 +136,7 @@ impl crate::Adapter for super::Adapter { | Tf::Rgba8Sint | Tf::Rgba16Uint | Tf::Rgba16Sint => { - read_write_tier2_if | Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT | msaa_count + read_write_tier2_if | Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | msaa_count } Tf::R16Unorm | Tf::R16Snorm @@ -143,65 +145,72 @@ impl crate::Adapter for super::Adapter { | Tf::Rgba16Unorm | Tf::Rgba16Snorm => { Tfc::SAMPLED_LINEAR - | Tfc::STORAGE_WRITE + | Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | Tfc::COLOR_ATTACHMENT_BLEND | msaa_count | msaa_resolve_desktop_if } Tf::Rg8Unorm | Tf::Rg16Float | Tf::Bgra8Unorm => all_caps, - Tf::Rg8Uint | Tf::Rg8Sint => Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT | msaa_count, + Tf::Rg8Uint | Tf::Rg8Sint => { + Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | msaa_count + } Tf::R32Uint | Tf::R32Sint => { - read_write_tier1_if | Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT | msaa_count + read_write_tier1_if | Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | msaa_count } Tf::R32Float => { let flags = if pc.format_r32float_all { all_caps } else { - Tfc::STORAGE_WRITE + Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | Tfc::COLOR_ATTACHMENT_BLEND | msaa_count }; read_write_tier1_if | flags } - Tf::Rg16Uint | Tf::Rg16Sint => Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT | msaa_count, + Tf::Rg16Uint | Tf::Rg16Sint => { + Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | msaa_count + } Tf::Rgba8UnormSrgb | Tf::Bgra8UnormSrgb => { let mut flags = all_caps; - flags.set(Tfc::STORAGE_WRITE, pc.format_rgba8_srgb_all); + flags.set(Tfc::STORAGE_WRITE_ONLY, pc.format_rgba8_srgb_all); flags } Tf::Rgb10a2Uint => { let mut flags = Tfc::COLOR_ATTACHMENT | msaa_count; - flags.set(Tfc::STORAGE_WRITE, pc.format_rgb10a2_uint_write); + flags.set(Tfc::STORAGE_WRITE_ONLY, pc.format_rgb10a2_uint_write); flags } Tf::Rgb10a2Unorm => { let mut flags = all_caps; - flags.set(Tfc::STORAGE_WRITE, pc.format_rgb10a2_unorm_all); + flags.set(Tfc::STORAGE_WRITE_ONLY, pc.format_rgb10a2_unorm_all); flags } Tf::Rg11b10Ufloat => { let mut flags = all_caps; - flags.set(Tfc::STORAGE_WRITE, pc.format_rg11b10_all); + flags.set(Tfc::STORAGE_WRITE_ONLY, pc.format_rg11b10_all); flags } - Tf::Rg32Uint | Tf::Rg32Sint => Tfc::COLOR_ATTACHMENT | Tfc::STORAGE_WRITE | msaa_count, + Tf::Rg32Uint | Tf::Rg32Sint => { + Tfc::COLOR_ATTACHMENT | Tfc::STORAGE_WRITE_ONLY | msaa_count + } Tf::Rg32Float => { if pc.format_rg32float_all { all_caps } else { - Tfc::STORAGE_WRITE + Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | Tfc::COLOR_ATTACHMENT_BLEND | msaa_count } } Tf::Rgba32Uint | Tf::Rgba32Sint => { - read_write_tier2_if | Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT | msaa_count + read_write_tier2_if | Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT | msaa_count } Tf::Rgba32Float => { - let mut flags = read_write_tier2_if | Tfc::STORAGE_WRITE | Tfc::COLOR_ATTACHMENT; + let mut flags = + read_write_tier2_if | Tfc::STORAGE_WRITE_ONLY | Tfc::COLOR_ATTACHMENT; if pc.format_rgba32float_all { flags |= all_caps } else if pc.msaa_apple7 { @@ -351,7 +360,7 @@ impl crate::Adapter for super::Adapter { usage: crate::TextureUses::COLOR_TARGET | crate::TextureUses::COPY_SRC | crate::TextureUses::COPY_DST - | crate::TextureUses::STORAGE_READ + //| crate::TextureUses::STORAGE_READ_ONLY | crate::TextureUses::STORAGE_READ_WRITE, }) } diff --git a/wgpu-hal/src/metal/conv.rs b/wgpu-hal/src/metal/conv.rs index 6ebabee1a6d..f56141d5a78 100644 --- a/wgpu-hal/src/metal/conv.rs +++ b/wgpu-hal/src/metal/conv.rs @@ -13,12 +13,12 @@ pub fn map_texture_usage( mtl_usage.set( metal::MTLTextureUsage::ShaderRead, usage.intersects( - Tu::RESOURCE | Tu::DEPTH_STENCIL_READ | Tu::STORAGE_READ | Tu::STORAGE_READ_WRITE, + Tu::RESOURCE | Tu::DEPTH_STENCIL_READ | Tu::STORAGE_READ_ONLY | Tu::STORAGE_READ_WRITE, ), ); mtl_usage.set( metal::MTLTextureUsage::ShaderWrite, - usage.intersects(Tu::STORAGE_READ_WRITE), + usage.intersects(Tu::STORAGE_WRITE_ONLY | Tu::STORAGE_READ_WRITE), ); // needed for combined depth/stencil formats since we might // create a stencil-only view from them diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs index 40340ca03b1..0b6bd1e4a46 100644 --- a/wgpu-hal/src/vulkan/adapter.rs +++ b/wgpu-hal/src/vulkan/adapter.rs @@ -2115,7 +2115,7 @@ impl crate::Adapter for super::Adapter { // features.contains(vk::FormatFeatureFlags::SAMPLED_IMAGE_FILTER_MINMAX), // ); flags.set( - Tfc::STORAGE_WRITE | Tfc::STORAGE_READ_WRITE, + Tfc::STORAGE_READ_WRITE | Tfc::STORAGE_WRITE_ONLY | Tfc::STORAGE_READ_ONLY, features.contains(vk::FormatFeatureFlags::STORAGE_IMAGE), ); flags.set( diff --git a/wgpu-hal/src/vulkan/conv.rs b/wgpu-hal/src/vulkan/conv.rs index 75b807a29ce..ce4aa410bf5 100644 --- a/wgpu-hal/src/vulkan/conv.rs +++ b/wgpu-hal/src/vulkan/conv.rs @@ -263,7 +263,11 @@ pub fn map_texture_usage(usage: crate::TextureUses) -> vk::ImageUsageFlags { ) { flags |= vk::ImageUsageFlags::DEPTH_STENCIL_ATTACHMENT; } - if usage.intersects(crate::TextureUses::STORAGE_READ | crate::TextureUses::STORAGE_READ_WRITE) { + if usage.intersects( + crate::TextureUses::STORAGE_READ_ONLY + | crate::TextureUses::STORAGE_WRITE_ONLY + | crate::TextureUses::STORAGE_READ_WRITE, + ) { flags |= vk::ImageUsageFlags::STORAGE; } flags @@ -305,13 +309,17 @@ pub fn map_texture_usage_to_barrier( access |= vk::AccessFlags::DEPTH_STENCIL_ATTACHMENT_READ | vk::AccessFlags::DEPTH_STENCIL_ATTACHMENT_WRITE; } - if usage.contains(crate::TextureUses::STORAGE_READ) { + if usage + .intersects(crate::TextureUses::STORAGE_READ_ONLY | crate::TextureUses::STORAGE_READ_WRITE) + { stages |= shader_stages; access |= vk::AccessFlags::SHADER_READ; } - if usage.contains(crate::TextureUses::STORAGE_READ_WRITE) { + if usage + .intersects(crate::TextureUses::STORAGE_WRITE_ONLY | crate::TextureUses::STORAGE_READ_WRITE) + { stages |= shader_stages; - access |= vk::AccessFlags::SHADER_READ | vk::AccessFlags::SHADER_WRITE; + access |= vk::AccessFlags::SHADER_WRITE; } if usage == crate::TextureUses::UNINITIALIZED || usage == crate::TextureUses::PRESENT { @@ -342,7 +350,7 @@ pub fn map_vk_image_usage(usage: vk::ImageUsageFlags) -> crate::TextureUses { bits |= crate::TextureUses::DEPTH_STENCIL_READ | crate::TextureUses::DEPTH_STENCIL_WRITE; } if usage.contains(vk::ImageUsageFlags::STORAGE) { - bits |= crate::TextureUses::STORAGE_READ | crate::TextureUses::STORAGE_READ_WRITE; + bits |= crate::TextureUses::STORAGE_READ_WRITE; } bits } @@ -507,7 +515,11 @@ pub fn map_buffer_usage(usage: crate::BufferUses) -> vk::BufferUsageFlags { if usage.contains(crate::BufferUses::UNIFORM) { flags |= vk::BufferUsageFlags::UNIFORM_BUFFER; } - if usage.intersects(crate::BufferUses::STORAGE_READ | crate::BufferUses::STORAGE_READ_WRITE) { + if usage.intersects( + crate::BufferUses::STORAGE_READ_ONLY + | crate::BufferUses::STORAGE_WRITE_ONLY + | crate::BufferUses::STORAGE_READ_WRITE, + ) { flags |= vk::BufferUsageFlags::STORAGE_BUFFER; } if usage.contains(crate::BufferUses::INDEX) { @@ -561,13 +573,17 @@ pub fn map_buffer_usage_to_barrier( stages |= shader_stages; access |= vk::AccessFlags::UNIFORM_READ; } - if usage.intersects(crate::BufferUses::STORAGE_READ) { + if usage + .intersects(crate::BufferUses::STORAGE_READ_ONLY | crate::BufferUses::STORAGE_READ_WRITE) + { stages |= shader_stages; access |= vk::AccessFlags::SHADER_READ; } - if usage.intersects(crate::BufferUses::STORAGE_READ_WRITE) { + if usage + .intersects(crate::BufferUses::STORAGE_WRITE_ONLY | crate::BufferUses::STORAGE_READ_WRITE) + { stages |= shader_stages; - access |= vk::AccessFlags::SHADER_READ | vk::AccessFlags::SHADER_WRITE; + access |= vk::AccessFlags::SHADER_WRITE; } if usage.contains(crate::BufferUses::INDEX) { stages |= vk::PipelineStageFlags::VERTEX_INPUT; diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs index 546a3fde0da..43676a6822d 100644 --- a/wgpu-types/src/lib.rs +++ b/wgpu-types/src/lib.rs @@ -2381,8 +2381,14 @@ bitflags::bitflags! { /// to a render pass for an automatic driver-implemented resolve. const MULTISAMPLE_RESOLVE = 1 << 5; /// When used as a STORAGE texture, then a texture with this format can be bound with - /// [`StorageTextureAccess::ReadOnly`] or [`StorageTextureAccess::ReadWrite`]. - const STORAGE_WRITE = 1 << 6; + /// [`StorageTextureAccess::ReadOnly`]. + const STORAGE_READ_ONLY = 1 << 8; + /// When used as a STORAGE texture, then a texture with this format can be bound with + /// [`StorageTextureAccess::WriteOnly`]. + const STORAGE_WRITE_ONLY = 1 << 6; + /// When used as a STORAGE texture, then a texture with this format can be bound with + /// any [`StorageTextureAccess`]. + const STORAGE_READ_WRITE = 1 << 9; /// If not present, the texture can't be blended into the render target. const BLENDABLE = 1 << 7; }