[ir] Add capability for multiple entry points
Most transforms and printers assume there is only one entry point, as
we require that SingleEntryPoint has run before entering the
backends. Capture this in a capability so that the fuzzers respect
this precondition.
Change-Id: I8b20ab1ae9cacfe3129c20197c1bebbdf2e367b5
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/244157
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Auto-Submit: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/cmd/fuzz/ir/fuzz.cc b/src/tint/cmd/fuzz/ir/fuzz.cc
index 948c6ef..56fa4f3 100644
--- a/src/tint/cmd/fuzz/ir/fuzz.cc
+++ b/src/tint/cmd/fuzz/ir/fuzz.cc
@@ -101,7 +101,11 @@
}
}
- if (auto val = core::ir::Validate(ir.Get()); val != Success) {
+ if (auto val = core::ir::Validate(ir.Get(),
+ core::ir::Capabilities{
+ core::ir::Capability::kAllowMultipleEntryPoints,
+ });
+ val != Success) {
TINT_ICE() << val.Failure();
}
// Copy relevant options from wgsl::Context to ir::Context
diff --git a/src/tint/lang/core/ir/transform/single_entry_point.cc b/src/tint/lang/core/ir/transform/single_entry_point.cc
index 4ec9c04..040f2a0 100644
--- a/src/tint/lang/core/ir/transform/single_entry_point.cc
+++ b/src/tint/lang/core/ir/transform/single_entry_point.cc
@@ -95,8 +95,11 @@
} // namespace
Result<SuccessType> SingleEntryPoint(Module& ir, std::string_view entry_point_name) {
- auto result = ValidateAndDumpIfNeeded(
- ir, "core.SingleEntryPoint", core::ir::Capabilities{core::ir::Capability::kAllowOverrides});
+ auto result = ValidateAndDumpIfNeeded(ir, "core.SingleEntryPoint",
+ Capabilities{
+ Capability::kAllowMultipleEntryPoints,
+ Capability::kAllowOverrides,
+ });
if (result != Success) {
return result.Failure();
}
diff --git a/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc b/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc
index d232617..eaa1a27 100644
--- a/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc
+++ b/src/tint/lang/core/ir/transform/zero_init_workgroup_memory_test.cc
@@ -1805,121 +1805,5 @@
EXPECT_EQ(expect, str());
}
-TEST_F(IR_ZeroInitWorkgroupMemoryTest, MultipleEntryPoints_SameVarViaHelper) {
- auto* var = MakeVar("wgvar", ty.bool_());
-
- auto* foo = b.Function("foo", ty.void_());
- b.Append(foo->Block(), [&] { //
- auto* loop = b.Loop();
- b.Append(loop->Body(), [&] { //
- b.Continue(loop);
- b.Append(loop->Continuing(), [&] { //
- auto* load = b.Load(var);
- b.BreakIf(loop, load);
- });
- });
- b.Return(foo);
- });
-
- auto* ep1 = MakeEntryPoint("ep1", 1, 1, 1);
- b.Append(ep1->Block(), [&] { //
- b.Call(ty.void_(), foo);
- b.Return(ep1);
- });
-
- auto* ep2 = MakeEntryPoint("ep2", 1, 1, 1);
- b.Append(ep2->Block(), [&] { //
- b.Call(ty.void_(), foo);
- b.Return(ep2);
- });
-
- auto* src = R"(
-$B1: { # root
- %wgvar:ptr<workgroup, bool, read_write> = var undef
-}
-
-%foo = func():void {
- $B2: {
- loop [b: $B3, c: $B4] { # loop_1
- $B3: { # body
- continue # -> $B4
- }
- $B4: { # continuing
- %3:bool = load %wgvar
- break_if %3 # -> [t: exit_loop loop_1, f: $B3]
- }
- }
- ret
- }
-}
-%ep1 = @compute @workgroup_size(1u, 1u, 1u) func():void {
- $B5: {
- %5:void = call %foo
- ret
- }
-}
-%ep2 = @compute @workgroup_size(1u, 1u, 1u) func():void {
- $B6: {
- %7:void = call %foo
- ret
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-$B1: { # root
- %wgvar:ptr<workgroup, bool, read_write> = var undef
-}
-
-%foo = func():void {
- $B2: {
- loop [b: $B3, c: $B4] { # loop_1
- $B3: { # body
- continue # -> $B4
- }
- $B4: { # continuing
- %3:bool = load %wgvar
- break_if %3 # -> [t: exit_loop loop_1, f: $B3]
- }
- }
- ret
- }
-}
-%ep1 = @compute @workgroup_size(1u, 1u, 1u) func(%tint_local_index:u32 [@local_invocation_index]):void {
- $B5: {
- %6:bool = lt %tint_local_index, 1u
- if %6 [t: $B6] { # if_1
- $B6: { # true
- store %wgvar, false
- exit_if # if_1
- }
- }
- %7:void = workgroupBarrier
- %8:void = call %foo
- ret
- }
-}
-%ep2 = @compute @workgroup_size(1u, 1u, 1u) func(%tint_local_index_1:u32 [@local_invocation_index]):void { # %tint_local_index_1: 'tint_local_index'
- $B7: {
- %11:bool = lt %tint_local_index_1, 1u
- if %11 [t: $B8] { # if_2
- $B8: { # true
- store %wgvar, false
- exit_if # if_2
- }
- }
- %12:void = workgroupBarrier
- %13:void = call %foo
- ret
- }
-}
-)";
-
- Run(ZeroInitWorkgroupMemory);
-
- EXPECT_EQ(expect, str());
-}
-
} // namespace
} // namespace tint::core::ir::transform
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 24fc0db..4b09228 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -2072,8 +2072,17 @@
scope_stack_.Push();
TINT_DEFER(scope_stack_.Pop());
- // Checking the name early, so its usage can be recorded, even if the function is malformed.
if (func->IsEntryPoint()) {
+ // Check that there is at most one entry point unless we allow multiple entry points.
+ if (!capabilities_.Contains(Capability::kAllowMultipleEntryPoints)) {
+ if (!entry_point_names_.IsEmpty()) {
+ AddError(func) << "a module with multiple entry points requires the "
+ "AllowMultipleEntryPoints capability";
+ return;
+ }
+ }
+
+ // Checking the name early, so its usage can be recorded, even if the function is malformed.
const auto name = mod_.NameOf(func).Name();
if (!entry_point_names_.Add(name)) {
AddError(func) << "entry point name " << style::Function(name) << " is not unique";
diff --git a/src/tint/lang/core/ir/validator.h b/src/tint/lang/core/ir/validator.h
index b7efa1d..444537c 100644
--- a/src/tint/lang/core/ir/validator.h
+++ b/src/tint/lang/core/ir/validator.h
@@ -50,6 +50,8 @@
kAllowHandleVarsWithoutBindings,
/// Allows module scoped lets
kAllowModuleScopeLets,
+ /// Allows multiple entry points in the module.
+ kAllowMultipleEntryPoints,
/// Allow overrides
kAllowOverrides,
/// Allows pointers and handle addressspace variables inside structures.
diff --git a/src/tint/lang/core/ir/validator_function_test.cc b/src/tint/lang/core/ir/validator_function_test.cc
index 0eafdec..796a42d 100644
--- a/src/tint/lang/core/ir/validator_function_test.cc
+++ b/src/tint/lang/core/ir/validator_function_test.cc
@@ -95,6 +95,37 @@
)")) << res.Failure();
}
+TEST_F(IR_ValidatorTest, Function_MultipleEntryPoints_WithCapability) {
+ auto* ep1 = ComputeEntryPoint("ep1");
+ ep1->Block()->Append(b.Return(ep1));
+
+ auto* ep2 = ComputeEntryPoint("ep2");
+ ep2->Block()->Append(b.Return(ep2));
+
+ auto res = ir::Validate(mod, Capabilities{
+ Capability::kAllowMultipleEntryPoints,
+ });
+ ASSERT_EQ(res, Success) << res.Failure();
+}
+
+TEST_F(IR_ValidatorTest, Function_MultipleEntryPoints_WithoutCapability) {
+ auto* ep1 = ComputeEntryPoint("ep1");
+ ep1->Block()->Append(b.Return(ep1));
+
+ auto* ep2 = ComputeEntryPoint("ep2");
+ ep2->Block()->Append(b.Return(ep2));
+
+ auto res = ir::Validate(mod);
+ ASSERT_NE(res, Success);
+ EXPECT_THAT(
+ res.Failure().reason,
+ testing::HasSubstr(
+ R"(:6:1 error: a module with multiple entry points requires the AllowMultipleEntryPoints capability
+%ep2 = @compute @workgroup_size(1u, 1u, 1u) func():void {
+^^^^
+)")) << res.Failure();
+}
+
TEST_F(IR_ValidatorTest, Function_DuplicateEntryPointNames) {
auto* c = ComputeEntryPoint("dup");
c->Block()->Append(b.Return(c));
@@ -102,7 +133,9 @@
auto* f = FragmentEntryPoint("dup");
f->Block()->Append(b.Return(f));
- auto res = ir::Validate(mod);
+ auto res = ir::Validate(mod, Capabilities{
+ Capability::kAllowMultipleEntryPoints,
+ });
ASSERT_NE(res, Success);
EXPECT_THAT(res.Failure().reason,
testing::HasSubstr(R"(:6:1 error: entry point name 'dup' is not unique
diff --git a/src/tint/lang/glsl/writer/raise/shader_io_test.cc b/src/tint/lang/glsl/writer/raise/shader_io_test.cc
index ad4134f..db661c6 100644
--- a/src/tint/lang/glsl/writer/raise/shader_io_test.cc
+++ b/src/tint/lang/glsl/writer/raise/shader_io_test.cc
@@ -679,147 +679,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(GlslWriter_ShaderIOTest, Struct_SharedByVertexAndFragment) {
- auto* vec4f = ty.vec4<f32>();
- auto* str_ty = ty.Struct(mod.symbols.New("Interface"),
- {
- {
- mod.symbols.New("position"),
- vec4f,
- core::IOAttributes{
- .builtin = core::BuiltinValue::kPosition,
- },
- },
- {
- mod.symbols.New("color"),
- vec4f,
- core::IOAttributes{
- .location = 0u,
- },
- },
- });
-
- // Vertex shader.
- {
- auto* ep = b.Function("vert", str_ty);
- ep->SetStage(core::ir::Function::PipelineStage::kVertex);
-
- b.Append(ep->Block(), [&] { //
- auto* position = b.Construct(vec4f, 0_f);
- auto* color = b.Construct(vec4f, 1_f);
- b.Return(ep, b.Construct(str_ty, position, color));
- });
- }
-
- // Fragment shader.
- {
- auto* ep = b.Function("frag", vec4f);
- auto* inputs = b.FunctionParam("inputs", str_ty);
- ep->SetStage(core::ir::Function::PipelineStage::kFragment);
- ep->SetParams({inputs});
- ep->SetReturnLocation(0u);
-
- b.Append(ep->Block(), [&] { //
- auto* position = b.Access(vec4f, inputs, 0_u);
- auto* color = b.Access(vec4f, inputs, 1_u);
- b.Return(ep, b.Add(vec4f, position, color));
- });
- }
-
- auto* src = R"(
-Interface = struct @align(16) {
- position:vec4<f32> @offset(0), @builtin(position)
- color:vec4<f32> @offset(16), @location(0)
-}
-
-%vert = @vertex func():Interface {
- $B1: {
- %2:vec4<f32> = construct 0.0f
- %3:vec4<f32> = construct 1.0f
- %4:Interface = construct %2, %3
- ret %4
- }
-}
-%frag = @fragment func(%inputs:Interface):vec4<f32> [@location(0)] {
- $B2: {
- %7:vec4<f32> = access %inputs, 0u
- %8:vec4<f32> = access %inputs, 1u
- %9:vec4<f32> = add %7, %8
- ret %9
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-Interface = struct @align(16) {
- position:vec4<f32> @offset(0)
- color:vec4<f32> @offset(16)
-}
-
-$B1: { # root
- %vert_position:ptr<__out, vec4<f32>, write> = var undef @builtin(position)
- %vert_loc0_Output:ptr<__out, vec4<f32>, write> = var undef @location(0)
- %vert___point_size:ptr<__out, f32, write> = var undef @builtin(__point_size)
- %frag_position:ptr<__in, vec4<f32>, read> = var undef @builtin(position)
- %frag_loc0_Input:ptr<__in, vec4<f32>, read> = var undef @location(0)
- %frag_loc0_Output:ptr<__out, vec4<f32>, write> = var undef @location(0)
-}
-
-%vert_inner = func():Interface {
- $B2: {
- %8:vec4<f32> = construct 0.0f
- %9:vec4<f32> = construct 1.0f
- %10:Interface = construct %8, %9
- ret %10
- }
-}
-%frag_inner = func(%inputs:Interface):vec4<f32> {
- $B3: {
- %13:vec4<f32> = access %inputs, 0u
- %14:vec4<f32> = access %inputs, 1u
- %15:vec4<f32> = add %13, %14
- ret %15
- }
-}
-%vert = @vertex func():void {
- $B4: {
- %17:Interface = call %vert_inner
- %18:vec4<f32> = access %17, 0u
- %19:f32 = swizzle %18, x
- %20:f32 = swizzle %18, y
- %21:f32 = negation %20
- %22:f32 = swizzle %18, z
- %23:f32 = swizzle %18, w
- %24:f32 = mul 2.0f, %22
- %25:f32 = sub %24, %23
- %26:vec4<f32> = construct %19, %21, %25, %23
- store %vert_position, %26
- %27:vec4<f32> = access %17, 1u
- store %vert_loc0_Output, %27
- store %vert___point_size, 1.0f
- ret
- }
-}
-%frag = @fragment func():void {
- $B5: {
- %29:vec4<f32> = load %frag_position
- %30:vec4<f32> = load %frag_loc0_Input
- %31:Interface = construct %29, %30
- %32:vec4<f32> = call %frag_inner, %31
- store %frag_loc0_Output, %32
- ret
- }
-}
-)";
-
- core::ir::transform::ImmediateDataLayout immediate_data;
- ShaderIOConfig config{immediate_data};
- Run(ShaderIO, config);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(GlslWriter_ShaderIOTest, Struct_SharedWithBuffer) {
auto* vec4f = ty.vec4<f32>();
auto* str_ty =
@@ -1005,8 +864,8 @@
EXPECT_EQ(expect, str());
}
-// Test that interpolation attributes are stripped from vertex inputs and fragment outputs.
-TEST_F(GlslWriter_ShaderIOTest, InterpolationOnVertexInputOrFragmentOutput) {
+// Test that interpolation attributes are stripped from vertex inputs.
+TEST_F(GlslWriter_ShaderIOTest, InterpolationOnVertexInput) {
auto* str_ty = ty.Struct(mod.symbols.New("MyStruct"),
{
{
@@ -1023,45 +882,20 @@
},
});
- // Vertex shader.
- {
- auto* ep = b.Function("vert", ty.vec4<f32>());
- ep->SetReturnBuiltin(core::BuiltinValue::kPosition);
- ep->SetReturnInvariant(true);
- ep->SetStage(core::ir::Function::PipelineStage::kVertex);
+ auto* ep = b.Function("vert", ty.vec4<f32>());
+ ep->SetReturnBuiltin(core::BuiltinValue::kPosition);
+ ep->SetReturnInvariant(true);
+ ep->SetStage(core::ir::Function::PipelineStage::kVertex);
- auto* str_param = b.FunctionParam("input", str_ty);
- auto* ival = b.FunctionParam("ival", ty.i32());
- ival->SetLocation(1);
- ival->SetInterpolation(core::Interpolation{core::InterpolationType::kFlat});
- ep->SetParams({str_param, ival});
+ auto* str_param = b.FunctionParam("input", str_ty);
+ auto* ival = b.FunctionParam("ival", ty.i32());
+ ival->SetLocation(1);
+ ival->SetInterpolation(core::Interpolation{core::InterpolationType::kFlat});
+ ep->SetParams({str_param, ival});
- b.Append(ep->Block(), [&] { //
- b.Return(ep, b.Construct(ty.vec4<f32>(), 0.5_f));
- });
- }
-
- // Fragment shader with struct output.
- {
- auto* ep = b.Function("frag1", str_ty);
- ep->SetStage(core::ir::Function::PipelineStage::kFragment);
-
- b.Append(ep->Block(), [&] { //
- b.Return(ep, b.Construct(str_ty, 0.5_f));
- });
- }
-
- // Fragment shader with non-struct output.
- {
- auto* ep = b.Function("frag2", ty.i32());
- ep->SetStage(core::ir::Function::PipelineStage::kFragment);
- ep->SetReturnLocation(0);
- ep->SetReturnInterpolation(core::Interpolation{core::InterpolationType::kFlat});
-
- b.Append(ep->Block(), [&] { //
- b.Return(ep, b.Constant(42_i));
- });
- }
+ b.Append(ep->Block(), [&] { //
+ b.Return(ep, b.Construct(ty.vec4<f32>(), 0.5_f));
+ });
auto* src = R"(
MyStruct = struct @align(4) {
@@ -1074,17 +908,6 @@
ret %4
}
}
-%frag1 = @fragment func():MyStruct {
- $B2: {
- %6:MyStruct = construct 0.5f
- ret %6
- }
-}
-%frag2 = @fragment func():i32 [@location(0), @interpolate(flat)] {
- $B3: {
- ret 42i
- }
-}
)";
EXPECT_EQ(src, str());
@@ -1098,58 +921,147 @@
%vert_loc1_Input_1:ptr<__in, i32, read> = var undef @location(1) # %vert_loc1_Input_1: 'vert_loc1_Input'
%vert_position:ptr<__out, vec4<f32>, write> = var undef @invariant @builtin(position)
%vert___point_size:ptr<__out, f32, write> = var undef @builtin(__point_size)
- %frag1_loc1_Output:ptr<__out, f32, write> = var undef @location(1)
- %frag2_loc0_Output:ptr<__out, i32, write> = var undef @location(0)
}
%vert_inner = func(%input:MyStruct, %ival:i32):vec4<f32> {
$B2: {
- %10:vec4<f32> = construct 0.5f
- ret %10
- }
-}
-%frag1_inner = func():MyStruct {
- $B3: {
- %12:MyStruct = construct 0.5f
- ret %12
- }
-}
-%frag2_inner = func():i32 {
- $B4: {
- ret 42i
+ %8:vec4<f32> = construct 0.5f
+ ret %8
}
}
%vert = @vertex func():void {
- $B5: {
- %15:f32 = load %vert_loc1_Input
- %16:MyStruct = construct %15
- %17:i32 = load %vert_loc1_Input_1
- %18:vec4<f32> = call %vert_inner, %16, %17
- %19:f32 = swizzle %18, x
- %20:f32 = swizzle %18, y
- %21:f32 = negation %20
- %22:f32 = swizzle %18, z
- %23:f32 = swizzle %18, w
- %24:f32 = mul 2.0f, %22
- %25:f32 = sub %24, %23
- %26:vec4<f32> = construct %19, %21, %25, %23
- store %vert_position, %26
+ $B3: {
+ %10:f32 = load %vert_loc1_Input
+ %11:MyStruct = construct %10
+ %12:i32 = load %vert_loc1_Input_1
+ %13:vec4<f32> = call %vert_inner, %11, %12
+ %14:f32 = swizzle %13, x
+ %15:f32 = swizzle %13, y
+ %16:f32 = negation %15
+ %17:f32 = swizzle %13, z
+ %18:f32 = swizzle %13, w
+ %19:f32 = mul 2.0f, %17
+ %20:f32 = sub %19, %18
+ %21:vec4<f32> = construct %14, %16, %20, %18
+ store %vert_position, %21
store %vert___point_size, 1.0f
ret
}
}
+)";
+
+ core::ir::transform::ImmediateDataLayout immediate_data;
+ ShaderIOConfig config{immediate_data};
+ Run(ShaderIO, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+// Test that interpolation attributes are stripped from fragment struct outputs
+TEST_F(GlslWriter_ShaderIOTest, InterpolationOnFragmentOutput_Struct) {
+ auto* str_ty = ty.Struct(mod.symbols.New("MyStruct"),
+ {
+ {
+ mod.symbols.New("color"),
+ ty.f32(),
+ core::IOAttributes{
+ .location = 1u,
+ .interpolation =
+ core::Interpolation{
+ core::InterpolationType::kLinear,
+ core::InterpolationSampling::kSample,
+ },
+ },
+ },
+ });
+
+ auto* ep = b.Function("frag1", str_ty);
+ ep->SetStage(core::ir::Function::PipelineStage::kFragment);
+
+ b.Append(ep->Block(), [&] { //
+ b.Return(ep, b.Construct(str_ty, 0.5_f));
+ });
+
+ auto* src = R"(
+MyStruct = struct @align(4) {
+ color:f32 @offset(0), @location(1), @interpolate(linear, sample)
+}
+
+%frag1 = @fragment func():MyStruct {
+ $B1: {
+ %2:MyStruct = construct 0.5f
+ ret %2
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+MyStruct = struct @align(4) {
+ color:f32 @offset(0)
+}
+
+$B1: { # root
+ %frag1_loc1_Output:ptr<__out, f32, write> = var undef @location(1)
+}
+
+%frag1_inner = func():MyStruct {
+ $B2: {
+ %3:MyStruct = construct 0.5f
+ ret %3
+ }
+}
%frag1 = @fragment func():void {
- $B6: {
- %28:MyStruct = call %frag1_inner
- %29:f32 = access %28, 0u
- store %frag1_loc1_Output, %29
+ $B3: {
+ %5:MyStruct = call %frag1_inner
+ %6:f32 = access %5, 0u
+ store %frag1_loc1_Output, %6
ret
}
}
+)";
+
+ core::ir::transform::ImmediateDataLayout immediate_data;
+ ShaderIOConfig config{immediate_data};
+ Run(ShaderIO, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+// Test that interpolation attributes are stripped from fragment struct outputs
+TEST_F(GlslWriter_ShaderIOTest, InterpolationOnFragmentOutput_NonStruct) {
+ auto* ep = b.Function("frag2", ty.i32());
+ ep->SetStage(core::ir::Function::PipelineStage::kFragment);
+ ep->SetReturnLocation(0);
+ ep->SetReturnInterpolation(core::Interpolation{core::InterpolationType::kFlat});
+
+ b.Append(ep->Block(), [&] { //
+ b.Return(ep, b.Constant(42_i));
+ });
+
+ auto* src = R"(
+%frag2 = @fragment func():i32 [@location(0), @interpolate(flat)] {
+ $B1: {
+ ret 42i
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+$B1: { # root
+ %frag2_loc0_Output:ptr<__out, i32, write> = var undef @location(0)
+}
+
+%frag2_inner = func():i32 {
+ $B2: {
+ ret 42i
+ }
+}
%frag2 = @fragment func():void {
- $B7: {
- %31:i32 = call %frag2_inner
- store %frag2_loc0_Output, %31
+ $B3: {
+ %4:i32 = call %frag2_inner
+ store %frag2_loc0_Output, %4
ret
}
}
@@ -1255,161 +1167,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(GlslWriter_ShaderIOTest, ClampFragDepth_MultipleFragmentShaders) {
- auto* str_ty =
- ty.Struct(mod.symbols.New("Outputs"), {
- {
- mod.symbols.New("color"),
- ty.f32(),
- core::IOAttributes{
- .location = 0u,
- },
- },
- {
- mod.symbols.New("depth"),
- ty.f32(),
- core::IOAttributes{
- .builtin = core::BuiltinValue::kFragDepth,
- },
- },
- });
-
- auto make_entry_point = [&](std::string_view name) {
- auto* ep = b.Function(name, str_ty);
- ep->SetStage(core::ir::Function::PipelineStage::kFragment);
- b.Append(ep->Block(), [&] { //
- b.Return(ep, b.Construct(str_ty, 0.5_f, 2_f));
- });
- };
- make_entry_point("ep1");
- make_entry_point("ep2");
- make_entry_point("ep3");
-
- auto* src = R"(
-Outputs = struct @align(4) {
- color:f32 @offset(0), @location(0)
- depth:f32 @offset(4), @builtin(frag_depth)
-}
-
-%ep1 = @fragment func():Outputs {
- $B1: {
- %2:Outputs = construct 0.5f, 2.0f
- ret %2
- }
-}
-%ep2 = @fragment func():Outputs {
- $B2: {
- %4:Outputs = construct 0.5f, 2.0f
- ret %4
- }
-}
-%ep3 = @fragment func():Outputs {
- $B3: {
- %6:Outputs = construct 0.5f, 2.0f
- ret %6
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-Outputs = struct @align(4) {
- color:f32 @offset(0)
- depth:f32 @offset(4)
-}
-
-tint_immediate_data_struct = struct @align(4), @block {
- depth_min:f32 @offset(4)
- depth_max:f32 @offset(8)
-}
-
-$B1: { # root
- %tint_immediate_data:ptr<immediate, tint_immediate_data_struct, read> = var undef
- %ep1_loc0_Output:ptr<__out, f32, write> = var undef @location(0)
- %ep1_frag_depth:ptr<__out, f32, write> = var undef @builtin(frag_depth)
- %ep2_loc0_Output:ptr<__out, f32, write> = var undef @location(0)
- %ep2_frag_depth:ptr<__out, f32, write> = var undef @builtin(frag_depth)
- %ep3_loc0_Output:ptr<__out, f32, write> = var undef @location(0)
- %ep3_frag_depth:ptr<__out, f32, write> = var undef @builtin(frag_depth)
-}
-
-%ep1_inner = func():Outputs {
- $B2: {
- %9:Outputs = construct 0.5f, 2.0f
- ret %9
- }
-}
-%ep2_inner = func():Outputs {
- $B3: {
- %11:Outputs = construct 0.5f, 2.0f
- ret %11
- }
-}
-%ep3_inner = func():Outputs {
- $B4: {
- %13:Outputs = construct 0.5f, 2.0f
- ret %13
- }
-}
-%ep1 = @fragment func():void {
- $B5: {
- %15:Outputs = call %ep1_inner
- %16:f32 = access %15, 0u
- store %ep1_loc0_Output, %16
- %17:f32 = access %15, 1u
- %18:ptr<immediate, f32, read> = access %tint_immediate_data, 0u
- %19:f32 = load %18
- %20:ptr<immediate, f32, read> = access %tint_immediate_data, 1u
- %21:f32 = load %20
- %22:f32 = clamp %17, %19, %21
- store %ep1_frag_depth, %22
- ret
- }
-}
-%ep2 = @fragment func():void {
- $B6: {
- %24:Outputs = call %ep2_inner
- %25:f32 = access %24, 0u
- store %ep2_loc0_Output, %25
- %26:f32 = access %24, 1u
- %27:ptr<immediate, f32, read> = access %tint_immediate_data, 0u
- %28:f32 = load %27
- %29:ptr<immediate, f32, read> = access %tint_immediate_data, 1u
- %30:f32 = load %29
- %31:f32 = clamp %26, %28, %30
- store %ep2_frag_depth, %31
- ret
- }
-}
-%ep3 = @fragment func():void {
- $B7: {
- %33:Outputs = call %ep3_inner
- %34:f32 = access %33, 0u
- store %ep3_loc0_Output, %34
- %35:f32 = access %33, 1u
- %36:ptr<immediate, f32, read> = access %tint_immediate_data, 0u
- %37:f32 = load %36
- %38:ptr<immediate, f32, read> = access %tint_immediate_data, 1u
- %39:f32 = load %38
- %40:f32 = clamp %35, %37, %39
- store %ep3_frag_depth, %40
- ret
- }
-}
-)";
-
- core::ir::transform::PrepareImmediateDataConfig immediate_data_config;
- immediate_data_config.AddInternalImmediateData(4, mod.symbols.New("depth_min"), ty.f32());
- immediate_data_config.AddInternalImmediateData(8, mod.symbols.New("depth_max"), ty.f32());
- auto immediate_data = PrepareImmediateData(mod, immediate_data_config);
- EXPECT_EQ(immediate_data, Success);
- ShaderIOConfig config{immediate_data.Get()};
- config.depth_range_offsets = {4, 8};
- Run(ShaderIO, config);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(GlslWriter_ShaderIOTest, BGRASwizzleSingleValue) {
auto* ep = b.Function("vert", ty.vec4<f32>());
ep->SetReturnBuiltin(core::BuiltinValue::kPosition);
diff --git a/src/tint/lang/hlsl/writer/function_test.cc b/src/tint/lang/hlsl/writer/function_test.cc
index ad8ead0..e200c46 100644
--- a/src/tint/lang/hlsl/writer/function_test.cc
+++ b/src/tint/lang/hlsl/writer/function_test.cc
@@ -227,196 +227,6 @@
)");
}
-TEST_F(HlslWriterTest, FunctionEntryPointSharedStructDifferentStages) {
- // struct Interface {
- // @builtin(position) pos : vec4<f32>;
- // @location(1) col1 : f32;
- // @location(2) col2 : f32;
- // };
- //
- // fn vert_main() -> Interface {
- // return Interface(vec4<f32>(), 0.4, 0.6);
- // }
- //
- // fn frag_main(inputs : Interface) {
- // const r = inputs.col1;
- // const g = inputs.col2;
- // const p = inputs.pos;
- // }
-
- core::IOAttributes pos_attrs{};
- pos_attrs.builtin = core::BuiltinValue::kPosition;
- core::IOAttributes col1_attrs{};
- col1_attrs.location = 1;
- core::IOAttributes col2_attrs{};
- col2_attrs.location = 2;
-
- Vector members{
- ty.Get<core::type::StructMember>(b.ir.symbols.New("pos"), ty.vec4<f32>(), 0u, 0u, 16u, 16u,
- pos_attrs),
- ty.Get<core::type::StructMember>(b.ir.symbols.New("col1"), ty.f32(), 1u, 16u, 4u, 4u,
- col1_attrs),
- ty.Get<core::type::StructMember>(b.ir.symbols.New("col2"), ty.f32(), 2u, 16u, 4u, 4u,
- col2_attrs),
- };
- auto* strct = ty.Struct(b.ir.symbols.New("Interface"), std::move(members));
-
- auto* vert_func = b.Function("vert_main", strct, core::ir::Function::PipelineStage::kVertex);
- b.Append(vert_func->Block(), [&] { //
- b.Return(vert_func, b.Construct(strct, b.Zero(ty.vec4<f32>()), 0.5_f, 0.25_f));
- });
-
- auto* frag_param = b.FunctionParam("inputs", strct);
- auto* frag_func =
- b.Function("frag_main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
- frag_func->SetParams({frag_param});
- b.Append(frag_func->Block(), [&] {
- auto* r = b.Access(ty.f32(), frag_param, 1_u);
- auto* g = b.Access(ty.f32(), frag_param, 2_u);
- auto* p = b.Access(ty.vec4<f32>(), frag_param, 0_u);
-
- b.Let("r", r);
- b.Let("g", g);
- b.Let("p", p);
- b.Return(frag_func);
- });
-
- ASSERT_TRUE(Generate()) << err_ << output_.hlsl;
- EXPECT_EQ(output_.hlsl, R"(struct Interface {
- float4 pos;
- float col1;
- float col2;
-};
-
-struct vert_main_outputs {
- float Interface_col1 : TEXCOORD1;
- float Interface_col2 : TEXCOORD2;
- float4 Interface_pos : SV_Position;
-};
-
-struct frag_main_inputs {
- float Interface_col1 : TEXCOORD1;
- float Interface_col2 : TEXCOORD2;
- float4 Interface_pos : SV_Position;
-};
-
-
-Interface vert_main_inner() {
- Interface v = {(0.0f).xxxx, 0.5f, 0.25f};
- return v;
-}
-
-void frag_main_inner(Interface inputs) {
- float r = inputs.col1;
- float g = inputs.col2;
- float4 p = inputs.pos;
-}
-
-vert_main_outputs vert_main() {
- Interface v_1 = vert_main_inner();
- vert_main_outputs v_2 = {v_1.col1, v_1.col2, v_1.pos};
- return v_2;
-}
-
-void frag_main(frag_main_inputs inputs) {
- Interface v_3 = {float4(inputs.Interface_pos.xyz, (1.0f / inputs.Interface_pos.w)), inputs.Interface_col1, inputs.Interface_col2};
- frag_main_inner(v_3);
-}
-
-)");
- EXPECT_EQ(0u, output_.workgroup_info.x);
- EXPECT_EQ(0u, output_.workgroup_info.y);
- EXPECT_EQ(0u, output_.workgroup_info.z);
-}
-
-TEST_F(HlslWriterTest, FunctionEntryPointSharedStructHelperFunction) {
- // struct VertexOutput {
- // @builtin(position) pos : vec4<f32>;
- // };
- // fn foo(x : f32) -> VertexOutput {
- // return VertexOutput(vec4<f32>(x, x, x, 1.0));
- // }
- // fn vert_main1() -> VertexOutput {
- // return foo(0.5);
- // }
- // fn vert_main2() -> VertexOutput {
- // return foo(0.25);
- // }
-
- core::IOAttributes pos_attrs{};
- pos_attrs.builtin = core::BuiltinValue::kPosition;
-
- Vector members{ty.Get<core::type::StructMember>(b.ir.symbols.New("pos"), ty.vec4<f32>(), 0u, 0u,
- 16u, 16u, pos_attrs)};
- auto* strct = ty.Struct(b.ir.symbols.New("VertexOutput"), std::move(members));
-
- auto* x = b.FunctionParam("x", ty.f32());
- auto* foo_func = b.Function("foo", strct);
- foo_func->SetParams({x});
- b.Append(foo_func->Block(), [&] { //
- b.Return(foo_func, b.Construct(strct, b.Construct(ty.vec4<f32>(), x, x, x, 1_f)));
- });
-
- {
- auto* vert1_func =
- b.Function("vert1_main1", strct, core::ir::Function::PipelineStage::kVertex);
- b.Append(vert1_func->Block(), [&] { //
- b.Return(vert1_func, b.Call(foo_func, 0.5_f));
- });
- }
-
- {
- auto* vert2_func =
- b.Function("vert2_main1", strct, core::ir::Function::PipelineStage::kVertex);
- b.Append(vert2_func->Block(), [&] { //
- b.Return(vert2_func, b.Call(foo_func, 0.25_f));
- });
- }
-
- ASSERT_TRUE(Generate()) << err_ << output_.hlsl;
- EXPECT_EQ(output_.hlsl, R"(struct VertexOutput {
- float4 pos;
-};
-
-struct vert1_main1_outputs {
- float4 VertexOutput_pos : SV_Position;
-};
-
-struct vert2_main1_outputs {
- float4 VertexOutput_pos : SV_Position;
-};
-
-
-VertexOutput foo(float x) {
- VertexOutput v = {float4(x, x, x, 1.0f)};
- return v;
-}
-
-VertexOutput vert1_main1_inner() {
- VertexOutput v_1 = foo(0.5f);
- return v_1;
-}
-
-VertexOutput vert2_main1_inner() {
- VertexOutput v_2 = foo(0.25f);
- return v_2;
-}
-
-vert1_main1_outputs vert1_main1() {
- VertexOutput v_3 = vert1_main1_inner();
- vert1_main1_outputs v_4 = {v_3.pos};
- return v_4;
-}
-
-vert2_main1_outputs vert2_main1() {
- VertexOutput v_5 = vert2_main1_inner();
- vert2_main1_outputs v_6 = {v_5.pos};
- return v_6;
-}
-
-)");
-}
-
TEST_F(HlslWriterTest, FunctionEntryPointWithUniform) {
// struct Uniforms {
// coord: vec4f,
@@ -1003,67 +813,6 @@
)");
}
-// https://crbug.com/tint/297
-TEST_F(HlslWriterTest, FunctionMultipleEntryPointWithSameModuleVar) {
- // struct Data {
- // d : f32;
- // };
- // @binding(0) @group(0) var<storage, read_write> data : Data;
- //
- // @compute @workgroup_size(1)
- // fn a() {
- // var v = data.d;
- // return;
- // }
- //
- // @compute @workgroup_size(1)
- // fn b() {
- // var v = data.d;
- // return;
- // }
-
- Vector members{ty.Get<core::type::StructMember>(b.ir.symbols.New("d"), ty.f32(), 0u, 0u, 4u, 4u,
- core::IOAttributes{})};
- auto* strct = ty.Struct(b.ir.symbols.New("Data"), std::move(members));
-
- auto* data = b.Var("data", storage, strct, read_write);
- data->SetBindingPoint(0, 0);
- b.ir.root_block->Append(data);
-
- {
- auto* func = b.ComputeFunction("a");
- b.Append(func->Block(), [&] { //
- auto* a = b.Access(ty.ptr<storage, f32>(), data, 0_u);
- b.Var("v", b.Load(a));
- b.Return(func);
- });
- }
-
- {
- auto* func = b.ComputeFunction("b");
- b.Append(func->Block(), [&] { //
- auto* a = b.Access(ty.ptr<storage, f32>(), data, 0_u);
- b.Var("v", b.Load(a));
- b.Return(func);
- });
- }
-
- ASSERT_TRUE(Generate()) << err_ << output_.hlsl;
- EXPECT_EQ(output_.hlsl, R"(
-RWByteAddressBuffer data : register(u0);
-[numthreads(1, 1, 1)]
-void a() {
- float v = asfloat(data.Load(0u));
-}
-
-[numthreads(1, 1, 1)]
-void b() {
- float v = asfloat(data.Load(0u));
-}
-
-)");
-}
-
TEST_F(HlslWriterTest, DuplicateConstant) {
auto* ret_arr = b.Function("ret_arr", ty.array<vec4<i32>, 4>());
b.Append(ret_arr->Block(), [&] { b.Return(ret_arr, b.Zero<array<vec4<i32>, 4>>()); });
diff --git a/src/tint/lang/hlsl/writer/raise/pixel_local_test.cc b/src/tint/lang/hlsl/writer/raise/pixel_local_test.cc
index aefddf0..a462370 100644
--- a/src/tint/lang/hlsl/writer/raise/pixel_local_test.cc
+++ b/src/tint/lang/hlsl/writer/raise/pixel_local_test.cc
@@ -512,172 +512,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(HlslWriterPixelLocalTest, MultipleEntryPoints) {
- auto* pixel_local_struct_ty =
- ty.Struct(mod.symbols.New("PixelLocal"), {
- {mod.symbols.New("a"), ty.u32()},
- });
- auto* pl = b.Var("pl", ty.ptr<pixel_local>(pixel_local_struct_ty));
- mod.root_block->Append(pl);
-
- Vector<core::type::Manager::StructMemberDesc, 3> members;
- core::IOAttributes attrs;
- attrs.builtin = core::BuiltinValue::kPosition;
- members.Emplace(mod.symbols.New("pos"), ty.vec4<f32>(), attrs);
- auto* param_struct_ty = ty.Struct(mod.symbols.New("params"), members);
-
- for (size_t i = 0; i < 3; ++i) {
- auto* func = b.Function("main" + std::to_string(i), ty.vec4<f32>(),
- core::ir::Function::PipelineStage::kFragment);
- func->SetReturnLocation(0_u);
- func->SetParams({b.FunctionParam(param_struct_ty)});
-
- b.Append(func->Block(), [&] {
- auto* access = b.Access(ty.ptr<pixel_local>(ty.u32()), pl, 0_u);
- auto* add = b.Add<u32>(b.Load(access), 42_u);
- b.Store(access, add);
- b.Return(func, b.Construct<vec4<f32>>(1_f, 0_f, 0_f, 1_f));
- });
- }
-
- auto* src = R"(
-PixelLocal = struct @align(4) {
- a:u32 @offset(0)
-}
-
-params = struct @align(16) {
- pos:vec4<f32> @offset(0), @builtin(position)
-}
-
-$B1: { # root
- %pl:ptr<pixel_local, PixelLocal, read_write> = var undef
-}
-
-%main0 = @fragment func(%3:params):vec4<f32> [@location(0)] {
- $B2: {
- %4:ptr<pixel_local, u32, read_write> = access %pl, 0u
- %5:u32 = load %4
- %6:u32 = add %5, 42u
- store %4, %6
- %7:vec4<f32> = construct 1.0f, 0.0f, 0.0f, 1.0f
- ret %7
- }
-}
-%main1 = @fragment func(%9:params):vec4<f32> [@location(0)] {
- $B3: {
- %10:ptr<pixel_local, u32, read_write> = access %pl, 0u
- %11:u32 = load %10
- %12:u32 = add %11, 42u
- store %10, %12
- %13:vec4<f32> = construct 1.0f, 0.0f, 0.0f, 1.0f
- ret %13
- }
-}
-%main2 = @fragment func(%15:params):vec4<f32> [@location(0)] {
- $B4: {
- %16:ptr<pixel_local, u32, read_write> = access %pl, 0u
- %17:u32 = load %16
- %18:u32 = add %17, 42u
- store %16, %18
- %19:vec4<f32> = construct 1.0f, 0.0f, 0.0f, 1.0f
- ret %19
- }
-}
-)";
-
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-PixelLocal = struct @align(4) {
- a:u32 @offset(0)
-}
-
-params = struct @align(16) {
- pos:vec4<f32> @offset(0), @builtin(position)
-}
-
-$B1: { # root
- %pl:ptr<private, PixelLocal, read_write> = var undef
- %pixel_local_a:ptr<handle, hlsl.rasterizer_ordered_texture_2d<r32uint>, read> = var undef @binding_point(7, 10)
-}
-
-%main0 = @fragment func(%4:params):vec4<f32> [@location(0)] {
- $B2: {
- %5:vec4<f32> = access %4, 0u
- %6:vec2<f32> = swizzle %5, xy
- %7:vec2<u32> = convert %6
- %8:hlsl.rasterizer_ordered_texture_2d<r32uint> = load %pixel_local_a
- %9:vec4<u32> = %8.Load %7
- %10:u32 = swizzle %9, x
- %11:ptr<private, u32, read_write> = access %pl, 0u
- store %11, %10
- %12:ptr<private, u32, read_write> = access %pl, 0u
- %13:u32 = load %12
- %14:u32 = add %13, 42u
- store %12, %14
- %15:vec4<f32> = construct 1.0f, 0.0f, 0.0f, 1.0f
- %16:ptr<private, u32, read_write> = access %pl, 0u
- %17:u32 = load %16
- %18:vec4<u32> = construct %17
- %19:hlsl.rasterizer_ordered_texture_2d<r32uint> = load %pixel_local_a
- %20:void = hlsl.textureStore %19, %7, %18
- ret %15
- }
-}
-%main1 = @fragment func(%22:params):vec4<f32> [@location(0)] {
- $B3: {
- %23:vec4<f32> = access %22, 0u
- %24:vec2<f32> = swizzle %23, xy
- %25:vec2<u32> = convert %24
- %26:hlsl.rasterizer_ordered_texture_2d<r32uint> = load %pixel_local_a
- %27:vec4<u32> = %26.Load %25
- %28:u32 = swizzle %27, x
- %29:ptr<private, u32, read_write> = access %pl, 0u
- store %29, %28
- %30:ptr<private, u32, read_write> = access %pl, 0u
- %31:u32 = load %30
- %32:u32 = add %31, 42u
- store %30, %32
- %33:vec4<f32> = construct 1.0f, 0.0f, 0.0f, 1.0f
- %34:ptr<private, u32, read_write> = access %pl, 0u
- %35:u32 = load %34
- %36:vec4<u32> = construct %35
- %37:hlsl.rasterizer_ordered_texture_2d<r32uint> = load %pixel_local_a
- %38:void = hlsl.textureStore %37, %25, %36
- ret %33
- }
-}
-%main2 = @fragment func(%40:params):vec4<f32> [@location(0)] {
- $B4: {
- %41:vec4<f32> = access %40, 0u
- %42:vec2<f32> = swizzle %41, xy
- %43:vec2<u32> = convert %42
- %44:hlsl.rasterizer_ordered_texture_2d<r32uint> = load %pixel_local_a
- %45:vec4<u32> = %44.Load %43
- %46:u32 = swizzle %45, x
- %47:ptr<private, u32, read_write> = access %pl, 0u
- store %47, %46
- %48:ptr<private, u32, read_write> = access %pl, 0u
- %49:u32 = load %48
- %50:u32 = add %49, 42u
- store %48, %50
- %51:vec4<f32> = construct 1.0f, 0.0f, 0.0f, 1.0f
- %52:ptr<private, u32, read_write> = access %pl, 0u
- %53:u32 = load %52
- %54:vec4<u32> = construct %53
- %55:hlsl.rasterizer_ordered_texture_2d<r32uint> = load %pixel_local_a
- %56:void = hlsl.textureStore %55, %43, %54
- ret %51
- }
-}
-)";
-
- auto config = OneArgConfig();
- Run(PixelLocal, config);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(HlslWriterPixelLocalTest, MultipleMembers) {
auto r = ThreeArgFunc();
b.Append(r.func->Block(), [&] {
diff --git a/src/tint/lang/hlsl/writer/raise/shader_io_test.cc b/src/tint/lang/hlsl/writer/raise/shader_io_test.cc
index 9f8e854..f265245 100644
--- a/src/tint/lang/hlsl/writer/raise/shader_io_test.cc
+++ b/src/tint/lang/hlsl/writer/raise/shader_io_test.cc
@@ -658,146 +658,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(HlslWriterTransformTest, ShaderIOStruct_SharedByVertexAndFragment) {
- auto* str_ty = ty.Struct(mod.symbols.New("Interface"),
- {
- {
- mod.symbols.New("position"),
- ty.vec4<f32>(),
- core::IOAttributes{
- .builtin = core::BuiltinValue::kPosition,
- },
- },
- {
- mod.symbols.New("color"),
- ty.vec3<f32>(),
- core::IOAttributes{
- .location = 0u,
- },
- },
- });
-
- // Vertex shader.
- {
- auto* ep = b.Function("vert", str_ty, core::ir::Function::PipelineStage::kVertex);
-
- b.Append(ep->Block(), [&] {
- auto* position = b.Construct(ty.vec4<f32>(), 0_f);
- auto* color = b.Construct(ty.vec3<f32>(), 1_f);
- b.Return(ep, b.Construct(str_ty, position, color));
- });
- }
-
- // Fragment shader.
- {
- auto* inputs = b.FunctionParam("inputs", str_ty);
-
- auto* ep = b.Function("frag", ty.vec4<f32>(), core::ir::Function::PipelineStage::kFragment);
- ep->SetParams({inputs});
- ep->SetReturnLocation(0u);
-
- b.Append(ep->Block(), [&] {
- auto* position = b.Access(ty.vec4<f32>(), inputs, 0_u);
- auto* color = b.Access(ty.vec3<f32>(), inputs, 1_u);
- b.Return(ep, b.Add(ty.vec4<f32>(), position, b.Construct(ty.vec4<f32>(), color, 1_f)));
- });
- }
-
- auto* src = R"(
-Interface = struct @align(16) {
- position:vec4<f32> @offset(0), @builtin(position)
- color:vec3<f32> @offset(16), @location(0)
-}
-
-%vert = @vertex func():Interface {
- $B1: {
- %2:vec4<f32> = construct 0.0f
- %3:vec3<f32> = construct 1.0f
- %4:Interface = construct %2, %3
- ret %4
- }
-}
-%frag = @fragment func(%inputs:Interface):vec4<f32> [@location(0)] {
- $B2: {
- %7:vec4<f32> = access %inputs, 0u
- %8:vec3<f32> = access %inputs, 1u
- %9:vec4<f32> = construct %8, 1.0f
- %10:vec4<f32> = add %7, %9
- ret %10
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-Interface = struct @align(16) {
- position:vec4<f32> @offset(0)
- color:vec3<f32> @offset(16)
-}
-
-vert_outputs = struct @align(16) {
- Interface_color:vec3<f32> @offset(0), @location(0)
- Interface_position:vec4<f32> @offset(16), @builtin(position)
-}
-
-frag_inputs = struct @align(16) {
- Interface_color:vec3<f32> @offset(0), @location(0)
- Interface_position:vec4<f32> @offset(16), @builtin(position)
-}
-
-frag_outputs = struct @align(16) {
- tint_symbol:vec4<f32> @offset(0), @location(0)
-}
-
-%vert_inner = func():Interface {
- $B1: {
- %2:vec4<f32> = construct 0.0f
- %3:vec3<f32> = construct 1.0f
- %4:Interface = construct %2, %3
- ret %4
- }
-}
-%frag_inner = func(%inputs:Interface):vec4<f32> {
- $B2: {
- %7:vec4<f32> = access %inputs, 0u
- %8:vec3<f32> = access %inputs, 1u
- %9:vec4<f32> = construct %8, 1.0f
- %10:vec4<f32> = add %7, %9
- ret %10
- }
-}
-%vert = @vertex func():vert_outputs {
- $B3: {
- %12:Interface = call %vert_inner
- %13:vec4<f32> = access %12, 0u
- %14:vec3<f32> = access %12, 1u
- %15:vert_outputs = construct %14, %13
- ret %15
- }
-}
-%frag = @fragment func(%inputs_1:frag_inputs):frag_outputs { # %inputs_1: 'inputs'
- $B4: {
- %18:vec4<f32> = access %inputs_1, 1u
- %19:f32 = access %18, 3u
- %20:f32 = div 1.0f, %19
- %21:vec3<f32> = swizzle %18, xyz
- %22:vec4<f32> = construct %21, %20
- %23:vec3<f32> = access %inputs_1, 0u
- %24:Interface = construct %22, %23
- %25:vec4<f32> = call %frag_inner, %24
- %26:frag_outputs = construct %25
- ret %26
- }
-}
-)";
-
- core::ir::transform::ImmediateDataLayout immediate_data;
- ShaderIOConfig config{immediate_data};
- Run(ShaderIO, config);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(HlslWriterTransformTest, ShaderIOStruct_SharedWithBuffer) {
auto* str_ty =
ty.Struct(mod.symbols.New("Outputs"), {
@@ -3649,279 +3509,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(HlslWriterTransformTest,
- ShaderIOParameters_TruncateInterstage_MultipleEntryPointsSharedStruct) {
- core::IOAttributes pos_attr;
- pos_attr.builtin = core::BuiltinValue::kPosition;
-
- core::IOAttributes loc0_attr;
- loc0_attr.location = 0;
- core::IOAttributes loc1_attr;
- loc1_attr.location = 1;
- core::IOAttributes loc2_attr;
- loc2_attr.location = 2;
-
- auto* str_ty = ty.Struct(mod.symbols.New("Outputs"),
- {
- {mod.symbols.New("position"), ty.vec4<f32>(), pos_attr},
- {mod.symbols.New("loc0"), ty.f32(), loc0_attr},
- {mod.symbols.New("loc1"), ty.i32(), loc1_attr},
- {mod.symbols.New("loc2"), ty.vec3<i32>(), loc2_attr},
- });
- auto* ep1 = b.Function("foo1", str_ty, core::ir::Function::PipelineStage::kVertex);
- b.Append(ep1->Block(), [&] {
- auto* pos = b.Construct(ty.vec4<f32>(), 0.5_f);
- auto* loc2 = b.Construct(ty.vec3<i32>(), 3_i);
- b.Return(ep1, b.Construct(str_ty, pos, 1_f, 2_i, loc2));
- });
- auto* ep2 = b.Function("foo2", str_ty, core::ir::Function::PipelineStage::kVertex);
- b.Append(ep2->Block(), [&] {
- auto* pos = b.Construct(ty.vec4<f32>(), 0.5_f);
- auto* loc2 = b.Construct(ty.vec3<i32>(), 3_i);
- b.Return(ep2, b.Construct(str_ty, pos, 1_f, 2_i, loc2));
- });
-
- auto* src = R"(
-Outputs = struct @align(16) {
- position:vec4<f32> @offset(0), @builtin(position)
- loc0:f32 @offset(16), @location(0)
- loc1:i32 @offset(20), @location(1)
- loc2:vec3<i32> @offset(32), @location(2)
-}
-
-%foo1 = @vertex func():Outputs {
- $B1: {
- %2:vec4<f32> = construct 0.5f
- %3:vec3<i32> = construct 3i
- %4:Outputs = construct %2, 1.0f, 2i, %3
- ret %4
- }
-}
-%foo2 = @vertex func():Outputs {
- $B2: {
- %6:vec4<f32> = construct 0.5f
- %7:vec3<i32> = construct 3i
- %8:Outputs = construct %6, 1.0f, 2i, %7
- ret %8
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-Outputs = struct @align(16) {
- position:vec4<f32> @offset(0)
- loc0:f32 @offset(16)
- loc1:i32 @offset(20)
- loc2:vec3<i32> @offset(32)
-}
-
-foo1_outputs = struct @align(16) {
- Outputs_loc0:f32 @offset(0), @location(0)
- Outputs_loc2:vec3<i32> @offset(16), @location(2)
- Outputs_position:vec4<f32> @offset(32), @builtin(position)
-}
-
-foo2_outputs = struct @align(16) {
- Outputs_loc0:f32 @offset(0), @location(0)
- Outputs_loc2:vec3<i32> @offset(16), @location(2)
- Outputs_position:vec4<f32> @offset(32), @builtin(position)
-}
-
-%foo1_inner = func():Outputs {
- $B1: {
- %2:vec4<f32> = construct 0.5f
- %3:vec3<i32> = construct 3i
- %4:Outputs = construct %2, 1.0f, 2i, %3
- ret %4
- }
-}
-%foo2_inner = func():Outputs {
- $B2: {
- %6:vec4<f32> = construct 0.5f
- %7:vec3<i32> = construct 3i
- %8:Outputs = construct %6, 1.0f, 2i, %7
- ret %8
- }
-}
-%foo1 = @vertex func():foo1_outputs {
- $B3: {
- %10:Outputs = call %foo1_inner
- %11:vec4<f32> = access %10, 0u
- %12:f32 = access %10, 1u
- %13:i32 = access %10, 2u
- %14:vec3<i32> = access %10, 3u
- %15:foo1_outputs = construct %12, %14, %11
- ret %15
- }
-}
-%foo2 = @vertex func():foo2_outputs {
- $B4: {
- %17:Outputs = call %foo2_inner
- %18:vec4<f32> = access %17, 0u
- %19:f32 = access %17, 1u
- %20:i32 = access %17, 2u
- %21:vec3<i32> = access %17, 3u
- %22:foo2_outputs = construct %19, %21, %18
- ret %22
- }
-}
-)";
-
- core::ir::transform::ImmediateDataLayout immediate_data;
- ShaderIOConfig config{immediate_data};
- config.truncate_interstage_variables = true;
- config.interstage_locations[0] = true;
- config.interstage_locations[2] = true;
- Run(ShaderIO, config);
-
- EXPECT_EQ(expect, str());
-}
-
-TEST_F(HlslWriterTransformTest,
- ShaderIOParameters_TruncateInterstage_MultipleEntryPointsSeparateStruct) {
- core::IOAttributes pos_attr;
- pos_attr.builtin = core::BuiltinValue::kPosition;
-
- core::IOAttributes loc1_attr;
- loc1_attr.location = 1;
- core::IOAttributes loc2_attr;
- loc2_attr.location = 2;
- core::IOAttributes loc3_attr;
- loc3_attr.location = 3;
- core::IOAttributes loc5_attr;
- loc5_attr.location = 5;
-
- auto* str_ty1 = ty.Struct(mod.symbols.New("Outputs1"),
- {
- {mod.symbols.New("position"), ty.vec4<f32>(), pos_attr},
- {mod.symbols.New("loc1"), ty.f32(), loc1_attr},
- {mod.symbols.New("loc3"), ty.i32(), loc3_attr},
- {mod.symbols.New("loc5"), ty.vec3<i32>(), loc5_attr},
- });
- auto* ep1 = b.Function("foo1", str_ty1, core::ir::Function::PipelineStage::kVertex);
- b.Append(ep1->Block(), [&] {
- auto* pos = b.Construct(ty.vec4<f32>(), 0.5_f);
- auto* loc5 = b.Construct(ty.vec3<i32>(), 3_i);
- b.Return(ep1, b.Construct(str_ty1, pos, 1_f, 2_i, loc5));
- });
-
- auto* str_ty2 = ty.Struct(mod.symbols.New("Outputs2"),
- {
- {mod.symbols.New("position"), ty.vec4<f32>(), pos_attr},
- {mod.symbols.New("loc2"), ty.vec3<i32>(), loc2_attr},
- });
-
- auto* ep2 = b.Function("foo2", str_ty2, core::ir::Function::PipelineStage::kVertex);
- b.Append(ep2->Block(), [&] {
- auto* pos = b.Construct(ty.vec4<f32>(), 0.5_f);
- auto* loc2 = b.Construct(ty.vec3<i32>(), 3_i);
- b.Return(ep2, b.Construct(str_ty2, pos, loc2));
- });
-
- auto* src = R"(
-Outputs1 = struct @align(16) {
- position:vec4<f32> @offset(0), @builtin(position)
- loc1:f32 @offset(16), @location(1)
- loc3:i32 @offset(20), @location(3)
- loc5:vec3<i32> @offset(32), @location(5)
-}
-
-Outputs2 = struct @align(16) {
- position_1:vec4<f32> @offset(0), @builtin(position)
- loc2:vec3<i32> @offset(16), @location(2)
-}
-
-%foo1 = @vertex func():Outputs1 {
- $B1: {
- %2:vec4<f32> = construct 0.5f
- %3:vec3<i32> = construct 3i
- %4:Outputs1 = construct %2, 1.0f, 2i, %3
- ret %4
- }
-}
-%foo2 = @vertex func():Outputs2 {
- $B2: {
- %6:vec4<f32> = construct 0.5f
- %7:vec3<i32> = construct 3i
- %8:Outputs2 = construct %6, %7
- ret %8
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-Outputs1 = struct @align(16) {
- position:vec4<f32> @offset(0)
- loc1:f32 @offset(16)
- loc3:i32 @offset(20)
- loc5:vec3<i32> @offset(32)
-}
-
-Outputs2 = struct @align(16) {
- position_1:vec4<f32> @offset(0)
- loc2:vec3<i32> @offset(16)
-}
-
-foo1_outputs = struct @align(16) {
- Outputs1_loc3:i32 @offset(0), @location(3)
- Outputs1_position:vec4<f32> @offset(16), @builtin(position)
-}
-
-foo2_outputs = struct @align(16) {
- Outputs2_loc2:vec3<i32> @offset(0), @location(2)
- Outputs2_position_1:vec4<f32> @offset(16), @builtin(position)
-}
-
-%foo1_inner = func():Outputs1 {
- $B1: {
- %2:vec4<f32> = construct 0.5f
- %3:vec3<i32> = construct 3i
- %4:Outputs1 = construct %2, 1.0f, 2i, %3
- ret %4
- }
-}
-%foo2_inner = func():Outputs2 {
- $B2: {
- %6:vec4<f32> = construct 0.5f
- %7:vec3<i32> = construct 3i
- %8:Outputs2 = construct %6, %7
- ret %8
- }
-}
-%foo1 = @vertex func():foo1_outputs {
- $B3: {
- %10:Outputs1 = call %foo1_inner
- %11:vec4<f32> = access %10, 0u
- %12:f32 = access %10, 1u
- %13:i32 = access %10, 2u
- %14:vec3<i32> = access %10, 3u
- %15:foo1_outputs = construct %13, %11
- ret %15
- }
-}
-%foo2 = @vertex func():foo2_outputs {
- $B4: {
- %17:Outputs2 = call %foo2_inner
- %18:vec4<f32> = access %17, 0u
- %19:vec3<i32> = access %17, 1u
- %20:foo2_outputs = construct %19, %18
- ret %20
- }
-}
-)";
-
- core::ir::transform::ImmediateDataLayout immediate_data;
- ShaderIOConfig config{immediate_data};
- config.truncate_interstage_variables = true;
- config.interstage_locations[2] = true;
- config.interstage_locations[3] = true;
- Run(ShaderIO, config);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(HlslWriterTransformTest, ShaderIOParameters_FirstIndexOffset_VertexIndex) {
auto* vert_idx = b.FunctionParam("vert_idx", ty.u32());
vert_idx->SetBuiltin(core::BuiltinValue::kVertexIndex);
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars.cc b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
index 87fe418..e4d709f 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
@@ -77,8 +77,7 @@
}
// Create the structure to hold all module-scope variables.
- // This includes all variables declared in the module, even those that are unused by one or
- // more entry points.
+ // This includes all variables declared in the module, even those that are unused.
CreateStruct();
// Process functions in reverse-dependency order (i.e. root to leaves).
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
index 37cc4f2..f0215bd1 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
@@ -1326,324 +1326,6 @@
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 undef @binding_point(1, 2)
- %b:ptr<storage, i32, read_write> = var undef @binding_point(3, 4)
- %c:ptr<private, i32, read_write> = var undef
-}
-
-%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 undef
- %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 undef # %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 undef @binding_point(1, 2)
- %b:ptr<storage, i32, read_write> = var undef @binding_point(3, 4)
- %c:ptr<private, i32, read_write> = var undef
-}
-
-%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, unused
- %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 undef
- %15:tint_module_vars_struct = construct %a_1, unused, %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 undef @binding_point(1, 2)
- %b:ptr<storage, i32, read_write> = var undef @binding_point(3, 4)
- %c:ptr<private, i32, read_write> = var undef
-}
-
-%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, unused
- %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 undef
- %18:tint_module_vars_struct = construct %a_1, unused, %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>());
diff --git a/src/tint/lang/msl/writer/raise/shader_io.cc b/src/tint/lang/msl/writer/raise/shader_io.cc
index 8c53bfd..1e18776 100644
--- a/src/tint/lang/msl/writer/raise/shader_io.cc
+++ b/src/tint/lang/msl/writer/raise/shader_io.cc
@@ -41,12 +41,6 @@
namespace {
-/// State that persists across the whole module and can be shared between entry points.
-struct PerModuleState {
- /// The frag_depth clamp arguments.
- core::ir::Value* frag_depth_clamp_args = nullptr;
-};
-
/// PIMPL state for the parts of the shader IO transform specific to MSL.
/// For MSL, we take builtin inputs as entry point parameters, move non-builtin inputs to a struct
/// passed as an entry point parameter, and wrap outputs in a structure returned by the entry point.
@@ -54,9 +48,6 @@
/// The configuration options.
const ShaderIOConfig& config;
- /// The per-module state object.
- PerModuleState& module_state;
-
/// The input parameters of the entry point.
Vector<core::ir::FunctionParam*, 4> input_params;
@@ -77,11 +68,8 @@
std::optional<uint32_t> fixed_sample_mask_index;
/// Constructor
- StateImpl(core::ir::Module& mod,
- core::ir::Function* f,
- const ShaderIOConfig& cfg,
- PerModuleState& mod_state)
- : ShaderIOBackendState(mod, f), config(cfg), module_state(mod_state) {}
+ StateImpl(core::ir::Module& mod, core::ir::Function* f, const ShaderIOConfig& cfg)
+ : ShaderIOBackendState(mod, f), config(cfg) {}
/// Destructor
~StateImpl() override {}
@@ -242,9 +230,8 @@
return result;
}
- PerModuleState module_state;
core::ir::transform::RunShaderIOBase(ir, [&](core::ir::Module& mod, core::ir::Function* func) {
- return std::make_unique<StateImpl>(mod, func, config, module_state);
+ return std::make_unique<StateImpl>(mod, func, config);
});
return Success;
diff --git a/src/tint/lang/msl/writer/raise/shader_io_test.cc b/src/tint/lang/msl/writer/raise/shader_io_test.cc
index 74e80d3..97adaf9 100644
--- a/src/tint/lang/msl/writer/raise/shader_io_test.cc
+++ b/src/tint/lang/msl/writer/raise/shader_io_test.cc
@@ -655,147 +655,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(MslWriter_ShaderIOTest, Struct_SharedByVertexAndFragment) {
- auto* vec4f = ty.vec4<f32>();
- auto* str_ty = ty.Struct(mod.symbols.New("Interface"),
- {
- {
- mod.symbols.New("position"),
- vec4f,
- core::IOAttributes{
- .builtin = core::BuiltinValue::kPosition,
- },
- },
- {
- mod.symbols.New("color"),
- vec4f,
- core::IOAttributes{
- .location = 0u,
- },
- },
- });
-
- // Vertex shader.
- {
- auto* ep = b.Function("vert", str_ty);
- ep->SetStage(core::ir::Function::PipelineStage::kVertex);
-
- b.Append(ep->Block(), [&] { //
- auto* position = b.Construct(vec4f, 0_f);
- auto* color = b.Construct(vec4f, 1_f);
- b.Return(ep, b.Construct(str_ty, position, color));
- });
- }
-
- // Fragment shader.
- {
- auto* ep = b.Function("frag", vec4f);
- auto* inputs = b.FunctionParam("inputs", str_ty);
- ep->SetStage(core::ir::Function::PipelineStage::kFragment);
- ep->SetParams({inputs});
- ep->SetReturnLocation(0u);
-
- b.Append(ep->Block(), [&] { //
- auto* position = b.Access(vec4f, inputs, 0_u);
- auto* color = b.Access(vec4f, inputs, 1_u);
- b.Return(ep, b.Add(vec4f, position, color));
- });
- }
-
- auto* src = R"(
-Interface = struct @align(16) {
- position:vec4<f32> @offset(0), @builtin(position)
- color:vec4<f32> @offset(16), @location(0)
-}
-
-%vert = @vertex func():Interface {
- $B1: {
- %2:vec4<f32> = construct 0.0f
- %3:vec4<f32> = construct 1.0f
- %4:Interface = construct %2, %3
- ret %4
- }
-}
-%frag = @fragment func(%inputs:Interface):vec4<f32> [@location(0)] {
- $B2: {
- %7:vec4<f32> = access %inputs, 0u
- %8:vec4<f32> = access %inputs, 1u
- %9:vec4<f32> = add %7, %8
- ret %9
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-Interface = struct @align(16) {
- position:vec4<f32> @offset(0)
- color:vec4<f32> @offset(16)
-}
-
-vert_outputs = struct @align(16) {
- Interface_position:vec4<f32> @offset(0), @builtin(position)
- Interface_color:vec4<f32> @offset(16), @location(0)
-}
-
-frag_inputs = struct @align(16) {
- Interface_color:vec4<f32> @offset(0), @location(0)
-}
-
-frag_outputs = struct @align(16) {
- tint_symbol:vec4<f32> @offset(0), @location(0)
-}
-
-%vert_inner = func():Interface {
- $B1: {
- %2:vec4<f32> = construct 0.0f
- %3:vec4<f32> = construct 1.0f
- %4:Interface = construct %2, %3
- ret %4
- }
-}
-%frag_inner = func(%inputs:Interface):vec4<f32> {
- $B2: {
- %7:vec4<f32> = access %inputs, 0u
- %8:vec4<f32> = access %inputs, 1u
- %9:vec4<f32> = add %7, %8
- ret %9
- }
-}
-%vert = @vertex func():vert_outputs {
- $B3: {
- %11:Interface = call %vert_inner
- %12:vec4<f32> = access %11, 0u
- %13:vec4<f32> = access %11, 1u
- %tint_wrapper_result:ptr<function, vert_outputs, read_write> = var undef
- %15:ptr<function, vec4<f32>, read_write> = access %tint_wrapper_result, 0u
- store %15, %12
- %16:ptr<function, vec4<f32>, read_write> = access %tint_wrapper_result, 1u
- store %16, %13
- %17:vert_outputs = load %tint_wrapper_result
- ret %17
- }
-}
-%frag = @fragment func(%Interface_position:vec4<f32> [@position], %inputs_1:frag_inputs):frag_outputs { # %inputs_1: 'inputs'
- $B4: {
- %21:vec4<f32> = access %inputs_1, 0u
- %22:Interface = construct %Interface_position, %21
- %23:vec4<f32> = call %frag_inner, %22
- %tint_wrapper_result_1:ptr<function, frag_outputs, read_write> = var undef # %tint_wrapper_result_1: 'tint_wrapper_result'
- %25:ptr<function, vec4<f32>, read_write> = access %tint_wrapper_result_1, 0u
- store %25, %23
- %26:frag_outputs = load %tint_wrapper_result_1
- ret %26
- }
-}
-)";
-
- ShaderIOConfig config;
- Run(ShaderIO, config);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(MslWriter_ShaderIOTest, Struct_SharedWithBuffer) {
auto* vec4f = ty.vec4<f32>();
auto* str_ty =
diff --git a/src/tint/lang/msl/writer/raise/simd_ballot_test.cc b/src/tint/lang/msl/writer/raise/simd_ballot_test.cc
index 848dd7b..634be7f 100644
--- a/src/tint/lang/msl/writer/raise/simd_ballot_test.cc
+++ b/src/tint/lang/msl/writer/raise/simd_ballot_test.cc
@@ -159,19 +159,10 @@
b.Return(foo, result);
});
- auto* ep1 = b.Function("ep1", ty.void_(), core::ir::Function::PipelineStage::kFragment);
- auto* subgroup_size = b.FunctionParam("user_subgroup_size", ty.u32());
- subgroup_size->SetLocation(0);
- ep1->SetParams({subgroup_size});
- b.Append(ep1->Block(), [&] { //
- b.Call<vec4<u32>>(foo, true);
- b.Return(ep1);
- });
-
- auto* ep2 = b.Function("ep2", ty.void_(), core::ir::Function::PipelineStage::kFragment);
- b.Append(ep2->Block(), [&] { //
+ auto* ep = b.Function("ep", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] { //
b.Call<vec4<u32>>(foo, false);
- b.Return(ep2);
+ b.Return(ep);
});
auto* src = R"(
@@ -181,15 +172,9 @@
ret %3
}
}
-%ep1 = @fragment func(%user_subgroup_size:u32 [@location(0)]):void {
+%ep = @fragment func():void {
$B2: {
- %6:vec4<u32> = call %foo, true
- ret
- }
-}
-%ep2 = @fragment func():void {
- $B3: {
- %8:vec4<u32> = call %foo, false
+ %5:vec4<u32> = call %foo, false
ret
}
}
@@ -207,43 +192,28 @@
ret %4
}
}
-%ep1 = @fragment func(%user_subgroup_size:u32 [@location(0)], %tint_subgroup_size:u32 [@subgroup_size]):void {
+%ep = @fragment func(%tint_subgroup_size:u32 [@subgroup_size]):void {
$B3: {
- %9:bool = gt %tint_subgroup_size, 32u
- %10:u32 = sub 32u, %tint_subgroup_size
- %11:u32 = shr 4294967295u, %10
- %12:u32 = select %11, 4294967295u, %9
- %13:u32 = sub 64u, %tint_subgroup_size
- %14:u32 = shr 4294967295u, %13
- %15:u32 = select 0u, %14, %9
- store_vector_element %tint_subgroup_size_mask, 0u, %12
- store_vector_element %tint_subgroup_size_mask, 1u, %15
- %16:vec4<u32> = call %foo, true
- ret
- }
-}
-%ep2 = @fragment func(%tint_subgroup_size_1:u32 [@subgroup_size]):void { # %tint_subgroup_size_1: 'tint_subgroup_size'
- $B4: {
- %19:bool = gt %tint_subgroup_size_1, 32u
- %20:u32 = sub 32u, %tint_subgroup_size_1
- %21:u32 = shr 4294967295u, %20
- %22:u32 = select %21, 4294967295u, %19
- %23:u32 = sub 64u, %tint_subgroup_size_1
- %24:u32 = shr 4294967295u, %23
- %25:u32 = select 0u, %24, %19
- store_vector_element %tint_subgroup_size_mask, 0u, %22
- store_vector_element %tint_subgroup_size_mask, 1u, %25
- %26:vec4<u32> = call %foo, false
+ %8:bool = gt %tint_subgroup_size, 32u
+ %9:u32 = sub 32u, %tint_subgroup_size
+ %10:u32 = shr 4294967295u, %9
+ %11:u32 = select %10, 4294967295u, %8
+ %12:u32 = sub 64u, %tint_subgroup_size
+ %13:u32 = shr 4294967295u, %12
+ %14:u32 = select 0u, %13, %8
+ store_vector_element %tint_subgroup_size_mask, 0u, %11
+ store_vector_element %tint_subgroup_size_mask, 1u, %14
+ %15:vec4<u32> = call %foo, false
ret
}
}
%tint_subgroup_ballot = func(%pred_1:bool):vec4<u32> { # %pred_1: 'pred'
- $B5: {
- %28:vec2<u32> = msl.simd_ballot %pred_1
- %29:vec2<u32> = load %tint_subgroup_size_mask
- %30:vec2<u32> = and %28, %29
- %31:vec4<u32> = construct %30, 0u, 0u
- ret %31
+ $B4: {
+ %17:vec2<u32> = msl.simd_ballot %pred_1
+ %18:vec2<u32> = load %tint_subgroup_size_mask
+ %19:vec2<u32> = and %17, %18
+ %20:vec4<u32> = construct %19, 0u, 0u
+ ret %20
}
}
)";
diff --git a/src/tint/lang/msl/writer/writer_test.cc b/src/tint/lang/msl/writer/writer_test.cc
index 1cb9e27..895d98f 100644
--- a/src/tint/lang/msl/writer/writer_test.cc
+++ b/src/tint/lang/msl/writer/writer_test.cc
@@ -35,6 +35,29 @@
using namespace tint::core::fluent_types; // NOLINT
using namespace tint::core::number_suffixes; // NOLINT
+TEST_F(MslWriterTest, WorkgroupAllocations_NoAllocations) {
+ 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);
+
+ // No allocations, but still needs an entry in the map.
+ auto* bar = b.ComputeFunction("bar");
+ b.Append(bar->Block(), [&] { b.Return(bar); });
+
+ ASSERT_TRUE(Generate()) << err_ << output_.msl;
+ EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
+using namespace metal;
+
+kernel void bar() {
+}
+)");
+
+ ASSERT_EQ(output_.workgroup_info.allocations.size(), 1u);
+ ASSERT_EQ(output_.workgroup_info.allocations.count("bar"), 1u);
+ EXPECT_THAT(output_.workgroup_info.allocations.at("bar"), testing::ElementsAre());
+}
+
TEST_F(MslWriterTest, WorkgroupAllocations) {
auto* var_a = b.Var("a", ty.ptr<workgroup, i32>());
auto* var_b = b.Var("b", ty.ptr<workgroup, i32>());
@@ -49,10 +72,6 @@
b.Return(foo);
});
- // No allocations, but still needs an entry in the map.
- auto* bar = b.ComputeFunction("bar");
- b.Append(bar->Block(), [&] { b.Return(bar); });
-
ASSERT_TRUE(Generate()) << err_ << output_.msl;
EXPECT_EQ(output_.msl, R"(#include <metal_stdlib>
using namespace metal;
@@ -76,19 +95,14 @@
(*tint_module_vars.a) = as_type<int>((as_type<uint>((*tint_module_vars.a)) + as_type<uint>((*tint_module_vars.b))));
}
-kernel void bar() {
-}
-
kernel void foo(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {
tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.a=(&(*v).tint_symbol), .b=(&(*v).tint_symbol_1)};
foo_inner(tint_local_index, tint_module_vars);
}
)");
- ASSERT_EQ(output_.workgroup_info.allocations.size(), 2u);
+ ASSERT_EQ(output_.workgroup_info.allocations.size(), 1u);
ASSERT_EQ(output_.workgroup_info.allocations.count("foo"), 1u);
- ASSERT_EQ(output_.workgroup_info.allocations.count("bar"), 1u);
EXPECT_THAT(output_.workgroup_info.allocations.at("foo"), testing::ElementsAre(8u));
- EXPECT_THAT(output_.workgroup_info.allocations.at("bar"), testing::ElementsAre());
}
TEST_F(MslWriterTest, NeedsStorageBufferSizes_False) {
diff --git a/src/tint/lang/spirv/reader/lower/shader_io.cc b/src/tint/lang/spirv/reader/lower/shader_io.cc
index 272d800..84382b4 100644
--- a/src/tint/lang/spirv/reader/lower/shader_io.cc
+++ b/src/tint/lang/spirv/reader/lower/shader_io.cc
@@ -466,6 +466,7 @@
Result<SuccessType> ShaderIO(core::ir::Module& ir) {
auto result = ValidateAndDumpIfNeeded(ir, "spirv.ShaderIO",
core::ir::Capabilities{
+ core::ir::Capability::kAllowMultipleEntryPoints,
core::ir::Capability::kAllowOverrides,
});
if (result != Success) {
diff --git a/src/tint/lang/spirv/reader/lower/shader_io_test.cc b/src/tint/lang/spirv/reader/lower/shader_io_test.cc
index a427870..5471411 100644
--- a/src/tint/lang/spirv/reader/lower/shader_io_test.cc
+++ b/src/tint/lang/spirv/reader/lower/shader_io_test.cc
@@ -38,6 +38,9 @@
using namespace tint::core::number_suffixes; // NOLINT
class SpirvReader_ShaderIOTest : public core::ir::transform::TransformTest {
+ public:
+ void SetUp() override { capabilities.Add(core::ir::Capability::kAllowMultipleEntryPoints); }
+
protected:
core::IOAttributes BuiltinAttrs(core::BuiltinValue builtin) {
core::IOAttributes attrs;
diff --git a/src/tint/lang/spirv/reader/parser/helper_test.h b/src/tint/lang/spirv/reader/parser/helper_test.h
index b4f2577..5a0959d 100644
--- a/src/tint/lang/spirv/reader/parser/helper_test.h
+++ b/src/tint/lang/spirv/reader/parser/helper_test.h
@@ -74,6 +74,7 @@
// Validate the IR module against the capabilities supported by the SPIR-V dialect.
auto validated =
core::ir::Validate(parsed.Get(), core::ir::Capabilities{
+ core::ir::Capability::kAllowMultipleEntryPoints,
core::ir::Capability::kAllowOverrides,
core::ir::Capability::kAllowVectorElementPointer,
});
diff --git a/src/tint/lang/spirv/writer/function_test.cc b/src/tint/lang/spirv/writer/function_test.cc
index f1d780f..8a9deae 100644
--- a/src/tint/lang/spirv/writer/function_test.cc
+++ b/src/tint/lang/spirv/writer/function_test.cc
@@ -205,60 +205,6 @@
EXPECT_EQ(workgroup_info.z, 0u);
}
-TEST_F(SpirvWriterTest, Function_EntryPoint_Multiple) {
- auto* f1 = b.ComputeFunction("main1", 32_u, 4_u, 1_u);
- b.Append(f1->Block(), [&] { //
- b.Return(f1);
- });
-
- auto* f2 = b.ComputeFunction("main2", 8_u, 2_u, 16_u);
- b.Append(f2->Block(), [&] { //
- b.Return(f2);
- });
-
- auto* f3 = b.Function("main3", ty.void_(), core::ir::Function::PipelineStage::kFragment);
- b.Append(f3->Block(), [&] { //
- b.Return(f3);
- });
-
- ASSERT_TRUE(Generate()) << Error() << output_;
- EXPECT_INST(R"(
- OpEntryPoint GLCompute %main1 "main1"
- OpEntryPoint GLCompute %main2 "main2"
- OpEntryPoint Fragment %main3 "main3"
- OpExecutionMode %main1 LocalSize 32 4 1
- OpExecutionMode %main2 LocalSize 8 2 16
- OpExecutionMode %main3 OriginUpperLeft
-
- ; Debug Information
- OpName %main1 "main1" ; id %1
- OpName %main2 "main2" ; id %5
- OpName %main3 "main3" ; id %7
-
- ; Types, variables and constants
- %void = OpTypeVoid
- %3 = OpTypeFunction %void
-
- ; Function main1
- %main1 = OpFunction %void None %3
- %4 = OpLabel
- OpReturn
- OpFunctionEnd
-
- ; Function main2
- %main2 = OpFunction %void None %3
- %6 = OpLabel
- OpReturn
- OpFunctionEnd
-
- ; Function main3
- %main3 = OpFunction %void None %3
- %8 = OpLabel
- OpReturn
- OpFunctionEnd
-)");
-}
-
TEST_F(SpirvWriterTest, Function_ReturnValue) {
auto* func = b.Function("foo", ty.i32());
b.Append(func->Block(), [&] { //
diff --git a/src/tint/lang/spirv/writer/raise/shader_io.cc b/src/tint/lang/spirv/writer/raise/shader_io.cc
index e12dfb9..0732fd5 100644
--- a/src/tint/lang/spirv/writer/raise/shader_io.cc
+++ b/src/tint/lang/spirv/writer/raise/shader_io.cc
@@ -42,12 +42,6 @@
namespace {
-/// State that persists across the whole module and can be shared between entry points.
-struct PerModuleState {
- /// The frag_depth clamp arguments.
- core::ir::Value* frag_depth_clamp_args = nullptr;
-};
-
/// PIMPL state for the parts of the shader IO transform specific to SPIR-V.
/// For SPIR-V, we declare a global variable for each input and output. The wrapper entry point then
/// loads from and stores to these variables. We also modify the type of the SampleMask builtin to
@@ -61,15 +55,9 @@
/// The configuration options.
const ShaderIOConfig& config;
- /// The per-module state object.
- PerModuleState& module_state;
-
/// Constructor
- StateImpl(core::ir::Module& mod,
- core::ir::Function* f,
- const ShaderIOConfig& cfg,
- PerModuleState& mod_state)
- : ShaderIOBackendState(mod, f), config(cfg), module_state(mod_state) {}
+ StateImpl(core::ir::Module& mod, core::ir::Function* f, const ShaderIOConfig& cfg)
+ : ShaderIOBackendState(mod, f), config(cfg) {}
/// Destructor
~StateImpl() override {}
@@ -212,9 +200,8 @@
return result;
}
- PerModuleState module_state;
core::ir::transform::RunShaderIOBase(ir, [&](core::ir::Module& mod, core::ir::Function* func) {
- return std::make_unique<StateImpl>(mod, func, config, module_state);
+ return std::make_unique<StateImpl>(mod, func, config);
});
return Success;
diff --git a/src/tint/lang/spirv/writer/raise/shader_io_test.cc b/src/tint/lang/spirv/writer/raise/shader_io_test.cc
index 5acc03a..9d9126a 100644
--- a/src/tint/lang/spirv/writer/raise/shader_io_test.cc
+++ b/src/tint/lang/spirv/writer/raise/shader_io_test.cc
@@ -660,137 +660,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvWriter_ShaderIOTest, Struct_SharedByVertexAndFragment) {
- auto* vec4f = ty.vec4<f32>();
- auto* str_ty = ty.Struct(mod.symbols.New("Interface"),
- {
- {
- mod.symbols.New("position"),
- vec4f,
- core::IOAttributes{
- .builtin = core::BuiltinValue::kPosition,
- },
- },
- {
- mod.symbols.New("color"),
- vec4f,
- core::IOAttributes{
- .location = 0u,
- },
- },
- });
-
- // Vertex shader.
- {
- auto* ep = b.Function("vert", str_ty);
- ep->SetStage(core::ir::Function::PipelineStage::kVertex);
-
- b.Append(ep->Block(), [&] { //
- auto* position = b.Construct(vec4f, 0_f);
- auto* color = b.Construct(vec4f, 1_f);
- b.Return(ep, b.Construct(str_ty, position, color));
- });
- }
-
- // Fragment shader.
- {
- auto* ep = b.Function("frag", vec4f);
- auto* inputs = b.FunctionParam("inputs", str_ty);
- ep->SetStage(core::ir::Function::PipelineStage::kFragment);
- ep->SetParams({inputs});
- ep->SetReturnLocation(0u);
-
- b.Append(ep->Block(), [&] { //
- auto* position = b.Access(vec4f, inputs, 0_u);
- auto* color = b.Access(vec4f, inputs, 1_u);
- b.Return(ep, b.Add(vec4f, position, color));
- });
- }
-
- auto* src = R"(
-Interface = struct @align(16) {
- position:vec4<f32> @offset(0), @builtin(position)
- color:vec4<f32> @offset(16), @location(0)
-}
-
-%vert = @vertex func():Interface {
- $B1: {
- %2:vec4<f32> = construct 0.0f
- %3:vec4<f32> = construct 1.0f
- %4:Interface = construct %2, %3
- ret %4
- }
-}
-%frag = @fragment func(%inputs:Interface):vec4<f32> [@location(0)] {
- $B2: {
- %7:vec4<f32> = access %inputs, 0u
- %8:vec4<f32> = access %inputs, 1u
- %9:vec4<f32> = add %7, %8
- ret %9
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-Interface = struct @align(16) {
- position:vec4<f32> @offset(0)
- color:vec4<f32> @offset(16)
-}
-
-$B1: { # root
- %vert_position_Output:ptr<__out, vec4<f32>, write> = var undef @builtin(position)
- %vert_loc0_Output:ptr<__out, vec4<f32>, write> = var undef @location(0)
- %frag_position_Input:ptr<__in, vec4<f32>, read> = var undef @builtin(position)
- %frag_loc0_Input:ptr<__in, vec4<f32>, read> = var undef @location(0)
- %frag_loc0_Output:ptr<__out, vec4<f32>, write> = var undef @location(0)
-}
-
-%vert_inner = func():Interface {
- $B2: {
- %7:vec4<f32> = construct 0.0f
- %8:vec4<f32> = construct 1.0f
- %9:Interface = construct %7, %8
- ret %9
- }
-}
-%frag_inner = func(%inputs:Interface):vec4<f32> {
- $B3: {
- %12:vec4<f32> = access %inputs, 0u
- %13:vec4<f32> = access %inputs, 1u
- %14:vec4<f32> = add %12, %13
- ret %14
- }
-}
-%vert = @vertex func():void {
- $B4: {
- %16:Interface = call %vert_inner
- %17:vec4<f32> = access %16, 0u
- store %vert_position_Output, %17
- %18:vec4<f32> = access %16, 1u
- store %vert_loc0_Output, %18
- ret
- }
-}
-%frag = @fragment func():void {
- $B5: {
- %20:vec4<f32> = load %frag_position_Input
- %21:vec4<f32> = load %frag_loc0_Input
- %22:Interface = construct %20, %21
- %23:vec4<f32> = call %frag_inner, %22
- store %frag_loc0_Output, %23
- ret
- }
-}
-)";
-
- core::ir::transform::ImmediateDataLayout immediate_data;
- ShaderIOConfig config{immediate_data};
- Run(ShaderIO, config);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(SpirvWriter_ShaderIOTest, Struct_SharedWithBuffer) {
auto* vec4f = ty.vec4<f32>();
auto* str_ty =
@@ -964,8 +833,8 @@
EXPECT_EQ(expect, str());
}
-// Test that interpolation attributes are stripped from vertex inputs and fragment outputs.
-TEST_F(SpirvWriter_ShaderIOTest, InterpolationOnVertexInputOrFragmentOutput) {
+// Test that interpolation attributes are stripped from vertex inputs.
+TEST_F(SpirvWriter_ShaderIOTest, InterpolationOnVertexInput) {
auto* str_ty = ty.Struct(mod.symbols.New("MyStruct"),
{
{
@@ -982,45 +851,20 @@
},
});
- // Vertex shader.
- {
- auto* ep = b.Function("vert", ty.vec4<f32>());
- ep->SetReturnBuiltin(core::BuiltinValue::kPosition);
- ep->SetReturnInvariant(true);
- ep->SetStage(core::ir::Function::PipelineStage::kVertex);
+ auto* ep = b.Function("vert", ty.vec4<f32>());
+ ep->SetReturnBuiltin(core::BuiltinValue::kPosition);
+ ep->SetReturnInvariant(true);
+ ep->SetStage(core::ir::Function::PipelineStage::kVertex);
- auto* str_param = b.FunctionParam("input", str_ty);
- auto* ival = b.FunctionParam("ival", ty.i32());
- ival->SetLocation(1);
- ival->SetInterpolation(core::Interpolation{core::InterpolationType::kFlat});
- ep->SetParams({str_param, ival});
+ auto* str_param = b.FunctionParam("input", str_ty);
+ auto* ival = b.FunctionParam("ival", ty.i32());
+ ival->SetLocation(1);
+ ival->SetInterpolation(core::Interpolation{core::InterpolationType::kFlat});
+ ep->SetParams({str_param, ival});
- b.Append(ep->Block(), [&] { //
- b.Return(ep, b.Construct(ty.vec4<f32>(), 0.5_f));
- });
- }
-
- // Fragment shader with struct output.
- {
- auto* ep = b.Function("frag1", str_ty);
- ep->SetStage(core::ir::Function::PipelineStage::kFragment);
-
- b.Append(ep->Block(), [&] { //
- b.Return(ep, b.Construct(str_ty, 0.5_f));
- });
- }
-
- // Fragment shader with non-struct output.
- {
- auto* ep = b.Function("frag2", ty.i32());
- ep->SetStage(core::ir::Function::PipelineStage::kFragment);
- ep->SetReturnLocation(0);
- ep->SetReturnInterpolation(core::Interpolation{core::InterpolationType::kFlat});
-
- b.Append(ep->Block(), [&] { //
- b.Return(ep, b.Constant(42_i));
- });
- }
+ b.Append(ep->Block(), [&] { //
+ b.Return(ep, b.Construct(ty.vec4<f32>(), 0.5_f));
+ });
auto* src = R"(
MyStruct = struct @align(4) {
@@ -1033,17 +877,6 @@
ret %4
}
}
-%frag1 = @fragment func():MyStruct {
- $B2: {
- %6:MyStruct = construct 0.5f
- ret %6
- }
-}
-%frag2 = @fragment func():i32 [@location(0), @interpolate(flat)] {
- $B3: {
- ret 42i
- }
-}
)";
EXPECT_EQ(src, str());
@@ -1056,49 +889,138 @@
%vert_loc1_Input:ptr<__in, f32, read> = var undef @location(1)
%vert_loc1_Input_1:ptr<__in, i32, read> = var undef @location(1) # %vert_loc1_Input_1: 'vert_loc1_Input'
%vert_position_Output:ptr<__out, vec4<f32>, write> = var undef @invariant @builtin(position)
- %frag1_loc1_Output:ptr<__out, f32, write> = var undef @location(1)
- %frag2_loc0_Output:ptr<__out, i32, write> = var undef @location(0)
}
%vert_inner = func(%input:MyStruct, %ival:i32):vec4<f32> {
$B2: {
- %9:vec4<f32> = construct 0.5f
- ret %9
- }
-}
-%frag1_inner = func():MyStruct {
- $B3: {
- %11:MyStruct = construct 0.5f
- ret %11
- }
-}
-%frag2_inner = func():i32 {
- $B4: {
- ret 42i
+ %7:vec4<f32> = construct 0.5f
+ ret %7
}
}
%vert = @vertex func():void {
- $B5: {
- %14:f32 = load %vert_loc1_Input
- %15:MyStruct = construct %14
- %16:i32 = load %vert_loc1_Input_1
- %17:vec4<f32> = call %vert_inner, %15, %16
- store %vert_position_Output, %17
+ $B3: {
+ %9:f32 = load %vert_loc1_Input
+ %10:MyStruct = construct %9
+ %11:i32 = load %vert_loc1_Input_1
+ %12:vec4<f32> = call %vert_inner, %10, %11
+ store %vert_position_Output, %12
ret
}
}
+)";
+
+ core::ir::transform::ImmediateDataLayout immediate_data;
+ ShaderIOConfig config{immediate_data};
+ Run(ShaderIO, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+// Test that interpolation attributes are stripped from fragment struct outputs.
+TEST_F(SpirvWriter_ShaderIOTest, InterpolationOnFragmentOutput_Struct) {
+ auto* str_ty = ty.Struct(mod.symbols.New("MyStruct"),
+ {
+ {
+ mod.symbols.New("color"),
+ ty.f32(),
+ core::IOAttributes{
+ .location = 1u,
+ .interpolation =
+ core::Interpolation{
+ core::InterpolationType::kLinear,
+ core::InterpolationSampling::kSample,
+ },
+ },
+ },
+ });
+
+ auto* ep = b.Function("frag1", str_ty);
+ ep->SetStage(core::ir::Function::PipelineStage::kFragment);
+
+ b.Append(ep->Block(), [&] { //
+ b.Return(ep, b.Construct(str_ty, 0.5_f));
+ });
+
+ auto* src = R"(
+MyStruct = struct @align(4) {
+ color:f32 @offset(0), @location(1), @interpolate(linear, sample)
+}
+
+%frag1 = @fragment func():MyStruct {
+ $B1: {
+ %2:MyStruct = construct 0.5f
+ ret %2
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+MyStruct = struct @align(4) {
+ color:f32 @offset(0)
+}
+
+$B1: { # root
+ %frag1_loc1_Output:ptr<__out, f32, write> = var undef @location(1)
+}
+
+%frag1_inner = func():MyStruct {
+ $B2: {
+ %3:MyStruct = construct 0.5f
+ ret %3
+ }
+}
%frag1 = @fragment func():void {
- $B6: {
- %19:MyStruct = call %frag1_inner
- %20:f32 = access %19, 0u
- store %frag1_loc1_Output, %20
+ $B3: {
+ %5:MyStruct = call %frag1_inner
+ %6:f32 = access %5, 0u
+ store %frag1_loc1_Output, %6
ret
}
}
+)";
+
+ core::ir::transform::ImmediateDataLayout immediate_data;
+ ShaderIOConfig config{immediate_data};
+ Run(ShaderIO, config);
+
+ EXPECT_EQ(expect, str());
+}
+
+// Test that interpolation attributes are stripped from fragment non-struct outputs.
+TEST_F(SpirvWriter_ShaderIOTest, InterpolationOnFragmentOutput_NonStruct) {
+ auto* ep = b.Function("frag2", ty.i32());
+ ep->SetStage(core::ir::Function::PipelineStage::kFragment);
+ ep->SetReturnLocation(0);
+ ep->SetReturnInterpolation(core::Interpolation{core::InterpolationType::kFlat});
+
+ b.Append(ep->Block(), [&] { //
+ b.Return(ep, b.Constant(42_i));
+ });
+
+ auto* src = R"(
+%frag2 = @fragment func():i32 [@location(0), @interpolate(flat)] {
+ $B1: {
+ ret 42i
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+$B1: { # root
+ %frag2_loc0_Output:ptr<__out, i32, write> = var undef @location(0)
+}
+
+%frag2_inner = func():i32 {
+ $B2: {
+ ret 42i
+ }
+}
%frag2 = @fragment func():void {
- $B7: {
- %22:i32 = call %frag2_inner
- store %frag2_loc0_Output, %22
+ $B3: {
+ %4:i32 = call %frag2_inner
+ store %frag2_loc0_Output, %4
ret
}
}
@@ -1205,162 +1127,6 @@
EXPECT_EQ(expect, str());
}
-TEST_F(SpirvWriter_ShaderIOTest, ClampFragDepth_MultipleFragmentShaders) {
- auto* str_ty =
- ty.Struct(mod.symbols.New("Outputs"), {
- {
- mod.symbols.New("color"),
- ty.f32(),
- core::IOAttributes{
- .location = 0u,
- },
- },
- {
- mod.symbols.New("depth"),
- ty.f32(),
- core::IOAttributes{
- .builtin = core::BuiltinValue::kFragDepth,
- },
- },
- });
-
- auto make_entry_point = [&](std::string_view name) {
- auto* ep = b.Function(name, str_ty);
- ep->SetStage(core::ir::Function::PipelineStage::kFragment);
- b.Append(ep->Block(), [&] { //
- b.Return(ep, b.Construct(str_ty, 0.5_f, 2_f));
- });
- };
- make_entry_point("ep1");
- make_entry_point("ep2");
- make_entry_point("ep3");
-
- auto* src = R"(
-Outputs = struct @align(4) {
- color:f32 @offset(0), @location(0)
- depth:f32 @offset(4), @builtin(frag_depth)
-}
-
-%ep1 = @fragment func():Outputs {
- $B1: {
- %2:Outputs = construct 0.5f, 2.0f
- ret %2
- }
-}
-%ep2 = @fragment func():Outputs {
- $B2: {
- %4:Outputs = construct 0.5f, 2.0f
- ret %4
- }
-}
-%ep3 = @fragment func():Outputs {
- $B3: {
- %6:Outputs = construct 0.5f, 2.0f
- ret %6
- }
-}
-)";
- EXPECT_EQ(src, str());
-
- auto* expect = R"(
-Outputs = struct @align(4) {
- color:f32 @offset(0)
- depth:f32 @offset(4)
-}
-
-tint_immediate_data_struct = struct @align(4), @block {
- depth_min:f32 @offset(4)
- depth_max:f32 @offset(8)
-}
-
-$B1: { # root
- %tint_immediate_data:ptr<immediate, tint_immediate_data_struct, read> = var undef
- %ep1_loc0_Output:ptr<__out, f32, write> = var undef @location(0)
- %ep1_frag_depth_Output:ptr<__out, f32, write> = var undef @builtin(frag_depth)
- %ep2_loc0_Output:ptr<__out, f32, write> = var undef @location(0)
- %ep2_frag_depth_Output:ptr<__out, f32, write> = var undef @builtin(frag_depth)
- %ep3_loc0_Output:ptr<__out, f32, write> = var undef @location(0)
- %ep3_frag_depth_Output:ptr<__out, f32, write> = var undef @builtin(frag_depth)
-}
-
-%ep1_inner = func():Outputs {
- $B2: {
- %9:Outputs = construct 0.5f, 2.0f
- ret %9
- }
-}
-%ep2_inner = func():Outputs {
- $B3: {
- %11:Outputs = construct 0.5f, 2.0f
- ret %11
- }
-}
-%ep3_inner = func():Outputs {
- $B4: {
- %13:Outputs = construct 0.5f, 2.0f
- ret %13
- }
-}
-%ep1 = @fragment func():void {
- $B5: {
- %15:Outputs = call %ep1_inner
- %16:f32 = access %15, 0u
- store %ep1_loc0_Output, %16
- %17:f32 = access %15, 1u
- %18:ptr<immediate, f32, read> = access %tint_immediate_data, 0u
- %19:f32 = load %18
- %20:ptr<immediate, f32, read> = access %tint_immediate_data, 1u
- %21:f32 = load %20
- %22:f32 = clamp %17, %19, %21
- store %ep1_frag_depth_Output, %22
- ret
- }
-}
-%ep2 = @fragment func():void {
- $B6: {
- %24:Outputs = call %ep2_inner
- %25:f32 = access %24, 0u
- store %ep2_loc0_Output, %25
- %26:f32 = access %24, 1u
- %27:ptr<immediate, f32, read> = access %tint_immediate_data, 0u
- %28:f32 = load %27
- %29:ptr<immediate, f32, read> = access %tint_immediate_data, 1u
- %30:f32 = load %29
- %31:f32 = clamp %26, %28, %30
- store %ep2_frag_depth_Output, %31
- ret
- }
-}
-%ep3 = @fragment func():void {
- $B7: {
- %33:Outputs = call %ep3_inner
- %34:f32 = access %33, 0u
- store %ep3_loc0_Output, %34
- %35:f32 = access %33, 1u
- %36:ptr<immediate, f32, read> = access %tint_immediate_data, 0u
- %37:f32 = load %36
- %38:ptr<immediate, f32, read> = access %tint_immediate_data, 1u
- %39:f32 = load %38
- %40:f32 = clamp %35, %37, %39
- store %ep3_frag_depth_Output, %40
- ret
- }
-}
-)";
-
- core::ir::transform::PrepareImmediateDataConfig immediate_data_config;
- immediate_data_config.AddInternalImmediateData(4, mod.symbols.New("depth_min"), ty.f32());
- immediate_data_config.AddInternalImmediateData(8, mod.symbols.New("depth_max"), ty.f32());
- auto immediate_data = PrepareImmediateData(mod, immediate_data_config);
- EXPECT_EQ(immediate_data, Success);
-
- ShaderIOConfig config{immediate_data.Get()};
- config.depth_range_offsets = {4, 8};
- Run(ShaderIO, config);
-
- EXPECT_EQ(expect, str());
-}
-
TEST_F(SpirvWriter_ShaderIOTest, EmitVertexPointSize) {
auto* ep = b.Function("foo", ty.vec4<f32>());
ep->SetStage(core::ir::Function::PipelineStage::kVertex);
diff --git a/src/tint/lang/wgsl/reader/lower/lower.cc b/src/tint/lang/wgsl/reader/lower/lower.cc
index 6ace4ff..a4fc881 100644
--- a/src/tint/lang/wgsl/reader/lower/lower.cc
+++ b/src/tint/lang/wgsl/reader/lower/lower.cc
@@ -208,8 +208,12 @@
} // namespace
Result<SuccessType> Lower(core::ir::Module& mod) {
- auto res = core::ir::ValidateAndDumpIfNeeded(
- mod, "wgsl.Lower", core::ir::Capabilities{core::ir::Capability::kAllowOverrides});
+ auto res =
+ core::ir::ValidateAndDumpIfNeeded(mod, "wgsl.Lower",
+ core::ir::Capabilities{
+ core::ir::Capability::kAllowMultipleEntryPoints,
+ core::ir::Capability::kAllowOverrides,
+ });
if (res != Success) {
return res.Failure();
}
