[SYCL] Support or diagnose use of namespace std types as kernel type names (#1963)

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: jtmott-intel

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10982,6 +10982,7 @@ def err_sycl_kernel_incorrectly_named : Error<
   "kernel %select{name is missing"
   "|needs to have a globally-visible name"
   "|name is invalid. Unscoped enum requires fixed underlying type"
+  "|name cannot be a type in the \"std\" namespace"
   "}0">;
 def err_sycl_kernel_not_function_object
     : Error<"kernel parameter must be a lambda or function object">;

clang/lib/Sema/SemaSYCL.cpp

@@ -2908,6 +2908,13 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
       }
       break;
     }
+
+    if (NS->isStdNamespace()) {
+      Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
+          << /* name cannot be a type in the std namespace */ 3;
+      return;
+    }
+
     ++NamespaceCnt;
     const StringRef NSInlinePrefix = NS->isInline() ? "inline " : "";
     NSStr.insert(
@@ -2990,8 +2997,13 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
     ;
   const CXXRecordDecl *RD = T->getAsCXXRecordDecl();
 
-  if (!RD)
+  if (!RD) {
+    if (T->isNullPtrType())
+      Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
+          << /* name cannot be a type in the std namespace */ 3;
+
     return;
+  }
 
   // see if this is a template specialization ...
   if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {

clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp

@@ -0,0 +1,63 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -DCHECK_ERROR -verify %s
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s
+// RUN: FileCheck -input-file=%t.h %s
+//
+// CHECK: #include <CL/sycl/detail/defines.hpp>
+// CHECK-NEXT: #include <CL/sycl/detail/kernel_desc.hpp>
+//
+// CHECK: static constexpr
+// CHECK-NEXT: const char* const kernel_names[] = {
+// CHECK-NEXT:   "_ZTSm",
+// CHECK-NEXT:   "_ZTSl"
+// CHECK-NEXT: };
+//
+// CHECK: static constexpr
+// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = {
+// 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, // _ZTSm
+// CHECK-NEXT:  1 // _ZTSl
+// CHECK-NEXT: };
+
+// 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;
+class T;
+class U;
+} // namespace std
+
+template <typename T>
+struct Templated_kernel_name;
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  kernelFunc();
+}
+
+int main() {
+#ifdef CHECK_ERROR
+  kernel_single_task<std::nullptr_t>([=]() {});                        // expected-error {{kernel name cannot be a type in the "std" namespace}}
+  kernel_single_task<std::T>([=]() {});                                // expected-error {{kernel name cannot be a type in the "std" namespace}}
+  kernel_single_task<Templated_kernel_name<std::nullptr_t>>([=]() {}); // expected-error {{kernel name cannot be a type in the "std" namespace}}
+  kernel_single_task<Templated_kernel_name<std::U>>([=]() {});         // expected-error {{kernel name cannot be a type in the "std" namespace}}
+#endif
+
+  // Although in the std namespace, these resolve to builtins such as `int` that are allowed in kernel names
+  kernel_single_task<std::size_t>([=]() {});
+  kernel_single_task<std::ptrdiff_t>([=]() {});
+
+  return 0;
+}

clang/test/SemaSYCL/unnamed-kernel.cpp

@@ -11,6 +11,11 @@ template <typename T>
 class KernelName;
 }
 
+namespace std {
+typedef struct {
+} max_align_t;
+} // namespace std
+
 struct MyWrapper {
 private:
   class InvalidKernelName0 {};
@@ -41,15 +46,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 +65,17 @@ struct MyWrapper {
       h.single_task<ValidAlias>([] {});
     });
 
+#ifndef __SYCL_UNNAMED_LAMBDA__
+    // expected-error@+3 {{kernel name cannot be a type in the "std" namespace}}
+#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 +84,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>>([] {});

sycl/test/group-algorithm/broadcast.cpp

@@ -15,14 +15,11 @@
 using namespace sycl;
 using namespace sycl::ONEAPI;
 
