| // Copyright 2024 The Dawn & Tint Authors |
| // |
| // Redistribution and use in source and binary forms, with or without |
| // modification, are permitted provided that the following conditions are met: |
| // |
| // 1. Redistributions of source code must retain the above copyright notice, this |
| // list of conditions and the following disclaimer. |
| // |
| // 2. Redistributions in binary form must reproduce the above copyright notice, |
| // this list of conditions and the following disclaimer in the documentation |
| // and/or other materials provided with the distribution. |
| // |
| // 3. Neither the name of the copyright holder nor the names of its |
| // contributors may be used to endorse or promote products derived from |
| // this software without specific prior written permission. |
| // |
| // THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" |
| // AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE |
| // IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE |
| // DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE |
| // FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL |
| // DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR |
| // SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER |
| // CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, |
| // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE |
| // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. |
| |
| #include "src/tint/lang/msl/writer/raise/module_scope_vars.h" |
| |
| #include <utility> |
| |
| #include "src/tint/lang/core/fluent_types.h" |
| #include "src/tint/lang/core/ir/transform/helper_test.h" |
| #include "src/tint/lang/core/type/sampled_texture.h" |
| |
| using namespace tint::core::fluent_types; // NOLINT |
| using namespace tint::core::number_suffixes; // NOLINT |
| |
| namespace tint::msl::writer::raise { |
| namespace { |
| |
| using MslWriter_ModuleScopeVarsTest = core::ir::transform::TransformTest; |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, NoModuleScopeVars) { |
| auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| auto* var = b.Var<function, i32>("v"); |
| b.Load(var); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| %foo = @fragment func():void { |
| $B1: { |
| %v:ptr<function, i32, read_write> = var |
| %3:i32 = load %v |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = src; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, Private) { |
| auto* var_a = b.Var("a", ty.ptr<private_, i32>()); |
| auto* var_b = b.Var("b", ty.ptr<private_, i32>()); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_a, b.Add<i32>(load_a, load_b)); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<private, i32, read_write> = var |
| %b:ptr<private, i32, read_write> = var |
| } |
| |
| %foo = @fragment func():void { |
| $B2: { |
| %4:i32 = load %a |
| %5:i32 = load %b |
| %6:i32 = add %4, %5 |
| store %a, %6 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<private, i32, read_write> @offset(0) |
| b:ptr<private, i32, read_write> @offset(0) |
| } |
| |
| %foo = @fragment func():void { |
| $B1: { |
| %a:ptr<private, i32, read_write> = var |
| %b:ptr<private, i32, read_write> = var |
| %4:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars:tint_module_vars_struct = let %4 |
| %6:ptr<private, i32, read_write> = access %tint_module_vars, 0u |
| %7:i32 = load %6 |
| %8:ptr<private, i32, read_write> = access %tint_module_vars, 1u |
| %9:i32 = load %8 |
| %10:i32 = add %7, %9 |
| %11:ptr<private, i32, read_write> = access %tint_module_vars, 0u |
| store %11, %10 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, Private_WithInitializers) { |
| auto* var_a = b.Var<private_>("a", 42_i); |
| auto* var_b = b.Var<private_>("b", -1_i); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_a, b.Add<i32>(load_a, load_b)); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<private, i32, read_write> = var, 42i |
| %b:ptr<private, i32, read_write> = var, -1i |
| } |
| |
| %foo = @fragment func():void { |
| $B2: { |
| %4:i32 = load %a |
| %5:i32 = load %b |
| %6:i32 = add %4, %5 |
| store %a, %6 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<private, i32, read_write> @offset(0) |
| b:ptr<private, i32, read_write> @offset(0) |
| } |
| |
| %foo = @fragment func():void { |
| $B1: { |
| %a:ptr<private, i32, read_write> = var, 42i |
| %b:ptr<private, i32, read_write> = var, -1i |
| %4:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars:tint_module_vars_struct = let %4 |
| %6:ptr<private, i32, read_write> = access %tint_module_vars, 0u |
| %7:i32 = load %6 |
| %8:ptr<private, i32, read_write> = access %tint_module_vars, 1u |
| %9:i32 = load %8 |
| %10:i32 = add %7, %9 |
| %11:ptr<private, i32, read_write> = access %tint_module_vars, 0u |
| store %11, %10 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, Storage) { |
| auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_b, b.Add<i32>(load_a, load_b)); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<storage, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| } |
| |
| %foo = @fragment func():void { |
| $B2: { |
| %4:i32 = load %a |
| %5:i32 = load %b |
| %6:i32 = add %4, %5 |
| store %b, %6 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<storage, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| } |
| |
| %foo = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B1: { |
| %4:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars:tint_module_vars_struct = let %4 |
| %6:ptr<storage, i32, read> = access %tint_module_vars, 0u |
| %7:i32 = load %6 |
| %8:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| %9:i32 = load %8 |
| %10:i32 = add %7, %9 |
| %11:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| store %11, %10 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, Uniform) { |
| auto* var_a = b.Var("a", ty.ptr<uniform, i32>()); |
| auto* var_b = b.Var("b", ty.ptr<uniform, i32>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Add<i32>(load_a, load_b); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<uniform, i32, read> = var @binding_point(1, 2) |
| %b:ptr<uniform, i32, read> = var @binding_point(3, 4) |
| } |
| |
| %foo = @fragment func():void { |
| $B2: { |
| %4:i32 = load %a |
| %5:i32 = load %b |
| %6:i32 = add %4, %5 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<uniform, i32, read> @offset(0) |
| b:ptr<uniform, i32, read> @offset(0) |
| } |
| |
| %foo = @fragment func(%a:ptr<uniform, i32, read> [@binding_point(1, 2)], %b:ptr<uniform, i32, read> [@binding_point(3, 4)]):void { |
| $B1: { |
| %4:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars:tint_module_vars_struct = let %4 |
| %6:ptr<uniform, i32, read> = access %tint_module_vars, 0u |
| %7:i32 = load %6 |
| %8:ptr<uniform, i32, read> = access %tint_module_vars, 1u |
| %9:i32 = load %8 |
| %10:i32 = add %7, %9 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, HandleTypes) { |
| auto* t = ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32()); |
| auto* var_t = b.Var("t", ty.ptr<handle>(t)); |
| auto* var_s = b.Var("s", ty.ptr<handle>(ty.sampler())); |
| var_t->SetBindingPoint(1, 2); |
| var_s->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_t); |
| mod.root_block->Append(var_s); |
| |
| auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| auto* load_t = b.Load(var_t); |
| auto* load_s = b.Load(var_s); |
| b.Call<vec4<f32>>(core::BuiltinFn::kTextureSample, load_t, load_s, b.Splat<vec2<f32>>(0_f)); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %t:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 2) |
| %s:ptr<handle, sampler, read> = var @binding_point(3, 4) |
| } |
| |
| %foo = @fragment func():void { |
| $B2: { |
| %4:texture_2d<f32> = load %t |
| %5:sampler = load %s |
| %6:vec4<f32> = textureSample %4, %5, vec2<f32>(0.0f) |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| t:texture_2d<f32> @offset(0) |
| s:sampler @offset(0) |
| } |
| |
| %foo = @fragment func(%t:texture_2d<f32> [@binding_point(1, 2)], %s:sampler [@binding_point(3, 4)]):void { |
| $B1: { |
| %4:tint_module_vars_struct = construct %t, %s |
| %tint_module_vars:tint_module_vars_struct = let %4 |
| %6:texture_2d<f32> = access %tint_module_vars, 0u |
| %7:sampler = access %tint_module_vars, 1u |
| %8:vec4<f32> = textureSample %6, %7, vec2<f32>(0.0f) |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, Workgroup) { |
| auto* var_a = b.Var("a", ty.ptr<workgroup, i32>()); |
| auto* var_b = b.Var("b", ty.ptr<workgroup, i32>()); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute, |
| std::array<uint32_t, 3>{1u, 1u, 1u}); |
| b.Append(func->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_a, b.Add<i32>(load_a, load_b)); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<workgroup, i32, read_write> = var |
| %b:ptr<workgroup, i32, read_write> = var |
| } |
| |
| %foo = @compute @workgroup_size(1, 1, 1) func():void { |
| $B2: { |
| %4:i32 = load %a |
| %5:i32 = load %b |
| %6:i32 = add %4, %5 |
| store %a, %6 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<workgroup, i32, read_write> @offset(0) |
| b:ptr<workgroup, i32, read_write> @offset(0) |
| } |
| |
| tint_symbol_2 = struct @align(4) { |
| tint_symbol:i32 @offset(0) |
| tint_symbol_1:i32 @offset(4) |
| } |
| |
| %foo = @compute @workgroup_size(1, 1, 1) func(%2:ptr<workgroup, tint_symbol_2, read_write>):void { |
| $B1: { |
| %a:ptr<workgroup, i32, read_write> = access %2, 0u |
| %b:ptr<workgroup, i32, read_write> = access %2, 1u |
| %5:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars:tint_module_vars_struct = let %5 |
| %7:ptr<workgroup, i32, read_write> = access %tint_module_vars, 0u |
| %8:i32 = load %7 |
| %9:ptr<workgroup, i32, read_write> = access %tint_module_vars, 1u |
| %10:i32 = load %9 |
| %11:i32 = add %8, %10 |
| %12:ptr<workgroup, i32, read_write> = access %tint_module_vars, 0u |
| store %12, %11 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, MultipleAddressSpaces) { |
| auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| auto* var_c = b.Var("c", ty.ptr<private_, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| mod.root_block->Append(var_c); |
| |
| auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| auto* load_c = b.Load(var_c); |
| b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, load_c))); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<uniform, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| %c:ptr<private, i32, read_write> = var |
| } |
| |
| %foo = @fragment func():void { |
| $B2: { |
| %5:i32 = load %a |
| %6:i32 = load %b |
| %7:i32 = load %c |
| %8:i32 = add %6, %7 |
| %9:i32 = add %5, %8 |
| store %b, %9 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<uniform, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| c:ptr<private, i32, read_write> @offset(0) |
| } |
| |
| %foo = @fragment func(%a:ptr<uniform, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B1: { |
| %c:ptr<private, i32, read_write> = var |
| %5:tint_module_vars_struct = construct %a, %b, %c |
| %tint_module_vars:tint_module_vars_struct = let %5 |
| %7:ptr<uniform, i32, read> = access %tint_module_vars, 0u |
| %8:i32 = load %7 |
| %9:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| %10:i32 = load %9 |
| %11:ptr<private, i32, read_write> = access %tint_module_vars, 2u |
| %12:i32 = load %11 |
| %13:i32 = add %10, %12 |
| %14:i32 = add %8, %13 |
| %15:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| store %15, %14 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, EntryPointHasExistingParameters) { |
| auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| auto* param = b.FunctionParam<i32>("param"); |
| param->SetLocation( |
| core::ir::Location{1_u, core::Interpolation{core::InterpolationType::kFlat}}); |
| func->SetParams({param}); |
| b.Append(func->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, param))); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<storage, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| } |
| |
| %foo = @fragment func(%param:i32 [@location(1), @interpolate(flat)]):void { |
| $B2: { |
| %5:i32 = load %a |
| %6:i32 = load %b |
| %7:i32 = add %6, %param |
| %8:i32 = add %5, %7 |
| store %b, %8 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<storage, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| } |
| |
| %foo = @fragment func(%param:i32 [@location(1), @interpolate(flat)], %a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B1: { |
| %5:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars:tint_module_vars_struct = let %5 |
| %7:ptr<storage, i32, read> = access %tint_module_vars, 0u |
| %8:i32 = load %7 |
| %9:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| %10:i32 = load %9 |
| %11:i32 = add %10, %param |
| %12:i32 = add %8, %11 |
| %13:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| store %13, %12 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatUsesVars_NoArgs) { |
| auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* foo = b.Function("foo", ty.void_()); |
| b.Append(foo->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_b, b.Add<i32>(load_a, load_b)); |
| b.Return(foo); |
| }); |
| |
| auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| b.Call(foo); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<storage, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| } |
| |
| %foo = func():void { |
| $B2: { |
| %4:i32 = load %a |
| %5:i32 = load %b |
| %6:i32 = add %4, %5 |
| store %b, %6 |
| ret |
| } |
| } |
| %main = @fragment func():void { |
| $B3: { |
| %8:void = call %foo |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<storage, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| } |
| |
| %foo = func(%tint_module_vars:tint_module_vars_struct):void { |
| $B1: { |
| %3:ptr<storage, i32, read> = access %tint_module_vars, 0u |
| %4:i32 = load %3 |
| %5:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| %6:i32 = load %5 |
| %7:i32 = add %4, %6 |
| %8:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| store %8, %7 |
| ret |
| } |
| } |
| %main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B2: { |
| %12:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars_1:tint_module_vars_struct = let %12 # %tint_module_vars_1: 'tint_module_vars' |
| %14:void = call %foo, %tint_module_vars_1 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatUsesVars_WithExistingParameters) { |
| auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* foo = b.Function("foo", ty.void_()); |
| auto* param = b.FunctionParam<i32>("param"); |
| foo->SetParams({param}); |
| b.Append(foo->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, param))); |
| b.Return(foo); |
| }); |
| |
| auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| b.Call(foo, 42_i); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<storage, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| } |
| |
| %foo = func(%param:i32):void { |
| $B2: { |
| %5:i32 = load %a |
| %6:i32 = load %b |
| %7:i32 = add %6, %param |
| %8:i32 = add %5, %7 |
| store %b, %8 |
| ret |
| } |
| } |
| %main = @fragment func():void { |
| $B3: { |
| %10:void = call %foo, 42i |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<storage, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| } |
| |
| %foo = func(%param:i32, %tint_module_vars:tint_module_vars_struct):void { |
| $B1: { |
| %4:ptr<storage, i32, read> = access %tint_module_vars, 0u |
| %5:i32 = load %4 |
| %6:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| %7:i32 = load %6 |
| %8:i32 = add %7, %param |
| %9:i32 = add %5, %8 |
| %10:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| store %10, %9 |
| ret |
| } |
| } |
| %main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B2: { |
| %14:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars_1:tint_module_vars_struct = let %14 # %tint_module_vars_1: 'tint_module_vars' |
| %16:void = call %foo, 42i, %tint_module_vars_1 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatUsesVars_HandleTypes) { |
| auto* t = ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32()); |
| auto* var_t = b.Var("t", ty.ptr<handle>(t)); |
| auto* var_s = b.Var("s", ty.ptr<handle>(ty.sampler())); |
| var_t->SetBindingPoint(1, 2); |
| var_s->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_t); |
| mod.root_block->Append(var_s); |
| |
| auto* foo = b.Function("foo", ty.vec4<f32>()); |
| auto* param = b.FunctionParam<i32>("param"); |
| foo->SetParams({param}); |
| b.Append(foo->Block(), [&] { |
| auto* load_t = b.Load(var_t); |
| auto* load_s = b.Load(var_s); |
| auto* result = b.Call<vec4<f32>>(core::BuiltinFn::kTextureSample, load_t, load_s, |
| b.Splat<vec2<f32>>(0_f)); |
| b.Return(foo, result); |
| }); |
| |
| auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| b.Call(foo, 42_i); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %t:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 2) |
| %s:ptr<handle, sampler, read> = var @binding_point(3, 4) |
| } |
| |
| %foo = func(%param:i32):vec4<f32> { |
| $B2: { |
| %5:texture_2d<f32> = load %t |
| %6:sampler = load %s |
| %7:vec4<f32> = textureSample %5, %6, vec2<f32>(0.0f) |
| ret %7 |
| } |
| } |
| %main = @fragment func():void { |
| $B3: { |
| %9:vec4<f32> = call %foo, 42i |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| t:texture_2d<f32> @offset(0) |
| s:sampler @offset(0) |
| } |
| |
| %foo = func(%param:i32, %tint_module_vars:tint_module_vars_struct):vec4<f32> { |
| $B1: { |
| %4:texture_2d<f32> = access %tint_module_vars, 0u |
| %5:sampler = access %tint_module_vars, 1u |
| %6:vec4<f32> = textureSample %4, %5, vec2<f32>(0.0f) |
| ret %6 |
| } |
| } |
| %main = @fragment func(%t:texture_2d<f32> [@binding_point(1, 2)], %s:sampler [@binding_point(3, 4)]):void { |
| $B2: { |
| %10:tint_module_vars_struct = construct %t, %s |
| %tint_module_vars_1:tint_module_vars_struct = let %10 # %tint_module_vars_1: 'tint_module_vars' |
| %12:vec4<f32> = call %foo, 42i, %tint_module_vars_1 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatUsesVars_OutOfOrder) { |
| auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| |
| auto* foo = b.Function("foo", ty.void_()); |
| b.Append(foo->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_b, b.Add<i32>(load_a, load_b)); |
| b.Return(foo); |
| }); |
| |
| b.Append(func->Block(), [&] { |
| b.Call(foo); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<storage, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| } |
| |
| %main = @fragment func():void { |
| $B2: { |
| %4:void = call %foo |
| ret |
| } |
| } |
| %foo = func():void { |
| $B3: { |
| %6:i32 = load %a |
| %7:i32 = load %b |
| %8:i32 = add %6, %7 |
| store %b, %8 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<storage, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| } |
| |
| %main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B1: { |
| %4:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars:tint_module_vars_struct = let %4 |
| %6:void = call %foo, %tint_module_vars |
| ret |
| } |
| } |
| %foo = func(%tint_module_vars_1:tint_module_vars_struct):void { # %tint_module_vars_1: 'tint_module_vars' |
| $B2: { |
| %9:ptr<storage, i32, read> = access %tint_module_vars_1, 0u |
| %10:i32 = load %9 |
| %11:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u |
| %12:i32 = load %11 |
| %13:i32 = add %10, %12 |
| %14:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u |
| store %14, %13 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| // Test that we do not add the structure to functions that do not need it. |
| TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionThatDoesNotUseVars) { |
| auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* foo = b.Function("foo", ty.i32()); |
| b.Append(foo->Block(), [&] { // |
| b.Return(foo, 42_i); |
| }); |
| |
| auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, b.Call(foo)))); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<storage, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| } |
| |
| %foo = func():i32 { |
| $B2: { |
| ret 42i |
| } |
| } |
| %main = @fragment func():void { |
| $B3: { |
| %5:i32 = load %a |
| %6:i32 = load %b |
| %7:i32 = call %foo |
| %8:i32 = add %6, %7 |
| %9:i32 = add %5, %8 |
| store %b, %9 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<storage, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| } |
| |
| %foo = func():i32 { |
| $B1: { |
| ret 42i |
| } |
| } |
| %main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B2: { |
| %5:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars:tint_module_vars_struct = let %5 |
| %7:ptr<storage, i32, read> = access %tint_module_vars, 0u |
| %8:i32 = load %7 |
| %9:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| %10:i32 = load %9 |
| %11:i32 = call %foo |
| %12:i32 = add %10, %11 |
| %13:i32 = add %8, %12 |
| %14:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| store %14, %13 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| // Test that we *do* add the structure to functions that only have transitive uses. |
| TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionWithOnlyTransitiveUses) { |
| auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* bar = b.Function("bar", ty.i32()); |
| b.Append(bar->Block(), [&] { // |
| b.Return(bar, b.Load(var_a)); |
| }); |
| |
| auto* foo = b.Function("foo", ty.i32()); |
| b.Append(foo->Block(), [&] { // |
| b.Return(foo, b.Call(bar)); |
| }); |
| |
| auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, b.Call(foo)))); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<storage, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| } |
| |
| %bar = func():i32 { |
| $B2: { |
| %4:i32 = load %a |
| ret %4 |
| } |
| } |
| %foo = func():i32 { |
| $B3: { |
| %6:i32 = call %bar |
| ret %6 |
| } |
| } |
| %main = @fragment func():void { |
| $B4: { |
| %8:i32 = load %a |
| %9:i32 = load %b |
| %10:i32 = call %foo |
| %11:i32 = add %9, %10 |
| %12:i32 = add %8, %11 |
| store %b, %12 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<storage, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| } |
| |
| %bar = func(%tint_module_vars:tint_module_vars_struct):i32 { |
| $B1: { |
| %3:ptr<storage, i32, read> = access %tint_module_vars, 0u |
| %4:i32 = load %3 |
| ret %4 |
| } |
| } |
| %foo = func(%tint_module_vars_1:tint_module_vars_struct):i32 { # %tint_module_vars_1: 'tint_module_vars' |
| $B2: { |
| %7:i32 = call %bar, %tint_module_vars_1 |
| ret %7 |
| } |
| } |
| %main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B3: { |
| %11:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars_2:tint_module_vars_struct = let %11 # %tint_module_vars_2: 'tint_module_vars' |
| %13:ptr<storage, i32, read> = access %tint_module_vars_2, 0u |
| %14:i32 = load %13 |
| %15:ptr<storage, i32, read_write> = access %tint_module_vars_2, 1u |
| %16:i32 = load %15 |
| %17:i32 = call %foo, %tint_module_vars_2 |
| %18:i32 = add %16, %17 |
| %19:i32 = add %14, %18 |
| %20:ptr<storage, i32, read_write> = access %tint_module_vars_2, 1u |
| store %20, %19 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| // Test that we *do* add the structure to functions that only have transitive uses, where that |
| // function is declared first. |
| TEST_F(MslWriter_ModuleScopeVarsTest, CallFunctionWithOnlyTransitiveUses_OutOfOrder) { |
| auto* var_a = b.Var("a", ty.ptr<storage, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| |
| auto* foo = b.Function("foo", ty.i32()); |
| |
| auto* bar = b.Function("bar", ty.i32()); |
| b.Append(bar->Block(), [&] { // |
| b.Return(bar, b.Load(var_a)); |
| }); |
| |
| b.Append(foo->Block(), [&] { // |
| b.Return(foo, b.Call(bar)); |
| }); |
| |
| auto* func = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, b.Call(foo)))); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<storage, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| } |
| |
| %foo = func():i32 { |
| $B2: { |
| %4:i32 = call %bar |
| ret %4 |
| } |
| } |
| %bar = func():i32 { |
| $B3: { |
| %6:i32 = load %a |
| ret %6 |
| } |
| } |
| %main = @fragment func():void { |
| $B4: { |
| %8:i32 = load %a |
| %9:i32 = load %b |
| %10:i32 = call %foo |
| %11:i32 = add %9, %10 |
| %12:i32 = add %8, %11 |
| store %b, %12 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<storage, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| } |
| |
| %foo = func(%tint_module_vars:tint_module_vars_struct):i32 { |
| $B1: { |
| %3:i32 = call %bar, %tint_module_vars |
| ret %3 |
| } |
| } |
| %bar = func(%tint_module_vars_1:tint_module_vars_struct):i32 { # %tint_module_vars_1: 'tint_module_vars' |
| $B2: { |
| %6:ptr<storage, i32, read> = access %tint_module_vars_1, 0u |
| %7:i32 = load %6 |
| ret %7 |
| } |
| } |
| %main = @fragment func(%a:ptr<storage, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B3: { |
| %11:tint_module_vars_struct = construct %a, %b |
| %tint_module_vars_2:tint_module_vars_struct = let %11 # %tint_module_vars_2: 'tint_module_vars' |
| %13:ptr<storage, i32, read> = access %tint_module_vars_2, 0u |
| %14:i32 = load %13 |
| %15:ptr<storage, i32, read_write> = access %tint_module_vars_2, 1u |
| %16:i32 = load %15 |
| %17:i32 = call %foo, %tint_module_vars_2 |
| %18:i32 = add %16, %17 |
| %19:i32 = add %14, %18 |
| %20:ptr<storage, i32, read_write> = access %tint_module_vars_2, 1u |
| store %20, %19 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, MultipleEntryPoints) { |
| auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| auto* var_c = b.Var("c", ty.ptr<private_, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| mod.root_block->Append(var_c); |
| |
| auto* main_a = b.Function("main_a", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(main_a->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| auto* load_c = b.Load(var_c); |
| b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, load_c))); |
| b.Return(main_a); |
| }); |
| |
| auto* main_b = b.Function("main_b", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(main_b->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| auto* load_c = b.Load(var_c); |
| b.Store(var_b, b.Multiply<i32>(load_a, b.Multiply<i32>(load_b, load_c))); |
| b.Return(main_b); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<uniform, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| %c:ptr<private, i32, read_write> = var |
| } |
| |
| %main_a = @fragment func():void { |
| $B2: { |
| %5:i32 = load %a |
| %6:i32 = load %b |
| %7:i32 = load %c |
| %8:i32 = add %6, %7 |
| %9:i32 = add %5, %8 |
| store %b, %9 |
| ret |
| } |
| } |
| %main_b = @fragment func():void { |
| $B3: { |
| %11:i32 = load %a |
| %12:i32 = load %b |
| %13:i32 = load %c |
| %14:i32 = mul %12, %13 |
| %15:i32 = mul %11, %14 |
| store %b, %15 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<uniform, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| c:ptr<private, i32, read_write> @offset(0) |
| } |
| |
| %main_a = @fragment func(%a:ptr<uniform, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B1: { |
| %c:ptr<private, i32, read_write> = var |
| %5:tint_module_vars_struct = construct %a, %b, %c |
| %tint_module_vars:tint_module_vars_struct = let %5 |
| %7:ptr<uniform, i32, read> = access %tint_module_vars, 0u |
| %8:i32 = load %7 |
| %9:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| %10:i32 = load %9 |
| %11:ptr<private, i32, read_write> = access %tint_module_vars, 2u |
| %12:i32 = load %11 |
| %13:i32 = add %10, %12 |
| %14:i32 = add %8, %13 |
| %15:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| store %15, %14 |
| ret |
| } |
| } |
| %main_b = @fragment func(%a_1:ptr<uniform, i32, read> [@binding_point(1, 2)], %b_1:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { # %a_1: 'a', %b_1: 'b' |
| $B2: { |
| %c_1:ptr<private, i32, read_write> = var # %c_1: 'c' |
| %20:tint_module_vars_struct = construct %a_1, %b_1, %c_1 |
| %tint_module_vars_1:tint_module_vars_struct = let %20 # %tint_module_vars_1: 'tint_module_vars' |
| %22:ptr<uniform, i32, read> = access %tint_module_vars_1, 0u |
| %23:i32 = load %22 |
| %24:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u |
| %25:i32 = load %24 |
| %26:ptr<private, i32, read_write> = access %tint_module_vars_1, 2u |
| %27:i32 = load %26 |
| %28:i32 = mul %25, %27 |
| %29:i32 = mul %23, %28 |
| %30:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u |
| store %30, %29 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, MultipleEntryPoints_DifferentUsageSets) { |
| auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| auto* var_c = b.Var("c", ty.ptr<private_, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| mod.root_block->Append(var_c); |
| |
| auto* main_a = b.Function("main_a", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(main_a->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| b.Store(var_b, b.Add<i32>(load_a, load_b)); |
| b.Return(main_a); |
| }); |
| |
| auto* main_b = b.Function("main_b", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(main_b->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_c = b.Load(var_c); |
| b.Store(var_c, b.Multiply<i32>(load_a, load_c)); |
| b.Return(main_b); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<uniform, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| %c:ptr<private, i32, read_write> = var |
| } |
| |
| %main_a = @fragment func():void { |
| $B2: { |
| %5:i32 = load %a |
| %6:i32 = load %b |
| %7:i32 = add %5, %6 |
| store %b, %7 |
| ret |
| } |
| } |
| %main_b = @fragment func():void { |
| $B3: { |
| %9:i32 = load %a |
| %10:i32 = load %c |
| %11:i32 = mul %9, %10 |
| store %c, %11 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<uniform, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| c:ptr<private, i32, read_write> @offset(0) |
| } |
| |
| %main_a = @fragment func(%a:ptr<uniform, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B1: { |
| %4:tint_module_vars_struct = construct %a, %b, undef |
| %tint_module_vars:tint_module_vars_struct = let %4 |
| %6:ptr<uniform, i32, read> = access %tint_module_vars, 0u |
| %7:i32 = load %6 |
| %8:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| %9:i32 = load %8 |
| %10:i32 = add %7, %9 |
| %11:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| store %11, %10 |
| ret |
| } |
| } |
| %main_b = @fragment func(%a_1:ptr<uniform, i32, read> [@binding_point(1, 2)]):void { # %a_1: 'a' |
| $B2: { |
| %c:ptr<private, i32, read_write> = var |
| %15:tint_module_vars_struct = construct %a_1, undef, %c |
| %tint_module_vars_1:tint_module_vars_struct = let %15 # %tint_module_vars_1: 'tint_module_vars' |
| %17:ptr<uniform, i32, read> = access %tint_module_vars_1, 0u |
| %18:i32 = load %17 |
| %19:ptr<private, i32, read_write> = access %tint_module_vars_1, 2u |
| %20:i32 = load %19 |
| %21:i32 = mul %18, %20 |
| %22:ptr<private, i32, read_write> = access %tint_module_vars_1, 2u |
| store %22, %21 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, MultipleEntryPoints_DifferentUsageSets_CommonHelper) { |
| auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>()); |
| auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| auto* var_c = b.Var("c", ty.ptr<private_, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| mod.root_block->Append(var_c); |
| |
| auto* foo = b.Function("foo", ty.i32()); |
| b.Append(foo->Block(), [&] { // |
| b.Return(foo, b.Load(var_a)); |
| }); |
| |
| auto* main_a = b.Function("main_a", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(main_a->Block(), [&] { |
| auto* load_b = b.Load(var_b); |
| b.Store(var_b, b.Add<i32>(b.Call(foo), load_b)); |
| b.Return(main_a); |
| }); |
| |
| auto* main_b = b.Function("main_b", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(main_b->Block(), [&] { |
| auto* load_c = b.Load(var_c); |
| b.Store(var_c, b.Multiply<i32>(b.Call(foo), load_c)); |
| b.Return(main_b); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %a:ptr<uniform, i32, read> = var @binding_point(1, 2) |
| %b:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| %c:ptr<private, i32, read_write> = var |
| } |
| |
| %foo = func():i32 { |
| $B2: { |
| %5:i32 = load %a |
| ret %5 |
| } |
| } |
| %main_a = @fragment func():void { |
| $B3: { |
| %7:i32 = load %b |
| %8:i32 = call %foo |
| %9:i32 = add %8, %7 |
| store %b, %9 |
| ret |
| } |
| } |
| %main_b = @fragment func():void { |
| $B4: { |
| %11:i32 = load %c |
| %12:i32 = call %foo |
| %13:i32 = mul %12, %11 |
| store %c, %13 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| a:ptr<uniform, i32, read> @offset(0) |
| b:ptr<storage, i32, read_write> @offset(0) |
| c:ptr<private, i32, read_write> @offset(0) |
| } |
| |
| %foo = func(%tint_module_vars:tint_module_vars_struct):i32 { |
| $B1: { |
| %3:ptr<uniform, i32, read> = access %tint_module_vars, 0u |
| %4:i32 = load %3 |
| ret %4 |
| } |
| } |
| %main_a = @fragment func(%a:ptr<uniform, i32, read> [@binding_point(1, 2)], %b:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B2: { |
| %8:tint_module_vars_struct = construct %a, %b, undef |
| %tint_module_vars_1:tint_module_vars_struct = let %8 # %tint_module_vars_1: 'tint_module_vars' |
| %10:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u |
| %11:i32 = load %10 |
| %12:i32 = call %foo, %tint_module_vars_1 |
| %13:i32 = add %12, %11 |
| %14:ptr<storage, i32, read_write> = access %tint_module_vars_1, 1u |
| store %14, %13 |
| ret |
| } |
| } |
| %main_b = @fragment func(%a_1:ptr<uniform, i32, read> [@binding_point(1, 2)]):void { # %a_1: 'a' |
| $B3: { |
| %c:ptr<private, i32, read_write> = var |
| %18:tint_module_vars_struct = construct %a_1, undef, %c |
| %tint_module_vars_2:tint_module_vars_struct = let %18 # %tint_module_vars_2: 'tint_module_vars' |
| %20:ptr<private, i32, read_write> = access %tint_module_vars_2, 2u |
| %21:i32 = load %20 |
| %22:i32 = call %foo, %tint_module_vars_2 |
| %23:i32 = mul %22, %21 |
| %24:ptr<private, i32, read_write> = access %tint_module_vars_2, 2u |
| store %24, %23 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| TEST_F(MslWriter_ModuleScopeVarsTest, VarsWithNoNames) { |
| auto* var_a = b.Var(ty.ptr<uniform, i32, core::Access::kRead>()); |
| auto* var_b = b.Var(ty.ptr<storage, i32, core::Access::kReadWrite>()); |
| auto* var_c = b.Var(ty.ptr<private_, i32, core::Access::kReadWrite>()); |
| var_a->SetBindingPoint(1, 2); |
| var_b->SetBindingPoint(3, 4); |
| mod.root_block->Append(var_a); |
| mod.root_block->Append(var_b); |
| mod.root_block->Append(var_c); |
| |
| auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment); |
| b.Append(func->Block(), [&] { |
| auto* load_a = b.Load(var_a); |
| auto* load_b = b.Load(var_b); |
| auto* load_c = b.Load(var_c); |
| b.Store(var_b, b.Add<i32>(load_a, b.Add<i32>(load_b, load_c))); |
| b.Return(func); |
| }); |
| |
| auto* src = R"( |
| $B1: { # root |
| %1:ptr<uniform, i32, read> = var @binding_point(1, 2) |
| %2:ptr<storage, i32, read_write> = var @binding_point(3, 4) |
| %3:ptr<private, i32, read_write> = var |
| } |
| |
| %foo = @fragment func():void { |
| $B2: { |
| %5:i32 = load %1 |
| %6:i32 = load %2 |
| %7:i32 = load %3 |
| %8:i32 = add %6, %7 |
| %9:i32 = add %5, %8 |
| store %2, %9 |
| ret |
| } |
| } |
| )"; |
| EXPECT_EQ(src, str()); |
| |
| auto* expect = R"( |
| tint_module_vars_struct = struct @align(1) { |
| tint_symbol:ptr<uniform, i32, read> @offset(0) |
| tint_symbol_1:ptr<storage, i32, read_write> @offset(0) |
| tint_symbol_2:ptr<private, i32, read_write> @offset(0) |
| } |
| |
| %foo = @fragment func(%2:ptr<uniform, i32, read> [@binding_point(1, 2)], %3:ptr<storage, i32, read_write> [@binding_point(3, 4)]):void { |
| $B1: { |
| %4:ptr<private, i32, read_write> = var |
| %5:tint_module_vars_struct = construct %2, %3, %4 |
| %tint_module_vars:tint_module_vars_struct = let %5 |
| %7:ptr<uniform, i32, read> = access %tint_module_vars, 0u |
| %8:i32 = load %7 |
| %9:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| %10:i32 = load %9 |
| %11:ptr<private, i32, read_write> = access %tint_module_vars, 2u |
| %12:i32 = load %11 |
| %13:i32 = add %10, %12 |
| %14:i32 = add %8, %13 |
| %15:ptr<storage, i32, read_write> = access %tint_module_vars, 1u |
| store %15, %14 |
| ret |
| } |
| } |
| )"; |
| |
| Run(ModuleScopeVars); |
| |
| EXPECT_EQ(expect, str()); |
| } |
| |
| } // namespace |
| } // namespace tint::msl::writer::raise |