archive-gh-me/wgpu
1
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-12-08 21:26:17 +00:00
Code Issues Projects Releases Wiki Activity

[naga] Support textureSampleBaseClampToEdge() for texture2d

Browse Source 
Adds a new flag to the IR indicating when image sample coordinates are
to be clamped. Adds wgsl-in support for parsing and lowering to
IR. Validation ensures this flag is only used when sampling a 2D
non-arrayed sampled texture, without offset, gather, or depth
comparison. This matches the WGSL requirements, with the exception of
supporting `texture_external` textures, which will follow in a later
patch.

SPIRV, HLSL, and Metal backends are supported so far, with GLSL left
for a follow up. (In GLSL the texture will simply be sampled without
the coordinates being clamped.)

It may seem unfortunate to have to handle this separately for each
backend, and indeed it would have been possible to implement this simply
in the WGSL frontend. However, future patches will add support for using
textureSampleBaseClampToEdge() with external textures, which will
actually have to be handled by each backend. This patch is laying the
groundwork for that.
This commit is contained in:
Jamie Nicol 2025-04-16 10:25:38 +01:00 committed by Jim Blandy
parent a5f328613d
commit 55a2c3095d
29 changed files with 815 additions and 478 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
naga/src/back/dot/mod.rs
View File

@ -600,6 +600,7 @@ fn write_function_expressions(
offset: _,
level,
depth_ref,
clamp_to_edge: _,
} => {
edges.insert("image", image);
edges.insert("sampler", sampler);

4
naga/src/back/glsl/mod.rs
View File

@ -3056,6 +3056,7 @@ impl<'a, W: Write> Writer<'a, W> {
offset,
level,
depth_ref,
clamp_to_edge: _,
} => {
let (dim, class, arrayed) = match *ctx.resolve_type(image, &self.module.types) {
TypeInner::Image {
@ -3114,6 +3115,9 @@ impl<'a, W: Write> Writer<'a, W> {
// The space here isn't required but it helps with readability
write!(self.out, ", ")?;
// TODO: handle clamp_to_edge
// https://github.com/gfx-rs/wgpu/issues/7791
// We need to get the coordinates vector size to later build a vector that's `size + 1`
// if `depth_ref` is some, if it isn't a vector we panic as that's not a valid expression
let mut coord_dim = match *ctx.resolve_type(coordinate, &self.module.types) {

40
naga/src/back/hlsl/help.rs
View File

@ -33,7 +33,8 @@ use super::{
super::FunctionCtx,
writer::{
ABS_FUNCTION, DIV_FUNCTION, EXTRACT_BITS_FUNCTION, F2I32_FUNCTION, F2I64_FUNCTION,
F2U32_FUNCTION, F2U64_FUNCTION, INSERT_BITS_FUNCTION, MOD_FUNCTION, NEG_FUNCTION,
F2U32_FUNCTION, F2U64_FUNCTION, IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION,
INSERT_BITS_FUNCTION, MOD_FUNCTION, NEG_FUNCTION,
},
BackendResult, WrappedType,
};
@ -44,6 +45,11 @@ pub(super) struct WrappedArrayLength {
pub(super) writable: bool,
}
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
pub(super) struct WrappedImageSample {
pub(super) clamp_to_edge: bool,
}
#[derive(Clone, Copy, Debug, Hash, Eq, Ord, PartialEq, PartialOrd)]
pub(super) struct WrappedImageQuery {
pub(super) dim: crate::ImageDimension,
@ -247,6 +253,32 @@ impl<W: Write> super::Writer<'_, W> {
Ok(())
}
pub(super) fn write_wrapped_image_sample_function(
&mut self,
sample: WrappedImageSample,
) -> BackendResult {
match sample {
WrappedImageSample {
clamp_to_edge: true,
} => {
writeln!(self.out, "float4 {IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION}(Texture2D<float4> tex, SamplerState samp, float2 coords) {{")?;
let l1 = crate::back::Level(1);
writeln!(self.out, "{l1}float2 size;")?;
writeln!(self.out, "{l1}tex.GetDimensions(size.x, size.y);")?;
writeln!(self.out, "{l1}float2 half_texel = float2(0.5, 0.5) / size;")?;
writeln!(
self.out,
"{l1}return tex.SampleLevel(samp, clamp(coords, half_texel, 1.0 - half_texel), 0.0);"
)?;
writeln!(self.out, "}}")?;
writeln!(self.out)?;
}
_ => {}
}
Ok(())
}
pub(super) fn write_wrapped_image_query_function_name(
&mut self,
query: WrappedImageQuery,
@ -1522,6 +1554,12 @@ impl<W: Write> super::Writer<'_, W> {
self.write_wrapped_array_length_function(wal)?;
}
}
crate::Expression::ImageSample { clamp_to_edge, .. } => {
let wrapped = WrappedImageSample { clamp_to_edge };
if self.wrapped.insert(WrappedType::ImageSample(wrapped)) {
self.write_wrapped_image_sample_function(wrapped)?;
}
}
crate::Expression::ImageQuery { image, query } => {
let wiq = match *func_ctx.resolve_type(image, &module.types) {
crate::TypeInner::Image {

1
naga/src/back/hlsl/keywords.rs
View File

@ -834,6 +834,7 @@ pub const RESERVED: &[&str] = &[
super::writer::F2U32_FUNCTION,
super::writer::F2I64_FUNCTION,
super::writer::F2U64_FUNCTION,
super::writer::IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION,
];
// DXC scalar types, from https://github.com/microsoft/DirectXShaderCompiler/blob/18c9e114f9c314f93e68fbc72ce207d4ed2e65ae/tools/clang/lib/AST/ASTContextHLSL.cpp#L48-L254

1
naga/src/back/hlsl/mod.rs
View File

@ -472,6 +472,7 @@ pub enum Error {
enum WrappedType {
ZeroValue(help::WrappedZeroValue),
ArrayLength(help::WrappedArrayLength),
ImageSample(help::WrappedImageSample),
ImageQuery(help::WrappedImageQuery),
ImageLoadScalar(crate::Scalar),
Constructor(help::WrappedConstructor),

28
naga/src/back/hlsl/writer.rs
View File

@ -42,6 +42,8 @@ pub(crate) const F2I32_FUNCTION: &str = "naga_f2i32";
pub(crate) const F2U32_FUNCTION: &str = "naga_f2u32";
pub(crate) const F2I64_FUNCTION: &str = "naga_f2i64";
pub(crate) const F2U64_FUNCTION: &str = "naga_f2u64";
pub(crate) const IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION: &str =
"nagaTextureSampleBaseClampToEdge";
struct EpStructMember {
name: String,
@ -3183,6 +3185,25 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
let name = &self.names[&key];
write!(self.out, "{name}")?;
}
Expression::ImageSample {
coordinate,
image,
sampler,
clamp_to_edge: true,
gather: None,
array_index: None,
offset: None,
level: crate::SampleLevel::Zero,
depth_ref: None,
} => {
write!(self.out, "{IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION}(")?;
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)?;
write!(self.out, ")")?;
}
Expression::ImageSample {
image,
sampler,
@ -3192,7 +3213,14 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
offset,
level,
depth_ref,
clamp_to_edge,
} => {
if clamp_to_edge {
return Err(Error::Custom(
"ImageSample::clamp_to_edge should have been validated out".to_string(),
));
}
use crate::SampleLevel as Sl;
const COMPONENTS: [&str; 4] = ["", "Green", "Blue", "Alpha"];

1
naga/src/back/msl/keywords.rs
View File

@ -353,6 +353,7 @@ pub const RESERVED: &[&str] = &[
super::writer::F2U32_FUNCTION,
super::writer::F2I64_FUNCTION,
super::writer::F2U64_FUNCTION,
super::writer::IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION,
super::writer::ARGUMENT_BUFFER_WRAPPER_STRUCT,
];

52
naga/src/back/msl/writer.rs
View File

@ -61,6 +61,8 @@ pub(crate) const F2I32_FUNCTION: &str = "naga_f2i32";
pub(crate) const F2U32_FUNCTION: &str = "naga_f2u32";
pub(crate) const F2I64_FUNCTION: &str = "naga_f2i64";
pub(crate) const F2U64_FUNCTION: &str = "naga_f2u64";
pub(crate) const IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION: &str =
"nagaTextureSampleBaseClampToEdge";
/// For some reason, Metal does not let you have `metal::texture<..>*` as a buffer argument.
/// However, if you put that texture inside a struct, everything is totally fine. This
/// baffles me to no end.
@ -446,6 +448,9 @@ enum WrappedFunction {
vector_size: Option<crate::VectorSize>,
dst_scalar: crate::Scalar,
},
ImageSample {
clamp_to_edge: bool,
},
}
pub struct Writer<W> {
@ -1940,6 +1945,25 @@ impl<W: Write> Writer<W> {
write!(self.out, "{name}")?;
}
crate::Expression::Load { pointer } => self.put_load(pointer, context, is_scoped)?,
crate::Expression::ImageSample {
coordinate,
image,
sampler,
clamp_to_edge: true,
gather: None,
array_index: None,
offset: None,
level: crate::SampleLevel::Zero,
depth_ref: None,
} => {
write!(self.out, "{IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION}(")?;
self.put_expression(image, context, true)?;
write!(self.out, ", ")?;
self.put_expression(sampler, context, true)?;
write!(self.out, ", ")?;
self.put_expression(coordinate, context, true)?;
write!(self.out, ")")?;
}
crate::Expression::ImageSample {
image,
sampler,
@ -1949,7 +1973,14 @@ impl<W: Write> Writer<W> {
offset,
level,
depth_ref,
clamp_to_edge,
} => {
if clamp_to_edge {
return Err(Error::GenericValidation(
"ImageSample::clamp_to_edge should have been validated out".to_string(),
));
}
let main_op = match gather {
Some(_) => "gather",
None => "sample",
@ -5830,6 +5861,27 @@ template <typename A>
writeln!(self.out, "}}")?;
writeln!(self.out)?;
}
crate::Expression::ImageSample {
clamp_to_edge: true,
..
} => {
let wrapped = WrappedFunction::ImageSample {
clamp_to_edge: true,
};
if !self.wrapped_functions.insert(wrapped) {
continue;
}
writeln!(self.out, "{NAMESPACE}::float4 {IMAGE_SAMPLE_BASE_CLAMP_TO_EDGE_FUNCTION}({NAMESPACE}::texture2d<float, {NAMESPACE}::access::sample> tex, {NAMESPACE}::sampler samp, {NAMESPACE}::float2 coords) {{")?;
let l1 = back::Level(1);
writeln!(self.out, "{l1}{NAMESPACE}::float2 half_texel = 0.5 / {NAMESPACE}::float2(tex.get_width(0u), tex.get_height(0u));")?;
writeln!(
self.out,
"{l1}return tex.sample(samp, {NAMESPACE}::clamp(coords, half_texel, 1.0 - half_texel), {NAMESPACE}::level(0.0));"
)?;
writeln!(self.out, "}}")?;
writeln!(self.out)?;
}
_ => {}
}
}

1
naga/src/back/pipeline_constants.rs
View File

@ -468,6 +468,7 @@ fn adjust_expr(new_pos: &HandleVec<Expression, Handle<Expression>>, expr: &mut E
ref mut level,
ref mut depth_ref,
gather: _,
clamp_to_edge: _,
} => {
adjust(image);
adjust(sampler);

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

@ -1661,6 +1661,7 @@ impl BlockContext<'_> {
offset,
level,
depth_ref,
clamp_to_edge,
} => self.write_image_sample(
result_type_id,
image,
@ -1671,6 +1672,7 @@ impl BlockContext<'_> {
offset,
level,
depth_ref,
clamp_to_edge,
block,
)?,
crate::Expression::Select {

97
naga/src/back/spv/image.rs
View File

@ -810,6 +810,7 @@ impl BlockContext<'_> {
offset: Option<Handle<crate::Expression>>,
level: crate::SampleLevel,
depth_ref: Option<Handle<crate::Expression>>,
clamp_to_edge: bool,
block: &mut Block,
) -> Result<Word, Error> {
use super::instructions::SampleLod;
@ -840,9 +841,99 @@ impl BlockContext<'_> {
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 = self.write_image_coordinates(coordinate, array_index, block)?;
let coordinates_id = if clamp_to_edge {
self.writer.require_any(
"clamp sample coordinates to edge",
&[spirv::Capability::ImageQuery],
)?;
// clamp_to_edge can only be used with Level 0, and no array offset, offset,
// depth_ref or gather. This should have been caught by validation. Rather
// than entirely duplicate validation code here just ensure the level is
// zero, as we rely on that to query the texture size in order to calculate
// the clamped coordinates.
if level != crate::SampleLevel::Zero {
return Err(Error::Validation(
"ImageSample::clamp_to_edge requires SampleLevel::Zero",
));
}
// Query the size of level 0 of the texture.
let image_size_id = self.gen_id();
let vec2u_type_id = self.writer.get_vec2u_type_id();
let const_zero_uint_id = self.writer.get_constant_scalar(crate::Literal::U32(0));
let mut query_inst = Instruction::image_query(
spirv::Op::ImageQuerySizeLod,
vec2u_type_id,
image_size_id,
image_id,
);
query_inst.add_operand(const_zero_uint_id);
block.body.push(query_inst);
let image_size_f_id = self.gen_id();
let vec2f_type_id = self.writer.get_vec2f_type_id();
block.body.push(Instruction::unary(
spirv::Op::ConvertUToF,
vec2f_type_id,
image_size_f_id,
image_size_id,
));
// Calculate the top-left and bottom-right margin for clamping to. I.e. a
// half-texel from each side.
let const_0_5_f32_id = self.writer.get_constant_scalar(crate::Literal::F32(0.5));
let const_0_5_vec2f_id = self.writer.get_constant_composite(
LookupType::Local(LocalType::Numeric(NumericType::Vector {
size: crate::VectorSize::Bi,
scalar: crate::Scalar::F32,
})),
&[const_0_5_f32_id, const_0_5_f32_id],
);
let margin_left_id = self.gen_id();
block.body.push(Instruction::binary(
spirv::Op::FDiv,
vec2f_type_id,
margin_left_id,
const_0_5_vec2f_id,
image_size_f_id,
));
let const_1_f32_id = self.writer.get_constant_scalar(crate::Literal::F32(1.0));
let const_1_vec2f_id = self.writer.get_constant_composite(
LookupType::Local(LocalType::Numeric(NumericType::Vector {
size: crate::VectorSize::Bi,
scalar: crate::Scalar::F32,
})),
&[const_1_f32_id, const_1_f32_id],
);
let margin_right_id = self.gen_id();
block.body.push(Instruction::binary(
spirv::Op::FSub,
vec2f_type_id,
margin_right_id,
const_1_vec2f_id,
margin_left_id,
));
// Clamp the coords to the calculated margins
let clamped_coords_id = self.gen_id();
block.body.push(Instruction::ext_inst(
self.writer.gl450_ext_inst_id,
spirv::GLOp::NClamp,
vec2f_type_id,
clamped_coords_id,
&[coordinates.value_id, margin_left_id, margin_right_id],
));
clamped_coords_id
} else {
coordinates.value_id
};
let sampled_image_id = self.gen_id();
block.body.push(Instruction::sampled_image(

7
naga/src/back/spv/writer.rs
View File

@ -338,6 +338,13 @@ impl Writer {
})
}
pub(super) fn get_vec2f_type_id(&mut self) -> Word {
self.get_numeric_type_id(NumericType::Vector {
size: crate::VectorSize::Bi,
scalar: crate::Scalar::F32,
})
}
pub(super) fn get_vec3u_type_id(&mut self) -> Word {
self.get_numeric_type_id(NumericType::Vector {
size: crate::VectorSize::Tri,

7
naga/src/back/wgsl/writer.rs
View File

@ -1303,6 +1303,7 @@ impl<W: Write> Writer<W> {
offset,
level,
depth_ref,
clamp_to_edge,
} => {
use crate::SampleLevel as Sl;
@ -1312,6 +1313,7 @@ impl<W: Write> Writer<W> {
};
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",
@ -1337,8 +1339,8 @@ impl<W: Write> Writer<W> {
match level {
Sl::Auto => {}
Sl::Zero => {
// Level 0 is implied for depth comparison
if depth_ref.is_none() {
// Level 0 is implied for depth comparison and BaseClampToEdge
if depth_ref.is_none() && !clamp_to_edge {
write!(self.out, ", 0.0")?;
}
}
@ -1375,6 +1377,7 @@ impl<W: Write> Writer<W> {
offset,
level: _,
depth_ref,
clamp_to_edge: _,
} => {
let suffix_cmp = match depth_ref {
Some(_) => "Compare",

2
naga/src/compact/expressions.rs
View File

@ -150,6 +150,7 @@ impl ExpressionTracer<'_> {
offset,
ref level,
depth_ref,
clamp_to_edge: _,
} => {
self.expressions_used
.insert_iter([image, sampler, coordinate]);
@ -323,6 +324,7 @@ impl ModuleMap {
ref mut offset,
ref mut level,
ref mut depth_ref,
clamp_to_edge: _,
} => {
adjust(image);
adjust(sampler);

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

@ -2085,6 +2085,7 @@ fn texture_call(
offset,
level,
depth_ref: comps.depth_ref,
clamp_to_edge: false,
},
meta,
)?)

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

@ -754,6 +754,7 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
offset,
level,
depth_ref,
clamp_to_edge: false,
};
let image_sample_handle = ctx.expressions.append(expr, self.span_from_with_op(start));
let handle = if is_depth && depth_ref.is_none() {

12
naga/src/front/wgsl/lower/mod.rs
View File

@ -1040,7 +1040,7 @@ enum Texture {
SampleCompareLevel,
SampleGrad,
SampleLevel,
// SampleBaseClampToEdge,
SampleBaseClampToEdge,
}
impl Texture {
@ -1055,7 +1055,7 @@ impl Texture {
"textureSampleCompareLevel" => Self::SampleCompareLevel,
"textureSampleGrad" => Self::SampleGrad,
"textureSampleLevel" => Self::SampleLevel,
// "textureSampleBaseClampToEdge" => Some(Self::SampleBaseClampToEdge),
"textureSampleBaseClampToEdge" => Self::SampleBaseClampToEdge,
_ => return None,
})
}
@ -1071,7 +1071,7 @@ impl Texture {
Self::SampleCompareLevel => 5,
Self::SampleGrad => 6,
Self::SampleLevel => 5,
// Self::SampleBaseClampToEdge => 3,
Self::SampleBaseClampToEdge => 3,
}
}
}
@ -3527,6 +3527,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
let sampler = self.expression_for_abstract(args.next()?, ctx)?;
let coordinate = self.expression_with_leaf_scalar(args.next()?, ir::Scalar::F32, ctx)?;
let clamp_to_edge = matches!(fun, Texture::SampleBaseClampToEdge);
let (class, arrayed) = ctx.image_data(image, image_span)?;
let array_index = arrayed
@ -3593,6 +3594,10 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
level = ir::SampleLevel::Exact(exact);
depth_ref = None;
}
Texture::SampleBaseClampToEdge => {
level = crate::SampleLevel::Zero;
depth_ref = None;
}
};
let offset = args
@ -3612,6 +3617,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
offset,
level,
depth_ref,
clamp_to_edge,
})
}

4
naga/src/ir/mod.rs
View File

@ -1509,6 +1509,10 @@ pub enum Expression {
offset: Option<Handle<Expression>>,
level: SampleLevel,
depth_ref: Option<Handle<Expression>>,
/// Whether the sampling operation should clamp each component of
/// `coordinate` to the range `[half_texel, 1 - half_texel]`, regardless
/// of `sampler`.
clamp_to_edge: bool,
},
/// Load a texel from an image.

1
naga/src/valid/analyzer.rs
View File

@ -664,6 +664,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)?;

49
naga/src/valid/expression.rs
View File

@ -123,6 +123,8 @@ pub enum ExpressionError {
InvalidSampleLevelBiasDimension(crate::ImageDimension),
#[error("Sample level (gradient) of {1:?} doesn't match the image dimension {0:?}")]
InvalidSampleLevelGradientType(crate::ImageDimension, Handle<crate::Expression>),
#[error("Clamping sample coordinate to edge is not supported with {0}")]
InvalidSampleClampCoordinateToEdge(alloc::string::String),
#[error("Unable to cast")]
InvalidCastArgument,
#[error("Invalid argument count for {0:?}")]
@ -413,6 +415,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)?;
@ -540,6 +543,52 @@ impl super::Validator {
}
}
// Clamping coordinate to edge is only supported with 2d non-arrayed, sampled images
// when sampling from level Zero without any offset, gather, or depth comparison.
if clamp_to_edge {
if !matches!(
class,
crate::ImageClass::Sampled {
kind: crate::ScalarKind::Float,
multi: false
}
) {
return Err(ExpressionError::InvalidSampleClampCoordinateToEdge(
alloc::format!("image class `{class:?}`"),
));
}
if dim != crate::ImageDimension::D2 {
return Err(ExpressionError::InvalidSampleClampCoordinateToEdge(
alloc::format!("image dimension `{dim:?}`"),
));
}
if gather.is_some() {
return Err(ExpressionError::InvalidSampleClampCoordinateToEdge(
"gather".into(),
));
}
if array_index.is_some() {
return Err(ExpressionError::InvalidSampleClampCoordinateToEdge(
"array index".into(),
));
}
if offset.is_some() {
return Err(ExpressionError::InvalidSampleClampCoordinateToEdge(
"offset".into(),
));
}
if level != crate::SampleLevel::Zero {
return Err(ExpressionError::InvalidSampleClampCoordinateToEdge(
"non-zero level".into(),
));
}
if depth_ref.is_some() {
return Err(ExpressionError::InvalidSampleClampCoordinateToEdge(
"depth comparison".into(),
));
}
}
// check level properties
match level {
crate::SampleLevel::Auto => ShaderStages::FRAGMENT,

1
naga/src/valid/handles.rs
View File

@ -529,6 +529,7 @@ impl super::Validator {
offset,
level,
depth_ref,
clamp_to_edge: _,
} => {
handle
.check_dep(image)?

1
naga/tests/in/wgsl/image.wgsl
View File

@ -127,6 +127,7 @@ fn texture_sample() -> @location(0) vec4<f32> {
a += textureSampleLevel(image_2d, sampler_reg, tc, level);
a += textureSampleLevel(image_2d, sampler_reg, tc, level, offset);
a += textureSampleBias(image_2d, sampler_reg, tc, 2.0, offset);
a += textureSampleBaseClampToEdge(image_2d, sampler_reg, tc);
a += textureSample(image_2d_array, sampler_reg, tc, 0u);
a += textureSample(image_2d_array, sampler_reg, tc, 0u, offset);
a += textureSampleLevel(image_2d_array, sampler_reg, tc, 0u, level);

103
naga/tests/out/glsl/wgsl-image.texture_sample.Fragment.glsl
View File

@ -32,56 +32,59 @@ void main() {
vec4 _e40 = textureOffset(_group_0_binding_1_fs, vec2(_e1), ivec2(3, 1), 2.0);
vec4 _e41 = a;
a = (_e41 + _e40);
vec4 _e46 = texture(_group_0_binding_4_fs, vec3(_e1, 0u));
vec4 _e47 = a;
a = (_e47 + _e46);
vec4 _e52 = textureOffset(_group_0_binding_4_fs, vec3(_e1, 0u), ivec2(3, 1));
vec4 _e53 = a;
a = (_e53 + _e52);
vec4 _e58 = textureLod(_group_0_binding_4_fs, vec3(_e1, 0u), 2.3);
vec4 _e59 = a;
a = (_e59 + _e58);
vec4 _e64 = textureLodOffset(_group_0_binding_4_fs, vec3(_e1, 0u), 2.3, ivec2(3, 1));
vec4 _e65 = a;
a = (_e65 + _e64);
vec4 _e71 = textureOffset(_group_0_binding_4_fs, vec3(_e1, 0u), ivec2(3, 1), 2.0);
vec4 _e72 = a;
a = (_e72 + _e71);
vec4 _e77 = texture(_group_0_binding_4_fs, vec3(_e1, 0));
vec4 _e78 = a;
a = (_e78 + _e77);
vec4 _e83 = textureOffset(_group_0_binding_4_fs, vec3(_e1, 0), ivec2(3, 1));
vec4 _e84 = a;
a = (_e84 + _e83);
vec4 _e89 = textureLod(_group_0_binding_4_fs, vec3(_e1, 0), 2.3);
vec4 _e90 = a;
a = (_e90 + _e89);
vec4 _e95 = textureLodOffset(_group_0_binding_4_fs, vec3(_e1, 0), 2.3, ivec2(3, 1));
vec4 _e96 = a;
a = (_e96 + _e95);
vec4 _e102 = textureOffset(_group_0_binding_4_fs, vec3(_e1, 0), ivec2(3, 1), 2.0);
vec4 _e103 = a;
a = (_e103 + _e102);
vec4 _e108 = texture(_group_0_binding_6_fs, vec4(_e3, 0u));
vec4 _e109 = a;
a = (_e109 + _e108);
vec4 _e114 = textureLod(_group_0_binding_6_fs, vec4(_e3, 0u), 2.3);
vec4 _e115 = a;
a = (_e115 + _e114);
vec4 _e121 = texture(_group_0_binding_6_fs, vec4(_e3, 0u), 2.0);
vec4 _e122 = a;
a = (_e122 + _e121);
vec4 _e127 = texture(_group_0_binding_6_fs, vec4(_e3, 0));
vec4 _e128 = a;
a = (_e128 + _e127);
vec4 _e133 = textureLod(_group_0_binding_6_fs, vec4(_e3, 0), 2.3);
vec4 _e134 = a;
a = (_e134 + _e133);
vec4 _e140 = texture(_group_0_binding_6_fs, vec4(_e3, 0), 2.0);
vec4 _e141 = a;
a = (_e141 + _e140);
vec4 _e143 = a;
_fs2p_location0 = _e143;
vec4 _e45 = textureLod(_group_0_binding_1_fs, vec2(_e1), 0.0);
vec4 _e46 = a;
a = (_e46 + _e45);
vec4 _e51 = texture(_group_0_binding_4_fs, vec3(_e1, 0u));
vec4 _e52 = a;
a = (_e52 + _e51);
vec4 _e57 = textureOffset(_group_0_binding_4_fs, vec3(_e1, 0u), ivec2(3, 1));
vec4 _e58 = a;
a = (_e58 + _e57);
vec4 _e63 = textureLod(_group_0_binding_4_fs, vec3(_e1, 0u), 2.3);
vec4 _e64 = a;
a = (_e64 + _e63);
vec4 _e69 = textureLodOffset(_group_0_binding_4_fs, vec3(_e1, 0u), 2.3, ivec2(3, 1));
vec4 _e70 = a;
a = (_e70 + _e69);
vec4 _e76 = textureOffset(_group_0_binding_4_fs, vec3(_e1, 0u), ivec2(3, 1), 2.0);
vec4 _e77 = a;
a = (_e77 + _e76);
vec4 _e82 = texture(_group_0_binding_4_fs, vec3(_e1, 0));
vec4 _e83 = a;
a = (_e83 + _e82);
vec4 _e88 = textureOffset(_group_0_binding_4_fs, vec3(_e1, 0), ivec2(3, 1));
vec4 _e89 = a;
a = (_e89 + _e88);
vec4 _e94 = textureLod(_group_0_binding_4_fs, vec3(_e1, 0), 2.3);
vec4 _e95 = a;
a = (_e95 + _e94);
vec4 _e100 = textureLodOffset(_group_0_binding_4_fs, vec3(_e1, 0), 2.3, ivec2(3, 1));
vec4 _e101 = a;
a = (_e101 + _e100);
vec4 _e107 = textureOffset(_group_0_binding_4_fs, vec3(_e1, 0), ivec2(3, 1), 2.0);
vec4 _e108 = a;
a = (_e108 + _e107);
vec4 _e113 = texture(_group_0_binding_6_fs, vec4(_e3, 0u));
vec4 _e114 = a;
a = (_e114 + _e113);
vec4 _e119 = textureLod(_group_0_binding_6_fs, vec4(_e3, 0u), 2.3);
vec4 _e120 = a;
a = (_e120 + _e119);
vec4 _e126 = texture(_group_0_binding_6_fs, vec4(_e3, 0u), 2.0);
vec4 _e127 = a;
a = (_e127 + _e126);
vec4 _e132 = texture(_group_0_binding_6_fs, vec4(_e3, 0));
vec4 _e133 = a;
a = (_e133 + _e132);
vec4 _e138 = textureLod(_group_0_binding_6_fs, vec4(_e3, 0), 2.3);
vec4 _e139 = a;
a = (_e139 + _e138);
vec4 _e145 = texture(_group_0_binding_6_fs, vec4(_e3, 0), 2.0);
vec4 _e146 = a;
a = (_e146 + _e145);
vec4 _e148 = a;
_fs2p_location0 = _e148;
return;
}

110
naga/tests/out/hlsl/wgsl-image.hlsl
View File

@ -258,6 +258,13 @@ float4 levels_queries() : SV_Position
return (float(sum_1)).xxxx;
}
float4 nagaTextureSampleBaseClampToEdge(Texture2D<float4> tex, SamplerState samp, float2 coords) {
float2 size;
tex.GetDimensions(size.x, size.y);
float2 half_texel = float2(0.5, 0.5) / size;
return tex.SampleLevel(samp, clamp(coords, half_texel, 1.0 - half_texel), 0.0);
}
float4 texture_sample() : SV_Target0
{
float4 a = (float4)0;
@ -283,56 +290,59 @@ float4 texture_sample() : SV_Target0
float4 _e40 = image_2d.SampleBias(sampler_reg, _e1, 2.0, int2(int2(int(3), int(1))));
float4 _e41 = a;
a = (_e41 + _e40);
float4 _e46 = image_2d_array.Sample(sampler_reg, float3(_e1, 0u));
float4 _e47 = a;
a = (_e47 + _e46);
float4 _e52 = image_2d_array.Sample(sampler_reg, float3(_e1, 0u), int2(int2(int(3), int(1))));
float4 _e53 = a;
a = (_e53 + _e52);
float4 _e58 = image_2d_array.SampleLevel(sampler_reg, float3(_e1, 0u), 2.3);
float4 _e59 = a;
a = (_e59 + _e58);
float4 _e64 = image_2d_array.SampleLevel(sampler_reg, float3(_e1, 0u), 2.3, int2(int2(int(3), int(1))));
float4 _e65 = a;
a = (_e65 + _e64);
float4 _e71 = image_2d_array.SampleBias(sampler_reg, float3(_e1, 0u), 2.0, int2(int2(int(3), int(1))));
float4 _e72 = a;
a = (_e72 + _e71);
float4 _e77 = image_2d_array.Sample(sampler_reg, float3(_e1, int(0)));
float4 _e78 = a;
a = (_e78 + _e77);
float4 _e83 = image_2d_array.Sample(sampler_reg, float3(_e1, int(0)), int2(int2(int(3), int(1))));
float4 _e84 = a;
a = (_e84 + _e83);
float4 _e89 = image_2d_array.SampleLevel(sampler_reg, float3(_e1, int(0)), 2.3);
float4 _e90 = a;
a = (_e90 + _e89);
float4 _e95 = image_2d_array.SampleLevel(sampler_reg, float3(_e1, int(0)), 2.3, int2(int2(int(3), int(1))));
float4 _e96 = a;
a = (_e96 + _e95);
float4 _e102 = image_2d_array.SampleBias(sampler_reg, float3(_e1, int(0)), 2.0, int2(int2(int(3), int(1))));
float4 _e103 = a;
a = (_e103 + _e102);
float4 _e108 = image_cube_array.Sample(sampler_reg, float4(_e3, 0u));
float4 _e109 = a;
a = (_e109 + _e108);
float4 _e114 = image_cube_array.SampleLevel(sampler_reg, float4(_e3, 0u), 2.3);
float4 _e115 = a;
a = (_e115 + _e114);
float4 _e121 = image_cube_array.SampleBias(sampler_reg, float4(_e3, 0u), 2.0);
float4 _e122 = a;
a = (_e122 + _e121);
float4 _e127 = image_cube_array.Sample(sampler_reg, float4(_e3, int(0)));
float4 _e128 = a;
a = (_e128 + _e127);
float4 _e133 = image_cube_array.SampleLevel(sampler_reg, float4(_e3, int(0)), 2.3);
float4 _e134 = a;
a = (_e134 + _e133);
float4 _e140 = image_cube_array.SampleBias(sampler_reg, float4(_e3, int(0)), 2.0);
float4 _e141 = a;
a = (_e141 + _e140);
float4 _e143 = a;
return _e143;
float4 _e45 = nagaTextureSampleBaseClampToEdge(image_2d, sampler_reg, _e1);
float4 _e46 = a;
a = (_e46 + _e45);
float4 _e51 = image_2d_array.Sample(sampler_reg, float3(_e1, 0u));
float4 _e52 = a;
a = (_e52 + _e51);
float4 _e57 = image_2d_array.Sample(sampler_reg, float3(_e1, 0u), int2(int2(int(3), int(1))));
float4 _e58 = a;
a = (_e58 + _e57);
float4 _e63 = image_2d_array.SampleLevel(sampler_reg, float3(_e1, 0u), 2.3);
float4 _e64 = a;
a = (_e64 + _e63);
float4 _e69 = image_2d_array.SampleLevel(sampler_reg, float3(_e1, 0u), 2.3, int2(int2(int(3), int(1))));
float4 _e70 = a;
a = (_e70 + _e69);
float4 _e76 = image_2d_array.SampleBias(sampler_reg, float3(_e1, 0u), 2.0, int2(int2(int(3), int(1))));
float4 _e77 = a;
a = (_e77 + _e76);
float4 _e82 = image_2d_array.Sample(sampler_reg, float3(_e1, int(0)));
float4 _e83 = a;
a = (_e83 + _e82);
float4 _e88 = image_2d_array.Sample(sampler_reg, float3(_e1, int(0)), int2(int2(int(3), int(1))));
float4 _e89 = a;
a = (_e89 + _e88);
float4 _e94 = image_2d_array.SampleLevel(sampler_reg, float3(_e1, int(0)), 2.3);
float4 _e95 = a;
a = (_e95 + _e94);
float4 _e100 = image_2d_array.SampleLevel(sampler_reg, float3(_e1, int(0)), 2.3, int2(int2(int(3), int(1))));
float4 _e101 = a;
a = (_e101 + _e100);
float4 _e107 = image_2d_array.SampleBias(sampler_reg, float3(_e1, int(0)), 2.0, int2(int2(int(3), int(1))));
float4 _e108 = a;
a = (_e108 + _e107);
float4 _e113 = image_cube_array.Sample(sampler_reg, float4(_e3, 0u));
float4 _e114 = a;
a = (_e114 + _e113);
float4 _e119 = image_cube_array.SampleLevel(sampler_reg, float4(_e3, 0u), 2.3);
float4 _e120 = a;
a = (_e120 + _e119);
float4 _e126 = image_cube_array.SampleBias(sampler_reg, float4(_e3, 0u), 2.0);
float4 _e127 = a;
a = (_e127 + _e126);
float4 _e132 = image_cube_array.Sample(sampler_reg, float4(_e3, int(0)));
float4 _e133 = a;
a = (_e133 + _e132);
float4 _e138 = image_cube_array.SampleLevel(sampler_reg, float4(_e3, int(0)), 2.3);
float4 _e139 = a;
a = (_e139 + _e138);
float4 _e145 = image_cube_array.SampleBias(sampler_reg, float4(_e3, int(0)), 2.0);
float4 _e146 = a;
a = (_e146 + _e145);
float4 _e148 = a;
return _e148;
}
float texture_sample_comparison() : SV_Target0

1
naga/tests/out/ir/spv-shadow.compact.ron
View File

@ -516,6 +516,7 @@
offset: None,
level: Zero,
depth_ref: Some(28),
clamp_to_edge: false,
),
],
named_expressions: {},

1
naga/tests/out/ir/spv-shadow.ron
View File

@ -771,6 +771,7 @@
offset: None,
level: Zero,
depth_ref: Some(64),
clamp_to_edge: false,
),
],
named_expressions: {},

108
naga/tests/out/msl/wgsl-image.msl
View File

@ -120,6 +120,11 @@ vertex levels_queriesOutput levels_queries(
return levels_queriesOutput { metal::float4(static_cast<float>(sum_1)) };
}
metal::float4 nagaTextureSampleBaseClampToEdge(metal::texture2d<float, metal::access::sample> tex, metal::sampler samp, metal::float2 coords) {
metal::float2 half_texel = 0.5 / metal::float2(tex.get_width(0u), tex.get_height(0u));
return tex.sample(samp, metal::clamp(coords, half_texel, 1.0 - half_texel), metal::level(0.0));
}
struct texture_sampleOutput {
metal::float4 member_4 [[color(0)]];
@ -153,56 +158,59 @@ fragment texture_sampleOutput texture_sample(
metal::float4 _e40 = image_2d.sample(sampler_reg, _e1, metal::bias(2.0), _e6);
metal::float4 _e41 = a;
a = _e41 + _e40;
metal::float4 _e46 = image_2d_array.sample(sampler_reg, _e1, 0u);
metal::float4 _e47 = a;
a = _e47 + _e46;
metal::float4 _e52 = image_2d_array.sample(sampler_reg, _e1, 0u, _e6);
metal::float4 _e53 = a;
a = _e53 + _e52;
metal::float4 _e58 = image_2d_array.sample(sampler_reg, _e1, 0u, metal::level(2.3));
metal::float4 _e59 = a;
a = _e59 + _e58;
metal::float4 _e64 = image_2d_array.sample(sampler_reg, _e1, 0u, metal::level(2.3), _e6);
metal::float4 _e65 = a;
a = _e65 + _e64;
metal::float4 _e71 = image_2d_array.sample(sampler_reg, _e1, 0u, metal::bias(2.0), _e6);
metal::float4 _e72 = a;
a = _e72 + _e71;
metal::float4 _e77 = image_2d_array.sample(sampler_reg, _e1, 0);
metal::float4 _e78 = a;
a = _e78 + _e77;
metal::float4 _e83 = image_2d_array.sample(sampler_reg, _e1, 0, _e6);
metal::float4 _e84 = a;
a = _e84 + _e83;
metal::float4 _e89 = image_2d_array.sample(sampler_reg, _e1, 0, metal::level(2.3));
metal::float4 _e90 = a;
a = _e90 + _e89;
metal::float4 _e95 = image_2d_array.sample(sampler_reg, _e1, 0, metal::level(2.3), _e6);
metal::float4 _e96 = a;
a = _e96 + _e95;
metal::float4 _e102 = image_2d_array.sample(sampler_reg, _e1, 0, metal::bias(2.0), _e6);
metal::float4 _e103 = a;
a = _e103 + _e102;
metal::float4 _e108 = image_cube_array.sample(sampler_reg, _e3, 0u);
metal::float4 _e109 = a;
a = _e109 + _e108;
metal::float4 _e114 = image_cube_array.sample(sampler_reg, _e3, 0u, metal::level(2.3));
metal::float4 _e115 = a;
a = _e115 + _e114;
metal::float4 _e121 = image_cube_array.sample(sampler_reg, _e3, 0u, metal::bias(2.0));
metal::float4 _e122 = a;
a = _e122 + _e121;
metal::float4 _e127 = image_cube_array.sample(sampler_reg, _e3, 0);
metal::float4 _e128 = a;
a = _e128 + _e127;
metal::float4 _e133 = image_cube_array.sample(sampler_reg, _e3, 0, metal::level(2.3));
metal::float4 _e134 = a;
a = _e134 + _e133;
metal::float4 _e140 = image_cube_array.sample(sampler_reg, _e3, 0, metal::bias(2.0));
metal::float4 _e141 = a;
a = _e141 + _e140;
metal::float4 _e143 = a;
return texture_sampleOutput { _e143 };
metal::float4 _e45 = nagaTextureSampleBaseClampToEdge(image_2d, sampler_reg, _e1);
metal::float4 _e46 = a;
a = _e46 + _e45;
metal::float4 _e51 = image_2d_array.sample(sampler_reg, _e1, 0u);
metal::float4 _e52 = a;
a = _e52 + _e51;
metal::float4 _e57 = image_2d_array.sample(sampler_reg, _e1, 0u, _e6);
metal::float4 _e58 = a;
a = _e58 + _e57;
metal::float4 _e63 = image_2d_array.sample(sampler_reg, _e1, 0u, metal::level(2.3));
metal::float4 _e64 = a;
a = _e64 + _e63;
metal::float4 _e69 = image_2d_array.sample(sampler_reg, _e1, 0u, metal::level(2.3), _e6);
metal::float4 _e70 = a;
a = _e70 + _e69;
metal::float4 _e76 = image_2d_array.sample(sampler_reg, _e1, 0u, metal::bias(2.0), _e6);
metal::float4 _e77 = a;
a = _e77 + _e76;
metal::float4 _e82 = image_2d_array.sample(sampler_reg, _e1, 0);
metal::float4 _e83 = a;
a = _e83 + _e82;
metal::float4 _e88 = image_2d_array.sample(sampler_reg, _e1, 0, _e6);
metal::float4 _e89 = a;
a = _e89 + _e88;
metal::float4 _e94 = image_2d_array.sample(sampler_reg, _e1, 0, metal::level(2.3));
metal::float4 _e95 = a;
a = _e95 + _e94;
metal::float4 _e100 = image_2d_array.sample(sampler_reg, _e1, 0, metal::level(2.3), _e6);
metal::float4 _e101 = a;
a = _e101 + _e100;
metal::float4 _e107 = image_2d_array.sample(sampler_reg, _e1, 0, metal::bias(2.0), _e6);
metal::float4 _e108 = a;
a = _e108 + _e107;
metal::float4 _e113 = image_cube_array.sample(sampler_reg, _e3, 0u);
metal::float4 _e114 = a;
a = _e114 + _e113;
metal::float4 _e119 = image_cube_array.sample(sampler_reg, _e3, 0u, metal::level(2.3));
metal::float4 _e120 = a;
a = _e120 + _e119;
metal::float4 _e126 = image_cube_array.sample(sampler_reg, _e3, 0u, metal::bias(2.0));
metal::float4 _e127 = a;
a = _e127 + _e126;
metal::float4 _e132 = image_cube_array.sample(sampler_reg, _e3, 0);
metal::float4 _e133 = a;
a = _e133 + _e132;
metal::float4 _e138 = image_cube_array.sample(sampler_reg, _e3, 0, metal::level(2.3));
metal::float4 _e139 = a;
a = _e139 + _e138;
metal::float4 _e145 = image_cube_array.sample(sampler_reg, _e3, 0, metal::bias(2.0));
metal::float4 _e146 = a;
a = _e146 + _e145;
metal::float4 _e148 = a;
return texture_sampleOutput { _e148 };
}

553
naga/tests/out/spv/wgsl-image.spvasm
View File

@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 556
; Bound: 567
OpCapability Shader
OpCapability Image1D
OpCapability Sampled1D
@ -14,15 +14,15 @@ OpEntryPoint GLCompute %197 "depth_load" %195
OpEntryPoint Vertex %220 "queries" %218
OpEntryPoint Vertex %272 "levels_queries" %271
OpEntryPoint Fragment %303 "texture_sample" %302
OpEntryPoint Fragment %450 "texture_sample_comparison" %448
OpEntryPoint Fragment %505 "gather" %504
OpEntryPoint Fragment %539 "depth_no_comparison" %538
OpEntryPoint Fragment %461 "texture_sample_comparison" %459
OpEntryPoint Fragment %516 "gather" %515
OpEntryPoint Fragment %550 "depth_no_comparison" %549
OpExecutionMode %96 LocalSize 16 1 1
OpExecutionMode %197 LocalSize 16 1 1
OpExecutionMode %303 OriginUpperLeft
OpExecutionMode %450 OriginUpperLeft
OpExecutionMode %505 OriginUpperLeft
OpExecutionMode %539 OriginUpperLeft
OpExecutionMode %461 OriginUpperLeft
OpExecutionMode %516 OriginUpperLeft
OpExecutionMode %550 OriginUpperLeft
%3 = OpString "image.wgsl"
OpSource Unknown 0 %3 "@group(0) @binding(0)
var image_mipmapped_src: texture_2d<u32>;
@ -153,6 +153,7 @@ fn texture_sample() -> @location(0) vec4<f32> {
a += textureSampleLevel(image_2d, sampler_reg, tc, level);
a += textureSampleLevel(image_2d, sampler_reg, tc, level, offset);
a += textureSampleBias(image_2d, sampler_reg, tc, 2.0, offset);
a += textureSampleBaseClampToEdge(image_2d, sampler_reg, tc);
a += textureSample(image_2d_array, sampler_reg, tc, 0u);
a += textureSample(image_2d_array, sampler_reg, tc, 0u, offset);
a += textureSampleLevel(image_2d_array, sampler_reg, tc, 0u, level);
@ -257,10 +258,10 @@ OpName %220 "queries"
OpName %272 "levels_queries"
OpName %303 "texture_sample"
OpName %318 "a"
OpName %450 "texture_sample_comparison"
OpName %455 "a"
OpName %505 "gather"
OpName %539 "depth_no_comparison"
OpName %461 "texture_sample_comparison"
OpName %466 "a"
OpName %516 "gather"
OpName %550 "depth_no_comparison"
OpDecorate %29 DescriptorSet 0
OpDecorate %29 Binding 0
OpDecorate %31 DescriptorSet 0
@ -313,9 +314,9 @@ OpDecorate %195 BuiltIn LocalInvocationId
OpDecorate %218 BuiltIn Position
OpDecorate %271 BuiltIn Position
OpDecorate %302 Location 0
OpDecorate %448 Location 0
OpDecorate %504 Location 0
OpDecorate %538 Location 0
OpDecorate %459 Location 0
OpDecorate %515 Location 0
OpDecorate %549 Location 0
%2 = OpTypeVoid
%5 = OpTypeInt 32 0
%4 = OpTypeImage %5 2D 0 0 0 1 Unknown
@ -424,22 +425,24 @@ OpDecorate %538 Location 0
%320 = OpConstantNull %24
%322 = OpTypeSampledImage %16
%327 = OpTypeSampledImage %17
%348 = OpTypeSampledImage %19
%409 = OpTypeSampledImage %21
%449 = OpTypePointer Output %8
%448 = OpVariable %449 Output
%456 = OpTypePointer Function %8
%457 = OpConstantNull %8
%459 = OpTypeSampledImage %26
%464 = OpTypeSampledImage %27
%477 = OpTypeSampledImage %28
%504 = OpVariable %219 Output
%515 = OpConstant %5 1
%518 = OpConstant %5 3
%523 = OpTypeSampledImage %4
%526 = OpTypeVector %15 4
%527 = OpTypeSampledImage %18
%538 = OpVariable %219 Output
%351 = OpConstant %8 1.0
%352 = OpConstantComposite %310 %351 %351
%359 = OpTypeSampledImage %19
%420 = OpTypeSampledImage %21
%460 = OpTypePointer Output %8
%459 = OpVariable %460 Output
%467 = OpTypePointer Function %8
%468 = OpConstantNull %8
%470 = OpTypeSampledImage %26
%475 = OpTypeSampledImage %27
%488 = OpTypeSampledImage %28
%515 = OpVariable %219 Output
%526 = OpConstant %5 1
%529 = OpConstant %5 3
%534 = OpTypeSampledImage %4
%537 = OpTypeVector %15 4
%538 = OpTypeSampledImage %18
%549 = OpVariable %219 Output
%70 = OpFunction %14 None %71
%72 = OpFunctionParameter %14
%73 = OpFunctionParameter %14
@ -778,299 +781,311 @@ OpLine %3 129 5
OpLine %3 129 5
OpStore %318 %347
OpLine %3 130 5
%349 = OpConvertUToF %8 %229
%350 = OpCompositeConstruct %312 %311 %349
%351 = OpSampledImage %348 %306 %308
%352 = OpImageSampleImplicitLod %24 %351 %350
%353 = OpLoad %24 %318
%354 = OpFAdd %24 %353 %352
%348 = OpImageQuerySizeLod %109 %305 %229
%349 = OpConvertUToF %310 %348
%350 = OpFDiv %310 %311 %349
%353 = OpFSub %310 %352 %350
%354 = OpExtInst %310 %1 NClamp %311 %350 %353
%355 = OpSampledImage %327 %305 %308
%356 = OpImageSampleExplicitLod %24 %355 %354 Lod %212
%357 = OpLoad %24 %318
%358 = OpFAdd %24 %357 %356
OpLine %3 130 5
OpStore %318 %354
OpStore %318 %358
OpLine %3 131 5
%355 = OpConvertUToF %8 %229
%356 = OpCompositeConstruct %312 %311 %355
%357 = OpSampledImage %348 %306 %308
%358 = OpImageSampleImplicitLod %24 %357 %356 ConstOffset %315
%359 = OpLoad %24 %318
%360 = OpFAdd %24 %359 %358
%360 = OpConvertUToF %8 %229
%361 = OpCompositeConstruct %312 %311 %360
%362 = OpSampledImage %359 %306 %308
%363 = OpImageSampleImplicitLod %24 %362 %361
%364 = OpLoad %24 %318
%365 = OpFAdd %24 %364 %363
OpLine %3 131 5
OpStore %318 %360
OpStore %318 %365
OpLine %3 132 5
%361 = OpConvertUToF %8 %229
%362 = OpCompositeConstruct %312 %311 %361
%363 = OpSampledImage %348 %306 %308
%364 = OpImageSampleExplicitLod %24 %363 %362 Lod %316
%365 = OpLoad %24 %318
%366 = OpFAdd %24 %365 %364
%366 = OpConvertUToF %8 %229
%367 = OpCompositeConstruct %312 %311 %366
%368 = OpSampledImage %359 %306 %308
%369 = OpImageSampleImplicitLod %24 %368 %367 ConstOffset %315
%370 = OpLoad %24 %318
%371 = OpFAdd %24 %370 %369
OpLine %3 132 5
OpStore %318 %366
OpStore %318 %371
OpLine %3 133 5
%367 = OpConvertUToF %8 %229
%368 = OpCompositeConstruct %312 %311 %367
%369 = OpSampledImage %348 %306 %308
%370 = OpImageSampleExplicitLod %24 %369 %368 Lod|ConstOffset %316 %315
%371 = OpLoad %24 %318
%372 = OpFAdd %24 %371 %370
%372 = OpConvertUToF %8 %229
%373 = OpCompositeConstruct %312 %311 %372
%374 = OpSampledImage %359 %306 %308
%375 = OpImageSampleExplicitLod %24 %374 %373 Lod %316
%376 = OpLoad %24 %318
%377 = OpFAdd %24 %376 %375
OpLine %3 133 5
OpStore %318 %372
OpStore %318 %377
OpLine %3 134 5
%373 = OpConvertUToF %8 %229
%374 = OpCompositeConstruct %312 %311 %373
%375 = OpSampledImage %348 %306 %308
%376 = OpImageSampleImplicitLod %24 %375 %374 Bias|ConstOffset %317 %315
%377 = OpLoad %24 %318
%378 = OpFAdd %24 %377 %376
%378 = OpConvertUToF %8 %229
%379 = OpCompositeConstruct %312 %311 %378
%380 = OpSampledImage %359 %306 %308
%381 = OpImageSampleExplicitLod %24 %380 %379 Lod|ConstOffset %316 %315
%382 = OpLoad %24 %318
%383 = OpFAdd %24 %382 %381
OpLine %3 134 5
OpStore %318 %378
OpStore %318 %383
OpLine %3 135 5
%379 = OpConvertSToF %8 %77
%380 = OpCompositeConstruct %312 %311 %379
%381 = OpSampledImage %348 %306 %308
%382 = OpImageSampleImplicitLod %24 %381 %380
%383 = OpLoad %24 %318
%384 = OpFAdd %24 %383 %382
%384 = OpConvertUToF %8 %229
%385 = OpCompositeConstruct %312 %311 %384
%386 = OpSampledImage %359 %306 %308
%387 = OpImageSampleImplicitLod %24 %386 %385 Bias|ConstOffset %317 %315
%388 = OpLoad %24 %318
%389 = OpFAdd %24 %388 %387
OpLine %3 135 5
OpStore %318 %384
OpStore %318 %389
OpLine %3 136 5
%385 = OpConvertSToF %8 %77
%386 = OpCompositeConstruct %312 %311 %385
%387 = OpSampledImage %348 %306 %308
%388 = OpImageSampleImplicitLod %24 %387 %386 ConstOffset %315
%389 = OpLoad %24 %318
%390 = OpFAdd %24 %389 %388
%390 = OpConvertSToF %8 %77
%391 = OpCompositeConstruct %312 %311 %390
%392 = OpSampledImage %359 %306 %308
%393 = OpImageSampleImplicitLod %24 %392 %391
%394 = OpLoad %24 %318
%395 = OpFAdd %24 %394 %393
OpLine %3 136 5
OpStore %318 %390
OpStore %318 %395
OpLine %3 137 5
%391 = OpConvertSToF %8 %77
%392 = OpCompositeConstruct %312 %311 %391
%393 = OpSampledImage %348 %306 %308
%394 = OpImageSampleExplicitLod %24 %393 %392 Lod %316
%395 = OpLoad %24 %318
%396 = OpFAdd %24 %395 %394
%396 = OpConvertSToF %8 %77
%397 = OpCompositeConstruct %312 %311 %396
%398 = OpSampledImage %359 %306 %308
%399 = OpImageSampleImplicitLod %24 %398 %397 ConstOffset %315
%400 = OpLoad %24 %318
%401 = OpFAdd %24 %400 %399
OpLine %3 137 5
OpStore %318 %396
OpStore %318 %401
OpLine %3 138 5
%397 = OpConvertSToF %8 %77
%398 = OpCompositeConstruct %312 %311 %397
%399 = OpSampledImage %348 %306 %308
%400 = OpImageSampleExplicitLod %24 %399 %398 Lod|ConstOffset %316 %315
%401 = OpLoad %24 %318
%402 = OpFAdd %24 %401 %400
%402 = OpConvertSToF %8 %77
%403 = OpCompositeConstruct %312 %311 %402
%404 = OpSampledImage %359 %306 %308
%405 = OpImageSampleExplicitLod %24 %404 %403 Lod %316
%406 = OpLoad %24 %318
%407 = OpFAdd %24 %406 %405
OpLine %3 138 5
OpStore %318 %402
OpStore %318 %407
OpLine %3 139 5
%403 = OpConvertSToF %8 %77
%404 = OpCompositeConstruct %312 %311 %403
%405 = OpSampledImage %348 %306 %308
%406 = OpImageSampleImplicitLod %24 %405 %404 Bias|ConstOffset %317 %315
%407 = OpLoad %24 %318
%408 = OpFAdd %24 %407 %406
%408 = OpConvertSToF %8 %77
%409 = OpCompositeConstruct %312 %311 %408
%410 = OpSampledImage %359 %306 %308
%411 = OpImageSampleExplicitLod %24 %410 %409 Lod|ConstOffset %316 %315
%412 = OpLoad %24 %318
%413 = OpFAdd %24 %412 %411
OpLine %3 139 5
OpStore %318 %408
OpStore %318 %413
OpLine %3 140 5
%410 = OpConvertUToF %8 %229
%411 = OpCompositeConstruct %24 %313 %410
%412 = OpSampledImage %409 %307 %308
%413 = OpImageSampleImplicitLod %24 %412 %411
%414 = OpLoad %24 %318
%415 = OpFAdd %24 %414 %413
%414 = OpConvertSToF %8 %77
%415 = OpCompositeConstruct %312 %311 %414
%416 = OpSampledImage %359 %306 %308
%417 = OpImageSampleImplicitLod %24 %416 %415 Bias|ConstOffset %317 %315
%418 = OpLoad %24 %318
%419 = OpFAdd %24 %418 %417
OpLine %3 140 5
OpStore %318 %415
OpStore %318 %419
OpLine %3 141 5
%416 = OpConvertUToF %8 %229
%417 = OpCompositeConstruct %24 %313 %416
%418 = OpSampledImage %409 %307 %308
%419 = OpImageSampleExplicitLod %24 %418 %417 Lod %316
%420 = OpLoad %24 %318
%421 = OpFAdd %24 %420 %419
%421 = OpConvertUToF %8 %229
%422 = OpCompositeConstruct %24 %313 %421
%423 = OpSampledImage %420 %307 %308
%424 = OpImageSampleImplicitLod %24 %423 %422
%425 = OpLoad %24 %318
%426 = OpFAdd %24 %425 %424
OpLine %3 141 5
OpStore %318 %421
OpStore %318 %426
OpLine %3 142 5
%422 = OpConvertUToF %8 %229
%423 = OpCompositeConstruct %24 %313 %422
%424 = OpSampledImage %409 %307 %308
%425 = OpImageSampleImplicitLod %24 %424 %423 Bias %317
%426 = OpLoad %24 %318
%427 = OpFAdd %24 %426 %425
%427 = OpConvertUToF %8 %229
%428 = OpCompositeConstruct %24 %313 %427
%429 = OpSampledImage %420 %307 %308
%430 = OpImageSampleExplicitLod %24 %429 %428 Lod %316
%431 = OpLoad %24 %318
%432 = OpFAdd %24 %431 %430
OpLine %3 142 5
OpStore %318 %427
OpStore %318 %432
OpLine %3 143 5
%428 = OpConvertSToF %8 %77
%429 = OpCompositeConstruct %24 %313 %428
%430 = OpSampledImage %409 %307 %308
%431 = OpImageSampleImplicitLod %24 %430 %429
%432 = OpLoad %24 %318
%433 = OpFAdd %24 %432 %431
%433 = OpConvertUToF %8 %229
%434 = OpCompositeConstruct %24 %313 %433
%435 = OpSampledImage %420 %307 %308
%436 = OpImageSampleImplicitLod %24 %435 %434 Bias %317
%437 = OpLoad %24 %318
%438 = OpFAdd %24 %437 %436
OpLine %3 143 5
OpStore %318 %433
OpStore %318 %438
OpLine %3 144 5
%434 = OpConvertSToF %8 %77
%435 = OpCompositeConstruct %24 %313 %434
%436 = OpSampledImage %409 %307 %308
%437 = OpImageSampleExplicitLod %24 %436 %435 Lod %316
%438 = OpLoad %24 %318
%439 = OpFAdd %24 %438 %437
%439 = OpConvertSToF %8 %77
%440 = OpCompositeConstruct %24 %313 %439
%441 = OpSampledImage %420 %307 %308
%442 = OpImageSampleImplicitLod %24 %441 %440
%443 = OpLoad %24 %318
%444 = OpFAdd %24 %443 %442
OpLine %3 144 5
OpStore %318 %439
OpStore %318 %444
OpLine %3 145 5
%440 = OpConvertSToF %8 %77
%441 = OpCompositeConstruct %24 %313 %440
%442 = OpSampledImage %409 %307 %308
%443 = OpImageSampleImplicitLod %24 %442 %441 Bias %317
%444 = OpLoad %24 %318
%445 = OpFAdd %24 %444 %443
%445 = OpConvertSToF %8 %77
%446 = OpCompositeConstruct %24 %313 %445
%447 = OpSampledImage %420 %307 %308
%448 = OpImageSampleExplicitLod %24 %447 %446 Lod %316
%449 = OpLoad %24 %318
%450 = OpFAdd %24 %449 %448
OpLine %3 145 5
OpStore %318 %445
OpStore %318 %450
OpLine %3 146 5
%451 = OpConvertSToF %8 %77
%452 = OpCompositeConstruct %24 %313 %451
%453 = OpSampledImage %420 %307 %308
%454 = OpImageSampleImplicitLod %24 %453 %452 Bias %317
%455 = OpLoad %24 %318
%456 = OpFAdd %24 %455 %454
OpLine %3 146 5
OpStore %318 %456
OpLine %3 1 1
%446 = OpLoad %24 %318
OpStore %302 %446
%457 = OpLoad %24 %318
OpStore %302 %457
OpReturn
OpFunctionEnd
%450 = OpFunction %2 None %97
%447 = OpLabel
%455 = OpVariable %456 Function %457
%451 = OpLoad %25 %63
%452 = OpLoad %26 %64
%453 = OpLoad %27 %66
%454 = OpLoad %28 %68
OpBranch %458
%461 = OpFunction %2 None %97
%458 = OpLabel
OpLine %3 160 14
OpLine %3 161 15
OpLine %3 164 5
%460 = OpSampledImage %459 %452 %451
%461 = OpImageSampleDrefImplicitLod %8 %460 %311 %309
%462 = OpLoad %8 %455
%463 = OpFAdd %8 %462 %461
OpLine %3 164 5
OpStore %455 %463
%466 = OpVariable %467 Function %468
%462 = OpLoad %25 %63
%463 = OpLoad %26 %64
%464 = OpLoad %27 %66
%465 = OpLoad %28 %68
OpBranch %469
%469 = OpLabel
OpLine %3 161 14
OpLine %3 162 15
OpLine %3 165 5
%465 = OpConvertUToF %8 %229
%466 = OpCompositeConstruct %312 %311 %465
%467 = OpSampledImage %464 %453 %451
%468 = OpImageSampleDrefImplicitLod %8 %467 %466 %309
%469 = OpLoad %8 %455
%470 = OpFAdd %8 %469 %468
%471 = OpSampledImage %470 %463 %462
%472 = OpImageSampleDrefImplicitLod %8 %471 %311 %309
%473 = OpLoad %8 %466
%474 = OpFAdd %8 %473 %472
OpLine %3 165 5
OpStore %455 %470
OpStore %466 %474
OpLine %3 166 5
%471 = OpConvertSToF %8 %77
%472 = OpCompositeConstruct %312 %311 %471
%473 = OpSampledImage %464 %453 %451
%474 = OpImageSampleDrefImplicitLod %8 %473 %472 %309
%475 = OpLoad %8 %455
%476 = OpFAdd %8 %475 %474
OpLine %3 166 5
OpStore %455 %476
OpLine %3 167 5
%478 = OpSampledImage %477 %454 %451
%479 = OpImageSampleDrefImplicitLod %8 %478 %313 %309
%480 = OpLoad %8 %455
%476 = OpConvertUToF %8 %229
%477 = OpCompositeConstruct %312 %311 %476
%478 = OpSampledImage %475 %464 %462
%479 = OpImageSampleDrefImplicitLod %8 %478 %477 %309
%480 = OpLoad %8 %466
%481 = OpFAdd %8 %480 %479
OpLine %3 166 5
OpStore %466 %481
OpLine %3 167 5
OpStore %455 %481
%482 = OpConvertSToF %8 %77
%483 = OpCompositeConstruct %312 %311 %482
%484 = OpSampledImage %475 %464 %462
%485 = OpImageSampleDrefImplicitLod %8 %484 %483 %309
%486 = OpLoad %8 %466
%487 = OpFAdd %8 %486 %485
OpLine %3 167 5
OpStore %466 %487
OpLine %3 168 5
%482 = OpSampledImage %459 %452 %451
%483 = OpImageSampleDrefExplicitLod %8 %482 %311 %309 Lod %212
%484 = OpLoad %8 %455
%485 = OpFAdd %8 %484 %483
%489 = OpSampledImage %488 %465 %462
%490 = OpImageSampleDrefImplicitLod %8 %489 %313 %309
%491 = OpLoad %8 %466
%492 = OpFAdd %8 %491 %490
OpLine %3 168 5
OpStore %455 %485
OpStore %466 %492
OpLine %3 169 5
%486 = OpConvertUToF %8 %229
%487 = OpCompositeConstruct %312 %311 %486
%488 = OpSampledImage %464 %453 %451
%489 = OpImageSampleDrefExplicitLod %8 %488 %487 %309 Lod %212
%490 = OpLoad %8 %455
%491 = OpFAdd %8 %490 %489
%493 = OpSampledImage %470 %463 %462
%494 = OpImageSampleDrefExplicitLod %8 %493 %311 %309 Lod %212
%495 = OpLoad %8 %466
%496 = OpFAdd %8 %495 %494
OpLine %3 169 5
OpStore %455 %491
OpStore %466 %496
OpLine %3 170 5
%492 = OpConvertSToF %8 %77
%493 = OpCompositeConstruct %312 %311 %492
%494 = OpSampledImage %464 %453 %451
%495 = OpImageSampleDrefExplicitLod %8 %494 %493 %309 Lod %212
%496 = OpLoad %8 %455
%497 = OpFAdd %8 %496 %495
%497 = OpConvertUToF %8 %229
%498 = OpCompositeConstruct %312 %311 %497
%499 = OpSampledImage %475 %464 %462
%500 = OpImageSampleDrefExplicitLod %8 %499 %498 %309 Lod %212
%501 = OpLoad %8 %466
%502 = OpFAdd %8 %501 %500
OpLine %3 170 5
OpStore %455 %497
OpStore %466 %502
OpLine %3 171 5
%498 = OpSampledImage %477 %454 %451
%499 = OpImageSampleDrefExplicitLod %8 %498 %313 %309 Lod %212
%500 = OpLoad %8 %455
%501 = OpFAdd %8 %500 %499
%503 = OpConvertSToF %8 %77
%504 = OpCompositeConstruct %312 %311 %503
%505 = OpSampledImage %475 %464 %462
%506 = OpImageSampleDrefExplicitLod %8 %505 %504 %309 Lod %212
%507 = OpLoad %8 %466
%508 = OpFAdd %8 %507 %506
OpLine %3 171 5
OpStore %455 %501
OpStore %466 %508
OpLine %3 172 5
%509 = OpSampledImage %488 %465 %462
%510 = OpImageSampleDrefExplicitLod %8 %509 %313 %309 Lod %212
%511 = OpLoad %8 %466
%512 = OpFAdd %8 %511 %510
OpLine %3 172 5
OpStore %466 %512
OpLine %3 1 1
%502 = OpLoad %8 %455
OpStore %448 %502
%513 = OpLoad %8 %466
OpStore %459 %513
OpReturn
OpFunctionEnd
%505 = OpFunction %2 None %97
%503 = OpLabel
%506 = OpLoad %17 %46
%507 = OpLoad %4 %48
%508 = OpLoad %18 %49
%509 = OpLoad %25 %61
%510 = OpLoad %25 %63
%511 = OpLoad %26 %64
OpBranch %512
%512 = OpLabel
OpLine %3 177 14
OpLine %3 179 15
%513 = OpSampledImage %327 %506 %509
%514 = OpImageGather %24 %513 %311 %515
OpLine %3 180 22
%516 = OpSampledImage %327 %506 %509
%517 = OpImageGather %24 %516 %311 %518 ConstOffset %315
OpLine %3 181 21
%519 = OpSampledImage %459 %511 %510
%520 = OpImageDrefGather %24 %519 %311 %309
OpLine %3 182 28
%521 = OpSampledImage %459 %511 %510
%522 = OpImageDrefGather %24 %521 %311 %309 ConstOffset %315
OpLine %3 184 13
%524 = OpSampledImage %523 %507 %509
%525 = OpImageGather %117 %524 %311 %229
%516 = OpFunction %2 None %97
%514 = OpLabel
%517 = OpLoad %17 %46
%518 = OpLoad %4 %48
%519 = OpLoad %18 %49
%520 = OpLoad %25 %61
%521 = OpLoad %25 %63
%522 = OpLoad %26 %64
OpBranch %523
%523 = OpLabel
OpLine %3 178 14
OpLine %3 180 15
%524 = OpSampledImage %327 %517 %520
%525 = OpImageGather %24 %524 %311 %526
OpLine %3 181 22
%527 = OpSampledImage %327 %517 %520
%528 = OpImageGather %24 %527 %311 %529 ConstOffset %315
OpLine %3 182 21
%530 = OpSampledImage %470 %522 %521
%531 = OpImageDrefGather %24 %530 %311 %309
OpLine %3 183 28
%532 = OpSampledImage %470 %522 %521
%533 = OpImageDrefGather %24 %532 %311 %309 ConstOffset %315
OpLine %3 185 13
%528 = OpSampledImage %527 %508 %509
%529 = OpImageGather %526 %528 %311 %229
%535 = OpSampledImage %534 %518 %520
%536 = OpImageGather %117 %535 %311 %229
OpLine %3 186 13
%530 = OpConvertUToF %24 %525
%531 = OpConvertSToF %24 %529
%532 = OpFAdd %24 %530 %531
OpLine %3 188 12
%533 = OpFAdd %24 %514 %517
%534 = OpFAdd %24 %533 %520
%535 = OpFAdd %24 %534 %522
%536 = OpFAdd %24 %535 %532
OpStore %504 %536
%539 = OpSampledImage %538 %519 %520
%540 = OpImageGather %537 %539 %311 %229
OpLine %3 187 13
%541 = OpConvertUToF %24 %536
%542 = OpConvertSToF %24 %540
%543 = OpFAdd %24 %541 %542
OpLine %3 189 12
%544 = OpFAdd %24 %525 %528
%545 = OpFAdd %24 %544 %531
%546 = OpFAdd %24 %545 %533
%547 = OpFAdd %24 %546 %543
OpStore %515 %547
OpReturn
OpFunctionEnd
%539 = OpFunction %2 None %97
%537 = OpLabel
%540 = OpLoad %25 %61
%541 = OpLoad %26 %64
OpBranch %542
%542 = OpLabel
OpLine %3 193 14
OpLine %3 195 15
%543 = OpSampledImage %459 %541 %540
%544 = OpImageSampleImplicitLod %24 %543 %311
%545 = OpCompositeExtract %8 %544 0
OpLine %3 196 22
%546 = OpSampledImage %459 %541 %540
%547 = OpImageGather %24 %546 %311 %229
OpLine %3 197 21
%548 = OpSampledImage %459 %541 %540
%550 = OpConvertSToF %8 %88
%549 = OpImageSampleExplicitLod %24 %548 %311 Lod %550
%551 = OpCompositeExtract %8 %549 0
OpLine %3 195 15
%552 = OpCompositeConstruct %24 %545 %545 %545 %545
%553 = OpFAdd %24 %552 %547
%554 = OpCompositeConstruct %24 %551 %551 %551 %551
%555 = OpFAdd %24 %553 %554
OpStore %538 %555
%550 = OpFunction %2 None %97
%548 = OpLabel
%551 = OpLoad %25 %61
%552 = OpLoad %26 %64
OpBranch %553
%553 = OpLabel
OpLine %3 194 14
OpLine %3 196 15
%554 = OpSampledImage %470 %552 %551
%555 = OpImageSampleImplicitLod %24 %554 %311
%556 = OpCompositeExtract %8 %555 0
OpLine %3 197 22
%557 = OpSampledImage %470 %552 %551
%558 = OpImageGather %24 %557 %311 %229
OpLine %3 198 21
%559 = OpSampledImage %470 %552 %551
%561 = OpConvertSToF %8 %88
%560 = OpImageSampleExplicitLod %24 %559 %311 Lod %561
%562 = OpCompositeExtract %8 %560 0
OpLine %3 196 15
%563 = OpCompositeConstruct %24 %556 %556 %556 %556
%564 = OpFAdd %24 %563 %558
%565 = OpCompositeConstruct %24 %562 %562 %562 %562
%566 = OpFAdd %24 %564 %565
OpStore %549 %566
OpReturn
OpFunctionEnd

103
naga/tests/out/wgsl/wgsl-image.wgsl
View File

@ -136,56 +136,59 @@ fn texture_sample() -> @location(0) vec4<f32> {
let _e40 = textureSampleBias(image_2d, sampler_reg, _e1, 2f, vec2<i32>(3i, 1i));
let _e41 = a;
a = (_e41 + _e40);
let _e46 = textureSample(image_2d_array, sampler_reg, _e1, 0u);
let _e47 = a;
a = (_e47 + _e46);
let _e52 = textureSample(image_2d_array, sampler_reg, _e1, 0u, vec2<i32>(3i, 1i));
let _e53 = a;
a = (_e53 + _e52);
let _e58 = textureSampleLevel(image_2d_array, sampler_reg, _e1, 0u, 2.3f);
let _e59 = a;
a = (_e59 + _e58);
let _e64 = textureSampleLevel(image_2d_array, sampler_reg, _e1, 0u, 2.3f, vec2<i32>(3i, 1i));
let _e65 = a;
a = (_e65 + _e64);
let _e71 = textureSampleBias(image_2d_array, sampler_reg, _e1, 0u, 2f, vec2<i32>(3i, 1i));
let _e72 = a;
a = (_e72 + _e71);
let _e77 = textureSample(image_2d_array, sampler_reg, _e1, 0i);
let _e78 = a;
a = (_e78 + _e77);
let _e83 = textureSample(image_2d_array, sampler_reg, _e1, 0i, vec2<i32>(3i, 1i));
let _e84 = a;
a = (_e84 + _e83);
let _e89 = textureSampleLevel(image_2d_array, sampler_reg, _e1, 0i, 2.3f);
let _e90 = a;
a = (_e90 + _e89);
let _e95 = textureSampleLevel(image_2d_array, sampler_reg, _e1, 0i, 2.3f, vec2<i32>(3i, 1i));
let _e96 = a;
a = (_e96 + _e95);
let _e102 = textureSampleBias(image_2d_array, sampler_reg, _e1, 0i, 2f, vec2<i32>(3i, 1i));
let _e103 = a;
a = (_e103 + _e102);
let _e108 = textureSample(image_cube_array, sampler_reg, _e3, 0u);
let _e109 = a;
a = (_e109 + _e108);
let _e114 = textureSampleLevel(image_cube_array, sampler_reg, _e3, 0u, 2.3f);
let _e115 = a;
a = (_e115 + _e114);
let _e121 = textureSampleBias(image_cube_array, sampler_reg, _e3, 0u, 2f);
let _e122 = a;
a = (_e122 + _e121);
let _e127 = textureSample(image_cube_array, sampler_reg, _e3, 0i);
let _e128 = a;
a = (_e128 + _e127);
let _e133 = textureSampleLevel(image_cube_array, sampler_reg, _e3, 0i, 2.3f);
let _e134 = a;
a = (_e134 + _e133);
let _e140 = textureSampleBias(image_cube_array, sampler_reg, _e3, 0i, 2f);
let _e141 = a;
a = (_e141 + _e140);
let _e143 = a;
return _e143;
let _e45 = textureSampleBaseClampToEdge(image_2d, sampler_reg, _e1);
let _e46 = a;
a = (_e46 + _e45);
let _e51 = textureSample(image_2d_array, sampler_reg, _e1, 0u);
let _e52 = a;
a = (_e52 + _e51);
let _e57 = textureSample(image_2d_array, sampler_reg, _e1, 0u, vec2<i32>(3i, 1i));
let _e58 = a;
a = (_e58 + _e57);
let _e63 = textureSampleLevel(image_2d_array, sampler_reg, _e1, 0u, 2.3f);
let _e64 = a;
a = (_e64 + _e63);
let _e69 = textureSampleLevel(image_2d_array, sampler_reg, _e1, 0u, 2.3f, vec2<i32>(3i, 1i));
let _e70 = a;
a = (_e70 + _e69);
let _e76 = textureSampleBias(image_2d_array, sampler_reg, _e1, 0u, 2f, vec2<i32>(3i, 1i));
let _e77 = a;
a = (_e77 + _e76);
let _e82 = textureSample(image_2d_array, sampler_reg, _e1, 0i);
let _e83 = a;
a = (_e83 + _e82);
let _e88 = textureSample(image_2d_array, sampler_reg, _e1, 0i, vec2<i32>(3i, 1i));
let _e89 = a;
a = (_e89 + _e88);
let _e94 = textureSampleLevel(image_2d_array, sampler_reg, _e1, 0i, 2.3f);
let _e95 = a;
a = (_e95 + _e94);
let _e100 = textureSampleLevel(image_2d_array, sampler_reg, _e1, 0i, 2.3f, vec2<i32>(3i, 1i));
let _e101 = a;
a = (_e101 + _e100);
let _e107 = textureSampleBias(image_2d_array, sampler_reg, _e1, 0i, 2f, vec2<i32>(3i, 1i));
let _e108 = a;
a = (_e108 + _e107);
let _e113 = textureSample(image_cube_array, sampler_reg, _e3, 0u);
let _e114 = a;
a = (_e114 + _e113);
let _e119 = textureSampleLevel(image_cube_array, sampler_reg, _e3, 0u, 2.3f);
let _e120 = a;
a = (_e120 + _e119);
let _e126 = textureSampleBias(image_cube_array, sampler_reg, _e3, 0u, 2f);
let _e127 = a;
a = (_e127 + _e126);
let _e132 = textureSample(image_cube_array, sampler_reg, _e3, 0i);
let _e133 = a;
a = (_e133 + _e132);
let _e138 = textureSampleLevel(image_cube_array, sampler_reg, _e3, 0i, 2.3f);
let _e139 = a;
a = (_e139 + _e138);
let _e145 = textureSampleBias(image_cube_array, sampler_reg, _e3, 0i, 2f);
let _e146 = a;
a = (_e146 + _e145);
let _e148 = a;
return _e148;
}
@fragment