Merge from 'sycl' to 'sycl-web' (#2)

Author: iclsrc

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10988,6 +10988,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/Driver/Driver.cpp

@@ -1791,20 +1791,21 @@ llvm::Triple Driver::MakeSYCLDeviceTriple(StringRef TargetArch) const {
 // Print the help from any of the given tools which are used for AOT
 // compilation for SYCL
 void Driver::PrintSYCLToolHelp(const Compilation &C) const {
-  SmallVector<std::tuple<llvm::Triple, StringRef, StringRef>, 4> HelpArgs;
+  SmallVector<std::tuple<llvm::Triple, StringRef, StringRef, StringRef>, 4>
+      HelpArgs;
   // Populate the vector with the tools and help options
   if (Arg *A = C.getArgs().getLastArg(options::OPT_fsycl_help_EQ)) {
     StringRef AV(A->getValue());
     llvm::Triple T;
     if (AV == "gen" || AV == "all")
       HelpArgs.push_back(std::make_tuple(MakeSYCLDeviceTriple("spir64_gen"),
-                                         "ocloc", "--help"));
+                                         "ocloc", "--help", ""));
     if (AV == "fpga" || AV == "all")
-      HelpArgs.push_back(
-          std::make_tuple(MakeSYCLDeviceTriple("spir64_fpga"), "aoc", "-help"));
+      HelpArgs.push_back(std::make_tuple(MakeSYCLDeviceTriple("spir64_fpga"),
+                                         "aoc", "-help", "-sycl"));
     if (AV == "x86_64" || AV == "all")
       HelpArgs.push_back(std::make_tuple(MakeSYCLDeviceTriple("spir64_x86_64"),
-                                         "opencl-aot", "--help"));
+                                         "opencl-aot", "--help", ""));
     if (HelpArgs.empty()) {
       C.getDriver().Diag(diag::err_drv_unsupported_option_argument)
                          << A->getOption().getName() << AV;
@@ -1817,7 +1818,8 @@ void Driver::PrintSYCLToolHelp(const Compilation &C) const {
     llvm::outs() << "Emitting help information for " << std::get<1>(HA) << '\n'
         << "Use triple of '" << std::get<0>(HA).normalize() <<
         "' to enable ahead of time compilation\n";
-    std::vector<StringRef> ToolArgs = { std::get<1>(HA), std::get<2>(HA) };
+    std::vector<StringRef> ToolArgs = {std::get<1>(HA), std::get<2>(HA),
+                                       std::get<3>(HA)};
     SmallString<128> ExecPath(
         C.getDefaultToolChain().GetProgramPath(std::get<1>(HA).data()));
     auto ToolBinary = llvm::sys::findProgramByName(ExecPath);
@@ -1827,7 +1829,10 @@ void Driver::PrintSYCLToolHelp(const Compilation &C) const {
     }
     // do not run the tools with -###.
     if (C.getArgs().hasArg(options::OPT__HASH_HASH_HASH)) {
-      llvm::errs() << "\"" << ExecPath << "\" \"" << ToolArgs[1] << "\"\n";
+      llvm::errs() << "\"" << ExecPath << "\" \"" << ToolArgs[1] << "\"";
+      if (!ToolArgs[2].empty())
+        llvm::errs() << " \"" << ToolArgs[2] << "\"";
+      llvm::errs() << "\n";
       continue;
     }
     // Run the Tool.

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/Driver/sycl.c

@@ -75,7 +75,7 @@
 // SYCL-HELP-BADARG: unsupported argument 'foo' to option 'fsycl-help='
 // SYCL-HELP-GEN: Emitting help information for ocloc
 // SYCL-HELP-GEN: Use triple of 'spir64_gen-unknown-unknown-sycldevice' to enable ahead of time compilation
-// SYCL-HELP-FPGA-OUT: "[[DIR]]{{[/\\]+}}aoc" "-help"
+// SYCL-HELP-FPGA-OUT: "[[DIR]]{{[/\\]+}}aoc" "-help" "-sycl"
 // SYCL-HELP-FPGA: Emitting help information for aoc
 // SYCL-HELP-FPGA: Use triple of 'spir64_fpga-unknown-unknown-sycldevice' to enable ahead of time compilation
 // SYCL-HELP-CPU: Emitting help information for opencl-aot

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;