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 ddef6a4c03..1a4cf51ad9 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 `SampleBaseClampToEdge` +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 `SampleBaseClampToEdge` 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)?; @@ -2482,6 +2497,9 @@ impl<'a, W: Write> Writer<'a, W> { ))) } crate::SampleLevel::Auto => {} + crate::SampleLevel::Base => { + unreachable!("SampleBaseClampToEdge should not have passed validation") + } } } @@ -2508,6 +2526,7 @@ impl<'a, W: Write> Writer<'a, W> { } } crate::SampleLevel::Gradient { .. } => "textureGrad", + crate::SampleLevel::Base => "textureLod", }; let offset_name = match offset { Some(_) => "Offset", @@ -2539,24 +2558,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) { @@ -2609,6 +2642,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 { @@ -3471,6 +3507,30 @@ impl<'a, W: Write> Writer<'a, W> { Ok(()) } + /// Helper function to write the local holding the half-texel + /// for use with `SampleBaseClampToEdge` + 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 e9809bd2d4..e04b70247d 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 `SampleBaseClampToEdge` +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 `SampleBaseClampToEdge` 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)?; @@ -2379,6 +2394,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { }; let level_str = match level { Sl::Zero if gather.is_none() => "LevelZero", + Sl::Base => "LevelZero", Sl::Auto | Sl::Zero => "", Sl::Exact(_) => "Level", Sl::Bias(_) => "Bias", @@ -2389,14 +2405,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, ", ")?; @@ -2404,7 +2432,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } match level { - Sl::Auto | Sl::Zero => {} + Sl::Auto | Sl::Zero | Sl::Base => {} Sl::Exact(expr) => { write!(self.out, ", ")?; self.write_expr(module, expr, func_ctx)?; @@ -3256,6 +3284,39 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { } Ok(()) } + + /// Helper function to write the locals holding the half-texel + /// for use with `SampleBaseClampToEdge` + 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, 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 77231a286d..3628333850 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 `SampleBaseClampToEdge` +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 @@ -766,6 +770,9 @@ impl Writer { self.put_expression(y, context, true)?; write!(self.out, ")")?; } + crate::SampleLevel::Base => { + write!(self.out, ", {NAMESPACE}::level(0.0)")?; + } } Ok(()) } @@ -1455,7 +1462,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)?; @@ -2628,6 +2650,17 @@ impl Writer { } }; + // If we are going to write a `SampleBaseClampToEdge` 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)?; @@ -4166,6 +4199,36 @@ impl Writer { } Ok(()) } + + /// Helper function to write the locals holding the half-texel + /// for use with `SampleBaseClampToEdge` + 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 b4d2fb9e16..2ddd07db82 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 `SampleBaseClampToEdge` + 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 428a413374..9ef4d00cde 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 4ccf559c4e..b65e95568d 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]), } @@ -367,7 +367,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 56d8709708..d58e502405 100644 --- a/src/front/wgsl/lower/mod.rs +++ b/src/front/wgsl/lower/mod.rs @@ -751,7 +751,7 @@ enum Texture { SampleCompareLevel, SampleGrad, SampleLevel, - // SampleBaseClampToEdge, + SampleBaseClampToEdge, } impl Texture { @@ -766,7 +766,7 @@ impl Texture { "textureSampleCompareLevel" => Self::SampleCompareLevel, "textureSampleGrad" => Self::SampleGrad, "textureSampleLevel" => Self::SampleLevel, - // "textureSampleBaseClampToEdge" => Some(Self::SampleBaseClampToEdge), + "textureSampleBaseClampToEdge" => Self::SampleBaseClampToEdge, _ => return None, }) } @@ -782,7 +782,7 @@ impl Texture { Self::SampleCompareLevel => 5, Self::SampleGrad => 6, Self::SampleLevel => 5, - // Self::SampleBaseClampToEdge => 3, + Self::SampleBaseClampToEdge => 3, } } } @@ -2327,6 +2327,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 e7abe1e1c5..e69d72ae88 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -1214,6 +1214,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 90fd9500ab..ad21385802 100644 --- a/src/proc/mod.rs +++ b/src/proc/mod.rs @@ -499,7 +499,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 9b21b7f732..d9ef78ce58 100644 --- a/src/valid/analyzer.rs +++ b/src/valid/analyzer.rs @@ -602,7 +602,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 3426bf008e..68660940db 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")] @@ -523,6 +527,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 da95f60842..befe3a5363 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 e284813a1b..40831b5da1 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, num); + float2 _expr140_half_texel = (0.5).xx / _expr140_dim; + float4 _expr140 = image_2d.SampleLevelZero(sampler_reg, clamp(tc, _expr140_half_texel, (1.0).xx - _expr140_half_texel)); + 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 1b6ef28879..e3d714d6ae 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: 523 +; Bound: 535 OpCapability Shader OpCapability Image1D OpCapability Sampled1D @@ -14,15 +14,15 @@ OpEntryPoint GLCompute %169 "depth_load" %167 OpEntryPoint Vertex %190 "queries" %188 OpEntryPoint Vertex %242 "levels_queries" %241 OpEntryPoint Fragment %274 "texture_sample" %273 -OpEntryPoint Fragment %421 "texture_sample_comparison" %419 -OpEntryPoint Fragment %476 "gather" %475 -OpEntryPoint Fragment %511 "depth_no_comparison" %510 +OpEntryPoint Fragment %434 "texture_sample_comparison" %432 +OpEntryPoint Fragment %488 "gather" %487 +OpEntryPoint Fragment %523 "depth_no_comparison" %522 OpExecutionMode %78 LocalSize 16 1 1 OpExecutionMode %169 LocalSize 16 1 1 OpExecutionMode %274 OriginUpperLeft -OpExecutionMode %421 OriginUpperLeft -OpExecutionMode %476 OriginUpperLeft -OpExecutionMode %511 OriginUpperLeft +OpExecutionMode %434 OriginUpperLeft +OpExecutionMode %488 OriginUpperLeft +OpExecutionMode %523 OriginUpperLeft OpName %31 "image_mipmapped_src" OpName %33 "image_multisampled_src" OpName %35 "image_depth_multisampled_src" @@ -53,10 +53,10 @@ OpName %190 "queries" OpName %242 "levels_queries" OpName %269 "a" OpName %274 "texture_sample" -OpName %415 "a" -OpName %421 "texture_sample_comparison" -OpName %476 "gather" -OpName %511 "depth_no_comparison" +OpName %428 "a" +OpName %434 "texture_sample_comparison" +OpName %488 "gather" +OpName %523 "depth_no_comparison" OpDecorate %31 DescriptorSet 0 OpDecorate %31 Binding 0 OpDecorate %33 DescriptorSet 0 @@ -109,9 +109,9 @@ OpDecorate %167 BuiltIn LocalInvocationId OpDecorate %188 BuiltIn Position OpDecorate %241 BuiltIn Position OpDecorate %273 Location 0 -OpDecorate %419 Location 0 -OpDecorate %475 Location 0 -OpDecorate %510 Location 0 +OpDecorate %432 Location 0 +OpDecorate %487 Location 0 +OpDecorate %522 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 0 %3 = OpTypeImage %4 2D 0 0 0 1 Unknown @@ -210,21 +210,24 @@ OpDecorate %510 Location 0 %295 = OpTypeSampledImage %16 %316 = OpTypeSampledImage %18 %377 = OpTypeSampledImage %20 -%416 = OpTypePointer Function %7 -%417 = OpConstantNull %7 -%420 = OpTypePointer Output %7 -%419 = OpVariable %420 Output -%429 = OpTypeSampledImage %25 -%434 = OpTypeSampledImage %26 -%447 = OpTypeSampledImage %27 -%454 = OpConstant %7 0.0 -%475 = OpVariable %189 Output -%487 = OpConstant %4 1 -%490 = OpConstant %4 3 -%495 = OpTypeSampledImage %3 -%498 = OpTypeVector %14 4 -%499 = OpTypeSampledImage %17 -%510 = OpVariable %189 Output +%416 = OpConstantComposite %285 %280 %280 +%418 = OpConstant %7 1.0 +%419 = OpConstantComposite %285 %418 %418 +%424 = OpConstant %7 0.0 +%429 = OpTypePointer Function %7 +%430 = OpConstantNull %7 +%433 = OpTypePointer Output %7 +%432 = OpVariable %433 Output +%442 = OpTypeSampledImage %25 +%447 = OpTypeSampledImage %26 +%460 = OpTypeSampledImage %27 +%487 = OpVariable %189 Output +%499 = OpConstant %4 1 +%502 = OpConstant %4 3 +%507 = OpTypeSampledImage %3 +%510 = OpTypeVector %14 4 +%511 = OpTypeSampledImage %17 +%522 = OpVariable %189 Output %78 = OpFunction %2 None %79 %74 = OpLabel %77 = OpLoad %12 %75 @@ -578,120 +581,130 @@ OpStore %269 %407 %412 = OpLoad %23 %269 %413 = OpFAdd %23 %412 %411 OpStore %269 %413 -%414 = OpLoad %23 %269 -OpStore %273 %414 +%414 = OpImageQuerySizeLod %13 %276 %199 +%415 = OpConvertSToF %285 %414 +%417 = OpFDiv %285 %416 %415 +%420 = OpFSub %285 %419 %417 +%421 = OpExtInst %285 %1 FClamp %286 %417 %420 +%422 = OpSampledImage %295 %276 %279 +%423 = OpImageSampleExplicitLod %23 %422 %421 Lod %424 +%425 = OpLoad %23 %269 +%426 = OpFAdd %23 %425 %423 +OpStore %269 %426 +%427 = OpLoad %23 %269 +OpStore %273 %427 OpReturn OpFunctionEnd -%421 = OpFunction %2 None %79 -%418 = OpLabel -%415 = OpVariable %416 Function %417 -%422 = OpLoad %24 %66 -%423 = OpLoad %25 %68 -%424 = OpLoad %26 %70 -%425 = OpLoad %27 %72 -OpBranch %426 -%426 = OpLabel -%427 = OpCompositeConstruct %285 %280 %280 -%428 = OpCompositeConstruct %287 %280 %280 %280 -%430 = OpSampledImage %429 %423 %422 -%431 = OpImageSampleDrefImplicitLod %7 %430 %427 %280 -%432 = OpLoad %7 %415 -%433 = OpFAdd %7 %432 %431 -OpStore %415 %433 -%435 = OpConvertUToF %7 %199 -%436 = OpCompositeConstruct %287 %427 %435 -%437 = OpSampledImage %434 %424 %422 -%438 = OpImageSampleDrefImplicitLod %7 %437 %436 %280 -%439 = OpLoad %7 %415 -%440 = OpFAdd %7 %439 %438 -OpStore %415 %440 -%441 = OpConvertSToF %7 %283 -%442 = OpCompositeConstruct %287 %427 %441 -%443 = OpSampledImage %434 %424 %422 -%444 = OpImageSampleDrefImplicitLod %7 %443 %442 %280 -%445 = OpLoad %7 %415 +%434 = OpFunction %2 None %79 +%431 = OpLabel +%428 = OpVariable %429 Function %430 +%435 = OpLoad %24 %66 +%436 = OpLoad %25 %68 +%437 = OpLoad %26 %70 +%438 = OpLoad %27 %72 +OpBranch %439 +%439 = OpLabel +%440 = OpCompositeConstruct %285 %280 %280 +%441 = OpCompositeConstruct %287 %280 %280 %280 +%443 = OpSampledImage %442 %436 %435 +%444 = OpImageSampleDrefImplicitLod %7 %443 %440 %280 +%445 = OpLoad %7 %428 %446 = OpFAdd %7 %445 %444 -OpStore %415 %446 -%448 = OpSampledImage %447 %425 %422 -%449 = OpImageSampleDrefImplicitLod %7 %448 %428 %280 -%450 = OpLoad %7 %415 -%451 = OpFAdd %7 %450 %449 -OpStore %415 %451 -%452 = OpSampledImage %429 %423 %422 -%453 = OpImageSampleDrefExplicitLod %7 %452 %427 %280 Lod %454 -%455 = OpLoad %7 %415 -%456 = OpFAdd %7 %455 %453 -OpStore %415 %456 -%457 = OpConvertUToF %7 %199 -%458 = OpCompositeConstruct %287 %427 %457 -%459 = OpSampledImage %434 %424 %422 -%460 = OpImageSampleDrefExplicitLod %7 %459 %458 %280 Lod %454 -%461 = OpLoad %7 %415 -%462 = OpFAdd %7 %461 %460 -OpStore %415 %462 -%463 = OpConvertSToF %7 %283 -%464 = OpCompositeConstruct %287 %427 %463 -%465 = OpSampledImage %434 %424 %422 -%466 = OpImageSampleDrefExplicitLod %7 %465 %464 %280 Lod %454 -%467 = OpLoad %7 %415 +OpStore %428 %446 +%448 = OpConvertUToF %7 %199 +%449 = OpCompositeConstruct %287 %440 %448 +%450 = OpSampledImage %447 %437 %435 +%451 = OpImageSampleDrefImplicitLod %7 %450 %449 %280 +%452 = OpLoad %7 %428 +%453 = OpFAdd %7 %452 %451 +OpStore %428 %453 +%454 = OpConvertSToF %7 %283 +%455 = OpCompositeConstruct %287 %440 %454 +%456 = OpSampledImage %447 %437 %435 +%457 = OpImageSampleDrefImplicitLod %7 %456 %455 %280 +%458 = OpLoad %7 %428 +%459 = OpFAdd %7 %458 %457 +OpStore %428 %459 +%461 = OpSampledImage %460 %438 %435 +%462 = OpImageSampleDrefImplicitLod %7 %461 %441 %280 +%463 = OpLoad %7 %428 +%464 = OpFAdd %7 %463 %462 +OpStore %428 %464 +%465 = OpSampledImage %442 %436 %435 +%466 = OpImageSampleDrefExplicitLod %7 %465 %440 %280 Lod %424 +%467 = OpLoad %7 %428 %468 = OpFAdd %7 %467 %466 -OpStore %415 %468 -%469 = OpSampledImage %447 %425 %422 -%470 = OpImageSampleDrefExplicitLod %7 %469 %428 %280 Lod %454 -%471 = OpLoad %7 %415 -%472 = OpFAdd %7 %471 %470 -OpStore %415 %472 -%473 = OpLoad %7 %415 -OpStore %419 %473 +OpStore %428 %468 +%469 = OpConvertUToF %7 %199 +%470 = OpCompositeConstruct %287 %440 %469 +%471 = OpSampledImage %447 %437 %435 +%472 = OpImageSampleDrefExplicitLod %7 %471 %470 %280 Lod %424 +%473 = OpLoad %7 %428 +%474 = OpFAdd %7 %473 %472 +OpStore %428 %474 +%475 = OpConvertSToF %7 %283 +%476 = OpCompositeConstruct %287 %440 %475 +%477 = OpSampledImage %447 %437 %435 +%478 = OpImageSampleDrefExplicitLod %7 %477 %476 %280 Lod %424 +%479 = OpLoad %7 %428 +%480 = OpFAdd %7 %479 %478 +OpStore %428 %480 +%481 = OpSampledImage %460 %438 %435 +%482 = OpImageSampleDrefExplicitLod %7 %481 %441 %280 Lod %424 +%483 = OpLoad %7 %428 +%484 = OpFAdd %7 %483 %482 +OpStore %428 %484 +%485 = OpLoad %7 %428 +OpStore %432 %485 OpReturn OpFunctionEnd -%476 = OpFunction %2 None %79 -%474 = OpLabel -%477 = OpLoad %16 %49 -%478 = OpLoad %3 %51 -%479 = OpLoad %17 %52 -%480 = OpLoad %24 %64 -%481 = OpLoad %24 %66 -%482 = OpLoad %25 %68 -OpBranch %483 -%483 = OpLabel -%484 = OpCompositeConstruct %285 %280 %280 -%485 = OpSampledImage %295 %477 %480 -%486 = OpImageGather %23 %485 %484 %487 -%488 = OpSampledImage %295 %477 %480 -%489 = OpImageGather %23 %488 %484 %490 ConstOffset %30 -%491 = OpSampledImage %429 %482 %481 -%492 = OpImageDrefGather %23 %491 %484 %280 -%493 = OpSampledImage %429 %482 %481 -%494 = OpImageDrefGather %23 %493 %484 %280 ConstOffset %30 -%496 = OpSampledImage %495 %478 %480 -%497 = OpImageGather %98 %496 %484 %199 -%500 = OpSampledImage %499 %479 %480 -%501 = OpImageGather %498 %500 %484 %199 -%502 = OpConvertUToF %23 %497 -%503 = OpConvertSToF %23 %501 -%504 = OpFAdd %23 %502 %503 -%505 = OpFAdd %23 %486 %489 -%506 = OpFAdd %23 %505 %492 -%507 = OpFAdd %23 %506 %494 -%508 = OpFAdd %23 %507 %504 -OpStore %475 %508 +%488 = OpFunction %2 None %79 +%486 = OpLabel +%489 = OpLoad %16 %49 +%490 = OpLoad %3 %51 +%491 = OpLoad %17 %52 +%492 = OpLoad %24 %64 +%493 = OpLoad %24 %66 +%494 = OpLoad %25 %68 +OpBranch %495 +%495 = OpLabel +%496 = OpCompositeConstruct %285 %280 %280 +%497 = OpSampledImage %295 %489 %492 +%498 = OpImageGather %23 %497 %496 %499 +%500 = OpSampledImage %295 %489 %492 +%501 = OpImageGather %23 %500 %496 %502 ConstOffset %30 +%503 = OpSampledImage %442 %494 %493 +%504 = OpImageDrefGather %23 %503 %496 %280 +%505 = OpSampledImage %442 %494 %493 +%506 = OpImageDrefGather %23 %505 %496 %280 ConstOffset %30 +%508 = OpSampledImage %507 %490 %492 +%509 = OpImageGather %98 %508 %496 %199 +%512 = OpSampledImage %511 %491 %492 +%513 = OpImageGather %510 %512 %496 %199 +%514 = OpConvertUToF %23 %509 +%515 = OpConvertSToF %23 %513 +%516 = OpFAdd %23 %514 %515 +%517 = OpFAdd %23 %498 %501 +%518 = OpFAdd %23 %517 %504 +%519 = OpFAdd %23 %518 %506 +%520 = OpFAdd %23 %519 %516 +OpStore %487 %520 OpReturn OpFunctionEnd -%511 = OpFunction %2 None %79 -%509 = OpLabel -%512 = OpLoad %24 %64 -%513 = OpLoad %25 %68 -OpBranch %514 -%514 = OpLabel -%515 = OpCompositeConstruct %285 %280 %280 -%516 = OpSampledImage %429 %513 %512 -%517 = OpImageSampleImplicitLod %23 %516 %515 -%518 = OpCompositeExtract %7 %517 0 -%519 = OpSampledImage %429 %513 %512 -%520 = OpImageGather %23 %519 %515 %199 -%521 = OpCompositeConstruct %23 %518 %518 %518 %518 -%522 = OpFAdd %23 %521 %520 -OpStore %510 %522 +%523 = OpFunction %2 None %79 +%521 = OpLabel +%524 = OpLoad %24 %64 +%525 = OpLoad %25 %68 +OpBranch %526 +%526 = OpLabel +%527 = OpCompositeConstruct %285 %280 %280 +%528 = OpSampledImage %442 %525 %524 +%529 = OpImageSampleImplicitLod %23 %528 %527 +%530 = OpCompositeExtract %7 %529 0 +%531 = OpSampledImage %442 %525 %524 +%532 = OpImageGather %23 %531 %527 %199 +%533 = OpCompositeConstruct %23 %530 %530 %530 %530 +%534 = OpFAdd %23 %533 %532 +OpStore %522 %534 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