From be83128031408a5f77f4ec26faa22db4f9d450e2 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Mon, 31 Oct 2022 17:33:35 +0000 Subject: [PATCH] tint/transform: Don't hoist textures / samplers in the PromoteSideEffectsToDecl transform Fixed: tint:1739 Change-Id: Iaa5a6e086708a6bba601c59962042650829795f6 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/107683 Kokoro: Kokoro Reviewed-by: Antonio Maiorano Commit-Queue: Ben Clayton --- .../transform/promote_side_effects_to_decl.cc | 8 +- .../promote_side_effects_to_decl_test.cc | 49 ++++ test/tint/bug/tint/1739.wgsl | 13 + .../tint/bug/tint/1739.wgsl.expected.dxc.hlsl | 98 +++++++ .../tint/bug/tint/1739.wgsl.expected.fxc.hlsl | 98 +++++++ test/tint/bug/tint/1739.wgsl.expected.glsl | 66 +++++ test/tint/bug/tint/1739.wgsl.expected.msl | 67 +++++ test/tint/bug/tint/1739.wgsl.expected.spvasm | 258 ++++++++++++++++++ test/tint/bug/tint/1739.wgsl.expected.wgsl | 58 ++++ 9 files changed, 713 insertions(+), 2 deletions(-) create mode 100644 test/tint/bug/tint/1739.wgsl create mode 100644 test/tint/bug/tint/1739.wgsl.expected.dxc.hlsl create mode 100644 test/tint/bug/tint/1739.wgsl.expected.fxc.hlsl create mode 100644 test/tint/bug/tint/1739.wgsl.expected.glsl create mode 100644 test/tint/bug/tint/1739.wgsl.expected.msl create mode 100644 test/tint/bug/tint/1739.wgsl.expected.spvasm create mode 100644 test/tint/bug/tint/1739.wgsl.expected.wgsl diff --git a/src/tint/transform/promote_side_effects_to_decl.cc b/src/tint/transform/promote_side_effects_to_decl.cc index a913838033..e36c416002 100644 --- a/src/tint/transform/promote_side_effects_to_decl.cc +++ b/src/tint/transform/promote_side_effects_to_decl.cc @@ -281,11 +281,15 @@ class DecomposeSideEffects::CollectHoistsState : public StateBase { if (var_user->ConstantValue()) { return false; } - // Don't hoist read-only variables as they cannot receive - // side-effects. + // Don't hoist read-only variables as they cannot receive side-effects. if (var_user->Variable()->Access() == ast::Access::kRead) { return false; } + // Don't hoist textures / samplers as they can't be placed into a let, nor + // can they have side effects. + if (var_user->Variable()->Type()->IsAnyOf()) { + return false; + } return true; } } diff --git a/src/tint/transform/promote_side_effects_to_decl_test.cc b/src/tint/transform/promote_side_effects_to_decl_test.cc index 899fcd8d48..f064834264 100644 --- a/src/tint/transform/promote_side_effects_to_decl_test.cc +++ b/src/tint/transform/promote_side_effects_to_decl_test.cc @@ -4081,5 +4081,54 @@ fn f() { EXPECT_EQ(expect, str(got)); } + +TEST_F(PromoteSideEffectsToDeclTest, TextureSamplerParameter) { + auto* src = R"( +@group(0) @binding(0) var T : texture_2d; +@group(0) @binding(1) var S : sampler; + +var P : vec2; +fn side_effects() -> vec2 { + P += vec2(1.0); + return P; +} + +fn f(t : texture_2d, s : sampler) -> vec4 { + return textureSample(t, s, side_effects()); +} + +fn m() -> vec4{ + return f(T, S); +} +)"; + + auto* expect = R"( +@group(0) @binding(0) var T : texture_2d; + +@group(0) @binding(1) var S : sampler; + +var P : vec2; + +fn side_effects() -> vec2 { + P += vec2(1.0); + return P; +} + +fn f(t : texture_2d, s : sampler) -> vec4 { + let tint_symbol = side_effects(); + return textureSample(t, s, tint_symbol); +} + +fn m() -> vec4 { + return f(T, S); +} +)"; + + DataMap data; + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); +} + } // namespace } // namespace tint::transform diff --git a/test/tint/bug/tint/1739.wgsl b/test/tint/bug/tint/1739.wgsl new file mode 100644 index 0000000000..d513d64794 --- /dev/null +++ b/test/tint/bug/tint/1739.wgsl @@ -0,0 +1,13 @@ +// flags: --transform robustness,multiplaner_external_texture +@group(0) @binding(0) var t : texture_external; + +@group(0) @binding(1) var outImage : texture_storage_2d; + +@compute @workgroup_size(1) +fn main() { + var red : vec4 = textureLoad(t, vec2(10, 10)); + textureStore(outImage, vec2(0, 0), red); + var green : vec4 = textureLoad(t, vec2(70, 118)); + textureStore(outImage, vec2(1, 0), green); + return; +} diff --git a/test/tint/bug/tint/1739.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1739.wgsl.expected.dxc.hlsl new file mode 100644 index 0000000000..d5d51439eb --- /dev/null +++ b/test/tint/bug/tint/1739.wgsl.expected.dxc.hlsl @@ -0,0 +1,98 @@ +struct GammaTransferParams { + float G; + float A; + float B; + float C; + float D; + float E; + float F; + uint padding; +}; +struct ExternalTextureParams { + uint numPlanes; + uint doYuvToRgbConversionOnly; + float3x4 yuvToRgbConversionMatrix; + GammaTransferParams gammaDecodeParams; + GammaTransferParams gammaEncodeParams; + float3x3 gamutConversionMatrix; +}; + +Texture2D ext_tex_plane_1 : register(t2, space0); +cbuffer cbuffer_ext_tex_params : register(b3, space0) { + uint4 ext_tex_params[11]; +}; +Texture2D t : register(t0, space0); +RWTexture2D outImage : register(u1, space0); + +float3 gammaCorrection(float3 v, GammaTransferParams params) { + const bool3 cond = (abs(v) < float3((params.D).xxx)); + const float3 t_1 = (sign(v) * ((params.C * abs(v)) + params.F)); + const float3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), float3((params.G).xxx)) + params.E)); + return (cond ? t_1 : f); +} + +float4 textureLoadExternal(Texture2D plane0, Texture2D plane1, int2 coord, ExternalTextureParams params) { + float3 color = float3(0.0f, 0.0f, 0.0f); + if ((params.numPlanes == 1u)) { + color = plane0.Load(int3(coord, 0)).rgb; + } else { + color = mul(params.yuvToRgbConversionMatrix, float4(plane0.Load(int3(coord, 0)).r, plane1.Load(int3(coord, 0)).rg, 1.0f)); + } + if ((params.doYuvToRgbConversionOnly == 0u)) { + color = gammaCorrection(color, params.gammaDecodeParams); + color = mul(color, params.gamutConversionMatrix); + color = gammaCorrection(color, params.gammaEncodeParams); + } + return float4(color, 1.0f); +} + +float3x4 tint_symbol_2(uint4 buffer[11], uint offset) { + const uint scalar_offset = ((offset + 0u)) / 4; + const uint scalar_offset_1 = ((offset + 16u)) / 4; + const uint scalar_offset_2 = ((offset + 32u)) / 4; + return float3x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4])); +} + +GammaTransferParams tint_symbol_4(uint4 buffer[11], uint offset) { + const uint scalar_offset_3 = ((offset + 0u)) / 4; + const uint scalar_offset_4 = ((offset + 4u)) / 4; + const uint scalar_offset_5 = ((offset + 8u)) / 4; + const uint scalar_offset_6 = ((offset + 12u)) / 4; + const uint scalar_offset_7 = ((offset + 16u)) / 4; + const uint scalar_offset_8 = ((offset + 20u)) / 4; + const uint scalar_offset_9 = ((offset + 24u)) / 4; + const uint scalar_offset_10 = ((offset + 28u)) / 4; + const GammaTransferParams tint_symbol_8 = {asfloat(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), asfloat(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4]), asfloat(buffer[scalar_offset_6 / 4][scalar_offset_6 % 4]), asfloat(buffer[scalar_offset_7 / 4][scalar_offset_7 % 4]), asfloat(buffer[scalar_offset_8 / 4][scalar_offset_8 % 4]), asfloat(buffer[scalar_offset_9 / 4][scalar_offset_9 % 4]), buffer[scalar_offset_10 / 4][scalar_offset_10 % 4]}; + return tint_symbol_8; +} + +float3x3 tint_symbol_6(uint4 buffer[11], uint offset) { + const uint scalar_offset_11 = ((offset + 0u)) / 4; + const uint scalar_offset_12 = ((offset + 16u)) / 4; + const uint scalar_offset_13 = ((offset + 32u)) / 4; + return float3x3(asfloat(buffer[scalar_offset_11 / 4].xyz), asfloat(buffer[scalar_offset_12 / 4].xyz), asfloat(buffer[scalar_offset_13 / 4].xyz)); +} + +ExternalTextureParams tint_symbol(uint4 buffer[11], uint offset) { + const uint scalar_offset_14 = ((offset + 0u)) / 4; + const uint scalar_offset_15 = ((offset + 4u)) / 4; + const ExternalTextureParams tint_symbol_9 = {buffer[scalar_offset_14 / 4][scalar_offset_14 % 4], buffer[scalar_offset_15 / 4][scalar_offset_15 % 4], tint_symbol_2(buffer, (offset + 16u)), tint_symbol_4(buffer, (offset + 64u)), tint_symbol_4(buffer, (offset + 96u)), tint_symbol_6(buffer, (offset + 128u))}; + return tint_symbol_9; +} + +[numthreads(1, 1, 1)] +void main() { + int2 tint_tmp; + t.GetDimensions(tint_tmp.x, tint_tmp.y); + float4 red = textureLoadExternal(t, ext_tex_plane_1, clamp((10).xx, (0).xx, int2((uint2(tint_tmp) - (1u).xx))), tint_symbol(ext_tex_params, 0u)); + int2 tint_tmp_1; + outImage.GetDimensions(tint_tmp_1.x, tint_tmp_1.y); + outImage[clamp((0).xx, (0).xx, int2((uint2(tint_tmp_1) - (1u).xx)))] = red; + int2 tint_tmp_2; + t.GetDimensions(tint_tmp_2.x, tint_tmp_2.y); + float4 green = textureLoadExternal(t, ext_tex_plane_1, clamp(int2(70, 118), (0).xx, int2((uint2(tint_tmp_2) - (1u).xx))), tint_symbol(ext_tex_params, 0u)); + int2 tint_tmp_3; + outImage.GetDimensions(tint_tmp_3.x, tint_tmp_3.y); + outImage[clamp(int2(1, 0), (0).xx, int2((uint2(tint_tmp_3) - (1u).xx)))] = green; + return; +} diff --git a/test/tint/bug/tint/1739.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1739.wgsl.expected.fxc.hlsl new file mode 100644 index 0000000000..d5d51439eb --- /dev/null +++ b/test/tint/bug/tint/1739.wgsl.expected.fxc.hlsl @@ -0,0 +1,98 @@ +struct GammaTransferParams { + float G; + float A; + float B; + float C; + float D; + float E; + float F; + uint padding; +}; +struct ExternalTextureParams { + uint numPlanes; + uint doYuvToRgbConversionOnly; + float3x4 yuvToRgbConversionMatrix; + GammaTransferParams gammaDecodeParams; + GammaTransferParams gammaEncodeParams; + float3x3 gamutConversionMatrix; +}; + +Texture2D ext_tex_plane_1 : register(t2, space0); +cbuffer cbuffer_ext_tex_params : register(b3, space0) { + uint4 ext_tex_params[11]; +}; +Texture2D t : register(t0, space0); +RWTexture2D outImage : register(u1, space0); + +float3 gammaCorrection(float3 v, GammaTransferParams params) { + const bool3 cond = (abs(v) < float3((params.D).xxx)); + const float3 t_1 = (sign(v) * ((params.C * abs(v)) + params.F)); + const float3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), float3((params.G).xxx)) + params.E)); + return (cond ? t_1 : f); +} + +float4 textureLoadExternal(Texture2D plane0, Texture2D plane1, int2 coord, ExternalTextureParams params) { + float3 color = float3(0.0f, 0.0f, 0.0f); + if ((params.numPlanes == 1u)) { + color = plane0.Load(int3(coord, 0)).rgb; + } else { + color = mul(params.yuvToRgbConversionMatrix, float4(plane0.Load(int3(coord, 0)).r, plane1.Load(int3(coord, 0)).rg, 1.0f)); + } + if ((params.doYuvToRgbConversionOnly == 0u)) { + color = gammaCorrection(color, params.gammaDecodeParams); + color = mul(color, params.gamutConversionMatrix); + color = gammaCorrection(color, params.gammaEncodeParams); + } + return float4(color, 1.0f); +} + +float3x4 tint_symbol_2(uint4 buffer[11], uint offset) { + const uint scalar_offset = ((offset + 0u)) / 4; + const uint scalar_offset_1 = ((offset + 16u)) / 4; + const uint scalar_offset_2 = ((offset + 32u)) / 4; + return float3x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4])); +} + +GammaTransferParams tint_symbol_4(uint4 buffer[11], uint offset) { + const uint scalar_offset_3 = ((offset + 0u)) / 4; + const uint scalar_offset_4 = ((offset + 4u)) / 4; + const uint scalar_offset_5 = ((offset + 8u)) / 4; + const uint scalar_offset_6 = ((offset + 12u)) / 4; + const uint scalar_offset_7 = ((offset + 16u)) / 4; + const uint scalar_offset_8 = ((offset + 20u)) / 4; + const uint scalar_offset_9 = ((offset + 24u)) / 4; + const uint scalar_offset_10 = ((offset + 28u)) / 4; + const GammaTransferParams tint_symbol_8 = {asfloat(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4]), asfloat(buffer[scalar_offset_4 / 4][scalar_offset_4 % 4]), asfloat(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4]), asfloat(buffer[scalar_offset_6 / 4][scalar_offset_6 % 4]), asfloat(buffer[scalar_offset_7 / 4][scalar_offset_7 % 4]), asfloat(buffer[scalar_offset_8 / 4][scalar_offset_8 % 4]), asfloat(buffer[scalar_offset_9 / 4][scalar_offset_9 % 4]), buffer[scalar_offset_10 / 4][scalar_offset_10 % 4]}; + return tint_symbol_8; +} + +float3x3 tint_symbol_6(uint4 buffer[11], uint offset) { + const uint scalar_offset_11 = ((offset + 0u)) / 4; + const uint scalar_offset_12 = ((offset + 16u)) / 4; + const uint scalar_offset_13 = ((offset + 32u)) / 4; + return float3x3(asfloat(buffer[scalar_offset_11 / 4].xyz), asfloat(buffer[scalar_offset_12 / 4].xyz), asfloat(buffer[scalar_offset_13 / 4].xyz)); +} + +ExternalTextureParams tint_symbol(uint4 buffer[11], uint offset) { + const uint scalar_offset_14 = ((offset + 0u)) / 4; + const uint scalar_offset_15 = ((offset + 4u)) / 4; + const ExternalTextureParams tint_symbol_9 = {buffer[scalar_offset_14 / 4][scalar_offset_14 % 4], buffer[scalar_offset_15 / 4][scalar_offset_15 % 4], tint_symbol_2(buffer, (offset + 16u)), tint_symbol_4(buffer, (offset + 64u)), tint_symbol_4(buffer, (offset + 96u)), tint_symbol_6(buffer, (offset + 128u))}; + return tint_symbol_9; +} + +[numthreads(1, 1, 1)] +void main() { + int2 tint_tmp; + t.GetDimensions(tint_tmp.x, tint_tmp.y); + float4 red = textureLoadExternal(t, ext_tex_plane_1, clamp((10).xx, (0).xx, int2((uint2(tint_tmp) - (1u).xx))), tint_symbol(ext_tex_params, 0u)); + int2 tint_tmp_1; + outImage.GetDimensions(tint_tmp_1.x, tint_tmp_1.y); + outImage[clamp((0).xx, (0).xx, int2((uint2(tint_tmp_1) - (1u).xx)))] = red; + int2 tint_tmp_2; + t.GetDimensions(tint_tmp_2.x, tint_tmp_2.y); + float4 green = textureLoadExternal(t, ext_tex_plane_1, clamp(int2(70, 118), (0).xx, int2((uint2(tint_tmp_2) - (1u).xx))), tint_symbol(ext_tex_params, 0u)); + int2 tint_tmp_3; + outImage.GetDimensions(tint_tmp_3.x, tint_tmp_3.y); + outImage[clamp(int2(1, 0), (0).xx, int2((uint2(tint_tmp_3) - (1u).xx)))] = green; + return; +} diff --git a/test/tint/bug/tint/1739.wgsl.expected.glsl b/test/tint/bug/tint/1739.wgsl.expected.glsl new file mode 100644 index 0000000000..913050b808 --- /dev/null +++ b/test/tint/bug/tint/1739.wgsl.expected.glsl @@ -0,0 +1,66 @@ +#version 310 es + +struct GammaTransferParams { + float G; + float A; + float B; + float C; + float D; + float E; + float F; + uint padding; +}; + +struct ExternalTextureParams { + uint numPlanes; + uint doYuvToRgbConversionOnly; + uint pad; + uint pad_1; + mat3x4 yuvToRgbConversionMatrix; + GammaTransferParams gammaDecodeParams; + GammaTransferParams gammaEncodeParams; + mat3 gamutConversionMatrix; +}; + +layout(binding = 3, std140) uniform ext_tex_params_block_ubo { + ExternalTextureParams inner; +} ext_tex_params; + +layout(rgba8) uniform highp writeonly image2D outImage; +vec3 gammaCorrection(vec3 v, GammaTransferParams params) { + bvec3 cond = lessThan(abs(v), vec3(params.D)); + vec3 t_1 = (sign(v) * ((params.C * abs(v)) + params.F)); + vec3 f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3(params.G)) + params.E)); + return mix(f, t_1, cond); +} + +vec4 textureLoadExternal(highp sampler2D plane0_1, highp sampler2D plane1_1, ivec2 coord, ExternalTextureParams params) { + vec3 color = vec3(0.0f, 0.0f, 0.0f); + if ((params.numPlanes == 1u)) { + color = texelFetch(plane0_1, coord, 0).rgb; + } else { + color = (vec4(texelFetch(plane0_1, coord, 0).r, texelFetch(plane1_1, coord, 0).rg, 1.0f) * params.yuvToRgbConversionMatrix); + } + if ((params.doYuvToRgbConversionOnly == 0u)) { + color = gammaCorrection(color, params.gammaDecodeParams); + color = (params.gamutConversionMatrix * color); + color = gammaCorrection(color, params.gammaEncodeParams); + } + return vec4(color, 1.0f); +} + +uniform highp sampler2D t_2; +uniform highp sampler2D ext_tex_plane_1_1; +void tint_symbol() { + vec4 red = textureLoadExternal(t_2, ext_tex_plane_1_1, clamp(ivec2(10), ivec2(0), ivec2((uvec2(uvec2(textureSize(t_2, 0))) - uvec2(1u)))), ext_tex_params.inner); + imageStore(outImage, clamp(ivec2(0), ivec2(0), ivec2((uvec2(uvec2(imageSize(outImage))) - uvec2(1u)))), red); + vec4 green = textureLoadExternal(t_2, ext_tex_plane_1_1, clamp(ivec2(70, 118), ivec2(0), ivec2((uvec2(uvec2(textureSize(t_2, 0))) - uvec2(1u)))), ext_tex_params.inner); + imageStore(outImage, clamp(ivec2(1, 0), ivec2(0), ivec2((uvec2(uvec2(imageSize(outImage))) - uvec2(1u)))), green); + return; +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + tint_symbol(); + return; +} diff --git a/test/tint/bug/tint/1739.wgsl.expected.msl b/test/tint/bug/tint/1739.wgsl.expected.msl new file mode 100644 index 0000000000..bcd3a46273 --- /dev/null +++ b/test/tint/bug/tint/1739.wgsl.expected.msl @@ -0,0 +1,67 @@ +#include + +using namespace metal; + +template +struct tint_array { + const constant T& operator[](size_t i) const constant { return elements[i]; } + device T& operator[](size_t i) device { return elements[i]; } + const device T& operator[](size_t i) const device { return elements[i]; } + thread T& operator[](size_t i) thread { return elements[i]; } + const thread T& operator[](size_t i) const thread { return elements[i]; } + threadgroup T& operator[](size_t i) threadgroup { return elements[i]; } + const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; } + T elements[N]; +}; + +struct GammaTransferParams { + /* 0x0000 */ float G; + /* 0x0004 */ float A; + /* 0x0008 */ float B; + /* 0x000c */ float C; + /* 0x0010 */ float D; + /* 0x0014 */ float E; + /* 0x0018 */ float F; + /* 0x001c */ uint padding; +}; + +struct ExternalTextureParams { + /* 0x0000 */ uint numPlanes; + /* 0x0004 */ uint doYuvToRgbConversionOnly; + /* 0x0008 */ tint_array tint_pad; + /* 0x0010 */ float3x4 yuvToRgbConversionMatrix; + /* 0x0040 */ GammaTransferParams gammaDecodeParams; + /* 0x0060 */ GammaTransferParams gammaEncodeParams; + /* 0x0080 */ float3x3 gamutConversionMatrix; +}; + +float3 gammaCorrection(float3 v, GammaTransferParams params) { + bool3 const cond = (fabs(v) < float3(params.D)); + float3 const t_1 = (sign(v) * ((params.C * fabs(v)) + params.F)); + float3 const f = (sign(v) * (pow(((params.A * fabs(v)) + params.B), float3(params.G)) + params.E)); + return select(f, t_1, cond); +} + +float4 textureLoadExternal(texture2d plane0, texture2d plane1, int2 coord, ExternalTextureParams params) { + float3 color = 0.0f; + if ((params.numPlanes == 1u)) { + color = float4(plane0.read(uint2(coord), 0)).rgb; + } else { + color = (float4(plane0.read(uint2(coord), 0)[0], float4(plane1.read(uint2(coord), 0)).rg, 1.0f) * params.yuvToRgbConversionMatrix); + } + if ((params.doYuvToRgbConversionOnly == 0u)) { + color = gammaCorrection(color, params.gammaDecodeParams); + color = (params.gamutConversionMatrix * color); + color = gammaCorrection(color, params.gammaEncodeParams); + } + return float4(color, 1.0f); +} + +kernel void tint_symbol(texture2d tint_symbol_1 [[texture(0)]], texture2d tint_symbol_2 [[texture(1)]], const constant ExternalTextureParams* tint_symbol_3 [[buffer(0)]], texture2d tint_symbol_4 [[texture(2)]]) { + float4 red = textureLoadExternal(tint_symbol_1, tint_symbol_2, clamp(int2(10), int2(0), int2((uint2(uint2(tint_symbol_1.get_width(), tint_symbol_1.get_height())) - uint2(1u)))), *(tint_symbol_3)); + tint_symbol_4.write(red, uint2(clamp(int2(0), int2(0), int2((uint2(uint2(tint_symbol_4.get_width(), tint_symbol_4.get_height())) - uint2(1u)))))); + float4 green = textureLoadExternal(tint_symbol_1, tint_symbol_2, clamp(int2(70, 118), int2(0), int2((uint2(uint2(tint_symbol_1.get_width(), tint_symbol_1.get_height())) - uint2(1u)))), *(tint_symbol_3)); + tint_symbol_4.write(green, uint2(clamp(int2(1, 0), int2(0), int2((uint2(uint2(tint_symbol_4.get_width(), tint_symbol_4.get_height())) - uint2(1u)))))); + return; +} + diff --git a/test/tint/bug/tint/1739.wgsl.expected.spvasm b/test/tint/bug/tint/1739.wgsl.expected.spvasm new file mode 100644 index 0000000000..d40805afc8 --- /dev/null +++ b/test/tint/bug/tint/1739.wgsl.expected.spvasm @@ -0,0 +1,258 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 169 +; Schema: 0 + OpCapability Shader + OpCapability ImageQuery + %25 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 1 1 1 + OpName %ext_tex_plane_1 "ext_tex_plane_1" + OpName %ext_tex_params_block "ext_tex_params_block" + OpMemberName %ext_tex_params_block 0 "inner" + OpName %ExternalTextureParams "ExternalTextureParams" + OpMemberName %ExternalTextureParams 0 "numPlanes" + OpMemberName %ExternalTextureParams 1 "doYuvToRgbConversionOnly" + OpMemberName %ExternalTextureParams 2 "yuvToRgbConversionMatrix" + OpMemberName %ExternalTextureParams 3 "gammaDecodeParams" + OpName %GammaTransferParams "GammaTransferParams" + OpMemberName %GammaTransferParams 0 "G" + OpMemberName %GammaTransferParams 1 "A" + OpMemberName %GammaTransferParams 2 "B" + OpMemberName %GammaTransferParams 3 "C" + OpMemberName %GammaTransferParams 4 "D" + OpMemberName %GammaTransferParams 5 "E" + OpMemberName %GammaTransferParams 6 "F" + OpMemberName %GammaTransferParams 7 "padding" + OpMemberName %ExternalTextureParams 4 "gammaEncodeParams" + OpMemberName %ExternalTextureParams 5 "gamutConversionMatrix" + OpName %ext_tex_params "ext_tex_params" + OpName %t "t" + OpName %outImage "outImage" + OpName %gammaCorrection "gammaCorrection" + OpName %v "v" + OpName %params "params" + OpName %textureLoadExternal "textureLoadExternal" + OpName %plane0 "plane0" + OpName %plane1 "plane1" + OpName %coord "coord" + OpName %params_0 "params" + OpName %color "color" + OpName %main "main" + OpName %red "red" + OpName %green "green" + OpDecorate %ext_tex_plane_1 DescriptorSet 0 + OpDecorate %ext_tex_plane_1 Binding 2 + OpDecorate %ext_tex_params_block Block + OpMemberDecorate %ext_tex_params_block 0 Offset 0 + OpMemberDecorate %ExternalTextureParams 0 Offset 0 + OpMemberDecorate %ExternalTextureParams 1 Offset 4 + OpMemberDecorate %ExternalTextureParams 2 Offset 16 + OpMemberDecorate %ExternalTextureParams 2 ColMajor + OpMemberDecorate %ExternalTextureParams 2 MatrixStride 16 + OpMemberDecorate %ExternalTextureParams 3 Offset 64 + OpMemberDecorate %GammaTransferParams 0 Offset 0 + OpMemberDecorate %GammaTransferParams 1 Offset 4 + OpMemberDecorate %GammaTransferParams 2 Offset 8 + OpMemberDecorate %GammaTransferParams 3 Offset 12 + OpMemberDecorate %GammaTransferParams 4 Offset 16 + OpMemberDecorate %GammaTransferParams 5 Offset 20 + OpMemberDecorate %GammaTransferParams 6 Offset 24 + OpMemberDecorate %GammaTransferParams 7 Offset 28 + OpMemberDecorate %ExternalTextureParams 4 Offset 96 + OpMemberDecorate %ExternalTextureParams 5 Offset 128 + OpMemberDecorate %ExternalTextureParams 5 ColMajor + OpMemberDecorate %ExternalTextureParams 5 MatrixStride 16 + OpDecorate %ext_tex_params NonWritable + OpDecorate %ext_tex_params DescriptorSet 0 + OpDecorate %ext_tex_params Binding 3 + OpDecorate %t DescriptorSet 0 + OpDecorate %t Binding 0 + OpDecorate %outImage NonReadable + OpDecorate %outImage DescriptorSet 0 + OpDecorate %outImage Binding 1 + %float = OpTypeFloat 32 + %3 = OpTypeImage %float 2D 0 0 0 1 Unknown +%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3 +%ext_tex_plane_1 = OpVariable %_ptr_UniformConstant_3 UniformConstant + %uint = OpTypeInt 32 0 + %v4float = OpTypeVector %float 4 +%mat3v4float = OpTypeMatrix %v4float 3 +%GammaTransferParams = OpTypeStruct %float %float %float %float %float %float %float %uint + %v3float = OpTypeVector %float 3 +%mat3v3float = OpTypeMatrix %v3float 3 +%ExternalTextureParams = OpTypeStruct %uint %uint %mat3v4float %GammaTransferParams %GammaTransferParams %mat3v3float +%ext_tex_params_block = OpTypeStruct %ExternalTextureParams +%_ptr_Uniform_ext_tex_params_block = OpTypePointer Uniform %ext_tex_params_block +%ext_tex_params = OpVariable %_ptr_Uniform_ext_tex_params_block Uniform + %t = OpVariable %_ptr_UniformConstant_3 UniformConstant + %18 = OpTypeImage %float 2D 0 0 0 2 Rgba8 +%_ptr_UniformConstant_18 = OpTypePointer UniformConstant %18 + %outImage = OpVariable %_ptr_UniformConstant_18 UniformConstant + %19 = OpTypeFunction %v3float %v3float %GammaTransferParams + %bool = OpTypeBool + %v3bool = OpTypeVector %bool 3 +%_ptr_Function_v3float = OpTypePointer Function %v3float + %39 = OpConstantNull %v3float + %int = OpTypeInt 32 1 + %v2int = OpTypeVector %int 2 + %59 = OpTypeFunction %v4float %3 %3 %v2int %ExternalTextureParams + %uint_1 = OpConstant %uint 1 + %76 = OpConstantNull %int + %v2float = OpTypeVector %float 2 + %float_1 = OpConstant %float 1 + %90 = OpConstantNull %uint + %void = OpTypeVoid + %108 = OpTypeFunction %void + %int_10 = OpConstant %int 10 + %117 = OpConstantComposite %v2int %int_10 %int_10 + %118 = OpConstantNull %v2int + %v2uint = OpTypeVector %uint 2 + %int_0 = OpConstant %int 0 + %125 = OpConstantComposite %v2uint %uint_1 %uint_1 + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_ExternalTextureParams = OpTypePointer Uniform %ExternalTextureParams +%_ptr_Function_v4float = OpTypePointer Function %v4float + %133 = OpConstantNull %v4float + %int_70 = OpConstant %int 70 + %int_118 = OpConstant %int 118 + %149 = OpConstantComposite %v2int %int_70 %int_118 + %int_1 = OpConstant %int 1 + %162 = OpConstantComposite %v2int %int_1 %76 +%gammaCorrection = OpFunction %v3float None %19 + %v = OpFunctionParameter %v3float + %params = OpFunctionParameter %GammaTransferParams + %23 = OpLabel + %37 = OpVariable %_ptr_Function_v3float Function %39 + %49 = OpVariable %_ptr_Function_v3float Function %39 + %55 = OpVariable %_ptr_Function_v3float Function %39 + %24 = OpExtInst %v3float %25 FAbs %v + %26 = OpCompositeExtract %float %params 4 + %27 = OpCompositeConstruct %v3float %26 %26 %26 + %28 = OpFOrdLessThan %v3bool %24 %27 + %31 = OpExtInst %v3float %25 FSign %v + %32 = OpCompositeExtract %float %params 3 + %33 = OpExtInst %v3float %25 FAbs %v + %34 = OpVectorTimesScalar %v3float %33 %32 + %35 = OpCompositeExtract %float %params 6 + %40 = OpCompositeConstruct %v3float %35 %35 %35 + %36 = OpFAdd %v3float %34 %40 + %41 = OpFMul %v3float %31 %36 + %42 = OpExtInst %v3float %25 FSign %v + %44 = OpCompositeExtract %float %params 1 + %45 = OpExtInst %v3float %25 FAbs %v + %46 = OpVectorTimesScalar %v3float %45 %44 + %47 = OpCompositeExtract %float %params 2 + %50 = OpCompositeConstruct %v3float %47 %47 %47 + %48 = OpFAdd %v3float %46 %50 + %51 = OpCompositeExtract %float %params 0 + %52 = OpCompositeConstruct %v3float %51 %51 %51 + %43 = OpExtInst %v3float %25 Pow %48 %52 + %53 = OpCompositeExtract %float %params 5 + %56 = OpCompositeConstruct %v3float %53 %53 %53 + %54 = OpFAdd %v3float %43 %56 + %57 = OpFMul %v3float %42 %54 + %58 = OpSelect %v3float %28 %41 %57 + OpReturnValue %58 + OpFunctionEnd +%textureLoadExternal = OpFunction %v4float None %59 + %plane0 = OpFunctionParameter %3 + %plane1 = OpFunctionParameter %3 + %coord = OpFunctionParameter %v2int + %params_0 = OpFunctionParameter %ExternalTextureParams + %67 = OpLabel + %color = OpVariable %_ptr_Function_v3float Function %39 + %69 = OpCompositeExtract %uint %params_0 0 + %71 = OpIEqual %bool %69 %uint_1 + OpSelectionMerge %72 None + OpBranchConditional %71 %73 %74 + %73 = OpLabel + %75 = OpImageFetch %v4float %plane0 %coord Lod %76 + %77 = OpVectorShuffle %v3float %75 %75 0 1 2 + OpStore %color %77 + OpBranch %72 + %74 = OpLabel + %78 = OpImageFetch %v4float %plane0 %coord Lod %76 + %79 = OpCompositeExtract %float %78 0 + %80 = OpImageFetch %v4float %plane1 %coord Lod %76 + %82 = OpVectorShuffle %v2float %80 %80 0 1 + %83 = OpCompositeExtract %float %82 0 + %84 = OpCompositeExtract %float %82 1 + %86 = OpCompositeConstruct %v4float %79 %83 %84 %float_1 + %87 = OpCompositeExtract %mat3v4float %params_0 2 + %88 = OpVectorTimesMatrix %v3float %86 %87 + OpStore %color %88 + OpBranch %72 + %72 = OpLabel + %89 = OpCompositeExtract %uint %params_0 1 + %91 = OpIEqual %bool %89 %90 + OpSelectionMerge %92 None + OpBranchConditional %91 %93 %92 + %93 = OpLabel + %95 = OpLoad %v3float %color + %96 = OpCompositeExtract %GammaTransferParams %params_0 3 + %94 = OpFunctionCall %v3float %gammaCorrection %95 %96 + OpStore %color %94 + %97 = OpCompositeExtract %mat3v3float %params_0 5 + %98 = OpLoad %v3float %color + %99 = OpMatrixTimesVector %v3float %97 %98 + OpStore %color %99 + %101 = OpLoad %v3float %color + %102 = OpCompositeExtract %GammaTransferParams %params_0 4 + %100 = OpFunctionCall %v3float %gammaCorrection %101 %102 + OpStore %color %100 + OpBranch %92 + %92 = OpLabel + %103 = OpLoad %v3float %color + %104 = OpCompositeExtract %float %103 0 + %105 = OpCompositeExtract %float %103 1 + %106 = OpCompositeExtract %float %103 2 + %107 = OpCompositeConstruct %v4float %104 %105 %106 %float_1 + OpReturnValue %107 + OpFunctionEnd + %main = OpFunction %void None %108 + %111 = OpLabel + %red = OpVariable %_ptr_Function_v4float Function %133 + %green = OpVariable %_ptr_Function_v4float Function %133 + %113 = OpLoad %3 %t + %114 = OpLoad %3 %ext_tex_plane_1 + %123 = OpLoad %3 %t + %122 = OpImageQuerySizeLod %v2uint %123 %int_0 + %126 = OpISub %v2uint %122 %125 + %119 = OpBitcast %v2int %126 + %115 = OpExtInst %v2int %25 SClamp %117 %118 %119 + %129 = OpAccessChain %_ptr_Uniform_ExternalTextureParams %ext_tex_params %uint_0 + %130 = OpLoad %ExternalTextureParams %129 + %112 = OpFunctionCall %v4float %textureLoadExternal %113 %114 %115 %130 + OpStore %red %112 + %135 = OpLoad %18 %outImage + %140 = OpLoad %18 %outImage + %139 = OpImageQuerySize %v2uint %140 + %141 = OpISub %v2uint %139 %125 + %137 = OpBitcast %v2int %141 + %136 = OpExtInst %v2int %25 SClamp %118 %118 %137 + %142 = OpLoad %v4float %red + OpImageWrite %135 %136 %142 + %144 = OpLoad %3 %t + %145 = OpLoad %3 %ext_tex_plane_1 + %153 = OpLoad %3 %t + %152 = OpImageQuerySizeLod %v2uint %153 %int_0 + %154 = OpISub %v2uint %152 %125 + %150 = OpBitcast %v2int %154 + %146 = OpExtInst %v2int %25 SClamp %149 %118 %150 + %155 = OpAccessChain %_ptr_Uniform_ExternalTextureParams %ext_tex_params %uint_0 + %156 = OpLoad %ExternalTextureParams %155 + %143 = OpFunctionCall %v4float %textureLoadExternal %144 %145 %146 %156 + OpStore %green %143 + %159 = OpLoad %18 %outImage + %166 = OpLoad %18 %outImage + %165 = OpImageQuerySize %v2uint %166 + %167 = OpISub %v2uint %165 %125 + %163 = OpBitcast %v2int %167 + %160 = OpExtInst %v2int %25 SClamp %162 %118 %163 + %168 = OpLoad %v4float %green + OpImageWrite %159 %160 %168 + OpReturn + OpFunctionEnd diff --git a/test/tint/bug/tint/1739.wgsl.expected.wgsl b/test/tint/bug/tint/1739.wgsl.expected.wgsl new file mode 100644 index 0000000000..d6bd056533 --- /dev/null +++ b/test/tint/bug/tint/1739.wgsl.expected.wgsl @@ -0,0 +1,58 @@ +struct GammaTransferParams { + G : f32, + A : f32, + B : f32, + C : f32, + D : f32, + E : f32, + F : f32, + padding : u32, +} + +struct ExternalTextureParams { + numPlanes : u32, + doYuvToRgbConversionOnly : u32, + yuvToRgbConversionMatrix : mat3x4, + gammaDecodeParams : GammaTransferParams, + gammaEncodeParams : GammaTransferParams, + gamutConversionMatrix : mat3x3, +} + +@group(0) @binding(2) var ext_tex_plane_1 : texture_2d; + +@group(0) @binding(3) var ext_tex_params : ExternalTextureParams; + +@group(0) @binding(0) var t : texture_2d; + +@group(0) @binding(1) var outImage : texture_storage_2d; + +fn gammaCorrection(v : vec3, params : GammaTransferParams) -> vec3 { + let cond = (abs(v) < vec3(params.D)); + let t = (sign(v) * ((params.C * abs(v)) + params.F)); + let f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3(params.G)) + params.E)); + return select(f, t, cond); +} + +fn textureLoadExternal(plane0 : texture_2d, plane1 : texture_2d, coord : vec2, params : ExternalTextureParams) -> vec4 { + var color : vec3; + if ((params.numPlanes == 1)) { + color = textureLoad(plane0, coord, 0).rgb; + } else { + color = (vec4(textureLoad(plane0, coord, 0).r, textureLoad(plane1, coord, 0).rg, 1) * params.yuvToRgbConversionMatrix); + } + if ((params.doYuvToRgbConversionOnly == 0)) { + color = gammaCorrection(color, params.gammaDecodeParams); + color = (params.gamutConversionMatrix * color); + color = gammaCorrection(color, params.gammaEncodeParams); + } + return vec4(color, 1); +} + +@compute @workgroup_size(1) +fn main() { + var red : vec4 = textureLoadExternal(t, ext_tex_plane_1, clamp(vec2(10, 10), vec2(0), vec2((vec2(textureDimensions(t)) - vec2(1)))), ext_tex_params); + textureStore(outImage, clamp(vec2(0, 0), vec2(0), vec2((vec2(textureDimensions(outImage)) - vec2(1)))), red); + var green : vec4 = textureLoadExternal(t, ext_tex_plane_1, clamp(vec2(70, 118), vec2(0), vec2((vec2(textureDimensions(t)) - vec2(1)))), ext_tex_params); + textureStore(outImage, clamp(vec2(1, 0), vec2(0), vec2((vec2(textureDimensions(outImage)) - vec2(1)))), green); + return; +}