[SYCL] Support passing bool as kernel param for SPIR-V target (#17427)

When targeting SPIR-V, the compiler was incorrectly lowering bool to i1.
This caused the translator to produce an OpTypeBool which is illegal for
an entry point in SPIR-V.

Instead, the patch coerce the type to follow the host size.

Fix #11531

Author: Victor Lomuller

clang/lib/CodeGen/Targets/SPIR.cpp

@@ -38,9 +38,20 @@ class CommonSPIRABIInfo : public DefaultABIInfo {
 ABIArgInfo CommonSPIRABIInfo::classifyKernelArgumentType(QualType Ty) const {
   Ty = useFirstFieldIfTransparentUnion(Ty);
 
-  if (getContext().getLangOpts().SYCLIsDevice && isAggregateTypeForABI(Ty)) {
+  if (getContext().getLangOpts().SYCLIsDevice) {
+    if (const BuiltinType *BT = Ty->getAs<BuiltinType>()) {
+      switch (BT->getKind()) {
+      case BuiltinType::Bool:
+        // Bool / i1 isn't a legal kernel argument in SPIR-V.
+        // Coerce the type to follow the host representation of bool.
+        return ABIArgInfo::getDirect(CGT.ConvertTypeForMem(Ty));
+      default:
+        break;
+      }
+    }
     // Pass all aggregate types allowed by Sema by value.
-    return getNaturalAlignIndirect(Ty);
+    if (isAggregateTypeForABI(Ty))
+      return getNaturalAlignIndirect(Ty);
   }
 
   return DefaultABIInfo::classifyArgumentType(Ty);

clang/test/CodeGenSYCL/bool-kernel-argument.cpp

@@ -0,0 +1,26 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+
+#include "sycl.hpp"
+
+void take_bool(bool) {}
+
+int main() {
+  bool test = false;
+  sycl::queue q;
+
+  // CHECK: @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E11test_kernel(i8 {{.*}} [[ARG:%[A-Za-z_0-9]*]]
+  // CHECK: %__SYCLKernel = alloca
+  // CHECK: %test = getelementptr inbounds nuw %class.anon, ptr addrspace(4) %__SYCLKernel.ascast
+  // CHECK: store i8 %{{.*}}, ptr addrspace(4) %test
+  // CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv
+  //
+  // CHECK: define {{.*}} @_Z9take_boolb(i1
+  q.submit([&](sycl::handler &h) {
+    h.single_task<class test_kernel>([=]() {
+      (void)test;
+      take_bool(test);
+    });
+  });
+
+  return 0;
+}

sycl/test-e2e/Basic/bool-test.cpp

@@ -0,0 +1,30 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Check booleans are promoted correctly
+
+#include <sycl/detail/core.hpp>
+#include <sycl/usm.hpp>
+
+void run_test(sycl::queue q, bool test, int *res) {
+  q.submit([&](sycl::handler &cgh) {
+     cgh.single_task([=]() {
+       if (test)
+         *res = 42;
+       else
+         *res = -42;
+     });
+   }).wait();
+}
+
+int main() {
+  sycl::queue q;
+  int *p = sycl::malloc_shared<int>(1, q);
+  *p = 0;
+  run_test(q, true, p);
+  assert(*p == 42);
+  *p = 0;
+  run_test(q, false, p);
+  assert(*p == -42);
+  sycl::free(p, q);
+}