diff --git a/src/back/dot/mod.rs b/src/back/dot/mod.rs index 1556371df1..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,8 +461,7 @@ 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::Exact(expr) => { edges.insert("level", expr); } diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 77e206d3c6..655e831422 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 { + clamp_to_edge: true, + 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)?; @@ -2470,6 +2485,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, @@ -2548,24 +2564,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 clamp_to_edge { + // 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) { @@ -3480,6 +3510,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..81253356b2 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 { + clamp_to_edge: true, + 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)?; @@ -2379,6 +2394,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"]; @@ -2392,6 +2408,7 @@ 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(_) => "Level", @@ -2403,14 +2420,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 clamp_to_edge { + // 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, ", ")?; @@ -2418,6 +2447,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, ", ")?; @@ -3259,6 +3291,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..d4de0ac74e 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 @@ -739,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? } @@ -1460,6 +1468,7 @@ impl Writer { offset, level, depth_ref, + clamp_to_edge, } => { let main_op = match gather { Some(_) => "gather", @@ -1473,7 +1482,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 clamp_to_edge { + // 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)?; @@ -1483,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, ", ")?; @@ -2646,6 +2670,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 { + clamp_to_edge: true, + 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 +4225,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/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 1630dc0ddb..9e9df15859 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) => { @@ -826,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; @@ -858,9 +948,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 clamp_to_edge { + 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( diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index 36380952af..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,6 +1255,7 @@ 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", @@ -1278,6 +1280,7 @@ impl Writer { match level { Sl::Auto => {} + Sl::Zero if clamp_to_edge => {} Sl::Zero => { // Level 0 is implied for depth comparison if depth_ref.is_none() { @@ -1317,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 c1326e92be..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); @@ -266,6 +267,7 @@ impl ModuleMap { ref mut offset, ref mut level, ref mut depth_ref, + clamp_to_edge: _, } => { adjust(image); adjust(sampler); 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 c883732e8b..b9d6e3c61e 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, } } } @@ -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,6 +2430,10 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { let level = self.expression(args.next()?, ctx.reborrow())?; (crate::SampleLevel::Exact(level), None) } + Texture::SampleBaseClampToEdge => { + clamp_to_edge = true; + (crate::SampleLevel::Zero, None) + } }; let offset = args @@ -2447,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 9d70190421..3b2b0233fc 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -1414,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/valid/analyzer.rs b/src/valid/analyzer.rs index ff1db071c8..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)?; diff --git a/src/valid/expression.rs b/src/valid/expression.rs index 95225a3926..3cad033021 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")] @@ -386,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)?; @@ -511,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] { diff --git a/src/valid/handles.rs b/src/valid/handles.rs index c68ded074b..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)?; 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/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: {}, 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