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 <noreply+kokoro@google.com>
Reviewed-by: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>

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
+                        // Don't hoist read-only variables as they cannot receive side-effects.
-                        // 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<sem::Texture, sem::Sampler>()) {
+                            return false;
+                        }
                         return true;
                     }
                 }

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<f32>;
+@group(0) @binding(1) var S : sampler;
+
+var<private> P : vec2<f32>;
+fn side_effects() -> vec2<f32> {
+  P += vec2(1.0);
+  return P;
+}
+
+fn f(t : texture_2d<f32>, s : sampler) -> vec4<f32> {
+  return textureSample(t, s, side_effects());
+}
+
+fn m() -> vec4<f32>{
+  return f(T, S);
+}
+)";
+
+    auto* expect = R"(
+@group(0) @binding(0) var T : texture_2d<f32>;
+
+@group(0) @binding(1) var S : sampler;
+
+var<private> P : vec2<f32>;
+
+fn side_effects() -> vec2<f32> {
+  P += vec2(1.0);
+  return P;
+}
+
+fn f(t : texture_2d<f32>, s : sampler) -> vec4<f32> {
+  let tint_symbol = side_effects();
+  return textureSample(t, s, tint_symbol);
+}
+
+fn m() -> vec4<f32> {
+  return f(T, S);
+}
+)";
+
+    DataMap data;
+    auto got = Run<PromoteSideEffectsToDecl>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace tint::transform

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<rgba8unorm, write>;
+
+@compute @workgroup_size(1)
+fn main() {
+  var red : vec4<f32> = textureLoad(t, vec2<i32>(10, 10));
+  textureStore(outImage, vec2<i32>(0, 0), red);
+  var green : vec4<f32> = textureLoad(t, vec2<i32>(70, 118));
+  textureStore(outImage, vec2<i32>(1, 0), green);
+  return;
+}

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<float4> ext_tex_plane_1 : register(t2, space0);
+cbuffer cbuffer_ext_tex_params : register(b3, space0) {
+  uint4 ext_tex_params[11];
+};
+Texture2D<float4> t : register(t0, space0);
+RWTexture2D<float4> 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<float4> plane0, Texture2D<float4> 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;
+}

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<float4> ext_tex_plane_1 : register(t2, space0);
+cbuffer cbuffer_ext_tex_params : register(b3, space0) {
+  uint4 ext_tex_params[11];
+};
+Texture2D<float4> t : register(t0, space0);
+RWTexture2D<float4> 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<float4> plane0, Texture2D<float4> 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;
+}

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

test/tint/bug/tint/1739.wgsl.expected.msl

@@ -0,0 +1,67 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+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<int8_t, 8> 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<float, access::sample> plane0, texture2d<float, access::sample> 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<float, access::sample> tint_symbol_1 [[texture(0)]], texture2d<float, access::sample> tint_symbol_2 [[texture(1)]], const constant ExternalTextureParams* tint_symbol_3 [[buffer(0)]], texture2d<float, access::write> 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;
+}
+

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

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<f32>,
+  gammaDecodeParams : GammaTransferParams,
+  gammaEncodeParams : GammaTransferParams,
+  gamutConversionMatrix : mat3x3<f32>,
+}
+
+@group(0) @binding(2) var ext_tex_plane_1 : texture_2d<f32>;
+
+@group(0) @binding(3) var<uniform> ext_tex_params : ExternalTextureParams;
+
+@group(0) @binding(0) var t : texture_2d<f32>;
+
+@group(0) @binding(1) var outImage : texture_storage_2d<rgba8unorm, write>;
+
+fn gammaCorrection(v : vec3<f32>, params : GammaTransferParams) -> vec3<f32> {
+  let cond = (abs(v) < vec3<f32>(params.D));
+  let t = (sign(v) * ((params.C * abs(v)) + params.F));
+  let f = (sign(v) * (pow(((params.A * abs(v)) + params.B), vec3<f32>(params.G)) + params.E));
+  return select(f, t, cond);
+}
+
+fn textureLoadExternal(plane0 : texture_2d<f32>, plane1 : texture_2d<f32>, coord : vec2<i32>, params : ExternalTextureParams) -> vec4<f32> {
+  var color : vec3<f32>;
+  if ((params.numPlanes == 1)) {
+    color = textureLoad(plane0, coord, 0).rgb;
+  } else {
+    color = (vec4<f32>(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<f32>(color, 1);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  var red : vec4<f32> = textureLoadExternal(t, ext_tex_plane_1, clamp(vec2<i32>(10, 10), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(t)) - vec2(1)))), ext_tex_params);
+  textureStore(outImage, clamp(vec2<i32>(0, 0), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(outImage)) - vec2(1)))), red);
+  var green : vec4<f32> = textureLoadExternal(t, ext_tex_plane_1, clamp(vec2<i32>(70, 118), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(t)) - vec2(1)))), ext_tex_params);
+  textureStore(outImage, clamp(vec2<i32>(1, 0), vec2(0), vec2<i32>((vec2<u32>(textureDimensions(outImage)) - vec2(1)))), green);
+  return;
+}