From 8caa2bd87e2e4b3bc3ba6276f5bd898d601f1c7d Mon Sep 17 00:00:00 2001 From: Dzmitry Malyshau Date: Wed, 15 Dec 2021 00:29:07 -0500 Subject: [PATCH] Gather operations --- src/back/dot/mod.rs | 7 +- src/back/glsl/mod.rs | 29 ++++--- src/back/hlsl/writer.rs | 34 +++++--- src/back/msl/writer.rs | 30 ++++++- src/back/spv/block.rs | 2 + src/back/spv/image.rs | 30 +++++-- src/back/spv/instructions.rs | 27 +++++++ src/back/wgsl/writer.rs | 49 ++++++++++++ src/front/glsl/builtins.rs | 1 + src/front/spv/image.rs | 1 + src/front/wgsl/mod.rs | 96 +++++++++++++++++++++++ src/lib.rs | 5 +- src/proc/typifier.rs | 18 +++++ src/valid/analyzer.rs | 1 + src/valid/expression.rs | 27 +++++++ tests/in/image.param.ron | 2 +- tests/in/image.wgsl | 19 +++++ tests/out/glsl/image.gather.Fragment.glsl | 22 ++++++ tests/out/hlsl/image.hlsl | 18 +++++ tests/out/hlsl/image.hlsl.config | 2 +- tests/out/ir/shadow.ron | 1 + tests/out/msl/image.msl | 32 ++++++++ tests/out/spv/image.spvasm | 54 ++++++++++++- tests/out/wgsl/image.wgsl | 18 +++++ 24 files changed, 484 insertions(+), 41 deletions(-) create mode 100644 tests/out/glsl/image.gather.Fragment.glsl diff --git a/src/back/dot/mod.rs b/src/back/dot/mod.rs index b63a163598..bee46704ce 100644 --- a/src/back/dot/mod.rs +++ b/src/back/dot/mod.rs @@ -231,6 +231,7 @@ fn write_fun( E::ImageSample { image, sampler, + gather, coordinate, array_index, offset: _, @@ -260,7 +261,11 @@ fn write_fun( if let Some(expr) = depth_ref { edges.insert("depth_ref", expr); } - ("ImageSample".into(), 5) + let string = match gather { + Some(component) => Cow::Owned(format!("ImageGather{:?}", component)), + _ => Cow::Borrowed("ImageSample"), + }; + (string, 5) } E::ImageLoad { image, diff --git a/src/back/glsl/mod.rs b/src/back/glsl/mod.rs index 9bea98eb16..e32c803f05 100644 --- a/src/back/glsl/mod.rs +++ b/src/back/glsl/mod.rs @@ -1928,6 +1928,7 @@ impl<'a, W: Write> Writer<'a, W> { Expression::ImageSample { image, sampler: _, //TODO? + gather, coordinate, array_index, offset, @@ -1962,6 +1963,7 @@ impl<'a, W: Write> Writer<'a, W> { let workaround_lod_array_shadow_as_grad = (array_index.is_some() || dim == crate::ImageDimension::Cube) && depth_ref.is_some() + && gather.is_none() && !self .options .writer_flags @@ -1969,6 +1971,7 @@ impl<'a, W: Write> Writer<'a, W> { //Write the function to be used depending on the sample level let fun_name = match level { + crate::SampleLevel::Zero if gather.is_some() => "textureGather", crate::SampleLevel::Auto | crate::SampleLevel::Bias(_) => "texture", crate::SampleLevel::Zero | crate::SampleLevel::Exact(_) => { if workaround_lod_array_shadow_as_grad { @@ -2002,8 +2005,8 @@ impl<'a, W: Write> Writer<'a, W> { if array_index.is_some() { coord_dim += 1; } - let cube_array_shadow = coord_dim == 4; - if depth_ref.is_some() && !cube_array_shadow { + let merge_depth_ref = depth_ref.is_some() && gather.is_none() && coord_dim < 4; + if merge_depth_ref { coord_dim += 1; } @@ -2021,21 +2024,17 @@ impl<'a, W: Write> Writer<'a, W> { write!(self.out, ", ")?; self.write_expr(expr, ctx)?; } - if !cube_array_shadow { - if let Some(expr) = depth_ref { - 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 cube_array_shadow { - if let Some(expr) = depth_ref { - write!(self.out, ", ")?; - self.write_expr(expr, ctx)?; - } + if let (Some(expr), false) = (depth_ref, merge_depth_ref) { + write!(self.out, ", ")?; + self.write_expr(expr, ctx)?; } match level { @@ -2045,7 +2044,7 @@ impl<'a, W: Write> Writer<'a, W> { crate::SampleLevel::Zero => { if workaround_lod_array_shadow_as_grad { write!(self.out, ", vec2(0,0), vec2(0,0)")?; - } else { + } else if gather.is_none() { write!(self.out, ", 0.0")?; } } @@ -2082,6 +2081,10 @@ impl<'a, W: Write> Writer<'a, W> { } } + if let (Some(component), None) = (gather, depth_ref) { + write!(self.out, ", {}", component as usize)?; + } + // End the function write!(self.out, ")")? } diff --git a/src/back/hlsl/writer.rs b/src/back/hlsl/writer.rs index d8cfde3cd0..ca00d889c0 100644 --- a/src/back/hlsl/writer.rs +++ b/src/back/hlsl/writer.rs @@ -1594,6 +1594,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { Expression::ImageSample { image, sampler, + gather, coordinate, array_index, offset, @@ -1601,23 +1602,30 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> { depth_ref, } => { use crate::SampleLevel as Sl; + const COMPONENTS: [&str; 4] = ["", "Green", "Blue", "Alpha"]; - let texture_func = match level { - Sl::Auto => { - if depth_ref.is_some() { - "SampleCmp" - } else { - "Sample" - } - } - Sl::Zero => "SampleCmpLevelZero", - Sl::Exact(_) => "SampleLevel", - Sl::Bias(_) => "SampleBias", - Sl::Gradient { .. } => "SampleGrad", + let (base_str, component_str) = match gather { + Some(component) => ("Gather", COMPONENTS[component as usize]), + None => ("Sample", ""), + }; + let cmp_str = match depth_ref { + Some(_) => "Cmp", + None => "", + }; + let level_str = match level { + Sl::Zero if gather.is_none() => "LevelZero", + Sl::Auto | Sl::Zero => "", + Sl::Exact(_) => "Level", + Sl::Bias(_) => "Bias", + Sl::Gradient { .. } => "Grad", }; self.write_expr(module, image, func_ctx)?; - write!(self.out, ".{}(", texture_func)?; + 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( diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 51f892d783..1050f1434c 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -896,18 +896,23 @@ impl Writer { crate::Expression::ImageSample { image, sampler, + gather, coordinate, array_index, offset, level, depth_ref, } => { - let op = match depth_ref { - Some(_) => "sample_compare", + let main_op = match gather { + Some(_) => "gather", None => "sample", }; + let comparison_op = match depth_ref { + Some(_) => "_compare", + None => "", + }; self.put_expression(image, context, false)?; - write!(self.out, ".{}(", op)?; + write!(self.out, ".{}{}(", main_op, comparison_op)?; self.put_expression(sampler, context, true)?; write!(self.out, ", ")?; self.put_expression(coordinate, context, true)?; @@ -931,6 +936,25 @@ impl Writer { }; write!(self.out, ", {}", coco)?; } + match gather { + None | Some(crate::SwizzleComponent::X) => {} + Some(component) => { + let is_cube_map = match *context.resolve_type(image) { + crate::TypeInner::Image { + dim: crate::ImageDimension::Cube, + .. + } => true, + _ => false, + }; + // Offset always comes before the gather, except + // in cube maps where it's not applicable + if offset.is_none() && !is_cube_map { + write!(self.out, ", int2(0)")?; + } + let letter = ['x', 'y', 'z', 'w'][component as usize]; + write!(self.out, ", {}::component::{}", NAMESPACE, letter)?; + } + } write!(self.out, ")")?; } crate::Expression::ImageLoad { diff --git a/src/back/spv/block.rs b/src/back/spv/block.rs index a59f854531..5f01ca028f 100644 --- a/src/back/spv/block.rs +++ b/src/back/spv/block.rs @@ -891,6 +891,7 @@ impl<'w> BlockContext<'w> { crate::Expression::ImageSample { image, sampler, + gather, coordinate, array_index, offset, @@ -900,6 +901,7 @@ impl<'w> BlockContext<'w> { result_type_id, image, sampler, + gather, coordinate, array_index, offset, diff --git a/src/back/spv/image.rs b/src/back/spv/image.rs index 442045ed2a..443ef8c57e 100644 --- a/src/back/spv/image.rs +++ b/src/back/spv/image.rs @@ -799,6 +799,7 @@ impl<'w> BlockContext<'w> { result_type_id: Word, image: Handle, sampler: Handle, + gather: Option, coordinate: Handle, array_index: Option>, offset: Option>, @@ -816,7 +817,7 @@ impl<'w> BlockContext<'w> { crate::TypeInner::Image { class: crate::ImageClass::Depth { .. }, .. - } => depth_ref.is_none(), + } => depth_ref.is_none() && gather.is_none(), _ => false, }; let sample_result_type_id = if needs_sub_access { @@ -853,8 +854,23 @@ impl<'w> BlockContext<'w> { let mut mask = spirv::ImageOperands::empty(); mask.set(spirv::ImageOperands::CONST_OFFSET, offset.is_some()); - let mut main_instruction = match level { - crate::SampleLevel::Zero => { + let mut main_instruction = match (level, gather) { + (_, Some(component)) => { + let component_id = self.get_index_constant(component as u32); + let mut inst = Instruction::image_gather( + sample_result_type_id, + id, + sampled_image_id, + coordinates_id, + component_id, + depth_id, + ); + if !mask.is_empty() { + inst.add_operand(mask.bits()); + } + inst + } + (crate::SampleLevel::Zero, None) => { let mut inst = Instruction::image_sample( sample_result_type_id, id, @@ -874,7 +890,7 @@ impl<'w> BlockContext<'w> { inst } - crate::SampleLevel::Auto => { + (crate::SampleLevel::Auto, None) => { let mut inst = Instruction::image_sample( sample_result_type_id, id, @@ -888,7 +904,7 @@ impl<'w> BlockContext<'w> { } inst } - crate::SampleLevel::Exact(lod_handle) => { + (crate::SampleLevel::Exact(lod_handle), None) => { let mut inst = Instruction::image_sample( sample_result_type_id, id, @@ -905,7 +921,7 @@ impl<'w> BlockContext<'w> { inst } - crate::SampleLevel::Bias(bias_handle) => { + (crate::SampleLevel::Bias(bias_handle), None) => { let mut inst = Instruction::image_sample( sample_result_type_id, id, @@ -922,7 +938,7 @@ impl<'w> BlockContext<'w> { inst } - crate::SampleLevel::Gradient { x, y } => { + (crate::SampleLevel::Gradient { x, y }, None) => { let mut inst = Instruction::image_sample( sample_result_type_id, id, diff --git a/src/back/spv/instructions.rs b/src/back/spv/instructions.rs index 211c51a1ed..79f54d4500 100644 --- a/src/back/spv/instructions.rs +++ b/src/back/spv/instructions.rs @@ -569,6 +569,33 @@ impl super::Instruction { instruction } + pub(super) fn image_gather( + result_type_id: Word, + id: Word, + sampled_image: Word, + coordinates: Word, + component_id: Word, + depth_ref: Option, + ) -> Self { + let op = match depth_ref { + None => Op::ImageGather, + Some(_) => Op::ImageDrefGather, + }; + + let mut instruction = Self::new(op); + instruction.set_type(result_type_id); + instruction.set_result(id); + instruction.add_operand(sampled_image); + instruction.add_operand(coordinates); + if let Some(dref) = depth_ref { + instruction.add_operand(dref); + } else { + instruction.add_operand(component_id); + } + + instruction + } + pub(super) fn image_fetch_or_read( op: Op, result_type_id: Word, diff --git a/src/back/wgsl/writer.rs b/src/back/wgsl/writer.rs index 4d06b7a941..6c226fe443 100644 --- a/src/back/wgsl/writer.rs +++ b/src/back/wgsl/writer.rs @@ -1239,6 +1239,7 @@ impl Writer { Expression::ImageSample { image, sampler, + gather: None, coordinate, array_index, offset, @@ -1306,6 +1307,54 @@ impl Writer { write!(self.out, ")")?; } + Expression::ImageSample { + image, + sampler, + gather: Some(component), + coordinate, + array_index, + offset, + level: _, + depth_ref, + } => { + let suffix_cmp = match depth_ref { + Some(_) => "Compare", + None => "", + }; + + write!(self.out, "textureGather{}(", suffix_cmp)?; + match *func_ctx.info[image].ty.inner_with(&module.types) { + TypeInner::Image { + class: crate::ImageClass::Depth { multi: _ }, + .. + } => {} + _ => { + write!(self.out, "{}, ", component as u8)?; + } + } + self.write_expr(module, image, func_ctx)?; + write!(self.out, ", ")?; + self.write_expr(module, sampler, func_ctx)?; + write!(self.out, ", ")?; + self.write_expr(module, coordinate, func_ctx)?; + + if let Some(array_index) = array_index { + write!(self.out, ", ")?; + self.write_expr(module, array_index, func_ctx)?; + } + + if let Some(depth_ref) = depth_ref { + write!(self.out, ", ")?; + self.write_expr(module, depth_ref, func_ctx)?; + } + + if let Some(offset) = offset { + write!(self.out, ", ")?; + self.write_constant(module, offset)?; + } + + write!(self.out, ")")?; + } Expression::ImageQuery { image, query } => { use crate::ImageQuery as Iq; diff --git a/src/front/glsl/builtins.rs b/src/front/glsl/builtins.rs index ed220bc98b..a7416ded2d 100644 --- a/src/front/glsl/builtins.rs +++ b/src/front/glsl/builtins.rs @@ -1930,6 +1930,7 @@ fn texture_call( Expression::ImageSample { image, sampler, + gather: None, //TODO coordinate: comps.coordinate, array_index, offset, diff --git a/src/front/spv/image.rs b/src/front/spv/image.rs index 1df6cb5f98..e797a516a1 100644 --- a/src/front/spv/image.rs +++ b/src/front/spv/image.rs @@ -606,6 +606,7 @@ impl> super::Parser { let expr = crate::Expression::ImageSample { image: si_lexp.image, sampler: si_lexp.sampler, + gather: None, //TODO coordinate, array_index, offset, diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs index f9ef1b981e..a58f63d423 100644 --- a/src/front/wgsl/mod.rs +++ b/src/front/wgsl/mod.rs @@ -148,6 +148,7 @@ pub enum Error<'a> { }, InvalidResolve(ResolveError), InvalidForInitializer(Span), + InvalidGatherComponent(Span, i32), ReservedIdentifierPrefix(Span), UnknownStorageClass(Span), UnknownAttribute(Span), @@ -342,6 +343,11 @@ impl<'a> Error<'a> { labels: vec![(bad_span.clone(), "not an assignment or function call".into())], notes: vec![], }, + Error::InvalidGatherComponent(ref bad_span, component) => ParseError { + message: format!("textureGather component {} doesn't exist, must be 0, 1, 2, or 3", component), + labels: vec![(bad_span.clone(), "invalid component".into())], + notes: vec![], + }, Error::ReservedIdentifierPrefix(ref bad_span) => ParseError { message: format!("Identifier starts with a reserved prefix: '{}'", &source[bad_span.clone()]), labels: vec![(bad_span.clone(), "invalid identifier".into())], @@ -1648,6 +1654,7 @@ impl Parser { crate::Expression::ImageSample { image: sc.image, sampler: ctx.lookup_ident.lookup(sampler_name, sampler_span)?.handle, + gather: None, coordinate, array_index, offset, @@ -1681,6 +1688,7 @@ impl Parser { crate::Expression::ImageSample { image: sc.image, sampler: ctx.lookup_ident.lookup(sampler_name, sampler_span)?.handle, + gather: None, coordinate, array_index, offset, @@ -1714,6 +1722,7 @@ impl Parser { crate::Expression::ImageSample { image: sc.image, sampler: ctx.lookup_ident.lookup(sampler_name, sampler_span)?.handle, + gather: None, coordinate, array_index, offset, @@ -1749,6 +1758,7 @@ impl Parser { crate::Expression::ImageSample { image: sc.image, sampler: ctx.lookup_ident.lookup(sampler_name, sampler_span)?.handle, + gather: None, coordinate, array_index, offset, @@ -1782,6 +1792,7 @@ impl Parser { crate::Expression::ImageSample { image: sc.image, sampler: ctx.lookup_ident.lookup(sampler_name, sampler_span)?.handle, + gather: None, coordinate, array_index, offset, @@ -1815,6 +1826,91 @@ impl Parser { crate::Expression::ImageSample { image: sc.image, sampler: ctx.lookup_ident.lookup(sampler_name, sampler_span)?.handle, + gather: None, + coordinate, + array_index, + offset, + level: crate::SampleLevel::Zero, + depth_ref: Some(reference), + } + } + "textureGather" => { + let _ = lexer.next(); + lexer.open_arguments()?; + let component = if let ( + Token::Number { + value, + ty: NumberType::Sint, + width: None, + }, + span, + ) = lexer.peek() + { + let _ = lexer.next(); + lexer.expect(Token::Separator(','))?; + let index = get_i32_literal(value, span.clone())?; + *crate::SwizzleComponent::XYZW + .get(index as usize) + .ok_or(Error::InvalidGatherComponent(span, index))? + } else { + crate::SwizzleComponent::X + }; + let (image_name, image_span) = lexer.next_ident_with_span()?; + lexer.expect(Token::Separator(','))?; + let (sampler_name, sampler_span) = lexer.next_ident_with_span()?; + lexer.expect(Token::Separator(','))?; + let coordinate = self.parse_general_expression(lexer, ctx.reborrow())?; + let sc = ctx.prepare_sampling(image_name, image_span)?; + let array_index = if sc.arrayed { + lexer.expect(Token::Separator(','))?; + Some(self.parse_general_expression(lexer, ctx.reborrow())?) + } else { + None + }; + let offset = if lexer.skip(Token::Separator(',')) { + Some(self.parse_const_expression(lexer, ctx.types, ctx.constants)?) + } else { + None + }; + lexer.close_arguments()?; + crate::Expression::ImageSample { + image: sc.image, + sampler: ctx.lookup_ident.lookup(sampler_name, sampler_span)?.handle, + gather: Some(component), + coordinate, + array_index, + offset, + level: crate::SampleLevel::Zero, + depth_ref: None, + } + } + "textureGatherCompare" => { + let _ = lexer.next(); + lexer.open_arguments()?; + let (image_name, image_span) = lexer.next_ident_with_span()?; + lexer.expect(Token::Separator(','))?; + let (sampler_name, sampler_span) = lexer.next_ident_with_span()?; + lexer.expect(Token::Separator(','))?; + let coordinate = self.parse_general_expression(lexer, ctx.reborrow())?; + let sc = ctx.prepare_sampling(image_name, image_span)?; + let array_index = if sc.arrayed { + lexer.expect(Token::Separator(','))?; + Some(self.parse_general_expression(lexer, ctx.reborrow())?) + } else { + None + }; + lexer.expect(Token::Separator(','))?; + let reference = self.parse_general_expression(lexer, ctx.reborrow())?; + let offset = if lexer.skip(Token::Separator(',')) { + Some(self.parse_const_expression(lexer, ctx.types, ctx.constants)?) + } else { + None + }; + lexer.close_arguments()?; + crate::Expression::ImageSample { + image: sc.image, + sampler: ctx.lookup_ident.lookup(sampler_name, sampler_span)?.handle, + gather: Some(crate::SwizzleComponent::X), coordinate, array_index, offset, diff --git a/src/lib.rs b/src/lib.rs index a2d2ec15a7..7be2b53773 100644 --- a/src/lib.rs +++ b/src/lib.rs @@ -975,7 +975,7 @@ pub enum ImageQuery { /// Component selection for a vector swizzle. #[repr(u8)] -#[derive(Clone, Copy, Debug, PartialEq)] +#[derive(Clone, Copy, Debug, PartialEq, PartialOrd)] #[cfg_attr(feature = "serialize", derive(Serialize))] #[cfg_attr(feature = "deserialize", derive(Deserialize))] pub enum SwizzleComponent { @@ -1128,6 +1128,9 @@ pub enum Expression { ImageSample { image: Handle, sampler: Handle, + /// If Some(), this operation is a gather operation + /// on the selected component. + gather: Option, coordinate: Handle, array_index: Option>, offset: Option>, diff --git a/src/proc/typifier.rs b/src/proc/typifier.rs index e1d2bd6b8d..cf6035802f 100644 --- a/src/proc/typifier.rs +++ b/src/proc/typifier.rs @@ -472,6 +472,24 @@ impl<'a> ResolveContext<'a> { return Err(ResolveError::InvalidPointer(pointer)); } }, + crate::Expression::ImageSample { + image, + gather: Some(_), + .. + } => match *past(image).inner_with(types) { + Ti::Image { class, .. } => TypeResolution::Value(Ti::Vector { + kind: match class { + crate::ImageClass::Sampled { kind, multi: _ } => kind, + _ => crate::ScalarKind::Float, + }, + width: 4, + size: crate::VectorSize::Quad, + }), + ref other => { + log::error!("Image type {:?}", other); + return Err(ResolveError::InvalidImage(image)); + } + }, crate::Expression::ImageSample { image, .. } | crate::Expression::ImageLoad { image, .. } => match *past(image).inner_with(types) { Ti::Image { class, .. } => TypeResolution::Value(match class { diff --git a/src/valid/analyzer.rs b/src/valid/analyzer.rs index c581a5681c..ab4d77c421 100644 --- a/src/valid/analyzer.rs +++ b/src/valid/analyzer.rs @@ -462,6 +462,7 @@ impl FunctionInfo { E::ImageSample { image, sampler, + gather: _, coordinate, array_index, offset: _, diff --git a/src/valid/expression.rs b/src/valid/expression.rs index 21da06baab..26ec967b19 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -99,6 +99,12 @@ pub enum ExpressionError { InvalidDepthReference(Handle), #[error("Depth sample level can only be Auto or Zero")] InvalidDepthSampleLevel, + #[error("Gather level can only be Zero")] + InvalidGatherLevel, + #[error("Gather component {0:?} doesn't exist in the image")] + InvalidGatherComponent(crate::SwizzleComponent), + #[error("Gather can't be done for image dimension {0:?}")] + InvalidGatherDimension(crate::ImageDimension), #[error("Sample level (exact) type {0:?} is not a scalar float")] InvalidSampleLevelExactType(Handle), #[error("Sample level (bias) type {0:?} is not a scalar float")] @@ -343,6 +349,7 @@ impl super::Validator { E::ImageSample { image, sampler, + gather, coordinate, array_index, offset, @@ -466,6 +473,26 @@ impl super::Validator { } } + if let Some(component) = gather { + match dim { + crate::ImageDimension::D2 | crate::ImageDimension::Cube => {} + crate::ImageDimension::D1 | crate::ImageDimension::D3 => { + return Err(ExpressionError::InvalidGatherDimension(dim)) + } + }; + let max_component = match class { + crate::ImageClass::Depth { .. } => crate::SwizzleComponent::X, + _ => crate::SwizzleComponent::W, + }; + if component > max_component { + return Err(ExpressionError::InvalidGatherComponent(component)); + } + match level { + crate::SampleLevel::Zero => {} + _ => return Err(ExpressionError::InvalidGatherLevel), + } + } + // check level properties match level { crate::SampleLevel::Auto => ShaderStages::FRAGMENT, diff --git a/tests/in/image.param.ron b/tests/in/image.param.ron index 76105e3352..5b6d71defa 100644 --- a/tests/in/image.param.ron +++ b/tests/in/image.param.ron @@ -3,5 +3,5 @@ version: (1, 1), debug: true, ), - glsl_exclude_list: ["depth_load", "levels_queries"] + glsl_exclude_list: ["depth_load", "depth_no_comparison", "levels_queries"] ) diff --git a/tests/in/image.wgsl b/tests/in/image.wgsl index e9fd4f0fa9..8d27b88208 100644 --- a/tests/in/image.wgsl +++ b/tests/in/image.wgsl @@ -116,3 +116,22 @@ fn sample_comparison() -> [[location(0)]] f32 { let s2d_depth_level = textureSampleCompareLevel(image_2d_depth, sampler_cmp, tc, dref); return s2d_depth + s2d_depth_level; } + +[[stage(fragment)]] +fn gather() -> [[location(0)]] vec4 { + let tc = vec2(0.5); + let dref = 0.5; + let s2d = textureGather(1, image_2d, sampler_reg, tc); + let s2d_offset = textureGather(3, image_2d, sampler_reg, tc, vec2(3, 1)); + let s2d_depth = textureGatherCompare(image_2d_depth, sampler_cmp, tc, dref); + let s2d_depth_offset = textureGatherCompare(image_2d_depth, sampler_cmp, tc, dref, vec2(3, 1)); + return s2d + s2d_offset + s2d_depth + s2d_depth_offset; +} + +[[stage(fragment)]] +fn depth_no_comparison() -> [[location(0)]] vec4 { + let tc = vec2(0.5); + let s2d = textureSample(image_2d_depth, sampler_reg, tc); + let s2d_gather = textureGather(image_2d_depth, sampler_reg, tc); + return s2d + s2d_gather; +} diff --git a/tests/out/glsl/image.gather.Fragment.glsl b/tests/out/glsl/image.gather.Fragment.glsl new file mode 100644 index 0000000000..89961fcff9 --- /dev/null +++ b/tests/out/glsl/image.gather.Fragment.glsl @@ -0,0 +1,22 @@ +#version 310 es +#extension GL_EXT_texture_cube_map_array : require + +precision highp float; +precision highp int; + +uniform highp sampler2D _group_0_binding_1; + +uniform highp sampler2DShadow _group_1_binding_2; + +layout(location = 0) out vec4 _fs2p_location0; + +void main() { + vec2 tc = vec2(0.5); + vec4 s2d = textureGather(_group_0_binding_1, vec2(tc), 1); + vec4 s2d_offset = textureGatherOffset(_group_0_binding_1, vec2(tc), ivec2(3, 1), 3); + vec4 s2d_depth = textureGather(_group_1_binding_2, vec2(tc), 0.5); + vec4 s2d_depth_offset = textureGatherOffset(_group_1_binding_2, vec2(tc), 0.5, ivec2(3, 1)); + _fs2p_location0 = (((s2d + s2d_offset) + s2d_depth) + s2d_depth_offset); + return; +} + diff --git a/tests/out/hlsl/image.hlsl b/tests/out/hlsl/image.hlsl index 796c748411..1d36273f49 100644 --- a/tests/out/hlsl/image.hlsl +++ b/tests/out/hlsl/image.hlsl @@ -229,3 +229,21 @@ float sample_comparison() : SV_Target0 float s2d_depth_level = image_2d_depth.SampleCmpLevelZero(sampler_cmp, tc_1, 0.5); return (s2d_depth + s2d_depth_level); } + +float4 gather() : SV_Target0 +{ + float2 tc_2 = float2(0.5.xx); + float4 s2d_1 = image_2d.GatherGreen(sampler_reg, tc_2); + float4 s2d_offset_1 = image_2d.GatherAlpha(sampler_reg, tc_2, int2(3, 1)); + float4 s2d_depth_1 = image_2d_depth.GatherCmp(sampler_cmp, tc_2, 0.5); + float4 s2d_depth_offset = image_2d_depth.GatherCmp(sampler_cmp, tc_2, 0.5, int2(3, 1)); + return (((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset); +} + +float4 depth_no_comparison() : SV_Target0 +{ + float2 tc_3 = float2(0.5.xx); + float s2d_2 = image_2d_depth.Sample(sampler_reg, tc_3); + float4 s2d_gather = image_2d_depth.Gather(sampler_reg, tc_3); + return (float4(s2d_2.xxxx) + s2d_gather); +} diff --git a/tests/out/hlsl/image.hlsl.config b/tests/out/hlsl/image.hlsl.config index 28af7d931d..ba8a24dd08 100644 --- a/tests/out/hlsl/image.hlsl.config +++ b/tests/out/hlsl/image.hlsl.config @@ -1,3 +1,3 @@ vertex=(queries:vs_5_1 levels_queries:vs_5_1 ) -fragment=(sample_:ps_5_1 sample_comparison:ps_5_1 ) +fragment=(sample_:ps_5_1 sample_comparison:ps_5_1 gather:ps_5_1 depth_no_comparison:ps_5_1 ) compute=(main:cs_5_1 depth_load:cs_5_1 ) diff --git a/tests/out/ir/shadow.ron b/tests/out/ir/shadow.ron index 0356f1371e..ab8c3eeb7b 100644 --- a/tests/out/ir/shadow.ron +++ b/tests/out/ir/shadow.ron @@ -870,6 +870,7 @@ ImageSample( image: 4, sampler: 5, + gather: None, coordinate: 72, array_index: Some(74), offset: None, diff --git a/tests/out/msl/image.msl b/tests/out/msl/image.msl index 449dc9c433..cbe68c6494 100644 --- a/tests/out/msl/image.msl +++ b/tests/out/msl/image.msl @@ -122,3 +122,35 @@ fragment sample_comparisonOutput sample_comparison( float s2d_depth_level = image_2d_depth.sample_compare(sampler_cmp, tc_1, 0.5); return sample_comparisonOutput { s2d_depth + s2d_depth_level }; } + + +struct gatherOutput { + metal::float4 member_6 [[color(0)]]; +}; +fragment gatherOutput gather( + metal::texture2d image_2d [[user(fake0)]] +, metal::sampler sampler_reg [[user(fake0)]] +, metal::sampler sampler_cmp [[user(fake0)]] +, metal::depth2d image_2d_depth [[user(fake0)]] +) { + metal::float2 tc_2 = metal::float2(0.5); + metal::float4 s2d_1 = image_2d.gather(sampler_reg, tc_2, int2(0), metal::component::y); + metal::float4 s2d_offset_1 = image_2d.gather(sampler_reg, tc_2, const_type_8_, metal::component::w); + metal::float4 s2d_depth_1 = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5); + metal::float4 s2d_depth_offset = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5, const_type_8_); + return gatherOutput { ((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset }; +} + + +struct depth_no_comparisonOutput { + metal::float4 member_7 [[color(0)]]; +}; +fragment depth_no_comparisonOutput depth_no_comparison( + metal::sampler sampler_reg [[user(fake0)]] +, metal::depth2d image_2d_depth [[user(fake0)]] +) { + metal::float2 tc_3 = metal::float2(0.5); + float s2d_2 = image_2d_depth.sample(sampler_reg, tc_3); + metal::float4 s2d_gather = image_2d_depth.gather(sampler_reg, tc_3); + return depth_no_comparisonOutput { metal::float4(s2d_2) + s2d_gather }; +} diff --git a/tests/out/spv/image.spvasm b/tests/out/spv/image.spvasm index 9fc7d518e7..82ab063c92 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: 244 +; Bound: 280 OpCapability SampledCubeArray OpCapability ImageQuery OpCapability Image1D @@ -15,10 +15,14 @@ OpEntryPoint Vertex %128 "queries" %126 OpEntryPoint Vertex %176 "levels_queries" %175 OpEntryPoint Fragment %205 "sample" %204 OpEntryPoint Fragment %232 "sample_comparison" %230 +OpEntryPoint Fragment %246 "gather" %245 +OpEntryPoint Fragment %268 "depth_no_comparison" %267 OpExecutionMode %69 LocalSize 16 1 1 OpExecutionMode %107 LocalSize 16 1 1 OpExecutionMode %205 OriginUpperLeft OpExecutionMode %232 OriginUpperLeft +OpExecutionMode %246 OriginUpperLeft +OpExecutionMode %268 OriginUpperLeft OpSource GLSL 450 OpName %31 "image_mipmapped_src" OpName %33 "image_multisampled_src" @@ -45,6 +49,8 @@ OpName %128 "queries" OpName %176 "levels_queries" OpName %205 "sample" OpName %232 "sample_comparison" +OpName %246 "gather" +OpName %268 "depth_no_comparison" OpDecorate %31 DescriptorSet 0 OpDecorate %31 Binding 0 OpDecorate %33 DescriptorSet 0 @@ -88,6 +94,8 @@ OpDecorate %126 BuiltIn Position OpDecorate %175 BuiltIn Position OpDecorate %204 Location 0 OpDecorate %230 Location 0 +OpDecorate %245 Location 0 +OpDecorate %267 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 10 @@ -170,6 +178,10 @@ OpDecorate %230 Location 0 %230 = OpVariable %231 Output %237 = OpTypeSampledImage %29 %242 = OpConstant %8 0.0 +%245 = OpVariable %127 Output +%255 = OpConstant %12 1 +%258 = OpConstant %12 3 +%267 = OpVariable %127 Output %69 = OpFunction %2 None %70 %65 = OpLabel %68 = OpLoad %18 %66 @@ -353,4 +365,44 @@ OpBranch %235 %243 = OpFAdd %8 %239 %241 OpStore %230 %243 OpReturn +OpFunctionEnd +%246 = OpFunction %2 None %70 +%244 = OpLabel +%247 = OpLoad %21 %47 +%248 = OpLoad %28 %59 +%249 = OpLoad %28 %61 +%250 = OpLoad %29 %63 +OpBranch %251 +%251 = OpLabel +%252 = OpCompositeConstruct %210 %7 %7 +%253 = OpSampledImage %216 %247 %248 +%254 = OpImageGather %27 %253 %252 %255 +%256 = OpSampledImage %216 %247 %248 +%257 = OpImageGather %27 %256 %252 %258 ConstOffset %30 +%259 = OpSampledImage %237 %250 %249 +%260 = OpImageDrefGather %27 %259 %252 %7 +%261 = OpSampledImage %237 %250 %249 +%262 = OpImageDrefGather %27 %261 %252 %7 ConstOffset %30 +%263 = OpFAdd %27 %254 %257 +%264 = OpFAdd %27 %263 %260 +%265 = OpFAdd %27 %264 %262 +OpStore %245 %265 +OpReturn +OpFunctionEnd +%268 = OpFunction %2 None %70 +%266 = OpLabel +%269 = OpLoad %28 %59 +%270 = OpLoad %29 %63 +OpBranch %271 +%271 = OpLabel +%272 = OpCompositeConstruct %210 %7 %7 +%273 = OpSampledImage %237 %270 %269 +%274 = OpImageSampleImplicitLod %27 %273 %272 +%275 = OpCompositeExtract %8 %274 0 +%276 = OpSampledImage %237 %270 %269 +%277 = OpImageGather %27 %276 %272 %136 +%278 = OpCompositeConstruct %27 %275 %275 %275 %275 +%279 = OpFAdd %27 %278 %277 +OpStore %267 %279 +OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/image.wgsl b/tests/out/wgsl/image.wgsl index 0bc840088c..0a7eda004b 100644 --- a/tests/out/wgsl/image.wgsl +++ b/tests/out/wgsl/image.wgsl @@ -103,3 +103,21 @@ fn sample_comparison() -> [[location(0)]] f32 { let s2d_depth_level = textureSampleCompareLevel(image_2d_depth, sampler_cmp, tc_1, 0.5); return (s2d_depth + s2d_depth_level); } + +[[stage(fragment)]] +fn gather() -> [[location(0)]] vec4 { + let tc_2 = vec2(0.5); + let s2d_1 = textureGather(1, image_2d, sampler_reg, tc_2); + let s2d_offset_1 = textureGather(3, image_2d, sampler_reg, tc_2, vec2(3, 1)); + let s2d_depth_1 = textureGatherCompare(image_2d_depth, sampler_cmp, tc_2, 0.5); + let s2d_depth_offset = textureGatherCompare(image_2d_depth, sampler_cmp, tc_2, 0.5, vec2(3, 1)); + return (((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset); +} + +[[stage(fragment)]] +fn depth_no_comparison() -> [[location(0)]] vec4 { + let tc_3 = vec2(0.5); + let s2d_2 = textureSample(image_2d_depth, sampler_reg, tc_3); + let s2d_gather = textureGather(image_2d_depth, sampler_reg, tc_3); + return (vec4(s2d_2) + s2d_gather); +}