[SYCL] Always set NoUnwind attribute for SYCL.

[hvdijk]

Summary

Like CUDA and OpenCL, the SYCL specification says that throwing and
catching exceptions in device functions is not supported, so this change
extends the logic for adding the NoUnwind attribute to SYCL.

The existing convergent.cpp test, which tests that the convergent
attribute is added to functions by default, is renamed and reused to
test that the nounwind attribute is added by default. This test now has
-fexceptions added to it, which the driver adds by default as well.

The obvious question here is why not simply change the driver to remove
-fexceptions. This change follows the direction given by the TODO
comment because removing -fexceptions would also disable the
__EXCEPTIONS macro, which should reflect whether exceptions are enabled
on the host, rather than on the device, to avoid conflicts in types
shared between host and device.

[hvdijk]

Is the rationale I gave in the description correct, or would it be better for SYCL device code to unconditionally build without -fexceptions and get the nounwind attribute added that way?

[bader]

@hvdijk, thanks a lot for fixing this.

In D147097#4229121, @hvdijk wrote:

Is the rationale I gave in the description correct, or would it be better for SYCL device code to unconditionally build without -fexceptions and get the nounwind attribute added that way?

That's a good question. I haven't looked into this issue deep enough, but I think using -fexceptions requires using delayed diagnostics to avoid false diagnostics during host code analysis.
Anyway, all GPU offloading single-source modes have the same restriction and design and we better have unified solution whether it's using -fexceptions or adding nounwind attribute in CodeGen.

[hvdijk]

That's a good question. I haven't looked into this issue deep enough, but I think using -fexceptions requires using delayed diagnostics to avoid false diagnostics during host code analysis.

I am assuming you mean -fno-exceptions (or, in clang -cc1, the absence of -fexceptions)? This is a good point. Delayed diagnostics would probably be good in general: we currently already emit warnings for host code when compiling for device, but as long as the generated warnings are identical for the host code as when compiling for host, it is easy enough to ignore; if the device compilation were to result in additional warnings for host code, and those warnings are incorrect, that would be a rather poor user experience. That sounds like a good additional reason to do it the way I had done here, the way the existing code comment had indicated.

[hvdijk]

Apparently the test was already passing without my change, so I updated the test to make it a function that did not previously get the nounwind attribute. This also revealed that the test was more fragile than I realised: it depends on the precise order in which functions are emitted. We always match "Function Attrs:" to the first emitted function, and then CHECK-NEXT can fail. We'd ideally first CHECK for void @_Z3foov() (now void @_Z3barv()) and then CHECK-PREV for the attributes, but FileCheck does not support that. What we can do instead though is find the matching attributes #n = { ... } line and verify that everything is in there; updated accordingly.

[hvdijk]

Forgot to update the call to foo() to call bar() instead.

[bader]

LGTM. Thanks!

clang/lib/CodeGen/CGCall.cpp

if (getLangOpts().assumeFunctionsAreConvergent()) {
     // convergent (meaning, they may call an intrinsically convergent op, such
     // as __syncthreads() / barrier(), and so can't have certain optimizations
     // applied around them).  LLVM will remove this attribute where it safely
     // can.
     FuncAttrs.addAttribute(llvm::Attribute::Convergent);
   }
 
   // TODO: NoUnwind attribute should be added for other GPU modes HIP,
-  // SYCL, OpenMP offload. AFAIK, none of them support exceptions in device
-  // code.
+  // OpenMP offload. AFAIK, neither of them support exceptions in device code.
   if ((getLangOpts().CUDA && getLangOpts().CUDAIsDevice) ||
-      getLangOpts().OpenCL) {
+      getLangOpts().OpenCL || getLangOpts().SYCLIsDevice) {
     FuncAttrs.addAttribute(llvm::Attribute::NoUnwind);
   }
 
   for (StringRef Attr : CodeGenOpts.DefaultFunctionAttrs) {
     StringRef Var, Value;
     std::tie(Var, Value) = Attr.split('=');
     FuncAttrs.addAttribute(Var, Value);
   }

clang/test/CodeGenSYCL/convergent.cpp

-// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -disable-llvm-passes \
-// RUN:  -triple spir64 -emit-llvm %s -o - | FileCheck %s
-
-// CHECK-DAG: Function Attrs:
-// CHECK-DAG-SAME: convergent
-// CHECK-DAG-NEXT: define void @_Z3foov
-void foo() {
-  int a = 1;
-}
-
-template <typename Name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
-  kernelFunc();
-}
-
-int main() {
-  kernel_single_task<class fake_kernel>([] { foo(); });
-  return 0;
-}

clang/test/CodeGenSYCL/function-attrs.cpp

+// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -disable-llvm-passes \
+// RUN:  -triple spir64 -fexceptions -emit-llvm %s -o - | FileCheck %s
+
+int foo();
+
+// CHECK: define dso_local spir_func void @_Z3barv() [[BAR:#[0-9]+]]
+// CHECK: attributes [[BAR]] =
+// CHECK-SAME: convergent
+// CHECK-SAME: nounwind
+void bar() {
+  int a = foo();
+}
+
+int foo() {
+  return 1;
+}
+
+template <typename Name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(const Func &kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  kernel_single_task<class fake_kernel>([] { bar(); });
+  return 0;
+}