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

glsl-out: texture function fixes

Browse Source 
- samplerCubeShadow requires vec3 gradients
- Bias must be written after offset
- sampler1D hack on GLES requires vec2 gradients
- textureSize doesn't accept a lod argument for MS samplers
- sampler1DArray hack on GLES requires a y component on imageLoad
This commit is contained in:
João Capucho
2022-02-22 16:26:37 +00:00
committed by Dzmitry Malyshau
parent 419a6c69f9
commit 98e7add6a2
10 changed files with 435 additions and 352 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -2192,7 +2192,11 @@ impl<'a, W: Write> Writer<'a, W> {
// Zero needs level set to 0
crate::SampleLevel::Zero => {
if workaround_lod_array_shadow_as_grad {
write!(self.out, ", vec2(0,0), vec2(0,0)")?;
let vec_dim = match dim {
crate::ImageDimension::Cube => 3,
_ => 2,
};
write!(self.out, ", vec{}(0.0), vec{}(0.0)", vec_dim, vec_dim)?;
} else if gather.is_none() {
write!(self.out, ", 0.0")?;
}
@@ -2207,15 +2211,25 @@ impl<'a, W: Write> Writer<'a, W> {
self.write_expr(expr, ctx)?;
}
}
crate::SampleLevel::Bias(expr) => {
write!(self.out, ", ")?;
self.write_expr(expr, ctx)?;
crate::SampleLevel::Bias(_) => {
// This needs to be done after the offset writing
}
crate::SampleLevel::Gradient { x, y } => {
write!(self.out, ", ")?;
self.write_expr(x, ctx)?;
write!(self.out, ", ")?;
self.write_expr(y, ctx)?;
// If we are using sampler2D to replace sampler1D, we also
// need to make sure to use vec2 gradients
if tex_1d_hack {
write!(self.out, ", vec2(")?;
self.write_expr(x, ctx)?;
write!(self.out, ", 0.0)")?;
write!(self.out, ", vec2(")?;
self.write_expr(y, ctx)?;
write!(self.out, ", 0.0)")?;
} else {
write!(self.out, ", ")?;
self.write_expr(x, ctx)?;
write!(self.out, ", ")?;
self.write_expr(y, ctx)?;
}
}
}
@@ -2230,6 +2244,12 @@ impl<'a, W: Write> Writer<'a, W> {
}
}
// Bias is always the last argument
if let crate::SampleLevel::Bias(expr) = level {
write!(self.out, ", ")?;
self.write_expr(expr, ctx)?;
}
if let (Some(component), None) = (gather, depth_ref) {
write!(self.out, ", {}", component as usize)?;
}
@@ -2307,14 +2327,16 @@ impl<'a, W: Write> Writer<'a, W> {
match query {
crate::ImageQuery::Size { level } => {
match class {
ImageClass::Sampled { .. } | ImageClass::Depth { .. } => {
ImageClass::Sampled { multi, .. } | ImageClass::Depth { multi } => {
write!(self.out, "textureSize(")?;
self.write_expr(image, ctx)?;
write!(self.out, ", ")?;
if let Some(expr) = level {
write!(self.out, ", ")?;
self.write_expr(expr, ctx)?;
} else {
write!(self.out, "0")?;
} else if !multi {
// All textureSize calls requires an lod argument
// except for multisampled samplers
write!(self.out, ", 0")?;
}
}
ImageClass::Storage { .. } => {
@@ -2339,8 +2361,14 @@ impl<'a, W: Write> Writer<'a, W> {
};
write!(self.out, "{}(", fun_name)?;
self.write_expr(image, ctx)?;
// All textureSize calls requires an lod argument
// except for multisampled samplers
if class.is_multisampled() {
write!(self.out, ", 0")?;
}
write!(self.out, ")")?;
if components != 1 || self.options.version.is_es() {
write!(self.out, ", 0).{}", back::COMPONENTS[components])?;
write!(self.out, ".{}", back::COMPONENTS[components])?;
}
}
crate::ImageQuery::NumSamples => {
@@ -2808,22 +2836,27 @@ impl<'a, W: Write> Writer<'a, W> {
) -> Result<(), Error> {
use crate::ImageDimension as IDim;
let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es();
match array_index {
Some(layer_expr) => {
let tex_coord_type = match dim {
IDim::D1 => "ivec2",
IDim::D2 => "ivec3",
IDim::D3 => "ivec4",
IDim::Cube => "ivec4",
let tex_coord_size = match dim {
IDim::D1 => 2,
IDim::D2 => 3,
IDim::D3 => 4,
IDim::Cube => 4,
};
write!(self.out, "{}(", tex_coord_type)?;
write!(self.out, "ivec{}(", tex_coord_size + tex_1d_hack as u8)?;
self.write_expr(coordinate, ctx)?;
write!(self.out, ", ")?;
// If we are replacing sampler1D with sampler2D we also need
// to add another zero to the coordinates vector for the y component
if tex_1d_hack {
write!(self.out, "0, ")?;
}
self.write_expr(layer_expr, ctx)?;
write!(self.out, ")")?;
}
None => {
let tex_1d_hack = dim == IDim::D1 && self.options.version.is_es();
if tex_1d_hack {
write!(self.out, "ivec2(")?;
}

5
tests/in/image.wgsl
View File

@@ -69,6 +69,7 @@ fn queries() -> @builtin(position) vec4<f32> {
let dim_cube_array_lod = textureDimensions(image_cube_array, 1);
let dim_3d = textureDimensions(image_3d);
let dim_3d_lod = textureDimensions(image_3d, 1);
let dim_2s_ms = textureDimensions(image_aa);
let sum = dim_1d + dim_2d.y + dim_2d_lod.y + dim_2d_array.y + dim_2d_array_lod.y +
dim_cube.y + dim_cube_lod.y + dim_cube_array.y + dim_cube_array_lod.y +
@@ -104,6 +105,7 @@ fn sample() -> @location(0) vec4<f32> {
let s2d_offset = textureSample(image_2d, sampler_reg, tc, vec2<i32>(3, 1));
let s2d_level = textureSampleLevel(image_2d, sampler_reg, tc, level);
let s2d_level_offset = textureSampleLevel(image_2d, sampler_reg, tc, level, vec2<i32>(3, 1));
let s2d_bias_offset = textureSampleBias(image_2d, sampler_reg, tc, 2.0, vec2<i32>(3, 1));
return s1d + s2d + s2d_offset + s2d_level + s2d_level_offset;
}
@@ -111,6 +113,8 @@ fn sample() -> @location(0) vec4<f32> {
var sampler_cmp: sampler_comparison;
@group(1) @binding(2)
var image_2d_depth: texture_depth_2d;
@group(1) @binding(3)
var image_cube_depth: texture_depth_cube;
@stage(fragment)
fn sample_comparison() -> @location(0) f32 {
@@ -118,6 +122,7 @@ fn sample_comparison() -> @location(0) f32 {
let dref = 0.5;
let s2d_depth = textureSampleCompare(image_2d_depth, sampler_cmp, tc, dref);
let s2d_depth_level = textureSampleCompareLevel(image_2d_depth, sampler_cmp, tc, dref);
let scube_depth_level = textureSampleCompareLevel(image_cube_depth, sampler_cmp, vec3<f32>(0.5), dref);
return s2d_depth + s2d_depth_level;
}

3
tests/out/glsl/image.queries.Vertex.glsl
View File

@@ -16,6 +16,8 @@ uniform highp samplerCubeArray _group_0_binding_4_vs;
uniform highp sampler3D _group_0_binding_5_vs;
uniform highp sampler2DMS _group_0_binding_6_vs;
void main() {
int dim_1d = textureSize(_group_0_binding_0_vs, 0).x;
@@ -30,6 +32,7 @@ void main() {
ivec2 dim_cube_array_lod = textureSize(_group_0_binding_4_vs, 1).xy;
ivec3 dim_3d = textureSize(_group_0_binding_5_vs, 0).xyz;
ivec3 dim_3d_lod = textureSize(_group_0_binding_5_vs, 1).xyz;
ivec2 dim_2s_ms = textureSize(_group_0_binding_6_vs).xy;
int sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
gl_Position = vec4(float(sum));
gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w);

1
tests/out/glsl/image.sample.Fragment.glsl
View File

@@ -17,6 +17,7 @@ void main() {
vec4 s2d_offset = textureOffset(_group_0_binding_1_fs, vec2(tc), ivec2(3, 1));
vec4 s2d_level = textureLod(_group_0_binding_1_fs, vec2(tc), 2.299999952316284);
vec4 s2d_level_offset = textureLodOffset(_group_0_binding_1_fs, vec2(tc), 2.299999952316284, ivec2(3, 1));
vec4 s2d_bias_offset = textureOffset(_group_0_binding_1_fs, vec2(tc), ivec2(3, 1), 2.0);
_fs2p_location0 = ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset);
return;
}

3
tests/out/glsl/image.sample_comparison.Fragment.glsl
View File

@@ -6,12 +6,15 @@ precision highp int;
uniform highp sampler2DShadow _group_1_binding_2_fs;
uniform highp samplerCubeShadow _group_1_binding_3_fs;
layout(location = 0) out float _fs2p_location0;
void main() {
vec2 tc = vec2(0.5);
float s2d_depth = texture(_group_1_binding_2_fs, vec3(tc, 0.5));
float s2d_depth_level = textureLod(_group_1_binding_2_fs, vec3(tc, 0.5), 0.0);
float scube_depth_level = textureGrad(_group_1_binding_3_fs, vec4(vec3(0.5), 0.5), vec3(0.0), vec3(0.0));
_fs2p_location0 = (s2d_depth + s2d_depth_level);
return;
}

2
tests/out/glsl/shadow.fs_main.Fragment.glsl
View File

@@ -29,7 +29,7 @@ float fetch_shadow(uint light_id, vec4 homogeneous_coords) {
}
vec2 flip_correction = vec2(0.5, -0.5);
vec2 light_local = (((homogeneous_coords.xy * flip_correction) / vec2(homogeneous_coords.w)) + vec2(0.5, 0.5));
float _e26 = textureGrad(_group_0_binding_2_fs, vec4(light_local, int(light_id), (homogeneous_coords.z / homogeneous_coords.w)), vec2(0,0), vec2(0,0));
float _e26 = textureGrad(_group_0_binding_2_fs, vec4(light_local, int(light_id), (homogeneous_coords.z / homogeneous_coords.w)), vec2(0.0), vec2(0.0));
return _e26;
}

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

@@ -17,6 +17,7 @@ Texture2DMS<float4> image_aa : register(t6);
SamplerState sampler_reg : register(s0, space1);
SamplerComparisonState sampler_cmp : register(s1, space1);
Texture2D<float> image_2d_depth : register(t2, space1);
TextureCube<float> image_cube_depth : register(t3, space1);
int2 NagaRWDimensions2D(RWTexture2D<uint4> tex)
{
@@ -133,6 +134,13 @@ int3 NagaMipDimensions3D(Texture3D<float4> tex, uint mip_level)
return ret.xyz;
}
int2 NagaMSDimensions2D(Texture2DMS<float4> tex)
{
uint4 ret;
tex.GetDimensions(ret.x, ret.y, ret.z);
return ret.xy;
}
float4 queries() : SV_Position
{
int dim_1d = NagaDimensions1D(image_1d);
@@ -147,6 +155,7 @@ float4 queries() : SV_Position
int2 dim_cube_array_lod = NagaMipDimensionsCubeArray(image_cube_array, 1);
int3 dim_3d = NagaDimensions3D(image_3d);
int3 dim_3d_lod = NagaMipDimensions3D(image_3d, 1);
int2 dim_2s_ms = NagaMSDimensions2D(image_aa);
int sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
return float4(float(sum).xxxx);
}
@@ -229,6 +238,7 @@ float4 sample_() : SV_Target0
float4 s2d_offset = image_2d.Sample(sampler_reg, tc, int2(3, 1));
float4 s2d_level = image_2d.SampleLevel(sampler_reg, tc, 2.299999952316284);
float4 s2d_level_offset = image_2d.SampleLevel(sampler_reg, tc, 2.299999952316284, int2(3, 1));
float4 s2d_bias_offset = image_2d.SampleBias(sampler_reg, tc, 2.0, int2(3, 1));
return ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset);
}
@@ -237,6 +247,7 @@ float sample_comparison() : SV_Target0
float2 tc_1 = float2(0.5.xx);
float s2d_depth = image_2d_depth.SampleCmp(sampler_cmp, tc_1, 0.5);
float s2d_depth_level = image_2d_depth.SampleCmpLevelZero(sampler_cmp, tc_1, 0.5);
float scube_depth_level = image_cube_depth.SampleCmpLevelZero(sampler_cmp, float3(0.5.xxx), 0.5);
return (s2d_depth + s2d_depth_level);
}

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

@@ -55,6 +55,7 @@ vertex queriesOutput queries(
, metal::texturecube<float, metal::access::sample> image_cube [[user(fake0)]]
, metal::texturecube_array<float, metal::access::sample> image_cube_array [[user(fake0)]]
, metal::texture3d<float, metal::access::sample> image_3d [[user(fake0)]]
, metal::texture2d_ms<float, metal::access::read> image_aa [[user(fake0)]]
) {
int dim_1d = int(image_1d.get_width());
int dim_1d_lod = int(image_1d.get_width());
@@ -68,6 +69,7 @@ vertex queriesOutput queries(
metal::int2 dim_cube_array_lod = metal::int2(image_cube_array.get_width(1));
metal::int3 dim_3d = metal::int3(image_3d.get_width(), image_3d.get_height(), image_3d.get_depth());
metal::int3 dim_3d_lod = metal::int3(image_3d.get_width(1), image_3d.get_height(1), image_3d.get_depth(1));
metal::int2 dim_2s_ms = metal::int2(image_aa.get_width(), image_aa.get_height());
int sum = (((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z;
return queriesOutput { metal::float4(static_cast<float>(sum)) };
}
@@ -111,6 +113,7 @@ fragment sampleOutput sample(
metal::float4 s2d_offset = image_2d.sample(sampler_reg, tc, const_type_9_);
metal::float4 s2d_level = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284));
metal::float4 s2d_level_offset = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284), const_type_9_);
metal::float4 s2d_bias_offset = image_2d.sample(sampler_reg, tc, metal::bias(2.0), const_type_9_);
return sampleOutput { (((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset };
}
@@ -121,10 +124,12 @@ struct sample_comparisonOutput {
fragment sample_comparisonOutput sample_comparison(
metal::sampler sampler_cmp [[user(fake0)]]
, metal::depth2d<float, metal::access::sample> image_2d_depth [[user(fake0)]]
, metal::depthcube<float, metal::access::sample> image_cube_depth [[user(fake0)]]
) {
metal::float2 tc_1 = metal::float2(0.5);
float s2d_depth = image_2d_depth.sample_compare(sampler_cmp, tc_1, 0.5);
float s2d_depth_level = image_2d_depth.sample_compare(sampler_cmp, tc_1, 0.5);
float scube_depth_level = image_cube_depth.sample_compare(sampler_cmp, metal::float3(0.5), 0.5);
return sample_comparisonOutput { s2d_depth + s2d_depth_level };
}

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

@@ -1,7 +1,7 @@
; SPIR-V
; Version: 1.1
; Generator: rspirv
; Bound: 292
; Bound: 306
OpCapability SampledCubeArray
OpCapability ImageQuery
OpCapability Image1D
@@ -9,96 +9,99 @@ OpCapability Shader
OpCapability Sampled1D
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %72 "main" %69
OpEntryPoint GLCompute %117 "depth_load" %115
OpEntryPoint Vertex %138 "queries" %136
OpEntryPoint Vertex %188 "levels_queries" %187
OpEntryPoint Fragment %217 "sample" %216
OpEntryPoint Fragment %244 "sample_comparison" %242
OpEntryPoint Fragment %258 "gather" %257
OpEntryPoint Fragment %280 "depth_no_comparison" %279
OpExecutionMode %72 LocalSize 16 1 1
OpExecutionMode %117 LocalSize 16 1 1
OpExecutionMode %217 OriginUpperLeft
OpExecutionMode %244 OriginUpperLeft
OpExecutionMode %258 OriginUpperLeft
OpExecutionMode %280 OriginUpperLeft
OpEntryPoint GLCompute %76 "main" %73
OpEntryPoint GLCompute %121 "depth_load" %119
OpEntryPoint Vertex %142 "queries" %140
OpEntryPoint Vertex %194 "levels_queries" %193
OpEntryPoint Fragment %223 "sample" %222
OpEntryPoint Fragment %252 "sample_comparison" %250
OpEntryPoint Fragment %272 "gather" %271
OpEntryPoint Fragment %294 "depth_no_comparison" %293
OpExecutionMode %76 LocalSize 16 1 1
OpExecutionMode %121 LocalSize 16 1 1
OpExecutionMode %223 OriginUpperLeft
OpExecutionMode %252 OriginUpperLeft
OpExecutionMode %272 OriginUpperLeft
OpExecutionMode %294 OriginUpperLeft
OpSource GLSL 450
OpName %32 "image_mipmapped_src"
OpName %34 "image_multisampled_src"
OpName %36 "image_depth_multisampled_src"
OpName %38 "image_storage_src"
OpName %40 "image_array_src"
OpName %42 "image_dup_src"
OpName %44 "image_1d_src"
OpName %46 "image_dst"
OpName %48 "image_1d"
OpName %50 "image_2d"
OpName %52 "image_2d_array"
OpName %54 "image_cube"
OpName %56 "image_cube_array"
OpName %58 "image_3d"
OpName %60 "image_aa"
OpName %62 "sampler_reg"
OpName %64 "sampler_cmp"
OpName %66 "image_2d_depth"
OpName %69 "local_id"
OpName %72 "main"
OpName %115 "local_id"
OpName %117 "depth_load"
OpName %138 "queries"
OpName %188 "levels_queries"
OpName %217 "sample"
OpName %244 "sample_comparison"
OpName %258 "gather"
OpName %280 "depth_no_comparison"
OpDecorate %32 DescriptorSet 0
OpDecorate %32 Binding 0
OpName %34 "image_mipmapped_src"
OpName %36 "image_multisampled_src"
OpName %38 "image_depth_multisampled_src"
OpName %40 "image_storage_src"
OpName %42 "image_array_src"
OpName %44 "image_dup_src"
OpName %46 "image_1d_src"
OpName %48 "image_dst"
OpName %50 "image_1d"
OpName %52 "image_2d"
OpName %54 "image_2d_array"
OpName %56 "image_cube"
OpName %58 "image_cube_array"
OpName %60 "image_3d"
OpName %62 "image_aa"
OpName %64 "sampler_reg"
OpName %66 "sampler_cmp"
OpName %68 "image_2d_depth"
OpName %70 "image_cube_depth"
OpName %73 "local_id"
OpName %76 "main"
OpName %119 "local_id"
OpName %121 "depth_load"
OpName %142 "queries"
OpName %194 "levels_queries"
OpName %223 "sample"
OpName %252 "sample_comparison"
OpName %272 "gather"
OpName %294 "depth_no_comparison"
OpDecorate %34 DescriptorSet 0
OpDecorate %34 Binding 3
OpDecorate %34 Binding 0
OpDecorate %36 DescriptorSet 0
OpDecorate %36 Binding 4
OpDecorate %38 NonWritable
OpDecorate %36 Binding 3
OpDecorate %38 DescriptorSet 0
OpDecorate %38 Binding 1
OpDecorate %38 Binding 4
OpDecorate %40 NonWritable
OpDecorate %40 DescriptorSet 0
OpDecorate %40 Binding 5
OpDecorate %42 NonWritable
OpDecorate %40 Binding 1
OpDecorate %42 DescriptorSet 0
OpDecorate %42 Binding 6
OpDecorate %42 Binding 5
OpDecorate %44 NonWritable
OpDecorate %44 DescriptorSet 0
OpDecorate %44 Binding 7
OpDecorate %46 NonReadable
OpDecorate %44 Binding 6
OpDecorate %46 DescriptorSet 0
OpDecorate %46 Binding 2
OpDecorate %46 Binding 7
OpDecorate %48 NonReadable
OpDecorate %48 DescriptorSet 0
OpDecorate %48 Binding 0
OpDecorate %48 Binding 2
OpDecorate %50 DescriptorSet 0
OpDecorate %50 Binding 1
OpDecorate %50 Binding 0
OpDecorate %52 DescriptorSet 0
OpDecorate %52 Binding 2
OpDecorate %52 Binding 1
OpDecorate %54 DescriptorSet 0
OpDecorate %54 Binding 3
OpDecorate %54 Binding 2
OpDecorate %56 DescriptorSet 0
OpDecorate %56 Binding 4
OpDecorate %56 Binding 3
OpDecorate %58 DescriptorSet 0
OpDecorate %58 Binding 5
OpDecorate %58 Binding 4
OpDecorate %60 DescriptorSet 0
OpDecorate %60 Binding 6
OpDecorate %62 DescriptorSet 1
OpDecorate %62 Binding 0
OpDecorate %60 Binding 5
OpDecorate %62 DescriptorSet 0
OpDecorate %62 Binding 6
OpDecorate %64 DescriptorSet 1
OpDecorate %64 Binding 1
OpDecorate %64 Binding 0
OpDecorate %66 DescriptorSet 1
OpDecorate %66 Binding 2
OpDecorate %69 BuiltIn LocalInvocationId
OpDecorate %115 BuiltIn LocalInvocationId
OpDecorate %136 BuiltIn Position
OpDecorate %187 BuiltIn Position
OpDecorate %216 Location 0
OpDecorate %242 Location 0
OpDecorate %257 Location 0
OpDecorate %279 Location 0
OpDecorate %66 Binding 1
OpDecorate %68 DescriptorSet 1
OpDecorate %68 Binding 2
OpDecorate %70 DescriptorSet 1
OpDecorate %70 Binding 3
OpDecorate %73 BuiltIn LocalInvocationId
OpDecorate %119 BuiltIn LocalInvocationId
OpDecorate %140 BuiltIn Position
OpDecorate %193 BuiltIn Position
OpDecorate %222 Location 0
OpDecorate %250 Location 0
OpDecorate %271 Location 0
OpDecorate %293 Location 0
%2 = OpTypeVoid
%4 = OpTypeInt 32 1
%3 = OpConstant %4 10
@@ -108,30 +111,30 @@ OpDecorate %279 Location 0
%7 = OpConstant %8 0.5
%9 = OpConstant %8 2.3
%10 = OpConstant %4 3
%12 = OpTypeInt 32 0
%11 = OpTypeImage %12 2D 0 0 0 1 Unknown
%13 = OpTypeImage %12 2D 0 0 1 1 Unknown
%14 = OpTypeImage %8 2D 1 0 1 1 Unknown
%15 = OpTypeImage %12 2D 0 0 0 2 Rgba8ui
%16 = OpTypeImage %12 2D 0 1 0 1 Unknown
%17 = OpTypeImage %12 1D 0 0 0 2 R32ui
%18 = OpTypeImage %12 1D 0 0 0 1 Unknown
%19 = OpTypeVector %12 3
%20 = OpTypeVector %4 2
%21 = OpTypeImage %8 1D 0 0 0 1 Unknown
%22 = OpTypeImage %8 2D 0 0 0 1 Unknown
%23 = OpTypeImage %8 2D 0 1 0 1 Unknown
%24 = OpTypeImage %8 Cube 0 0 0 1 Unknown
%25 = OpTypeImage %8 Cube 0 1 0 1 Unknown
%26 = OpTypeImage %8 3D 0 0 0 1 Unknown
%27 = OpTypeImage %8 2D 0 0 1 1 Unknown
%28 = OpTypeVector %8 4
%29 = OpTypeSampler
%30 = OpTypeImage %8 2D 1 0 0 1 Unknown
%31 = OpConstantComposite %20 %10 %6
%33 = OpTypePointer UniformConstant %11
%32 = OpVariable %33 UniformConstant
%35 = OpTypePointer UniformConstant %13
%11 = OpConstant %8 2.0
%13 = OpTypeInt 32 0
%12 = OpTypeImage %13 2D 0 0 0 1 Unknown
%14 = OpTypeImage %13 2D 0 0 1 1 Unknown
%15 = OpTypeImage %8 2D 1 0 1 1 Unknown
%16 = OpTypeImage %13 2D 0 0 0 2 Rgba8ui
%17 = OpTypeImage %13 2D 0 1 0 1 Unknown
%18 = OpTypeImage %13 1D 0 0 0 2 R32ui
%19 = OpTypeImage %13 1D 0 0 0 1 Unknown
%20 = OpTypeVector %13 3
%21 = OpTypeVector %4 2
%22 = OpTypeImage %8 1D 0 0 0 1 Unknown
%23 = OpTypeImage %8 2D 0 0 0 1 Unknown
%24 = OpTypeImage %8 2D 0 1 0 1 Unknown
%25 = OpTypeImage %8 Cube 0 0 0 1 Unknown
%26 = OpTypeImage %8 Cube 0 1 0 1 Unknown
%27 = OpTypeImage %8 3D 0 0 0 1 Unknown
%28 = OpTypeImage %8 2D 0 0 1 1 Unknown
%29 = OpTypeVector %8 4
%30 = OpTypeSampler
%31 = OpTypeImage %8 2D 1 0 0 1 Unknown
%32 = OpTypeImage %8 Cube 1 0 0 1 Unknown
%33 = OpConstantComposite %21 %10 %6
%35 = OpTypePointer UniformConstant %12
%34 = OpVariable %35 UniformConstant
%37 = OpTypePointer UniformConstant %14
%36 = OpVariable %37 UniformConstant
@@ -143,9 +146,9 @@ OpDecorate %279 Location 0
%42 = OpVariable %43 UniformConstant
%45 = OpTypePointer UniformConstant %18
%44 = OpVariable %45 UniformConstant
%47 = OpTypePointer UniformConstant %17
%47 = OpTypePointer UniformConstant %19
%46 = OpVariable %47 UniformConstant
%49 = OpTypePointer UniformConstant %21
%49 = OpTypePointer UniformConstant %18
%48 = OpVariable %49 UniformConstant
%51 = OpTypePointer UniformConstant %22
%50 = OpVariable %51 UniformConstant
@@ -159,265 +162,279 @@ OpDecorate %279 Location 0
%58 = OpVariable %59 UniformConstant
%61 = OpTypePointer UniformConstant %27
%60 = OpVariable %61 UniformConstant
%63 = OpTypePointer UniformConstant %29
%63 = OpTypePointer UniformConstant %28
%62 = OpVariable %63 UniformConstant
%65 = OpTypePointer UniformConstant %29
%65 = OpTypePointer UniformConstant %30
%64 = OpVariable %65 UniformConstant
%67 = OpTypePointer UniformConstant %30
%66 = OpVariable %67 UniformConstant
%70 = OpTypePointer Input %19
%69 = OpVariable %70 Input
%73 = OpTypeFunction %2
%82 = OpTypeVector %12 2
%90 = OpTypeVector %12 4
%101 = OpTypeVector %4 3
%115 = OpVariable %70 Input
%137 = OpTypePointer Output %28
%136 = OpVariable %137 Output
%146 = OpConstant %12 0
%187 = OpVariable %137 Output
%216 = OpVariable %137 Output
%222 = OpTypeVector %8 2
%225 = OpTypeSampledImage %21
%228 = OpTypeSampledImage %22
%243 = OpTypePointer Output %8
%242 = OpVariable %243 Output
%249 = OpTypeSampledImage %30
%254 = OpConstant %8 0.0
%257 = OpVariable %137 Output
%267 = OpConstant %12 1
%270 = OpConstant %12 3
%279 = OpVariable %137 Output
%72 = OpFunction %2 None %73
%68 = OpLabel
%71 = OpLoad %19 %69
%74 = OpLoad %11 %32
%75 = OpLoad %13 %34
%76 = OpLoad %15 %38
%77 = OpLoad %16 %40
%78 = OpLoad %18 %44
%79 = OpLoad %17 %46
OpBranch %80
%80 = OpLabel
%81 = OpImageQuerySize %20 %76
%83 = OpVectorShuffle %82 %71 %71 0 1
%84 = OpBitcast %20 %83
%85 = OpIMul %20 %81 %84
%86 = OpCompositeConstruct %20 %3 %5
%87 = OpSMod %20 %85 %86
%88 = OpCompositeExtract %12 %71 2
%89 = OpBitcast %4 %88
%91 = OpImageFetch %90 %74 %87 Lod %89
%92 = OpCompositeExtract %12 %71 2
%69 = OpTypePointer UniformConstant %31
%68 = OpVariable %69 UniformConstant
%71 = OpTypePointer UniformConstant %32
%70 = OpVariable %71 UniformConstant
%74 = OpTypePointer Input %20
%73 = OpVariable %74 Input
%77 = OpTypeFunction %2
%86 = OpTypeVector %13 2
%94 = OpTypeVector %13 4
%105 = OpTypeVector %4 3
%119 = OpVariable %74 Input
%141 = OpTypePointer Output %29
%140 = OpVariable %141 Output
%151 = OpConstant %13 0
%193 = OpVariable %141 Output
%222 = OpVariable %141 Output
%228 = OpTypeVector %8 2
%231 = OpTypeSampledImage %22
%234 = OpTypeSampledImage %23
%251 = OpTypePointer Output %8
%250 = OpVariable %251 Output
%258 = OpTypeSampledImage %31
%263 = OpConstant %8 0.0
%264 = OpTypeVector %8 3
%266 = OpTypeSampledImage %32
%271 = OpVariable %141 Output
%281 = OpConstant %13 1
%284 = OpConstant %13 3
%293 = OpVariable %141 Output
%76 = OpFunction %2 None %77
%72 = OpLabel
%75 = OpLoad %20 %73
%78 = OpLoad %12 %34
%79 = OpLoad %14 %36
%80 = OpLoad %16 %40
%81 = OpLoad %17 %42
%82 = OpLoad %19 %46
%83 = OpLoad %18 %48
OpBranch %84
%84 = OpLabel
%85 = OpImageQuerySize %21 %80
%87 = OpVectorShuffle %86 %75 %75 0 1
%88 = OpBitcast %21 %87
%89 = OpIMul %21 %85 %88
%90 = OpCompositeConstruct %21 %3 %5
%91 = OpSMod %21 %89 %90
%92 = OpCompositeExtract %13 %75 2
%93 = OpBitcast %4 %92
%94 = OpImageFetch %90 %75 %87 Sample %93
%95 = OpImageRead %90 %76 %87
%96 = OpCompositeExtract %12 %71 2
%95 = OpImageFetch %94 %78 %91 Lod %93
%96 = OpCompositeExtract %13 %75 2
%97 = OpBitcast %4 %96
%98 = OpCompositeExtract %12 %71 2
%99 = OpBitcast %4 %98
%100 = OpIAdd %4 %99 %6
%102 = OpCompositeConstruct %101 %87 %97
%103 = OpImageFetch %90 %77 %102 Lod %100
%104 = OpCompositeExtract %12 %71 0
%105 = OpBitcast %4 %104
%106 = OpCompositeExtract %12 %71 2
%107 = OpBitcast %4 %106
%108 = OpImageFetch %90 %78 %105 Lod %107
%109 = OpCompositeExtract %4 %87 0
%110 = OpIAdd %90 %91 %94
%111 = OpIAdd %90 %110 %95
%112 = OpIAdd %90 %111 %103
%113 = OpIAdd %90 %112 %108
OpImageWrite %79 %109 %113
%98 = OpImageFetch %94 %79 %91 Sample %97
%99 = OpImageRead %94 %80 %91
%100 = OpCompositeExtract %13 %75 2
%101 = OpBitcast %4 %100
%102 = OpCompositeExtract %13 %75 2
%103 = OpBitcast %4 %102
%104 = OpIAdd %4 %103 %6
%106 = OpCompositeConstruct %105 %91 %101
%107 = OpImageFetch %94 %81 %106 Lod %104
%108 = OpCompositeExtract %13 %75 0
%109 = OpBitcast %4 %108
%110 = OpCompositeExtract %13 %75 2
%111 = OpBitcast %4 %110
%112 = OpImageFetch %94 %82 %109 Lod %111
%113 = OpCompositeExtract %4 %91 0
%114 = OpIAdd %94 %95 %98
%115 = OpIAdd %94 %114 %99
%116 = OpIAdd %94 %115 %107
%117 = OpIAdd %94 %116 %112
OpImageWrite %83 %113 %117
OpReturn
OpFunctionEnd
%117 = OpFunction %2 None %73
%114 = OpLabel
%116 = OpLoad %19 %115
%118 = OpLoad %14 %36
%119 = OpLoad %15 %38
%120 = OpLoad %17 %46
OpBranch %121
%121 = OpLabel
%122 = OpImageQuerySize %20 %119
%123 = OpVectorShuffle %82 %116 %116 0 1
%124 = OpBitcast %20 %123
%125 = OpIMul %20 %122 %124
%126 = OpCompositeConstruct %20 %3 %5
%127 = OpSMod %20 %125 %126
%128 = OpCompositeExtract %12 %116 2
%129 = OpBitcast %4 %128
%130 = OpImageFetch %28 %118 %127 Sample %129
%131 = OpCompositeExtract %8 %130 0
%132 = OpCompositeExtract %4 %127 0
%133 = OpConvertFToU %12 %131
%134 = OpCompositeConstruct %90 %133 %133 %133 %133
OpImageWrite %120 %132 %134
%121 = OpFunction %2 None %77
%118 = OpLabel
%120 = OpLoad %20 %119
%122 = OpLoad %15 %38
%123 = OpLoad %16 %40
%124 = OpLoad %18 %48
OpBranch %125
%125 = OpLabel
%126 = OpImageQuerySize %21 %123
%127 = OpVectorShuffle %86 %120 %120 0 1
%128 = OpBitcast %21 %127
%129 = OpIMul %21 %126 %128
%130 = OpCompositeConstruct %21 %3 %5
%131 = OpSMod %21 %129 %130
%132 = OpCompositeExtract %13 %120 2
%133 = OpBitcast %4 %132
%134 = OpImageFetch %29 %122 %131 Sample %133
%135 = OpCompositeExtract %8 %134 0
%136 = OpCompositeExtract %4 %131 0
%137 = OpConvertFToU %13 %135
%138 = OpCompositeConstruct %94 %137 %137 %137 %137
OpImageWrite %124 %136 %138
OpReturn
OpFunctionEnd
%138 = OpFunction %2 None %73
%135 = OpLabel
%139 = OpLoad %21 %48
%140 = OpLoad %22 %50
%141 = OpLoad %23 %52
%142 = OpLoad %24 %54
%143 = OpLoad %25 %56
%144 = OpLoad %26 %58
OpBranch %145
%145 = OpLabel
%147 = OpImageQuerySizeLod %4 %139 %146
%148 = OpBitcast %4 %147
%149 = OpImageQuerySizeLod %4 %139 %148
%150 = OpImageQuerySizeLod %20 %140 %146
%151 = OpImageQuerySizeLod %20 %140 %6
%152 = OpImageQuerySizeLod %101 %141 %146
%153 = OpVectorShuffle %20 %152 %152 0 1
%154 = OpImageQuerySizeLod %101 %141 %6
%155 = OpVectorShuffle %20 %154 %154 0 1
%156 = OpImageQuerySizeLod %20 %142 %146
%157 = OpImageQuerySizeLod %20 %142 %6
%158 = OpImageQuerySizeLod %101 %143 %146
%159 = OpVectorShuffle %20 %158 %158 0 0
%160 = OpImageQuerySizeLod %101 %143 %6
%161 = OpVectorShuffle %20 %160 %160 0 0
%162 = OpImageQuerySizeLod %101 %144 %146
%163 = OpImageQuerySizeLod %101 %144 %6
%164 = OpCompositeExtract %4 %150 1
%165 = OpIAdd %4 %147 %164
%166 = OpCompositeExtract %4 %151 1
%167 = OpIAdd %4 %165 %166
%168 = OpCompositeExtract %4 %153 1
%169 = OpIAdd %4 %167 %168
%142 = OpFunction %2 None %77
%139 = OpLabel
%143 = OpLoad %22 %50
%144 = OpLoad %23 %52
%145 = OpLoad %24 %54
%146 = OpLoad %25 %56
%147 = OpLoad %26 %58
%148 = OpLoad %27 %60
%149 = OpLoad %28 %62
OpBranch %150
%150 = OpLabel
%152 = OpImageQuerySizeLod %4 %143 %151
%153 = OpBitcast %4 %152
%154 = OpImageQuerySizeLod %4 %143 %153
%155 = OpImageQuerySizeLod %21 %144 %151
%156 = OpImageQuerySizeLod %21 %144 %6
%157 = OpImageQuerySizeLod %105 %145 %151
%158 = OpVectorShuffle %21 %157 %157 0 1
%159 = OpImageQuerySizeLod %105 %145 %6
%160 = OpVectorShuffle %21 %159 %159 0 1
%161 = OpImageQuerySizeLod %21 %146 %151
%162 = OpImageQuerySizeLod %21 %146 %6
%163 = OpImageQuerySizeLod %105 %147 %151
%164 = OpVectorShuffle %21 %163 %163 0 0
%165 = OpImageQuerySizeLod %105 %147 %6
%166 = OpVectorShuffle %21 %165 %165 0 0
%167 = OpImageQuerySizeLod %105 %148 %151
%168 = OpImageQuerySizeLod %105 %148 %6
%169 = OpImageQuerySize %21 %149
%170 = OpCompositeExtract %4 %155 1
%171 = OpIAdd %4 %169 %170
%171 = OpIAdd %4 %152 %170
%172 = OpCompositeExtract %4 %156 1
%173 = OpIAdd %4 %171 %172
%174 = OpCompositeExtract %4 %157 1
%174 = OpCompositeExtract %4 %158 1
%175 = OpIAdd %4 %173 %174
%176 = OpCompositeExtract %4 %159 1
%176 = OpCompositeExtract %4 %160 1
%177 = OpIAdd %4 %175 %176
%178 = OpCompositeExtract %4 %161 1
%179 = OpIAdd %4 %177 %178
%180 = OpCompositeExtract %4 %162 2
%180 = OpCompositeExtract %4 %162 1
%181 = OpIAdd %4 %179 %180
%182 = OpCompositeExtract %4 %163 2
%182 = OpCompositeExtract %4 %164 1
%183 = OpIAdd %4 %181 %182
%184 = OpConvertSToF %8 %183
%185 = OpCompositeConstruct %28 %184 %184 %184 %184
OpStore %136 %185
%184 = OpCompositeExtract %4 %166 1
%185 = OpIAdd %4 %183 %184
%186 = OpCompositeExtract %4 %167 2
%187 = OpIAdd %4 %185 %186
%188 = OpCompositeExtract %4 %168 2
%189 = OpIAdd %4 %187 %188
%190 = OpConvertSToF %8 %189
%191 = OpCompositeConstruct %29 %190 %190 %190 %190
OpStore %140 %191
OpReturn
OpFunctionEnd
%188 = OpFunction %2 None %73
%186 = OpLabel
%189 = OpLoad %22 %50
%190 = OpLoad %23 %52
%191 = OpLoad %24 %54
%192 = OpLoad %25 %56
%193 = OpLoad %26 %58
%194 = OpLoad %27 %60
OpBranch %195
%195 = OpLabel
%196 = OpImageQueryLevels %4 %189
%197 = OpImageQueryLevels %4 %190
%198 = OpImageQuerySizeLod %101 %190 %146
%199 = OpCompositeExtract %4 %198 2
%200 = OpImageQueryLevels %4 %191
%201 = OpImageQueryLevels %4 %192
%202 = OpImageQuerySizeLod %101 %192 %146
%203 = OpCompositeExtract %4 %202 2
%204 = OpImageQueryLevels %4 %193
%205 = OpImageQuerySamples %4 %194
%206 = OpIAdd %4 %199 %203
%207 = OpIAdd %4 %206 %205
%208 = OpIAdd %4 %207 %196
%209 = OpIAdd %4 %208 %197
%210 = OpIAdd %4 %209 %204
%211 = OpIAdd %4 %210 %200
%212 = OpIAdd %4 %211 %201
%213 = OpConvertSToF %8 %212
%214 = OpCompositeConstruct %28 %213 %213 %213 %213
OpStore %187 %214
%194 = OpFunction %2 None %77
%192 = OpLabel
%195 = OpLoad %23 %52
%196 = OpLoad %24 %54
%197 = OpLoad %25 %56
%198 = OpLoad %26 %58
%199 = OpLoad %27 %60
%200 = OpLoad %28 %62
OpBranch %201
%201 = OpLabel
%202 = OpImageQueryLevels %4 %195
%203 = OpImageQueryLevels %4 %196
%204 = OpImageQuerySizeLod %105 %196 %151
%205 = OpCompositeExtract %4 %204 2
%206 = OpImageQueryLevels %4 %197
%207 = OpImageQueryLevels %4 %198
%208 = OpImageQuerySizeLod %105 %198 %151
%209 = OpCompositeExtract %4 %208 2
%210 = OpImageQueryLevels %4 %199
%211 = OpImageQuerySamples %4 %200
%212 = OpIAdd %4 %205 %209
%213 = OpIAdd %4 %212 %211
%214 = OpIAdd %4 %213 %202
%215 = OpIAdd %4 %214 %203
%216 = OpIAdd %4 %215 %210
%217 = OpIAdd %4 %216 %206
%218 = OpIAdd %4 %217 %207
%219 = OpConvertSToF %8 %218
%220 = OpCompositeConstruct %29 %219 %219 %219 %219
OpStore %193 %220
OpReturn
OpFunctionEnd
%217 = OpFunction %2 None %73
%215 = OpLabel
%218 = OpLoad %21 %48
%219 = OpLoad %22 %50
%220 = OpLoad %29 %62
OpBranch %221
%223 = OpFunction %2 None %77
%221 = OpLabel
%223 = OpCompositeConstruct %222 %7 %7
%224 = OpCompositeExtract %8 %223 0
%226 = OpSampledImage %225 %218 %220
%227 = OpImageSampleImplicitLod %28 %226 %224
%229 = OpSampledImage %228 %219 %220
%230 = OpImageSampleImplicitLod %28 %229 %223
%231 = OpSampledImage %228 %219 %220
%232 = OpImageSampleImplicitLod %28 %231 %223 ConstOffset %31
%233 = OpSampledImage %228 %219 %220
%234 = OpImageSampleExplicitLod %28 %233 %223 Lod %9
%235 = OpSampledImage %228 %219 %220
%236 = OpImageSampleExplicitLod %28 %235 %223 Lod|ConstOffset %9 %31
%237 = OpFAdd %28 %227 %230
%238 = OpFAdd %28 %237 %232
%239 = OpFAdd %28 %238 %234
%240 = OpFAdd %28 %239 %236
OpStore %216 %240
%224 = OpLoad %22 %50
%225 = OpLoad %23 %52
%226 = OpLoad %30 %64
OpBranch %227
%227 = OpLabel
%229 = OpCompositeConstruct %228 %7 %7
%230 = OpCompositeExtract %8 %229 0
%232 = OpSampledImage %231 %224 %226
%233 = OpImageSampleImplicitLod %29 %232 %230
%235 = OpSampledImage %234 %225 %226
%236 = OpImageSampleImplicitLod %29 %235 %229
%237 = OpSampledImage %234 %225 %226
%238 = OpImageSampleImplicitLod %29 %237 %229 ConstOffset %33
%239 = OpSampledImage %234 %225 %226
%240 = OpImageSampleExplicitLod %29 %239 %229 Lod %9
%241 = OpSampledImage %234 %225 %226
%242 = OpImageSampleExplicitLod %29 %241 %229 Lod|ConstOffset %9 %33
%243 = OpSampledImage %234 %225 %226
%244 = OpImageSampleImplicitLod %29 %243 %229 Bias|ConstOffset %11 %33
%245 = OpFAdd %29 %233 %236
%246 = OpFAdd %29 %245 %238
%247 = OpFAdd %29 %246 %240
%248 = OpFAdd %29 %247 %242
OpStore %222 %248
OpReturn
OpFunctionEnd
%244 = OpFunction %2 None %73
%241 = OpLabel
%245 = OpLoad %29 %64
%246 = OpLoad %30 %66
OpBranch %247
%247 = OpLabel
%248 = OpCompositeConstruct %222 %7 %7
%250 = OpSampledImage %249 %246 %245
%251 = OpImageSampleDrefImplicitLod %8 %250 %248 %7
%252 = OpSampledImage %249 %246 %245
%253 = OpImageSampleDrefExplicitLod %8 %252 %248 %7 Lod %254
%255 = OpFAdd %8 %251 %253
OpStore %242 %255
OpReturn
OpFunctionEnd
%258 = OpFunction %2 None %73
%252 = OpFunction %2 None %77
%249 = OpLabel
%253 = OpLoad %30 %66
%254 = OpLoad %31 %68
%255 = OpLoad %32 %70
OpBranch %256
%256 = OpLabel
%259 = OpLoad %22 %50
%260 = OpLoad %29 %62
%261 = OpLoad %29 %64
%262 = OpLoad %30 %66
OpBranch %263
%263 = OpLabel
%264 = OpCompositeConstruct %222 %7 %7
%265 = OpSampledImage %228 %259 %260
%266 = OpImageGather %28 %265 %264 %267
%268 = OpSampledImage %228 %259 %260
%269 = OpImageGather %28 %268 %264 %270 ConstOffset %31
%271 = OpSampledImage %249 %262 %261
%272 = OpImageDrefGather %28 %271 %264 %7
%273 = OpSampledImage %249 %262 %261
%274 = OpImageDrefGather %28 %273 %264 %7 ConstOffset %31
%275 = OpFAdd %28 %266 %269
%276 = OpFAdd %28 %275 %272
%277 = OpFAdd %28 %276 %274
OpStore %257 %277
%257 = OpCompositeConstruct %228 %7 %7
%259 = OpSampledImage %258 %254 %253
%260 = OpImageSampleDrefImplicitLod %8 %259 %257 %7
%261 = OpSampledImage %258 %254 %253
%262 = OpImageSampleDrefExplicitLod %8 %261 %257 %7 Lod %263
%265 = OpCompositeConstruct %264 %7 %7 %7
%267 = OpSampledImage %266 %255 %253
%268 = OpImageSampleDrefExplicitLod %8 %267 %265 %7 Lod %263
%269 = OpFAdd %8 %260 %262
OpStore %250 %269
OpReturn
OpFunctionEnd
%280 = OpFunction %2 None %73
%278 = OpLabel
%281 = OpLoad %29 %62
%282 = OpLoad %30 %66
OpBranch %283
%283 = OpLabel
%284 = OpCompositeConstruct %222 %7 %7
%285 = OpSampledImage %249 %282 %281
%286 = OpImageSampleImplicitLod %28 %285 %284
%287 = OpCompositeExtract %8 %286 0
%288 = OpSampledImage %249 %282 %281
%289 = OpImageGather %28 %288 %284 %146
%290 = OpCompositeConstruct %28 %287 %287 %287 %287
%291 = OpFAdd %28 %290 %289
OpStore %279 %291
%272 = OpFunction %2 None %77
%270 = OpLabel
%273 = OpLoad %23 %52
%274 = OpLoad %30 %64
%275 = OpLoad %30 %66
%276 = OpLoad %31 %68
OpBranch %277
%277 = OpLabel
%278 = OpCompositeConstruct %228 %7 %7
%279 = OpSampledImage %234 %273 %274
%280 = OpImageGather %29 %279 %278 %281
%282 = OpSampledImage %234 %273 %274
%283 = OpImageGather %29 %282 %278 %284 ConstOffset %33
%285 = OpSampledImage %258 %276 %275
%286 = OpImageDrefGather %29 %285 %278 %7
%287 = OpSampledImage %258 %276 %275
%288 = OpImageDrefGather %29 %287 %278 %7 ConstOffset %33
%289 = OpFAdd %29 %280 %283
%290 = OpFAdd %29 %289 %286
%291 = OpFAdd %29 %290 %288
OpStore %271 %291
OpReturn
OpFunctionEnd
%294 = OpFunction %2 None %77
%292 = OpLabel
%295 = OpLoad %30 %64
%296 = OpLoad %31 %68
OpBranch %297
%297 = OpLabel
%298 = OpCompositeConstruct %228 %7 %7
%299 = OpSampledImage %258 %296 %295
%300 = OpImageSampleImplicitLod %29 %299 %298
%301 = OpCompositeExtract %8 %300 0
%302 = OpSampledImage %258 %296 %295
%303 = OpImageGather %29 %302 %298 %151
%304 = OpCompositeConstruct %29 %301 %301 %301 %301
%305 = OpFAdd %29 %304 %303
OpStore %293 %305
OpReturn
OpFunctionEnd

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

@@ -34,6 +34,8 @@ var sampler_reg: sampler;
var sampler_cmp: sampler_comparison;
@group(1) @binding(2)
var image_2d_depth: texture_depth_2d;
@group(1) @binding(3)
var image_cube_depth: texture_depth_cube;
@stage(compute) @workgroup_size(16, 1, 1)
fn main(@builtin(local_invocation_id) local_id: vec3<u32>) {
@@ -71,6 +73,7 @@ fn queries() -> @builtin(position) vec4<f32> {
let dim_cube_array_lod = textureDimensions(image_cube_array, 1);
let dim_3d = textureDimensions(image_3d);
let dim_3d_lod = textureDimensions(image_3d, 1);
let dim_2s_ms = textureDimensions(image_aa);
let sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z);
return vec4<f32>(f32(sum));
}
@@ -97,6 +100,7 @@ fn sample() -> @location(0) vec4<f32> {
let s2d_offset = textureSample(image_2d, sampler_reg, tc, vec2<i32>(3, 1));
let s2d_level = textureSampleLevel(image_2d, sampler_reg, tc, 2.299999952316284);
let s2d_level_offset = textureSampleLevel(image_2d, sampler_reg, tc, 2.299999952316284, vec2<i32>(3, 1));
let s2d_bias_offset = textureSampleBias(image_2d, sampler_reg, tc, 2.0, vec2<i32>(3, 1));
return ((((s1d + s2d) + s2d_offset) + s2d_level) + s2d_level_offset);
}
@@ -105,6 +109,7 @@ fn sample_comparison() -> @location(0) f32 {
let tc_1 = vec2<f32>(0.5);
let s2d_depth = textureSampleCompare(image_2d_depth, sampler_cmp, tc_1, 0.5);
let s2d_depth_level = textureSampleCompareLevel(image_2d_depth, sampler_cmp, tc_1, 0.5);
let scube_depth_level = textureSampleCompareLevel(image_cube_depth, sampler_cmp, vec3<f32>(0.5), 0.5);
return (s2d_depth + s2d_depth_level);
}