From c3190d59f34dccdda465c8a10333010654b9114c Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Mon, 16 Oct 2023 14:10:03 -0400 Subject: [PATCH 1/3] Run snapshots --- src/back/dot/mod.rs | 5 +- src/back/glsl/mod.rs | 98 ++++-- src/back/hlsl/writer.rs | 81 ++++- src/back/msl/writer.rs | 65 +++- src/back/spv/image.rs | 100 ++++++- src/back/wgsl/writer.rs | 3 +- src/compact/expressions.rs | 4 +- src/front/wgsl/lower/mod.rs | 7 +- src/lib.rs | 1 + src/proc/mod.rs | 2 +- src/valid/analyzer.rs | 2 +- src/valid/expression.rs | 34 +++ src/valid/handles.rs | 4 +- tests/in/image.wgsl | 1 + .../glsl/image.texture_sample.Fragment.glsl | 8 +- tests/out/hlsl/image.hlsl | 11 +- tests/out/msl/image.msl | 9 +- tests/out/spv/image.spvasm | 280 +++++++++--------- tests/out/wgsl/image.wgsl | 7 +- 19 files changed, 537 insertions(+), 185 deletions(-) diff --git a/src/back/dot/mod.rs b/src/back/dot/mod.rs index 1556371df1..eed7c5bd99 100644 --- a/src/back/dot/mod.rs +++ b/src/back/dot/mod.rs @@ -460,8 +460,9 @@ fn write_function_expressions( edges.insert("array_index", expr); } match level { - crate::SampleLevel::Auto => {} - crate::SampleLevel::Zero => {} + crate::SampleLevel::Auto + | crate::SampleLevel::Zero + | crate::SampleLevel::Base => {} crate::SampleLevel::Exact(expr) => { edges.insert("level", expr); } diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 77e206d3c6..267d484ea2 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -72,6 +72,10 @@ pub const SUPPORTED_ES_VERSIONS: &[u16] = &[300, 310, 320]; /// of detail for bounds checking in `ImageLoad` const CLAMPED_LOD_SUFFIX: &str = "_clamped_lod"; +/// The suffix of the variable that will hold the calculated half-texel +/// for use with `textureSampleBaseClampToEdge` +const HALF_TEXEL_SUFFIX: &str = "_half_texel"; + pub(crate) const MODF_FUNCTION: &str = "naga_modf"; pub(crate) const FREXP_FUNCTION: &str = "naga_frexp"; @@ -1855,6 +1859,17 @@ impl<'a, W: Write> Writer<'a, W> { } } + // If we are going to write a `textureSampleBaseClampToEdge` next, + // precompute the half-texel before clamping the coordinates. + if let crate::Expression::ImageSample { + level: crate::SampleLevel::Base, + image, + .. + } = ctx.expressions[handle] + { + self.write_half_texel(ctx, handle, image, level)? + } + if let Some(name) = expr_name { write!(self.out, "{level}")?; self.write_named_expr(handle, name, handle, ctx)?; @@ -2491,6 +2506,11 @@ impl<'a, W: Write> Writer<'a, W> { ))) } crate::SampleLevel::Auto => {} + crate::SampleLevel::Base => { + unreachable!( + "textureSampleBaseClampToEdge should not have passed validation" + ) + } } } @@ -2517,6 +2537,7 @@ impl<'a, W: Write> Writer<'a, W> { } } crate::SampleLevel::Gradient { .. } => "textureGrad", + crate::SampleLevel::Base => "textureLod", }; let offset_name = match offset { Some(_) => "Offset", @@ -2548,24 +2569,38 @@ impl<'a, W: Write> Writer<'a, W> { let tex_1d_hack = dim == crate::ImageDimension::D1 && self.options.version.is_es(); let is_vec = tex_1d_hack || coord_dim != 1; - // Compose a new texture coordinates vector - if is_vec { - write!(self.out, "vec{}(", coord_dim + tex_1d_hack as u8)?; - } - self.write_expr(coordinate, ctx)?; - if tex_1d_hack { - write!(self.out, ", 0.0")?; - } - if let Some(expr) = array_index { - write!(self.out, ", ")?; - self.write_expr(expr, ctx)?; - } - if merge_depth_ref { - write!(self.out, ", ")?; - self.write_expr(depth_ref.unwrap(), ctx)?; - } - if is_vec { - write!(self.out, ")")?; + + if level == crate::SampleLevel::Base { + // clamp the coordinates to [ half_texel, 1 - half_texel ] + write!(self.out, "clamp(")?; + self.write_expr(coordinate, ctx)?; + write!( + self.out, + ", {}{}{}, vec2(1.0) - {0}{1}{2})", + back::BAKE_PREFIX, + expr.index(), + HALF_TEXEL_SUFFIX + )? + } else { + // Compose a new texture coordinates vector + if is_vec { + write!(self.out, "vec{}(", coord_dim + tex_1d_hack as u8)?; + } + self.write_expr(coordinate, ctx)?; + if tex_1d_hack { + write!(self.out, ", 0.0")?; + } + if let Some(expr) = array_index { + write!(self.out, ", ")?; + self.write_expr(expr, ctx)?; + } + if merge_depth_ref { + write!(self.out, ", ")?; + self.write_expr(depth_ref.unwrap(), ctx)?; + } + if is_vec { + write!(self.out, ")")?; + } } if let (Some(expr), false) = (depth_ref, merge_depth_ref) { @@ -2618,6 +2653,9 @@ impl<'a, W: Write> Writer<'a, W> { self.write_expr(y, ctx)?; } } + crate::SampleLevel::Base => { + write!(self.out, ", 0.0")?; + } } if let Some(constant) = offset { @@ -3480,6 +3518,30 @@ impl<'a, W: Write> Writer<'a, W> { Ok(()) } + /// Helper function to write the local holding the half-texel + /// for use with `textureSampleBaseClampToEdge` + fn write_half_texel( + &mut self, + ctx: &back::FunctionCtx, + expr: Handle, + image: Handle, + level: back::Level, + ) -> Result<(), Error> { + write!( + self.out, + "{level}vec2 {}{}{} = vec2(0.5) / vec2(textureSize(", + back::BAKE_PREFIX, + expr.index(), + HALF_TEXEL_SUFFIX, + )?; + + self.write_expr(image, ctx)?; + + writeln!(self.out, ", 0));")?; + + Ok(()) + } + // Helper method used to retrieve how many elements a coordinate vector // for the images operations need. fn get_coordinate_vector_size(&self, dim: crate::ImageDimension, arrayed: bool) -> u8 { diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 590850fc4d..aed1e4d53f 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -17,6 +17,10 @@ const SPECIAL_BASE_VERTEX: &str = "base_vertex"; const SPECIAL_BASE_INSTANCE: &str = "base_instance"; const SPECIAL_OTHER: &str = "other"; +/// The suffix of the variable that will hold the calculated half-texel +/// for use with `textureSampleBaseClampToEdge` +const HALF_TEXEL_SUFFIX: &str = "_half_texel"; + pub(crate) const MODF_FUNCTION: &str = "naga_modf"; pub(crate) const FREXP_FUNCTION: &str = "naga_frexp"; @@ -1339,6 +1343,17 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { None }; + // If we are going to write a `textureSampleBaseClampToEdge` next, + // precompute the half-texel before clamping the coordinates. + if let crate::Expression::ImageSample { + level: crate::SampleLevel::Base, + image, + .. + } = func_ctx.expressions[handle] + { + self.write_half_texel(module, handle, image, level, func_ctx)? + } + if let Some(name) = expr_name { write!(self.out, "{level}")?; self.write_named_expr(module, handle, name, handle, func_ctx)?; @@ -2394,7 +2409,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { let level_str = match level { Sl::Zero if gather.is_none() => "LevelZero", Sl::Auto | Sl::Zero => "", - Sl::Exact(_) => "Level", + Sl::Exact(_) | Sl::Base => "Level", Sl::Bias(_) => "Bias", Sl::Gradient { .. } => "Grad", }; @@ -2403,14 +2418,26 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { write!(self.out, ".{base_str}{cmp_str}{component_str}{level_str}(")?; self.write_expr(module, sampler, func_ctx)?; write!(self.out, ", ")?; - self.write_texture_coordinates( - "float", - coordinate, - array_index, - None, - module, - func_ctx, - )?; + if level == Sl::Base { + // clamp the coordinates to [ half_texel, 1 - half_texel ] + write!(self.out, "clamp(")?; + self.write_expr(module, coordinate, func_ctx)?; + write!( + self.out, + ", _expr{}{}, (1.0).xx - _expr{0}{1})", + expr.index(), + HALF_TEXEL_SUFFIX + )? + } else { + self.write_texture_coordinates( + "float", + coordinate, + array_index, + None, + module, + func_ctx, + )?; + } if let Some(depth_ref) = depth_ref { write!(self.out, ", ")?; @@ -2433,6 +2460,9 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { write!(self.out, ", ")?; self.write_expr(module, y, func_ctx)?; } + Sl::Base => { + write!(self.out, ", 0.0")?; + } } if let Some(offset) = offset { @@ -3259,6 +3289,39 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } Ok(()) } + + /// Helper function to write the locals holding the half-texel + /// for use with `textureSampleBaseClampToEdge` + fn write_half_texel( + &mut self, + module: &Module, + expr: Handle, + image: Handle, + level: back::Level, + func_ctx: &back::FunctionCtx, + ) -> Result<(), Error> { + let prefix = format!("_expr{}", expr.index()); + + writeln!(self.out, "{level}float2 {prefix}_dim;")?; + + // will not be used, but required for the method call + writeln!(self.out, "{level}float {prefix}_num;")?; + + write!(self.out, "{level}")?; + self.write_expr(module, image, func_ctx)?; + + writeln!( + self.out, + ".GetDimensions(0u, {prefix}_dim.x, {prefix}_dim.y, {prefix}_num);" + )?; + + writeln!( + self.out, + "{level}float2 {prefix}{HALF_TEXEL_SUFFIX} = (0.5).xx / {prefix}_dim;" + )?; + + Ok(()) + } } pub(super) struct MatrixType { diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index d0f5413136..d41cc23f25 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -35,6 +35,10 @@ const RAY_QUERY_FUN_MAP_INTERSECTION: &str = "_map_intersection_type"; pub(crate) const MODF_FUNCTION: &str = "naga_modf"; pub(crate) const FREXP_FUNCTION: &str = "naga_frexp"; +/// The suffix of the variable that will hold the calculated half-texel +/// for use with `textureSampleBaseClampToEdge` +const HALF_TEXEL_SUFFIX: &str = "_half_texel"; + /// Write the Metal name for a Naga numeric type: scalar, vector, or matrix. /// /// The `sizes` slice determines whether this function writes a @@ -767,6 +771,9 @@ impl Writer { self.put_expression(y, context, true)?; write!(self.out, ")")?; } + crate::SampleLevel::Base => { + write!(self.out, ", {NAMESPACE}::level(0.0)")?; + } } Ok(()) } @@ -1473,7 +1480,22 @@ impl Writer { write!(self.out, ".{main_op}{comparison_op}(")?; self.put_expression(sampler, context, true)?; write!(self.out, ", ")?; - self.put_expression(coordinate, context, true)?; + + if level == crate::SampleLevel::Base { + // clamp the coordinates to [ half_texel, 1 - half_texel ] + write!(self.out, "{NAMESPACE}::clamp(")?; + self.put_expression(coordinate, context, true)?; + write!( + self.out, + ", {}{}{}, {NAMESPACE}::float2(1.0) - {0}{1}{2})", + back::BAKE_PREFIX, + expr_handle.index(), + HALF_TEXEL_SUFFIX + )? + } else { + self.put_expression(coordinate, context, true)?; + } + if let Some(expr) = array_index { write!(self.out, ", ")?; self.put_expression(expr, context, true)?; @@ -2646,6 +2668,17 @@ impl Writer { } }; + // If we are going to write a `textureSampleBaseClampToEdge` next, + // precompute the half-texel before clamping the coordinates. + if let crate::Expression::ImageSample { + level: crate::SampleLevel::Base, + image, + .. + } = context.expression.function.expressions[handle] + { + self.write_half_texel(handle, image, level, &context.expression)? + } + if let Some(name) = expr_name { write!(self.out, "{level}")?; self.start_baking_expression(handle, &context.expression, &name)?; @@ -4190,6 +4223,36 @@ impl Writer { } Ok(()) } + + /// Helper function to write the locals holding the half-texel + /// for use with `textureSampleBaseClampToEdge` + fn write_half_texel( + &mut self, + expr_handle: Handle, + image: Handle, + level: back::Level, + context: &ExpressionContext, + ) -> Result<(), Error> { + let prefix = format!("{}{}", back::BAKE_PREFIX, expr_handle.index()); + + write!( + self.out, + "{level}{NAMESPACE}::float2 {prefix}_dim = {NAMESPACE}::float2(" + )?; + + self.put_expression(image, context, true)?; + write!(self.out, ".get_width(), ")?; + + self.put_expression(image, context, true)?; + writeln!(self.out, ".get_height());")?; + + writeln!( + self.out, + "{level}{NAMESPACE}::float2 {prefix}{HALF_TEXEL_SUFFIX} = {NAMESPACE}::float2(0.5) / {prefix}_dim;" + )?; + + Ok(()) + } } /// Initializing workgroup variables is more tricky for Metal because we have to deal diff --git a/src/back/spv/image.rs b/src/back/spv/image.rs index 1630dc0ddb..003b821835 100644 --- a/src/back/spv/image.rs +++ b/src/back/spv/image.rs @@ -379,6 +379,95 @@ impl<'w> BlockContext<'w> { }) } + /// Helper function to clamp the image coordinates to [ half_texel, 1 - half_texel ] + /// for use with `textureSampleBaseClampToEdge` + fn write_clamped_image_coordinates( + &mut self, + image_id: Word, + coordinates: Handle, + block: &mut Block, + ) -> Result { + // query image size + let dim_type_id = self.get_type_id( + LocalType::Value { + vector_size: Some(crate::VectorSize::Bi), + kind: crate::ScalarKind::Sint, + width: 4, + pointer_space: None, + } + .into(), + ); + let dim_id = self.gen_id(); + let mut inst = + Instruction::image_query(spirv::Op::ImageQuerySizeLod, dim_type_id, dim_id, image_id); + inst.add_operand(self.writer.get_constant_scalar(crate::Literal::U32(0))); + block.body.push(inst); + + let vec_type = LocalType::Value { + vector_size: Some(crate::VectorSize::Bi), + kind: crate::ScalarKind::Float, + width: 4, + pointer_space: None, + } + .into(); + let vec_type_id = self.get_type_id(vec_type); + + // conv vec2i to vec2f + let conv_id = self.gen_id(); + block.body.push(Instruction::unary( + spirv::Op::ConvertSToF, + vec_type_id, + conv_id, + dim_id, + )); + + // vec2(0.5) / dim + self.temp_list.clear(); + self.temp_list + .resize(2, self.writer.get_constant_scalar(crate::Literal::F32(0.5))); + let vec_id = self + .writer + .get_constant_composite(vec_type, &self.temp_list); + let half_texel_id = self.gen_id(); + block.body.push(Instruction::binary( + spirv::Op::FDiv, + vec_type_id, + half_texel_id, + vec_id, + conv_id, + )); + + // vec2(1.0) - half_texel + self.temp_list + .fill(self.writer.get_constant_scalar(crate::Literal::F32(1.0))); + let vec_id = self + .writer + .get_constant_composite(vec_type, &self.temp_list); + let sub_id = self.gen_id(); + block.body.push(Instruction::binary( + spirv::Op::FSub, + vec_type_id, + sub_id, + vec_id, + half_texel_id, + )); + + // clamp coords + let coord_id = self + .write_image_coordinates(coordinates, None, block)? + .value_id; + let clamp_id = self.gen_id(); + block.body.push(Instruction::ext_inst( + self.writer.gl450_ext_inst_id, + spirv::GLOp::FClamp, + vec_type_id, + clamp_id, + &[coord_id, half_texel_id, sub_id], + )); + + Ok(clamp_id) + } + pub(super) fn get_handle_id(&mut self, expr_handle: Handle) -> Word { let id = match self.ir_function.expressions[expr_handle] { crate::Expression::GlobalVariable(handle) => { @@ -858,9 +947,12 @@ impl<'w> BlockContext<'w> { self.get_type_id(LookupType::Local(LocalType::SampledImage { image_type_id })); let sampler_id = self.get_handle_id(sampler); - let coordinates_id = self - .write_image_coordinates(coordinate, array_index, block)? - .value_id; + let coordinates_id = if level == crate::SampleLevel::Base { + self.write_clamped_image_coordinates(image_id, coordinate, block)? + } else { + self.write_image_coordinates(coordinate, array_index, block)? + .value_id + }; let sampled_image_id = self.gen_id(); block.body.push(Instruction::sampled_image( @@ -891,7 +983,7 @@ impl<'w> BlockContext<'w> { } inst } - (crate::SampleLevel::Zero, None) => { + (crate::SampleLevel::Zero | crate::SampleLevel::Base, None) => { let mut inst = Instruction::image_sample( sample_result_type_id, id, diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index 36380952af..cb84309fd2 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -1257,6 +1257,7 @@ impl Writer { Sl::Zero | Sl::Exact(_) => "Level", Sl::Bias(_) => "Bias", Sl::Gradient { .. } => "Grad", + Sl::Base => "BaseClampToEdge", }; write!(self.out, "textureSample{suffix_cmp}{suffix_level}(")?; @@ -1277,7 +1278,7 @@ impl Writer { } match level { - Sl::Auto => {} + Sl::Auto | Sl::Base => {} Sl::Zero => { // Level 0 is implied for depth comparison if depth_ref.is_none() { diff --git a/src/compact/expressions.rs b/src/compact/expressions.rs index c1326e92be..b008680a6e 100644 --- a/src/compact/expressions.rs +++ b/src/compact/expressions.rs @@ -96,7 +96,7 @@ impl<'tracer> ExpressionTracer<'tracer> { } use crate::SampleLevel as Sl; match *level { - Sl::Auto | Sl::Zero => {} + Sl::Auto | Sl::Zero | Sl::Base => {} Sl::Exact(expr) | Sl::Bias(expr) => work_list.push(expr), Sl::Gradient { x, y } => work_list.extend([x, y]), } @@ -366,7 +366,7 @@ impl ModuleMap { use crate::SampleLevel as Sl; match *level { - Sl::Auto | Sl::Zero => {} + Sl::Auto | Sl::Zero | Sl::Base => {} Sl::Exact(ref mut expr) => adjust(expr), Sl::Bias(ref mut expr) => adjust(expr), Sl::Gradient { diff --git a/src/front/wgsl/lower/mod.rs b/src/front/wgsl/lower/mod.rs index c883732e8b..0627427ecd 100644 --- a/src/front/wgsl/lower/mod.rs +++ b/src/front/wgsl/lower/mod.rs @@ -800,7 +800,7 @@ enum Texture { SampleCompareLevel, SampleGrad, SampleLevel, - // SampleBaseClampToEdge, + SampleBaseClampToEdge, } impl Texture { @@ -815,7 +815,7 @@ impl Texture { "textureSampleCompareLevel" => Self::SampleCompareLevel, "textureSampleGrad" => Self::SampleGrad, "textureSampleLevel" => Self::SampleLevel, - // "textureSampleBaseClampToEdge" => Some(Self::SampleBaseClampToEdge), + "textureSampleBaseClampToEdge" => Self::SampleBaseClampToEdge, _ => return None, }) } @@ -831,7 +831,7 @@ impl Texture { Self::SampleCompareLevel => 5, Self::SampleGrad => 6, Self::SampleLevel => 5, - // Self::SampleBaseClampToEdge => 3, + Self::SampleBaseClampToEdge => 3, } } } @@ -2428,6 +2428,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let level = self.expression(args.next()?, ctx.reborrow())?; (crate::SampleLevel::Exact(level), None) } + Texture::SampleBaseClampToEdge => (crate::SampleLevel::Base, None), }; let offset = args diff --git a/src/lib.rs b/src/lib.rs index 9d70190421..69dfc7d980 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -1221,6 +1221,7 @@ pub enum SampleLevel { x: Handle, y: Handle, }, + Base, } /// Type of an image query. diff --git a/src/proc/mod.rs b/src/proc/mod.rs index d50f165a56..9b84779a92 100644 --- a/src/proc/mod.rs +++ b/src/proc/mod.rs @@ -508,7 +508,7 @@ impl crate::SampleLevel { pub const fn implicit_derivatives(&self) -> bool { match *self { Self::Auto | Self::Bias(_) => true, - Self::Zero | Self::Exact(_) | Self::Gradient { .. } => false, + Self::Zero | Self::Exact(_) | Self::Gradient { .. } | Self::Base => false, } } } diff --git a/src/valid/analyzer.rs b/src/valid/analyzer.rs index ff1db071c8..0fa474da5e 100644 --- a/src/valid/analyzer.rs +++ b/src/valid/analyzer.rs @@ -619,7 +619,7 @@ impl FunctionInfo { // "nur" == "Non-Uniform Result" let array_nur = array_index.and_then(|h| self.add_ref(h)); let level_nur = match level { - Sl::Auto | Sl::Zero => None, + Sl::Auto | Sl::Zero | Sl::Base => None, Sl::Exact(h) | Sl::Bias(h) => self.add_ref(h), Sl::Gradient { x, y } => self.add_ref(x).or(self.add_ref(y)), }; diff --git a/src/valid/expression.rs b/src/valid/expression.rs index 95225a3926..e47a7872d6 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -74,6 +74,10 @@ pub enum ExpressionError { ExpectedSamplerType(Handle), #[error("Unable to operate on image class {0:?}")] InvalidImageClass(crate::ImageClass), + #[error("Needs to be texture_2d")] + InvalidSampleBaseClampToEdgeImageType, + #[error("Needs to be vec2")] + InvalidSampleBaseClampToEdgeCoordinateType, #[error("Derivatives can only be taken from scalar and vector floats")] InvalidDerivative, #[error("Image array index parameter is misplaced")] @@ -559,6 +563,36 @@ impl super::Validator { } ShaderStages::all() } + crate::SampleLevel::Base => { + // TODO: handle external textures + match module.types[image_ty].inner { + Ti::Image { + dim: crate::ImageDimension::D2, + arrayed: false, + class: + crate::ImageClass::Sampled { + kind: Sk::Float, + multi: false, + }, + } => {} + _ => { + return Err(ExpressionError::InvalidSampleBaseClampToEdgeImageType) + } + } + match resolver[coordinate] { + Ti::Vector { + size: crate::VectorSize::Bi, + kind: Sk::Float, + .. + } => {} + _ => { + return Err( + ExpressionError::InvalidSampleBaseClampToEdgeCoordinateType, + ) + } + } + ShaderStages::all() + } } } E::ImageLoad { diff --git a/src/valid/handles.rs b/src/valid/handles.rs index c68ded074b..49b32cc8e5 100644 --- a/src/valid/handles.rs +++ b/src/valid/handles.rs @@ -302,7 +302,9 @@ impl super::Validator { .check_dep_opt(array_index)?; match level { - crate::SampleLevel::Auto | crate::SampleLevel::Zero => (), + crate::SampleLevel::Auto + | crate::SampleLevel::Zero + | crate::SampleLevel::Base => (), crate::SampleLevel::Exact(expr) => { handle.check_dep(expr)?; } diff --git a/tests/in/image.wgsl b/tests/in/image.wgsl index 2bae8f9d80..b6031574d1 100644 --- a/tests/in/image.wgsl +++ b/tests/in/image.wgsl @@ -136,6 +136,7 @@ fn texture_sample() -> @location(0) vec4 { a += textureSample(image_cube_array, sampler_reg, tc3, 0); a += textureSampleLevel(image_cube_array, sampler_reg, tc3, 0, level); a += textureSampleBias(image_cube_array, sampler_reg, tc3, 0, 2.0); + a += textureSampleBaseClampToEdge(image_2d, sampler_reg, tc); return a; } diff --git a/tests/out/glsl/image.texture_sample.Fragment.glsl b/tests/out/glsl/image.texture_sample.Fragment.glsl index 97be5a59d0..6e94ca94ae 100644 --- a/tests/out/glsl/image.texture_sample.Fragment.glsl +++ b/tests/out/glsl/image.texture_sample.Fragment.glsl @@ -84,8 +84,12 @@ void main() { vec4 _e135 = texture(_group_0_binding_6_fs, vec4(tc3_, 0), 2.0); vec4 _e136 = a; a = (_e136 + _e135); - vec4 _e138 = a; - _fs2p_location0 = _e138; + vec2 _e140_half_texel = vec2(0.5) / vec2(textureSize(_group_0_binding_1_fs, 0)); + vec4 _e140 = textureLod(_group_0_binding_1_fs, clamp(tc, _e140_half_texel, vec2(1.0) - _e140_half_texel), 0.0); + vec4 _e141 = a; + a = (_e141 + _e140); + vec4 _e143 = a; + _fs2p_location0 = _e143; return; } diff --git a/tests/out/hlsl/image.hlsl b/tests/out/hlsl/image.hlsl index 7fbd68b105..6b339a55ad 100644 --- a/tests/out/hlsl/image.hlsl +++ b/tests/out/hlsl/image.hlsl @@ -312,8 +312,15 @@ float4 texture_sample() : SV_Target0 float4 _expr135 = image_cube_array.SampleBias(sampler_reg, float4(tc3_, 0), 2.0); float4 _expr136 = a; a = (_expr136 + _expr135); - float4 _expr138 = a; - return _expr138; + float2 _expr140_dim; + float _expr140_num; + image_2d.GetDimensions(0u, _expr140_dim.x, _expr140_dim.y, _expr140_num); + float2 _expr140_half_texel = (0.5).xx / _expr140_dim; + float4 _expr140 = image_2d.SampleLevel(sampler_reg, clamp(tc, _expr140_half_texel, (1.0).xx - _expr140_half_texel), 0.0); + float4 _expr141 = a; + a = (_expr141 + _expr140); + float4 _expr143 = a; + return _expr143; } float texture_sample_comparison() : SV_Target0 diff --git a/tests/out/msl/image.msl b/tests/out/msl/image.msl index e390c2e0fc..8dafbdf940 100644 --- a/tests/out/msl/image.msl +++ b/tests/out/msl/image.msl @@ -185,8 +185,13 @@ fragment texture_sampleOutput texture_sample( metal::float4 _e135 = image_cube_array.sample(sampler_reg, tc3_, 0, metal::bias(2.0)); metal::float4 _e136 = a; a = _e136 + _e135; - metal::float4 _e138 = a; - return texture_sampleOutput { _e138 }; + metal::float2 _e140_dim = metal::float2(image_2d.get_width(), image_2d.get_height()); + metal::float2 _e140_half_texel = metal::float2(0.5) / _e140_dim; + metal::float4 _e140 = image_2d.sample(sampler_reg, metal::clamp(tc, _e140_half_texel, metal::float2(1.0) - _e140_half_texel), metal::level(0.0)); + metal::float4 _e141 = a; + a = _e141 + _e140; + metal::float4 _e143 = a; + return texture_sampleOutput { _e143 }; } diff --git a/tests/out/spv/image.spvasm b/tests/out/spv/image.spvasm index 708cd65f28..d873695aac 100644 --- a/tests/out/spv/image.spvasm +++ b/tests/out/spv/image.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 518 +; Bound: 529 OpCapability Shader OpCapability Image1D OpCapability Sampled1D @@ -14,15 +14,15 @@ OpEntryPoint GLCompute %169 "depth_load" %167 OpEntryPoint Vertex %189 "queries" %187 OpEntryPoint Vertex %241 "levels_queries" %240 OpEntryPoint Fragment %270 "texture_sample" %269 -OpEntryPoint Fragment %417 "texture_sample_comparison" %415 -OpEntryPoint Fragment %473 "gather" %472 -OpEntryPoint Fragment %507 "depth_no_comparison" %506 +OpEntryPoint Fragment %429 "texture_sample_comparison" %427 +OpEntryPoint Fragment %484 "gather" %483 +OpEntryPoint Fragment %518 "depth_no_comparison" %517 OpExecutionMode %78 LocalSize 16 1 1 OpExecutionMode %169 LocalSize 16 1 1 OpExecutionMode %270 OriginUpperLeft -OpExecutionMode %417 OriginUpperLeft -OpExecutionMode %473 OriginUpperLeft -OpExecutionMode %507 OriginUpperLeft +OpExecutionMode %429 OriginUpperLeft +OpExecutionMode %484 OriginUpperLeft +OpExecutionMode %518 OriginUpperLeft OpName %31 "image_mipmapped_src" OpName %33 "image_multisampled_src" OpName %35 "image_depth_multisampled_src" @@ -53,10 +53,10 @@ OpName %189 "queries" OpName %241 "levels_queries" OpName %270 "texture_sample" OpName %284 "a" -OpName %417 "texture_sample_comparison" -OpName %422 "a" -OpName %473 "gather" -OpName %507 "depth_no_comparison" +OpName %429 "texture_sample_comparison" +OpName %434 "a" +OpName %484 "gather" +OpName %518 "depth_no_comparison" OpDecorate %31 DescriptorSet 0 OpDecorate %31 Binding 0 OpDecorate %33 DescriptorSet 0 @@ -109,9 +109,9 @@ OpDecorate %167 BuiltIn LocalInvocationId OpDecorate %187 BuiltIn Position OpDecorate %240 BuiltIn Position OpDecorate %269 Location 0 -OpDecorate %415 Location 0 -OpDecorate %472 Location 0 -OpDecorate %506 Location 0 +OpDecorate %427 Location 0 +OpDecorate %483 Location 0 +OpDecorate %517 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 0 %3 = OpTypeImage %4 2D 0 0 0 1 Unknown @@ -213,21 +213,23 @@ OpDecorate %506 Location 0 %294 = OpTypeSampledImage %16 %315 = OpTypeSampledImage %18 %376 = OpTypeSampledImage %20 -%416 = OpTypePointer Output %7 -%415 = OpVariable %416 Output -%423 = OpTypePointer Function %7 -%424 = OpConstantNull %7 -%426 = OpTypeSampledImage %25 -%431 = OpTypeSampledImage %26 -%444 = OpTypeSampledImage %27 -%451 = OpConstant %7 0.0 -%472 = OpVariable %188 Output -%483 = OpConstant %4 1 -%486 = OpConstant %4 3 -%491 = OpTypeSampledImage %3 -%494 = OpTypeVector %14 4 -%495 = OpTypeSampledImage %17 -%506 = OpVariable %188 Output +%416 = OpConstant %7 1.0 +%417 = OpConstantComposite %277 %416 %416 +%422 = OpConstant %7 0.0 +%428 = OpTypePointer Output %7 +%427 = OpVariable %428 Output +%435 = OpTypePointer Function %7 +%436 = OpConstantNull %7 +%438 = OpTypeSampledImage %25 +%443 = OpTypeSampledImage %26 +%456 = OpTypeSampledImage %27 +%483 = OpVariable %188 Output +%494 = OpConstant %4 1 +%497 = OpConstant %4 3 +%502 = OpTypeSampledImage %3 +%505 = OpTypeVector %14 4 +%506 = OpTypeSampledImage %17 +%517 = OpVariable %188 Output %78 = OpFunction %2 None %79 %74 = OpLabel %77 = OpLoad %12 %75 @@ -577,116 +579,126 @@ OpStore %284 %406 %411 = OpLoad %23 %284 %412 = OpFAdd %23 %411 %410 OpStore %284 %412 -%413 = OpLoad %23 %284 -OpStore %269 %413 +%413 = OpImageQuerySizeLod %13 %272 %198 +%414 = OpConvertSToF %277 %413 +%415 = OpFDiv %277 %278 %414 +%418 = OpFSub %277 %417 %415 +%419 = OpExtInst %277 %1 FClamp %278 %415 %418 +%420 = OpSampledImage %294 %272 %275 +%421 = OpImageSampleExplicitLod %23 %420 %419 Lod %422 +%423 = OpLoad %23 %284 +%424 = OpFAdd %23 %423 %421 +OpStore %284 %424 +%425 = OpLoad %23 %284 +OpStore %269 %425 OpReturn OpFunctionEnd -%417 = OpFunction %2 None %79 -%414 = OpLabel -%422 = OpVariable %423 Function %424 -%418 = OpLoad %24 %66 -%419 = OpLoad %25 %68 -%420 = OpLoad %26 %70 -%421 = OpLoad %27 %72 -OpBranch %425 -%425 = OpLabel -%427 = OpSampledImage %426 %419 %418 -%428 = OpImageSampleDrefImplicitLod %7 %427 %278 %276 -%429 = OpLoad %7 %422 -%430 = OpFAdd %7 %429 %428 -OpStore %422 %430 -%432 = OpConvertUToF %7 %198 -%433 = OpCompositeConstruct %279 %278 %432 -%434 = OpSampledImage %431 %420 %418 -%435 = OpImageSampleDrefImplicitLod %7 %434 %433 %276 -%436 = OpLoad %7 %422 -%437 = OpFAdd %7 %436 %435 -OpStore %422 %437 -%438 = OpConvertSToF %7 %283 -%439 = OpCompositeConstruct %279 %278 %438 -%440 = OpSampledImage %431 %420 %418 -%441 = OpImageSampleDrefImplicitLod %7 %440 %439 %276 -%442 = OpLoad %7 %422 -%443 = OpFAdd %7 %442 %441 -OpStore %422 %443 -%445 = OpSampledImage %444 %421 %418 -%446 = OpImageSampleDrefImplicitLod %7 %445 %280 %276 -%447 = OpLoad %7 %422 -%448 = OpFAdd %7 %447 %446 -OpStore %422 %448 -%449 = OpSampledImage %426 %419 %418 -%450 = OpImageSampleDrefExplicitLod %7 %449 %278 %276 Lod %451 -%452 = OpLoad %7 %422 -%453 = OpFAdd %7 %452 %450 -OpStore %422 %453 -%454 = OpConvertUToF %7 %198 -%455 = OpCompositeConstruct %279 %278 %454 -%456 = OpSampledImage %431 %420 %418 -%457 = OpImageSampleDrefExplicitLod %7 %456 %455 %276 Lod %451 -%458 = OpLoad %7 %422 -%459 = OpFAdd %7 %458 %457 -OpStore %422 %459 -%460 = OpConvertSToF %7 %283 -%461 = OpCompositeConstruct %279 %278 %460 -%462 = OpSampledImage %431 %420 %418 -%463 = OpImageSampleDrefExplicitLod %7 %462 %461 %276 Lod %451 -%464 = OpLoad %7 %422 -%465 = OpFAdd %7 %464 %463 -OpStore %422 %465 -%466 = OpSampledImage %444 %421 %418 -%467 = OpImageSampleDrefExplicitLod %7 %466 %280 %276 Lod %451 -%468 = OpLoad %7 %422 -%469 = OpFAdd %7 %468 %467 -OpStore %422 %469 -%470 = OpLoad %7 %422 -OpStore %415 %470 +%429 = OpFunction %2 None %79 +%426 = OpLabel +%434 = OpVariable %435 Function %436 +%430 = OpLoad %24 %66 +%431 = OpLoad %25 %68 +%432 = OpLoad %26 %70 +%433 = OpLoad %27 %72 +OpBranch %437 +%437 = OpLabel +%439 = OpSampledImage %438 %431 %430 +%440 = OpImageSampleDrefImplicitLod %7 %439 %278 %276 +%441 = OpLoad %7 %434 +%442 = OpFAdd %7 %441 %440 +OpStore %434 %442 +%444 = OpConvertUToF %7 %198 +%445 = OpCompositeConstruct %279 %278 %444 +%446 = OpSampledImage %443 %432 %430 +%447 = OpImageSampleDrefImplicitLod %7 %446 %445 %276 +%448 = OpLoad %7 %434 +%449 = OpFAdd %7 %448 %447 +OpStore %434 %449 +%450 = OpConvertSToF %7 %283 +%451 = OpCompositeConstruct %279 %278 %450 +%452 = OpSampledImage %443 %432 %430 +%453 = OpImageSampleDrefImplicitLod %7 %452 %451 %276 +%454 = OpLoad %7 %434 +%455 = OpFAdd %7 %454 %453 +OpStore %434 %455 +%457 = OpSampledImage %456 %433 %430 +%458 = OpImageSampleDrefImplicitLod %7 %457 %280 %276 +%459 = OpLoad %7 %434 +%460 = OpFAdd %7 %459 %458 +OpStore %434 %460 +%461 = OpSampledImage %438 %431 %430 +%462 = OpImageSampleDrefExplicitLod %7 %461 %278 %276 Lod %422 +%463 = OpLoad %7 %434 +%464 = OpFAdd %7 %463 %462 +OpStore %434 %464 +%465 = OpConvertUToF %7 %198 +%466 = OpCompositeConstruct %279 %278 %465 +%467 = OpSampledImage %443 %432 %430 +%468 = OpImageSampleDrefExplicitLod %7 %467 %466 %276 Lod %422 +%469 = OpLoad %7 %434 +%470 = OpFAdd %7 %469 %468 +OpStore %434 %470 +%471 = OpConvertSToF %7 %283 +%472 = OpCompositeConstruct %279 %278 %471 +%473 = OpSampledImage %443 %432 %430 +%474 = OpImageSampleDrefExplicitLod %7 %473 %472 %276 Lod %422 +%475 = OpLoad %7 %434 +%476 = OpFAdd %7 %475 %474 +OpStore %434 %476 +%477 = OpSampledImage %456 %433 %430 +%478 = OpImageSampleDrefExplicitLod %7 %477 %280 %276 Lod %422 +%479 = OpLoad %7 %434 +%480 = OpFAdd %7 %479 %478 +OpStore %434 %480 +%481 = OpLoad %7 %434 +OpStore %427 %481 OpReturn OpFunctionEnd -%473 = OpFunction %2 None %79 -%471 = OpLabel -%474 = OpLoad %16 %49 -%475 = OpLoad %3 %51 -%476 = OpLoad %17 %52 -%477 = OpLoad %24 %64 -%478 = OpLoad %24 %66 -%479 = OpLoad %25 %68 -OpBranch %480 -%480 = OpLabel -%481 = OpSampledImage %294 %474 %477 -%482 = OpImageGather %23 %481 %278 %483 -%484 = OpSampledImage %294 %474 %477 -%485 = OpImageGather %23 %484 %278 %486 ConstOffset %30 -%487 = OpSampledImage %426 %479 %478 -%488 = OpImageDrefGather %23 %487 %278 %276 -%489 = OpSampledImage %426 %479 %478 -%490 = OpImageDrefGather %23 %489 %278 %276 ConstOffset %30 -%492 = OpSampledImage %491 %475 %477 -%493 = OpImageGather %98 %492 %278 %198 -%496 = OpSampledImage %495 %476 %477 -%497 = OpImageGather %494 %496 %278 %198 -%498 = OpConvertUToF %23 %493 -%499 = OpConvertSToF %23 %497 -%500 = OpFAdd %23 %498 %499 -%501 = OpFAdd %23 %482 %485 -%502 = OpFAdd %23 %501 %488 -%503 = OpFAdd %23 %502 %490 -%504 = OpFAdd %23 %503 %500 -OpStore %472 %504 +%484 = OpFunction %2 None %79 +%482 = OpLabel +%485 = OpLoad %16 %49 +%486 = OpLoad %3 %51 +%487 = OpLoad %17 %52 +%488 = OpLoad %24 %64 +%489 = OpLoad %24 %66 +%490 = OpLoad %25 %68 +OpBranch %491 +%491 = OpLabel +%492 = OpSampledImage %294 %485 %488 +%493 = OpImageGather %23 %492 %278 %494 +%495 = OpSampledImage %294 %485 %488 +%496 = OpImageGather %23 %495 %278 %497 ConstOffset %30 +%498 = OpSampledImage %438 %490 %489 +%499 = OpImageDrefGather %23 %498 %278 %276 +%500 = OpSampledImage %438 %490 %489 +%501 = OpImageDrefGather %23 %500 %278 %276 ConstOffset %30 +%503 = OpSampledImage %502 %486 %488 +%504 = OpImageGather %98 %503 %278 %198 +%507 = OpSampledImage %506 %487 %488 +%508 = OpImageGather %505 %507 %278 %198 +%509 = OpConvertUToF %23 %504 +%510 = OpConvertSToF %23 %508 +%511 = OpFAdd %23 %509 %510 +%512 = OpFAdd %23 %493 %496 +%513 = OpFAdd %23 %512 %499 +%514 = OpFAdd %23 %513 %501 +%515 = OpFAdd %23 %514 %511 +OpStore %483 %515 OpReturn OpFunctionEnd -%507 = OpFunction %2 None %79 -%505 = OpLabel -%508 = OpLoad %24 %64 -%509 = OpLoad %25 %68 -OpBranch %510 -%510 = OpLabel -%511 = OpSampledImage %426 %509 %508 -%512 = OpImageSampleImplicitLod %23 %511 %278 -%513 = OpCompositeExtract %7 %512 0 -%514 = OpSampledImage %426 %509 %508 -%515 = OpImageGather %23 %514 %278 %198 -%516 = OpCompositeConstruct %23 %513 %513 %513 %513 -%517 = OpFAdd %23 %516 %515 -OpStore %506 %517 +%518 = OpFunction %2 None %79 +%516 = OpLabel +%519 = OpLoad %24 %64 +%520 = OpLoad %25 %68 +OpBranch %521 +%521 = OpLabel +%522 = OpSampledImage %438 %520 %519 +%523 = OpImageSampleImplicitLod %23 %522 %278 +%524 = OpCompositeExtract %7 %523 0 +%525 = OpSampledImage %438 %520 %519 +%526 = OpImageGather %23 %525 %278 %198 +%527 = OpCompositeConstruct %23 %524 %524 %524 %524 +%528 = OpFAdd %23 %527 %526 +OpStore %517 %528 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/image.wgsl b/tests/out/wgsl/image.wgsl index 062d377139..2ea1130c9e 100644 --- a/tests/out/wgsl/image.wgsl +++ b/tests/out/wgsl/image.wgsl @@ -178,8 +178,11 @@ fn texture_sample() -> @location(0) vec4 { let _e135 = textureSampleBias(image_cube_array, sampler_reg, tc3_, 0, 2.0); let _e136 = a; a = (_e136 + _e135); - let _e138 = a; - return _e138; + let _e140 = textureSampleBaseClampToEdge(image_2d, sampler_reg, tc); + let _e141 = a; + a = (_e141 + _e140); + let _e143 = a; + return _e143; } @fragment From a152303e59567038ed9037bbb13e6187396ee4b3 Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Mon, 16 Oct 2023 14:08:32 -0400 Subject: [PATCH 2/3] Replace Base with clamp_to_edge --- src/back/dot/mod.rs | 5 ++- src/back/glsl/mod.rs | 15 +++----- src/back/hlsl/writer.rs | 15 ++++---- src/back/msl/writer.rs | 15 ++++---- src/back/spv/block.rs | 2 ++ src/back/spv/image.rs | 5 +-- src/back/wgsl/writer.rs | 7 ++-- src/compact/expressions.rs | 6 ++-- src/front/glsl/builtins.rs | 1 + src/front/spv/image.rs | 1 + src/front/wgsl/lower/mod.rs | 8 ++++- src/lib.rs | 2 +- src/proc/mod.rs | 2 +- src/valid/analyzer.rs | 3 +- src/valid/expression.rs | 61 +++++++++++++++++---------------- src/valid/handles.rs | 5 ++- tests/out/ir/shadow.compact.ron | 1 + tests/out/ir/shadow.ron | 1 + 18 files changed, 86 insertions(+), 69 deletions(-) diff --git a/src/back/dot/mod.rs b/src/back/dot/mod.rs index eed7c5bd99..dbeb1fa71c 100644 --- a/src/back/dot/mod.rs +++ b/src/back/dot/mod.rs @@ -452,6 +452,7 @@ fn write_function_expressions( offset: _, level, depth_ref, + clamp_to_edge: _, } => { edges.insert("image", image); edges.insert("sampler", sampler); @@ -460,9 +461,7 @@ fn write_function_expressions( edges.insert("array_index", expr); } match level { - crate::SampleLevel::Auto - | crate::SampleLevel::Zero - | crate::SampleLevel::Base => {} + crate::SampleLevel::Auto | crate::SampleLevel::Zero => {} crate::SampleLevel::Exact(expr) => { edges.insert("level", expr); } diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 267d484ea2..bfeef7d88c 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -1862,7 +1862,8 @@ impl<'a, W: Write> Writer<'a, W> { // If we are going to write a `textureSampleBaseClampToEdge` next, // precompute the half-texel before clamping the coordinates. if let crate::Expression::ImageSample { - level: crate::SampleLevel::Base, + level: crate::SampleLevel::Zero, + clamp_to_edge: true, image, .. } = ctx.expressions[handle] @@ -2485,6 +2486,7 @@ impl<'a, W: Write> Writer<'a, W> { offset, level, depth_ref, + clamp_to_edge, } => { let dim = match *ctx.info[image].ty.inner_with(&self.module.types) { TypeInner::Image { dim, .. } => dim, @@ -2506,11 +2508,6 @@ impl<'a, W: Write> Writer<'a, W> { ))) } crate::SampleLevel::Auto => {} - crate::SampleLevel::Base => { - unreachable!( - "textureSampleBaseClampToEdge should not have passed validation" - ) - } } } @@ -2537,7 +2534,6 @@ impl<'a, W: Write> Writer<'a, W> { } } crate::SampleLevel::Gradient { .. } => "textureGrad", - crate::SampleLevel::Base => "textureLod", }; let offset_name = match offset { Some(_) => "Offset", @@ -2570,7 +2566,7 @@ impl<'a, W: Write> Writer<'a, W> { let tex_1d_hack = dim == crate::ImageDimension::D1 && self.options.version.is_es(); let is_vec = tex_1d_hack || coord_dim != 1; - if level == crate::SampleLevel::Base { + if clamp_to_edge { // clamp the coordinates to [ half_texel, 1 - half_texel ] write!(self.out, "clamp(")?; self.write_expr(coordinate, ctx)?; @@ -2653,9 +2649,6 @@ impl<'a, W: Write> Writer<'a, W> { self.write_expr(y, ctx)?; } } - crate::SampleLevel::Base => { - write!(self.out, ", 0.0")?; - } } if let Some(constant) = offset { diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index aed1e4d53f..64980ce3f3 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -1346,7 +1346,8 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // If we are going to write a `textureSampleBaseClampToEdge` next, // precompute the half-texel before clamping the coordinates. if let crate::Expression::ImageSample { - level: crate::SampleLevel::Base, + level: crate::SampleLevel::Zero, + clamp_to_edge: true, image, .. } = func_ctx.expressions[handle] @@ -2394,6 +2395,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { offset, level, depth_ref, + clamp_to_edge, } => { use crate::SampleLevel as Sl; const COMPONENTS: [&str; 4] = ["", "Green", "Blue", "Alpha"]; @@ -2407,9 +2409,10 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { None => "", }; let level_str = match level { + Sl::Zero if clamp_to_edge => "Level", Sl::Zero if gather.is_none() => "LevelZero", Sl::Auto | Sl::Zero => "", - Sl::Exact(_) | Sl::Base => "Level", + Sl::Exact(_) => "Level", Sl::Bias(_) => "Bias", Sl::Gradient { .. } => "Grad", }; @@ -2418,7 +2421,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { write!(self.out, ".{base_str}{cmp_str}{component_str}{level_str}(")?; self.write_expr(module, sampler, func_ctx)?; write!(self.out, ", ")?; - if level == Sl::Base { + if clamp_to_edge { // clamp the coordinates to [ half_texel, 1 - half_texel ] write!(self.out, "clamp(")?; self.write_expr(module, coordinate, func_ctx)?; @@ -2445,6 +2448,9 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } match level { + Sl::Zero if clamp_to_edge => { + write!(self.out, ", 0.0")?; + } Sl::Auto | Sl::Zero => {} Sl::Exact(expr) => { write!(self.out, ", ")?; @@ -2460,9 +2466,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { write!(self.out, ", ")?; self.write_expr(module, y, func_ctx)?; } - Sl::Base => { - write!(self.out, ", 0.0")?; - } } if let Some(offset) = offset { diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index d41cc23f25..c5a81bb931 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -743,11 +743,15 @@ impl Writer { &mut self, image: Handle, level: crate::SampleLevel, + clamp_to_edge: bool, context: &ExpressionContext, ) -> BackendResult { let has_levels = context.image_needs_lod(image); match level { crate::SampleLevel::Auto => {} + crate::SampleLevel::Zero if clamp_to_edge => { + write!(self.out, ", {NAMESPACE}::level(0.0)")?; + } crate::SampleLevel::Zero => { //TODO: do we support Zero on `Sampled` image classes? } @@ -771,9 +775,6 @@ impl Writer { self.put_expression(y, context, true)?; write!(self.out, ")")?; } - crate::SampleLevel::Base => { - write!(self.out, ", {NAMESPACE}::level(0.0)")?; - } } Ok(()) } @@ -1467,6 +1468,7 @@ impl Writer { offset, level, depth_ref, + clamp_to_edge, } => { let main_op = match gather { Some(_) => "gather", @@ -1481,7 +1483,7 @@ impl Writer { self.put_expression(sampler, context, true)?; write!(self.out, ", ")?; - if level == crate::SampleLevel::Base { + if clamp_to_edge { // clamp the coordinates to [ half_texel, 1 - half_texel ] write!(self.out, "{NAMESPACE}::clamp(")?; self.put_expression(coordinate, context, true)?; @@ -1505,7 +1507,7 @@ impl Writer { self.put_expression(dref, context, true)?; } - self.put_image_sample_level(image, level, context)?; + self.put_image_sample_level(image, level, clamp_to_edge, context)?; if let Some(offset) = offset { write!(self.out, ", ")?; @@ -2671,7 +2673,8 @@ impl Writer { // If we are going to write a `textureSampleBaseClampToEdge` next, // precompute the half-texel before clamping the coordinates. if let crate::Expression::ImageSample { - level: crate::SampleLevel::Base, + level: crate::SampleLevel::Zero, + clamp_to_edge: true, image, .. } = context.expression.function.expressions[handle] diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index 4dba7ea0ca..5d7ee9d1af 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -1312,6 +1312,7 @@ impl<'w> BlockContext<'w> { offset, level, depth_ref, + clamp_to_edge, } => self.write_image_sample( result_type_id, image, @@ -1322,6 +1323,7 @@ impl<'w> BlockContext<'w> { offset, level, depth_ref, + clamp_to_edge, block, )?, crate::Expression::Select { diff --git a/src/back/spv/image.rs b/src/back/spv/image.rs index 003b821835..9e9df15859 100644 --- a/src/back/spv/image.rs +++ b/src/back/spv/image.rs @@ -915,6 +915,7 @@ impl<'w> BlockContext<'w> { offset: Option>, level: crate::SampleLevel, depth_ref: Option>, + clamp_to_edge: bool, block: &mut Block, ) -> Result { use super::instructions::SampleLod; @@ -947,7 +948,7 @@ impl<'w> BlockContext<'w> { self.get_type_id(LookupType::Local(LocalType::SampledImage { image_type_id })); let sampler_id = self.get_handle_id(sampler); - let coordinates_id = if level == crate::SampleLevel::Base { + let coordinates_id = if clamp_to_edge { self.write_clamped_image_coordinates(image_id, coordinate, block)? } else { self.write_image_coordinates(coordinate, array_index, block)? @@ -983,7 +984,7 @@ impl<'w> BlockContext<'w> { } inst } - (crate::SampleLevel::Zero | crate::SampleLevel::Base, None) => { + (crate::SampleLevel::Zero, None) => { let mut inst = Instruction::image_sample( sample_result_type_id, id, diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index cb84309fd2..96f954fce1 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -1245,6 +1245,7 @@ impl Writer { offset, level, depth_ref, + clamp_to_edge, } => { use crate::SampleLevel as Sl; @@ -1254,10 +1255,10 @@ impl Writer { }; let suffix_level = match level { Sl::Auto => "", + Sl::Zero if clamp_to_edge => "BaseClampToEdge", Sl::Zero | Sl::Exact(_) => "Level", Sl::Bias(_) => "Bias", Sl::Gradient { .. } => "Grad", - Sl::Base => "BaseClampToEdge", }; write!(self.out, "textureSample{suffix_cmp}{suffix_level}(")?; @@ -1278,7 +1279,8 @@ impl Writer { } match level { - Sl::Auto | Sl::Base => {} + Sl::Auto => {} + Sl::Zero if clamp_to_edge => {} Sl::Zero => { // Level 0 is implied for depth comparison if depth_ref.is_none() { @@ -1318,6 +1320,7 @@ impl Writer { offset, level: _, depth_ref, + clamp_to_edge: _, } => { let suffix_cmp = match depth_ref { Some(_) => "Compare", diff --git a/src/compact/expressions.rs b/src/compact/expressions.rs index b008680a6e..7dc93a8121 100644 --- a/src/compact/expressions.rs +++ b/src/compact/expressions.rs @@ -86,6 +86,7 @@ impl<'tracer> ExpressionTracer<'tracer> { offset, ref level, depth_ref, + clamp_to_edge: _, } => { work_list.push(image); work_list.push(sampler); @@ -96,7 +97,7 @@ impl<'tracer> ExpressionTracer<'tracer> { } use crate::SampleLevel as Sl; match *level { - Sl::Auto | Sl::Zero | Sl::Base => {} + Sl::Auto | Sl::Zero => {} Sl::Exact(expr) | Sl::Bias(expr) => work_list.push(expr), Sl::Gradient { x, y } => work_list.extend([x, y]), } @@ -266,6 +267,7 @@ impl ModuleMap { ref mut offset, ref mut level, ref mut depth_ref, + clamp_to_edge: _, } => { adjust(image); adjust(sampler); @@ -366,7 +368,7 @@ impl ModuleMap { use crate::SampleLevel as Sl; match *level { - Sl::Auto | Sl::Zero | Sl::Base => {} + Sl::Auto | Sl::Zero => {} Sl::Exact(ref mut expr) => adjust(expr), Sl::Bias(ref mut expr) => adjust(expr), Sl::Gradient { diff --git a/src/front/glsl/builtins.rs b/src/front/glsl/builtins.rs index 8d57b66da2..e203860903 100644 --- a/src/front/glsl/builtins.rs +++ b/src/front/glsl/builtins.rs @@ -2167,6 +2167,7 @@ fn texture_call( offset, level, depth_ref: comps.depth_ref, + clamp_to_edge: false, }, meta, )?) diff --git a/src/front/spv/image.rs b/src/front/spv/image.rs index ee58c7ba14..6289f0123b 100644 --- a/src/front/spv/image.rs +++ b/src/front/spv/image.rs @@ -647,6 +647,7 @@ impl> super::Frontend { offset, level, depth_ref, + clamp_to_edge: false, }; self.lookup_expression.insert( result_id, diff --git a/src/front/wgsl/lower/mod.rs b/src/front/wgsl/lower/mod.rs index 0627427ecd..b9d6e3c61e 100644 --- a/src/front/wgsl/lower/mod.rs +++ b/src/front/wgsl/lower/mod.rs @@ -2399,6 +2399,8 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { .then(|| self.expression(args.next()?, ctx.reborrow())) .transpose()?; + let mut clamp_to_edge = false; + let (level, depth_ref) = match fun { Texture::Gather => (crate::SampleLevel::Zero, None), Texture::GatherCompare => { @@ -2428,7 +2430,10 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let level = self.expression(args.next()?, ctx.reborrow())?; (crate::SampleLevel::Exact(level), None) } - Texture::SampleBaseClampToEdge => (crate::SampleLevel::Base, None), + Texture::SampleBaseClampToEdge => { + clamp_to_edge = true; + (crate::SampleLevel::Zero, None) + } }; let offset = args @@ -2448,6 +2453,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { offset, level, depth_ref, + clamp_to_edge, }) } diff --git a/src/lib.rs b/src/lib.rs index 69dfc7d980..3b2b0233fc 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -1221,7 +1221,6 @@ pub enum SampleLevel { x: Handle, y: Handle, }, - Base, } /// Type of an image query. @@ -1415,6 +1414,7 @@ pub enum Expression { offset: Option>, level: SampleLevel, depth_ref: Option>, + clamp_to_edge: bool, }, /// Load a texel from an image. diff --git a/src/proc/mod.rs b/src/proc/mod.rs index 9b84779a92..d50f165a56 100644 --- a/src/proc/mod.rs +++ b/src/proc/mod.rs @@ -508,7 +508,7 @@ impl crate::SampleLevel { pub const fn implicit_derivatives(&self) -> bool { match *self { Self::Auto | Self::Bias(_) => true, - Self::Zero | Self::Exact(_) | Self::Gradient { .. } | Self::Base => false, + Self::Zero | Self::Exact(_) | Self::Gradient { .. } => false, } } } diff --git a/src/valid/analyzer.rs b/src/valid/analyzer.rs index 0fa474da5e..fe8cb5e3e9 100644 --- a/src/valid/analyzer.rs +++ b/src/valid/analyzer.rs @@ -600,6 +600,7 @@ impl FunctionInfo { offset: _, level, depth_ref, + clamp_to_edge: _, } => { let image_storage = GlobalOrArgument::from_expression(expression_arena, image)?; let sampler_storage = GlobalOrArgument::from_expression(expression_arena, sampler)?; @@ -619,7 +620,7 @@ impl FunctionInfo { // "nur" == "Non-Uniform Result" let array_nur = array_index.and_then(|h| self.add_ref(h)); let level_nur = match level { - Sl::Auto | Sl::Zero | Sl::Base => None, + Sl::Auto | Sl::Zero => None, Sl::Exact(h) | Sl::Bias(h) => self.add_ref(h), Sl::Gradient { x, y } => self.add_ref(x).or(self.add_ref(y)), }; diff --git a/src/valid/expression.rs b/src/valid/expression.rs index e47a7872d6..3cad033021 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -390,6 +390,7 @@ impl super::Validator { offset, level, depth_ref, + clamp_to_edge, } => { // check the validity of expressions let image_ty = Self::global_var_ty(module, function, image)?; @@ -515,6 +516,36 @@ impl super::Validator { // check level properties match level { crate::SampleLevel::Auto => ShaderStages::FRAGMENT, + crate::SampleLevel::Zero if clamp_to_edge => { + // TODO: handle external textures + match module.types[image_ty].inner { + Ti::Image { + dim: crate::ImageDimension::D2, + arrayed: false, + class: + crate::ImageClass::Sampled { + kind: Sk::Float, + multi: false, + }, + } => {} + _ => { + return Err(ExpressionError::InvalidSampleBaseClampToEdgeImageType) + } + } + match resolver[coordinate] { + Ti::Vector { + size: crate::VectorSize::Bi, + kind: Sk::Float, + .. + } => {} + _ => { + return Err( + ExpressionError::InvalidSampleBaseClampToEdgeCoordinateType, + ) + } + } + ShaderStages::all() + } crate::SampleLevel::Zero => ShaderStages::all(), crate::SampleLevel::Exact(expr) => { match resolver[expr] { @@ -563,36 +594,6 @@ impl super::Validator { } ShaderStages::all() } - crate::SampleLevel::Base => { - // TODO: handle external textures - match module.types[image_ty].inner { - Ti::Image { - dim: crate::ImageDimension::D2, - arrayed: false, - class: - crate::ImageClass::Sampled { - kind: Sk::Float, - multi: false, - }, - } => {} - _ => { - return Err(ExpressionError::InvalidSampleBaseClampToEdgeImageType) - } - } - match resolver[coordinate] { - Ti::Vector { - size: crate::VectorSize::Bi, - kind: Sk::Float, - .. - } => {} - _ => { - return Err( - ExpressionError::InvalidSampleBaseClampToEdgeCoordinateType, - ) - } - } - ShaderStages::all() - } } } E::ImageLoad { diff --git a/src/valid/handles.rs b/src/valid/handles.rs index 49b32cc8e5..0894061265 100644 --- a/src/valid/handles.rs +++ b/src/valid/handles.rs @@ -290,6 +290,7 @@ impl super::Validator { offset, level, depth_ref, + clamp_to_edge: _, } => { if let Some(offset) = offset { validate_const_expr(offset)?; @@ -302,9 +303,7 @@ impl super::Validator { .check_dep_opt(array_index)?; match level { - crate::SampleLevel::Auto - | crate::SampleLevel::Zero - | crate::SampleLevel::Base => (), + crate::SampleLevel::Auto | crate::SampleLevel::Zero => (), crate::SampleLevel::Exact(expr) => { handle.check_dep(expr)?; } diff --git a/tests/out/ir/shadow.compact.ron b/tests/out/ir/shadow.compact.ron index 9ca6799c21..c379459028 100644 --- a/tests/out/ir/shadow.compact.ron +++ b/tests/out/ir/shadow.compact.ron @@ -522,6 +522,7 @@ offset: None, level: Zero, depth_ref: Some(29), + clamp_to_edge: false, ), ], named_expressions: {}, diff --git a/tests/out/ir/shadow.ron b/tests/out/ir/shadow.ron index 07bf66fcc8..56e91e6f09 100644 --- a/tests/out/ir/shadow.ron +++ b/tests/out/ir/shadow.ron @@ -793,6 +793,7 @@ offset: None, level: Zero, depth_ref: Some(65), + clamp_to_edge: false, ), ], named_expressions: {}, From d546786f4f015f5224563b5212706bbc5570fbb2 Mon Sep 17 00:00:00 2001 From: Evan Mark Hopkins Date: Tue, 17 Oct 2023 00:32:13 -0400 Subject: [PATCH 3/3] Remove extra level checks --- src/back/glsl/mod.rs | 1 - src/back/hlsl/writer.rs | 1 - src/back/msl/writer.rs | 1 - 3 files changed, 3 deletions(-) diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index bfeef7d88c..655e831422 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -1862,7 +1862,6 @@ impl<'a, W: Write> Writer<'a, W> { // If we are going to write a `textureSampleBaseClampToEdge` next, // precompute the half-texel before clamping the coordinates. if let crate::Expression::ImageSample { - level: crate::SampleLevel::Zero, clamp_to_edge: true, image, .. diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index 64980ce3f3..81253356b2 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -1346,7 +1346,6 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { // If we are going to write a `textureSampleBaseClampToEdge` next, // precompute the half-texel before clamping the coordinates. if let crate::Expression::ImageSample { - level: crate::SampleLevel::Zero, clamp_to_edge: true, image, .. diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index c5a81bb931..d4de0ac74e 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -2673,7 +2673,6 @@ impl Writer { // If we are going to write a `textureSampleBaseClampToEdge` next, // precompute the half-texel before clamping the coordinates. if let crate::Expression::ImageSample { - level: crate::SampleLevel::Zero, clamp_to_edge: true, image, ..