Skip to content

Commit

Permalink
unfuck things
Browse files Browse the repository at this point in the history
  • Loading branch information
@atlv24
atlv24 committed Sep 28, 2024
1 parent b800087 commit e323b14
Show file tree
Hide file tree
Showing 10 changed files with 74 additions and 66 deletions.
18 changes: 11 additions & 7 deletions naga/src/back/glsl/mod.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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(), "", ""),
};

Expand All @@ -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,
Expand Down
15 changes: 10 additions & 5 deletions naga/src/back/msl/writer.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -214,22 +214,27 @@ 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 {
("_ms", "read")
} 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
Expand All @@ -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,
Expand Down Expand Up @@ -1202,7 +1207,7 @@ impl<W: Write> Writer<W> {
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(())
}
Expand Down
11 changes: 7 additions & 4 deletions naga/src/back/spv/mod.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down Expand Up @@ -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(),
Expand Down
5 changes: 1 addition & 4 deletions naga/src/back/spv/writer.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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));
Expand Down
2 changes: 1 addition & 1 deletion naga/src/front/glsl/functions.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}
Expand Down
11 changes: 8 additions & 3 deletions naga/src/proc/mod.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -20,10 +20,10 @@ pub use namer::{EntryPointIndex, NameKey, Namer};
pub use terminator::ensure_block_returns;
pub use typifier::{ResolveContext, ResolveError, TypeResolution};

impl From<super::StorageFormat> for super::ScalarKind {
impl From<super::StorageFormat> 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,
Expand Down Expand Up @@ -65,7 +65,12 @@ impl From<super::StorageFormat> 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 }
}
}

Expand Down
5 changes: 1 addition & 4 deletions naga/src/proc/typifier.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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,
},
}),
Expand Down
5 changes: 1 addition & 4 deletions naga/src/valid/function.rs
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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(),
}
}
_ => {
Expand Down
2 changes: 1 addition & 1 deletion naga/tests/out/msl/atomicTexture-int64.msl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@ struct cs_mainInput {
};
kernel void cs_main(
metal::uint3 id [[thread_position_in_threadgroup]]
, metal::texture2d<uint, metal::access::write> image [[user(fake0)]]
, metal::texture2d<ulong, metal::access::write> 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);
Expand Down
66 changes: 33 additions & 33 deletions naga/tests/out/spv/atomicTexture-int64.spvasm
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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

0 comments on commit e323b14

Please sign in to comment.