diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index 12b3b0f1c2..2730c398ee 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -2948,7 +2948,44 @@ bool GeneratorImpl::EmitStatement(const ast::Statement* stmt) { return false; } +bool GeneratorImpl::EmitDefaultOnlySwitch(const ast::SwitchStatement* stmt) { + TINT_ASSERT(Writer, stmt->body.size() == 1 && stmt->body[0]->IsDefault()); + + // FXC fails to compile a switch with just a default case, ignoring the + // default case body. We work around this here by emitting the default case + // without the switch. + + // Emit the switch condition as-is in case it has side-effects (e.g. + // function call). Note that's it's fine not to assign the result of the + // expression. + { + auto out = line(); + if (!EmitExpression(out, stmt->condition)) { + return false; + } + out << ";"; + } + + // Emit "do { } while(false);". We use a 'do' loop so + // that break statements work as expected, and make it 'while (false)' in + // case there isn't a break statement. + line() << "do {"; + { + ScopedIndent si(this); + if (!EmitStatements(stmt->body[0]->body->statements)) { + return false; + } + } + line() << "} while (false);"; + return true; +} + bool GeneratorImpl::EmitSwitch(const ast::SwitchStatement* stmt) { + // BUG(crbug.com/tint/1188): work around default-only switches + if (stmt->body.size() == 1 && stmt->body[0]->IsDefault()) { + return EmitDefaultOnlySwitch(stmt); + } + { // switch(expr) { auto out = line(); out << "switch("; diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h index 57d9cb848a..61b895dcff 100644 --- a/src/writer/hlsl/generator_impl.h +++ b/src/writer/hlsl/generator_impl.h @@ -327,6 +327,10 @@ class GeneratorImpl : public TextGenerator { /// @param stmt the statement to emit /// @returns true if the statement was emitted bool EmitSwitch(const ast::SwitchStatement* stmt); + // Handles generating a switch statement with only a default case + /// @param stmt the statement to emit + /// @returns true if the statement was emitted + bool EmitDefaultOnlySwitch(const ast::SwitchStatement* stmt); /// Handles generating type /// @param out the output stream /// @param type the type to generate diff --git a/src/writer/hlsl/generator_impl_switch_test.cc b/src/writer/hlsl/generator_impl_switch_test.cc index c0a92cd3fb..89d69bcc16 100644 --- a/src/writer/hlsl/generator_impl_switch_test.cc +++ b/src/writer/hlsl/generator_impl_switch_test.cc @@ -23,23 +23,10 @@ using HlslGeneratorImplTest_Switch = TestHelper; TEST_F(HlslGeneratorImplTest_Switch, Emit_Switch) { Global("cond", ty.i32(), ast::StorageClass::kPrivate); - - auto* def_body = Block(create()); - auto* def = create(ast::CaseSelectorList{}, def_body); - - ast::CaseSelectorList case_val; - case_val.push_back(Literal(5)); - - auto* case_body = Block(create()); - - auto* case_stmt = create(case_val, case_body); - - ast::CaseStatementList body; - body.push_back(case_stmt); - body.push_back(def); - - auto* cond = Expr("cond"); - auto* s = create(cond, body); + auto* s = Switch( // + Expr("cond"), // + Case(Literal(5), Block(Break())), // + DefaultCase()); WrapInFunction(s); GeneratorImpl& gen = Build(); @@ -58,6 +45,26 @@ TEST_F(HlslGeneratorImplTest_Switch, Emit_Switch) { )"); } +TEST_F(HlslGeneratorImplTest_Switch, Emit_Switch_OnlyDefaultCase) { + Global("cond", ty.i32(), ast::StorageClass::kPrivate); + Global("a", ty.i32(), ast::StorageClass::kPrivate); + auto* s = Switch( // + Expr("cond"), // + DefaultCase(Block(Assign(Expr("a"), Expr(42))))); + WrapInFunction(s); + + GeneratorImpl& gen = Build(); + + gen.increment_indent(); + + ASSERT_TRUE(gen.EmitStatement(s)) << gen.error(); + EXPECT_EQ(gen.result(), R"( cond; + do { + a = 42; + } while (false); +)"); +} + } // namespace } // namespace hlsl } // namespace writer diff --git a/test/statements/switch/common.wgsl b/test/statements/switch/common.wgsl new file mode 100644 index 0000000000..ada8f5e56d --- /dev/null +++ b/test/statements/switch/common.wgsl @@ -0,0 +1,20 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var i : i32; + var result : i32; + switch(i) { + case 0: { + result = 10; + } + case 1: { + result = 22; + } + case 2: { + result = 33; + } + default: { + result = 44; + break; + } + } +} diff --git a/test/statements/switch/common.wgsl.expected.hlsl b/test/statements/switch/common.wgsl.expected.hlsl new file mode 100644 index 0000000000..7aa505cf7e --- /dev/null +++ b/test/statements/switch/common.wgsl.expected.hlsl @@ -0,0 +1,24 @@ +[numthreads(1, 1, 1)] +void f() { + int i = 0; + int result = 0; + switch(i) { + case 0: { + result = 10; + break; + } + case 1: { + result = 22; + break; + } + case 2: { + result = 33; + break; + } + default: { + result = 44; + break; + } + } + return; +} diff --git a/test/statements/switch/common.wgsl.expected.msl b/test/statements/switch/common.wgsl.expected.msl new file mode 100644 index 0000000000..155514582a --- /dev/null +++ b/test/statements/switch/common.wgsl.expected.msl @@ -0,0 +1,27 @@ +#include + +using namespace metal; +kernel void f() { + int i = 0; + int result = 0; + switch(i) { + case 0: { + result = 10; + break; + } + case 1: { + result = 22; + break; + } + case 2: { + result = 33; + break; + } + default: { + result = 44; + break; + } + } + return; +} + diff --git a/test/statements/switch/common.wgsl.expected.spvasm b/test/statements/switch/common.wgsl.expected.spvasm new file mode 100644 index 0000000000..3e1b8e1a18 --- /dev/null +++ b/test/statements/switch/common.wgsl.expected.spvasm @@ -0,0 +1,43 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 20 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %i "i" + OpName %result "result" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 +%_ptr_Function_int = OpTypePointer Function %int + %8 = OpConstantNull %int + %int_10 = OpConstant %int 10 + %int_22 = OpConstant %int 22 + %int_33 = OpConstant %int 33 + %int_44 = OpConstant %int 44 + %f = OpFunction %void None %1 + %4 = OpLabel + %i = OpVariable %_ptr_Function_int Function %8 + %result = OpVariable %_ptr_Function_int Function %8 + %11 = OpLoad %int %i + OpSelectionMerge %10 None + OpSwitch %11 %12 0 %13 1 %14 2 %15 + %13 = OpLabel + OpStore %result %int_10 + OpBranch %10 + %14 = OpLabel + OpStore %result %int_22 + OpBranch %10 + %15 = OpLabel + OpStore %result %int_33 + OpBranch %10 + %12 = OpLabel + OpStore %result %int_44 + OpBranch %10 + %10 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/statements/switch/common.wgsl.expected.wgsl b/test/statements/switch/common.wgsl.expected.wgsl new file mode 100644 index 0000000000..9aad3f1fb3 --- /dev/null +++ b/test/statements/switch/common.wgsl.expected.wgsl @@ -0,0 +1,20 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var i : i32; + var result : i32; + switch(i) { + case 0: { + result = 10; + } + case 1: { + result = 22; + } + case 2: { + result = 33; + } + default: { + result = 44; + break; + } + } +} diff --git a/test/switch/fallthrough.wgsl b/test/statements/switch/fallthrough.wgsl similarity index 100% rename from test/switch/fallthrough.wgsl rename to test/statements/switch/fallthrough.wgsl diff --git a/test/switch/fallthrough.wgsl.expected.hlsl b/test/statements/switch/fallthrough.wgsl.expected.hlsl similarity index 100% rename from test/switch/fallthrough.wgsl.expected.hlsl rename to test/statements/switch/fallthrough.wgsl.expected.hlsl diff --git a/test/switch/fallthrough.wgsl.expected.msl b/test/statements/switch/fallthrough.wgsl.expected.msl similarity index 100% rename from test/switch/fallthrough.wgsl.expected.msl rename to test/statements/switch/fallthrough.wgsl.expected.msl diff --git a/test/switch/fallthrough.wgsl.expected.spvasm b/test/statements/switch/fallthrough.wgsl.expected.spvasm similarity index 100% rename from test/switch/fallthrough.wgsl.expected.spvasm rename to test/statements/switch/fallthrough.wgsl.expected.spvasm diff --git a/test/switch/fallthrough.wgsl.expected.wgsl b/test/statements/switch/fallthrough.wgsl.expected.wgsl similarity index 100% rename from test/switch/fallthrough.wgsl.expected.wgsl rename to test/statements/switch/fallthrough.wgsl.expected.wgsl diff --git a/test/statements/switch/only_default_case.wgsl b/test/statements/switch/only_default_case.wgsl new file mode 100644 index 0000000000..cb12fb9c11 --- /dev/null +++ b/test/statements/switch/only_default_case.wgsl @@ -0,0 +1,11 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var i : i32; + var result : i32; + switch(i) { + default: { + result = 44; + break; + } + } +} diff --git a/test/statements/switch/only_default_case.wgsl.expected.hlsl b/test/statements/switch/only_default_case.wgsl.expected.hlsl new file mode 100644 index 0000000000..128b8a2393 --- /dev/null +++ b/test/statements/switch/only_default_case.wgsl.expected.hlsl @@ -0,0 +1,11 @@ +[numthreads(1, 1, 1)] +void f() { + int i = 0; + int result = 0; + i; + do { + result = 44; + break; + } while (false); + return; +} diff --git a/test/statements/switch/only_default_case.wgsl.expected.msl b/test/statements/switch/only_default_case.wgsl.expected.msl new file mode 100644 index 0000000000..c74e332e4c --- /dev/null +++ b/test/statements/switch/only_default_case.wgsl.expected.msl @@ -0,0 +1,15 @@ +#include + +using namespace metal; +kernel void f() { + int i = 0; + int result = 0; + switch(i) { + default: { + result = 44; + break; + } + } + return; +} + diff --git a/test/statements/switch/only_default_case.wgsl.expected.spvasm b/test/statements/switch/only_default_case.wgsl.expected.spvasm new file mode 100644 index 0000000000..35237ab5be --- /dev/null +++ b/test/statements/switch/only_default_case.wgsl.expected.spvasm @@ -0,0 +1,31 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 14 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %f "f" + OpName %i "i" + OpName %result "result" + %void = OpTypeVoid + %1 = OpTypeFunction %void + %int = OpTypeInt 32 1 +%_ptr_Function_int = OpTypePointer Function %int + %8 = OpConstantNull %int + %int_44 = OpConstant %int 44 + %f = OpFunction %void None %1 + %4 = OpLabel + %i = OpVariable %_ptr_Function_int Function %8 + %result = OpVariable %_ptr_Function_int Function %8 + %11 = OpLoad %int %i + OpSelectionMerge %10 None + OpSwitch %11 %12 + %12 = OpLabel + OpStore %result %int_44 + OpBranch %10 + %10 = OpLabel + OpReturn + OpFunctionEnd diff --git a/test/statements/switch/only_default_case.wgsl.expected.wgsl b/test/statements/switch/only_default_case.wgsl.expected.wgsl new file mode 100644 index 0000000000..bd050b7592 --- /dev/null +++ b/test/statements/switch/only_default_case.wgsl.expected.wgsl @@ -0,0 +1,11 @@ +[[stage(compute), workgroup_size(1)]] +fn f() { + var i : i32; + var result : i32; + switch(i) { + default: { + result = 44; + break; + } + } +}