-template <typename InputContainer, typename OutputContainer>
-class broadcast_kernel;
-
-template <typename InputContainer, typename OutputContainer>
+template <typename kernel_name, typename InputContainer,
+          typename OutputContainer>
 void test(queue q, InputContainer input, OutputContainer output) {
   typedef typename InputContainer::value_type InputT;
   typedef typename OutputContainer::value_type OutputT;
-  typedef class broadcast_kernel<InputContainer, OutputContainer> kernel_name;
   size_t N = input.size();
   size_t G = 4;
   range<2> R(G, G);
@@ -63,7 +60,7 @@ int main() {
     std::array<int, 3> output;
     std::iota(input.begin(), input.end(), 1);
     std::fill(output.begin(), output.end(), false);
-    test(q, input, output);
+    test<class KernelName_EFL>(q, input, output);
   }
 
   // Test pointer type
@@ -74,7 +71,7 @@ int main() {
       input[i] = static_cast<int *>(0x0) + i;
     }
     std::fill(output.begin(), output.end(), static_cast<int *>(0x0));
-    test(q, input, output);
+    test<class KernelName_NrqELzFQToOSPsRNMi>(q, input, output);
   }
 
   // Test user-defined type
@@ -88,7 +85,7 @@ int main() {
           std::complex<float>(0, 1) + (float)i * std::complex<float>(2, 2);
     }
     std::fill(output.begin(), output.end(), std::complex<float>(0, 0));
-    test(q, input, output);
+    test<class KernelName_rCblcml>(q, input, output);
   }
   {
     std::array<std::complex<double>, N> input;
@@ -98,7 +95,7 @@ int main() {
           std::complex<double>(0, 1) + (double)i * std::complex<double>(2, 2);
     }
     std::fill(output.begin(), output.end(), std::complex<float>(0, 0));
-    test(q, input, output);
+    test<class KernelName_NCWhjnQ>(q, input, output);
   }
   std::cout << "Test passed." << std::endl;
 }

sycl/test/group-algorithm/exclusive_scan.cpp

@@ -24,7 +24,7 @@
 using namespace sycl;
 using namespace sycl::ONEAPI;
 
-template <class BinaryOperation, int TestNumber>
+template <class SpecializationKernelName, int TestNumber>
 class exclusive_scan_kernel;
 
 // std::exclusive_scan isn't implemented yet, so use serial implementation
@@ -44,17 +44,17 @@ OutputIterator exclusive_scan(InputIterator first, InputIterator last,
 }
 } // namespace emu
 
-template <typename InputContainer, typename OutputContainer,
-          class BinaryOperation>
+template <typename SpecializationKernelName, typename InputContainer,
+          typename OutputContainer, class BinaryOperation>
 void test(queue q, InputContainer input, OutputContainer output,
           BinaryOperation binary_op,
           typename OutputContainer::value_type identity) {
   typedef typename InputContainer::value_type InputT;
   typedef typename OutputContainer::value_type OutputT;
-  typedef class exclusive_scan_kernel<BinaryOperation, 0> kernel_name0;
-  typedef class exclusive_scan_kernel<BinaryOperation, 1> kernel_name1;
-  typedef class exclusive_scan_kernel<BinaryOperation, 2> kernel_name2;
-  typedef class exclusive_scan_kernel<BinaryOperation, 3> kernel_name3;
+  typedef class exclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
+  typedef class exclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
+  typedef class exclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
+  typedef class exclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
   OutputT init = 42;
   size_t N = input.size();
   size_t G = 16;
@@ -159,19 +159,24 @@ int main() {
   std::iota(input.begin(), input.end(), 0);
   std::fill(output.begin(), output.end(), 0);
 
-  test(q, input, output, plus<>(), 0);
-  test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
-  test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
+  test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
+  test<class KernelNameMinimumV>(q, input, output, minimum<>(),
+                                 std::numeric_limits<int>::max());
+  test<class KernelNameMaximumV>(q, input, output, maximum<>(),
+                                 std::numeric_limits<int>::lowest());
 
-  test(q, input, output, plus<int>(), 0);
-  test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
-  test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
+  test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
+  test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
+                                 std::numeric_limits<int>::max());
+  test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
+                                 std::numeric_limits<int>::lowest());
 
 #ifdef SPIRV_1_3
