// Copyright 2020 The Tint Authors. // // Licensed under the Apache License, Version 2.0 (the "License"); // you may not use this file except in compliance with the License. // You may obtain a copy of the License at // // http://www.apache.org/licenses/LICENSE-2.0 // // Unless required by applicable law or agreed to in writing, software // distributed under the License is distributed on an "AS IS" BASIS, // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. // See the License for the specific language governing permissions and // limitations under the License. #include "src/tint/transform/robustness.h" #include "src/tint/transform/test_helper.h" namespace tint::transform { static std::ostream& operator<<(std::ostream& out, Robustness::Action action) { switch (action) { case Robustness::Action::kIgnore: return out << "ignore"; case Robustness::Action::kClamp: return out << "clamp"; case Robustness::Action::kPredicate: return out << "predicate"; } return out << "unknown"; } namespace { DataMap Config(Robustness::Action action) { Robustness::Config cfg; cfg.value_action = action; cfg.texture_action = action; cfg.function_action = action; cfg.private_action = action; cfg.push_constant_action = action; cfg.storage_action = action; cfg.uniform_action = action; cfg.workgroup_action = action; DataMap data; data.Add(cfg); return data; } const char* Expect(Robustness::Action action, const char* expect_ignore, const char* expect_clamp, const char* expect_predicate) { switch (action) { case Robustness::Action::kIgnore: return expect_ignore; case Robustness::Action::kClamp: return expect_clamp; case Robustness::Action::kPredicate: return expect_predicate; } return ""; } using RobustnessTest = TransformTestWithParam; //////////////////////////////////////////////////////////////////////////////// // Constant sized array //////////////////////////////////////////////////////////////////////////////// TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithLiteral) { auto* src = R"( fn f() { var b : f32 = array()[1i]; } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithConst) { auto* src = R"( const c : u32 = 1u; fn f() { let b : f32 = array()[c]; } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithLet) { auto* src = R"( fn f() { let l : u32 = 1u; let b : f32 = array()[l]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( fn f() { let l : u32 = 1u; let b : f32 = array()[min(l, 2u)]; } )", /* predicate */ R"( fn f() { let l : u32 = 1u; let index = l; let predicate = (u32(index) <= 2u); var predicated_expr : f32; if (predicate) { predicated_expr = array()[index]; } let b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithRuntimeArrayIndex) { auto* src = R"( var i : u32; fn f() { let a = array(); let b = array(); var c : f32 = a[b[i]]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var i : u32; fn f() { let a = array(); let b = array(); var c : f32 = a[min(u32(b[min(i, 4u)]), 2u)]; } )", /* predicate */ R"( var i : u32; fn f() { let a = array(); let b = array(); let index = i; let predicate = (u32(index) <= 4u); var predicated_expr : i32; if (predicate) { predicated_expr = b[index]; } let index_1 = predicated_expr; let predicate_1 = (u32(index_1) <= 2u); var predicated_expr_1 : f32; if (predicate_1) { predicated_expr_1 = a[index_1]; } var c : f32 = predicated_expr_1; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithRuntimeExpression) { auto* src = R"( var c : i32; fn f() { var b : f32 = array()[((c + 2) - 3)]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var c : i32; fn f() { var b : f32 = array()[min(u32(((c + 2) - 3)), 2u)]; } )", /* predicate */ R"( var c : i32; fn f() { let index = ((c + 2) - 3); let predicate = (u32(index) <= 2u); var predicated_expr : f32; if (predicate) { predicated_expr = array()[index]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_NestedConstantSizedArraysVal_IndexWithRuntimeExpressions) { auto* src = R"( var x : i32; var y : i32; var z : i32; fn f() { let a = array, 2>, 3>(); var r = a[x][y][z]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var x : i32; var y : i32; var z : i32; fn f() { let a = array, 2>, 3>(); var r = a[min(u32(x), 2u)][min(u32(y), 1u)][min(u32(z), 0u)]; } )", /* predicate */ R"( var x : i32; var y : i32; var z : i32; fn f() { let a = array, 2>, 3>(); let index = x; let predicate = (u32(index) <= 2u); var predicated_expr : array, 2u>; if (predicate) { predicated_expr = a[index]; } let index_1 = y; let predicate_1 = (u32(index_1) <= 1u); var predicated_expr_1 : array; if (predicate_1) { predicated_expr_1 = predicated_expr[index_1]; } let index_2 = z; let predicate_2 = (u32(index_2) <= 0u); var predicated_expr_2 : f32; if (predicate_2) { predicated_expr_2 = predicated_expr_1[index_2]; } var r = predicated_expr_2; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithOverride) { auto* src = R"( @id(1300) override idx : i32; fn f() { let a = array(); var b : f32 = a[idx]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @id(1300) override idx : i32; fn f() { let a = array(); var b : f32 = a[min(u32(idx), 3u)]; } )", /* predicate */ R"( @id(1300) override idx : i32; fn f() { let a = array(); let index = idx; let predicate = (u32(index) <= 3u); var predicated_expr : f32; if (predicate) { predicated_expr = a[index]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithLiteral) { auto* src = R"( var a : array; fn f() { var b : f32 = a[1i]; } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithConst) { auto* src = R"( var a : array; const c : u32 = 1u; fn f() { let b : f32 = a[c]; } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithLet) { auto* src = R"( var a : array; fn f() { let l : u32 = 1u; let b : f32 = a[l]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array; fn f() { let l : u32 = 1u; let b : f32 = a[min(l, 2u)]; } )", /* predicate */ R"( var a : array; fn f() { let l : u32 = 1u; let index = l; let predicate = (u32(index) <= 2u); var predicated_expr : f32; if (predicate) { predicated_expr = a[index]; } let b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithRuntimeArrayIndex) { auto* src = R"( var a : array; var b : array; var i : u32; fn f() { var c : f32 = a[b[i]]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array; var b : array; var i : u32; fn f() { var c : f32 = a[min(u32(b[min(i, 4u)]), 2u)]; } )", /* predicate */ R"( var a : array; var b : array; var i : u32; fn f() { let index = i; let predicate = (u32(index) <= 4u); var predicated_expr : i32; if (predicate) { predicated_expr = b[index]; } let index_1 = predicated_expr; let predicate_1 = (u32(index_1) <= 2u); var predicated_expr_1 : f32; if (predicate_1) { predicated_expr_1 = a[index_1]; } var c : f32 = predicated_expr_1; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithRuntimeExpression) { auto* src = R"( var a : array; var c : i32; fn f() { var b : f32 = a[((c + 2) - 3)]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array; var c : i32; fn f() { var b : f32 = a[min(u32(((c + 2) - 3)), 2u)]; } )", /* predicate */ R"( var a : array; var c : i32; fn f() { let index = ((c + 2) - 3); let predicate = (u32(index) <= 2u); var predicated_expr : f32; if (predicate) { predicated_expr = a[index]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_NestedConstantSizedArraysRef_IndexWithRuntimeExpressions) { auto* src = R"( var a : array, 2>, 3>; var x : i32; var y : i32; var z : i32; fn f() { var r = a[x][y][z]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array, 2>, 3>; var x : i32; var y : i32; var z : i32; fn f() { var r = a[min(u32(x), 2u)][min(u32(y), 1u)][min(u32(z), 0u)]; } )", /* predicate */ R"( var a : array, 2>, 3>; var x : i32; var y : i32; var z : i32; fn f() { let index = x; let predicate = (u32(index) <= 2u); let index_1 = y; let predicate_1 = (predicate & (u32(index_1) <= 1u)); let index_2 = z; let predicate_2 = (predicate_1 & (u32(index_2) <= 0u)); var predicated_expr : f32; if (predicate_2) { predicated_expr = a[index][index_1][index_2]; } var r = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithOverride) { auto* src = R"( @id(1300) override idx : i32; fn f() { var a : array; var b : f32 = a[idx]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @id(1300) override idx : i32; fn f() { var a : array; var b : f32 = a[min(u32(idx), 3u)]; } )", /* predicate */ R"( @id(1300) override idx : i32; fn f() { var a : array; let index = idx; let predicate = (u32(index) <= 3u); var predicated_expr : f32; if (predicate) { predicated_expr = a[index]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayPtr_IndexWithLet) { auto* src = R"( var a : array; fn f() { let l : u32 = 1u; let p = &(a[l]); let f : f32 = *(p); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array; fn f() { let l : u32 = 1u; let p = &(a[min(l, 2u)]); let f : f32 = *(p); } )", /* predicate */ R"( var a : array; fn f() { let l : u32 = 1u; let index = l; let predicate = (u32(index) <= 2u); let p = &(a[index]); var predicated_expr : f32; if (predicate) { predicated_expr = *(p); } let f : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_ConstantSizedArrayPtr_IndexWithRuntimeArrayIndex) { auto* src = R"( var a : array; var b : array; var i : u32; fn f() { let pa = &(a); let pb = &(b); let p0 = &((*(pb))[i]); let p1 = &(a[*(p0)]); var x : f32 = *(p1); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array; var b : array; var i : u32; fn f() { let pa = &(a); let pb = &(b); let p0 = &((*(pb))[min(i, 4u)]); let p1 = &(a[min(u32(*(p0)), 2u)]); var x : f32 = *(p1); } )", /* predicate */ R"( var a : array; var b : array; var i : u32; fn f() { let pa = &(a); let pb = &(b); let index = i; let predicate = (u32(index) <= 4u); let p0 = &((*(pb))[index]); var predicated_expr : i32; if (predicate) { predicated_expr = *(p0); } let index_1 = predicated_expr; let predicate_1 = (u32(index_1) <= 2u); let p1 = &(a[index_1]); var predicated_expr_1 : f32; if (predicate_1) { predicated_expr_1 = *(p1); } var x : f32 = predicated_expr_1; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_NestedConstantSizedArraysPtr_IndexWithRuntimeExpressions) { auto* src = R"( var a : array, 2>, 3>; var x : i32; var y : i32; var z : i32; fn f() { let p0 = &(a); let p1 = &((*(p0))[x]); let p2 = &((*(p1))[y]); let p3 = &((*(p2))[z]); var r = *(p3); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array, 2>, 3>; var x : i32; var y : i32; var z : i32; fn f() { let p0 = &(a); let p1 = &((*(p0))[min(u32(x), 2u)]); let p2 = &((*(p1))[min(u32(y), 1u)]); let p3 = &((*(p2))[min(u32(z), 0u)]); var r = *(p3); } )", /* predicate */ R"( var a : array, 2>, 3>; var x : i32; var y : i32; var z : i32; fn f() { let p0 = &(a); let index = x; let predicate = (u32(index) <= 2u); let p1 = &((*(p0))[index]); let index_1 = y; let predicate_1 = (predicate & (u32(index_1) <= 1u)); let p2 = &((*(p1))[index_1]); let index_2 = z; let predicate_2 = (predicate_1 & (u32(index_2) <= 0u)); let p3 = &((*(p2))[index_2]); var predicated_expr : f32; if (predicate_2) { predicated_expr = *(p3); } var r = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_NestedConstantSizedArrays_MixedAccess) { auto* src = R"( var a : array, 2>, 3>; var x : i32; const y = 1; override z : i32; fn f() { let p0 = &(a); let p1 = &((*(p0))[x]); let p2 = &((*(p1))[y]); let p3 = &((*(p2))[z]); var r = *(p3); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array, 2>, 3>; var x : i32; const y = 1; override z : i32; fn f() { let p0 = &(a); let p1 = &((*(p0))[min(u32(x), 2u)]); let p2 = &((*(p1))[y]); let p3 = &((*(p2))[min(u32(z), 0u)]); var r = *(p3); } )", /* predicate */ R"( var a : array, 2>, 3>; var x : i32; const y = 1; override z : i32; fn f() { let p0 = &(a); let index = x; let predicate = (u32(index) <= 2u); let p1 = &((*(p0))[index]); let p2 = &((*(p1))[y]); let index_1 = z; let predicate_1 = (predicate & (u32(index_1) <= 0u)); let p3 = &((*(p2))[index_1]); var predicated_expr : f32; if (predicate_1) { predicated_expr = *(p3); } var r = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Assign_ConstantSizedArray_IndexWithLet) { auto* src = R"( var a : array; fn f() { let l : u32 = 1u; a[l] = 42.0f; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array; fn f() { let l : u32 = 1u; a[min(l, 2u)] = 42.0f; } )", /* predicate */ R"( var a : array; fn f() { let l : u32 = 1u; let index = l; let predicate = (u32(index) <= 2u); if (predicate) { a[index] = 42.0f; } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Assign_ConstantSizedArrayPtr_IndexWithLet) { auto* src = R"( var a : array; fn f() { let l : u32 = 1u; let p = &(a[l]); *(p) = 42.0f; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array; fn f() { let l : u32 = 1u; let p = &(a[min(l, 2u)]); *(p) = 42.0f; } )", /* predicate */ R"( var a : array; fn f() { let l : u32 = 1u; let index = l; let predicate = (u32(index) <= 2u); let p = &(a[index]); if (predicate) { *(p) = 42.0f; } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Assign_ConstantSizedArrayPtr_IndexWithRuntimeArrayIndex) { auto* src = R"( var a : array; var b : array; var i : u32; fn f() { let pa = &(a); let pb = &(b); let p0 = &((*(pb))[i]); let p1 = &(a[*(p0)]); *(p1) = 42.0f; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array; var b : array; var i : u32; fn f() { let pa = &(a); let pb = &(b); let p0 = &((*(pb))[min(i, 4u)]); let p1 = &(a[min(u32(*(p0)), 2u)]); *(p1) = 42.0f; } )", /* predicate */ R"( var a : array; var b : array; var i : u32; fn f() { let pa = &(a); let pb = &(b); let index = i; let predicate = (u32(index) <= 4u); let p0 = &((*(pb))[index]); var predicated_expr : i32; if (predicate) { predicated_expr = *(p0); } let index_1 = predicated_expr; let predicate_1 = (u32(index_1) <= 2u); let p1 = &(a[index_1]); if (predicate_1) { *(p1) = 42.0f; } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Assign_NestedConstantSizedArraysPtr_IndexWithRuntimeExpressions) { auto* src = R"( var a : array, 2>, 3>; var x : i32; var y : i32; var z : i32; fn f() { let p0 = &(a); let p1 = &((*(p0))[x]); let p2 = &((*(p1))[y]); let p3 = &((*(p2))[z]); *(p3) = 42.0f; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array, 2>, 3>; var x : i32; var y : i32; var z : i32; fn f() { let p0 = &(a); let p1 = &((*(p0))[min(u32(x), 2u)]); let p2 = &((*(p1))[min(u32(y), 1u)]); let p3 = &((*(p2))[min(u32(z), 0u)]); *(p3) = 42.0f; } )", /* predicate */ R"( var a : array, 2>, 3>; var x : i32; var y : i32; var z : i32; fn f() { let p0 = &(a); let index = x; let predicate = (u32(index) <= 2u); let p1 = &((*(p0))[index]); let index_1 = y; let predicate_1 = (predicate & (u32(index_1) <= 1u)); let p2 = &((*(p1))[index_1]); let index_2 = z; let predicate_2 = (predicate_1 & (u32(index_2) <= 0u)); let p3 = &((*(p2))[index_2]); if (predicate_2) { *(p3) = 42.0f; } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Assign_NestedConstantSizedArrays_MixedAccess) { auto* src = R"( var a : array, 2>, 3>; var x : i32; const y = 1; override z : i32; fn f() { let p0 = &(a); let p1 = &((*(p0))[x]); let p2 = &((*(p1))[y]); let p3 = &((*(p2))[z]); *(p3) = 42.0f; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array, 2>, 3>; var x : i32; const y = 1; override z : i32; fn f() { let p0 = &(a); let p1 = &((*(p0))[min(u32(x), 2u)]); let p2 = &((*(p1))[y]); let p3 = &((*(p2))[min(u32(z), 0u)]); *(p3) = 42.0f; } )", /* predicate */ R"( var a : array, 2>, 3>; var x : i32; const y = 1; override z : i32; fn f() { let p0 = &(a); let index = x; let predicate = (u32(index) <= 2u); let p1 = &((*(p0))[index]); let p2 = &((*(p1))[y]); let index_1 = z; let predicate_1 = (predicate & (u32(index_1) <= 0u)); let p3 = &((*(p2))[index_1]); if (predicate_1) { *(p3) = 42.0f; } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, CompoundAssign_ConstantSizedArray_IndexWithLet) { auto* src = R"( var a : array; fn f() { let l : u32 = 1u; a[l] += 42.0f; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array; fn f() { let l : u32 = 1u; a[min(l, 2u)] += 42.0f; } )", /* predicate */ R"( var a : array; fn f() { let l : u32 = 1u; let index = l; let predicate = (u32(index) <= 2u); if (predicate) { a[index] += 42.0f; } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Increment_ConstantSizedArray_IndexWithLet) { auto* src = R"( var a : array; fn f() { let l : u32 = 1u; a[l]++; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array; fn f() { let l : u32 = 1u; a[min(l, 2u)]++; } )", /* predicate */ R"( var a : array; fn f() { let l : u32 = 1u; let index = l; let predicate = (u32(index) <= 2u); if (predicate) { a[index]++; } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } //////////////////////////////////////////////////////////////////////////////// // Runtime sized array //////////////////////////////////////////////////////////////////////////////// TEST_P(RobustnessTest, Read_RuntimeArray_IndexWithLiteral) { auto* src = R"( struct S { a : f32, b : array, } @group(0) @binding(0) var s : S; fn f() { var d : f32 = s.b[25]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( struct S { a : f32, b : array, } @group(0) @binding(0) var s : S; fn f() { var d : f32 = s.b[min(u32(25), (arrayLength(&(s.b)) - 1u))]; } )", /* predicate */ R"( struct S { a : f32, b : array, } @group(0) @binding(0) var s : S; fn f() { let index = 25; let predicate = (u32(index) <= (arrayLength(&(s.b)) - 1u)); var predicated_expr : f32; if (predicate) { predicated_expr = s.b[index]; } var d : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } //////////////////////////////////////////////////////////////////////////////// // Vector //////////////////////////////////////////////////////////////////////////////// TEST_P(RobustnessTest, Read_Vector_IndexWithLiteral) { auto* src = R"( var a : vec3; fn f() { var b : f32 = a[1i]; } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_Vector_IndexWithConst) { auto* src = R"( var a : vec3; fn f() { const i = 1; var b : f32 = a[i]; } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_Vector_IndexWithLet) { auto* src = R"( fn f() { let i = 99; let v = vec4()[i]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( fn f() { let i = 99; let v = vec4()[min(u32(i), 3u)]; } )", /* predicate */ R"( fn f() { let i = 99; let index = i; let predicate = (u32(index) <= 3u); var predicated_expr : f32; if (predicate) { predicated_expr = vec4()[index]; } let v = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_Vector_IndexWithRuntimeExpression) { auto* src = R"( var a : vec3; var c : i32; fn f() { var b : f32 = a[((c + 2) - 3)]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : vec3; var c : i32; fn f() { var b : f32 = a[min(u32(((c + 2) - 3)), 2u)]; } )", /* predicate */ R"( var a : vec3; var c : i32; fn f() { let index = ((c + 2) - 3); let predicate = (u32(index) <= 2u); var predicated_expr : f32; if (predicate) { predicated_expr = a[index]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_Vector_SwizzleIndexWithGlobalVar) { auto* src = R"( var a : vec3; var c : i32; fn f() { var b : f32 = a.xy[c]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : vec3; var c : i32; fn f() { var b : f32 = a.xy[min(u32(c), 1u)]; } )", /* predicate */ R"( var a : vec3; var c : i32; fn f() { let index = c; let predicate = (u32(index) <= 1u); var predicated_expr : f32; if (predicate) { predicated_expr = a.xy[index]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_Vector_SwizzleIndexWithRuntimeExpression) { auto* src = R"( var a : vec3; var c : i32; fn f() { var b : f32 = a.xy[((c + 2) - 3)]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : vec3; var c : i32; fn f() { var b : f32 = a.xy[min(u32(((c + 2) - 3)), 1u)]; } )", /* predicate */ R"( var a : vec3; var c : i32; fn f() { let index = ((c + 2) - 3); let predicate = (u32(index) <= 1u); var predicated_expr : f32; if (predicate) { predicated_expr = a.xy[index]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_Vector_IndexWithOverride) { auto* src = R"( @id(1300) override idx : i32; fn f() { var a : vec3; var b : f32 = a[idx]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @id(1300) override idx : i32; fn f() { var a : vec3; var b : f32 = a[min(u32(idx), 2u)]; } )", /* predicate */ R"( @id(1300) override idx : i32; fn f() { var a : vec3; let index = idx; let predicate = (u32(index) <= 2u); var predicated_expr : f32; if (predicate) { predicated_expr = a[index]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } //////////////////////////////////////////////////////////////////////////////// // Matrix //////////////////////////////////////////////////////////////////////////////// TEST_P(RobustnessTest, Read_MatrixRef_IndexingWithLiterals) { auto* src = R"( var a : mat3x2; fn f() { var b : f32 = a[2i][1i]; } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_MatrixRef_IndexWithRuntimeExpressionThenLiteral) { auto* src = R"( var a : mat3x2; var c : i32; fn f() { var b : f32 = a[((c + 2) - 3)][1]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : mat3x2; var c : i32; fn f() { var b : f32 = a[min(u32(((c + 2) - 3)), 2u)][1]; } )", /* predicate */ R"( var a : mat3x2; var c : i32; fn f() { let index = ((c + 2) - 3); let predicate = (u32(index) <= 2u); var predicated_expr : f32; if (predicate) { predicated_expr = a[index][1]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_MatrixRef_IndexWithLiteralThenRuntimeExpression) { auto* src = R"( var a : mat3x2; var c : i32; fn f() { var b : f32 = a[1][((c + 2) - 3)]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : mat3x2; var c : i32; fn f() { var b : f32 = a[1][min(u32(((c + 2) - 3)), 1u)]; } )", /* predicate */ R"( var a : mat3x2; var c : i32; fn f() { let index = ((c + 2) - 3); let predicate = (u32(index) <= 1u); var predicated_expr : f32; if (predicate) { predicated_expr = a[1][index]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_MatrixRef_IndexWithOverrideThenLiteral) { auto* src = R"( @id(1300) override idx : i32; fn f() { var a : mat3x2; var b : f32 = a[idx][1]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @id(1300) override idx : i32; fn f() { var a : mat3x2; var b : f32 = a[min(u32(idx), 2u)][1]; } )", /* predicate */ R"( @id(1300) override idx : i32; fn f() { var a : mat3x2; let index = idx; let predicate = (u32(index) <= 2u); var predicated_expr : f32; if (predicate) { predicated_expr = a[index][1]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_MatrixRef_IndexWithLetThenSwizzle) { auto* src = R"( fn f() { let i = 1; var m = mat3x2(); var v = m[i].yx; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( fn f() { let i = 1; var m = mat3x2(); var v = m[min(u32(i), 2u)].yx; } )", /* predicate */ R"( fn f() { let i = 1; var m = mat3x2(); let index = i; let predicate = (u32(index) <= 2u); var predicated_expr : vec2; if (predicate) { predicated_expr = m[index]; } var v = predicated_expr.yx; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_MatrixRef_IndexWithLiteralThenOverride) { auto* src = R"( @id(1300) override idx : i32; fn f() { var a : mat3x2; var b : f32 = a[1][idx]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @id(1300) override idx : i32; fn f() { var a : mat3x2; var b : f32 = a[1][min(u32(idx), 1u)]; } )", /* predicate */ R"( @id(1300) override idx : i32; fn f() { var a : mat3x2; let index = idx; let predicate = (u32(index) <= 1u); var predicated_expr : f32; if (predicate) { predicated_expr = a[1][index]; } var b : f32 = predicated_expr; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Assign_Matrix_IndexWithLet) { auto* src = R"( var m : mat3x4f; fn f() { let c = 1; m[c] = vec4f(1); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var m : mat3x4f; fn f() { let c = 1; m[min(u32(c), 2u)] = vec4f(1); } )", /* predicate */ R"( var m : mat3x4f; fn f() { let c = 1; let index = c; let predicate = (u32(index) <= 2u); if (predicate) { m[index] = vec4f(1); } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, CompoundAssign_Matrix_IndexWithLet) { auto* src = R"( var m : mat3x4f; fn f() { let c = 1; m[c] += vec4f(1); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var m : mat3x4f; fn f() { let c = 1; m[min(u32(c), 2u)] += vec4f(1); } )", /* predicate */ R"( var m : mat3x4f; fn f() { let c = 1; let index = c; let predicate = (u32(index) <= 2u); if (predicate) { m[index] += vec4f(1); } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } //////////////////////////////////////////////////////////////////////////////// // Texture //////////////////////////////////////////////////////////////////////////////// TEST_P(RobustnessTest, TextureDimensions) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; fn dimensions() { let l = textureDimensions(t); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureDimensions_Level) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; fn dimensions_signed(level : i32) { let l = textureDimensions(t, level); } fn dimensions_unsigned(level : u32) { let l = textureDimensions(t, level); } @fragment fn main(@builtin(position) non_uniform : vec4f) { dimensions_signed(i32(non_uniform.x)); dimensions_unsigned(u32(non_uniform.x)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_2d; fn dimensions_signed(level : i32) { let level_idx = min(u32(level), (textureNumLevels(t) - 1)); let l = textureDimensions(t, level_idx); } fn dimensions_unsigned(level : u32) { let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureDimensions(t, level_idx_1); } @fragment fn main(@builtin(position) non_uniform : vec4f) { dimensions_signed(i32(non_uniform.x)); dimensions_unsigned(u32(non_uniform.x)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_2d; fn dimensions_signed(level : i32) { let level_idx = u32(level); let num_levels = textureNumLevels(t); var predicated_value : vec2; if ((level_idx < num_levels)) { predicated_value = textureDimensions(t, level_idx); } let l = predicated_value; } fn dimensions_unsigned(level : u32) { let level_idx_1 = u32(level); let num_levels_1 = textureNumLevels(t); var predicated_value_1 : vec2; if ((level_idx_1 < num_levels_1)) { predicated_value_1 = textureDimensions(t, level_idx_1); } let l = predicated_value_1; } @fragment fn main(@builtin(position) non_uniform : vec4f) { dimensions_signed(i32(non_uniform.x)); dimensions_unsigned(u32(non_uniform.x)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureGather) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; @group(0) @binding(1) var s : sampler; fn gather(coords : vec2f) { let l = textureGather(0, t, s, coords); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureGather_Array) { auto* src = R"( @group(0) @binding(0) var t : texture_2d_array; @group(0) @binding(1) var s : sampler; fn gather_signed(coords : vec2f, array : i32) { let l = textureGather(1, t, s, coords, array); } fn gather_unsigned(coords : vec2f, array : u32) { let l = textureGather(1, t, s, coords, array); } @fragment fn main(@builtin(position) non_uniform : vec4f) { gather_signed(non_uniform.xy, i32(non_uniform.x)); gather_unsigned(non_uniform.xy, u32(non_uniform.x)); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureGatherCompare) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d; @group(0) @binding(1) var s : sampler_comparison; fn gather(coords : vec2f, depth_ref : f32) { let l = textureGatherCompare(t, s, coords, depth_ref); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureGatherCompare_Array) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d_array; @group(0) @binding(1) var s : sampler_comparison; fn gather_signed(coords : vec2f, array : i32, depth_ref : f32) { let l = textureGatherCompare(t, s, coords, array, depth_ref); } fn gather_unsigned(coords : vec2f, array : u32, depth_ref : f32) { let l = textureGatherCompare(t, s, coords, array, depth_ref); } @fragment fn main(@builtin(position) non_uniform : vec4f) { gather_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.x); gather_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.x); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureLoad_1D) { auto* src = R"( @group(0) @binding(0) var t : texture_1d; fn load_signed(coords : i32, level : i32) { let l = textureLoad(t, coords, level); } fn load_unsigned(coords : u32, level : u32) { let l = textureLoad(t, coords, level); } fn load_mixed(coords : i32, level : u32) { let l = textureLoad(t, coords, level); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(i32(non_uniform.x), i32(non_uniform.x)); load_unsigned(u32(non_uniform.x), u32(non_uniform.x)); load_mixed(i32(non_uniform.x), u32(non_uniform.x)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_1d; fn load_signed(coords : i32, level : i32) { let level_idx = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, clamp(coords, 0, i32((textureDimensions(t, level_idx) - 1))), level_idx); } fn load_unsigned(coords : u32, level : u32) { let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - 1)), level_idx_1); } fn load_mixed(coords : i32, level : u32) { let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, clamp(coords, 0, i32((textureDimensions(t, level_idx_2) - 1))), level_idx_2); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(i32(non_uniform.x), i32(non_uniform.x)); load_unsigned(u32(non_uniform.x), u32(non_uniform.x)); load_mixed(i32(non_uniform.x), u32(non_uniform.x)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_1d; fn load_signed(coords : i32, level : i32) { let level_idx = u32(level); let num_levels = textureNumLevels(t); let coords_1 = u32(coords); var predicated_value : vec4; if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) { predicated_value = textureLoad(t, coords_1, level_idx); } let l = predicated_value; } fn load_unsigned(coords : u32, level : u32) { let level_idx_1 = u32(level); let num_levels_1 = textureNumLevels(t); let coords_2 = u32(coords); var predicated_value_1 : vec4; if (((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1))))))) { predicated_value_1 = textureLoad(t, coords_2, level_idx_1); } let l = predicated_value_1; } fn load_mixed(coords : i32, level : u32) { let level_idx_2 = u32(level); let num_levels_2 = textureNumLevels(t); let coords_3 = u32(coords); var predicated_value_2 : vec4; if (((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1))))))) { predicated_value_2 = textureLoad(t, coords_3, level_idx_2); } let l = predicated_value_2; } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(i32(non_uniform.x), i32(non_uniform.x)); load_unsigned(u32(non_uniform.x), u32(non_uniform.x)); load_mixed(i32(non_uniform.x), u32(non_uniform.x)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureLoad_2D) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; fn load_signed(coords : vec2i, level : i32) { let l = textureLoad(t, coords, level); } fn load_unsigned(coords : vec2u, level : u32) { let l = textureLoad(t, coords, level); } fn load_mixed(coords : vec2u, level : i32) { let l = textureLoad(t, coords, level); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x)); load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_2d; fn load_signed(coords : vec2i, level : i32) { let level_idx = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, clamp(coords, vec2(0), vec2((textureDimensions(t, level_idx) - vec2(1)))), level_idx); } fn load_unsigned(coords : vec2u, level : u32) { let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - vec2(1))), level_idx_1); } fn load_mixed(coords : vec2u, level : i32) { let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_2) - vec2(1))), level_idx_2); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x)); load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_2d; fn load_signed(coords : vec2i, level : i32) { let level_idx = u32(level); let num_levels = textureNumLevels(t); let coords_1 = vec2(coords); var predicated_value : vec4; if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) { predicated_value = textureLoad(t, coords_1, level_idx); } let l = predicated_value; } fn load_unsigned(coords : vec2u, level : u32) { let level_idx_1 = u32(level); let num_levels_1 = textureNumLevels(t); let coords_2 = vec2(coords); var predicated_value_1 : vec4; if (((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1))))))) { predicated_value_1 = textureLoad(t, coords_2, level_idx_1); } let l = predicated_value_1; } fn load_mixed(coords : vec2u, level : i32) { let level_idx_2 = u32(level); let num_levels_2 = textureNumLevels(t); let coords_3 = vec2(coords); var predicated_value_2 : vec4; if (((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1))))))) { predicated_value_2 = textureLoad(t, coords_3, level_idx_2); } let l = predicated_value_2; } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x)); load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureLoad_2DArray) { auto* src = R"( @group(0) @binding(0) var t : texture_2d_array; fn load_signed(coords : vec2i, array : i32, level : i32) { let l = textureLoad(t, coords, array, level); } fn load_unsigned(coords : vec2u, array : u32, level : u32) { let l = textureLoad(t, coords, array, level); } fn load_mixed(coords : vec2u, array : i32, level : u32) { let l = textureLoad(t, coords, array, level); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x)); load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_2d_array; fn load_signed(coords : vec2i, array : i32, level : i32) { let level_idx = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, clamp(coords, vec2(0), vec2((textureDimensions(t, level_idx) - vec2(1)))), clamp(array, 0, i32((textureNumLayers(t) - 1))), level_idx); } fn load_unsigned(coords : vec2u, array : u32, level : u32) { let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - vec2(1))), min(array, (textureNumLayers(t) - 1)), level_idx_1); } fn load_mixed(coords : vec2u, array : i32, level : u32) { let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_2) - vec2(1))), clamp(array, 0, i32((textureNumLayers(t) - 1))), level_idx_2); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x)); load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_2d_array; fn load_signed(coords : vec2i, array : i32, level : i32) { let level_idx = u32(level); let num_levels = textureNumLevels(t); let coords_1 = vec2(coords); let array_idx = u32(array); var predicated_value : vec4; if ((((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1)))))) & (array_idx < textureNumLayers(t)))) { predicated_value = textureLoad(t, coords_1, array_idx, level_idx); } let l = predicated_value; } fn load_unsigned(coords : vec2u, array : u32, level : u32) { let level_idx_1 = u32(level); let num_levels_1 = textureNumLevels(t); let coords_2 = vec2(coords); let array_idx_1 = u32(array); var predicated_value_1 : vec4; if ((((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1)))))) & (array_idx_1 < textureNumLayers(t)))) { predicated_value_1 = textureLoad(t, coords_2, array_idx_1, level_idx_1); } let l = predicated_value_1; } fn load_mixed(coords : vec2u, array : i32, level : u32) { let level_idx_2 = u32(level); let num_levels_2 = textureNumLevels(t); let coords_3 = vec2(coords); let array_idx_2 = u32(array); var predicated_value_2 : vec4; if ((((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1)))))) & (array_idx_2 < textureNumLayers(t)))) { predicated_value_2 = textureLoad(t, coords_3, array_idx_2, level_idx_2); } let l = predicated_value_2; } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x)); load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureLoad_3D) { auto* src = R"( @group(0) @binding(0) var t : texture_3d; fn load_signed(coords : vec3i, level : i32) { let l = textureLoad(t, coords, level); } fn load_unsigned(coords : vec3u, level : u32) { let l = textureLoad(t, coords, level); } fn load_mixed(coords : vec3u, level : i32) { let l = textureLoad(t, coords, level); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec3i(non_uniform.xyz), i32(non_uniform.x)); load_unsigned(vec3u(non_uniform.xyz), u32(non_uniform.x)); load_mixed(vec3u(non_uniform.xyz), i32(non_uniform.x)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_3d; fn load_signed(coords : vec3i, level : i32) { let level_idx = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, clamp(coords, vec3(0), vec3((textureDimensions(t, level_idx) - vec3(1)))), level_idx); } fn load_unsigned(coords : vec3u, level : u32) { let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - vec3(1))), level_idx_1); } fn load_mixed(coords : vec3u, level : i32) { let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_2) - vec3(1))), level_idx_2); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec3i(non_uniform.xyz), i32(non_uniform.x)); load_unsigned(vec3u(non_uniform.xyz), u32(non_uniform.x)); load_mixed(vec3u(non_uniform.xyz), i32(non_uniform.x)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_3d; fn load_signed(coords : vec3i, level : i32) { let level_idx = u32(level); let num_levels = textureNumLevels(t); let coords_1 = vec3(coords); var predicated_value : vec4; if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) { predicated_value = textureLoad(t, coords_1, level_idx); } let l = predicated_value; } fn load_unsigned(coords : vec3u, level : u32) { let level_idx_1 = u32(level); let num_levels_1 = textureNumLevels(t); let coords_2 = vec3(coords); var predicated_value_1 : vec4; if (((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1))))))) { predicated_value_1 = textureLoad(t, coords_2, level_idx_1); } let l = predicated_value_1; } fn load_mixed(coords : vec3u, level : i32) { let level_idx_2 = u32(level); let num_levels_2 = textureNumLevels(t); let coords_3 = vec3(coords); var predicated_value_2 : vec4; if (((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1))))))) { predicated_value_2 = textureLoad(t, coords_3, level_idx_2); } let l = predicated_value_2; } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec3i(non_uniform.xyz), i32(non_uniform.x)); load_unsigned(vec3u(non_uniform.xyz), u32(non_uniform.x)); load_mixed(vec3u(non_uniform.xyz), i32(non_uniform.x)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureLoad_Multisampled2D) { auto* src = R"( @group(0) @binding(0) var t : texture_multisampled_2d; fn load_signed(coords : vec2i, sample : i32) { let l = textureLoad(t, coords, sample); } fn load_unsigned(coords : vec2u, sample : u32) { let l = textureLoad(t, coords, sample); } fn load_mixed(coords : vec2i, sample : u32) { let l = textureLoad(t, coords, sample); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x)); load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_multisampled_2d; fn load_signed(coords : vec2i, sample : i32) { let l = textureLoad(t, clamp(coords, vec2(0), vec2((textureDimensions(t) - vec2(1)))), sample); } fn load_unsigned(coords : vec2u, sample : u32) { let l = textureLoad(t, min(coords, (textureDimensions(t) - vec2(1))), sample); } fn load_mixed(coords : vec2i, sample : u32) { let l = textureLoad(t, clamp(coords, vec2(0), vec2((textureDimensions(t) - vec2(1)))), sample); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x)); load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_multisampled_2d; fn load_signed(coords : vec2i, sample : i32) { let coords_1 = vec2(coords); var predicated_value : vec4; if (all((coords_1 < textureDimensions(t)))) { predicated_value = textureLoad(t, coords_1, sample); } let l = predicated_value; } fn load_unsigned(coords : vec2u, sample : u32) { let coords_2 = vec2(coords); var predicated_value_1 : vec4; if (all((coords_2 < textureDimensions(t)))) { predicated_value_1 = textureLoad(t, coords_2, sample); } let l = predicated_value_1; } fn load_mixed(coords : vec2i, sample : u32) { let coords_3 = vec2(coords); var predicated_value_2 : vec4; if (all((coords_3 < textureDimensions(t)))) { predicated_value_2 = textureLoad(t, coords_3, sample); } let l = predicated_value_2; } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x)); load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureLoad_Depth2D) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d; fn load_signed(coords : vec2i, level : i32) { let l = textureLoad(t, coords, level); } fn load_unsigned(coords : vec2u, level : u32) { let l = textureLoad(t, coords, level); } fn load_mixed(coords : vec2i, level : u32) { let l = textureLoad(t, coords, level); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x)); load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_depth_2d; fn load_signed(coords : vec2i, level : i32) { let level_idx = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, clamp(coords, vec2(0), vec2((textureDimensions(t, level_idx) - vec2(1)))), level_idx); } fn load_unsigned(coords : vec2u, level : u32) { let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - vec2(1))), level_idx_1); } fn load_mixed(coords : vec2i, level : u32) { let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, clamp(coords, vec2(0), vec2((textureDimensions(t, level_idx_2) - vec2(1)))), level_idx_2); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x)); load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_depth_2d; fn load_signed(coords : vec2i, level : i32) { let level_idx = u32(level); let num_levels = textureNumLevels(t); let coords_1 = vec2(coords); var predicated_value : f32; if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) { predicated_value = textureLoad(t, coords_1, level_idx); } let l = predicated_value; } fn load_unsigned(coords : vec2u, level : u32) { let level_idx_1 = u32(level); let num_levels_1 = textureNumLevels(t); let coords_2 = vec2(coords); var predicated_value_1 : f32; if (((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1))))))) { predicated_value_1 = textureLoad(t, coords_2, level_idx_1); } let l = predicated_value_1; } fn load_mixed(coords : vec2i, level : u32) { let level_idx_2 = u32(level); let num_levels_2 = textureNumLevels(t); let coords_3 = vec2(coords); var predicated_value_2 : f32; if (((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1))))))) { predicated_value_2 = textureLoad(t, coords_3, level_idx_2); } let l = predicated_value_2; } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x)); load_mixed(vec2i(non_uniform.xy), u32(non_uniform.x)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureLoad_Depth2DArray) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d_array; fn load_signed(coords : vec2i, array : i32, level : i32) { let l = textureLoad(t, coords, array, level); } fn load_unsigned(coords : vec2u, array : u32, level : u32) { let l = textureLoad(t, coords, array, level); } fn load_mixed(coords : vec2u, array : i32, level : u32) { let l = textureLoad(t, coords, array, level); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x)); load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_depth_2d_array; fn load_signed(coords : vec2i, array : i32, level : i32) { let level_idx = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, clamp(coords, vec2(0), vec2((textureDimensions(t, level_idx) - vec2(1)))), clamp(array, 0, i32((textureNumLayers(t) - 1))), level_idx); } fn load_unsigned(coords : vec2u, array : u32, level : u32) { let level_idx_1 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_1) - vec2(1))), min(array, (textureNumLayers(t) - 1)), level_idx_1); } fn load_mixed(coords : vec2u, array : i32, level : u32) { let level_idx_2 = min(u32(level), (textureNumLevels(t) - 1)); let l = textureLoad(t, min(coords, (textureDimensions(t, level_idx_2) - vec2(1))), clamp(array, 0, i32((textureNumLayers(t) - 1))), level_idx_2); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x)); load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_depth_2d_array; fn load_signed(coords : vec2i, array : i32, level : i32) { let level_idx = u32(level); let num_levels = textureNumLevels(t); let coords_1 = vec2(coords); let array_idx = u32(array); var predicated_value : f32; if ((((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1)))))) & (array_idx < textureNumLayers(t)))) { predicated_value = textureLoad(t, coords_1, array_idx, level_idx); } let l = predicated_value; } fn load_unsigned(coords : vec2u, array : u32, level : u32) { let level_idx_1 = u32(level); let num_levels_1 = textureNumLevels(t); let coords_2 = vec2(coords); let array_idx_1 = u32(array); var predicated_value_1 : f32; if ((((level_idx_1 < num_levels_1) & all((coords_2 < textureDimensions(t, min(level_idx_1, (num_levels_1 - 1)))))) & (array_idx_1 < textureNumLayers(t)))) { predicated_value_1 = textureLoad(t, coords_2, array_idx_1, level_idx_1); } let l = predicated_value_1; } fn load_mixed(coords : vec2u, array : i32, level : u32) { let level_idx_2 = u32(level); let num_levels_2 = textureNumLevels(t); let coords_3 = vec2(coords); let array_idx_2 = u32(array); var predicated_value_2 : f32; if ((((level_idx_2 < num_levels_2) & all((coords_3 < textureDimensions(t, min(level_idx_2, (num_levels_2 - 1)))))) & (array_idx_2 < textureNumLayers(t)))) { predicated_value_2 = textureLoad(t, coords_3, array_idx_2, level_idx_2); } let l = predicated_value_2; } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy), i32(non_uniform.x), i32(non_uniform.x)); load_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), u32(non_uniform.x)); load_mixed(vec2u(non_uniform.xy), i32(non_uniform.x), u32(non_uniform.x)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureLoad_External) { auto* src = R"( @group(0) @binding(0) var t : texture_external; fn load_signed(coords : vec2i) { let l = textureLoad(t, coords); } fn load_unsigned(coords : vec2u) { let l = textureLoad(t, coords); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy)); load_unsigned(vec2u(non_uniform.xy)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_external; fn load_signed(coords : vec2i) { let l = textureLoad(t, clamp(coords, vec2(0), vec2((textureDimensions(t) - vec2(1))))); } fn load_unsigned(coords : vec2u) { let l = textureLoad(t, min(coords, (textureDimensions(t) - vec2(1)))); } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy)); load_unsigned(vec2u(non_uniform.xy)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_external; fn load_signed(coords : vec2i) { let coords_1 = vec2(coords); var predicated_value : vec4; if (all((coords_1 < textureDimensions(t)))) { predicated_value = textureLoad(t, coords_1); } let l = predicated_value; } fn load_unsigned(coords : vec2u) { let coords_2 = vec2(coords); var predicated_value_1 : vec4; if (all((coords_2 < textureDimensions(t)))) { predicated_value_1 = textureLoad(t, coords_2); } let l = predicated_value_1; } @fragment fn main(@builtin(position) non_uniform : vec4f) { load_signed(vec2i(non_uniform.xy)); load_unsigned(vec2u(non_uniform.xy)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureNumLayers) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d_array; fn num_layers(coords : vec2f, depth_ref : f32) { let l = textureNumLayers(t); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureNumLevels) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d; fn num_levels(coords : vec2f, depth_ref : f32) { let l = textureNumLevels(t); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureNumSamples) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_multisampled_2d; fn num_levels(coords : vec2f, depth_ref : f32) { let l = textureNumSamples(t); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSample) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; @group(0) @binding(1) var s : sampler; fn sample(coords : vec2f) { let l = textureSample(t, s, coords); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSample_Offset) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; @group(0) @binding(1) var s : sampler; fn sample(coords : vec2f) { const offset = vec2i(1); let l = textureSample(t, s, coords, offset); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSample_ArrayIndex) { auto* src = R"( @group(0) @binding(0) var t : texture_2d_array; @group(0) @binding(1) var s : sampler; fn sample_signed(coords : vec2f, array : i32) { let l = textureSample(t, s, coords, array); } fn sample_unsigned(coords : vec2f, array : u32) { let l = textureSample(t, s, coords, array); } @fragment fn main(@builtin(position) non_uniform : vec4f) { sample_signed(non_uniform.xy, i32(non_uniform.x)); sample_unsigned(non_uniform.xy, u32(non_uniform.x)); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleBias) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; @group(0) @binding(1) var s : sampler; fn sample_bias(coords : vec2f, bias : f32) { let l = textureSampleBias(t, s, coords, bias); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleBias_Offset) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; @group(0) @binding(1) var s : sampler; fn sample_bias(coords : vec2f, bias : f32) { const offset = vec2i(1); let l = textureSampleBias(t, s, coords, bias, offset); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleBias_ArrayIndex) { auto* src = R"( @group(0) @binding(0) var t : texture_2d_array; @group(0) @binding(1) var s : sampler; fn sample_bias_signed(coords : vec2f, array : i32, bias : f32) { let l = textureSampleBias(t, s, coords, array, bias); } fn sample_bias_unsigned(coords : vec2f, array : u32, bias : f32) { let l = textureSampleBias(t, s, coords, array, bias); } @fragment fn main(@builtin(position) non_uniform : vec4f) { sample_bias_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.x); sample_bias_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.x); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleCompare) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d; @group(0) @binding(1) var s : sampler_comparison; fn sample_compare(coords : vec2f, depth_ref : f32) { let l = textureSampleCompare(t, s, coords, depth_ref); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleCompare_Offset) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d; @group(0) @binding(1) var s : sampler_comparison; fn sample_compare(coords : vec2f, depth_ref : f32) { const offset = vec2i(1); let l = textureSampleCompare(t, s, coords, depth_ref, offset); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleCompare_ArrayIndex) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d_array; @group(0) @binding(1) var s : sampler_comparison; fn sample_compare_signed(coords : vec2f, array : i32, depth_ref : f32) { let l = textureSampleCompare(t, s, coords, array, depth_ref); } fn sample_compare_unsigned(coords : vec2f, array : u32, depth_ref : f32) { let l = textureSampleCompare(t, s, coords, array, depth_ref); } @fragment fn main(@builtin(position) non_uniform : vec4f) { sample_compare_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.x); sample_compare_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.x); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleCompareLevel) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d; @group(0) @binding(1) var s : sampler_comparison; fn sample_compare_level(coords : vec2f, depth_ref : f32) { let l = textureSampleCompareLevel(t, s, coords, depth_ref); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleCompareLevel_Offset) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d; @group(0) @binding(1) var s : sampler_comparison; fn sample_compare_level(coords : vec2f, depth_ref : f32) { const offset = vec2i(1); let l = textureSampleCompareLevel(t, s, coords, depth_ref, offset); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleCompareLevel_ArrayIndex) { auto* src = R"( @group(0) @binding(0) var t : texture_depth_2d_array; @group(0) @binding(1) var s : sampler_comparison; fn sample_compare_level_signed(coords : vec2f, array : i32, depth_ref : f32) { let l = textureSampleCompareLevel(t, s, coords, array, depth_ref); } fn sample_compare_level_unsigned(coords : vec2f, array : u32, depth_ref : f32) { let l = textureSampleCompareLevel(t, s, coords, array, depth_ref); } @fragment fn main(@builtin(position) non_uniform : vec4f) { sample_compare_level_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.x); sample_compare_level_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.x); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleGrad) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; @group(0) @binding(1) var s : sampler; fn sample_compare_level(coords : vec2f, ddx : vec2f, ddy : vec2f) { let l = textureSampleGrad(t, s, coords, ddx, ddy); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleGrad_Offset) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; @group(0) @binding(1) var s : sampler; fn sample_compare_level(coords : vec2f, ddx : vec2f, ddy : vec2f) { const offset = vec2i(1); let l = textureSampleGrad(t, s, coords, ddx, ddy, offset); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleGrad_ArrayIndex) { auto* src = R"( @group(0) @binding(0) var t : texture_2d_array; @group(0) @binding(1) var s : sampler; fn sample_grad_signed(coords : vec2f, array : i32, ddx : vec2f, ddy : vec2f) { let l = textureSampleGrad(t, s, coords, array, ddx, ddy); } fn sample_grad_unsigned(coords : vec2f, array : u32, ddx : vec2f, ddy : vec2f) { let l = textureSampleGrad(t, s, coords, array, ddx, ddy); } @fragment fn main(@builtin(position) non_uniform : vec4f) { sample_grad_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.xy, non_uniform.xy); sample_grad_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.xy, non_uniform.xy); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleLevel) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; @group(0) @binding(1) var s : sampler; fn sample_compare_level(coords : vec2f, level : f32) { let l = textureSampleLevel(t, s, coords, level); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleLevel_Offset) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; @group(0) @binding(1) var s : sampler; fn sample_compare_level(coords : vec2f, level : f32) { const offset = vec2i(1); let l = textureSampleLevel(t, s, coords, level, offset); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleLevel_ArrayIndex) { auto* src = R"( @group(0) @binding(0) var t : texture_2d_array; @group(0) @binding(1) var s : sampler; fn sample_compare_level_signed(coords : vec2f, array : i32, level : f32) { let l = textureSampleLevel(t, s, coords, array, level); } fn sample_compare_level_unsigned(coords : vec2f, array : u32, level : f32) { let l = textureSampleLevel(t, s, coords, array, level); } @fragment fn main(@builtin(position) non_uniform : vec4f) { sample_compare_level_signed(non_uniform.xy, i32(non_uniform.x), non_uniform.x); sample_compare_level_unsigned(non_uniform.xy, u32(non_uniform.x), non_uniform.x); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureSampleBaseClampToEdge) { auto* src = R"( @group(0) @binding(0) var t : texture_external; @group(0) @binding(1) var s : sampler; fn sample_base_clamp_to_edge(coords : vec2f) { let l = textureSampleBaseClampToEdge(t, s, coords); } )"; auto* expect = src; auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureStore_1D) { auto* src = R"( @group(0) @binding(0) var t : texture_storage_1d; fn store_signed(coords : i32, value : vec4i) { textureStore(t, coords, value); } fn store_unsigned(coords : u32, value : vec4i) { textureStore(t, coords, value); } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(i32(non_uniform.x), vec4i(non_uniform)); store_unsigned(u32(non_uniform.x), vec4i(non_uniform)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_storage_1d; fn store_signed(coords : i32, value : vec4i) { textureStore(t, clamp(coords, 0, i32((textureDimensions(t) - 1))), value); } fn store_unsigned(coords : u32, value : vec4i) { textureStore(t, min(coords, (textureDimensions(t) - 1)), value); } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(i32(non_uniform.x), vec4i(non_uniform)); store_unsigned(u32(non_uniform.x), vec4i(non_uniform)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_storage_1d; fn store_signed(coords : i32, value : vec4i) { let coords_1 = u32(coords); if (all((coords_1 < textureDimensions(t)))) { textureStore(t, coords_1, value); } } fn store_unsigned(coords : u32, value : vec4i) { let coords_2 = u32(coords); if (all((coords_2 < textureDimensions(t)))) { textureStore(t, coords_2, value); } } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(i32(non_uniform.x), vec4i(non_uniform)); store_unsigned(u32(non_uniform.x), vec4i(non_uniform)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureStore_2D) { auto* src = R"( @group(0) @binding(0) var t : texture_storage_2d; fn store_signed(coords : vec2i, value : vec4i) { textureStore(t, coords, value); } fn store_unsigned(coords : vec2u, value : vec4i) { textureStore(t, coords, value); } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(vec2i(non_uniform.xy), vec4i(non_uniform)); store_unsigned(vec2u(non_uniform.xy), vec4i(non_uniform)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_storage_2d; fn store_signed(coords : vec2i, value : vec4i) { textureStore(t, clamp(coords, vec2(0), vec2((textureDimensions(t) - vec2(1)))), value); } fn store_unsigned(coords : vec2u, value : vec4i) { textureStore(t, min(coords, (textureDimensions(t) - vec2(1))), value); } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(vec2i(non_uniform.xy), vec4i(non_uniform)); store_unsigned(vec2u(non_uniform.xy), vec4i(non_uniform)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_storage_2d; fn store_signed(coords : vec2i, value : vec4i) { let coords_1 = vec2(coords); if (all((coords_1 < textureDimensions(t)))) { textureStore(t, coords_1, value); } } fn store_unsigned(coords : vec2u, value : vec4i) { let coords_2 = vec2(coords); if (all((coords_2 < textureDimensions(t)))) { textureStore(t, coords_2, value); } } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(vec2i(non_uniform.xy), vec4i(non_uniform)); store_unsigned(vec2u(non_uniform.xy), vec4i(non_uniform)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureStore_2DArray) { auto* src = R"( @group(0) @binding(0) var t : texture_storage_2d_array; fn store_signed(coords : vec2i, array : i32, value : vec4i) { textureStore(t, coords, array, value); } fn store_unsigned(coords : vec2u, array : u32, value : vec4i) { textureStore(t, coords, array, value); } fn store_mixed(coords : vec2i, array : u32, value : vec4i) { textureStore(t, coords, array, value); } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(vec2i(non_uniform.xy), i32(non_uniform.x), vec4i(non_uniform)); store_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform)); store_mixed(vec2i(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_storage_2d_array; fn store_signed(coords : vec2i, array : i32, value : vec4i) { textureStore(t, clamp(coords, vec2(0), vec2((textureDimensions(t) - vec2(1)))), clamp(array, 0, i32((textureNumLayers(t) - 1))), value); } fn store_unsigned(coords : vec2u, array : u32, value : vec4i) { textureStore(t, min(coords, (textureDimensions(t) - vec2(1))), min(array, (textureNumLayers(t) - 1)), value); } fn store_mixed(coords : vec2i, array : u32, value : vec4i) { textureStore(t, clamp(coords, vec2(0), vec2((textureDimensions(t) - vec2(1)))), min(array, (textureNumLayers(t) - 1)), value); } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(vec2i(non_uniform.xy), i32(non_uniform.x), vec4i(non_uniform)); store_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform)); store_mixed(vec2i(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_storage_2d_array; fn store_signed(coords : vec2i, array : i32, value : vec4i) { let coords_1 = vec2(coords); let array_idx = u32(array); if ((all((coords_1 < textureDimensions(t))) & (array_idx < textureNumLayers(t)))) { textureStore(t, coords_1, array_idx, value); } } fn store_unsigned(coords : vec2u, array : u32, value : vec4i) { let coords_2 = vec2(coords); let array_idx_1 = u32(array); if ((all((coords_2 < textureDimensions(t))) & (array_idx_1 < textureNumLayers(t)))) { textureStore(t, coords_2, array_idx_1, value); } } fn store_mixed(coords : vec2i, array : u32, value : vec4i) { let coords_3 = vec2(coords); let array_idx_2 = u32(array); if ((all((coords_3 < textureDimensions(t))) & (array_idx_2 < textureNumLayers(t)))) { textureStore(t, coords_3, array_idx_2, value); } } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(vec2i(non_uniform.xy), i32(non_uniform.x), vec4i(non_uniform)); store_unsigned(vec2u(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform)); store_mixed(vec2i(non_uniform.xy), u32(non_uniform.x), vec4i(non_uniform)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureStore_3D) { auto* src = R"( @group(0) @binding(0) var t : texture_storage_3d; fn store_signed(coords : vec3i, value : vec4i) { textureStore(t, coords, value); } fn store_unsigned(coords : vec3u, value : vec4i) { textureStore(t, coords, value); } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(vec3i(non_uniform.xyz), vec4i(non_uniform)); store_unsigned(vec3u(non_uniform.xyz), vec4i(non_uniform)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_storage_3d; fn store_signed(coords : vec3i, value : vec4i) { textureStore(t, clamp(coords, vec3(0), vec3((textureDimensions(t) - vec3(1)))), value); } fn store_unsigned(coords : vec3u, value : vec4i) { textureStore(t, min(coords, (textureDimensions(t) - vec3(1))), value); } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(vec3i(non_uniform.xyz), vec4i(non_uniform)); store_unsigned(vec3u(non_uniform.xyz), vec4i(non_uniform)); } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_storage_3d; fn store_signed(coords : vec3i, value : vec4i) { let coords_1 = vec3(coords); if (all((coords_1 < textureDimensions(t)))) { textureStore(t, coords_1, value); } } fn store_unsigned(coords : vec3u, value : vec4i) { let coords_2 = vec3(coords); if (all((coords_2 < textureDimensions(t)))) { textureStore(t, coords_2, value); } } @fragment fn main(@builtin(position) non_uniform : vec4f) { store_signed(vec3i(non_uniform.xyz), vec4i(non_uniform)); store_unsigned(vec3u(non_uniform.xyz), vec4i(non_uniform)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } //////////////////////////////////////////////////////////////////////////////// // Other //////////////////////////////////////////////////////////////////////////////// TEST_P(RobustnessTest, ShadowedVariable) { auto* src = R"( fn f() { var a : array; var i : u32; { var a : array; var b : f32 = a[i]; } var c : f32 = a[i]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( fn f() { var a : array; var i : u32; { var a : array; var b : f32 = a[min(i, 4u)]; } var c : f32 = a[min(i, 2u)]; } )", /* predicate */ R"( fn f() { var a : array; var i : u32; { var a : array; let index = i; let predicate = (u32(index) <= 4u); var predicated_expr : f32; if (predicate) { predicated_expr = a[index]; } var b : f32 = predicated_expr; } let index_1 = i; let predicate_1 = (u32(index_1) <= 2u); var predicated_expr_1 : f32; if (predicate_1) { predicated_expr_1 = a[index_1]; } var c : f32 = predicated_expr_1; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } // Check that existing use of min() and arrayLength() do not get renamed. TEST_P(RobustnessTest, DontRenameSymbols) { auto* src = R"( struct S { a : f32, b : array, } @group(0) @binding(0) var s : S; const c : u32 = 1u; fn f() { let b : f32 = s.b[c]; let x : i32 = min(1, 2); let y : u32 = arrayLength(&(s.b)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( struct S { a : f32, b : array, } @group(0) @binding(0) var s : S; const c : u32 = 1u; fn f() { let b : f32 = s.b[min(c, (arrayLength(&(s.b)) - 1u))]; let x : i32 = min(1, 2); let y : u32 = arrayLength(&(s.b)); } )", /* predicate */ R"( struct S { a : f32, b : array, } @group(0) @binding(0) var s : S; const c : u32 = 1u; fn f() { let index = c; let predicate = (u32(index) <= (arrayLength(&(s.b)) - 1u)); var predicated_expr : f32; if (predicate) { predicated_expr = s.b[index]; } let b : f32 = predicated_expr; let x : i32 = min(1, 2); let y : u32 = arrayLength(&(s.b)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, WorkgroupOverrideCount) { auto* src = R"( override N = 123; var w : array; fn f() { var b : f32 = w[1i]; } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"(error: array size is an override-expression, when expected a constant-expression. Was the SubstituteOverride transform run?)", /* predicate */ R"(error: array size is an override-expression, when expected a constant-expression. Was the SubstituteOverride transform run?)"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } //////////////////////////////////////////////////////////////////////////////// // atomic predication //////////////////////////////////////////////////////////////////////////////// TEST_P(RobustnessTest, AtomicLoad) { auto* src = R"( @group(0) @binding(0) var a : array, 4>; fn f() { let i = 0; let r = atomicLoad(&(a[i])); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var a : array, 4>; fn f() { let i = 0; let r = atomicLoad(&(a[min(u32(i), 3u)])); } )", /* predicate */ R"( @group(0) @binding(0) var a : array, 4>; fn f() { let i = 0; let index = i; let predicate = (u32(index) <= 3u); var predicated_value : i32; if (predicate) { predicated_value = atomicLoad(&(a[index])); } let r = predicated_value; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, AtomicAnd) { auto* src = R"( @group(0) @binding(0) var a : array, 4>; fn f() { let i = 0; atomicAnd(&(a[i]), 10); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var a : array, 4>; fn f() { let i = 0; atomicAnd(&(a[min(u32(i), 3u)]), 10); } )", /* predicate */ R"( @group(0) @binding(0) var a : array, 4>; fn f() { let i = 0; let index = i; let predicate = (u32(index) <= 3u); if (predicate) { atomicAnd(&(a[index]), 10); } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } //////////////////////////////////////////////////////////////////////////////// // workgroupUniformLoad //////////////////////////////////////////////////////////////////////////////// TEST_P(RobustnessTest, WorkgroupUniformLoad) { auto* src = R"( var a : array; @compute @workgroup_size(1) fn f() { let i = 1; let p = &(a[i]); let v = workgroupUniformLoad(p); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( var a : array; @compute @workgroup_size(1) fn f() { let i = 1; let p = &(a[min(u32(i), 31u)]); let v = workgroupUniformLoad(p); } )", /* predicate */ R"( var a : array; @compute @workgroup_size(1) fn f() { let i = 1; let index = i; let predicate = (u32(index) <= 31u); let p = &(a[index]); var predicated_value : u32; if (predicate) { predicated_value = workgroupUniformLoad(p); } else { workgroupBarrier(); } let v = predicated_value; } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } //////////////////////////////////////////////////////////////////////////////// // Usage in loops //////////////////////////////////////////////////////////////////////////////// TEST_P(RobustnessTest, ArrayVal_ForLoopInit) { auto* src = R"( fn f() { let a = array(); var v = 1; for(var i = a[v]; (i < 3); i++) { var in_loop = 42; } } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( fn f() { let a = array(); var v = 1; for(var i = a[min(u32(v), 2u)]; (i < 3); i++) { var in_loop = 42; } } )", /* predicate */ R"( fn f() { let a = array(); var v = 1; { let index = v; let predicate = (u32(index) <= 2u); var predicated_expr : i32; if (predicate) { predicated_expr = a[index]; } var i = predicated_expr; loop { if (!((i < 3))) { break; } { var in_loop = 42; } continuing { i++; } } } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, ArrayVal_ForLoopCond) { auto* src = R"( fn f() { let a = array(); var v = 1; for(var i = 0; (i < a[v]); i++) { var in_loop = 42; } } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( fn f() { let a = array(); var v = 1; for(var i = 0; (i < a[min(u32(v), 2u)]); i++) { var in_loop = 42; } } )", /* predicate */ R"( fn f() { let a = array(); var v = 1; { var i = 0; loop { let index = v; let predicate = (u32(index) <= 2u); var predicated_expr : i32; if (predicate) { predicated_expr = a[index]; } if (!((i < predicated_expr))) { break; } { var in_loop = 42; } continuing { i++; } } } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, ArrayVal_ForLoopCont) { auto* src = R"( fn f() { let a = array(); for(var i = 0; (i < 5); i = a[i]) { var in_loop = 42; } } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( fn f() { let a = array(); for(var i = 0; (i < 5); i = a[min(u32(i), 2u)]) { var in_loop = 42; } } )", /* predicate */ R"( fn f() { let a = array(); { var i = 0; loop { if (!((i < 5))) { break; } { var in_loop = 42; } continuing { let index = i; let predicate = (u32(index) <= 2u); var predicated_expr : i32; if (predicate) { predicated_expr = a[index]; } i = predicated_expr; } } } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureLoad_ForLoopInit) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; fn f() { var coords = vec2(1); var level = 1; for(var i = textureLoad(t, coords, level).x; (i < 3); i++) { var in_loop = 42; } } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_2d; fn f() { var coords = vec2(1); var level = 1; { let level_idx = min(u32(level), (textureNumLevels(t) - 1)); var i = textureLoad(t, clamp(coords, vec2(0), vec2((textureDimensions(t, level_idx) - vec2(1)))), level_idx).x; loop { if (!((i < 3))) { break; } { var in_loop = 42; } continuing { i++; } } } } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_2d; fn f() { var coords = vec2(1); var level = 1; { let level_idx = u32(level); let num_levels = textureNumLevels(t); let coords_1 = vec2(coords); var predicated_value : vec4; if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) { predicated_value = textureLoad(t, coords_1, level_idx); } var i = predicated_value.x; loop { if (!((i < 3))) { break; } { var in_loop = 42; } continuing { i++; } } } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureLoad_ForLoopCond) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; fn f() { var coords = vec2(1); var level = 1; for(var i = 0; (i < textureLoad(t, coords, level).x); i++) { var in_loop = 42; } } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_2d; fn f() { var coords = vec2(1); var level = 1; { var i = 0; loop { let level_idx = min(u32(level), (textureNumLevels(t) - 1)); if (!((i < textureLoad(t, clamp(coords, vec2(0), vec2((textureDimensions(t, level_idx) - vec2(1)))), level_idx).x))) { break; } { var in_loop = 42; } continuing { i++; } } } } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_2d; fn f() { var coords = vec2(1); var level = 1; { var i = 0; loop { let level_idx = u32(level); let num_levels = textureNumLevels(t); let coords_1 = vec2(coords); var predicated_value : vec4; if (((level_idx < num_levels) & all((coords_1 < textureDimensions(t, min(level_idx, (num_levels - 1))))))) { predicated_value = textureLoad(t, coords_1, level_idx); } if (!((i < predicated_value.x))) { break; } { var in_loop = 42; } continuing { i++; } } } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureLoad_ForLoopCont) { auto* src = R"( @group(0) @binding(0) var t : texture_2d; fn f() { var level = 1; for(var i = 0; (i < 5); i = textureLoad(t, vec2(i), i).x) { var in_loop = 42; } } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_2d; fn f() { var level = 1; { var i = 0; loop { if (!((i < 5))) { break; } { var in_loop = 42; } continuing { let level_idx = min(u32(i), (textureNumLevels(t) - 1)); i = textureLoad(t, clamp(vec2(i), vec2(0), vec2((textureDimensions(t, level_idx) - vec2(1)))), level_idx).x; } } } } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_2d; fn f() { var level = 1; { var i = 0; loop { if (!((i < 5))) { break; } { var in_loop = 42; } continuing { let level_idx = u32(i); let num_levels = textureNumLevels(t); let coords = vec2(vec2(i)); var predicated_value : vec4; if (((level_idx < num_levels) & all((coords < textureDimensions(t, min(level_idx, (num_levels - 1))))))) { predicated_value = textureLoad(t, coords, level_idx); } i = predicated_value.x; } } } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureStore_ForLoopInit) { auto* src = R"( @group(0) @binding(0) var t : texture_storage_1d; fn f() { var coords = 1; var value = vec4(1); var i = 0; for(textureStore(t, coords, value); (i < 3); i++) { var in_loop = 42; } } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_storage_1d; fn f() { var coords = 1; var value = vec4(1); var i = 0; for(textureStore(t, clamp(coords, 0, i32((textureDimensions(t) - 1))), value); (i < 3); i++) { var in_loop = 42; } } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_storage_1d; fn f() { var coords = 1; var value = vec4(1); var i = 0; { let coords_1 = u32(coords); if (all((coords_1 < textureDimensions(t)))) { textureStore(t, coords_1, value); } loop { if (!((i < 3))) { break; } { var in_loop = 42; } continuing { i++; } } } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, TextureStore_ForLoopCont) { auto* src = R"( @group(0) @binding(0) var t : texture_storage_1d; fn f() { var level = 1; var value = vec4(1); for(var i = 0; (i < 3); textureStore(t, i, value)) { var in_loop = 42; } } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var t : texture_storage_1d; fn f() { var level = 1; var value = vec4(1); for(var i = 0; (i < 3); textureStore(t, clamp(i, 0, i32((textureDimensions(t) - 1))), value)) { var in_loop = 42; } } )", /* predicate */ R"( @group(0) @binding(0) var t : texture_storage_1d; fn f() { var level = 1; var value = vec4(1); { var i = 0; loop { if (!((i < 3))) { break; } { var in_loop = 42; } continuing { let coords = u32(i); if (all((coords < textureDimensions(t)))) { textureStore(t, coords, value); } } } } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, AtomicXor_ForLoopInit) { auto* src = R"( @group(0) @binding(0) var a : array, 4>; fn f() { var i = 0; for(atomicXor(&(a[i]), 4); (i < 3); i++) { var in_loop = 42; } } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var a : array, 4>; fn f() { var i = 0; for(atomicXor(&(a[min(u32(i), 3u)]), 4); (i < 3); i++) { var in_loop = 42; } } )", /* predicate */ R"( @group(0) @binding(0) var a : array, 4>; fn f() { var i = 0; { let index = i; let predicate = (u32(index) <= 3u); if (predicate) { atomicXor(&(a[index]), 4); } loop { if (!((i < 3))) { break; } { var in_loop = 42; } continuing { i++; } } } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, AtomicXor_ForLoopCond) { auto* src = R"( @group(0) @binding(0) var a : array, 4>; fn f() { for(var i = 0; (i < atomicXor(&(a[i]), 4)); i++) { var in_loop = 42; } } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var a : array, 4>; fn f() { for(var i = 0; (i < atomicXor(&(a[min(u32(i), 3u)]), 4)); i++) { var in_loop = 42; } } )", /* predicate */ R"( @group(0) @binding(0) var a : array, 4>; fn f() { { var i = 0; loop { let index = i; let predicate = (u32(index) <= 3u); var predicated_value : i32; if (predicate) { predicated_value = atomicXor(&(a[index]), 4); } if (!((i < predicated_value))) { break; } { var in_loop = 42; } continuing { i++; } } } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, AtomicXor_ForLoopCont) { auto* src = R"( @group(0) @binding(0) var a : array, 4>; fn f() { for(var i = 0; (i < 5); i = atomicXor(&(a[i]), 4)) { var in_loop = 42; } } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( @group(0) @binding(0) var a : array, 4>; fn f() { for(var i = 0; (i < 5); i = atomicXor(&(a[min(u32(i), 3u)]), 4)) { var in_loop = 42; } } )", /* predicate */ R"( @group(0) @binding(0) var a : array, 4>; fn f() { { var i = 0; loop { if (!((i < 5))) { break; } { var in_loop = 42; } continuing { let index = i; let predicate = (u32(index) <= 3u); var predicated_value : i32; if (predicate) { predicated_value = atomicXor(&(a[index]), 4); } i = predicated_value; } } } } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } //////////////////////////////////////////////////////////////////////////////// // User pointer parameters //////////////////////////////////////////////////////////////////////////////// TEST_P(RobustnessTest, Read_PrivatePointerParameter_IndexWithConstant) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) -> i32 { return ((pre + *(p)) + post); } fn y(pre : i32, p : ptr, post : i32) -> i32 { return x(pre, p, post); } fn z() { y(1, &(a[1]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ src, /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) -> i32 { var predicated_expr : i32; if (p_predicate) { predicated_expr = *(p); } return ((pre + predicated_expr) + post); } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) -> i32 { return x(pre, p, p_predicate_1, post); } fn z() { y(1, &(a[1]), true, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_WorkgroupPointerParameter_IndexWithConstant) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) -> i32 { return ((pre + *(p)) + post); } fn y(pre : i32, p : ptr, post : i32) -> i32 { return x(pre, p, post); } fn z() { y(1, &(a[1]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ src, /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) -> i32 { var predicated_expr : i32; if (p_predicate) { predicated_expr = *(p); } return ((pre + predicated_expr) + post); } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) -> i32 { return x(pre, p, p_predicate_1, post); } fn z() { y(1, &(a[1]), true, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_UniformPointerParameter_IndexWithConstant) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return ((pre + *(p)) + post); } fn y(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return x(pre, p, post); } fn z() { y(vec4(1), &(a[1]), vec4(2)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ src, /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : vec4i, p : ptr, p_predicate : bool, post : vec4i) -> vec4i { var predicated_expr : vec4; if (p_predicate) { predicated_expr = *(p); } return ((pre + predicated_expr) + post); } fn y(pre : vec4i, p : ptr, p_predicate_1 : bool, post : vec4i) -> vec4i { return x(pre, p, p_predicate_1, post); } fn z() { y(vec4(1), &(a[1]), true, vec4(2)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_StoragePointerParameter_IndexWithConstant) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return ((pre + *(p)) + post); } fn y(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return x(pre, p, post); } fn z() { y(vec4(1), &(a[1]), vec4(2)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ src, /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : vec4i, p : ptr, p_predicate : bool, post : vec4i) -> vec4i { var predicated_expr : vec4; if (p_predicate) { predicated_expr = *(p); } return ((pre + predicated_expr) + post); } fn y(pre : vec4i, p : ptr, p_predicate_1 : bool, post : vec4i) -> vec4i { return x(pre, p, p_predicate_1, post); } fn z() { y(vec4(1), &(a[1]), true, vec4(2)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_FunctionPointerParameter_IndexWithConstant) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; fn x(pre : i32, p : ptr, post : i32) -> i32 { return ((pre + *(p)) + post); } fn y(pre : i32, p : ptr, post : i32) -> i32 { return x(pre, p, post); } fn z() { var a : array; y(1, &(a[1]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ src, /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) -> i32 { var predicated_expr : i32; if (p_predicate) { predicated_expr = *(p); } return ((pre + predicated_expr) + post); } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) -> i32 { return x(pre, p, p_predicate_1, post); } fn z() { var a : array; y(1, &(a[1]), true, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_PrivatePointerParameter_IndexWithLet) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) -> i32 { return ((pre + *(p)) + post); } fn y(pre : i32, p : ptr, post : i32) -> i32 { return x(pre, p, post); } fn z() { let i = 0; y(1, &(a[i]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) -> i32 { return ((pre + *(p)) + post); } fn y(pre : i32, p : ptr, post : i32) -> i32 { return x(pre, p, post); } fn z() { let i = 0; y(1, &(a[min(u32(i), 3u)]), 2); } )", /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) -> i32 { var predicated_expr : i32; if (p_predicate) { predicated_expr = *(p); } return ((pre + predicated_expr) + post); } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) -> i32 { return x(pre, p, p_predicate_1, post); } fn z() { let i = 0; let index = i; let predicate = (u32(index) <= 3u); y(1, &(a[index]), predicate, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_WorkgroupPointerParameter_IndexWithLet) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) -> i32 { return ((pre + *(p)) + post); } fn y(pre : i32, p : ptr, post : i32) -> i32 { return x(pre, p, post); } fn z() { let i = 0; y(1, &(a[i]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) -> i32 { return ((pre + *(p)) + post); } fn y(pre : i32, p : ptr, post : i32) -> i32 { return x(pre, p, post); } fn z() { let i = 0; y(1, &(a[min(u32(i), 3u)]), 2); } )", /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) -> i32 { var predicated_expr : i32; if (p_predicate) { predicated_expr = *(p); } return ((pre + predicated_expr) + post); } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) -> i32 { return x(pre, p, p_predicate_1, post); } fn z() { let i = 0; let index = i; let predicate = (u32(index) <= 3u); y(1, &(a[index]), predicate, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_UniformPointerParameter_IndexWithLet) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return ((pre + *(p)) + post); } fn y(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return x(pre, p, post); } fn z() { let i = 0; y(vec4(1), &(a[i]), vec4(2)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return ((pre + *(p)) + post); } fn y(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return x(pre, p, post); } fn z() { let i = 0; y(vec4(1), &(a[min(u32(i), 3u)]), vec4(2)); } )", /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : vec4i, p : ptr, p_predicate : bool, post : vec4i) -> vec4i { var predicated_expr : vec4; if (p_predicate) { predicated_expr = *(p); } return ((pre + predicated_expr) + post); } fn y(pre : vec4i, p : ptr, p_predicate_1 : bool, post : vec4i) -> vec4i { return x(pre, p, p_predicate_1, post); } fn z() { let i = 0; let index = i; let predicate = (u32(index) <= 3u); y(vec4(1), &(a[index]), predicate, vec4(2)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_StoragePointerParameter_IndexWithLet) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return ((pre + *(p)) + post); } fn y(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return x(pre, p, post); } fn z() { let i = 0; y(vec4(1), &(a[i]), vec4(2)); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return ((pre + *(p)) + post); } fn y(pre : vec4i, p : ptr, post : vec4i) -> vec4i { return x(pre, p, post); } fn z() { let i = 0; y(vec4(1), &(a[min(u32(i), 3u)]), vec4(2)); } )", /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : vec4i, p : ptr, p_predicate : bool, post : vec4i) -> vec4i { var predicated_expr : vec4; if (p_predicate) { predicated_expr = *(p); } return ((pre + predicated_expr) + post); } fn y(pre : vec4i, p : ptr, p_predicate_1 : bool, post : vec4i) -> vec4i { return x(pre, p, p_predicate_1, post); } fn z() { let i = 0; let index = i; let predicate = (u32(index) <= 3u); y(vec4(1), &(a[index]), predicate, vec4(2)); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Read_FunctionPointerParameter_IndexWithLet) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; fn x(pre : i32, p : ptr, post : i32) -> i32 { return ((pre + *(p)) + post); } fn y(pre : i32, p : ptr, post : i32) -> i32 { return x(pre, p, post); } fn z() { var a : array; let i = 0; y(1, &(a[i]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( enable chromium_experimental_full_ptr_parameters; fn x(pre : i32, p : ptr, post : i32) -> i32 { return ((pre + *(p)) + post); } fn y(pre : i32, p : ptr, post : i32) -> i32 { return x(pre, p, post); } fn z() { var a : array; let i = 0; y(1, &(a[min(u32(i), 3u)]), 2); } )", /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) -> i32 { var predicated_expr : i32; if (p_predicate) { predicated_expr = *(p); } return ((pre + predicated_expr) + post); } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) -> i32 { return x(pre, p, p_predicate_1, post); } fn z() { var a : array; let i = 0; let index = i; let predicate = (u32(index) <= 3u); y(1, &(a[index]), predicate, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Write_PrivatePointerParameter_IndexWithConstant) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { y(1, &(a[1]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ src, /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) { if (p_predicate) { *(p) = (pre + post); } } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) { x(pre, p, p_predicate_1, post); } fn z() { y(1, &(a[1]), true, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Write_WorkgroupPointerParameter_IndexWithConstant) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { y(1, &(a[1]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ src, /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) { if (p_predicate) { *(p) = (pre + post); } } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) { x(pre, p, p_predicate_1, post); } fn z() { y(1, &(a[1]), true, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Write_StoragePointerParameter_IndexWithConstant) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { y(1, &(a[1]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ src, /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) { if (p_predicate) { *(p) = (pre + post); } } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) { x(pre, p, p_predicate_1, post); } fn z() { y(1, &(a[1]), true, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Write_FunctionPointerParameter_IndexWithConstant) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { var a : array; y(1, &(a[1]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ src, /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) { if (p_predicate) { *(p) = (pre + post); } } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) { x(pre, p, p_predicate_1, post); } fn z() { var a : array; y(1, &(a[1]), true, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Write_PrivatePointerParameter_IndexWithLet) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { let i = 0; y(1, &(a[i]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { let i = 0; y(1, &(a[min(u32(i), 3u)]), 2); } )", /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) { if (p_predicate) { *(p) = (pre + post); } } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) { x(pre, p, p_predicate_1, post); } fn z() { let i = 0; let index = i; let predicate = (u32(index) <= 3u); y(1, &(a[index]), predicate, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Write_WorkgroupPointerParameter_IndexWithLet) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { let i = 0; y(1, &(a[i]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { let i = 0; y(1, &(a[min(u32(i), 3u)]), 2); } )", /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; var a : array; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) { if (p_predicate) { *(p) = (pre + post); } } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) { x(pre, p, p_predicate_1, post); } fn z() { let i = 0; let index = i; let predicate = (u32(index) <= 3u); y(1, &(a[index]), predicate, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Write_StoragePointerParameter_IndexWithLet) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { let i = 0; y(1, &(a[i]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { let i = 0; y(1, &(a[min(u32(i), 3u)]), 2); } )", /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; @group(0) @binding(0) var a : array; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) { if (p_predicate) { *(p) = (pre + post); } } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) { x(pre, p, p_predicate_1, post); } fn z() { let i = 0; let index = i; let predicate = (u32(index) <= 3u); y(1, &(a[index]), predicate, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } TEST_P(RobustnessTest, Write_FunctionPointerParameter_IndexWithLet) { auto* src = R"( enable chromium_experimental_full_ptr_parameters; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { var a : array; let i = 0; y(1, &(a[i]), 2); } )"; auto* expect = Expect(GetParam(), /* ignore */ src, /* clamp */ R"( enable chromium_experimental_full_ptr_parameters; fn x(pre : i32, p : ptr, post : i32) { *(p) = (pre + post); } fn y(pre : i32, p : ptr, post : i32) { x(pre, p, post); } fn z() { var a : array; let i = 0; y(1, &(a[min(u32(i), 3u)]), 2); } )", /* predicate */ R"( enable chromium_experimental_full_ptr_parameters; fn x(pre : i32, p : ptr, p_predicate : bool, post : i32) { if (p_predicate) { *(p) = (pre + post); } } fn y(pre : i32, p : ptr, p_predicate_1 : bool, post : i32) { x(pre, p, p_predicate_1, post); } fn z() { var a : array; let i = 0; let index = i; let predicate = (u32(index) <= 3u); y(1, &(a[index]), predicate, 2); } )"); auto got = Run(src, Config(GetParam())); EXPECT_EQ(expect, str(got)); } INSTANTIATE_TEST_SUITE_P(, RobustnessTest, testing::Values(Robustness::Action::kIgnore, Robustness::Action::kClamp, Robustness::Action::kPredicate)); } // namespace } // namespace tint::transform