From e323b14c903dae8d41addb38f5d722ebf7e541d5 Mon Sep 17 00:00:00 2001 From: atlas dostal Date: Fri, 27 Sep 2024 20:48:02 -0400 Subject: [PATCH] unfuck things --- naga/src/back/glsl/mod.rs | 18 +++-- naga/src/back/msl/writer.rs | 15 +++-- 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 +- naga/tests/out/msl/atomicTexture-int64.msl | 2 +- naga/tests/out/spv/atomicTexture-int64.spvasm | 66 +++++++++---------- 10 files changed, 74 insertions(+), 66 deletions(-) diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index f85e9d4bd3..2b0c18abc8 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 b41b98673d..cdaab1eba7 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, @@ -1202,7 +1207,7 @@ impl Writer { 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, ", metal::memory_order_relaxed);")?; + writeln!(self.out, ", {NAMESPACE}::memory_order_relaxed);")?; Ok(()) } diff --git a/naga/src/back/spv/mod.rs b/naga/src/back/spv/mod.rs index 32bd1fcecf..fdc42c7310 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: crate::Scalar::from(format), 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 e4324c2c08..066f0b8fef 100644 --- a/naga/src/valid/function.rs +++ b/naga/src/valid/function.rs @@ -1093,10 +1093,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(), } } _ => { diff --git a/naga/tests/out/msl/atomicTexture-int64.msl b/naga/tests/out/msl/atomicTexture-int64.msl index 1ee9b9b8b6..0c19d2011b 100644 --- a/naga/tests/out/msl/atomicTexture-int64.msl +++ b/naga/tests/out/msl/atomicTexture-int64.msl @@ -9,7 +9,7 @@ struct cs_mainInput { }; kernel void cs_main( metal::uint3 id [[thread_position_in_threadgroup]] -, metal::texture2d image [[user(fake0)]] +, metal::texture2d image [[user(fake0)]] ) { image.atomic_max(metal::uint2(metal::int2(0, 0)), 1uL, metal::memory_order_relaxed); metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup); diff --git a/naga/tests/out/spv/atomicTexture-int64.spvasm b/naga/tests/out/spv/atomicTexture-int64.spvasm index 3be12a5777..24b3c8dec1 100644 --- a/naga/tests/out/spv/atomicTexture-int64.spvasm +++ b/naga/tests/out/spv/atomicTexture-int64.spvasm @@ -8,43 +8,43 @@ OpCapability Int64 OpExtension "SPV_EXT_shader_image_int64" %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %14 "cs_main" %11 -OpExecutionMode %14 LocalSize 2 1 1 -OpDecorate %8 NonReadable -OpDecorate %8 DescriptorSet 0 -OpDecorate %8 Binding 0 -OpDecorate %11 BuiltIn LocalInvocationId +OpEntryPoint GLCompute %15 "cs_main" %12 +OpExecutionMode %15 LocalSize 2 1 1 +OpDecorate %9 NonReadable +OpDecorate %9 DescriptorSet 0 +OpDecorate %9 Binding 0 +OpDecorate %12 BuiltIn LocalInvocationId %2 = OpTypeVoid -%4 = OpTypeInt 32 0 +%4 = OpTypeInt 64 0 %3 = OpTypeImage %4 2D 0 0 0 2 R64ui -%5 = OpTypeVector %4 3 -%6 = OpTypeInt 32 1 -%7 = OpTypeVector %6 2 -%9 = OpTypePointer UniformConstant %3 -%8 = OpVariable %9 UniformConstant -%12 = OpTypePointer Input %5 -%11 = OpVariable %12 Input -%15 = OpTypeFunction %2 -%17 = OpConstant %6 0 -%18 = OpConstantComposite %7 %17 %17 -%19 = OpConstantNull %6 -%20 = OpTypeInt 64 0 -%21 = OpConstant %20 1 -%23 = OpTypePointer Image %20 -%24 = OpConstant %6 4 -%25 = OpConstant %4 0 -%28 = OpConstant %4 2 -%29 = OpConstant %4 264 -%14 = OpFunction %2 None %15 -%10 = OpLabel -%13 = OpLoad %5 %11 -%16 = OpLoad %3 %8 +%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 +%24 = OpConstant %7 4 +%25 = 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 -%26 = OpImageTexelPointer %23 %16 %18 %19 -%27 = OpAtomicUMax %20 %26 %24 %25 %21 +%26 = OpImageTexelPointer %23 %17 %19 %20 +%27 = OpAtomicUMax %4 %26 %24 %25 %21 OpControlBarrier %28 %28 %29 -%30 = OpImageTexelPointer %23 %16 %18 %19 -%31 = OpAtomicUMin %20 %30 %24 %25 %21 +%30 = OpImageTexelPointer %23 %17 %19 %20 +%31 = OpAtomicUMin %4 %30 %24 %25 %21 OpReturn OpFunctionEnd \ No newline at end of file