-  test(q, input, output, multiplies<int>(), 1);
-  test(q, input, output, bit_or<int>(), 0);
-  test(q, input, output, bit_xor<int>(), 0);
-  test(q, input, output, bit_and<int>(), ~0);
+  test<class KernelName_VzAPutpBRRJrQPB>(q, input, output, multiplies<int>(),
+                                         1);
+  test<class KernelName_UXdGbr>(q, input, output, bit_or<int>(), 0);
+  test<class KernelName_saYaodNyJknrPW>(q, input, output, bit_xor<int>(), 0);
+  test<class KernelName_GPcuAlvAOjrDyP>(q, input, output, bit_and<int>(), ~0);
 #endif // SPIRV_1_3
 
   std::cout << "Test passed." << std::endl;

sycl/test/group-algorithm/inclusive_scan.cpp

@@ -24,7 +24,7 @@
 using namespace sycl;
 using namespace sycl::ONEAPI;
 
-template <class BinaryOperation, int TestNumber>
+template <class SpecializationKernelName, int TestNumber>
 class inclusive_scan_kernel;
 
 // std::inclusive_scan isn't implemented yet, so use serial implementation
@@ -44,17 +44,17 @@ OutputIterator inclusive_scan(InputIterator first, InputIterator last,
 }
 } // namespace emu
 
-template <typename InputContainer, typename OutputContainer,
-          class BinaryOperation>
+template <typename SpecializationKernelName, typename InputContainer,
+          typename OutputContainer, class BinaryOperation>
 void test(queue q, InputContainer input, OutputContainer output,
           BinaryOperation binary_op,
           typename OutputContainer::value_type identity) {
   typedef typename InputContainer::value_type InputT;
   typedef typename OutputContainer::value_type OutputT;
-  typedef class inclusive_scan_kernel<BinaryOperation, 0> kernel_name0;
-  typedef class inclusive_scan_kernel<BinaryOperation, 1> kernel_name1;
-  typedef class inclusive_scan_kernel<BinaryOperation, 2> kernel_name2;
-  typedef class inclusive_scan_kernel<BinaryOperation, 3> kernel_name3;
+  typedef class inclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
+  typedef class inclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
+  typedef class inclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
+  typedef class inclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
   OutputT init = 42;
   size_t N = input.size();
   size_t G = 16;
@@ -159,19 +159,25 @@ int main() {
   std::iota(input.begin(), input.end(), 0);
   std::fill(output.begin(), output.end(), 0);
 
-  test(q, input, output, plus<>(), 0);
-  test(q, input, output, minimum<>(), std::numeric_limits<int>::max());
-  test(q, input, output, maximum<>(), std::numeric_limits<int>::lowest());
+  test<class KernelNamePlusV>(q, input, output, plus<>(), 0);
+  test<class KernelNameMinimumV>(q, input, output, minimum<>(),
+                                 std::numeric_limits<int>::max());
+  test<class KernelNameMaximumV>(q, input, output, maximum<>(),
+                                 std::numeric_limits<int>::lowest());
 
-  test(q, input, output, plus<int>(), 0);
-  test(q, input, output, minimum<int>(), std::numeric_limits<int>::max());
-  test(q, input, output, maximum<int>(), std::numeric_limits<int>::lowest());
+  test<class KernelNamePlusI>(q, input, output, plus<int>(), 0);
+  test<class KernelNameMinimumI>(q, input, output, minimum<int>(),
+                                 std::numeric_limits<int>::max());
+  test<class KernelNameMaximumI>(q, input, output, maximum<int>(),
+                                 std::numeric_limits<int>::lowest());
 
 #ifdef SPIRV_1_3
-  test(q, input, output, multiplies<int>(), 1);
-  test(q, input, output, bit_or<int>(), 0);
-  test(q, input, output, bit_xor<int>(), 0);
-  test(q, input, output, bit_and<int>(), ~0);
+  test<class KernelName_zMyjxUrBgeUGoxmDwhvJ>(q, input, output,
+                                              multiplies<int>(), 1);
+  test<class KernelName_SljjtroxNRaAXoVnT>(q, input, output, bit_or<int>(), 0);
+  test<class KernelName_yXIZfjwjxQGiPeQAnc>(q, input, output, bit_xor<int>(),
+                                            0);
+  test<class KernelName_xGnAnMYHvqekCk>(q, input, output, bit_and<int>(), ~0);
 #endif // SPIRV_1_3
 
   std::cout << "Test passed." << std::endl;