From 5ca5175dc5f7e1c184b8a7d698ee5c4cea08baf8 Mon Sep 17 00:00:00 2001 From: atlas dostal Date: Wed, 2 Oct 2024 17:07:42 -0400 Subject: [PATCH 1/9] add r64uint texture format, wgpu feature, naga capability, and adapter support detection for metal, directx, and vulkan --- Cargo.lock | 3 +- Cargo.toml | 2 +- naga/src/back/glsl/features.rs | 1 + naga/src/back/glsl/mod.rs | 1 + naga/src/back/hlsl/conv.rs | 2 +- naga/src/back/spv/instructions.rs | 1 + naga/src/back/spv/writer.rs | 11 +++-- naga/src/back/wgsl/writer.rs | 1 + naga/src/front/glsl/parser/types.rs | 1 + naga/src/front/spv/convert.rs | 1 + naga/src/front/wgsl/parse/conv.rs | 1 + naga/src/front/wgsl/to_wgsl.rs | 1 + naga/src/lib.rs | 1 + naga/src/proc/mod.rs | 1 + wgpu-core/src/conv.rs | 8 ++++ wgpu-core/src/instance.rs | 4 ++ wgpu-core/src/validation.rs | 3 ++ wgpu-hal/src/auxil/dxgi/conv.rs | 1 + wgpu-hal/src/dx12/adapter.rs | 2 + wgpu-hal/src/gles/adapter.rs | 3 ++ wgpu-hal/src/gles/conv.rs | 1 + wgpu-hal/src/lib.rs | 5 ++ wgpu-hal/src/metal/adapter.rs | 13 +++++ wgpu-hal/src/metal/conv.rs | 5 ++ wgpu-hal/src/metal/mod.rs | 1 + wgpu-hal/src/vulkan/adapter.rs | 47 ++++++++++++++++++ wgpu-hal/src/vulkan/conv.rs | 1 + wgpu-info/src/texture.rs | 3 +- wgpu-types/src/lib.rs | 48 ++++++++++++++++++- wgpu/src/backend/webgpu.rs | 1 + .../webgpu/webgpu_sys/gen_GpuTextureFormat.rs | 1 + 31 files changed, 164 insertions(+), 11 deletions(-) diff --git a/Cargo.lock b/Cargo.lock index 3c8d1aacce..9cba952a62 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -1814,8 +1814,7 @@ dependencies = [ [[package]] name = "metal" version = "0.29.0" -source = "registry+https://github.com/rust-lang/crates.io-index" -checksum = "7ecfd3296f8c56b7c1f6fbac3c71cefa9d78ce009850c45000015f206dc7fa21" +source = "git+https://github.com/gfx-rs/metal-rs.git?rev=ae7030be0edff4cda88ece74137e5bcd28ef48fa#ae7030be0edff4cda88ece74137e5bcd28ef48fa" dependencies = [ "bitflags 2.6.0", "block", diff --git a/Cargo.toml b/Cargo.toml index fbf06524c1..d3dd225735 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -136,7 +136,7 @@ winit = { version = "0.29", features = ["android-native-activity"] } # Metal dependencies block = "0.1" core-graphics-types = "0.1" -metal = { version = "0.29.0" } +metal = { version = "0.29.0", git = "https://github.com/gfx-rs/metal-rs.git", rev = "ae7030be0edff4cda88ece74137e5bcd28ef48fa" } objc = "0.2.5" # Vulkan dependencies diff --git a/naga/src/back/glsl/features.rs b/naga/src/back/glsl/features.rs index 362f4bb4f3..6780cfd990 100644 --- a/naga/src/back/glsl/features.rs +++ b/naga/src/back/glsl/features.rs @@ -400,6 +400,7 @@ impl<'a, W> Writer<'a, W> { | StorageFormat::Rgb10a2Uint | StorageFormat::Rgb10a2Unorm | StorageFormat::Rg11b10Ufloat + | StorageFormat::R64Uint | StorageFormat::Rg32Uint | StorageFormat::Rg32Sint | StorageFormat::Rg32Float => { diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index 2ce9f22f27..66a929302a 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -4826,6 +4826,7 @@ fn glsl_storage_format(format: crate::StorageFormat) -> Result<&'static str, Err Sf::Rgb10a2Uint => "rgb10_a2ui", Sf::Rgb10a2Unorm => "rgb10_a2", Sf::Rg11b10Ufloat => "r11f_g11f_b10f", + Sf::R64Uint => "r64ui", Sf::Rg32Uint => "rg32ui", Sf::Rg32Sint => "rg32i", Sf::Rg32Float => "rg32f", diff --git a/naga/src/back/hlsl/conv.rs b/naga/src/back/hlsl/conv.rs index 473fc9476d..9c3b82db5d 100644 --- a/naga/src/back/hlsl/conv.rs +++ b/naga/src/back/hlsl/conv.rs @@ -122,7 +122,7 @@ impl crate::StorageFormat { Self::R16Float | Self::R32Float => "float", Self::R8Unorm | Self::R16Unorm => "unorm float", Self::R8Snorm | Self::R16Snorm => "snorm float", - Self::R8Uint | Self::R16Uint | Self::R32Uint => "uint", + Self::R8Uint | Self::R16Uint | Self::R32Uint | Self::R64Uint => "uint", Self::R8Sint | Self::R16Sint | Self::R32Sint => "int", Self::Rg16Float | Self::Rg32Float => "float2", diff --git a/naga/src/back/spv/instructions.rs b/naga/src/back/spv/instructions.rs index 9bd58508a1..61a38e86e5 100644 --- a/naga/src/back/spv/instructions.rs +++ b/naga/src/back/spv/instructions.rs @@ -1171,6 +1171,7 @@ impl From for spirv::ImageFormat { Sf::Rgb10a2Uint => Self::Rgb10a2ui, Sf::Rgb10a2Unorm => Self::Rgb10A2, Sf::Rg11b10Ufloat => Self::R11fG11fB10f, + Sf::R64Uint => Self::R64ui, Sf::Rg32Uint => Self::Rg32ui, Sf::Rg32Sint => Self::Rg32i, Sf::Rg32Float => Self::Rg32f, diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index 678dcb4246..df71206322 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -1116,10 +1116,13 @@ impl Writer { "storage image format", &[spirv::Capability::StorageImageExtendedFormats], ), - If::R64ui | If::R64i => self.require_any( - "64-bit integer storage image format", - &[spirv::Capability::Int64ImageEXT], - ), + If::R64ui | If::R64i => { + self.use_extension("SPV_EXT_shader_image_int64"); + self.require_any( + "64-bit integer storage image format", + &[spirv::Capability::Int64ImageEXT], + ) + } If::Unknown | If::Rgba32f | If::Rgba16f diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs index e8b942a62c..eb842da497 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -2024,6 +2024,7 @@ const fn storage_format_str(format: crate::StorageFormat) -> &'static str { Sf::Rgb10a2Uint => "rgb10a2uint", Sf::Rgb10a2Unorm => "rgb10a2unorm", Sf::Rg11b10Ufloat => "rg11b10float", + Sf::R64Uint => "r64uint", Sf::Rg32Uint => "rg32uint", Sf::Rg32Sint => "rg32sint", Sf::Rg32Float => "rg32float", diff --git a/naga/src/front/glsl/parser/types.rs b/naga/src/front/glsl/parser/types.rs index 829b9fd897..2f3c18d3e4 100644 --- a/naga/src/front/glsl/parser/types.rs +++ b/naga/src/front/glsl/parser/types.rs @@ -428,6 +428,7 @@ fn map_image_format(word: &str) -> Option { "rgba32ui" => Sf::Rgba32Uint, "rgba16ui" => Sf::Rgba16Uint, "rgba8ui" => Sf::Rgba8Uint, + "r64ui" => Sf::R64Uint, "rg32ui" => Sf::Rg32Uint, "rg16ui" => Sf::Rg16Uint, "rg8ui" => Sf::Rg8Uint, diff --git a/naga/src/front/spv/convert.rs b/naga/src/front/spv/convert.rs index f131db616e..2619eb7f44 100644 --- a/naga/src/front/spv/convert.rs +++ b/naga/src/front/spv/convert.rs @@ -105,6 +105,7 @@ pub(super) fn map_image_format(word: spirv::Word) -> Result Ok(crate::StorageFormat::Rgb10a2Uint), Some(spirv::ImageFormat::Rgb10A2) => Ok(crate::StorageFormat::Rgb10a2Unorm), Some(spirv::ImageFormat::R11fG11fB10f) => Ok(crate::StorageFormat::Rg11b10Ufloat), + Some(spirv::ImageFormat::R64ui) => Ok(crate::StorageFormat::R64Uint), Some(spirv::ImageFormat::Rg32ui) => Ok(crate::StorageFormat::Rg32Uint), Some(spirv::ImageFormat::Rg32i) => Ok(crate::StorageFormat::Rg32Sint), Some(spirv::ImageFormat::Rg32f) => Ok(crate::StorageFormat::Rg32Float), diff --git a/naga/src/front/wgsl/parse/conv.rs b/naga/src/front/wgsl/parse/conv.rs index 3ba71b07cc..ecbf84fd54 100644 --- a/naga/src/front/wgsl/parse/conv.rs +++ b/naga/src/front/wgsl/parse/conv.rs @@ -95,6 +95,7 @@ pub fn map_storage_format(word: &str, span: Span) -> Result Sf::Rgb10a2Uint, "rgb10a2unorm" => Sf::Rgb10a2Unorm, "rg11b10float" => Sf::Rg11b10Ufloat, + "r64uint" => Sf::R64Uint, "rg32uint" => Sf::Rg32Uint, "rg32sint" => Sf::Rg32Sint, "rg32float" => Sf::Rg32Float, diff --git a/naga/src/front/wgsl/to_wgsl.rs b/naga/src/front/wgsl/to_wgsl.rs index 0884e0003b..189010e537 100644 --- a/naga/src/front/wgsl/to_wgsl.rs +++ b/naga/src/front/wgsl/to_wgsl.rs @@ -176,6 +176,7 @@ impl crate::StorageFormat { Sf::Rgb10a2Uint => "rgb10a2uint", Sf::Rgb10a2Unorm => "rgb10a2unorm", Sf::Rg11b10Ufloat => "rg11b10float", + Sf::R64Uint => "r64uint", Sf::Rg32Uint => "rg32uint", Sf::Rg32Sint => "rg32sint", Sf::Rg32Float => "rg32float", diff --git a/naga/src/lib.rs b/naga/src/lib.rs index 85fd7a4508..5f681b3399 100644 --- a/naga/src/lib.rs +++ b/naga/src/lib.rs @@ -625,6 +625,7 @@ pub enum StorageFormat { Rg11b10Ufloat, // 64-bit formats + R64Uint, Rg32Uint, Rg32Sint, Rg32Float, diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs index a5b3ea4e38..4ae38f08e1 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -49,6 +49,7 @@ impl From for super::ScalarKind { Sf::Rgb10a2Uint => Sk::Uint, Sf::Rgb10a2Unorm => Sk::Float, Sf::Rg11b10Ufloat => Sk::Float, + Sf::R64Uint => Sk::Uint, Sf::Rg32Uint => Sk::Uint, Sf::Rg32Sint => Sk::Sint, Sf::Rg32Float => Sk::Float, diff --git a/wgpu-core/src/conv.rs b/wgpu-core/src/conv.rs index d27583b02a..dcaf4c22ab 100644 --- a/wgpu-core/src/conv.rs +++ b/wgpu-core/src/conv.rs @@ -126,6 +126,10 @@ pub fn map_texture_usage( hal::TextureUses::DEPTH_STENCIL_READ | hal::TextureUses::DEPTH_STENCIL_WRITE, usage.contains(wgt::TextureUsages::RENDER_ATTACHMENT) && !is_color, ); + u.set( + hal::TextureUses::SHADER_ATOMIC, + usage.contains(wgt::TextureUsages::SHADER_ATOMIC), + ); u } @@ -177,6 +181,10 @@ pub fn map_texture_usage_from_hal(uses: hal::TextureUses) -> wgt::TextureUsages wgt::TextureUsages::RENDER_ATTACHMENT, uses.contains(hal::TextureUses::COLOR_TARGET), ); + u.set( + wgt::TextureUsages::SHADER_ATOMIC, + uses.contains(hal::TextureUses::SHADER_ATOMIC), + ); u } diff --git a/wgpu-core/src/instance.rs b/wgpu-core/src/instance.rs index 581c5ce0d9..0abcd4c234 100644 --- a/wgpu-core/src/instance.rs +++ b/wgpu-core/src/instance.rs @@ -216,6 +216,10 @@ impl Adapter { wgt::TextureUsages::RENDER_ATTACHMENT, caps.intersects(Tfc::COLOR_ATTACHMENT | Tfc::DEPTH_STENCIL_ATTACHMENT), ); + allowed_usages.set( + wgt::TextureUsages::SHADER_ATOMIC, + caps.intersects(Tfc::SHADER_ATOMIC), + ); let mut flags = wgt::TextureFormatFeatureFlags::empty(); flags.set( diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs index c1cbdaf183..3134930eb7 100644 --- a/wgpu-core/src/validation.rs +++ b/wgpu-core/src/validation.rs @@ -277,6 +277,7 @@ fn map_storage_format_to_naga(format: wgt::TextureFormat) -> Option Sf::Rgb10a2Unorm, Tf::Rg11b10Ufloat => Sf::Rg11b10Ufloat, + Tf::R64Uint => Sf::R64Uint, Tf::Rg32Uint => Sf::Rg32Uint, Tf::Rg32Sint => Sf::Rg32Sint, Tf::Rg32Float => Sf::Rg32Float, @@ -333,6 +334,7 @@ fn map_storage_format_from_naga(format: naga::StorageFormat) -> wgt::TextureForm Sf::Rgb10a2Unorm => Tf::Rgb10a2Unorm, Sf::Rg11b10Ufloat => Tf::Rg11b10Ufloat, + Sf::R64Uint => Tf::R64Uint, Sf::Rg32Uint => Tf::Rg32Uint, Sf::Rg32Sint => Tf::Rg32Sint, Sf::Rg32Float => Tf::Rg32Float, @@ -635,6 +637,7 @@ impl NumericType { Tf::Rg8Unorm | Tf::Rg8Snorm | Tf::Rg16Float | Tf::Rg32Float => { (NumericDimension::Vector(Vs::Bi), Scalar::F32) } + Tf::R64Uint => (NumericDimension::Scalar, Scalar::U64), Tf::Rg8Uint | Tf::Rg16Uint | Tf::Rg32Uint => { (NumericDimension::Vector(Vs::Bi), Scalar::U32) } diff --git a/wgpu-hal/src/auxil/dxgi/conv.rs b/wgpu-hal/src/auxil/dxgi/conv.rs index 878dab39e9..10a584d565 100644 --- a/wgpu-hal/src/auxil/dxgi/conv.rs +++ b/wgpu-hal/src/auxil/dxgi/conv.rs @@ -48,6 +48,7 @@ pub fn map_texture_format_failable( Tf::Rgb10a2Uint => DXGI_FORMAT_R10G10B10A2_UINT, Tf::Rgb10a2Unorm => DXGI_FORMAT_R10G10B10A2_UNORM, Tf::Rg11b10Ufloat => DXGI_FORMAT_R11G11B10_FLOAT, + Tf::R64Uint => DXGI_FORMAT_R32G32_UINT, // R64 emulated by R32G32 Tf::Rg32Uint => DXGI_FORMAT_R32G32_UINT, Tf::Rg32Sint => DXGI_FORMAT_R32G32_SINT, Tf::Rg32Float => DXGI_FORMAT_R32G32_FLOAT, diff --git a/wgpu-hal/src/dx12/adapter.rs b/wgpu-hal/src/dx12/adapter.rs index 45d69f5584..4b510290d9 100644 --- a/wgpu-hal/src/dx12/adapter.rs +++ b/wgpu-hal/src/dx12/adapter.rs @@ -402,6 +402,8 @@ impl super::Adapter { wgt::Features::SHADER_INT64_ATOMIC_ALL_OPS | wgt::Features::SHADER_INT64_ATOMIC_MIN_MAX, atomic_int64_on_typed_resource_supported, ); + // Ruint64 textures are always emulated on d3d12 + features.set(wgt::Features::TEXTURE_INT64_ATOMIC, true); // float32-filterable should always be available on d3d12 features.set(wgt::Features::FLOAT32_FILTERABLE, true); diff --git a/wgpu-hal/src/gles/adapter.rs b/wgpu-hal/src/gles/adapter.rs index a654215d21..6a04e3bbf6 100644 --- a/wgpu-hal/src/gles/adapter.rs +++ b/wgpu-hal/src/gles/adapter.rs @@ -1077,6 +1077,8 @@ impl crate::Adapter for super::Adapter { let texture_float_linear = feature_fn(wgt::Features::FLOAT32_FILTERABLE, filterable); + let image_atomic = feature_fn(wgt::Features::TEXTURE_INT64_ATOMIC, Tfc::SHADER_ATOMIC); + match format { Tf::R8Unorm => filterable_renderable, Tf::R8Snorm => filterable, @@ -1108,6 +1110,7 @@ impl crate::Adapter for super::Adapter { Tf::Rgb10a2Uint => renderable, Tf::Rgb10a2Unorm => filterable_renderable, Tf::Rg11b10Ufloat => filterable | float_renderable, + Tf::R64Uint => image_atomic, Tf::Rg32Uint => renderable, Tf::Rg32Sint => renderable, Tf::Rg32Float => unfilterable | float_renderable | texture_float_linear, diff --git a/wgpu-hal/src/gles/conv.rs b/wgpu-hal/src/gles/conv.rs index 3a6d5ebb2e..be6cfea203 100644 --- a/wgpu-hal/src/gles/conv.rs +++ b/wgpu-hal/src/gles/conv.rs @@ -50,6 +50,7 @@ impl super::AdapterShared { glow::RGB, glow::UNSIGNED_INT_10F_11F_11F_REV, ), + Tf::R64Uint => unreachable!(), Tf::Rg32Uint => (glow::RG32UI, glow::RG_INTEGER, glow::UNSIGNED_INT), Tf::Rg32Sint => (glow::RG32I, glow::RG_INTEGER, glow::INT), Tf::Rg32Float => (glow::RG32F, glow::RG, glow::FLOAT), diff --git a/wgpu-hal/src/lib.rs b/wgpu-hal/src/lib.rs index 6578252c1a..f69bc88fee 100644 --- a/wgpu-hal/src/lib.rs +++ b/wgpu-hal/src/lib.rs @@ -1570,6 +1570,9 @@ bitflags!( const COPY_SRC = 1 << 14; /// Format can be copied to. const COPY_DST = 1 << 15; + + /// Format can be used with image atomics + const SHADER_ATOMIC = 1 << 16; } ); @@ -1733,6 +1736,8 @@ bitflags::bitflags! { /// Flag used by the wgpu-core texture tracker to say that the tracker does not know the state of the sub-resource. /// This is different from UNINITIALIZED as that says the tracker does know, but the texture has not been initialized. const UNKNOWN = 1 << 11; + /// Image atomic enabled storage + const SHADER_ATOMIC = 1 << 12; } } diff --git a/wgpu-hal/src/metal/adapter.rs b/wgpu-hal/src/metal/adapter.rs index e7db97a1f9..92f21fda46 100644 --- a/wgpu-hal/src/metal/adapter.rs +++ b/wgpu-hal/src/metal/adapter.rs @@ -109,6 +109,12 @@ impl crate::Adapter for super::Adapter { ], ); + let image_atomic = if pc.int64_image_atomics { + Tfc::SHADER_ATOMIC + } else { + Tfc::empty() + }; + // Metal defined pixel format capabilities let all_caps = Tfc::SAMPLED_LINEAR | Tfc::STORAGE @@ -183,6 +189,7 @@ impl crate::Adapter for super::Adapter { flags.set(Tfc::STORAGE, pc.format_rg11b10_all); flags } + Tf::R64Uint => Tfc::COLOR_ATTACHMENT | Tfc::STORAGE | image_atomic, Tf::Rg32Uint | Tf::Rg32Sint => Tfc::COLOR_ATTACHMENT | Tfc::STORAGE | msaa_count, Tf::Rg32Float => { if pc.format_rg32float_all { @@ -828,6 +835,7 @@ impl super::PrivateCapabilities { && ((device.supports_family(MTLGPUFamily::Apple8) && device.supports_family(MTLGPUFamily::Mac2)) || device.supports_family(MTLGPUFamily::Apple9)), + int64_image_atomics: family_check && device.supports_family(MTLGPUFamily::Apple6), } } @@ -908,6 +916,10 @@ impl super::PrivateCapabilities { F::SHADER_INT64_ATOMIC_MIN_MAX, self.int64_atomics && self.msl_version >= MTLLanguageVersion::V2_4, ); + features.set( + F::TEXTURE_INT64_ATOMIC, + self.int64_image_atomics && self.msl_version >= MTLLanguageVersion::V3_1, + ); features.set( F::ADDRESS_MODE_CLAMP_TO_BORDER, @@ -1041,6 +1053,7 @@ impl super::PrivateCapabilities { Tf::Rgb10a2Uint => RGB10A2Uint, Tf::Rgb10a2Unorm => RGB10A2Unorm, Tf::Rg11b10Ufloat => RG11B10Float, + Tf::R64Uint => RG32Uint, Tf::Rg32Uint => RG32Uint, Tf::Rg32Sint => RG32Sint, Tf::Rg32Float => RG32Float, diff --git a/wgpu-hal/src/metal/conv.rs b/wgpu-hal/src/metal/conv.rs index 6ebabee1a6..e26b1e85b4 100644 --- a/wgpu-hal/src/metal/conv.rs +++ b/wgpu-hal/src/metal/conv.rs @@ -27,6 +27,11 @@ pub fn map_texture_usage( format.is_combined_depth_stencil_format(), ); + mtl_usage.set( + metal::MTLTextureUsage::ShaderAtomic, + usage.intersects(Tu::SHADER_ATOMIC), + ); + mtl_usage } diff --git a/wgpu-hal/src/metal/mod.rs b/wgpu-hal/src/metal/mod.rs index 1935e843ec..61fd05909e 100644 --- a/wgpu-hal/src/metal/mod.rs +++ b/wgpu-hal/src/metal/mod.rs @@ -289,6 +289,7 @@ struct PrivateCapabilities { supports_simd_scoped_operations: bool, int64: bool, int64_atomics: bool, + int64_image_atomics: bool, } #[derive(Clone, Debug)] diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs index ab6ae02c6f..0ce145baf4 100644 --- a/wgpu-hal/src/vulkan/adapter.rs +++ b/wgpu-hal/src/vulkan/adapter.rs @@ -109,6 +109,9 @@ pub struct PhysicalDeviceFeatures { /// Features provided by `VK_KHR_shader_atomic_int64`, promoted to Vulkan 1.2. shader_atomic_int64: Option>, + /// Features provided by `VK_EXT_shader_image_atomic_int64` + shader_image_atomic_int64: Option>, + /// Features provided by `VK_EXT_subgroup_size_control`, promoted to Vulkan 1.3. subgroup_size_control: Option>, } @@ -157,6 +160,9 @@ impl PhysicalDeviceFeatures { if let Some(ref mut feature) = self.shader_atomic_int64 { info = info.push_next(feature); } + if let Some(ref mut feature) = self.shader_image_atomic_int64 { + info = info.push_next(feature); + } if let Some(ref mut feature) = self.subgroup_size_control { info = info.push_next(feature); } @@ -438,6 +444,17 @@ impl PhysicalDeviceFeatures { } else { None }, + shader_image_atomic_int64: if enabled_extensions + .contains(&ext::shader_image_atomic_int64::NAME) + { + let needed = requested_features.intersects(wgt::Features::TEXTURE_INT64_ATOMIC); + Some( + vk::PhysicalDeviceShaderImageAtomicInt64FeaturesEXT::default() + .shader_image_int64_atomics(needed), + ) + } else { + None + }, subgroup_size_control: if device_api_version >= vk::API_VERSION_1_3 || enabled_extensions.contains(&ext::subgroup_size_control::NAME) { @@ -590,6 +607,16 @@ impl PhysicalDeviceFeatures { ); } + if let Some(ref shader_image_atomic_int64) = self.shader_image_atomic_int64 { + features.set( + F::TEXTURE_INT64_ATOMIC, + shader_image_atomic_int64 + .shader_image_int64_atomics(true) + .shader_image_int64_atomics + != 0, + ); + } + //if caps.supports_extension(khr::sampler_mirror_clamp_to_edge::NAME) { //if caps.supports_extension(ext::sampler_filter_minmax::NAME) { features.set( @@ -1010,6 +1037,11 @@ impl PhysicalDeviceProperties { extensions.push(khr::shader_atomic_int64::NAME); } + // Require `VK_EXT_shader_image_atomic_int64` if the associated feature was requested + if requested_features.intersects(wgt::Features::TEXTURE_INT64_ATOMIC) { + extensions.push(ext::shader_image_atomic_int64::NAME); + } + // Require VK_GOOGLE_display_timing if the associated feature was requested if requested_features.contains(wgt::Features::VULKAN_GOOGLE_DISPLAY_TIMING) { extensions.push(google::display_timing::NAME); @@ -1293,6 +1325,13 @@ impl super::InstanceShared { features2 = features2.push_next(next); } + if capabilities.supports_extension(ext::shader_image_atomic_int64::NAME) { + let next = features + .shader_image_atomic_int64 + .insert(vk::PhysicalDeviceShaderImageAtomicInt64FeaturesEXT::default()); + features2 = features2.push_next(next); + } + if capabilities.supports_extension(ext::image_robustness::NAME) { let next = features .image_robustness @@ -1786,6 +1825,10 @@ impl super::Adapter { capabilities.push(spv::Capability::Int64Atomics); } + if features.intersects(wgt::Features::TEXTURE_INT64_ATOMIC) { + capabilities.push(spv::Capability::Int64ImageEXT); + } + let mut flags = spv::WriterFlags::empty(); flags.set( spv::WriterFlags::DEBUG, @@ -2119,6 +2162,10 @@ impl crate::Adapter for super::Adapter { Tfc::COPY_DST, features.intersects(vk::FormatFeatureFlags::TRANSFER_DST), ); + flags.set( + Tfc::SHADER_ATOMIC, + features.intersects(vk::FormatFeatureFlags::STORAGE_IMAGE_ATOMIC), + ); // Vulkan is very permissive about MSAA flags.set(Tfc::MULTISAMPLE_RESOLVE, !format.is_compressed()); diff --git a/wgpu-hal/src/vulkan/conv.rs b/wgpu-hal/src/vulkan/conv.rs index b829307068..f7cce466cb 100644 --- a/wgpu-hal/src/vulkan/conv.rs +++ b/wgpu-hal/src/vulkan/conv.rs @@ -37,6 +37,7 @@ impl super::PrivateCapabilities { Tf::Rgb10a2Uint => F::A2B10G10R10_UINT_PACK32, Tf::Rgb10a2Unorm => F::A2B10G10R10_UNORM_PACK32, Tf::Rg11b10Ufloat => F::B10G11R11_UFLOAT_PACK32, + Tf::R64Uint => F::R64_UINT, Tf::Rg32Uint => F::R32G32_UINT, Tf::Rg32Sint => F::R32G32_SINT, Tf::Rg32Float => F::R32G32_SFLOAT, diff --git a/wgpu-info/src/texture.rs b/wgpu-info/src/texture.rs index 2487bf350f..64325f0e5b 100644 --- a/wgpu-info/src/texture.rs +++ b/wgpu-info/src/texture.rs @@ -1,6 +1,6 @@ // Lets keep these on one line #[rustfmt::skip] -pub const TEXTURE_FORMAT_LIST: [wgpu::TextureFormat; 116] = [ +pub const TEXTURE_FORMAT_LIST: [wgpu::TextureFormat; 117] = [ wgpu::TextureFormat::R8Unorm, wgpu::TextureFormat::R8Snorm, wgpu::TextureFormat::R8Uint, @@ -33,6 +33,7 @@ pub const TEXTURE_FORMAT_LIST: [wgpu::TextureFormat; 116] = [ wgpu::TextureFormat::Rgb10a2Uint, wgpu::TextureFormat::Rgb10a2Unorm, wgpu::TextureFormat::Rg11b10Ufloat, + wgpu::TextureFormat::R64Uint, wgpu::TextureFormat::Rg32Uint, wgpu::TextureFormat::Rg32Sint, wgpu::TextureFormat::Rg32Float, diff --git a/wgpu-types/src/lib.rs b/wgpu-types/src/lib.rs index c7167f826f..0650c00abc 100644 --- a/wgpu-types/src/lib.rs +++ b/wgpu-types/src/lib.rs @@ -969,6 +969,15 @@ bitflags::bitflags! { /// [VK_GOOGLE_display_timing]: https://registry.khronos.org/vulkan/specs/1.3-extensions/man/html/VK_GOOGLE_display_timing.html /// [`Surface::as_hal()`]: https://docs.rs/wgpu/latest/wgpu/struct.Surface.html#method.as_hal const VULKAN_GOOGLE_DISPLAY_TIMING = 1 << 62; + /// Enables R64Uint texture atomic min and max. + /// + /// Supported platforms: + /// - Vulkan (with VK_EXT_shader_image_atomic_int64) + /// - DX12 (with SM 6.6+ emulated via Rg32Uint texture) + /// - Metal (with MSL 3.1+ emulated via RG32Uint texture) + /// + /// This is a native only feature. + const TEXTURE_INT64_ATOMIC = 1 << 63; } } @@ -2557,6 +2566,10 @@ pub enum TextureFormat { Rg11b10Ufloat, // Normal 64 bit formats + /// Red channel only. 64 bit integer per channel. Unsigned in shader. + /// + /// [`Features::TEXTURE_INT64_ATOMIC`] must be enabled to use this texture format. + R64Uint, /// Red and green channels. 32 bit integer per channel. Unsigned in shader. Rg32Uint, /// Red and green channels. 32 bit integer per channel. Signed in shader. @@ -2843,6 +2856,7 @@ impl<'de> Deserialize<'de> for TextureFormat { "rgb10a2uint" => TextureFormat::Rgb10a2Uint, "rgb10a2unorm" => TextureFormat::Rgb10a2Unorm, "rg11b10ufloat" => TextureFormat::Rg11b10Ufloat, + "r64uint" => TextureFormat::R64Uint, "rg32uint" => TextureFormat::Rg32Uint, "rg32sint" => TextureFormat::Rg32Sint, "rg32float" => TextureFormat::Rg32Float, @@ -2971,6 +2985,7 @@ impl Serialize for TextureFormat { TextureFormat::Rgb10a2Uint => "rgb10a2uint", TextureFormat::Rgb10a2Unorm => "rgb10a2unorm", TextureFormat::Rg11b10Ufloat => "rg11b10ufloat", + TextureFormat::R64Uint => "r64uint", TextureFormat::Rg32Uint => "rg32uint", TextureFormat::Rg32Sint => "rg32sint", TextureFormat::Rg32Float => "rg32float", @@ -3213,6 +3228,7 @@ impl TextureFormat { | Self::Rgb10a2Uint | Self::Rgb10a2Unorm | Self::Rg11b10Ufloat + | Self::R64Uint | Self::Rg32Uint | Self::Rg32Sint | Self::Rg32Float @@ -3336,6 +3352,8 @@ impl TextureFormat { | Self::Depth24PlusStencil8 | Self::Depth32Float => Features::empty(), + Self::R64Uint => Features::TEXTURE_INT64_ATOMIC, + Self::Depth32FloatStencil8 => Features::DEPTH32FLOAT_STENCIL8, Self::NV12 => Features::TEXTURE_FORMAT_NV12, @@ -3440,6 +3458,7 @@ impl TextureFormat { Self::Rgb10a2Uint => ( msaa, attachment), Self::Rgb10a2Unorm => (msaa_resolve, attachment), Self::Rg11b10Ufloat => ( msaa, rg11b10f), + Self::R64Uint => ( noaa, attachment), Self::Rg32Uint => ( noaa, all_flags), Self::Rg32Sint => ( noaa, all_flags), Self::Rg32Float => ( noaa, all_flags), @@ -3561,6 +3580,7 @@ impl TextureFormat { | Self::Rg16Uint | Self::Rgba16Uint | Self::R32Uint + | Self::R64Uint | Self::Rg32Uint | Self::Rgba32Uint | Self::Rgb10a2Uint => Some(uint), @@ -3691,7 +3711,7 @@ impl TextureFormat { | Self::Rgba16Uint | Self::Rgba16Sint | Self::Rgba16Float => Some(8), - Self::Rg32Uint | Self::Rg32Sint | Self::Rg32Float => Some(8), + Self::R64Uint | Self::Rg32Uint | Self::Rg32Sint | Self::Rg32Float => Some(8), Self::Rgba32Uint | Self::Rgba32Sint | Self::Rgba32Float => Some(16), @@ -3780,6 +3800,7 @@ impl TextureFormat { | Self::Rgba16Unorm | Self::Rgba16Snorm | Self::Rgba16Float + | Self::R64Uint | Self::Rg32Uint | Self::Rg32Sint | Self::Rg32Float @@ -3860,6 +3881,7 @@ impl TextureFormat { Self::R32Uint | Self::R32Sint | Self::R32Float + | Self::R64Uint | Self::Rg32Uint | Self::Rg32Sint | Self::Rg32Float @@ -3928,7 +3950,8 @@ impl TextureFormat { | Self::R16Float | Self::R32Uint | Self::R32Sint - | Self::R32Float => 1, + | Self::R32Float + | Self::R64Uint => 1, Self::Rg8Unorm | Self::Rg8Snorm @@ -4180,6 +4203,10 @@ fn texture_format_serialize() { serde_json::to_string(&TextureFormat::Rg11b10Ufloat).unwrap(), "\"rg11b10ufloat\"".to_string() ); + assert_eq!( + serde_json::to_string(&TextureFormat::R64Uint).unwrap(), + "\"r64uint\"".to_string() + ); assert_eq!( serde_json::to_string(&TextureFormat::Rg32Uint).unwrap(), "\"rg32uint\"".to_string() @@ -4476,6 +4503,10 @@ fn texture_format_deserialize() { serde_json::from_str::("\"rg11b10ufloat\"").unwrap(), TextureFormat::Rg11b10Ufloat ); + assert_eq!( + serde_json::from_str::("\"r64uint\"").unwrap(), + TextureFormat::R64Uint + ); assert_eq!( serde_json::from_str::("\"rg32uint\"").unwrap(), TextureFormat::Rg32Uint @@ -5485,6 +5516,11 @@ bitflags::bitflags! { #[repr(transparent)] #[derive(Debug, Copy, Clone, PartialEq, Eq, Hash)] pub struct TextureUsages: u32 { + // + // ---- Start numbering at 1 << 0 ---- + // + // WebGPU features: + // /// Allows a texture to be the source in a [`CommandEncoder::copy_texture_to_buffer`] or /// [`CommandEncoder::copy_texture_to_texture`] operation. const COPY_SRC = 1 << 0; @@ -5497,6 +5533,14 @@ bitflags::bitflags! { const STORAGE_BINDING = 1 << 3; /// Allows a texture to be an output attachment of a render pass. const RENDER_ATTACHMENT = 1 << 4; + + // + // ---- Restart Numbering for Native Features --- + // + // Native Features: + // + /// Allows a texture to be used with image atomics. Requires [`Features::TEXTURE_INT64_ATOMIC`] + const SHADER_ATOMIC = 1 << 16; } } diff --git a/wgpu/src/backend/webgpu.rs b/wgpu/src/backend/webgpu.rs index e982300e70..a1fb459fb9 100644 --- a/wgpu/src/backend/webgpu.rs +++ b/wgpu/src/backend/webgpu.rs @@ -264,6 +264,7 @@ fn map_texture_format(texture_format: wgt::TextureFormat) -> webgpu_sys::GpuText TextureFormat::Rgb10a2Unorm => tf::Rgb10a2unorm, TextureFormat::Rg11b10Ufloat => tf::Rg11b10ufloat, // 64-bit formats + TextureFormat::R64Uint => tf::R64uint, TextureFormat::Rg32Uint => tf::Rg32uint, TextureFormat::Rg32Sint => tf::Rg32sint, TextureFormat::Rg32Float => tf::Rg32float, diff --git a/wgpu/src/backend/webgpu/webgpu_sys/gen_GpuTextureFormat.rs b/wgpu/src/backend/webgpu/webgpu_sys/gen_GpuTextureFormat.rs index cf07abd384..d96e5a94b1 100644 --- a/wgpu/src/backend/webgpu/webgpu_sys/gen_GpuTextureFormat.rs +++ b/wgpu/src/backend/webgpu/webgpu_sys/gen_GpuTextureFormat.rs @@ -46,6 +46,7 @@ pub enum GpuTextureFormat { Rgb10a2uint = "rgb10a2uint", Rgb10a2unorm = "rgb10a2unorm", Rg11b10ufloat = "rg11b10ufloat", + R64uint = "r64uint", Rg32uint = "rg32uint", Rg32sint = "rg32sint", Rg32float = "rg32float", From ed99b857c7c2fb19714eca8be5e6cf0eba718b0a Mon Sep 17 00:00:00 2001 From: atlas dostal Date: Wed, 2 Oct 2024 17:15:08 -0400 Subject: [PATCH 2/9] refactor texture format to scalar conversion --- naga/src/back/glsl/mod.rs | 18 +++++++++++------- naga/src/back/msl/writer.rs | 13 +++++++++---- naga/src/back/spv/mod.rs | 11 +++++++---- naga/src/back/spv/writer.rs | 5 +---- naga/src/front/glsl/functions.rs | 2 +- naga/src/proc/mod.rs | 11 ++++++++--- naga/src/proc/typifier.rs | 5 +---- naga/src/valid/function.rs | 5 +---- 8 files changed, 39 insertions(+), 31 deletions(-) diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index 66a929302a..2e62f6c285 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -1095,12 +1095,16 @@ impl<'a, W: Write> Writer<'a, W> { // - Array - used if it's an image array // - Shadow - used if it's a depth image use crate::ImageClass as Ic; - - let (base, kind, ms, comparison) = match class { - Ic::Sampled { kind, multi: true } => ("sampler", kind, "MS", ""), - Ic::Sampled { kind, multi: false } => ("sampler", kind, "", ""), - Ic::Depth { multi: true } => ("sampler", crate::ScalarKind::Float, "MS", ""), - Ic::Depth { multi: false } => ("sampler", crate::ScalarKind::Float, "", "Shadow"), + use crate::Scalar as S; + let float = S { + kind: crate::ScalarKind::Float, + width: 4, + }; + let (base, scalar, ms, comparison) = match class { + Ic::Sampled { kind, multi: true } => ("sampler", S { kind, width: 4 }, "MS", ""), + Ic::Sampled { kind, multi: false } => ("sampler", S { kind, width: 4 }, "", ""), + Ic::Depth { multi: true } => ("sampler", float, "MS", ""), + Ic::Depth { multi: false } => ("sampler", float, "", "Shadow"), Ic::Storage { format, .. } => ("image", format.into(), "", ""), }; @@ -1114,7 +1118,7 @@ impl<'a, W: Write> Writer<'a, W> { self.out, "{}{}{}{}{}{}{}", precision, - glsl_scalar(crate::Scalar { kind, width: 4 })?.prefix, + glsl_scalar(scalar)?.prefix, base, glsl_dimension(dim), ms, diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index 7ab97f491c..99ca6c3e88 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -214,14 +214,15 @@ impl<'a> Display for TypeContext<'a> { crate::ImageDimension::D3 => "3d", crate::ImageDimension::Cube => "cube", }; - let (texture_str, msaa_str, kind, access) = match class { + let (texture_str, msaa_str, scalar, access) = match class { crate::ImageClass::Sampled { kind, multi } => { let (msaa_str, access) = if multi { ("_ms", "read") } else { ("", "sample") }; - ("texture", msaa_str, kind, access) + let scalar = crate::Scalar { kind, width: 4 }; + ("texture", msaa_str, scalar, access) } crate::ImageClass::Depth { multi } => { let (msaa_str, access) = if multi { @@ -229,7 +230,11 @@ impl<'a> Display for TypeContext<'a> { } else { ("", "sample") }; - ("depth", msaa_str, crate::ScalarKind::Float, access) + let scalar = crate::Scalar { + kind: crate::ScalarKind::Float, + width: 4, + }; + ("depth", msaa_str, scalar, access) } crate::ImageClass::Storage { format, .. } => { let access = if self @@ -253,7 +258,7 @@ impl<'a> Display for TypeContext<'a> { ("texture", "", format.into(), access) } }; - let base_name = crate::Scalar { kind, width: 4 }.to_msl_name(); + let base_name = scalar.to_msl_name(); let array_str = if arrayed { "_array" } else { "" }; write!( out, diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index 32bd1fcecf..1cbc49b712 100644 --- a/naga/src/back/spv/mod.rs +++ b/naga/src/back/spv/mod.rs @@ -178,7 +178,7 @@ impl Function { /// where practical. #[derive(Debug, PartialEq, Hash, Eq, Copy, Clone)] struct LocalImageType { - sampled_type: crate::ScalarKind, + sampled_type: crate::Scalar, dim: spirv::Dim, flags: ImageTypeFlags, image_format: spirv::ImageFormat, @@ -209,19 +209,22 @@ impl LocalImageType { match class { crate::ImageClass::Sampled { kind, multi } => LocalImageType { - sampled_type: kind, + sampled_type: crate::Scalar { kind, width: 4 }, dim, flags: make_flags(multi, ImageTypeFlags::SAMPLED), image_format: spirv::ImageFormat::Unknown, }, crate::ImageClass::Depth { multi } => LocalImageType { - sampled_type: crate::ScalarKind::Float, + sampled_type: crate::Scalar { + kind: crate::ScalarKind::Float, + width: 4, + }, dim, flags: make_flags(multi, ImageTypeFlags::DEPTH | ImageTypeFlags::SAMPLED), image_format: spirv::ImageFormat::Unknown, }, crate::ImageClass::Storage { format, access: _ } => LocalImageType { - sampled_type: crate::ScalarKind::from(format), + sampled_type: format.into(), dim, flags: make_flags(false, ImageTypeFlags::empty()), image_format: format.into(), diff --git a/naga/src/back/spv/writer.rs b/naga/src/back/spv/writer.rs index df71206322..431114af5b 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -936,10 +936,7 @@ impl Writer { LocalType::Image(image) => { let local_type = LocalType::Value { vector_size: None, - scalar: crate::Scalar { - kind: image.sampled_type, - width: 4, - }, + scalar: image.sampled_type, pointer_space: None, }; let type_id = self.get_type_id(LookupType::Local(local_type)); diff --git a/naga/src/front/glsl/functions.rs b/naga/src/front/glsl/functions.rs index a1a6038263..c02081fdd1 100644 --- a/naga/src/front/glsl/functions.rs +++ b/naga/src/front/glsl/functions.rs @@ -622,7 +622,7 @@ impl Frontend { // check that the format scalar kind matches let good_format = overload_format == call_format || (overload.internal - && ScalarKind::from(overload_format) == ScalarKind::from(call_format)); + && Scalar::from(overload_format) == Scalar::from(call_format)); if !(good_size && good_format) { continue 'outer; } diff --git a/naga/src/proc/mod.rs b/naga/src/proc/mod.rs index 4ae38f08e1..3916d19fec 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -20,10 +20,10 @@ pub use namer::{EntryPointIndex, NameKey, Namer}; pub use terminator::ensure_block_returns; pub use typifier::{ResolveContext, ResolveError, TypeResolution}; -impl From for super::ScalarKind { +impl From for super::Scalar { fn from(format: super::StorageFormat) -> Self { use super::{ScalarKind as Sk, StorageFormat as Sf}; - match format { + let kind = match format { Sf::R8Unorm => Sk::Float, Sf::R8Snorm => Sk::Float, Sf::R8Uint => Sk::Uint, @@ -65,7 +65,12 @@ impl From for super::ScalarKind { Sf::Rg16Snorm => Sk::Float, Sf::Rgba16Unorm => Sk::Float, Sf::Rgba16Snorm => Sk::Float, - } + }; + let width = match format { + Sf::R64Uint => 8, + _ => 4, + }; + super::Scalar { kind, width } } } diff --git a/naga/src/proc/typifier.rs b/naga/src/proc/typifier.rs index d8af0cd236..3336c18dda 100644 --- a/naga/src/proc/typifier.rs +++ b/naga/src/proc/typifier.rs @@ -487,10 +487,7 @@ impl<'a> ResolveContext<'a> { size: crate::VectorSize::Quad, }, crate::ImageClass::Storage { format, .. } => Ti::Vector { - scalar: crate::Scalar { - kind: format.into(), - width: 4, - }, + scalar: format.into(), size: crate::VectorSize::Quad, }, }), diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index 23e6204ccb..c695d65144 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -1089,10 +1089,7 @@ impl super::Validator { crate::ImageClass::Storage { format, .. } => { crate::TypeInner::Vector { size: crate::VectorSize::Quad, - scalar: crate::Scalar { - kind: format.into(), - width: 4, - }, + scalar: format.into(), } } _ => { From e13ce215a8e70a34297fb7c35596f1735dc6099d Mon Sep 17 00:00:00 2001 From: atlas dostal Date: Wed, 2 Oct 2024 17:16:38 -0400 Subject: [PATCH 3/9] naga valid back and front for image atomics --- naga/src/back/dot/mod.rs | 13 ++++ naga/src/back/glsl/mod.rs | 61 +++++++++++++++ naga/src/back/hlsl/writer.rs | 19 +++++ naga/src/back/msl/writer.rs | 36 +++++++++ naga/src/back/pipeline_constants.rs | 12 +++ naga/src/back/spv/block.rs | 9 +++ naga/src/back/spv/image.rs | 68 +++++++++++++++++ naga/src/back/spv/instructions.rs | 35 +++++++++ naga/src/back/wgsl/writer.rs | 18 +++++ naga/src/compact/statements.rs | 24 ++++++ naga/src/front/spv/mod.rs | 1 + naga/src/front/wgsl/lower/mod.rs | 36 +++++++++ naga/src/front/wgsl/parse/mod.rs | 4 + naga/src/lib.rs | 48 ++++++++++++ naga/src/proc/terminator.rs | 1 + naga/src/valid/analyzer.rs | 13 ++++ naga/src/valid/expression.rs | 4 + naga/src/valid/function.rs | 111 ++++++++++++++++++++++++++++ naga/src/valid/handles.rs | 13 ++++ 19 files changed, 526 insertions(+) diff --git a/naga/src/back/dot/mod.rs b/naga/src/back/dot/mod.rs index 2780879657..6733ac115c 100644 --- a/naga/src/back/dot/mod.rs +++ b/naga/src/back/dot/mod.rs @@ -254,6 +254,19 @@ impl StatementGraph { } "Atomic" } + S::ImageAtomic { + image, + coordinate, + sample, + fun: _, + value, + } => { + self.dependencies.push((id, image, "image")); + self.dependencies.push((id, coordinate, "coordinate")); + self.dependencies.push((id, sample, "sample")); + self.dependencies.push((id, value, "value")); + "ImageAtomic" + } S::WorkGroupUniformLoad { pointer, result } => { self.emits.push((id, result)); self.dependencies.push((id, pointer, "pointer")); diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index 2e62f6c285..2b0c18abc8 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -2469,6 +2469,17 @@ impl<'a, W: Write> Writer<'a, W> { self.write_expr(value, ctx)?; writeln!(self.out, ");")?; } + // Stores a value into an image. + Statement::ImageAtomic { + image, + coordinate, + sample, + fun, + value, + } => { + write!(self.out, "{level}")?; + self.write_image_atomic(ctx, image, coordinate, sample, fun, value)? + } Statement::RayQuery { .. } => unreachable!(), Statement::SubgroupBallot { result, predicate } => { write!(self.out, "{level}")?; @@ -4089,6 +4100,56 @@ impl<'a, W: Write> Writer<'a, W> { Ok(()) } + /// Helper method to write the `ImageStore` statement + fn write_image_atomic( + &mut self, + ctx: &back::FunctionCtx, + image: Handle, + coordinate: Handle, + _sample: Handle, + fun: crate::AtomicFunction, + value: Handle, + ) -> Result<(), Error> { + use crate::ImageDimension as IDim; + + // NOTE: openGL requires that `imageStore`s have no effets when the texel is invalid + // so we don't need to generate bounds checks (OpenGL 4.2 Core ยง3.9.20) + + // This will only panic if the module is invalid + let dim = match *ctx.resolve_type(image, &self.module.types) { + TypeInner::Image { dim, .. } => dim, + _ => unreachable!(), + }; + + // Begin our call to `imageStore` + let fun_str = fun.to_glsl(); + write!(self.out, "imageAtomic{fun_str}(")?; + self.write_expr(image, ctx)?; + // Separate the image argument from the coordinates + write!(self.out, ", ")?; + + // openGL es doesn't have 1D images so we need workaround it + let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es(); + // Write the coordinate vector + self.write_texture_coord( + ctx, + // Get the size of the coordinate vector + self.get_coordinate_vector_size(dim, false), + coordinate, + None, + tex_1d_hack, + )?; + + // Separate the coordinate from the value to write and write the expression + // of the value to write. + write!(self.out, ", ")?; + self.write_expr(value, ctx)?; + // End the call to `imageStore` and the statement. + writeln!(self.out, ");")?; + + Ok(()) + } + /// Helper method for writing an `ImageLoad` expression. #[allow(clippy::too_many_arguments)] fn write_image_load( diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index 0eb18f0e16..ea90a76011 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/naga/src/back/hlsl/writer.rs @@ -2209,6 +2209,25 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { writeln!(self.out, ");")?; } + crate::Statement::ImageAtomic { + image, + coordinate, + sample: _, + fun, + value, + } => { + write!(self.out, "{level}")?; + + let fun_str = fun.to_hlsl_suffix(); + write!(self.out, "Interlocked{fun_str}(")?; + self.write_expr(module, image, func_ctx)?; + write!(self.out, "[")?; + self.write_expr(module, coordinate, func_ctx)?; + write!(self.out, "],")?; + + self.write_expr(module, value, func_ctx)?; + writeln!(self.out, ");")?; + } Statement::WorkGroupUniformLoad { pointer, result } => { self.write_barrier(crate::Barrier::WORK_GROUP, level)?; write!(self.out, "{level}")?; diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index 99ca6c3e88..4aa440dc1a 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -1191,6 +1191,27 @@ impl Writer { Ok(()) } + fn put_image_atomic( + &mut self, + level: back::Level, + image: Handle, + address: &TexelAddress, + fun: crate::AtomicFunction, + value: Handle, + context: &StatementContext, + ) -> BackendResult { + write!(self.out, "{level}")?; + self.put_expression(image, &context.expression, false)?; + write!(self.out, ".atomic_{}(", fun.to_msl_64_bit()?)?; + // coordinates in IR are int, but Metal expects uint + self.put_cast_to_uint_scalar_or_vector(address.coordinate, &context.expression)?; + write!(self.out, ", ")?; + self.put_expression(value, &context.expression, true)?; + writeln!(self.out, ");")?; + + Ok(()) + } + fn put_image_store( &mut self, level: back::Level, @@ -3207,6 +3228,21 @@ impl Writer { // done writeln!(self.out, ";")?; } + crate::Statement::ImageAtomic { + image, + coordinate, + sample, + fun, + value, + } => { + let address = TexelAddress { + coordinate, + array_index: None, + sample: Some(sample), + level: None, + }; + self.put_image_atomic(level, image, &address, fun, value, context)? + } crate::Statement::WorkGroupUniformLoad { pointer, result } => { self.write_barrier(crate::Barrier::WORK_GROUP, level)?; diff --git a/naga/src/back/pipeline_constants.rs b/naga/src/back/pipeline_constants.rs index 5f82862f72..55dc935696 100644 --- a/naga/src/back/pipeline_constants.rs +++ b/naga/src/back/pipeline_constants.rs @@ -635,6 +635,18 @@ fn adjust_stmt(new_pos: &HandleVec>, stmt: &mut S | crate::AtomicFunction::Exchange { compare: None } => {} } } + Statement::ImageAtomic { + ref mut image, + ref mut coordinate, + ref mut sample, + fun: _, + ref mut value, + } => { + adjust(image); + adjust(coordinate); + adjust(sample); + adjust(value); + } Statement::WorkGroupUniformLoad { ref mut pointer, ref mut result, diff --git a/naga/src/back/spv/block.rs b/naga/src/back/spv/block.rs index f0c3bfa848..6f57b287e4 100644 --- a/naga/src/back/spv/block.rs +++ b/naga/src/back/spv/block.rs @@ -2599,6 +2599,15 @@ impl<'w> BlockContext<'w> { block.body.push(instruction); } + Statement::ImageAtomic { + image, + coordinate, + sample, + fun, + value, + } => { + self.write_image_atomic(image, coordinate, sample, fun, value, &mut block)?; + } Statement::WorkGroupUniformLoad { pointer, result } => { self.writer .write_barrier(crate::Barrier::WORK_GROUP, &mut block); diff --git a/naga/src/back/spv/image.rs b/naga/src/back/spv/image.rs index 769971d136..9f27039449 100644 --- a/naga/src/back/spv/image.rs +++ b/naga/src/back/spv/image.rs @@ -1188,4 +1188,72 @@ impl<'w> BlockContext<'w> { Ok(()) } + + pub(super) fn write_image_atomic( + &mut self, + image: Handle, + coordinate: Handle, + sample: Handle, + fun: crate::AtomicFunction, + value: Handle, + block: &mut Block, + ) -> Result<(), Error> { + let image_id = match self.ir_function.originating_global(image) { + Some(handle) => self.writer.global_variables[handle].var_id, + _ => return Err(Error::Validation("Unexpected image type")), + }; + let crate::TypeInner::Image { class, .. } = + *self.fun_info[image].ty.inner_with(&self.ir_module.types) + else { + return Err(Error::Validation("Invalid image type")); + }; + let crate::ImageClass::Storage { format, .. } = class else { + return Err(Error::Validation("Invalid image class")); + }; + let scalar = format.into(); + let pointer_type_id = self.get_type_id(LookupType::Local(LocalType::Value { + vector_size: None, + scalar, + pointer_space: Some(spirv::StorageClass::Image), + })); + if scalar.width == 8 { + self.writer + .require_any("64 bit image atomics", &[spirv::Capability::Int64Atomics])?; + } + let pointer_id = self.gen_id(); + let coordinates = self.write_image_coordinates(coordinate, None, block)?; + let sample_id = self.cached[sample]; + block.body.push(Instruction::image_texel_pointer( + pointer_type_id, + pointer_id, + image_id, + coordinates.value_id, + sample_id, + )); + + let op = match fun { + crate::AtomicFunction::Max => spirv::Op::AtomicUMax, + crate::AtomicFunction::Min => spirv::Op::AtomicUMin, + _ => return Err(Error::Validation("Invalid image atomic operation")), + }; + let result_type_id = self.get_expression_type_id(&self.fun_info[value].ty); + let id = self.gen_id(); + let space = crate::AddressSpace::Handle; + let (semantics, scope) = space.to_spirv_semantics_and_scope(); + let scope_constant_id = self.get_scope_constant(scope as u32); + let semantics_id = self.get_index_constant(semantics.bits()); + let value_id = self.cached[value]; + + block.body.push(Instruction::image_atomic( + op, + result_type_id, + id, + pointer_id, + scope_constant_id, + semantics_id, + value_id, + )); + + Ok(()) + } } diff --git a/naga/src/back/spv/instructions.rs b/naga/src/back/spv/instructions.rs index 61a38e86e5..38aed8c351 100644 --- a/naga/src/back/spv/instructions.rs +++ b/naga/src/back/spv/instructions.rs @@ -702,6 +702,41 @@ impl super::Instruction { instruction } + pub(super) fn image_texel_pointer( + result_type_id: Word, + id: Word, + image: Word, + coordinates: Word, + sample: Word, + ) -> Self { + let mut instruction = Self::new(Op::ImageTexelPointer); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(image); + instruction.add_operand(coordinates); + instruction.add_operand(sample); + instruction + } + + pub(super) fn image_atomic( + op: Op, + result_type_id: Word, + id: Word, + pointer: Word, + scope_id: Word, + semantics_id: Word, + value: Word, + ) -> Self { + let mut instruction = Self::new(op); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(pointer); + instruction.add_operand(scope_id); + instruction.add_operand(semantics_id); + instruction.add_operand(value); + instruction + } + pub(super) fn image_query(op: Op, result_type_id: Word, id: Word, image: Word) -> Self { let mut instruction = Self::new(op); instruction.set_type(result_type_id); diff --git a/naga/src/back/wgsl/writer.rs b/naga/src/back/wgsl/writer.rs index eb842da497..55d8dac8d6 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -773,6 +773,24 @@ impl Writer { self.write_expr(module, value, func_ctx)?; writeln!(self.out, ");")? } + Statement::ImageAtomic { + image, + coordinate, + sample: _, + ref fun, + value, + } => { + write!(self.out, "{level}")?; + let fun_str = fun.to_wgsl(); + write!(self.out, "imageAtomic{fun_str}(")?; + self.write_expr(module, image, func_ctx)?; + write!(self.out, ", ")?; + self.write_expr(module, coordinate, func_ctx)?; + // We do not write sample because it is unsupported + write!(self.out, ", ")?; + self.write_expr(module, value, func_ctx)?; + writeln!(self.out, ");")?; + } Statement::WorkGroupUniformLoad { pointer, result } => { write!(self.out, "{level}")?; // TODO: Obey named expressions here. diff --git a/naga/src/compact/statements.rs b/naga/src/compact/statements.rs index 759dcc2eda..07ab70d456 100644 --- a/naga/src/compact/statements.rs +++ b/naga/src/compact/statements.rs @@ -79,6 +79,18 @@ impl FunctionTracer<'_> { self.expressions_used.insert(result); } } + St::ImageAtomic { + image, + coordinate, + sample, + fun: _, + value, + } => { + self.expressions_used.insert(image); + self.expressions_used.insert(coordinate); + self.expressions_used.insert(sample); + self.expressions_used.insert(value); + } St::WorkGroupUniformLoad { pointer, result } => { self.expressions_used.insert(pointer); self.expressions_used.insert(result); @@ -261,6 +273,18 @@ impl FunctionMap { adjust(result); } } + St::ImageAtomic { + ref mut image, + ref mut coordinate, + ref mut sample, + fun: _, + ref mut value, + } => { + adjust(image); + adjust(coordinate); + adjust(sample); + adjust(value); + } St::WorkGroupUniformLoad { ref mut pointer, ref mut result, diff --git a/naga/src/front/spv/mod.rs b/naga/src/front/spv/mod.rs index 5ad063a6b6..2e5c918d30 100644 --- a/naga/src/front/spv/mod.rs +++ b/naga/src/front/spv/mod.rs @@ -4377,6 +4377,7 @@ impl> Frontend { | S::Store { .. } | S::ImageStore { .. } | S::Atomic { .. } + | S::ImageAtomic { .. } | S::RayQuery { .. } | S::SubgroupBallot { .. } | S::SubgroupCollectiveOperation { .. } diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 78e81350b4..b786cfc98e 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -2349,6 +2349,42 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { ); return Ok(Some(result)); } + "imageAtomicMin" | "imageAtomicMax" => { + let mut args = ctx.prepare_args(arguments, 3, span); + + let image = args.next()?; + let image = self.expression(image, ctx)?; + + let coordinate = self.expression(args.next()?, ctx)?; + + let ty = ctx + .ensure_type_exists(crate::TypeInner::Scalar(crate::Scalar::I32)); + // We fib in a zero value for sample because it is not supported + let sample = + ctx.append_expression(crate::Expression::ZeroValue(ty), span)?; + + let value = self.expression(args.next()?, ctx)?; + + args.finish()?; + + let rctx = ctx.runtime_expression_ctx(span)?; + rctx.block + .extend(rctx.emitter.finish(&rctx.function.expressions)); + rctx.emitter.start(&rctx.function.expressions); + let stmt = crate::Statement::ImageAtomic { + image, + coordinate, + sample, + fun: match function.name { + "imageAtomicMin" => crate::AtomicFunction::Min, + "imageAtomicMax" => crate::AtomicFunction::Max, + _ => unreachable!(), + }, + value, + }; + rctx.block.push(stmt, span); + return Ok(None); + } "storageBarrier" => { ctx.prepare_args(arguments, 0, span).finish()?; diff --git a/naga/src/front/wgsl/parse/mod.rs b/naga/src/front/wgsl/parse/mod.rs index 3b1d60620b..3ff58b09e2 100644 --- a/naga/src/front/wgsl/parse/mod.rs +++ b/naga/src/front/wgsl/parse/mod.rs @@ -1502,6 +1502,10 @@ impl Parser { kind: Float | Sint | Uint, width: 4, } => Ok(()), + Scalar { + kind: Uint, + width: 8, + } => Ok(()), _ => Err(Error::BadTextureSampleType { span, scalar }), } } diff --git a/naga/src/lib.rs b/naga/src/lib.rs index 5f681b3399..c39982574f 100644 --- a/naga/src/lib.rs +++ b/naga/src/lib.rs @@ -1989,6 +1989,54 @@ pub enum Statement { /// [`SHADER_INT64_ATOMIC_ALL_OPS`]: crate::valid::Capabilities::SHADER_INT64_ATOMIC_ALL_OPS result: Option>, }, + /// Performs an atomic operation on a texel value of an image. + /// + /// Doing atomics on images with mipmaps is not supported, so there is no + /// `level` operand. + /// + /// This statement is a barrier for any operations on the corresponding + /// [`Expression::GlobalVariable`] for this image. + ImageAtomic { + /// The image to perform an atomic operation on. This must have type + /// [`Image`]. (This will necessarily be a [`GlobalVariable`] or + /// [`FunctionArgument`] expression, since no other expressions are + /// allowed to have that type.) + /// + /// [`Image`]: TypeInner::Image + /// [`GlobalVariable`]: Expression::GlobalVariable + /// [`FunctionArgument`]: Expression::FunctionArgument + image: Handle, + + /// The coordinate of the texel we wish to load. This must be a scalar + /// for [`D1`] images, a [`Bi`] vector for [`D2`] images, and a [`Tri`] + /// vector for [`D3`] images. (sample indices are supplied separately.) + /// Its component type must be [`Sint`]. + /// + /// If this image is arrayed, [`D1`] images require a [`Bi`] vector and + /// [`D2`] images require a [`Tri`] vector. + /// + /// Explicit level-of-detail values are unsupported. + /// + /// [`D1`]: ImageDimension::D1 + /// [`D2`]: ImageDimension::D2 + /// [`D3`]: ImageDimension::D3 + /// [`Bi`]: VectorSize::Bi + /// [`Tri`]: VectorSize::Tri + /// [`Sint`]: ScalarKind::Sint + coordinate: Handle, + + /// A sample index, for multisampled [`Sampled`] and [`Depth`] images. + /// + /// [`Sampled`]: ImageClass::Sampled + /// [`Depth`]: ImageClass::Depth + sample: Handle, + + /// The kind of atomic operation to perform on the texel. + fun: AtomicFunction, + + // The value with which to perform the atomic operation. + value: Handle, + }, /// Load uniformly from a uniform pointer in the workgroup address space. /// /// Corresponds to the [`workgroupUniformLoad`](https://www.w3.org/TR/WGSL/#workgroupUniformLoad-builtin) diff --git a/naga/src/proc/terminator.rs b/naga/src/proc/terminator.rs index 5edf55cb73..19c37294ec 100644 --- a/naga/src/proc/terminator.rs +++ b/naga/src/proc/terminator.rs @@ -36,6 +36,7 @@ pub fn ensure_block_returns(block: &mut crate::Block) { | S::Call { .. } | S::RayQuery { .. } | S::Atomic { .. } + | S::ImageAtomic { .. } | S::WorkGroupUniformLoad { .. } | S::SubgroupBallot { .. } | S::SubgroupCollectiveOperation { .. } diff --git a/naga/src/valid/analyzer.rs b/naga/src/valid/analyzer.rs index af95fd098f..b2d0ead8cf 100644 --- a/naga/src/valid/analyzer.rs +++ b/naga/src/valid/analyzer.rs @@ -1022,6 +1022,19 @@ impl FunctionInfo { } FunctionUniformity::new() } + S::ImageAtomic { + image, + coordinate, + sample, + fun: _, + value, + } => { + let _ = self.add_ref_impl(image, GlobalUse::WRITE); + let _ = self.add_ref(coordinate); + let _ = self.add_ref(sample); + let _ = self.add_ref(value); + FunctionUniformity::new() + } S::RayQuery { query, ref fun } => { let _ = self.add_ref(query); if let crate::RayQueryFunction::Initialize { diff --git a/naga/src/valid/expression.rs b/naga/src/valid/expression.rs index 0b0d115c57..f61f9f205c 100644 --- a/naga/src/valid/expression.rs +++ b/naga/src/valid/expression.rs @@ -67,6 +67,10 @@ pub enum ExpressionError { ExpectedSamplerType(Handle), #[error("Unable to operate on image class {0:?}")] InvalidImageClass(crate::ImageClass), + #[error("Image atomics are not supported for storage format {0:?}")] + InvalidImageFormat(crate::StorageFormat), + #[error("Image atomics require read/write storage access, {0:?} is insufficient")] + InvalidImageStorageAccess(crate::StorageAccess), #[error("Derivatives can only be taken from scalar and vector floats")] InvalidDerivative, #[error("Image array index parameter is misplaced")] diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index c695d65144..2cf87e460c 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -146,6 +146,10 @@ pub enum FunctionError { }, #[error("Image store parameters are invalid")] InvalidImageStore(#[source] ExpressionError), + #[error("Image atomic parameters are invalid")] + InvalidImageAtomic(#[source] ExpressionError), + #[error("Image atomic value is invalid")] + InvalidAtomicValue(Handle), #[error("Call to {function:?} is invalid")] InvalidCall { function: Handle, @@ -1136,6 +1140,113 @@ impl super::Validator { } => { self.validate_atomic(pointer, fun, value, result, span, context)?; } + S::ImageAtomic { + image, + coordinate, + sample: _, + fun: _, + value, + } => { + // Note: this code uses a lot of `FunctionError::InvalidImageAtomic`, + // and could probably be refactored. + let var = match *context.get_expression(image) { + crate::Expression::GlobalVariable(var_handle) => { + &context.global_vars[var_handle] + } + // We're looking at a binding index situation, so punch through the index and look at the global behind it. + crate::Expression::Access { base, .. } + | crate::Expression::AccessIndex { base, .. } => { + match *context.get_expression(base) { + crate::Expression::GlobalVariable(var_handle) => { + &context.global_vars[var_handle] + } + _ => { + return Err(FunctionError::InvalidImageAtomic( + ExpressionError::ExpectedGlobalVariable, + ) + .with_span_handle(image, context.expressions)) + } + } + } + _ => { + return Err(FunctionError::InvalidImageAtomic( + ExpressionError::ExpectedGlobalVariable, + ) + .with_span_handle(image, context.expressions)) + } + }; + + // Punch through a binding array to get the underlying type + let global_ty = match context.types[var.ty].inner { + Ti::BindingArray { base, .. } => &context.types[base].inner, + ref inner => inner, + }; + + let value_ty = match *global_ty { + Ti::Image { + class, + arrayed: _, + dim, + } => { + match context + .resolve_type(coordinate, &self.valid_expression_set)? + .image_storage_coordinates() + { + Some(coord_dim) if coord_dim == dim => {} + _ => { + return Err(FunctionError::InvalidImageAtomic( + ExpressionError::InvalidImageCoordinateType( + dim, coordinate, + ), + ) + .with_span_handle(coordinate, context.expressions)); + } + }; + + match class { + crate::ImageClass::Storage { format, access } => { + if access + != crate::StorageAccess::LOAD | crate::StorageAccess::STORE + { + return Err(FunctionError::InvalidImageAtomic( + ExpressionError::InvalidImageStorageAccess(access), + ) + .with_span_handle(image, context.expressions)); + } + match format { + crate::StorageFormat::R64Uint => {} + _ => { + return Err(FunctionError::InvalidImageAtomic( + ExpressionError::InvalidImageFormat(format), + ) + .with_span_handle(image, context.expressions)); + } + } + crate::TypeInner::Scalar(format.into()) + } + _ => { + return Err(FunctionError::InvalidImageAtomic( + ExpressionError::InvalidImageClass(class), + ) + .with_span_handle(image, context.expressions)); + } + } + } + _ => { + return Err(FunctionError::InvalidImageAtomic( + ExpressionError::ExpectedImageType(var.ty), + ) + .with_span() + .with_handle(var.ty, context.types) + .with_handle(image, context.expressions)) + } + }; + + if *context.resolve_type(value, &self.valid_expression_set)? != value_ty { + return Err(FunctionError::InvalidAtomicValue(value) + .with_span_handle(value, context.expressions)); + } + } S::WorkGroupUniformLoad { pointer, result } => { stages &= super::ShaderStages::COMPUTE; let pointer_inner = diff --git a/naga/src/valid/handles.rs b/naga/src/valid/handles.rs index f8be76d026..6420d36057 100644 --- a/naga/src/valid/handles.rs +++ b/naga/src/valid/handles.rs @@ -536,6 +536,19 @@ impl super::Validator { } Ok(()) } + crate::Statement::ImageAtomic { + image, + coordinate, + sample, + fun: _, + value, + } => { + validate_expr(image)?; + validate_expr(coordinate)?; + validate_expr(sample)?; + validate_expr(value)?; + Ok(()) + } crate::Statement::WorkGroupUniformLoad { pointer, result } => { validate_expr(pointer)?; validate_expr(result)?; From 6cc4b5a194ed6b7f3997e99394951d7cad776c3e Mon Sep 17 00:00:00 2001 From: atlas dostal Date: Wed, 2 Oct 2024 17:17:09 -0400 Subject: [PATCH 4/9] naga snapshot tests --- naga/tests/in/atomicTexture-int64.param.ron | 23 +++++++++ naga/tests/in/atomicTexture-int64.wgsl | 12 +++++ naga/tests/out/hlsl/atomicTexture-int64.hlsl | 21 ++++++++ naga/tests/out/hlsl/atomicTexture-int64.ron | 12 +++++ naga/tests/out/msl/atomicTexture-int64.msl | 18 +++++++ naga/tests/out/spv/atomicTexture-int64.spvasm | 50 +++++++++++++++++++ naga/tests/out/wgsl/atomicTexture-int64.wgsl | 10 ++++ naga/tests/snapshots.rs | 4 ++ 8 files changed, 150 insertions(+) create mode 100644 naga/tests/in/atomicTexture-int64.param.ron create mode 100644 naga/tests/in/atomicTexture-int64.wgsl create mode 100644 naga/tests/out/hlsl/atomicTexture-int64.hlsl create mode 100644 naga/tests/out/hlsl/atomicTexture-int64.ron create mode 100644 naga/tests/out/msl/atomicTexture-int64.msl create mode 100644 naga/tests/out/spv/atomicTexture-int64.spvasm create mode 100644 naga/tests/out/wgsl/atomicTexture-int64.wgsl diff --git a/naga/tests/in/atomicTexture-int64.param.ron b/naga/tests/in/atomicTexture-int64.param.ron new file mode 100644 index 0000000000..ed1025fe5c --- /dev/null +++ b/naga/tests/in/atomicTexture-int64.param.ron @@ -0,0 +1,23 @@ +( + god_mode: true, + spv: ( + version: (1, 0), + capabilities: [ Int64, Int64ImageEXT, Int64Atomics ], + ), + hlsl: ( + shader_model: V6_6, + binding_map: {}, + fake_missing_bindings: true, + special_constants_binding: Some((space: 1, register: 0)), + push_constants_target: Some((space: 0, register: 0)), + zero_initialize_workgroup_memory: true, + ), + msl: ( + lang_version: (3, 1), + per_entry_point_map: {}, + inline_samplers: [], + spirv_cross_compatibility: false, + fake_missing_bindings: true, + zero_initialize_workgroup_memory: true, + ), +) diff --git a/naga/tests/in/atomicTexture-int64.wgsl b/naga/tests/in/atomicTexture-int64.wgsl new file mode 100644 index 0000000000..581cfac57d --- /dev/null +++ b/naga/tests/in/atomicTexture-int64.wgsl @@ -0,0 +1,12 @@ +@group(0) @binding(0) +var image: texture_storage_2d; + +@compute +@workgroup_size(2) +fn cs_main(@builtin(local_invocation_id) id: vec3) { + imageAtomicMax(image, vec2(0, 0), 1lu); + + workgroupBarrier(); + + imageAtomicMin(image, vec2(0, 0), 1lu); +} diff --git a/naga/tests/out/hlsl/atomicTexture-int64.hlsl b/naga/tests/out/hlsl/atomicTexture-int64.hlsl new file mode 100644 index 0000000000..beeec0451e --- /dev/null +++ b/naga/tests/out/hlsl/atomicTexture-int64.hlsl @@ -0,0 +1,21 @@ +struct NagaConstants { + int first_vertex; + int first_instance; + uint other; +}; +ConstantBuffer _NagaConstants: register(b0, space1); + +RWTexture2D image : register(u0); + +int ZeroValueint() { + return (int)0; +} + +[numthreads(2, 1, 1)] +void cs_main(uint3 id : SV_GroupThreadID) +{ + InterlockedMax(image[int2(0, 0)],1uL); + GroupMemoryBarrierWithGroupSync(); + InterlockedMin(image[int2(0, 0)],1uL); + return; +} diff --git a/naga/tests/out/hlsl/atomicTexture-int64.ron b/naga/tests/out/hlsl/atomicTexture-int64.ron new file mode 100644 index 0000000000..67a9035512 --- /dev/null +++ b/naga/tests/out/hlsl/atomicTexture-int64.ron @@ -0,0 +1,12 @@ +( + vertex:[ + ], + fragment:[ + ], + compute:[ + ( + entry_point:"cs_main", + target_profile:"cs_6_6", + ), + ], +) diff --git a/naga/tests/out/msl/atomicTexture-int64.msl b/naga/tests/out/msl/atomicTexture-int64.msl new file mode 100644 index 0000000000..c00d8b7654 --- /dev/null +++ b/naga/tests/out/msl/atomicTexture-int64.msl @@ -0,0 +1,18 @@ +// language: metal3.1 +#include +#include + +using metal::uint; + + +struct cs_mainInput { +}; +kernel void cs_main( + metal::uint3 id [[thread_position_in_threadgroup]] +, metal::texture2d image [[user(fake0)]] +) { + image.atomic_max(metal::uint2(metal::int2(0, 0)), 1uL); + metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); + image.atomic_min(metal::uint2(metal::int2(0, 0)), 1uL); + return; +} diff --git a/naga/tests/out/spv/atomicTexture-int64.spvasm b/naga/tests/out/spv/atomicTexture-int64.spvasm new file mode 100644 index 0000000000..745414864d --- /dev/null +++ b/naga/tests/out/spv/atomicTexture-int64.spvasm @@ -0,0 +1,50 @@ +; SPIR-V +; Version: 1.0 +; Generator: rspirv +; Bound: 32 +OpCapability Shader +OpCapability Int64ImageEXT +OpCapability Int64 +OpCapability Int64Atomics +OpExtension "SPV_EXT_shader_image_int64" +%1 = OpExtInstImport "GLSL.std.450" +OpMemoryModel Logical GLSL450 +OpEntryPoint GLCompute %15 "cs_main" %12 +OpExecutionMode %15 LocalSize 2 1 1 +OpDecorate %9 DescriptorSet 0 +OpDecorate %9 Binding 0 +OpDecorate %12 BuiltIn LocalInvocationId +%2 = OpTypeVoid +%4 = OpTypeInt 64 0 +%3 = OpTypeImage %4 2D 0 0 0 2 R64ui +%6 = OpTypeInt 32 0 +%5 = OpTypeVector %6 3 +%7 = OpTypeInt 32 1 +%8 = OpTypeVector %7 2 +%10 = OpTypePointer UniformConstant %3 +%9 = OpVariable %10 UniformConstant +%13 = OpTypePointer Input %5 +%12 = OpVariable %13 Input +%16 = OpTypeFunction %2 +%18 = OpConstant %7 0 +%19 = OpConstantComposite %8 %18 %18 +%20 = OpConstantNull %7 +%21 = OpConstant %4 1 +%23 = OpTypePointer Image %4 +%26 = OpConstant %7 4 +%27 = OpConstant %6 0 +%28 = OpConstant %6 2 +%29 = OpConstant %6 264 +%15 = OpFunction %2 None %16 +%11 = OpLabel +%14 = OpLoad %5 %12 +%17 = OpLoad %3 %9 +OpBranch %22 +%22 = OpLabel +%24 = OpImageTexelPointer %23 %9 %19 %20 +%25 = OpAtomicUMax %4 %24 %26 %27 %21 +OpControlBarrier %28 %28 %29 +%30 = OpImageTexelPointer %23 %9 %19 %20 +%31 = OpAtomicUMin %4 %30 %26 %27 %21 +OpReturn +OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/atomicTexture-int64.wgsl b/naga/tests/out/wgsl/atomicTexture-int64.wgsl new file mode 100644 index 0000000000..4ac35b510d --- /dev/null +++ b/naga/tests/out/wgsl/atomicTexture-int64.wgsl @@ -0,0 +1,10 @@ +@group(0) @binding(0) +var image: texture_storage_2d; + +@compute @workgroup_size(2, 1, 1) +fn cs_main(@builtin(local_invocation_id) id: vec3) { + imageAtomicMax(image, vec2(0i, 0i), 1lu); + workgroupBarrier(); + imageAtomicMin(image, vec2(0i, 0i), 1lu); + return; +} diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index 596e4cea14..544bfb68d8 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -786,6 +786,10 @@ fn convert_wgsl() { "atomicOps-int64-min-max", Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL, ), + ( + "atomicTexture-int64", + Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL, + ), ( "atomicCompareExchange-int64", Targets::SPIRV | Targets::WGSL, From deb89bd64f3d6d33363712dbffd068456952badb Mon Sep 17 00:00:00 2001 From: atlas dostal Date: Wed, 2 Oct 2024 17:17:30 -0400 Subject: [PATCH 5/9] wgpu runtime tests --- tests/tests/image_atomics/image_atomics.wgsl | 14 +++ tests/tests/image_atomics/mod.rs | 106 +++++++++++++++++++ tests/tests/root.rs | 1 + 3 files changed, 121 insertions(+) create mode 100644 tests/tests/image_atomics/image_atomics.wgsl create mode 100644 tests/tests/image_atomics/mod.rs diff --git a/tests/tests/image_atomics/image_atomics.wgsl b/tests/tests/image_atomics/image_atomics.wgsl new file mode 100644 index 0000000000..1c43b6d180 --- /dev/null +++ b/tests/tests/image_atomics/image_atomics.wgsl @@ -0,0 +1,14 @@ +@group(0) @binding(0) +var image: texture_storage_2d; + +@compute +@workgroup_size(4, 4, 4) +fn cs_main(@builtin(local_invocation_id) id: vec3) { + let l = u64(textureLoad(image, id.xy).r); + + imageAtomicMin(image, id.xy, u64(id.z) + l); + + workgroupBarrier(); + + imageAtomicMax(image, id.xy, u64(id.z) + l); +} \ No newline at end of file diff --git a/tests/tests/image_atomics/mod.rs b/tests/tests/image_atomics/mod.rs new file mode 100644 index 0000000000..69413844a6 --- /dev/null +++ b/tests/tests/image_atomics/mod.rs @@ -0,0 +1,106 @@ +//! Tests for image atomics. + +use wgpu_test::{gpu_test, GpuTestConfiguration, TestParameters}; + +#[gpu_test] +static IMAGE_ATOMICS: GpuTestConfiguration = GpuTestConfiguration::new() + .parameters( + TestParameters::default() + .limits(wgt::Limits { + max_storage_textures_per_shader_stage: 1, + max_compute_invocations_per_workgroup: 64, + max_compute_workgroup_size_x: 4, + max_compute_workgroup_size_y: 4, + max_compute_workgroup_size_z: 4, + max_compute_workgroups_per_dimension: 64, + ..wgt::Limits::downlevel_webgl2_defaults() + }) + .features( + wgpu::Features::TEXTURE_INT64_ATOMIC + | wgpu::Features::SHADER_INT64 + | wgpu::Features::SHADER_INT64_ATOMIC_ALL_OPS + | wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES, + ), + ) + .run_sync(|ctx| { + let size = wgpu::Extent3d { + width: 256, + height: 256, + depth_or_array_layers: 1, + }; + let bind_group_layout_entries = vec![wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::StorageTexture { + access: wgpu::StorageTextureAccess::ReadWrite, + format: wgpu::TextureFormat::R64Uint, + view_dimension: wgpu::TextureViewDimension::D2, + }, + count: None, + }]; + + let bind_group_layout = + ctx.device + .create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: None, + entries: &bind_group_layout_entries, + }); + + let pipeline_layout = ctx + .device + .create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: None, + bind_group_layouts: &[&bind_group_layout], + push_constant_ranges: &[], + }); + let shader = ctx + .device + .create_shader_module(wgpu::include_wgsl!("image_atomics.wgsl")); + let pipeline = ctx + .device + .create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("image atomics pipeline"), + layout: Some(&pipeline_layout), + module: &shader, + entry_point: Some("cs_main"), + compilation_options: wgpu::PipelineCompilationOptions::default(), + cache: None, + }); + + let tex = ctx.device.create_texture(&wgpu::TextureDescriptor { + label: None, + dimension: wgpu::TextureDimension::D2, + size, + format: wgpu::TextureFormat::R64Uint, + usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::SHADER_ATOMIC, + mip_level_count: 1, + sample_count: 1, + view_formats: &[], + }); + let view = tex.create_view(&wgpu::TextureViewDescriptor { + format: Some(wgpu::TextureFormat::R64Uint), + aspect: wgpu::TextureAspect::All, + ..wgpu::TextureViewDescriptor::default() + }); + let bind_group = ctx.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &pipeline.get_bind_group_layout(0), + entries: &[wgpu::BindGroupEntry { + binding: 0, + resource: wgpu::BindingResource::TextureView(&view), + }], + }); + + let mut encoder = ctx + .device + .create_command_encoder(&wgpu::CommandEncoderDescriptor::default()); + let mut rpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: None, + timestamp_writes: None, + }); + rpass.set_pipeline(&pipeline); + rpass.set_bind_group(0, Some(&bind_group), &[]); + rpass.dispatch_workgroups(1, 1, 1); + drop(rpass); + ctx.queue.submit(Some(encoder.finish())); + }); diff --git a/tests/tests/root.rs b/tests/tests/root.rs index 3bb8e14a90..8c2c9e79e4 100644 --- a/tests/tests/root.rs +++ b/tests/tests/root.rs @@ -22,6 +22,7 @@ mod device; mod encoder; mod external_texture; mod float32_filterable; +mod image_atomics; mod instance; mod life_cycle; mod mem_leaks; From 2a813b8dfa4e8ab9d564de2e0bf1e8ff7b59a765 Mon Sep 17 00:00:00 2001 From: atlas dostal Date: Wed, 2 Oct 2024 17:17:39 -0400 Subject: [PATCH 6/9] changelog --- CHANGELOG.md | 2 ++ 1 file changed, 2 insertions(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index 9d564bb082..d13d306dca 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -77,6 +77,8 @@ By @bradwerth [#6216](https://github.com/gfx-rs/wgpu/pull/6216). ### New Features +64 bit image atomic support in shaders. By @atlv24 in [#5537](https://github.com/gfx-rs/wgpu/pull/5537) + #### Naga - Support constant evaluation for `firstLeadingBit` and `firstTrailingBit` numeric built-ins in WGSL. Front-ends that translate to these built-ins also benefit from constant evaluation. By @ErichDonGubler in [#5101](https://github.com/gfx-rs/wgpu/pull/5101). From 136607b49b5f193974481f91ae372bda632ad86d Mon Sep 17 00:00:00 2001 From: atlas dostal Date: Fri, 4 Oct 2024 14:16:52 -0400 Subject: [PATCH 7/9] try to improve tests --- tests/tests/image_atomics/image_atomics.wgsl | 9 ++- tests/tests/image_atomics/mod.rs | 62 ++++++++++++++++++-- wgpu-hal/src/dx12/adapter.rs | 6 ++ 3 files changed, 68 insertions(+), 9 deletions(-) diff --git a/tests/tests/image_atomics/image_atomics.wgsl b/tests/tests/image_atomics/image_atomics.wgsl index 1c43b6d180..027e6957bf 100644 --- a/tests/tests/image_atomics/image_atomics.wgsl +++ b/tests/tests/image_atomics/image_atomics.wgsl @@ -2,13 +2,12 @@ var image: texture_storage_2d; @compute -@workgroup_size(4, 4, 4) +@workgroup_size(4, 2, 4) fn cs_main(@builtin(local_invocation_id) id: vec3) { - let l = u64(textureLoad(image, id.xy).r); - - imageAtomicMin(image, id.xy, u64(id.z) + l); + let data = u64((id.x << 16) | (id.y << 8) | id.z); + imageAtomicMax(image, id.xy, (u64(100 - id.z) << 32) | data); workgroupBarrier(); - imageAtomicMax(image, id.xy, u64(id.z) + l); + imageAtomicMin(image, id.xy, (u64(10 - id.z) << 32) | data); } \ No newline at end of file diff --git a/tests/tests/image_atomics/mod.rs b/tests/tests/image_atomics/mod.rs index 69413844a6..5a0ba4a11d 100644 --- a/tests/tests/image_atomics/mod.rs +++ b/tests/tests/image_atomics/mod.rs @@ -12,7 +12,7 @@ static IMAGE_ATOMICS: GpuTestConfiguration = GpuTestConfiguration::new() max_compute_workgroup_size_x: 4, max_compute_workgroup_size_y: 4, max_compute_workgroup_size_z: 4, - max_compute_workgroups_per_dimension: 64, + max_compute_workgroups_per_dimension: 4, ..wgt::Limits::downlevel_webgl2_defaults() }) .features( @@ -22,7 +22,7 @@ static IMAGE_ATOMICS: GpuTestConfiguration = GpuTestConfiguration::new() | wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES, ), ) - .run_sync(|ctx| { + .run_async(|ctx| async move { let size = wgpu::Extent3d { width: 256, height: 256, @@ -72,7 +72,9 @@ static IMAGE_ATOMICS: GpuTestConfiguration = GpuTestConfiguration::new() dimension: wgpu::TextureDimension::D2, size, format: wgpu::TextureFormat::R64Uint, - usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::SHADER_ATOMIC, + usage: wgpu::TextureUsages::STORAGE_BINDING + | wgpu::TextureUsages::SHADER_ATOMIC + | wgpu::TextureUsages::COPY_SRC, mip_level_count: 1, sample_count: 1, view_formats: &[], @@ -100,7 +102,59 @@ static IMAGE_ATOMICS: GpuTestConfiguration = GpuTestConfiguration::new() }); rpass.set_pipeline(&pipeline); rpass.set_bind_group(0, Some(&bind_group), &[]); - rpass.dispatch_workgroups(1, 1, 1); + rpass.dispatch_workgroups(4, 2, 4); drop(rpass); ctx.queue.submit(Some(encoder.finish())); + + let read_buffer = ctx.device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size: (size.height * size.width * size.depth_or_array_layers) as u64, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + + let mut encoder = ctx + .device + .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None }); + + encoder.copy_texture_to_buffer( + wgpu::ImageCopyTexture { + texture: &tex, + mip_level: 0, + origin: wgpu::Origin3d::ZERO, + aspect: wgpu::TextureAspect::All, + }, + wgpu::ImageCopyBuffer { + buffer: &read_buffer, + layout: wgpu::ImageDataLayout { + offset: 0, + bytes_per_row: Some(size.width * 8), + rows_per_image: Some(size.height), + }, + }, + size, + ); + + ctx.queue.submit(Some(encoder.finish())); + + let slice = read_buffer.slice(..); + slice.map_async(wgpu::MapMode::Read, |_| ()); + ctx.async_poll(wgpu::Maintain::wait()) + .await + .panic_on_timeout(); + let data: Vec = slice.get_mapped_range().to_vec(); + + assert_eq!(data.len(), 256 * 256 * 8); + for (i, long) in data.chunks(8).into_iter().enumerate() { + let x = (i as u32 % size.width) as u8; + let y = (i as u32 / size.width) as u8; + assert_eq!(long[0], 3); + assert_eq!(long[1], y); + assert_eq!(long[2], x); + assert_eq!(long[3], 0); + assert_eq!(long[4], 7); + assert_eq!(long[5], 0); + assert_eq!(long[6], 0); + assert_eq!(long[7], 0); + } }); diff --git a/wgpu-hal/src/dx12/adapter.rs b/wgpu-hal/src/dx12/adapter.rs index 4b510290d9..5902ff2825 100644 --- a/wgpu-hal/src/dx12/adapter.rs +++ b/wgpu-hal/src/dx12/adapter.rs @@ -682,6 +682,12 @@ impl crate::Adapter for super::Adapter { .Support2 .contains(Direct3D12::D3D12_FORMAT_SUPPORT2_UAV_TYPED_LOAD), ); + caps.set( + Tfc::SHADER_ATOMIC, + data_srv_uav + .Support2 + .contains(Direct3D12::D3D12_FORMAT_SUPPORT2_UAV_ATOMIC_UNSIGNED_MIN_OR_MAX), + ); // We load via UAV/SRV so use srv_uav_format let no_msaa_load = caps.contains(Tfc::SAMPLED) From 5b4ccdf5ca5ce5f9d8edd6f9b6e064b3482e7bfa Mon Sep 17 00:00:00 2001 From: atlas dostal Date: Fri, 4 Oct 2024 14:26:28 -0400 Subject: [PATCH 8/9] clippy --- tests/tests/image_atomics/mod.rs | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/tests/image_atomics/mod.rs b/tests/tests/image_atomics/mod.rs index 5a0ba4a11d..3ebddb8df3 100644 --- a/tests/tests/image_atomics/mod.rs +++ b/tests/tests/image_atomics/mod.rs @@ -145,7 +145,7 @@ static IMAGE_ATOMICS: GpuTestConfiguration = GpuTestConfiguration::new() let data: Vec = slice.get_mapped_range().to_vec(); assert_eq!(data.len(), 256 * 256 * 8); - for (i, long) in data.chunks(8).into_iter().enumerate() { + for (i, long) in data.chunks(8).enumerate() { let x = (i as u32 % size.width) as u8; let y = (i as u32 / size.width) as u8; assert_eq!(long[0], 3); From dc9cc86e3f006993ae3b1f07b2ac5b8471ba96db Mon Sep 17 00:00:00 2001 From: atlas dostal Date: Mon, 21 Oct 2024 12:30:25 -0400 Subject: [PATCH 9/9] Address feedback --- wgpu-hal/src/dx12/adapter.rs | 11 +++++------ wgpu/src/backend/webgpu.rs | 1 - .../backend/webgpu/webgpu_sys/gen_GpuTextureFormat.rs | 1 - 3 files changed, 5 insertions(+), 8 deletions(-) diff --git a/wgpu-hal/src/dx12/adapter.rs b/wgpu-hal/src/dx12/adapter.rs index 5902ff2825..77092f2c90 100644 --- a/wgpu-hal/src/dx12/adapter.rs +++ b/wgpu-hal/src/dx12/adapter.rs @@ -315,7 +315,11 @@ impl super::Adapter { | wgt::Features::SHADER_PRIMITIVE_INDEX | wgt::Features::RG11B10UFLOAT_RENDERABLE | wgt::Features::DUAL_SOURCE_BLENDING - | wgt::Features::TEXTURE_FORMAT_NV12; + | wgt::Features::TEXTURE_FORMAT_NV12 + // Ruint64 textures are always emulated on d3d12 + | wgt::Features::TEXTURE_INT64_ATOMIC + // float32-filterable should always be available on d3d12 + | wgt::Features::FLOAT32_FILTERABLE; //TODO: in order to expose this, we need to run a compute shader // that extract the necessary statistics out of the D3D12 result. @@ -402,11 +406,6 @@ impl super::Adapter { wgt::Features::SHADER_INT64_ATOMIC_ALL_OPS | wgt::Features::SHADER_INT64_ATOMIC_MIN_MAX, atomic_int64_on_typed_resource_supported, ); - // Ruint64 textures are always emulated on d3d12 - features.set(wgt::Features::TEXTURE_INT64_ATOMIC, true); - - // float32-filterable should always be available on d3d12 - features.set(wgt::Features::FLOAT32_FILTERABLE, true); // TODO: Determine if IPresentationManager is supported let presentation_timer = auxil::dxgi::time::PresentationTimer::new_dxgi(); diff --git a/wgpu/src/backend/webgpu.rs b/wgpu/src/backend/webgpu.rs index 68368803ce..e0cf006e6e 100644 --- a/wgpu/src/backend/webgpu.rs +++ b/wgpu/src/backend/webgpu.rs @@ -264,7 +264,6 @@ fn map_texture_format(texture_format: wgt::TextureFormat) -> webgpu_sys::GpuText TextureFormat::Rgb10a2Unorm => tf::Rgb10a2unorm, TextureFormat::Rg11b10Ufloat => tf::Rg11b10ufloat, // 64-bit formats - TextureFormat::R64Uint => tf::R64uint, TextureFormat::Rg32Uint => tf::Rg32uint, TextureFormat::Rg32Sint => tf::Rg32sint, TextureFormat::Rg32Float => tf::Rg32float, diff --git a/wgpu/src/backend/webgpu/webgpu_sys/gen_GpuTextureFormat.rs b/wgpu/src/backend/webgpu/webgpu_sys/gen_GpuTextureFormat.rs index fcd9063518..d554598559 100644 --- a/wgpu/src/backend/webgpu/webgpu_sys/gen_GpuTextureFormat.rs +++ b/wgpu/src/backend/webgpu/webgpu_sys/gen_GpuTextureFormat.rs @@ -60,7 +60,6 @@ pub enum GpuTextureFormat { Rgb10a2uint = "rgb10a2uint", Rgb10a2unorm = "rgb10a2unorm", Rg11b10ufloat = "rg11b10ufloat", - R64uint = "r64uint", Rg32uint = "rg32uint", Rg32sint = "rg32sint", Rg32float = "rg32float",