diff --git a/CHANGELOG.md b/CHANGELOG.md index 501ee4b713..1c1ee765c6 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -106,6 +106,7 @@ By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148] ### New Features Image atomic support in shaders. By @atlv24 in [#6706](https://github.com/gfx-rs/wgpu/pull/6706) +64 bit image atomic support in shaders. By @atlv24 in [#5537](https://github.com/gfx-rs/wgpu/pull/5537) #### Naga diff --git a/naga/src/back/glsl/features.rs b/naga/src/back/glsl/features.rs index 0065db2f58..3eb0c42439 100644 --- a/naga/src/back/glsl/features.rs +++ b/naga/src/back/glsl/features.rs @@ -400,6 +400,7 @@ impl Writer<'_, 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 fda7e9aa73..ec724325bb 100644 --- a/naga/src/back/glsl/mod.rs +++ b/naga/src/back/glsl/mod.rs @@ -4941,6 +4941,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 83c7667eab..9573fce2a8 100644 --- a/naga/src/back/hlsl/conv.rs +++ b/naga/src/back/hlsl/conv.rs @@ -125,6 +125,7 @@ impl crate::StorageFormat { Self::R8Snorm | Self::R16Snorm => "snorm float", Self::R8Uint | Self::R16Uint | Self::R32Uint => "uint", Self::R8Sint | Self::R16Sint | Self::R32Sint => "int", + Self::R64Uint => "uint64_t", Self::Rg16Float | Self::Rg32Float => "float2", Self::Rg8Unorm | Self::Rg16Unorm => "unorm float2", diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs index a227c65601..fdc5eb1695 100644 --- a/naga/src/back/msl/writer.rs +++ b/naga/src/back/msl/writer.rs @@ -1212,7 +1212,11 @@ impl Writer { ) -> BackendResult { write!(self.out, "{level}")?; self.put_expression(image, &context.expression, false)?; - let op = fun.to_msl(); + let op = if context.expression.resolve_type(value).scalar_width() == Some(8) { + fun.to_msl_64_bit()? + } else { + fun.to_msl() + }; write!(self.out, ".atomic_{}(", op)?; // coordinates in IR are int, but Metal expects uint self.put_cast_to_uint_scalar_or_vector(address.coordinate, &context.expression)?; diff --git a/naga/src/back/spv/image.rs b/naga/src/back/spv/image.rs index 5b50c71ee2..889c633d8f 100644 --- a/naga/src/back/spv/image.rs +++ b/naga/src/back/spv/image.rs @@ -1252,6 +1252,10 @@ impl BlockContext<'_> { base: NumericType::Scalar(scalar), class: 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]; diff --git a/naga/src/back/spv/instructions.rs b/naga/src/back/spv/instructions.rs index 32b8113c69..38aed8c351 100644 --- a/naga/src/back/spv/instructions.rs +++ b/naga/src/back/spv/instructions.rs @@ -1206,6 +1206,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 a0b696e6ca..dc25c4aee5 100644 --- a/naga/src/back/spv/writer.rs +++ b/naga/src/back/spv/writer.rs @@ -1079,10 +1079,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 9c7feef211..064e2e295b 100644 --- a/naga/src/back/wgsl/writer.rs +++ b/naga/src/back/wgsl/writer.rs @@ -2076,6 +2076,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 c7d61222f8..eee53b6ca6 100644 --- a/naga/src/front/glsl/parser/types.rs +++ b/naga/src/front/glsl/parser/types.rs @@ -420,6 +420,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 33ed4793cf..6baf74225c 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 0c9341eb62..00c19d877d 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/parse/mod.rs b/naga/src/front/wgsl/parse/mod.rs index 0233347c36..808c453490 100644 --- a/naga/src/front/wgsl/parse/mod.rs +++ b/naga/src/front/wgsl/parse/mod.rs @@ -1633,6 +1633,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/front/wgsl/to_wgsl.rs b/naga/src/front/wgsl/to_wgsl.rs index 4d401b0708..7d4c17f5f0 100644 --- a/naga/src/front/wgsl/to_wgsl.rs +++ b/naga/src/front/wgsl/to_wgsl.rs @@ -178,6 +178,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 1e0955b012..441c3864fc 100644 --- a/naga/src/lib.rs +++ b/naga/src/lib.rs @@ -642,6 +642,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 76698fd102..fafac8cb30 100644 --- a/naga/src/proc/mod.rs +++ b/naga/src/proc/mod.rs @@ -49,6 +49,7 @@ impl From for super::Scalar { 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, @@ -65,7 +66,11 @@ impl From for super::Scalar { Sf::Rgba16Unorm => Sk::Float, Sf::Rgba16Snorm => Sk::Float, }; - super::Scalar { kind, width: 4 } + let width = match format { + Sf::R64Uint => 8, + _ => 4, + }; + super::Scalar { kind, width } } } diff --git a/naga/src/valid/function.rs b/naga/src/valid/function.rs index 30a09bf849..12ef77abee 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -1214,6 +1214,34 @@ impl super::Validator { .with_span_handle(image, context.expressions)); } match format { + crate::StorageFormat::R64Uint => { + if !self.capabilities.intersects( + super::Capabilities::TEXTURE_INT64_ATOMIC, + ) { + return Err(FunctionError::MissingCapability( + super::Capabilities::TEXTURE_INT64_ATOMIC, + ) + .with_span_static( + span, + "missing capability for this operation", + )); + } + match fun { + crate::AtomicFunction::Min + | crate::AtomicFunction::Max => {} + _ => { + return Err( + FunctionError::InvalidImageAtomicFunction( + fun, + ) + .with_span_handle( + image, + context.expressions, + ), + ); + } + } + } crate::StorageFormat::R32Sint | crate::StorageFormat::R32Uint => { if !self diff --git a/naga/src/valid/mod.rs b/naga/src/valid/mod.rs index 1de5eb9872..c3d8a7f521 100644 --- a/naga/src/valid/mod.rs +++ b/naga/src/valid/mod.rs @@ -145,6 +145,8 @@ bitflags::bitflags! { const SHADER_INT64_ATOMIC_ALL_OPS = 0x100000; /// Support for atomic operations on images. const TEXTURE_ATOMIC = 0x200000; + /// Support for atomic operations on 64-bit images. + const TEXTURE_INT64_ATOMIC = 0x400000; } } diff --git a/naga/tests/in/atomicTexture-int64.param.ron b/naga/tests/in/atomicTexture-int64.param.ron new file mode 100644 index 0000000000..ffc7fb4cb7 --- /dev/null +++ b/naga/tests/in/atomicTexture-int64.param.ron @@ -0,0 +1,24 @@ +( + 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, + restrict_indexing: 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..2e63cb7de6 --- /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..3d27e8ca60 --- /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..622600cf4b --- /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 93d131b739..95edec5277 100644 --- a/naga/tests/snapshots.rs +++ b/naga/tests/snapshots.rs @@ -792,6 +792,10 @@ fn convert_wgsl() { "atomicTexture", Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL, ), + ( + "atomicTexture-int64", + Targets::SPIRV | Targets::METAL | Targets::HLSL | Targets::WGSL, + ), ( "atomicCompareExchange-int64", Targets::SPIRV | Targets::WGSL, diff --git a/tests/tests/image_atomics/image_64_atomics.wgsl b/tests/tests/image_atomics/image_64_atomics.wgsl new file mode 100644 index 0000000000..65c857c1ad --- /dev/null +++ b/tests/tests/image_atomics/image_64_atomics.wgsl @@ -0,0 +1,13 @@ +@group(0) @binding(0) +var image: texture_storage_2d; + +@compute +@workgroup_size(4, 4, 1) +fn cs_main(@builtin(local_invocation_id) id: vec3, @builtin(workgroup_id) group_id: vec3) { + let pixel = id + group_id * 4; + imageAtomicMax(image, pixel.xy, u64(pixel.x)); + + storageBarrier(); + + imageAtomicMin(image, pixel.xy, u64(pixel.y)); +} \ No newline at end of file diff --git a/tests/tests/image_atomics/mod.rs b/tests/tests/image_atomics/mod.rs index 3bbdef33c4..5e9826688d 100644 --- a/tests/tests/image_atomics/mod.rs +++ b/tests/tests/image_atomics/mod.rs @@ -3,6 +3,34 @@ use wgpu::ShaderModuleDescriptor; use wgpu_test::{gpu_test, GpuTestConfiguration, TestParameters, TestingContext}; +#[gpu_test] +static IMAGE_64_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: wgt::COPY_BYTES_PER_ROW_ALIGNMENT, + ..wgt::Limits::downlevel_webgl2_defaults() + }) + .features( + wgpu::Features::TEXTURE_INT64_ATOMIC + | wgpu::Features::SHADER_INT64 + | wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES, + ), + ) + .run_async(|ctx| async move { + test_format( + ctx, + wgpu::TextureFormat::R64Uint, + wgpu::include_wgsl!("image_64_atomics.wgsl"), + ) + .await; + }); + #[gpu_test] static IMAGE_32_ATOMICS: GpuTestConfiguration = GpuTestConfiguration::new() .parameters( diff --git a/wgpu-core/src/device/mod.rs b/wgpu-core/src/device/mod.rs index bf202702e9..ad0ed199f1 100644 --- a/wgpu-core/src/device/mod.rs +++ b/wgpu-core/src/device/mod.rs @@ -430,6 +430,10 @@ pub fn create_validator( Caps::TEXTURE_ATOMIC, features.contains(wgt::Features::TEXTURE_ATOMIC), ); + caps.set( + Caps::TEXTURE_INT64_ATOMIC, + features.contains(wgt::Features::TEXTURE_INT64_ATOMIC), + ); caps.set( Caps::MULTISAMPLED_SHADING, downlevel.contains(wgt::DownlevelFlags::MULTISAMPLED_SHADING), diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs index 7077fa2016..c8ceb21e17 100644 --- a/wgpu-core/src/validation.rs +++ b/wgpu-core/src/validation.rs @@ -312,6 +312,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, @@ -368,6 +369,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, @@ -705,6 +707,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 ad64f044cc..0ab30feec2 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 0e58cb730b..4cc79f24c1 100644 --- a/wgpu-hal/src/dx12/adapter.rs +++ b/wgpu-hal/src/dx12/adapter.rs @@ -382,6 +382,13 @@ impl super::Adapter { && features1.Int64ShaderOps.as_bool(), ); + features.set( + wgt::Features::TEXTURE_INT64_ATOMIC, + shader_model >= naga::back::hlsl::ShaderModel::V6_6 + && hr.is_ok() + && features1.Int64ShaderOps.as_bool(), + ); + features.set( wgt::Features::TEXTURE_ATOMIC, shader_model >= naga::back::hlsl::ShaderModel::V5_0, diff --git a/wgpu-hal/src/gles/adapter.rs b/wgpu-hal/src/gles/adapter.rs index 908e3107af..60c091931a 100644 --- a/wgpu-hal/src/gles/adapter.rs +++ b/wgpu-hal/src/gles/adapter.rs @@ -1081,6 +1081,7 @@ 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_ATOMIC, Tfc::SHADER_ATOMIC); + let image_64_atomic = feature_fn(wgt::Features::TEXTURE_INT64_ATOMIC, Tfc::SHADER_ATOMIC); match format { Tf::R8Unorm => filterable_renderable, @@ -1113,6 +1114,7 @@ impl crate::Adapter for super::Adapter { Tf::Rgb10a2Uint => renderable, Tf::Rgb10a2Unorm => filterable_renderable, Tf::Rg11b10Ufloat => filterable | float_renderable, + Tf::R64Uint => image_64_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 59bcf43b83..f87b7f706c 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 => (glow::RG32UI, glow::RED_INTEGER, glow::UNSIGNED_INT), //TODO? 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/metal/adapter.rs b/wgpu-hal/src/metal/adapter.rs index 2420f9f30a..28447a41c3 100644 --- a/wgpu-hal/src/metal/adapter.rs +++ b/wgpu-hal/src/metal/adapter.rs @@ -115,6 +115,12 @@ impl crate::Adapter for super::Adapter { Tfc::empty() }; + let image_64_atomic_if = if pc.int64_atomics { + Tfc::SHADER_ATOMIC + } else { + Tfc::empty() + }; + // Metal defined pixel format capabilities let all_caps = Tfc::SAMPLED_LINEAR | Tfc::STORAGE_WRITE_ONLY @@ -200,6 +206,12 @@ impl crate::Adapter for super::Adapter { flags.set(Tfc::STORAGE_WRITE_ONLY, pc.format_rg11b10_all); flags } + Tf::R64Uint => { + Tfc::COLOR_ATTACHMENT + | Tfc::STORAGE_WRITE_ONLY + | image_64_atomic_if + | read_write_tier1_if + } Tf::Rg32Uint | Tf::Rg32Sint => { Tfc::COLOR_ATTACHMENT | Tfc::STORAGE_WRITE_ONLY | msaa_count } @@ -935,6 +947,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_atomics && self.msl_version >= MTLLanguageVersion::V3_1, + ); features.set( F::TEXTURE_ATOMIC, self.msl_version >= MTLLanguageVersion::V3_1, @@ -1074,6 +1090,8 @@ impl super::PrivateCapabilities { Tf::Rgb10a2Uint => RGB10A2Uint, Tf::Rgb10a2Unorm => RGB10A2Unorm, Tf::Rg11b10Ufloat => RG11B10Float, + // Ruint64 textures are emulated on metal + Tf::R64Uint => RG32Uint, Tf::Rg32Uint => RG32Uint, Tf::Rg32Sint => RG32Sint, Tf::Rg32Float => RG32Float, diff --git a/wgpu-hal/src/vulkan/adapter.rs b/wgpu-hal/src/vulkan/adapter.rs index f0bf22d51c..a1369f0229 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) { @@ -591,6 +608,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( @@ -1021,6 +1048,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); @@ -1311,6 +1343,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 @@ -1807,6 +1846,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, diff --git a/wgpu-hal/src/vulkan/conv.rs b/wgpu-hal/src/vulkan/conv.rs index 21ebd6c7b5..076ae7b04c 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 c94b7c20c6..23f4324527 100644 --- a/wgpu-types/src/lib.rs +++ b/wgpu-types/src/lib.rs @@ -948,6 +948,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 image 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; } } @@ -2578,6 +2587,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. @@ -2864,6 +2877,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, @@ -2992,6 +3006,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", @@ -3234,6 +3249,7 @@ impl TextureFormat { | Self::Rgb10a2Uint | Self::Rgb10a2Unorm | Self::Rg11b10Ufloat + | Self::R64Uint | Self::Rg32Uint | Self::Rg32Sint | Self::Rg32Float @@ -3357,6 +3373,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, @@ -3469,6 +3487,7 @@ impl TextureFormat { Self::Rgb10a2Uint => ( msaa, attachment), Self::Rgb10a2Unorm => (msaa_resolve, attachment), Self::Rg11b10Ufloat => ( msaa, rg11b10f), + Self::R64Uint => ( s_ro_wo, atomic), Self::Rg32Uint => ( s_ro_wo, all_flags), Self::Rg32Sint => ( s_ro_wo, all_flags), Self::Rg32Float => ( s_ro_wo, all_flags), @@ -3590,6 +3609,7 @@ impl TextureFormat { | Self::Rg16Uint | Self::Rgba16Uint | Self::R32Uint + | Self::R64Uint | Self::Rg32Uint | Self::Rgba32Uint | Self::Rgb10a2Uint => Some(uint), @@ -3720,7 +3740,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), @@ -3809,6 +3829,7 @@ impl TextureFormat { | Self::Rgba16Unorm | Self::Rgba16Snorm | Self::Rgba16Float + | Self::R64Uint | Self::Rg32Uint | Self::Rg32Sint | Self::Rg32Float @@ -3889,6 +3910,7 @@ impl TextureFormat { Self::R32Uint | Self::R32Sint | Self::R32Float + | Self::R64Uint | Self::Rg32Uint | Self::Rg32Sint | Self::Rg32Float @@ -3957,7 +3979,8 @@ impl TextureFormat { | Self::R16Float | Self::R32Uint | Self::R32Sint - | Self::R32Float => 1, + | Self::R32Float + | Self::R64Uint => 1, Self::Rg8Unorm | Self::Rg8Snorm @@ -4209,6 +4232,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() @@ -4505,6 +4532,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