github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

Gather operations

Browse Source
This commit is contained in:
Dzmitry Malyshau
2021-12-15 00:29:07 -05:00
parent db98743a5a
commit 8caa2bd87e
24 changed files with 484 additions and 41 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
src/back/dot/mod.rs
View File

@@ -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,

29
src/back/glsl/mod.rs
View File

@@ -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, ")")?
}

34
src/back/hlsl/writer.rs
View File

@@ -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(

30
src/back/msl/writer.rs
View File

@@ -896,18 +896,23 @@ impl<W: Write> Writer<W> {
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<W: Write> Writer<W> {
};
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 {

2
src/back/spv/block.rs
View File

@@ -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,

30
src/back/spv/image.rs
View File

@@ -799,6 +799,7 @@ impl<'w> BlockContext<'w> {
result_type_id: Word,
image: Handle<crate::Expression>,
sampler: Handle<crate::Expression>,
gather: Option<crate::SwizzleComponent>,
coordinate: Handle<crate::Expression>,
array_index: Option<Handle<crate::Expression>>,
offset: Option<Handle<crate::Constant>>,
@@ -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,

27
src/back/spv/instructions.rs
View File

@@ -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<Word>,
) -> 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,

49
src/back/wgsl/writer.rs
View File

@@ -1239,6 +1239,7 @@ impl<W: Write> Writer<W> {
Expression::ImageSample {
image,
sampler,
gather: None,
coordinate,
array_index,
offset,
@@ -1306,6 +1307,54 @@ impl<W: Write> Writer<W> {
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;

1
src/front/glsl/builtins.rs
View File

@@ -1930,6 +1930,7 @@ fn texture_call(
Expression::ImageSample {
image,
sampler,
gather: None, //TODO
coordinate: comps.coordinate,
array_index,
offset,

1
src/front/spv/image.rs
View File

@@ -606,6 +606,7 @@ impl<I: Iterator<Item = u32>> super::Parser<I> {
let expr = crate::Expression::ImageSample {
image: si_lexp.image,
sampler: si_lexp.sampler,
gather: None, //TODO
coordinate,
array_index,
offset,

96
src/front/wgsl/mod.rs
View File

@@ -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,

5
src/lib.rs
View File

@@ -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<Expression>,
sampler: Handle<Expression>,
/// If Some(), this operation is a gather operation
/// on the selected component.
gather: Option<SwizzleComponent>,
coordinate: Handle<Expression>,
array_index: Option<Handle<Expression>>,
offset: Option<Handle<Constant>>,

18
src/proc/typifier.rs
View File

@@ -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 {

1
src/valid/analyzer.rs
View File

@@ -462,6 +462,7 @@ impl FunctionInfo {
E::ImageSample {
image,
sampler,
gather: _,
coordinate,
array_index,
offset: _,

27
src/valid/expression.rs
View File

@@ -99,6 +99,12 @@ pub enum ExpressionError {
InvalidDepthReference(Handle<crate::Expression>),
#[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<crate::Expression>),
#[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,

2
tests/in/image.param.ron
View File

@@ -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"]
)

19
tests/in/image.wgsl
View File

@@ -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<f32> {
let tc = vec2<f32>(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<i32>(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<i32>(3, 1));
return s2d + s2d_offset + s2d_depth + s2d_depth_offset;
}
[[stage(fragment)]]
fn depth_no_comparison() -> [[location(0)]] vec4<f32> {
let tc = vec2<f32>(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;
}

22
tests/out/glsl/image.gather.Fragment.glsl Normal file
View File

@@ -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;
}

18
tests/out/hlsl/image.hlsl
View File

@@ -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);
}

2
tests/out/hlsl/image.hlsl.config
View File

@@ -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 )

1
tests/out/ir/shadow.ron
View File

@@ -870,6 +870,7 @@
ImageSample(
image: 4,
sampler: 5,
gather: None,
coordinate: 72,
array_index: Some(74),
offset: None,

32
tests/out/msl/image.msl
View File

@@ -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<float, metal::access::sample> image_2d [[user(fake0)]]
, metal::sampler sampler_reg [[user(fake0)]]
, metal::sampler sampler_cmp [[user(fake0)]]
, metal::depth2d<float, metal::access::sample> 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<float, metal::access::sample> 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 };
}

54
tests/out/spv/image.spvasm
View File

@@ -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

18
tests/out/wgsl/image.wgsl
View File

@@ -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<f32> {
let tc_2 = vec2<f32>(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<i32>(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<i32>(3, 1));
return (((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset);
}
[[stage(fragment)]]
fn depth_no_comparison() -> [[location(0)]] vec4<f32> {
let tc_3 = vec2<f32>(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<f32>(s2d_2) + s2d_gather);
}