| // 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<Robustness::Config>(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 "<invalid action>"; |
| } |
| |
| using RobustnessTest = TransformTestWithParam<Robustness::Action>; |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| // Constant sized array |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithLiteral) { |
| auto* src = R"( |
| fn f() { |
| var b : f32 = array<f32, 3>()[1i]; |
| } |
| )"; |
| |
| auto* expect = src; |
| |
| auto got = Run<Robustness>(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<f32, 3>()[c]; |
| } |
| )"; |
| |
| auto* expect = src; |
| |
| auto got = Run<Robustness>(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<f32, 3>()[l]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| fn f() { |
| let l : u32 = 1u; |
| let b : f32 = array<f32, 3>()[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<f32, 3>()[index]; |
| } |
| let b : f32 = predicated_expr; |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithRuntimeArrayIndex) { |
| auto* src = R"( |
| var<private> i : u32; |
| |
| fn f() { |
| let a = array<f32, 3>(); |
| let b = array<i32, 5>(); |
| var c : f32 = a[b[i]]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> i : u32; |
| |
| fn f() { |
| let a = array<f32, 3>(); |
| let b = array<i32, 5>(); |
| var c : f32 = a[min(u32(b[min(i, 4u)]), 2u)]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> i : u32; |
| |
| fn f() { |
| let a = array<f32, 3>(); |
| let b = array<i32, 5>(); |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_ConstantSizedArrayVal_IndexWithRuntimeExpression) { |
| auto* src = R"( |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = array<f32, 3>()[((c + 2) - 3)]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = array<f32, 3>()[min(u32(((c + 2) - 3)), 2u)]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> c : i32; |
| |
| fn f() { |
| let index = ((c + 2) - 3); |
| let predicate = (u32(index) <= 2u); |
| var predicated_expr : f32; |
| if (predicate) { |
| predicated_expr = array<f32, 3>()[index]; |
| } |
| var b : f32 = predicated_expr; |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_NestedConstantSizedArraysVal_IndexWithRuntimeExpressions) { |
| auto* src = R"( |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> z : i32; |
| |
| fn f() { |
| let a = array<array<array<f32, 1>, 2>, 3>(); |
| var r = a[x][y][z]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> z : i32; |
| |
| fn f() { |
| let a = array<array<array<f32, 1>, 2>, 3>(); |
| var r = a[min(u32(x), 2u)][min(u32(y), 1u)][min(u32(z), 0u)]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> z : i32; |
| |
| fn f() { |
| let a = array<array<array<f32, 1>, 2>, 3>(); |
| let index = x; |
| let predicate = (u32(index) <= 2u); |
| var predicated_expr : array<array<f32, 1u>, 2u>; |
| if (predicate) { |
| predicated_expr = a[index]; |
| } |
| let index_1 = y; |
| let predicate_1 = (u32(index_1) <= 1u); |
| var predicated_expr_1 : array<f32, 1u>; |
| 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<Robustness>(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<f32, 4>(); |
| var b : f32 = a[idx]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| @id(1300) override idx : i32; |
| |
| fn f() { |
| let a = array<f32, 4>(); |
| var b : f32 = a[min(u32(idx), 3u)]; |
| } |
| )", |
| /* predicate */ R"( |
| @id(1300) override idx : i32; |
| |
| fn f() { |
| let a = array<f32, 4>(); |
| 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<Robustness>(src, Config(GetParam())); |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithLiteral) { |
| auto* src = R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| var b : f32 = a[1i]; |
| } |
| )"; |
| |
| auto* expect = src; |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithConst) { |
| auto* src = R"( |
| var<private> a : array<f32, 3>; |
| |
| const c : u32 = 1u; |
| |
| fn f() { |
| let b : f32 = a[c]; |
| } |
| )"; |
| |
| auto* expect = src; |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithLet) { |
| auto* src = R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| let b : f32 = a[l]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| let b : f32 = a[min(l, 2u)]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : array<f32, 3>; |
| |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithRuntimeArrayIndex) { |
| auto* src = R"( |
| var<private> a : array<f32, 3>; |
| |
| var<private> b : array<i32, 5>; |
| |
| var<private> i : u32; |
| |
| fn f() { |
| var c : f32 = a[b[i]]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : array<f32, 3>; |
| |
| var<private> b : array<i32, 5>; |
| |
| var<private> i : u32; |
| |
| fn f() { |
| var c : f32 = a[min(u32(b[min(i, 4u)]), 2u)]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : array<f32, 3>; |
| |
| var<private> b : array<i32, 5>; |
| |
| var<private> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_ConstantSizedArrayRef_IndexWithRuntimeExpression) { |
| auto* src = R"( |
| var<private> a : array<f32, 3>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a[((c + 2) - 3)]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : array<f32, 3>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a[min(u32(((c + 2) - 3)), 2u)]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : array<f32, 3>; |
| |
| var<private> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_NestedConstantSizedArraysRef_IndexWithRuntimeExpressions) { |
| auto* src = R"( |
| var<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> z : i32; |
| |
| fn f() { |
| var r = a[x][y][z]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> z : i32; |
| |
| fn f() { |
| var r = a[min(u32(x), 2u)][min(u32(y), 1u)][min(u32(z), 0u)]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> 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<Robustness>(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<f32, 4>; |
| var b : f32 = a[idx]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| @id(1300) override idx : i32; |
| |
| fn f() { |
| var a : array<f32, 4>; |
| var b : f32 = a[min(u32(idx), 3u)]; |
| } |
| )", |
| /* predicate */ R"( |
| @id(1300) override idx : i32; |
| |
| fn f() { |
| var a : array<f32, 4>; |
| 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<Robustness>(src, Config(GetParam())); |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_ConstantSizedArrayPtr_IndexWithLet) { |
| auto* src = R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| let p = &(a[l]); |
| let f : f32 = *(p); |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| let p = &(a[min(l, 2u)]); |
| let f : f32 = *(p); |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : array<f32, 3>; |
| |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_ConstantSizedArrayPtr_IndexWithRuntimeArrayIndex) { |
| auto* src = R"( |
| var<private> a : array<f32, 3>; |
| |
| var<private> b : array<i32, 5>; |
| |
| var<private> 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<private> a : array<f32, 3>; |
| |
| var<private> b : array<i32, 5>; |
| |
| var<private> 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<private> a : array<f32, 3>; |
| |
| var<private> b : array<i32, 5>; |
| |
| var<private> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_NestedConstantSizedArraysPtr_IndexWithRuntimeExpressions) { |
| auto* src = R"( |
| var<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> 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<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> 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<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_NestedConstantSizedArrays_MixedAccess) { |
| auto* src = R"( |
| var<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> 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<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> 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<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Assign_ConstantSizedArray_IndexWithLet) { |
| auto* src = R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| a[l] = 42.0f; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| a[min(l, 2u)] = 42.0f; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| let index = l; |
| let predicate = (u32(index) <= 2u); |
| if (predicate) { |
| a[index] = 42.0f; |
| } |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Assign_ConstantSizedArrayPtr_IndexWithLet) { |
| auto* src = R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| let p = &(a[l]); |
| *(p) = 42.0f; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| let p = &(a[min(l, 2u)]); |
| *(p) = 42.0f; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : array<f32, 3>; |
| |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Assign_ConstantSizedArrayPtr_IndexWithRuntimeArrayIndex) { |
| auto* src = R"( |
| var<private> a : array<f32, 3>; |
| |
| var<private> b : array<i32, 5>; |
| |
| var<private> 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<private> a : array<f32, 3>; |
| |
| var<private> b : array<i32, 5>; |
| |
| var<private> 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<private> a : array<f32, 3>; |
| |
| var<private> b : array<i32, 5>; |
| |
| var<private> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Assign_NestedConstantSizedArraysPtr_IndexWithRuntimeExpressions) { |
| auto* src = R"( |
| var<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> 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<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> 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<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> x : i32; |
| |
| var<private> y : i32; |
| |
| var<private> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Assign_NestedConstantSizedArrays_MixedAccess) { |
| auto* src = R"( |
| var<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> 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<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> 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<private> a : array<array<array<f32, 1>, 2>, 3>; |
| |
| var<private> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, CompoundAssign_ConstantSizedArray_IndexWithLet) { |
| auto* src = R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| a[l] += 42.0f; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| a[min(l, 2u)] += 42.0f; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : array<f32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| let index = l; |
| let predicate = (u32(index) <= 2u); |
| if (predicate) { |
| a[index] += 42.0f; |
| } |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Increment_ConstantSizedArray_IndexWithLet) { |
| auto* src = R"( |
| var<private> a : array<i32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| a[l]++; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : array<i32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| a[min(l, 2u)]++; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : array<i32, 3>; |
| |
| fn f() { |
| let l : u32 = 1u; |
| let index = l; |
| let predicate = (u32(index) <= 2u); |
| if (predicate) { |
| a[index]++; |
| } |
| } |
| )"); |
| |
| auto got = Run<Robustness>(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<f32>, |
| } |
| |
| @group(0) @binding(0) var<storage, read> s : S; |
| |
| fn f() { |
| var d : f32 = s.b[25]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| struct S { |
| a : f32, |
| b : array<f32>, |
| } |
| |
| @group(0) @binding(0) var<storage, read> s : S; |
| |
| fn f() { |
| var d : f32 = s.b[min(u32(25), (arrayLength(&(s.b)) - 1u))]; |
| } |
| )", |
| /* predicate */ R"( |
| struct S { |
| a : f32, |
| b : array<f32>, |
| } |
| |
| @group(0) @binding(0) var<storage, read> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| // Vector |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_P(RobustnessTest, Read_Vector_IndexWithLiteral) { |
| auto* src = R"( |
| var<private> a : vec3<f32>; |
| |
| fn f() { |
| var b : f32 = a[1i]; |
| } |
| )"; |
| |
| auto* expect = src; |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_Vector_IndexWithConst) { |
| auto* src = R"( |
| var<private> a : vec3<f32>; |
| |
| fn f() { |
| const i = 1; |
| var b : f32 = a[i]; |
| } |
| )"; |
| |
| auto* expect = src; |
| auto got = Run<Robustness>(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<f32>()[i]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| fn f() { |
| let i = 99; |
| let v = vec4<f32>()[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<f32>()[index]; |
| } |
| let v = predicated_expr; |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_Vector_IndexWithRuntimeExpression) { |
| auto* src = R"( |
| var<private> a : vec3<f32>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a[((c + 2) - 3)]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : vec3<f32>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a[min(u32(((c + 2) - 3)), 2u)]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : vec3<f32>; |
| |
| var<private> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_Vector_SwizzleIndexWithGlobalVar) { |
| auto* src = R"( |
| var<private> a : vec3<f32>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a.xy[c]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : vec3<f32>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a.xy[min(u32(c), 1u)]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : vec3<f32>; |
| |
| var<private> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_Vector_SwizzleIndexWithRuntimeExpression) { |
| auto* src = R"( |
| var<private> a : vec3<f32>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a.xy[((c + 2) - 3)]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : vec3<f32>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a.xy[min(u32(((c + 2) - 3)), 1u)]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : vec3<f32>; |
| |
| var<private> 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<Robustness>(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<f32>; |
| var b : f32 = a[idx]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| @id(1300) override idx : i32; |
| |
| fn f() { |
| var a : vec3<f32>; |
| var b : f32 = a[min(u32(idx), 2u)]; |
| } |
| )", |
| /* predicate */ R"( |
| @id(1300) override idx : i32; |
| |
| fn f() { |
| var a : vec3<f32>; |
| 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<Robustness>(src, Config(GetParam())); |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| // Matrix |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_P(RobustnessTest, Read_MatrixRef_IndexingWithLiterals) { |
| auto* src = R"( |
| var<private> a : mat3x2<f32>; |
| |
| fn f() { |
| var b : f32 = a[2i][1i]; |
| } |
| )"; |
| |
| auto* expect = src; |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_MatrixRef_IndexWithRuntimeExpressionThenLiteral) { |
| auto* src = R"( |
| var<private> a : mat3x2<f32>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a[((c + 2) - 3)][1]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : mat3x2<f32>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a[min(u32(((c + 2) - 3)), 2u)][1]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : mat3x2<f32>; |
| |
| var<private> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_MatrixRef_IndexWithLiteralThenRuntimeExpression) { |
| auto* src = R"( |
| var<private> a : mat3x2<f32>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a[1][((c + 2) - 3)]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> a : mat3x2<f32>; |
| |
| var<private> c : i32; |
| |
| fn f() { |
| var b : f32 = a[1][min(u32(((c + 2) - 3)), 1u)]; |
| } |
| )", |
| /* predicate */ R"( |
| var<private> a : mat3x2<f32>; |
| |
| var<private> 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<Robustness>(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<f32>; |
| var b : f32 = a[idx][1]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| @id(1300) override idx : i32; |
| |
| fn f() { |
| var a : mat3x2<f32>; |
| var b : f32 = a[min(u32(idx), 2u)][1]; |
| } |
| )", |
| /* predicate */ R"( |
| @id(1300) override idx : i32; |
| |
| fn f() { |
| var a : mat3x2<f32>; |
| 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<Robustness>(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<f32>(); |
| var v = m[i].yx; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| fn f() { |
| let i = 1; |
| var m = mat3x2<f32>(); |
| var v = m[min(u32(i), 2u)].yx; |
| } |
| )", |
| /* predicate */ R"( |
| fn f() { |
| let i = 1; |
| var m = mat3x2<f32>(); |
| let index = i; |
| let predicate = (u32(index) <= 2u); |
| var predicated_expr : vec2<f32>; |
| if (predicate) { |
| predicated_expr = m[index]; |
| } |
| var v = predicated_expr.yx; |
| } |
| )"); |
| |
| auto got = Run<Robustness>(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<f32>; |
| var b : f32 = a[1][idx]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| @id(1300) override idx : i32; |
| |
| fn f() { |
| var a : mat3x2<f32>; |
| var b : f32 = a[1][min(u32(idx), 1u)]; |
| } |
| )", |
| /* predicate */ R"( |
| @id(1300) override idx : i32; |
| |
| fn f() { |
| var a : mat3x2<f32>; |
| 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<Robustness>(src, Config(GetParam())); |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Assign_Matrix_IndexWithLet) { |
| auto* src = R"( |
| var<private> m : mat3x4f; |
| |
| fn f() { |
| let c = 1; |
| m[c] = vec4f(1); |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> m : mat3x4f; |
| |
| fn f() { |
| let c = 1; |
| m[min(u32(c), 2u)] = vec4f(1); |
| } |
| )", |
| /* predicate */ R"( |
| var<private> m : mat3x4f; |
| |
| fn f() { |
| let c = 1; |
| let index = c; |
| let predicate = (u32(index) <= 2u); |
| if (predicate) { |
| m[index] = vec4f(1); |
| } |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, CompoundAssign_Matrix_IndexWithLet) { |
| auto* src = R"( |
| var<private> m : mat3x4f; |
| |
| fn f() { |
| let c = 1; |
| m[c] += vec4f(1); |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| var<private> m : mat3x4f; |
| |
| fn f() { |
| let c = 1; |
| m[min(u32(c), 2u)] += vec4f(1); |
| } |
| )", |
| /* predicate */ R"( |
| var<private> m : mat3x4f; |
| |
| fn f() { |
| let c = 1; |
| let index = c; |
| let predicate = (u32(index) <= 2u); |
| if (predicate) { |
| m[index] += vec4f(1); |
| } |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| // Texture |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_P(RobustnessTest, TextureDimensions) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| fn dimensions() { |
| let l = textureDimensions(t); |
| } |
| )"; |
| |
| auto* expect = src; |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureDimensions_Level) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| 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<f32>; |
| |
| 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<f32>; |
| |
| fn dimensions_signed(level : i32) { |
| let level_idx = u32(level); |
| let num_levels = textureNumLevels(t); |
| var predicated_value : vec2<u32>; |
| 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<u32>; |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureGather) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| @group(0) @binding(1) var s : sampler; |
| |
| fn gather(coords : vec2f) { |
| let l = textureGather(0, t, s, coords); |
| } |
| )"; |
| |
| auto* expect = src; |
| |
| auto got = Run<Robustness>(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<f32>; |
| |
| @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<Robustness>(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<Robustness>(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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureLoad_1D) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_1d<f32>; |
| |
| 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<f32>; |
| |
| 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<f32>; |
| |
| 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<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 : 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<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 : 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<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(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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureLoad_2D) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| 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<f32>; |
| |
| 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<i32>((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<f32>; |
| |
| fn load_signed(coords : vec2i, level : i32) { |
| let level_idx = u32(level); |
| let num_levels = textureNumLevels(t); |
| let coords_1 = vec2<u32>(coords); |
| var predicated_value : vec4<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<u32>(coords); |
| var predicated_value_1 : vec4<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 : vec2u, level : i32) { |
| let level_idx_2 = u32(level); |
| let num_levels_2 = textureNumLevels(t); |
| let coords_3 = vec2<u32>(coords); |
| var predicated_value_2 : vec4<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(vec2u(non_uniform.xy), i32(non_uniform.x)); |
| } |
| )"); |
| |
| auto got = Run<Robustness>(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<f32>; |
| |
| 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<f32>; |
| |
| 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<i32>((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<f32>; |
| |
| fn load_signed(coords : vec2i, array : i32, level : i32) { |
| let level_idx = u32(level); |
| let num_levels = textureNumLevels(t); |
| let coords_1 = vec2<u32>(coords); |
| let array_idx = u32(array); |
| var predicated_value : vec4<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<u32>(coords); |
| let array_idx_1 = u32(array); |
| var predicated_value_1 : vec4<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<u32>(coords); |
| let array_idx_2 = u32(array); |
| var predicated_value_2 : vec4<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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureLoad_3D) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_3d<f32>; |
| |
| 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<f32>; |
| |
| 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<i32>((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<f32>; |
| |
| fn load_signed(coords : vec3i, level : i32) { |
| let level_idx = u32(level); |
| let num_levels = textureNumLevels(t); |
| let coords_1 = vec3<u32>(coords); |
| var predicated_value : vec4<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 : vec3u, level : u32) { |
| let level_idx_1 = u32(level); |
| let num_levels_1 = textureNumLevels(t); |
| let coords_2 = vec3<u32>(coords); |
| var predicated_value_1 : vec4<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 : vec3u, level : i32) { |
| let level_idx_2 = u32(level); |
| let num_levels_2 = textureNumLevels(t); |
| let coords_3 = vec3<u32>(coords); |
| var predicated_value_2 : vec4<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(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<Robustness>(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<f32>; |
| |
| 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<f32>; |
| |
| fn load_signed(coords : vec2i, sample : i32) { |
| let l = textureLoad(t, clamp(coords, vec2(0), vec2<i32>((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<i32>((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<f32>; |
| |
| fn load_signed(coords : vec2i, sample : i32) { |
| let coords_1 = vec2<u32>(coords); |
| var predicated_value : vec4<f32>; |
| 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<u32>(coords); |
| var predicated_value_1 : vec4<f32>; |
| 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<u32>(coords); |
| var predicated_value_2 : vec4<f32>; |
| 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<Robustness>(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<i32>((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<i32>((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<u32>(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<u32>(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<u32>(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<Robustness>(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<i32>((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<u32>(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<u32>(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<u32>(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<Robustness>(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<i32>((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<u32>(coords); |
| var predicated_value : vec4<f32>; |
| 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<u32>(coords); |
| var predicated_value_1 : vec4<f32>; |
| 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<Robustness>(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<Robustness>(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<Robustness>(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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureSample) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| @group(0) @binding(1) var s : sampler; |
| |
| fn sample(coords : vec2f) { |
| let l = textureSample(t, s, coords); |
| } |
| )"; |
| |
| auto* expect = src; |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureSample_Offset) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| @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<Robustness>(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<f32>; |
| |
| @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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureSampleBias) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| @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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureSampleBias_Offset) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| @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<Robustness>(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<f32>; |
| |
| @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<Robustness>(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<Robustness>(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<Robustness>(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<Robustness>(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<Robustness>(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<Robustness>(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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureSampleGrad) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| @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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureSampleGrad_Offset) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| @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<Robustness>(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<f32>; |
| |
| @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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureSampleLevel) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| @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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureSampleLevel_Offset) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<f32>; |
| |
| @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<Robustness>(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<f32>; |
| |
| @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<Robustness>(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<Robustness>(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<rgba8sint, write>; |
| |
| 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<rgba8sint, write>; |
| |
| 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<rgba8sint, write>; |
| |
| 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<Robustness>(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<rgba8sint, write>; |
| |
| 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<rgba8sint, write>; |
| |
| fn store_signed(coords : vec2i, value : vec4i) { |
| textureStore(t, clamp(coords, vec2(0), vec2<i32>((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<rgba8sint, write>; |
| |
| fn store_signed(coords : vec2i, value : vec4i) { |
| let coords_1 = vec2<u32>(coords); |
| if (all((coords_1 < textureDimensions(t)))) { |
| textureStore(t, coords_1, value); |
| } |
| } |
| |
| fn store_unsigned(coords : vec2u, value : vec4i) { |
| let coords_2 = vec2<u32>(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<Robustness>(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<rgba8sint, write>; |
| |
| 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<rgba8sint, write>; |
| |
| fn store_signed(coords : vec2i, array : i32, value : vec4i) { |
| textureStore(t, clamp(coords, vec2(0), vec2<i32>((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<i32>((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<rgba8sint, write>; |
| |
| fn store_signed(coords : vec2i, array : i32, value : vec4i) { |
| let coords_1 = vec2<u32>(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<u32>(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<u32>(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<Robustness>(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<rgba8sint, write>; |
| |
| 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<rgba8sint, write>; |
| |
| fn store_signed(coords : vec3i, value : vec4i) { |
| textureStore(t, clamp(coords, vec3(0), vec3<i32>((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<rgba8sint, write>; |
| |
| fn store_signed(coords : vec3i, value : vec4i) { |
| let coords_1 = vec3<u32>(coords); |
| if (all((coords_1 < textureDimensions(t)))) { |
| textureStore(t, coords_1, value); |
| } |
| } |
| |
| fn store_unsigned(coords : vec3u, value : vec4i) { |
| let coords_2 = vec3<u32>(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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| // Other |
| //////////////////////////////////////////////////////////////////////////////// |
| TEST_P(RobustnessTest, ShadowedVariable) { |
| auto* src = R"( |
| fn f() { |
| var a : array<f32, 3>; |
| var i : u32; |
| { |
| var a : array<f32, 5>; |
| var b : f32 = a[i]; |
| } |
| var c : f32 = a[i]; |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| fn f() { |
| var a : array<f32, 3>; |
| var i : u32; |
| { |
| var a : array<f32, 5>; |
| var b : f32 = a[min(i, 4u)]; |
| } |
| var c : f32 = a[min(i, 2u)]; |
| } |
| )", |
| /* predicate */ R"( |
| fn f() { |
| var a : array<f32, 3>; |
| var i : u32; |
| { |
| var a : array<f32, 5>; |
| 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<Robustness>(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<f32>, |
| } |
| |
| @group(0) @binding(0) var<storage, read> 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<f32>, |
| } |
| |
| @group(0) @binding(0) var<storage, read> 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<f32>, |
| } |
| |
| @group(0) @binding(0) var<storage, read> 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, WorkgroupOverrideCount) { |
| auto* src = R"( |
| override N = 123; |
| |
| var<workgroup> w : array<f32, N>; |
| |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| // atomic predication |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_P(RobustnessTest, AtomicLoad) { |
| auto* src = R"( |
| @group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>; |
| |
| fn f() { |
| let i = 0; |
| let r = atomicLoad(&(a[i])); |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| @group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>; |
| |
| fn f() { |
| let i = 0; |
| let r = atomicLoad(&(a[min(u32(i), 3u)])); |
| } |
| )", |
| /* predicate */ R"( |
| @group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, AtomicAnd) { |
| auto* src = R"( |
| @group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>; |
| |
| fn f() { |
| let i = 0; |
| atomicAnd(&(a[i]), 10); |
| } |
| )"; |
| |
| auto* expect = Expect(GetParam(), |
| /* ignore */ src, |
| /* clamp */ R"( |
| @group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>; |
| |
| fn f() { |
| let i = 0; |
| atomicAnd(&(a[min(u32(i), 3u)]), 10); |
| } |
| )", |
| /* predicate */ R"( |
| @group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 4>; |
| |
| fn f() { |
| let i = 0; |
| let index = i; |
| let predicate = (u32(index) <= 3u); |
| if (predicate) { |
| atomicAnd(&(a[index]), 10); |
| } |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| // workgroupUniformLoad |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_P(RobustnessTest, WorkgroupUniformLoad) { |
| auto* src = R"( |
| var<workgroup> a : array<u32, 32>; |
| |
| @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<workgroup> a : array<u32, 32>; |
| |
| @compute @workgroup_size(1) |
| fn f() { |
| let i = 1; |
| let p = &(a[min(u32(i), 31u)]); |
| let v = workgroupUniformLoad(p); |
| } |
| )", |
| /* predicate */ R"( |
| var<workgroup> a : array<u32, 32>; |
| |
| @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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| //////////////////////////////////////////////////////////////////////////////// |
| // Usage in loops |
| //////////////////////////////////////////////////////////////////////////////// |
| |
| TEST_P(RobustnessTest, ArrayVal_ForLoopInit) { |
| auto* src = R"( |
| fn f() { |
| let a = array<i32, 3>(); |
| 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<i32, 3>(); |
| 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<i32, 3>(); |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, ArrayVal_ForLoopCond) { |
| auto* src = R"( |
| fn f() { |
| let a = array<i32, 3>(); |
| 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<i32, 3>(); |
| 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<i32, 3>(); |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, ArrayVal_ForLoopCont) { |
| auto* src = R"( |
| fn f() { |
| let a = array<i32, 3>(); |
| 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<i32, 3>(); |
| for(var i = 0; (i < 5); i = a[min(u32(i), 2u)]) { |
| var in_loop = 42; |
| } |
| } |
| )", |
| /* predicate */ R"( |
| fn f() { |
| let a = array<i32, 3>(); |
| { |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureLoad_ForLoopInit) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<i32>; |
| |
| 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<i32>; |
| |
| 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<i32>((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<i32>; |
| |
| fn f() { |
| var coords = vec2(1); |
| var level = 1; |
| { |
| let level_idx = u32(level); |
| let num_levels = textureNumLevels(t); |
| let coords_1 = vec2<u32>(coords); |
| var predicated_value : vec4<i32>; |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureLoad_ForLoopCond) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<i32>; |
| |
| 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<i32>; |
| |
| 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<i32>((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<i32>; |
| |
| 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<u32>(coords); |
| var predicated_value : vec4<i32>; |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, TextureLoad_ForLoopCont) { |
| auto* src = R"( |
| @group(0) @binding(0) var t : texture_2d<i32>; |
| |
| 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<i32>; |
| |
| 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<i32>((textureDimensions(t, level_idx) - vec2(1)))), level_idx).x; |
| } |
| } |
| } |
| } |
| )", |
| /* predicate */ R"( |
| @group(0) @binding(0) var t : texture_2d<i32>; |
| |
| 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<u32>(vec2(i)); |
| var predicated_value : vec4<i32>; |
| 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<Robustness>(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<rgba8sint, write>; |
| |
| 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<rgba8sint, write>; |
| |
| 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<rgba8sint, write>; |
| |
| 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<Robustness>(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<rgba8sint, write>; |
| |
| 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<rgba8sint, write>; |
| |
| 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<rgba8sint, write>; |
| |
| 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, AtomicXor_ForLoopInit) { |
| auto* src = R"( |
| @group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 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<storage, read_write> a : array<atomic<i32>, 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<storage, read_write> a : array<atomic<i32>, 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, AtomicXor_ForLoopCond) { |
| auto* src = R"( |
| @group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 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<storage, read_write> a : array<atomic<i32>, 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<storage, read_write> a : array<atomic<i32>, 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, AtomicXor_ForLoopCont) { |
| auto* src = R"( |
| @group(0) @binding(0) var<storage, read_write> a : array<atomic<i32>, 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<storage, read_write> a : array<atomic<i32>, 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<storage, read_write> a : array<atomic<i32>, 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<Robustness>(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<private> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<private, i32>, post : i32) -> i32 { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : i32, p : ptr<private, i32>, 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<private> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<private, i32>, 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<private, i32>, 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_WorkgroupPointerParameter_IndexWithConstant) { |
| auto* src = R"( |
| enable chromium_experimental_full_ptr_parameters; |
| |
| var<workgroup> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : i32, p : ptr<workgroup, i32>, 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<workgroup> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<workgroup, i32>, 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<workgroup, i32>, 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<Robustness>(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<uniform> a : array<vec4i, 4>; |
| |
| fn x(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : vec4i, p : ptr<uniform, vec4i>, 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<uniform> a : array<vec4i, 4>; |
| |
| fn x(pre : vec4i, p : ptr<uniform, vec4i>, p_predicate : bool, post : vec4i) -> vec4i { |
| var predicated_expr : vec4<i32>; |
| if (p_predicate) { |
| predicated_expr = *(p); |
| } |
| return ((pre + predicated_expr) + post); |
| } |
| |
| fn y(pre : vec4i, p : ptr<uniform, vec4i>, 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<Robustness>(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<storage> a : array<vec4i, 4>; |
| |
| fn x(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : vec4i, p : ptr<storage, vec4i>, 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<storage> a : array<vec4i, 4>; |
| |
| fn x(pre : vec4i, p : ptr<storage, vec4i>, p_predicate : bool, post : vec4i) -> vec4i { |
| var predicated_expr : vec4<i32>; |
| if (p_predicate) { |
| predicated_expr = *(p); |
| } |
| return ((pre + predicated_expr) + post); |
| } |
| |
| fn y(pre : vec4i, p : ptr<storage, vec4i>, 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<Robustness>(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<function, i32>, post : i32) -> i32 { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : i32, p : ptr<function, i32>, post : i32) -> i32 { |
| return x(pre, p, post); |
| } |
| |
| fn z() { |
| var a : array<i32, 4>; |
| 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<function, i32>, 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<function, i32>, p_predicate_1 : bool, post : i32) -> i32 { |
| return x(pre, p, p_predicate_1, post); |
| } |
| |
| fn z() { |
| var a : array<i32, 4>; |
| y(1, &(a[1]), true, 2); |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_PrivatePointerParameter_IndexWithLet) { |
| auto* src = R"( |
| enable chromium_experimental_full_ptr_parameters; |
| |
| var<private> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<private, i32>, post : i32) -> i32 { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : i32, p : ptr<private, i32>, 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<private> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<private, i32>, post : i32) -> i32 { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : i32, p : ptr<private, i32>, 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<private> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<private, i32>, 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<private, i32>, 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Read_WorkgroupPointerParameter_IndexWithLet) { |
| auto* src = R"( |
| enable chromium_experimental_full_ptr_parameters; |
| |
| var<workgroup> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : i32, p : ptr<workgroup, i32>, 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<workgroup> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) -> i32 { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : i32, p : ptr<workgroup, i32>, 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<workgroup> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<workgroup, i32>, 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<workgroup, i32>, 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<Robustness>(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<uniform> a : array<vec4i, 4>; |
| |
| fn x(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : vec4i, p : ptr<uniform, vec4i>, 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<uniform> a : array<vec4i, 4>; |
| |
| fn x(pre : vec4i, p : ptr<uniform, vec4i>, post : vec4i) -> vec4i { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : vec4i, p : ptr<uniform, vec4i>, 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<uniform> a : array<vec4i, 4>; |
| |
| fn x(pre : vec4i, p : ptr<uniform, vec4i>, p_predicate : bool, post : vec4i) -> vec4i { |
| var predicated_expr : vec4<i32>; |
| if (p_predicate) { |
| predicated_expr = *(p); |
| } |
| return ((pre + predicated_expr) + post); |
| } |
| |
| fn y(pre : vec4i, p : ptr<uniform, vec4i>, 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<Robustness>(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<storage> a : array<vec4i, 4>; |
| |
| fn x(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : vec4i, p : ptr<storage, vec4i>, 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<storage> a : array<vec4i, 4>; |
| |
| fn x(pre : vec4i, p : ptr<storage, vec4i>, post : vec4i) -> vec4i { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : vec4i, p : ptr<storage, vec4i>, 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<storage> a : array<vec4i, 4>; |
| |
| fn x(pre : vec4i, p : ptr<storage, vec4i>, p_predicate : bool, post : vec4i) -> vec4i { |
| var predicated_expr : vec4<i32>; |
| if (p_predicate) { |
| predicated_expr = *(p); |
| } |
| return ((pre + predicated_expr) + post); |
| } |
| |
| fn y(pre : vec4i, p : ptr<storage, vec4i>, 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<Robustness>(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<function, i32>, post : i32) -> i32 { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : i32, p : ptr<function, i32>, post : i32) -> i32 { |
| return x(pre, p, post); |
| } |
| |
| fn z() { |
| var a : array<i32, 4>; |
| 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<function, i32>, post : i32) -> i32 { |
| return ((pre + *(p)) + post); |
| } |
| |
| fn y(pre : i32, p : ptr<function, i32>, post : i32) -> i32 { |
| return x(pre, p, post); |
| } |
| |
| fn z() { |
| var a : array<i32, 4>; |
| 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<function, i32>, 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<function, i32>, p_predicate_1 : bool, post : i32) -> i32 { |
| return x(pre, p, p_predicate_1, post); |
| } |
| |
| fn z() { |
| var a : array<i32, 4>; |
| let i = 0; |
| let index = i; |
| let predicate = (u32(index) <= 3u); |
| y(1, &(a[index]), predicate, 2); |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Write_PrivatePointerParameter_IndexWithConstant) { |
| auto* src = R"( |
| enable chromium_experimental_full_ptr_parameters; |
| |
| var<private> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<private, i32>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<private, i32>, 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<private> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<private, i32>, p_predicate : bool, post : i32) { |
| if (p_predicate) { |
| *(p) = (pre + post); |
| } |
| } |
| |
| fn y(pre : i32, p : ptr<private, i32>, p_predicate_1 : bool, post : i32) { |
| x(pre, p, p_predicate_1, post); |
| } |
| |
| fn z() { |
| y(1, &(a[1]), true, 2); |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Write_WorkgroupPointerParameter_IndexWithConstant) { |
| auto* src = R"( |
| enable chromium_experimental_full_ptr_parameters; |
| |
| var<workgroup> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<workgroup, i32>, 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<workgroup> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<workgroup, i32>, p_predicate : bool, post : i32) { |
| if (p_predicate) { |
| *(p) = (pre + post); |
| } |
| } |
| |
| fn y(pre : i32, p : ptr<workgroup, i32>, p_predicate_1 : bool, post : i32) { |
| x(pre, p, p_predicate_1, post); |
| } |
| |
| fn z() { |
| y(1, &(a[1]), true, 2); |
| } |
| )"); |
| |
| auto got = Run<Robustness>(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<storage, read_write> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<storage, i32, read_write>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<storage, i32, read_write>, 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<storage, read_write> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<storage, i32, read_write>, p_predicate : bool, post : i32) { |
| if (p_predicate) { |
| *(p) = (pre + post); |
| } |
| } |
| |
| fn y(pre : i32, p : ptr<storage, i32, read_write>, p_predicate_1 : bool, post : i32) { |
| x(pre, p, p_predicate_1, post); |
| } |
| |
| fn z() { |
| y(1, &(a[1]), true, 2); |
| } |
| )"); |
| |
| auto got = Run<Robustness>(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<function, i32>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<function, i32>, post : i32) { |
| x(pre, p, post); |
| } |
| |
| fn z() { |
| var a : array<i32, 4>; |
| 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<function, i32>, p_predicate : bool, post : i32) { |
| if (p_predicate) { |
| *(p) = (pre + post); |
| } |
| } |
| |
| fn y(pre : i32, p : ptr<function, i32>, p_predicate_1 : bool, post : i32) { |
| x(pre, p, p_predicate_1, post); |
| } |
| |
| fn z() { |
| var a : array<i32, 4>; |
| y(1, &(a[1]), true, 2); |
| } |
| )"); |
| |
| auto got = Run<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Write_PrivatePointerParameter_IndexWithLet) { |
| auto* src = R"( |
| enable chromium_experimental_full_ptr_parameters; |
| |
| var<private> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<private, i32>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<private, i32>, 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<private> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<private, i32>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<private, i32>, 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<private> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<private, i32>, p_predicate : bool, post : i32) { |
| if (p_predicate) { |
| *(p) = (pre + post); |
| } |
| } |
| |
| fn y(pre : i32, p : ptr<private, i32>, 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<Robustness>(src, Config(GetParam())); |
| |
| EXPECT_EQ(expect, str(got)); |
| } |
| |
| TEST_P(RobustnessTest, Write_WorkgroupPointerParameter_IndexWithLet) { |
| auto* src = R"( |
| enable chromium_experimental_full_ptr_parameters; |
| |
| var<workgroup> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<workgroup, i32>, 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<workgroup> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<workgroup, i32>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<workgroup, i32>, 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<workgroup> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<workgroup, i32>, p_predicate : bool, post : i32) { |
| if (p_predicate) { |
| *(p) = (pre + post); |
| } |
| } |
| |
| fn y(pre : i32, p : ptr<workgroup, i32>, 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<Robustness>(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<storage, read_write> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<storage, i32, read_write>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<storage, i32, read_write>, 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<storage, read_write> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<storage, i32, read_write>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<storage, i32, read_write>, 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<storage, read_write> a : array<i32, 4>; |
| |
| fn x(pre : i32, p : ptr<storage, i32, read_write>, p_predicate : bool, post : i32) { |
| if (p_predicate) { |
| *(p) = (pre + post); |
| } |
| } |
| |
| fn y(pre : i32, p : ptr<storage, i32, read_write>, 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<Robustness>(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<function, i32>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<function, i32>, post : i32) { |
| x(pre, p, post); |
| } |
| |
| fn z() { |
| var a : array<i32, 4>; |
| 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<function, i32>, post : i32) { |
| *(p) = (pre + post); |
| } |
| |
| fn y(pre : i32, p : ptr<function, i32>, post : i32) { |
| x(pre, p, post); |
| } |
| |
| fn z() { |
| var a : array<i32, 4>; |
| 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<function, i32>, p_predicate : bool, post : i32) { |
| if (p_predicate) { |
| *(p) = (pre + post); |
| } |
| } |
| |
| fn y(pre : i32, p : ptr<function, i32>, p_predicate_1 : bool, post : i32) { |
| x(pre, p, p_predicate_1, post); |
| } |
| |
| fn z() { |
| var a : array<i32, 4>; |
| let i = 0; |
| let index = i; |
| let predicate = (u32(index) <= 3u); |
| y(1, &(a[index]), predicate, 2); |
| } |
| )"); |
| |
| auto got = Run<Robustness>(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 |