HLSL backend: work around FXC ignoring switch with only default case

In HLSL code, if a switch statement has only a default case, FXC will
effectively ignore the code in that case. In this change, we detect this
and work around it by emitting the code in the default block without the
switch.

Bug: tint:1188
Change-Id: I69b405cdb4c669fb093eb49aa138923419dcf8f8
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/68440
Kokoro: Antonio Maiorano <amaiorano@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@chromium.org>

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 { <default case body> } 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(";

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

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* s = Switch(                      //
-  auto* def_body = Block(create<ast::BreakStatement>());
+      Expr("cond"),                      //
-  auto* def = create<ast::CaseStatement>(ast::CaseSelectorList{}, def_body);
+      Case(Literal(5), Block(Break())),  //
-
+      DefaultCase());
-  ast::CaseSelectorList case_val;
-  case_val.push_back(Literal(5));
-
-  auto* case_body = Block(create<ast::BreakStatement>());
-
-  auto* case_stmt = create<ast::CaseStatement>(case_val, case_body);
-
-  ast::CaseStatementList body;
-  body.push_back(case_stmt);
-  body.push_back(def);
-
-  auto* cond = Expr("cond");
-  auto* s = create<ast::SwitchStatement>(cond, body);
   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

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;
+        }
+    }
+}

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;
+}

test/statements/switch/common.wgsl.expected.msl

@@ -0,0 +1,27 @@
+#include <metal_stdlib>
+
+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;
+}
+

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

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;
+    }
+  }
+}

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;
+        }
+    }
+}

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;
+}

test/statements/switch/only_default_case.wgsl.expected.msl

@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  int i = 0;
+  int result = 0;
+  switch(i) {
+    default: {
+      result = 44;
+      break;
+    }
+  }
+  return;
+}
+

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

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;
+    }
+  }
+}