Skip to content

Commit

Permalink
Merge #3559
Browse files Browse the repository at this point in the history 
3559: Distinguish STORAGE from STORAGE_READ_WRITE image feature r=kvark a=Wumpf

Implemented for Vulkan/DX12/DX11.
In Vulkan, storage image implies simultaneous read/write access.
Other APIs in contrast may allow STORAGE writeonly while disallowing reading.

Related to gfx-rs/wgpu#1109 - native extension for for read+write storage textures in wgpu native; following comments on the PR I'm redoing it and would need to add this to the hal api.

PR checklist:
- [x] `make` succeeds (on *nix)
- [x] `make reftests` succeeds
- [x] tested examples with the following backends:
    - Vulkan/DX12/DX11
    - did a dump of storage/storage-atomic/storage-rw flags each and eyeballed it as a sanity check



Co-authored-by: Andreas Reich <[email protected]>
  • Loading branch information
@bors @Wumpf
bors[bot] and Wumpf authored Jan 3, 2021
2 parents 83cebfd + 8b902f5 commit 187197d
Show file tree
Hide file tree
Showing 6 changed files with 93 additions and 27 deletions.
7 changes: 7 additions & 0 deletions src/backend/dx11/src/lib.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -410,7 +410,14 @@ fn get_format_properties(
props.buffer_features |= format::BufferFeature::STORAGE_TEXEL;
}
if can_image {
// Since read-only storage is exposed as SRV, we can guarantee read-only storage without checking D3D11_FORMAT_SUPPORT2_UAV_TYPED_LOAD first.
props.optimal_tiling |= format::ImageFeature::STORAGE;

if support_2.OutFormatSupport2 & d3d11::D3D11_FORMAT_SUPPORT2_UAV_TYPED_LOAD
!= 0
{
props.optimal_tiling |= format::ImageFeature::STORAGE_READ_WRITE;
}
}
}
}
Expand Down
5 changes: 5 additions & 0 deletions src/backend/dx12/src/lib.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1456,7 +1456,12 @@ impl FormatProperties {
props.buffer_features |= f::BufferFeature::STORAGE_TEXEL;
}
if can_image {
// Since read-only storage is exposed as SRV, we can guarantee read-only storage without checking D3D11_FORMAT_SUPPORT2_UAV_TYPED_LOAD first.
props.optimal_tiling |= f::ImageFeature::STORAGE;

if data.Support2 & d3d12::D3D12_FORMAT_SUPPORT2_UAV_TYPED_LOAD != 0 {
props.optimal_tiling |= f::ImageFeature::STORAGE_READ_WRITE;
}
}
}
//TODO: blits, linear tiling
Expand Down
41 changes: 28 additions & 13 deletions src/backend/metal/src/conversions.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -179,6 +179,15 @@ impl PrivateCapabilities {
let compressed_if = color_if | If::SAMPLED_LINEAR;
let depth_if = color_if | If::DEPTH_STENCIL_ATTACHMENT;

// Affected formats documented at:
// https://developer.apple.com/documentation/metal/mtlreadwritetexturetier/mtlreadwritetexturetier1?language=objc
// https://developer.apple.com/documentation/metal/mtlreadwritetexturetier/mtlreadwritetexturetier2?language=objc
let (read_write_tier1_if, read_write_tier2_if) = match self.read_write_texture_tier {
MTLReadWriteTextureTier::TierNone => (If::empty(), If::empty()),
MTLReadWriteTextureTier::Tier1 => (If::STORAGE_READ_WRITE, If::empty()),
MTLReadWriteTextureTier::Tier2 => (If::STORAGE_READ_WRITE, If::STORAGE_READ_WRITE),
};

match self.map_format(format) {
Some(A8Unorm) => Properties {
optimal_tiling: compressed_if,
Expand All @@ -187,6 +196,7 @@ impl PrivateCapabilities {
},
Some(R8Unorm) => Properties {
optimal_tiling: color_if
| read_write_tier2_if
| If::SAMPLED_LINEAR
| If::STORAGE
| If::COLOR_ATTACHMENT
Expand Down Expand Up @@ -221,12 +231,12 @@ impl PrivateCapabilities {
..Properties::default()
},
Some(R8Uint) => Properties {
optimal_tiling: color_if | If::STORAGE | If::COLOR_ATTACHMENT,
optimal_tiling: color_if | read_write_tier2_if | If::STORAGE | If::COLOR_ATTACHMENT,
buffer_features,
..Properties::default()
},
Some(R8Sint) => Properties {
optimal_tiling: color_if | If::STORAGE | If::COLOR_ATTACHMENT,
optimal_tiling: color_if | read_write_tier2_if | If::STORAGE | If::COLOR_ATTACHMENT,
buffer_features,
..Properties::default()
},
Expand All @@ -249,17 +259,18 @@ impl PrivateCapabilities {
..Properties::default()
},
Some(R16Uint) => Properties {
optimal_tiling: color_if | If::STORAGE | If::COLOR_ATTACHMENT,
optimal_tiling: color_if | read_write_tier2_if | If::STORAGE | If::COLOR_ATTACHMENT,
buffer_features,
..Properties::default()
},
Some(R16Sint) => Properties {
optimal_tiling: color_if | If::STORAGE | If::COLOR_ATTACHMENT,
optimal_tiling: color_if | read_write_tier2_if | If::STORAGE | If::COLOR_ATTACHMENT,
buffer_features,
..Properties::default()
},
Some(R16Float) => Properties {
optimal_tiling: color_if
| read_write_tier2_if
| If::SAMPLED_LINEAR
| If::STORAGE
| If::COLOR_ATTACHMENT
Expand Down Expand Up @@ -345,7 +356,7 @@ impl PrivateCapabilities {
..Properties::default()
},
Some(R32Uint) if self.format_r32_all => Properties {
optimal_tiling: color_if | If::STORAGE | If::COLOR_ATTACHMENT,
optimal_tiling: color_if | read_write_tier1_if | If::STORAGE | If::COLOR_ATTACHMENT,
buffer_features,
..Properties::default()
},
Expand All @@ -355,7 +366,7 @@ impl PrivateCapabilities {
..Properties::default()
},
Some(R32Sint) if self.format_r32_all => Properties {
optimal_tiling: color_if | If::STORAGE | If::COLOR_ATTACHMENT,
optimal_tiling: color_if | read_write_tier1_if | If::STORAGE | If::COLOR_ATTACHMENT,
buffer_features,
..Properties::default()
},
Expand All @@ -379,6 +390,7 @@ impl PrivateCapabilities {
},
Some(R32Float) if self.format_r32float_all => Properties {
optimal_tiling: color_if
| read_write_tier1_if
| If::SAMPLED_LINEAR
| If::STORAGE
| If::COLOR_ATTACHMENT
Expand Down Expand Up @@ -415,6 +427,7 @@ impl PrivateCapabilities {
},
Some(RGBA8Unorm) => Properties {
optimal_tiling: color_if
| read_write_tier2_if
| If::SAMPLED_LINEAR
| If::STORAGE
| If::COLOR_ATTACHMENT
Expand Down Expand Up @@ -449,12 +462,12 @@ impl PrivateCapabilities {
..Properties::default()
},
Some(RGBA8Uint) => Properties {
optimal_tiling: color_if | If::STORAGE | If::COLOR_ATTACHMENT,
optimal_tiling: color_if | read_write_tier2_if | If::STORAGE | If::COLOR_ATTACHMENT,
buffer_features,
..Properties::default()
},
Some(RGBA8Sint) => Properties {
optimal_tiling: color_if | If::STORAGE | If::COLOR_ATTACHMENT,
optimal_tiling: color_if | read_write_tier2_if | If::STORAGE | If::COLOR_ATTACHMENT,
buffer_features,
..Properties::default()
},
Expand Down Expand Up @@ -611,17 +624,18 @@ impl PrivateCapabilities {
..Properties::default()
},
Some(RGBA16Uint) => Properties {
optimal_tiling: color_if | If::STORAGE | If::COLOR_ATTACHMENT,
optimal_tiling: color_if | read_write_tier2_if | If::STORAGE | If::COLOR_ATTACHMENT,
buffer_features,
..Properties::default()
},
Some(RGBA16Sint) => Properties {
optimal_tiling: color_if | If::STORAGE | If::COLOR_ATTACHMENT,
optimal_tiling: color_if | read_write_tier2_if | If::STORAGE | If::COLOR_ATTACHMENT,
buffer_features,
..Properties::default()
},
Some(RGBA16Float) => Properties {
optimal_tiling: color_if
| read_write_tier2_if
| If::SAMPLED_LINEAR
| If::STORAGE
| If::COLOR_ATTACHMENT
Expand All @@ -635,7 +649,7 @@ impl PrivateCapabilities {
..Properties::default()
},
Some(RGBA32Uint) if self.format_rgba32int_color_write => Properties {
optimal_tiling: color_if | If::COLOR_ATTACHMENT | If::STORAGE,
optimal_tiling: color_if | read_write_tier2_if | If::COLOR_ATTACHMENT | If::STORAGE,
buffer_features,
..Properties::default()
},
Expand All @@ -645,12 +659,13 @@ impl PrivateCapabilities {
..Properties::default()
},
Some(RGBA32Sint) if self.format_rgba32int_color_write => Properties {
optimal_tiling: color_if | If::COLOR_ATTACHMENT | If::STORAGE,
optimal_tiling: color_if | read_write_tier2_if | If::COLOR_ATTACHMENT | If::STORAGE,
buffer_features,
..Properties::default()
},
Some(RGBA32Float) if self.format_rgba32float_all => Properties {
optimal_tiling: color_if
| read_write_tier2_if
| If::SAMPLED_LINEAR
| If::STORAGE
| If::COLOR_ATTACHMENT
Expand All @@ -664,7 +679,7 @@ impl PrivateCapabilities {
..Properties::default()
},
Some(RGBA32Float) if self.format_rgba32float_color_write => Properties {
optimal_tiling: color_if | If::COLOR_ATTACHMENT | If::STORAGE,
optimal_tiling: color_if | read_write_tier2_if | If::COLOR_ATTACHMENT | If::STORAGE,
buffer_features,
..Properties::default()
},
Expand Down
2 changes: 2 additions & 0 deletions src/backend/metal/src/lib.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -680,6 +680,7 @@ struct PrivateCapabilities {
os_version: (u32, u32),
msl_version: metal::MTLLanguageVersion,
exposed_queues: usize,
read_write_texture_tier: metal::MTLReadWriteTextureTier,
// if TRUE, we'll report `NON_FILL_POLYGON_MODE` feature without the points support
expose_line_mode: bool,
resource_heaps: bool,
Expand Down Expand Up @@ -823,6 +824,7 @@ impl PrivateCapabilities {
MTLLanguageVersion::V1_0
},
exposed_queues: 1,
read_write_texture_tier: device.read_write_texture_support(),
expose_line_mode: true,
resource_heaps: Self::supports_any(&device, RESOURCE_HEAP_SUPPORT),
argument_buffers: experiments.argument_buffers
Expand Down
35 changes: 34 additions & 1 deletion src/backend/vulkan/src/conv.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -378,7 +378,40 @@ pub fn map_query_result_flags(flags: query::ResultFlags) -> vk::QueryResultFlags
}

pub fn map_image_features(features: vk::FormatFeatureFlags) -> format::ImageFeature {
format::ImageFeature::from_bits_truncate(features.as_raw())
let mut mapped_flags = format::ImageFeature::empty();
if features.contains(vk::FormatFeatureFlags::SAMPLED_IMAGE) {
mapped_flags |= format::ImageFeature::SAMPLED;
}
if features.contains(vk::FormatFeatureFlags::SAMPLED_IMAGE_FILTER_LINEAR) {
mapped_flags |= format::ImageFeature::SAMPLED_LINEAR;
}

if features.contains(vk::FormatFeatureFlags::STORAGE_IMAGE) {
mapped_flags |= format::ImageFeature::STORAGE;
mapped_flags |= format::ImageFeature::STORAGE_READ_WRITE;
}
if features.contains(vk::FormatFeatureFlags::STORAGE_IMAGE_ATOMIC) {
mapped_flags |= format::ImageFeature::STORAGE_ATOMIC;
}

if features.contains(vk::FormatFeatureFlags::COLOR_ATTACHMENT) {
mapped_flags |= format::ImageFeature::COLOR_ATTACHMENT;
}
if features.contains(vk::FormatFeatureFlags::COLOR_ATTACHMENT_BLEND) {
mapped_flags |= format::ImageFeature::COLOR_ATTACHMENT_BLEND;
}
if features.contains(vk::FormatFeatureFlags::DEPTH_STENCIL_ATTACHMENT) {
mapped_flags |= format::ImageFeature::DEPTH_STENCIL_ATTACHMENT;
}

if features.contains(vk::FormatFeatureFlags::BLIT_SRC) {
mapped_flags |= format::ImageFeature::BLIT_SRC;
}
if features.contains(vk::FormatFeatureFlags::BLIT_DST) {
mapped_flags |= format::ImageFeature::BLIT_DST;
}

mapped_flags
}

pub fn map_buffer_features(features: vk::FormatFeatureFlags) -> format::BufferFeature {
Expand Down
30 changes: 17 additions & 13 deletions src/hal/src/format.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -148,24 +148,28 @@ bitflags!(
pub struct ImageFeature: u32 {
/// Image view can be sampled.
const SAMPLED = 0x1;
/// Image view can be used as storage image.
const STORAGE = 0x2;
/// Image view can be used as storage image (with atomics).
const STORAGE_ATOMIC = 0x4;
/// Image can be sampled with a (mipmap) linear sampler or as blit source with linear sampling.
/// (implies SAMPLED and BLIT_SRC support)
const SAMPLED_LINEAR = 0x2;

/// Image view can be used as storage image with exclusive read & write access.
const STORAGE = 0x10;
/// Image view can be used as storage image with simultaneous read/write access.
const STORAGE_READ_WRITE = 0x20;
/// Image view can be used as storage image with atomics.
const STORAGE_ATOMIC = 0x40;

/// Image view can be used as color and input attachment.
const COLOR_ATTACHMENT = 0x80;
const COLOR_ATTACHMENT = 0x100;
/// Image view can be used as color (with blending) and input attachment.
const COLOR_ATTACHMENT_BLEND = 0x100;
const COLOR_ATTACHMENT_BLEND = 0x200;
/// Image view can be used as depth-stencil and input attachment.
const DEPTH_STENCIL_ATTACHMENT = 0x200;
const DEPTH_STENCIL_ATTACHMENT = 0x400;

/// Image can be used as source for blit commands.
const BLIT_SRC = 0x400;
const BLIT_SRC = 0x1000;
/// Image can be used as destination for blit commands.
const BLIT_DST = 0x800;
/// Image can be sampled with a (mipmap) linear sampler or as blit source
/// with linear sampling.
/// Requires `SAMPLED` or `BLIT_SRC` flag.
const SAMPLED_LINEAR = 0x1000;
const BLIT_DST = 0x2000;
}
);

Expand Down

0 comments on commit 187197d

Please sign in to comment.