[SYCL] Issue deferred diagnostic for incorrect asm register usage (#10782)

The device compiler previously issued an asm diagnostic for the usage of
a register
that was incorrect for the device, even though such usage was not on the
device.
Now fixed by converting the diagnostic to a deferred one.

Author: premanandrao

clang/lib/Sema/SemaDecl.cpp

@@ -8106,8 +8106,14 @@ NamedDecl *Sema::ActOnVariableDeclarator(
       case SC_Register:
         // Local Named register
         if (!Context.getTargetInfo().isValidGCCRegisterName(Label) &&
-            DeclAttrsMatchCUDAMode(getLangOpts(), getCurFunctionDecl()))
-          Diag(E->getExprLoc(), diag::err_asm_unknown_register_name) << Label;
+            DeclAttrsMatchCUDAMode(getLangOpts(), getCurFunctionDecl())) {
+          if (getLangOpts().SYCLIsDevice)
+            SYCLDiagIfDeviceCode(E->getExprLoc(),
+                                 diag::err_asm_unknown_register_name)
+                << Label;
+          else
+            Diag(E->getExprLoc(), diag::err_asm_unknown_register_name) << Label;
+        }
         break;
       case SC_Static:
       case SC_Extern:

clang/test/SemaSYCL/sycl-device-asm.cpp

@@ -0,0 +1,29 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -sycl-std=2020 %s -verify
+
+// This test checks that the device compiler does not issue an asm
+// diagnostic about a register incorrect for it, unless it is a routine
+// called on the device.
+#include "sycl.hpp"
+
+void non_device_func(int value) {
+  // expected-no-diagnostic@+1
+  { register int v asm ("eax") = value; }
+}
+
+void device_func(int value) {
+  // expected-error@+1 {{unknown register name 'eax' in asm}}
+  { register int v asm ("eax") = value; }
+}
+
+int main() {
+  sycl::queue q;
+
+  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()'}}
+          device_func(5);
+        });
+  });
+}