blob: 2892d6054a6555e43a8762b68443388851bea6d9 [file] [log] [blame]
// 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