IR: Fix robustness transform on textureLoad of sampled and depth textures
For sampled and depth textures, which contain a 'level' argument, the
robustness transform is supposed to clamp 'coords' using the dimensions
at the clamped level, but it was looking up dimensions at level 0
instead.
Bug: 42250751
Bug: 42251045
Bug: 378541479
Change-Id: I0e7fd6148417b248a9b584ae19818e9027306b63
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/214514
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/core/ir/transform/robustness.cc b/src/tint/lang/core/ir/transform/robustness.cc
index 079a6cd..b9dd071 100644
--- a/src/tint/lang/core/ir/transform/robustness.cc
+++ b/src/tint/lang/core/ir/transform/robustness.cc
@@ -309,7 +309,6 @@
break;
}
case core::BuiltinFn::kTextureLoad: {
- clamp_coords(1u);
uint32_t next_arg = 2u;
if (type::IsTextureArray(texture->Dim())) {
clamp_array_index(next_arg++);
@@ -317,6 +316,7 @@
if (texture->IsAnyOf<type::SampledTexture, type::DepthTexture>()) {
clamp_level(next_arg++);
}
+ clamp_coords(1u); // Must run after clamp_level
break;
}
case core::BuiltinFn::kTextureStore: {
diff --git a/src/tint/lang/core/ir/transform/robustness_test.cc b/src/tint/lang/core/ir/transform/robustness_test.cc
index a9ecbf8..791101e 100644
--- a/src/tint/lang/core/ir/transform/robustness_test.cc
+++ b/src/tint/lang/core/ir/transform/robustness_test.cc
@@ -2146,28 +2146,28 @@
%load_signed = func(%coords:i32, %level:i32):vec4<f32> {
$B2: {
%5:texture_1d<f32> = load %texture
- %6:u32 = textureDimensions %5
+ %6:u32 = textureNumLevels %5
%7:u32 = sub %6, 1u
- %8:u32 = convert %coords
+ %8:u32 = convert %level
%9:u32 = min %8, %7
- %10:u32 = textureNumLevels %5
+ %10:u32 = textureDimensions %5, %9
%11:u32 = sub %10, 1u
- %12:u32 = convert %level
+ %12:u32 = convert %coords
%13:u32 = min %12, %11
- %14:vec4<f32> = textureLoad %5, %9, %13
+ %14:vec4<f32> = textureLoad %5, %13, %9
ret %14
}
}
%load_unsigned = func(%coords_1:u32, %level_1:u32):vec4<f32> { # %coords_1: 'coords', %level_1: 'level'
$B3: {
%18:texture_1d<f32> = load %texture
- %19:u32 = textureDimensions %18
+ %19:u32 = textureNumLevels %18
%20:u32 = sub %19, 1u
- %21:u32 = min %coords_1, %20
- %22:u32 = textureNumLevels %18
+ %21:u32 = min %level_1, %20
+ %22:u32 = textureDimensions %18, %21
%23:u32 = sub %22, 1u
- %24:u32 = min %level_1, %23
- %25:vec4<f32> = textureLoad %18, %21, %24
+ %24:u32 = min %coords_1, %23
+ %25:vec4<f32> = textureLoad %18, %24, %21
ret %25
}
}
@@ -2243,28 +2243,28 @@
%load_signed = func(%coords:vec2<i32>, %level:i32):vec4<f32> {
$B2: {
%5:texture_2d<f32> = load %texture
- %6:vec2<u32> = textureDimensions %5
- %7:vec2<u32> = sub %6, vec2<u32>(1u)
- %8:vec2<u32> = convert %coords
- %9:vec2<u32> = min %8, %7
- %10:u32 = textureNumLevels %5
- %11:u32 = sub %10, 1u
- %12:u32 = convert %level
- %13:u32 = min %12, %11
- %14:vec4<f32> = textureLoad %5, %9, %13
+ %6:u32 = textureNumLevels %5
+ %7:u32 = sub %6, 1u
+ %8:u32 = convert %level
+ %9:u32 = min %8, %7
+ %10:vec2<u32> = textureDimensions %5, %9
+ %11:vec2<u32> = sub %10, vec2<u32>(1u)
+ %12:vec2<u32> = convert %coords
+ %13:vec2<u32> = min %12, %11
+ %14:vec4<f32> = textureLoad %5, %13, %9
ret %14
}
}
%load_unsigned = func(%coords_1:vec2<u32>, %level_1:u32):vec4<f32> { # %coords_1: 'coords', %level_1: 'level'
$B3: {
%18:texture_2d<f32> = load %texture
- %19:vec2<u32> = textureDimensions %18
- %20:vec2<u32> = sub %19, vec2<u32>(1u)
- %21:vec2<u32> = min %coords_1, %20
- %22:u32 = textureNumLevels %18
- %23:u32 = sub %22, 1u
- %24:u32 = min %level_1, %23
- %25:vec4<f32> = textureLoad %18, %21, %24
+ %19:u32 = textureNumLevels %18
+ %20:u32 = sub %19, 1u
+ %21:u32 = min %level_1, %20
+ %22:vec2<u32> = textureDimensions %18, %21
+ %23:vec2<u32> = sub %22, vec2<u32>(1u)
+ %24:vec2<u32> = min %coords_1, %23
+ %25:vec4<f32> = textureLoad %18, %24, %21
ret %25
}
}
@@ -2343,35 +2343,35 @@
%load_signed = func(%coords:vec2<i32>, %layer:i32, %level:i32):vec4<f32> {
$B2: {
%6:texture_2d_array<f32> = load %texture
- %7:vec2<u32> = textureDimensions %6
- %8:vec2<u32> = sub %7, vec2<u32>(1u)
- %9:vec2<u32> = convert %coords
- %10:vec2<u32> = min %9, %8
- %11:u32 = textureNumLayers %6
+ %7:u32 = textureNumLayers %6
+ %8:u32 = sub %7, 1u
+ %9:u32 = convert %layer
+ %10:u32 = min %9, %8
+ %11:u32 = textureNumLevels %6
%12:u32 = sub %11, 1u
- %13:u32 = convert %layer
+ %13:u32 = convert %level
%14:u32 = min %13, %12
- %15:u32 = textureNumLevels %6
- %16:u32 = sub %15, 1u
- %17:u32 = convert %level
- %18:u32 = min %17, %16
- %19:vec4<f32> = textureLoad %6, %10, %14, %18
+ %15:vec2<u32> = textureDimensions %6, %14
+ %16:vec2<u32> = sub %15, vec2<u32>(1u)
+ %17:vec2<u32> = convert %coords
+ %18:vec2<u32> = min %17, %16
+ %19:vec4<f32> = textureLoad %6, %18, %10, %14
ret %19
}
}
%load_unsigned = func(%coords_1:vec2<u32>, %layer_1:u32, %level_1:u32):vec4<f32> { # %coords_1: 'coords', %layer_1: 'layer', %level_1: 'level'
$B3: {
%24:texture_2d_array<f32> = load %texture
- %25:vec2<u32> = textureDimensions %24
- %26:vec2<u32> = sub %25, vec2<u32>(1u)
- %27:vec2<u32> = min %coords_1, %26
- %28:u32 = textureNumLayers %24
+ %25:u32 = textureNumLayers %24
+ %26:u32 = sub %25, 1u
+ %27:u32 = min %layer_1, %26
+ %28:u32 = textureNumLevels %24
%29:u32 = sub %28, 1u
- %30:u32 = min %layer_1, %29
- %31:u32 = textureNumLevels %24
- %32:u32 = sub %31, 1u
- %33:u32 = min %level_1, %32
- %34:vec4<f32> = textureLoad %24, %27, %30, %33
+ %30:u32 = min %level_1, %29
+ %31:vec2<u32> = textureDimensions %24, %30
+ %32:vec2<u32> = sub %31, vec2<u32>(1u)
+ %33:vec2<u32> = min %coords_1, %32
+ %34:vec4<f32> = textureLoad %24, %33, %27, %30
ret %34
}
}
@@ -2447,28 +2447,28 @@
%load_signed = func(%coords:vec3<i32>, %level:i32):vec4<f32> {
$B2: {
%5:texture_3d<f32> = load %texture
- %6:vec3<u32> = textureDimensions %5
- %7:vec3<u32> = sub %6, vec3<u32>(1u)
- %8:vec3<u32> = convert %coords
- %9:vec3<u32> = min %8, %7
- %10:u32 = textureNumLevels %5
- %11:u32 = sub %10, 1u
- %12:u32 = convert %level
- %13:u32 = min %12, %11
- %14:vec4<f32> = textureLoad %5, %9, %13
+ %6:u32 = textureNumLevels %5
+ %7:u32 = sub %6, 1u
+ %8:u32 = convert %level
+ %9:u32 = min %8, %7
+ %10:vec3<u32> = textureDimensions %5, %9
+ %11:vec3<u32> = sub %10, vec3<u32>(1u)
+ %12:vec3<u32> = convert %coords
+ %13:vec3<u32> = min %12, %11
+ %14:vec4<f32> = textureLoad %5, %13, %9
ret %14
}
}
%load_unsigned = func(%coords_1:vec3<u32>, %level_1:u32):vec4<f32> { # %coords_1: 'coords', %level_1: 'level'
$B3: {
%18:texture_3d<f32> = load %texture
- %19:vec3<u32> = textureDimensions %18
- %20:vec3<u32> = sub %19, vec3<u32>(1u)
- %21:vec3<u32> = min %coords_1, %20
- %22:u32 = textureNumLevels %18
- %23:u32 = sub %22, 1u
- %24:u32 = min %level_1, %23
- %25:vec4<f32> = textureLoad %18, %21, %24
+ %19:u32 = textureNumLevels %18
+ %20:u32 = sub %19, 1u
+ %21:u32 = min %level_1, %20
+ %22:vec3<u32> = textureDimensions %18, %21
+ %23:vec3<u32> = sub %22, vec3<u32>(1u)
+ %24:vec3<u32> = min %coords_1, %23
+ %25:vec4<f32> = textureLoad %18, %24, %21
ret %25
}
}
@@ -2632,28 +2632,28 @@
%load_signed = func(%coords:vec2<i32>, %level:i32):f32 {
$B2: {
%5:texture_depth_2d = load %texture
- %6:vec2<u32> = textureDimensions %5
- %7:vec2<u32> = sub %6, vec2<u32>(1u)
- %8:vec2<u32> = convert %coords
- %9:vec2<u32> = min %8, %7
- %10:u32 = textureNumLevels %5
- %11:u32 = sub %10, 1u
- %12:u32 = convert %level
- %13:u32 = min %12, %11
- %14:f32 = textureLoad %5, %9, %13
+ %6:u32 = textureNumLevels %5
+ %7:u32 = sub %6, 1u
+ %8:u32 = convert %level
+ %9:u32 = min %8, %7
+ %10:vec2<u32> = textureDimensions %5, %9
+ %11:vec2<u32> = sub %10, vec2<u32>(1u)
+ %12:vec2<u32> = convert %coords
+ %13:vec2<u32> = min %12, %11
+ %14:f32 = textureLoad %5, %13, %9
ret %14
}
}
%load_unsigned = func(%coords_1:vec2<u32>, %level_1:u32):f32 { # %coords_1: 'coords', %level_1: 'level'
$B3: {
%18:texture_depth_2d = load %texture
- %19:vec2<u32> = textureDimensions %18
- %20:vec2<u32> = sub %19, vec2<u32>(1u)
- %21:vec2<u32> = min %coords_1, %20
- %22:u32 = textureNumLevels %18
- %23:u32 = sub %22, 1u
- %24:u32 = min %level_1, %23
- %25:f32 = textureLoad %18, %21, %24
+ %19:u32 = textureNumLevels %18
+ %20:u32 = sub %19, 1u
+ %21:u32 = min %level_1, %20
+ %22:vec2<u32> = textureDimensions %18, %21
+ %23:vec2<u32> = sub %22, vec2<u32>(1u)
+ %24:vec2<u32> = min %coords_1, %23
+ %25:f32 = textureLoad %18, %24, %21
ret %25
}
}
@@ -2731,35 +2731,35 @@
%load_signed = func(%coords:vec2<i32>, %layer:i32, %level:i32):f32 {
$B2: {
%6:texture_depth_2d_array = load %texture
- %7:vec2<u32> = textureDimensions %6
- %8:vec2<u32> = sub %7, vec2<u32>(1u)
- %9:vec2<u32> = convert %coords
- %10:vec2<u32> = min %9, %8
- %11:u32 = textureNumLayers %6
+ %7:u32 = textureNumLayers %6
+ %8:u32 = sub %7, 1u
+ %9:u32 = convert %layer
+ %10:u32 = min %9, %8
+ %11:u32 = textureNumLevels %6
%12:u32 = sub %11, 1u
- %13:u32 = convert %layer
+ %13:u32 = convert %level
%14:u32 = min %13, %12
- %15:u32 = textureNumLevels %6
- %16:u32 = sub %15, 1u
- %17:u32 = convert %level
- %18:u32 = min %17, %16
- %19:f32 = textureLoad %6, %10, %14, %18
+ %15:vec2<u32> = textureDimensions %6, %14
+ %16:vec2<u32> = sub %15, vec2<u32>(1u)
+ %17:vec2<u32> = convert %coords
+ %18:vec2<u32> = min %17, %16
+ %19:f32 = textureLoad %6, %18, %10, %14
ret %19
}
}
%load_unsigned = func(%coords_1:vec2<u32>, %layer_1:u32, %level_1:u32):f32 { # %coords_1: 'coords', %layer_1: 'layer', %level_1: 'level'
$B3: {
%24:texture_depth_2d_array = load %texture
- %25:vec2<u32> = textureDimensions %24
- %26:vec2<u32> = sub %25, vec2<u32>(1u)
- %27:vec2<u32> = min %coords_1, %26
- %28:u32 = textureNumLayers %24
+ %25:u32 = textureNumLayers %24
+ %26:u32 = sub %25, 1u
+ %27:u32 = min %layer_1, %26
+ %28:u32 = textureNumLevels %24
%29:u32 = sub %28, 1u
- %30:u32 = min %layer_1, %29
- %31:u32 = textureNumLevels %24
- %32:u32 = sub %31, 1u
- %33:u32 = min %level_1, %32
- %34:f32 = textureLoad %24, %27, %30, %33
+ %30:u32 = min %level_1, %29
+ %31:vec2<u32> = textureDimensions %24, %30
+ %32:vec2<u32> = sub %31, vec2<u32>(1u)
+ %33:vec2<u32> = min %coords_1, %32
+ %34:f32 = textureLoad %24, %33, %27, %30
ret %34
}
}
@@ -3191,28 +3191,28 @@
%load_signed = func(%coords:vec2<i32>, %layer:i32):vec4<f32> {
$B2: {
%5:texture_storage_2d_array<rgba8unorm, read_write> = load %texture
- %6:vec2<u32> = textureDimensions %5
- %7:vec2<u32> = sub %6, vec2<u32>(1u)
- %8:vec2<u32> = convert %coords
- %9:vec2<u32> = min %8, %7
- %10:u32 = textureNumLayers %5
- %11:u32 = sub %10, 1u
- %12:u32 = convert %layer
- %13:u32 = min %12, %11
- %14:vec4<f32> = textureLoad %5, %9, %13
+ %6:u32 = textureNumLayers %5
+ %7:u32 = sub %6, 1u
+ %8:u32 = convert %layer
+ %9:u32 = min %8, %7
+ %10:vec2<u32> = textureDimensions %5
+ %11:vec2<u32> = sub %10, vec2<u32>(1u)
+ %12:vec2<u32> = convert %coords
+ %13:vec2<u32> = min %12, %11
+ %14:vec4<f32> = textureLoad %5, %13, %9
ret %14
}
}
%load_unsigned = func(%coords_1:vec2<u32>, %layer_1:u32):vec4<f32> { # %coords_1: 'coords', %layer_1: 'layer'
$B3: {
%18:texture_storage_2d_array<rgba8unorm, read_write> = load %texture
- %19:vec2<u32> = textureDimensions %18
- %20:vec2<u32> = sub %19, vec2<u32>(1u)
- %21:vec2<u32> = min %coords_1, %20
- %22:u32 = textureNumLayers %18
- %23:u32 = sub %22, 1u
- %24:u32 = min %layer_1, %23
- %25:vec4<f32> = textureLoad %18, %21, %24
+ %19:u32 = textureNumLayers %18
+ %20:u32 = sub %19, 1u
+ %21:u32 = min %layer_1, %20
+ %22:vec2<u32> = textureDimensions %18
+ %23:vec2<u32> = sub %22, vec2<u32>(1u)
+ %24:vec2<u32> = min %coords_1, %23
+ %25:vec4<f32> = textureLoad %18, %24, %21
ret %25
}
}
diff --git a/src/tint/lang/spirv/writer/texture_builtin_test.cc b/src/tint/lang/spirv/writer/texture_builtin_test.cc
index 432dc0d..0a4da78 100644
--- a/src/tint/lang/spirv/writer/texture_builtin_test.cc
+++ b/src/tint/lang/spirv/writer/texture_builtin_test.cc
@@ -2006,14 +2006,14 @@
ASSERT_TRUE(Generate()) << Error() << output_;
EXPECT_INST(R"(
- %13 = OpImageQuerySizeLod %v2uint %texture %uint_0
- %15 = OpISub %v2uint %13 %16
- %18 = OpExtInst %v2uint %19 UMin %coords %15
- %20 = OpImageQueryLevels %uint %texture
- %21 = OpISub %uint %20 %uint_1
- %22 = OpBitcast %uint %level
- %23 = OpExtInst %uint %19 UMin %22 %21
- %result = OpImageFetch %v4float %texture %18 Lod %23
+ %13 = OpImageQueryLevels %uint %texture
+ %14 = OpISub %uint %13 %uint_1
+ %16 = OpBitcast %uint %level
+ %17 = OpExtInst %uint %18 UMin %16 %14
+ %19 = OpImageQuerySizeLod %v2uint %texture %17
+ %20 = OpISub %v2uint %19 %21
+ %22 = OpExtInst %v2uint %18 UMin %coords %20
+ %result = OpImageFetch %v4float %texture %22 Lod %17
)");
}
diff --git a/test/tint/bug/chromium/378541479.wgsl b/test/tint/bug/chromium/378541479.wgsl
new file mode 100644
index 0000000..8badf52
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl
@@ -0,0 +1,10 @@
+// flags: --transform robustness
+
+@group(0) @binding(0) var<uniform> level : u32;
+@group(0) @binding(1) var<uniform> coords : vec2<u32>;
+@group(0) @binding(2) var tex: texture_depth_2d;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ var res: f32 = textureLoad(tex, coords, level);
+}
diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.dxc.hlsl b/test/tint/bug/chromium/378541479.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..6ceba9a
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.dxc.hlsl
@@ -0,0 +1,13 @@
+cbuffer cbuffer_level : register(b0) {
+ uint4 level[1];
+};
+cbuffer cbuffer_coords : register(b1) {
+ uint4 coords[1];
+};
+Texture2D tex : register(t2);
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ float res = tex.Load(uint3(coords[0].xy, level[0].x)).x;
+ return;
+}
diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.fxc.hlsl b/test/tint/bug/chromium/378541479.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..6ceba9a
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.fxc.hlsl
@@ -0,0 +1,13 @@
+cbuffer cbuffer_level : register(b0) {
+ uint4 level[1];
+};
+cbuffer cbuffer_coords : register(b1) {
+ uint4 coords[1];
+};
+Texture2D tex : register(t2);
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+ float res = tex.Load(uint3(coords[0].xy, level[0].x)).x;
+ return;
+}
diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.glsl b/test/tint/bug/chromium/378541479.wgsl.expected.glsl
new file mode 100644
index 0000000..bb0ee27
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.glsl
@@ -0,0 +1,27 @@
+#version 310 es
+
+
+struct TintTextureUniformData {
+ uint tint_builtin_value_0;
+};
+
+layout(binding = 0, std140)
+uniform level_block_1_ubo {
+ uint inner;
+} v;
+layout(binding = 1, std140)
+uniform coords_block_1_ubo {
+ uvec2 inner;
+} v_1;
+layout(binding = 0, std140)
+uniform tint_symbol_1_ubo {
+ TintTextureUniformData inner;
+} v_2;
+uniform highp sampler2D tex;
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ uvec2 v_3 = v_1.inner;
+ uint v_4 = min(v.inner, (v_2.inner.tint_builtin_value_0 - 1u));
+ ivec2 v_5 = ivec2(min(v_3, (uvec2(textureSize(tex, int(v_4))) - uvec2(1u))));
+ float res = texelFetch(tex, v_5, int(v_4)).x;
+}
diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/chromium/378541479.wgsl.expected.ir.dxc.hlsl
new file mode 100644
index 0000000..72b8b84
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.ir.dxc.hlsl
@@ -0,0 +1,21 @@
+
+cbuffer cbuffer_level : register(b0) {
+ uint4 level[1];
+};
+cbuffer cbuffer_coords : register(b1) {
+ uint4 coords[1];
+};
+Texture2D tex : register(t2);
+[numthreads(1, 1, 1)]
+void compute_main() {
+ uint2 v = coords[0u].xy;
+ uint v_1 = level[0u].x;
+ uint3 v_2 = (0u).xxx;
+ tex.GetDimensions(0u, v_2[0u], v_2[1u], v_2[2u]);
+ uint v_3 = min(v_1, (v_2.z - 1u));
+ uint3 v_4 = (0u).xxx;
+ tex.GetDimensions(uint(v_3), v_4[0u], v_4[1u], v_4[2u]);
+ int2 v_5 = int2(min(v, (v_4.xy - (1u).xx)));
+ float res = tex.Load(int3(v_5, int(v_3))).x;
+}
+
diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.ir.fxc.hlsl b/test/tint/bug/chromium/378541479.wgsl.expected.ir.fxc.hlsl
new file mode 100644
index 0000000..72b8b84
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.ir.fxc.hlsl
@@ -0,0 +1,21 @@
+
+cbuffer cbuffer_level : register(b0) {
+ uint4 level[1];
+};
+cbuffer cbuffer_coords : register(b1) {
+ uint4 coords[1];
+};
+Texture2D tex : register(t2);
+[numthreads(1, 1, 1)]
+void compute_main() {
+ uint2 v = coords[0u].xy;
+ uint v_1 = level[0u].x;
+ uint3 v_2 = (0u).xxx;
+ tex.GetDimensions(0u, v_2[0u], v_2[1u], v_2[2u]);
+ uint v_3 = min(v_1, (v_2.z - 1u));
+ uint3 v_4 = (0u).xxx;
+ tex.GetDimensions(uint(v_3), v_4[0u], v_4[1u], v_4[2u]);
+ int2 v_5 = int2(min(v, (v_4.xy - (1u).xx)));
+ float res = tex.Load(int3(v_5, int(v_3))).x;
+}
+
diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.ir.msl b/test/tint/bug/chromium/378541479.wgsl.expected.ir.msl
new file mode 100644
index 0000000..b1edc57
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.ir.msl
@@ -0,0 +1,17 @@
+#include <metal_stdlib>
+using namespace metal;
+
+struct tint_module_vars_struct {
+ const constant uint* level;
+ const constant uint2* coords;
+ depth2d<float, access::sample> tex;
+};
+
+kernel void compute_main(const constant uint* level [[buffer(1)]], const constant uint2* coords [[buffer(0)]], depth2d<float, access::sample> tex [[texture(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.level=level, .coords=coords, .tex=tex};
+ uint2 const v = (*tint_module_vars.coords);
+ uint const v_1 = (*tint_module_vars.level);
+ uint const v_2 = min(v_1, (tint_module_vars.tex.get_num_mip_levels() - 1u));
+ uint const v_3 = tint_module_vars.tex.get_width(v_2);
+ float res = tint_module_vars.tex.read(min(v, (uint2(v_3, tint_module_vars.tex.get_height(v_2)) - uint2(1u))), v_2);
+}
diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.msl b/test/tint/bug/chromium/378541479.wgsl.expected.msl
new file mode 100644
index 0000000..0ac2d2a
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void compute_main(const constant uint* tint_symbol [[buffer(1)]], depth2d<float, access::sample> tint_symbol_1 [[texture(0)]], const constant uint2* tint_symbol_2 [[buffer(0)]]) {
+ uint const level_idx = min(uint(*(tint_symbol)), (tint_symbol_1.get_num_mip_levels() - 1u));
+ float res = tint_symbol_1.read(uint2(min(*(tint_symbol_2), (uint2(tint_symbol_1.get_width(level_idx), tint_symbol_1.get_height(level_idx)) - uint2(1u)))), level_idx);
+ return;
+}
+
diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.spvasm b/test/tint/bug/chromium/378541479.wgsl.expected.spvasm
new file mode 100644
index 0000000..b8a3b11
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.spvasm
@@ -0,0 +1,70 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 1
+; Bound: 39
+; Schema: 0
+ OpCapability Shader
+ OpCapability ImageQuery
+ %29 = OpExtInstImport "GLSL.std.450"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %compute_main "compute_main"
+ OpExecutionMode %compute_main LocalSize 1 1 1
+ OpMemberName %level_block 0 "inner"
+ OpName %level_block "level_block"
+ OpMemberName %coords_block 0 "inner"
+ OpName %coords_block "coords_block"
+ OpName %tex "tex"
+ OpName %compute_main "compute_main"
+ OpName %res "res"
+ OpMemberDecorate %level_block 0 Offset 0
+ OpDecorate %level_block Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ OpDecorate %1 NonWritable
+ OpMemberDecorate %coords_block 0 Offset 0
+ OpDecorate %coords_block Block
+ OpDecorate %5 DescriptorSet 0
+ OpDecorate %5 Binding 1
+ OpDecorate %5 NonWritable
+ OpDecorate %tex DescriptorSet 0
+ OpDecorate %tex Binding 2
+ %uint = OpTypeInt 32 0
+%level_block = OpTypeStruct %uint
+%_ptr_Uniform_level_block = OpTypePointer Uniform %level_block
+ %1 = OpVariable %_ptr_Uniform_level_block Uniform
+ %v2uint = OpTypeVector %uint 2
+%coords_block = OpTypeStruct %v2uint
+%_ptr_Uniform_coords_block = OpTypePointer Uniform %coords_block
+ %5 = OpVariable %_ptr_Uniform_coords_block Uniform
+ %float = OpTypeFloat 32
+ %11 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_11 = OpTypePointer UniformConstant %11
+ %tex = OpVariable %_ptr_UniformConstant_11 UniformConstant
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void
+%_ptr_Uniform_v2uint = OpTypePointer Uniform %v2uint
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+ %uint_1 = OpConstant %uint 1
+ %32 = OpConstantComposite %v2uint %uint_1 %uint_1
+ %v4float = OpTypeVector %float 4
+%_ptr_Function_float = OpTypePointer Function %float
+%compute_main = OpFunction %void None %15
+ %16 = OpLabel
+ %res = OpVariable %_ptr_Function_float Function
+ %17 = OpLoad %11 %tex None
+ %18 = OpAccessChain %_ptr_Uniform_v2uint %5 %uint_0
+ %21 = OpLoad %v2uint %18 None
+ %22 = OpAccessChain %_ptr_Uniform_uint %1 %uint_0
+ %24 = OpLoad %uint %22 None
+ %25 = OpImageQueryLevels %uint %17
+ %26 = OpISub %uint %25 %uint_1
+ %28 = OpExtInst %uint %29 UMin %24 %26
+ %30 = OpImageQuerySizeLod %v2uint %17 %28
+ %31 = OpISub %v2uint %30 %32
+ %33 = OpExtInst %v2uint %29 UMin %21 %31
+ %34 = OpImageFetch %v4float %17 %33 Lod %28
+ %36 = OpCompositeExtract %float %34 0
+ OpStore %res %36
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/bug/chromium/378541479.wgsl.expected.wgsl b/test/tint/bug/chromium/378541479.wgsl.expected.wgsl
new file mode 100644
index 0000000..705ce78
--- /dev/null
+++ b/test/tint/bug/chromium/378541479.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+@group(0) @binding(0) var<uniform> level : u32;
+
+@group(0) @binding(1) var<uniform> coords : vec2<u32>;
+
+@group(0) @binding(2) var tex : texture_depth_2d;
+
+@compute @workgroup_size(1)
+fn compute_main() {
+ var res : f32 = textureLoad(tex, coords, level);
+}
diff --git a/test/tint/bug/tint/1739.wgsl.expected.ir.dxc.hlsl b/test/tint/bug/tint/1739.wgsl.expected.ir.dxc.hlsl
index e6e322dc..32a4cf0 100644
--- a/test/tint/bug/tint/1739.wgsl.expected.ir.dxc.hlsl
+++ b/test/tint/bug/tint/1739.wgsl.expected.ir.dxc.hlsl
@@ -51,112 +51,109 @@
float3 v_7 = (0.0f).xxx;
float v_8 = 0.0f;
if ((params.numPlanes == 1u)) {
- uint2 v_9 = (0u).xx;
- plane_0.GetDimensions(v_9[0u], v_9[1u]);
- uint2 v_10 = min(v_6, (v_9 - (1u).xx));
+ uint3 v_9 = (0u).xxx;
+ plane_0.GetDimensions(0u, v_9[0u], v_9[1u], v_9[2u]);
+ uint v_10 = min(0u, (v_9.z - 1u));
uint3 v_11 = (0u).xxx;
- plane_0.GetDimensions(0u, v_11[0u], v_11[1u], v_11[2u]);
- uint v_12 = min(0u, (v_11.z - 1u));
- int2 v_13 = int2(v_10);
- float4 v_14 = float4(plane_0.Load(int3(v_13, int(v_12))));
- v_7 = v_14.xyz;
- v_8 = v_14[3u];
+ plane_0.GetDimensions(uint(v_10), v_11[0u], v_11[1u], v_11[2u]);
+ int2 v_12 = int2(min(v_6, (v_11.xy - (1u).xx)));
+ float4 v_13 = float4(plane_0.Load(int3(v_12, int(v_10))));
+ v_7 = v_13.xyz;
+ v_8 = v_13[3u];
} else {
- uint2 v_15 = (0u).xx;
- plane_0.GetDimensions(v_15[0u], v_15[1u]);
- uint2 v_16 = min(v_6, (v_15 - (1u).xx));
- uint3 v_17 = (0u).xxx;
- plane_0.GetDimensions(0u, v_17[0u], v_17[1u], v_17[2u]);
- uint v_18 = min(0u, (v_17.z - 1u));
- int2 v_19 = int2(v_16);
- float v_20 = float4(plane_0.Load(int3(v_19, int(v_18))))[0u];
- uint2 v_21 = tint_v2f32_to_v2u32((v_5 * params.plane1CoordFactor));
- uint2 v_22 = (0u).xx;
- plane_1.GetDimensions(v_22[0u], v_22[1u]);
- uint2 v_23 = min(v_21, (v_22 - (1u).xx));
- uint3 v_24 = (0u).xxx;
- plane_1.GetDimensions(0u, v_24[0u], v_24[1u], v_24[2u]);
- uint v_25 = min(0u, (v_24.z - 1u));
- int2 v_26 = int2(v_23);
- v_7 = mul(params.yuvToRgbConversionMatrix, float4(v_20, float4(plane_1.Load(int3(v_26, int(v_25)))).xy, 1.0f));
+ uint3 v_14 = (0u).xxx;
+ plane_0.GetDimensions(0u, v_14[0u], v_14[1u], v_14[2u]);
+ uint v_15 = min(0u, (v_14.z - 1u));
+ uint3 v_16 = (0u).xxx;
+ plane_0.GetDimensions(uint(v_15), v_16[0u], v_16[1u], v_16[2u]);
+ int2 v_17 = int2(min(v_6, (v_16.xy - (1u).xx)));
+ float v_18 = float4(plane_0.Load(int3(v_17, int(v_15))))[0u];
+ uint2 v_19 = tint_v2f32_to_v2u32((v_5 * params.plane1CoordFactor));
+ uint3 v_20 = (0u).xxx;
+ plane_1.GetDimensions(0u, v_20[0u], v_20[1u], v_20[2u]);
+ uint v_21 = min(0u, (v_20.z - 1u));
+ uint3 v_22 = (0u).xxx;
+ plane_1.GetDimensions(uint(v_21), v_22[0u], v_22[1u], v_22[2u]);
+ int2 v_23 = int2(min(v_19, (v_22.xy - (1u).xx)));
+ v_7 = mul(params.yuvToRgbConversionMatrix, float4(v_18, float4(plane_1.Load(int3(v_23, int(v_21)))).xy, 1.0f));
v_8 = 1.0f;
}
- float3 v_27 = v_7;
- float3 v_28 = (0.0f).xxx;
+ float3 v_24 = v_7;
+ float3 v_25 = (0.0f).xxx;
if ((params.doYuvToRgbConversionOnly == 0u)) {
- tint_GammaTransferParams v_29 = params.gammaDecodeParams;
- tint_GammaTransferParams v_30 = params.gammaEncodeParams;
- v_28 = tint_GammaCorrection(mul(tint_GammaCorrection(v_27, v_29), params.gamutConversionMatrix), v_30);
+ tint_GammaTransferParams v_26 = params.gammaDecodeParams;
+ tint_GammaTransferParams v_27 = params.gammaEncodeParams;
+ v_25 = tint_GammaCorrection(mul(tint_GammaCorrection(v_24, v_26), params.gamutConversionMatrix), v_27);
} else {
- v_28 = v_27;
+ v_25 = v_24;
}
- return float4(v_28, v_8);
+ return float4(v_25, v_8);
}
-float3x2 v_31(uint start_byte_offset) {
- uint4 v_32 = t_params[(start_byte_offset / 16u)];
- float2 v_33 = asfloat((((((start_byte_offset % 16u) / 4u) == 2u)) ? (v_32.zw) : (v_32.xy)));
- uint4 v_34 = t_params[((8u + start_byte_offset) / 16u)];
- float2 v_35 = asfloat(((((((8u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_34.zw) : (v_34.xy)));
- uint4 v_36 = t_params[((16u + start_byte_offset) / 16u)];
- return float3x2(v_33, v_35, asfloat(((((((16u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_36.zw) : (v_36.xy))));
+float3x2 v_28(uint start_byte_offset) {
+ uint4 v_29 = t_params[(start_byte_offset / 16u)];
+ float2 v_30 = asfloat((((((start_byte_offset % 16u) / 4u) == 2u)) ? (v_29.zw) : (v_29.xy)));
+ uint4 v_31 = t_params[((8u + start_byte_offset) / 16u)];
+ float2 v_32 = asfloat(((((((8u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_31.zw) : (v_31.xy)));
+ uint4 v_33 = t_params[((16u + start_byte_offset) / 16u)];
+ return float3x2(v_30, v_32, asfloat(((((((16u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_33.zw) : (v_33.xy))));
}
-float3x3 v_37(uint start_byte_offset) {
- float3 v_38 = asfloat(t_params[(start_byte_offset / 16u)].xyz);
- float3 v_39 = asfloat(t_params[((16u + start_byte_offset) / 16u)].xyz);
- return float3x3(v_38, v_39, asfloat(t_params[((32u + start_byte_offset) / 16u)].xyz));
+float3x3 v_34(uint start_byte_offset) {
+ float3 v_35 = asfloat(t_params[(start_byte_offset / 16u)].xyz);
+ float3 v_36 = asfloat(t_params[((16u + start_byte_offset) / 16u)].xyz);
+ return float3x3(v_35, v_36, asfloat(t_params[((32u + start_byte_offset) / 16u)].xyz));
}
-tint_GammaTransferParams v_40(uint start_byte_offset) {
- float v_41 = asfloat(t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)]);
- float v_42 = asfloat(t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)]);
- float v_43 = asfloat(t_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) % 16u) / 4u)]);
- float v_44 = asfloat(t_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) % 16u) / 4u)]);
- float v_45 = asfloat(t_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) % 16u) / 4u)]);
- float v_46 = asfloat(t_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) % 16u) / 4u)]);
- float v_47 = asfloat(t_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) % 16u) / 4u)]);
- tint_GammaTransferParams v_48 = {v_41, v_42, v_43, v_44, v_45, v_46, v_47, t_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) % 16u) / 4u)]};
- return v_48;
+tint_GammaTransferParams v_37(uint start_byte_offset) {
+ float v_38 = asfloat(t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)]);
+ float v_39 = asfloat(t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)]);
+ float v_40 = asfloat(t_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) % 16u) / 4u)]);
+ float v_41 = asfloat(t_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) % 16u) / 4u)]);
+ float v_42 = asfloat(t_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) % 16u) / 4u)]);
+ float v_43 = asfloat(t_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) % 16u) / 4u)]);
+ float v_44 = asfloat(t_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) % 16u) / 4u)]);
+ tint_GammaTransferParams v_45 = {v_38, v_39, v_40, v_41, v_42, v_43, v_44, t_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) % 16u) / 4u)]};
+ return v_45;
}
-float3x4 v_49(uint start_byte_offset) {
- float4 v_50 = asfloat(t_params[(start_byte_offset / 16u)]);
- float4 v_51 = asfloat(t_params[((16u + start_byte_offset) / 16u)]);
- return float3x4(v_50, v_51, asfloat(t_params[((32u + start_byte_offset) / 16u)]));
+float3x4 v_46(uint start_byte_offset) {
+ float4 v_47 = asfloat(t_params[(start_byte_offset / 16u)]);
+ float4 v_48 = asfloat(t_params[((16u + start_byte_offset) / 16u)]);
+ return float3x4(v_47, v_48, asfloat(t_params[((32u + start_byte_offset) / 16u)]));
}
-tint_ExternalTextureParams v_52(uint start_byte_offset) {
- uint v_53 = t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)];
- uint v_54 = t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)];
- float3x4 v_55 = v_49((16u + start_byte_offset));
- tint_GammaTransferParams v_56 = v_40((64u + start_byte_offset));
- tint_GammaTransferParams v_57 = v_40((96u + start_byte_offset));
- float3x3 v_58 = v_37((128u + start_byte_offset));
- float3x2 v_59 = v_31((176u + start_byte_offset));
- float3x2 v_60 = v_31((200u + start_byte_offset));
- uint4 v_61 = t_params[((224u + start_byte_offset) / 16u)];
- float2 v_62 = asfloat(((((((224u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_61.zw) : (v_61.xy)));
- uint4 v_63 = t_params[((232u + start_byte_offset) / 16u)];
- float2 v_64 = asfloat(((((((232u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_63.zw) : (v_63.xy)));
- uint4 v_65 = t_params[((240u + start_byte_offset) / 16u)];
- float2 v_66 = asfloat(((((((240u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_65.zw) : (v_65.xy)));
- uint4 v_67 = t_params[((248u + start_byte_offset) / 16u)];
- float2 v_68 = asfloat(((((((248u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_67.zw) : (v_67.xy)));
- uint4 v_69 = t_params[((256u + start_byte_offset) / 16u)];
- uint2 v_70 = ((((((256u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_69.zw) : (v_69.xy));
- uint4 v_71 = t_params[((264u + start_byte_offset) / 16u)];
- tint_ExternalTextureParams v_72 = {v_53, v_54, v_55, v_56, v_57, v_58, v_59, v_60, v_62, v_64, v_66, v_68, v_70, asfloat(((((((264u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_71.zw) : (v_71.xy)))};
- return v_72;
+tint_ExternalTextureParams v_49(uint start_byte_offset) {
+ uint v_50 = t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)];
+ uint v_51 = t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)];
+ float3x4 v_52 = v_46((16u + start_byte_offset));
+ tint_GammaTransferParams v_53 = v_37((64u + start_byte_offset));
+ tint_GammaTransferParams v_54 = v_37((96u + start_byte_offset));
+ float3x3 v_55 = v_34((128u + start_byte_offset));
+ float3x2 v_56 = v_28((176u + start_byte_offset));
+ float3x2 v_57 = v_28((200u + start_byte_offset));
+ uint4 v_58 = t_params[((224u + start_byte_offset) / 16u)];
+ float2 v_59 = asfloat(((((((224u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_58.zw) : (v_58.xy)));
+ uint4 v_60 = t_params[((232u + start_byte_offset) / 16u)];
+ float2 v_61 = asfloat(((((((232u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_60.zw) : (v_60.xy)));
+ uint4 v_62 = t_params[((240u + start_byte_offset) / 16u)];
+ float2 v_63 = asfloat(((((((240u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_62.zw) : (v_62.xy)));
+ uint4 v_64 = t_params[((248u + start_byte_offset) / 16u)];
+ float2 v_65 = asfloat(((((((248u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_64.zw) : (v_64.xy)));
+ uint4 v_66 = t_params[((256u + start_byte_offset) / 16u)];
+ uint2 v_67 = ((((((256u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_66.zw) : (v_66.xy));
+ uint4 v_68 = t_params[((264u + start_byte_offset) / 16u)];
+ tint_ExternalTextureParams v_69 = {v_50, v_51, v_52, v_53, v_54, v_55, v_56, v_57, v_59, v_61, v_63, v_65, v_67, asfloat(((((((264u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_68.zw) : (v_68.xy)))};
+ return v_69;
}
[numthreads(1, 1, 1)]
void main() {
- tint_ExternalTextureParams v_73 = v_52(0u);
- float4 red = tint_TextureLoadExternal(t_plane0, t_plane1, v_73, uint2((int(10)).xx));
+ tint_ExternalTextureParams v_70 = v_49(0u);
+ float4 red = tint_TextureLoadExternal(t_plane0, t_plane1, v_70, uint2((int(10)).xx));
outImage[(int(0)).xx] = red;
- tint_ExternalTextureParams v_74 = v_52(0u);
- float4 green = tint_TextureLoadExternal(t_plane0, t_plane1, v_74, uint2(int2(int(70), int(118))));
+ tint_ExternalTextureParams v_71 = v_49(0u);
+ float4 green = tint_TextureLoadExternal(t_plane0, t_plane1, v_71, uint2(int2(int(70), int(118))));
outImage[int2(int(1), int(0))] = green;
}
diff --git a/test/tint/bug/tint/1739.wgsl.expected.ir.fxc.hlsl b/test/tint/bug/tint/1739.wgsl.expected.ir.fxc.hlsl
index e6e322dc..32a4cf0 100644
--- a/test/tint/bug/tint/1739.wgsl.expected.ir.fxc.hlsl
+++ b/test/tint/bug/tint/1739.wgsl.expected.ir.fxc.hlsl
@@ -51,112 +51,109 @@
float3 v_7 = (0.0f).xxx;
float v_8 = 0.0f;
if ((params.numPlanes == 1u)) {
- uint2 v_9 = (0u).xx;
- plane_0.GetDimensions(v_9[0u], v_9[1u]);
- uint2 v_10 = min(v_6, (v_9 - (1u).xx));
+ uint3 v_9 = (0u).xxx;
+ plane_0.GetDimensions(0u, v_9[0u], v_9[1u], v_9[2u]);
+ uint v_10 = min(0u, (v_9.z - 1u));
uint3 v_11 = (0u).xxx;
- plane_0.GetDimensions(0u, v_11[0u], v_11[1u], v_11[2u]);
- uint v_12 = min(0u, (v_11.z - 1u));
- int2 v_13 = int2(v_10);
- float4 v_14 = float4(plane_0.Load(int3(v_13, int(v_12))));
- v_7 = v_14.xyz;
- v_8 = v_14[3u];
+ plane_0.GetDimensions(uint(v_10), v_11[0u], v_11[1u], v_11[2u]);
+ int2 v_12 = int2(min(v_6, (v_11.xy - (1u).xx)));
+ float4 v_13 = float4(plane_0.Load(int3(v_12, int(v_10))));
+ v_7 = v_13.xyz;
+ v_8 = v_13[3u];
} else {
- uint2 v_15 = (0u).xx;
- plane_0.GetDimensions(v_15[0u], v_15[1u]);
- uint2 v_16 = min(v_6, (v_15 - (1u).xx));
- uint3 v_17 = (0u).xxx;
- plane_0.GetDimensions(0u, v_17[0u], v_17[1u], v_17[2u]);
- uint v_18 = min(0u, (v_17.z - 1u));
- int2 v_19 = int2(v_16);
- float v_20 = float4(plane_0.Load(int3(v_19, int(v_18))))[0u];
- uint2 v_21 = tint_v2f32_to_v2u32((v_5 * params.plane1CoordFactor));
- uint2 v_22 = (0u).xx;
- plane_1.GetDimensions(v_22[0u], v_22[1u]);
- uint2 v_23 = min(v_21, (v_22 - (1u).xx));
- uint3 v_24 = (0u).xxx;
- plane_1.GetDimensions(0u, v_24[0u], v_24[1u], v_24[2u]);
- uint v_25 = min(0u, (v_24.z - 1u));
- int2 v_26 = int2(v_23);
- v_7 = mul(params.yuvToRgbConversionMatrix, float4(v_20, float4(plane_1.Load(int3(v_26, int(v_25)))).xy, 1.0f));
+ uint3 v_14 = (0u).xxx;
+ plane_0.GetDimensions(0u, v_14[0u], v_14[1u], v_14[2u]);
+ uint v_15 = min(0u, (v_14.z - 1u));
+ uint3 v_16 = (0u).xxx;
+ plane_0.GetDimensions(uint(v_15), v_16[0u], v_16[1u], v_16[2u]);
+ int2 v_17 = int2(min(v_6, (v_16.xy - (1u).xx)));
+ float v_18 = float4(plane_0.Load(int3(v_17, int(v_15))))[0u];
+ uint2 v_19 = tint_v2f32_to_v2u32((v_5 * params.plane1CoordFactor));
+ uint3 v_20 = (0u).xxx;
+ plane_1.GetDimensions(0u, v_20[0u], v_20[1u], v_20[2u]);
+ uint v_21 = min(0u, (v_20.z - 1u));
+ uint3 v_22 = (0u).xxx;
+ plane_1.GetDimensions(uint(v_21), v_22[0u], v_22[1u], v_22[2u]);
+ int2 v_23 = int2(min(v_19, (v_22.xy - (1u).xx)));
+ v_7 = mul(params.yuvToRgbConversionMatrix, float4(v_18, float4(plane_1.Load(int3(v_23, int(v_21)))).xy, 1.0f));
v_8 = 1.0f;
}
- float3 v_27 = v_7;
- float3 v_28 = (0.0f).xxx;
+ float3 v_24 = v_7;
+ float3 v_25 = (0.0f).xxx;
if ((params.doYuvToRgbConversionOnly == 0u)) {
- tint_GammaTransferParams v_29 = params.gammaDecodeParams;
- tint_GammaTransferParams v_30 = params.gammaEncodeParams;
- v_28 = tint_GammaCorrection(mul(tint_GammaCorrection(v_27, v_29), params.gamutConversionMatrix), v_30);
+ tint_GammaTransferParams v_26 = params.gammaDecodeParams;
+ tint_GammaTransferParams v_27 = params.gammaEncodeParams;
+ v_25 = tint_GammaCorrection(mul(tint_GammaCorrection(v_24, v_26), params.gamutConversionMatrix), v_27);
} else {
- v_28 = v_27;
+ v_25 = v_24;
}
- return float4(v_28, v_8);
+ return float4(v_25, v_8);
}
-float3x2 v_31(uint start_byte_offset) {
- uint4 v_32 = t_params[(start_byte_offset / 16u)];
- float2 v_33 = asfloat((((((start_byte_offset % 16u) / 4u) == 2u)) ? (v_32.zw) : (v_32.xy)));
- uint4 v_34 = t_params[((8u + start_byte_offset) / 16u)];
- float2 v_35 = asfloat(((((((8u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_34.zw) : (v_34.xy)));
- uint4 v_36 = t_params[((16u + start_byte_offset) / 16u)];
- return float3x2(v_33, v_35, asfloat(((((((16u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_36.zw) : (v_36.xy))));
+float3x2 v_28(uint start_byte_offset) {
+ uint4 v_29 = t_params[(start_byte_offset / 16u)];
+ float2 v_30 = asfloat((((((start_byte_offset % 16u) / 4u) == 2u)) ? (v_29.zw) : (v_29.xy)));
+ uint4 v_31 = t_params[((8u + start_byte_offset) / 16u)];
+ float2 v_32 = asfloat(((((((8u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_31.zw) : (v_31.xy)));
+ uint4 v_33 = t_params[((16u + start_byte_offset) / 16u)];
+ return float3x2(v_30, v_32, asfloat(((((((16u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_33.zw) : (v_33.xy))));
}
-float3x3 v_37(uint start_byte_offset) {
- float3 v_38 = asfloat(t_params[(start_byte_offset / 16u)].xyz);
- float3 v_39 = asfloat(t_params[((16u + start_byte_offset) / 16u)].xyz);
- return float3x3(v_38, v_39, asfloat(t_params[((32u + start_byte_offset) / 16u)].xyz));
+float3x3 v_34(uint start_byte_offset) {
+ float3 v_35 = asfloat(t_params[(start_byte_offset / 16u)].xyz);
+ float3 v_36 = asfloat(t_params[((16u + start_byte_offset) / 16u)].xyz);
+ return float3x3(v_35, v_36, asfloat(t_params[((32u + start_byte_offset) / 16u)].xyz));
}
-tint_GammaTransferParams v_40(uint start_byte_offset) {
- float v_41 = asfloat(t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)]);
- float v_42 = asfloat(t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)]);
- float v_43 = asfloat(t_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) % 16u) / 4u)]);
- float v_44 = asfloat(t_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) % 16u) / 4u)]);
- float v_45 = asfloat(t_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) % 16u) / 4u)]);
- float v_46 = asfloat(t_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) % 16u) / 4u)]);
- float v_47 = asfloat(t_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) % 16u) / 4u)]);
- tint_GammaTransferParams v_48 = {v_41, v_42, v_43, v_44, v_45, v_46, v_47, t_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) % 16u) / 4u)]};
- return v_48;
+tint_GammaTransferParams v_37(uint start_byte_offset) {
+ float v_38 = asfloat(t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)]);
+ float v_39 = asfloat(t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)]);
+ float v_40 = asfloat(t_params[((8u + start_byte_offset) / 16u)][(((8u + start_byte_offset) % 16u) / 4u)]);
+ float v_41 = asfloat(t_params[((12u + start_byte_offset) / 16u)][(((12u + start_byte_offset) % 16u) / 4u)]);
+ float v_42 = asfloat(t_params[((16u + start_byte_offset) / 16u)][(((16u + start_byte_offset) % 16u) / 4u)]);
+ float v_43 = asfloat(t_params[((20u + start_byte_offset) / 16u)][(((20u + start_byte_offset) % 16u) / 4u)]);
+ float v_44 = asfloat(t_params[((24u + start_byte_offset) / 16u)][(((24u + start_byte_offset) % 16u) / 4u)]);
+ tint_GammaTransferParams v_45 = {v_38, v_39, v_40, v_41, v_42, v_43, v_44, t_params[((28u + start_byte_offset) / 16u)][(((28u + start_byte_offset) % 16u) / 4u)]};
+ return v_45;
}
-float3x4 v_49(uint start_byte_offset) {
- float4 v_50 = asfloat(t_params[(start_byte_offset / 16u)]);
- float4 v_51 = asfloat(t_params[((16u + start_byte_offset) / 16u)]);
- return float3x4(v_50, v_51, asfloat(t_params[((32u + start_byte_offset) / 16u)]));
+float3x4 v_46(uint start_byte_offset) {
+ float4 v_47 = asfloat(t_params[(start_byte_offset / 16u)]);
+ float4 v_48 = asfloat(t_params[((16u + start_byte_offset) / 16u)]);
+ return float3x4(v_47, v_48, asfloat(t_params[((32u + start_byte_offset) / 16u)]));
}
-tint_ExternalTextureParams v_52(uint start_byte_offset) {
- uint v_53 = t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)];
- uint v_54 = t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)];
- float3x4 v_55 = v_49((16u + start_byte_offset));
- tint_GammaTransferParams v_56 = v_40((64u + start_byte_offset));
- tint_GammaTransferParams v_57 = v_40((96u + start_byte_offset));
- float3x3 v_58 = v_37((128u + start_byte_offset));
- float3x2 v_59 = v_31((176u + start_byte_offset));
- float3x2 v_60 = v_31((200u + start_byte_offset));
- uint4 v_61 = t_params[((224u + start_byte_offset) / 16u)];
- float2 v_62 = asfloat(((((((224u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_61.zw) : (v_61.xy)));
- uint4 v_63 = t_params[((232u + start_byte_offset) / 16u)];
- float2 v_64 = asfloat(((((((232u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_63.zw) : (v_63.xy)));
- uint4 v_65 = t_params[((240u + start_byte_offset) / 16u)];
- float2 v_66 = asfloat(((((((240u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_65.zw) : (v_65.xy)));
- uint4 v_67 = t_params[((248u + start_byte_offset) / 16u)];
- float2 v_68 = asfloat(((((((248u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_67.zw) : (v_67.xy)));
- uint4 v_69 = t_params[((256u + start_byte_offset) / 16u)];
- uint2 v_70 = ((((((256u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_69.zw) : (v_69.xy));
- uint4 v_71 = t_params[((264u + start_byte_offset) / 16u)];
- tint_ExternalTextureParams v_72 = {v_53, v_54, v_55, v_56, v_57, v_58, v_59, v_60, v_62, v_64, v_66, v_68, v_70, asfloat(((((((264u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_71.zw) : (v_71.xy)))};
- return v_72;
+tint_ExternalTextureParams v_49(uint start_byte_offset) {
+ uint v_50 = t_params[(start_byte_offset / 16u)][((start_byte_offset % 16u) / 4u)];
+ uint v_51 = t_params[((4u + start_byte_offset) / 16u)][(((4u + start_byte_offset) % 16u) / 4u)];
+ float3x4 v_52 = v_46((16u + start_byte_offset));
+ tint_GammaTransferParams v_53 = v_37((64u + start_byte_offset));
+ tint_GammaTransferParams v_54 = v_37((96u + start_byte_offset));
+ float3x3 v_55 = v_34((128u + start_byte_offset));
+ float3x2 v_56 = v_28((176u + start_byte_offset));
+ float3x2 v_57 = v_28((200u + start_byte_offset));
+ uint4 v_58 = t_params[((224u + start_byte_offset) / 16u)];
+ float2 v_59 = asfloat(((((((224u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_58.zw) : (v_58.xy)));
+ uint4 v_60 = t_params[((232u + start_byte_offset) / 16u)];
+ float2 v_61 = asfloat(((((((232u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_60.zw) : (v_60.xy)));
+ uint4 v_62 = t_params[((240u + start_byte_offset) / 16u)];
+ float2 v_63 = asfloat(((((((240u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_62.zw) : (v_62.xy)));
+ uint4 v_64 = t_params[((248u + start_byte_offset) / 16u)];
+ float2 v_65 = asfloat(((((((248u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_64.zw) : (v_64.xy)));
+ uint4 v_66 = t_params[((256u + start_byte_offset) / 16u)];
+ uint2 v_67 = ((((((256u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_66.zw) : (v_66.xy));
+ uint4 v_68 = t_params[((264u + start_byte_offset) / 16u)];
+ tint_ExternalTextureParams v_69 = {v_50, v_51, v_52, v_53, v_54, v_55, v_56, v_57, v_59, v_61, v_63, v_65, v_67, asfloat(((((((264u + start_byte_offset) % 16u) / 4u) == 2u)) ? (v_68.zw) : (v_68.xy)))};
+ return v_69;
}
[numthreads(1, 1, 1)]
void main() {
- tint_ExternalTextureParams v_73 = v_52(0u);
- float4 red = tint_TextureLoadExternal(t_plane0, t_plane1, v_73, uint2((int(10)).xx));
+ tint_ExternalTextureParams v_70 = v_49(0u);
+ float4 red = tint_TextureLoadExternal(t_plane0, t_plane1, v_70, uint2((int(10)).xx));
outImage[(int(0)).xx] = red;
- tint_ExternalTextureParams v_74 = v_52(0u);
- float4 green = tint_TextureLoadExternal(t_plane0, t_plane1, v_74, uint2(int2(int(70), int(118))));
+ tint_ExternalTextureParams v_71 = v_49(0u);
+ float4 green = tint_TextureLoadExternal(t_plane0, t_plane1, v_71, uint2(int2(int(70), int(118))));
outImage[int2(int(1), int(0))] = green;
}
