[SYCL] Support or diagnose use of namespace std types as kernel type names

When std::nullptr_t is used as a kernel type, the generated
integration header uses 'nullptr_t'.  This causes lookup
errors.  Use 'std::nullptr_t' instead.

std::max_align_t is defined (in one implementation) as a typedef
of an anonymous struct.  This causes errors when attempting to
forward declare the type in the integration header.  Diagnose
such cases earlier.

Signed-off-by: Premanand M Rao <[email protected]>

Author: premanandrao

clang/lib/Sema/SemaSYCL.cpp

@@ -1649,30 +1649,27 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
     auto *NS = dyn_cast_or_null<NamespaceDecl>(DC);
 
     if (!NS) {
-      if (!DC->isTranslationUnit()) {
-        const TagDecl *TD = isa<ClassTemplateDecl>(D)
-                                ? cast<ClassTemplateDecl>(D)->getTemplatedDecl()
-                                : dyn_cast<TagDecl>(D);
-
-        if (TD && !UnnamedLambdaSupport) {
-          // defined class constituting the kernel name is not globally
-          // accessible - contradicts the spec
-          const bool KernelNameIsMissing = TD->getName().empty();
-          if (KernelNameIsMissing) {
+      const TagDecl *TD = isa<ClassTemplateDecl>(D)
+                              ? cast<ClassTemplateDecl>(D)->getTemplatedDecl()
+                              : dyn_cast<TagDecl>(D);
+      if (!TD)
+        break;
+
+      const bool KernelNameIsMissing = TD->getName().empty();
+      if (KernelNameIsMissing)
+        Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
+            << /* kernel name is missing */ 0;
+      else if (!DC->isTranslationUnit()) {
+        // defined class constituting the kernel name is not globally
+        // accessible - contradicts the spec
+        if (!UnnamedLambdaSupport) {
+          if (TD->isCompleteDefinition())
             Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
-                << /* kernel name is missing */ 0;
-            // Don't emit note if kernel name was completely omitted
-          } else {
-            if (TD->isCompleteDefinition())
-              Diag.Report(KernelLocation,
-                          diag::err_sycl_kernel_incorrectly_named)
-                  << /* kernel name is not globally-visible */ 1;
-            else
-              Diag.Report(KernelLocation, diag::warn_sycl_implicit_decl);
-            Diag.Report(D->getSourceRange().getBegin(),
-                        diag::note_previous_decl)
-                << TD->getName();
-          }
+                << /* kernel name is not globally-visible */ 1;
+          else
+            Diag.Report(KernelLocation, diag::warn_sycl_implicit_decl);
+          Diag.Report(D->getSourceRange().getBegin(), diag::note_previous_decl)
+              << TD->getName();
         }
       }
       break;
@@ -1826,8 +1823,10 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "// This is auto-generated SYCL integration header.\n";
   O << "\n";
 
+  O << "#include <cstddef>\n";
   O << "#include <CL/sycl/detail/defines.hpp>\n";
   O << "#include <CL/sycl/detail/kernel_desc.hpp>\n";
+  O << "using nullptr_t = std::nullptr_t;\n";
 
   O << "\n";
 

clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp

@@ -0,0 +1,63 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h %s
+// RUN: FileCheck -input-file=%t.h %s
+//
+// CHECK: #include <cstddef>
+// CHECK-NEXT: #include <CL/sycl/detail/defines.hpp>
+// CHECK-NEXT: #include <CL/sycl/detail/kernel_desc.hpp>
+// CHECK-NEXT: using nullptr_t = std::nullptr_t;
+//
+// CHECK: static constexpr
+// CHECK-NEXT: const char* const kernel_names[] = {
+// CHECK-NEXT:   "_ZTSDn"
+// CHECK-NEXT:   "_ZTSSt4byte"
+// CHECK-NEXT:   "_ZTSm",
+// CHECK-NEXT:   "_ZTSl"
+// CHECK-NEXT: };
+//
+// CHECK: static constexpr
+// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
+// CHECK-NEXT:   //--- _ZTSDn
+// CHECK-EMPTY:
+// CHECK-NEXT:   //--- _ZTSSt4byte
+// CHECK-EMPTY:
+// CHECK-NEXT:   //--- _ZTSm
+// CHECK-EMPTY:
+// CHECK-NEXT:   //--- _ZTSl
+// CHECK-EMPTY:
+// CHECK-NEXT: };
+//
+// CHECK: static constexpr
+// CHECK-NEXT: const unsigned kernel_signature_start[] = {
+// CHECK-NEXT:  0, // _ZTSDn
+// CHECK-NEXT:  1, // _ZTSSt4byte
+// CHECK-NEXT:  2, // _ZTSm
+// CHECK-NEXT:  3 // _ZTSl
+// CHECK-NEXT: };
+
+// CHECK: template <> struct KernelInfo<nullptr_t> {
+// CHECK: template <> struct KernelInfo<::std::byte> {
+// CHECK: template <> struct KernelInfo<unsigned long> {
+// CHECK: template <> struct KernelInfo<long> {
+
+void usage() {
+}
+
+namespace std {
+typedef long unsigned int size_t;
+typedef long int ptrdiff_t;
+typedef decltype(nullptr) nullptr_t;
+enum class byte : unsigned char {};
+} // namespace std
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+  kernel_single_task<std::nullptr_t>([]() { usage(); });
+  kernel_single_task<std::byte>([=]() {});
+  kernel_single_task<std::size_t>([=]() {});
+  kernel_single_task<std::ptrdiff_t>([=]() {});
+  return 0;
+}

clang/test/SemaSYCL/unnamed-kernel.cpp

@@ -1,16 +1,22 @@
 // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s
 // RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s
-#include <sycl.hpp>
 
 #ifdef __SYCL_UNNAMED_LAMBDA__
 // expected-no-diagnostics
 #endif
 
+#include <sycl.hpp>
+
 namespace namespace1 {
 template <typename T>
 class KernelName;
 }
 
+namespace std {
+typedef struct {
+} max_align_t;
+} // namespace std
+
 struct MyWrapper {
 private:
   class InvalidKernelName0 {};
@@ -41,15 +47,15 @@ struct MyWrapper {
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
     // expected-error@+4 {{kernel needs to have a globally-visible name}}
-    // expected-note@16 {{InvalidKernelName0 declared here}}
+    // expected-note@21 {{InvalidKernelName0 declared here}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<InvalidKernelName0>([] {});
     });
 
 #ifndef __SYCL_UNNAMED_LAMBDA__
     // expected-error@+4 {{kernel needs to have a globally-visible name}}
-    // expected-note@17 {{InvalidKernelName3 declared here}}
+    // expected-note@22 {{InvalidKernelName3 declared here}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<namespace1::KernelName<InvalidKernelName3>>([] {});
@@ -60,10 +66,17 @@ struct MyWrapper {
       h.single_task<ValidAlias>([] {});
     });
 
+#ifndef __SYCL_UNNAMED_LAMBDA__
+    // expected-error@+3 {{kernel name is missing}}
+#endif
+    q.submit([&](cl::sycl::handler &h) {
+      h.single_task<std::max_align_t>([] {});
+    });
+
     using InvalidAlias = InvalidKernelName4;
 #ifndef __SYCL_UNNAMED_LAMBDA__
     // expected-error@+4 {{kernel needs to have a globally-visible name}}
-    // expected-note@18 {{InvalidKernelName4 declared here}}
+    // expected-note@23 {{InvalidKernelName4 declared here}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<InvalidAlias>([] {});
@@ -72,7 +85,7 @@ struct MyWrapper {
     using InvalidAlias1 = InvalidKernelName5;
 #ifndef __SYCL_UNNAMED_LAMBDA__
     // expected-error@+4 {{kernel needs to have a globally-visible name}}
-    // expected-note@19 {{InvalidKernelName5 declared here}}
+    // expected-note@24 {{InvalidKernelName5 declared here}}
 #endif
     q.submit([&](cl::sycl::handler &h) {
       h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});