[SYCL] Defer diagnostics about usage of SEH (intel#11619)

Rather issue an immediate diagnostic about the use of SEH, emit them
deferred similar to what is done for C++ exceptions.

clang/lib/Sema/SemaStmt.cpp

@@ -4626,8 +4626,12 @@ StmtResult Sema::ActOnSEHTryBlock(bool IsCXXTry, SourceLocation TryLoc,
     Diag(TryLoc, diag::err_seh_try_outside_functions);
 
   // Reject __try on unsupported targets.
-  if (!Context.getTargetInfo().isSEHTrySupported())
-    Diag(TryLoc, diag::err_seh_try_unsupported);
+  if (!Context.getTargetInfo().isSEHTrySupported()) {
+    if (getLangOpts().SYCLIsDevice)
+      SYCLDiagIfDeviceCode(TryLoc, diag::err_seh_try_unsupported);
+    else
+      Diag(TryLoc, diag::err_seh_try_unsupported);
+  }
 
   return SEHTryStmt::Create(Context, IsCXXTry, TryLoc, TryBlock, Handler);
 }

clang/test/SemaSYCL/sycl-seh.cpp

@@ -0,0 +1,34 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -aux-triple x86_64 -fms-compatibility %s -verify
+//
+// This test checks that we issue a diagnostic about the use of SEH
+// features only when they are used on the device.
+//
+#include "sycl.hpp"
+
+void seh_okay_on_host() {
+  __try {
+  } __except(0) {
+  }
+}
+
+void seh_not_okay_on_device() {
+  // expected-error@+2 {{SEH '__try' is not supported on this target}}
+  // expected-error@+1 {{SYCL kernel cannot use exceptions}}
+  __try {
+  } __except(0) {
+  }
+}
+
+void foo() {
+  sycl::queue q;
+
+  seh_okay_on_host();
+  q.submit([&](sycl::handler &h) {
+    // expected-note@#KernelSingleTaskKernelFuncCall {{called by 'kernel_single_task<kernel_wrapper}}
+    h.single_task<class kernel_wrapper>(
+        [=]() {
+          // expected-note@+1 {{called by 'operator()'}}
+          seh_not_okay_on_device();
+        });
+  });
+}