Merge remote-tracking branch 'upstream/sycl' into compress_img

.github/CODEOWNERS

@@ -129,6 +129,7 @@ clang/test/Driver/sycl-native-cpu*.cpp @intel/dpcpp-nativecpu-pi-reviewers
 sycl/**/native_cpu/ @intel/dpcpp-nativecpu-pi-reviewers
 sycl/doc/design/SYCLNativeCPU.md @intel/dpcpp-nativecpu-pi-reviewers
 sycl/include/sycl/detail/native_cpu.hpp @intel/dpcpp-nativecpu-pi-reviewers
+libdevice/nativecpu* @intel/dpcpp-nativecpu-pi-reviewers
 
 # SYCL-Graphs extensions
 sycl/include/sycl/ext/oneapi/experimental/graph.hpp @intel/sycl-graphs-reviewers

clang/include/clang/Basic/DiagnosticDriverKinds.td

@@ -217,7 +217,8 @@ def warn_drv_mismatch_fpga_archive : Warning<
 def warn_drv_sycl_native_cpu_and_targets: Warning<
   "-fsycl-targets=native_cpu overrides SYCL targets option">,
   InGroup<SyclNativeCPUTargets>;
-def err_drv_unsupported_opt_dpcpp : Error<"option '%0' unsupported with DPC++">;
+def err_drv_unsupported_opt_sycl : Error<
+  "option '%0' not supported with SYCL compilation">;
 def err_drv_argument_only_allowed_with : Error<
   "invalid argument '%0' only allowed with '%1'">;
 def err_drv_opt_unsupported_input_type : Error<
@@ -389,8 +390,6 @@ def err_drv_fsycl_with_c_type : Error<
   "'%0' must not be used in conjunction with '-fsycl', which expects C++ source">;
 def err_drv_fsycl_with_pch : Error<
   "precompiled header generation is not supported with '-fsycl'">;
-def err_drv_fsycl_unsupported_with_opt
-  : Error<"'%0' is not supported with '-fsycl'">;
 def warn_drv_opt_requires_opt
   : Warning<"'%0' should be used only in conjunction with '%1'">, InGroup<UnusedCommandLineArgument>;
 def err_drv_sycl_missing_amdgpu_arch : Error<

clang/lib/Driver/Driver.cpp

@@ -841,17 +841,16 @@ static bool addSYCLDefaultTriple(Compilation &C,
     return false;
   if (C.getInputArgs().hasArg(options::OPT_fsycl_force_target_EQ))
     return false;
+  llvm::Triple DefaultTriple =
+      C.getDriver().MakeSYCLDeviceTriple(getDefaultSYCLArch(C));
   for (const auto &SYCLTriple : SYCLTriples) {
-    if (SYCLTriple.getSubArch() == llvm::Triple::NoSubArch &&
-        SYCLTriple.isSPIROrSPIRV())
+    if (SYCLTriple == DefaultTriple)
       return false;
     // If we encounter a known non-spir* target, do not add the default triple.
     if (SYCLTriple.isNVPTX() || SYCLTriple.isAMDGCN())
       return false;
   }
   // Add the default triple as it was not found.
-  llvm::Triple DefaultTriple =
-      C.getDriver().MakeSYCLDeviceTriple(getDefaultSYCLArch(C));
   SYCLTriples.insert(SYCLTriples.begin(), DefaultTriple);
   return true;
 }
@@ -1069,9 +1068,8 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
   //
   // SYCL
   //
-  // We need to generate a SYCL toolchain if the user specified targets with
-  // the -fsycl-targets. If -fsycl is supplied without any of these we will
-  // assume SPIR-V.
+  // We need to generate a SYCL toolchain if the user specified -fsycl.
+  // If -fsycl is supplied without any of these we will assume SPIR-V.
   // Use of -fsycl-device-only overrides -fsycl.
   bool HasValidSYCLRuntime =
       C.getInputArgs().hasFlag(options::OPT_fsycl, options::OPT_fno_sycl,
@@ -1099,7 +1097,6 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
   };
 
   Arg *SYCLTargets = getArgRequiringSYCLRuntime(options::OPT_fsycl_targets_EQ);
-  Arg *SYCLLink = getArgRequiringSYCLRuntime(options::OPT_fsycl_link_EQ);
 
   // Check if -fsycl-host-compiler is used in conjunction with -fsycl.
   Arg *SYCLHostCompiler =
@@ -1121,7 +1118,7 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
     if (!HasValidSYCLRuntime)
       return;
     if (Arg *IncompatArg = C.getInputArgs().getLastArg(OptId))
-      Diag(clang::diag::err_drv_fsycl_unsupported_with_opt)
+      Diag(clang::diag::err_drv_unsupported_opt_sycl)
           << IncompatArg->getSpelling();
   };
   // -static-libstdc++ is not compatible with -fsycl.
@@ -1147,9 +1144,12 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
     Diag(clang::diag::err_drv_invalid_argument_to_option)
         << ArgValue << A->getOption().getName();
   };
+
+  Arg *SYCLLink = getArgRequiringSYCLRuntime(options::OPT_fsycl_link_EQ);
+  checkSingleArgValidity(SYCLLink, {"early", "image"});
+
   Arg *DeviceCodeSplit =
       C.getInputArgs().getLastArg(options::OPT_fsycl_device_code_split_EQ);
-  checkSingleArgValidity(SYCLLink, {"early", "image"});
   checkSingleArgValidity(DeviceCodeSplit,
                          {"per_kernel", "per_source", "auto", "off"});
 
@@ -1276,8 +1276,7 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
             << SYCLTargetsValues->getAsString(C.getInputArgs());
     }
   } else {
-    // If -fsycl is supplied without -fsycl-*targets we will assume SPIR-V
-    // unless -fintelfpga is supplied, which uses SPIR-V with fpga AOT.
+    // If -fsycl is supplied without -fsycl-targets we will assume SPIR-V.
     // For -fsycl-device-only, we also setup the implied triple as needed.
     if (HasValidSYCLRuntime) {
       StringRef SYCLTargetArch = getDefaultSYCLArch(C);
@@ -1295,10 +1294,10 @@ void Driver::CreateOffloadingDeviceToolChains(Compilation &C,
     for (auto &TT : UniqueSYCLTriplesVec) {
       if (TT.isNVPTX() || TT.isAMDGCN()) {
         Diag(diag::warn_flag_no_sycl_libspirv) << TT.getTriple();
-      } else {
-        Diag(diag::warn_drv_unsupported_option_for_target)
-            << "-fno-sycl-libspirv" << TT.getTriple();
+        continue;
       }
+      Diag(diag::warn_drv_unsupported_option_for_target)
+          << "-fno-sycl-libspirv" << TT.getTriple();
     }
   }
   // -fsycl-fp64-conv-emu is valid only for AOT compilation with an Intel GPU

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5117,7 +5117,7 @@ static void ProcessVSRuntimeLibrary(const ToolChain &TC, const ArgList &Args,
       (RTOptionID == options::OPT__SLASH_MT ||
        RTOptionID == options::OPT__SLASH_MTd))
     // Use of /MT or /MTd is not supported for SYCL.
-    TC.getDriver().Diag(diag::err_drv_unsupported_opt_dpcpp)
+    TC.getDriver().Diag(diag::err_drv_unsupported_opt_sycl)
         << SetArg->getOption().getName();
 
   enum { addDEBUG = 0x1, addMT = 0x2, addDLL = 0x4 };
@@ -5427,8 +5427,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
                      options::OPT_fno_sycl_dead_args_optimization,
                      isSYCLOptimizationO2orHigher(Args)))
       CmdArgs.push_back("-fenable-sycl-dae");
-    bool IsMSVC = AuxT.isWindowsMSVCEnvironment();
-    if (IsMSVC) {
+    if (IsWindowsMSVC) {
       CmdArgs.push_back("-fms-extensions");
       CmdArgs.push_back("-fms-compatibility");
       CmdArgs.push_back("-fdelayed-template-parsing");
@@ -5634,14 +5633,12 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
       // doing the host pass.
       CmdArgs.push_back("-fsycl-is-host");
 
-      if (!D.IsCLMode()) {
+      if (!D.IsCLMode() && IsWindowsMSVC &&
+          !Args.hasArg(options::OPT_fms_runtime_lib_EQ)) {
         // SYCL library is guaranteed to work correctly only with dynamic
         // MSVC runtime.
-        llvm::Triple AuxT = C.getDefaultToolChain().getTriple();
-        if (AuxT.isWindowsMSVCEnvironment()) {
-          CmdArgs.push_back("-D_MT");
-          CmdArgs.push_back("-D_DLL");
-        }
+        CmdArgs.push_back("-D_MT");
+        CmdArgs.push_back("-D_DLL");
       }
     }
     // Add any predefined macros associated with intel_gpu* type targets
@@ -6725,16 +6722,17 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   // the command line.  This is to allow for CMake based builds using the
   // Linux based driver on Windows to correctly pull in the expected debug
   // library.
-  if (Args.hasArg(options::OPT_fsycl) && !Args.hasArg(options::OPT_nolibsycl)) {
-    if (!D.IsCLMode() && TC.getTriple().isWindowsMSVCEnvironment()) {
+  if (Args.hasArg(options::OPT_fsycl) && !Args.hasArg(options::OPT_nolibsycl) &&
+      !D.IsCLMode()) {
+    if (TC.getTriple().isWindowsMSVCEnvironment()) {
       if (isDependentLibAdded(Args, "msvcrtd")) {
         if (Args.hasArg(options::OPT_fpreview_breaking_changes))
           CmdArgs.push_back("--dependent-lib=sycl" SYCL_MAJOR_VERSION
                             "-previewd");
         else
           CmdArgs.push_back("--dependent-lib=sycl" SYCL_MAJOR_VERSION "d");
       }
-    } else if (!D.IsCLMode() && TC.getTriple().isWindowsGNUEnvironment()) {
+    } else if (TC.getTriple().isWindowsGNUEnvironment()) {
       if (Args.hasArg(options::OPT_fpreview_breaking_changes))
         CmdArgs.push_back("--dependent-lib=sycl" SYCL_MAJOR_VERSION
                           "-preview.dll");
@@ -10551,7 +10549,8 @@ static void getTripleBasedSPIRVTransOpts(Compilation &C,
               ",+SPV_INTEL_tensor_float32_conversion"
               ",+SPV_INTEL_optnone"
               ",+SPV_KHR_non_semantic_info"
-              ",+SPV_KHR_cooperative_matrix";
+              ",+SPV_KHR_cooperative_matrix"
+              ",+SPV_INTEL_memory_access_aliasing";
   if (IsCPU)
     ExtArg += ",+SPV_INTEL_fp_max_error";
 

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -1798,8 +1798,10 @@ SYCLToolChain::GetCXXStdlibType(const ArgList &Args) const {
 void SYCLToolChain::AddSYCLIncludeArgs(const clang::driver::Driver &Driver,
                                        const ArgList &DriverArgs,
                                        ArgStringList &CC1Args) {
-  // Add ../include/sycl, ../include/sycl/stl_wrappers and ../include (in that
-  // order).
+  // Add the SYCL header search locations in the specified order.
+  //   ../include/sycl
+  //   ../include/sycl/stl_wrappers
+  //   ../include
   SmallString<128> IncludePath(Driver.Dir);
   llvm::sys::path::append(IncludePath, "..");
   llvm::sys::path::append(IncludePath, "include");

clang/lib/Sema/SemaSYCL.cpp

@@ -6029,6 +6029,23 @@ static void OutputStableNameInChars(raw_ostream &O, StringRef Name) {
   }
 }
 
+static void EmitPragmaDiagnosticPush(raw_ostream &O, StringRef DiagName) {
+  O << "\n";
+  O << "#ifdef __clang__\n";
+  O << "#pragma clang diagnostic push\n";
+  O << "#pragma clang diagnostic ignored \"" << DiagName.str() << "\"\n";
+  O << "#endif // defined(__clang__)\n";
+  O << "\n";
+}
+
+static void EmitPragmaDiagnosticPop(raw_ostream &O) {
+  O << "\n";
+  O << "#ifdef __clang__\n";
+  O << "#pragma clang diagnostic pop\n";
+  O << "#endif // defined(__clang__)\n";
+  O << "\n";
+}
+
 void SYCLIntegrationHeader::emit(raw_ostream &O) {
   O << "// This is auto-generated SYCL integration header.\n";
   O << "\n";
@@ -6127,6 +6144,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
   // main() function.
 
   if (NeedToEmitDeviceGlobalRegistration) {
+    // Supress the reserved identifier diagnostic that clang generates
+    // for the construct below.
+    EmitPragmaDiagnosticPush(O, "-Wreserved-identifier");
     O << "namespace {\n";
 
     O << "class __sycl_device_global_registration {\n";
@@ -6138,12 +6158,16 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     O << "} // namespace\n";
 
     O << "\n";
+    EmitPragmaDiagnosticPop(O);
   }
 
   // Generate declaration of variable of type __sycl_host_pipe_registration
   // whose sole purpose is to run its constructor before the application's
   // main() function.
   if (NeedToEmitHostPipeRegistration) {
+    // Supress the reserved identifier diagnostic that clang generates
+    // for the construct below.
+    EmitPragmaDiagnosticPush(O, "-Wreserved-identifier");
     O << "namespace {\n";
 
     O << "class __sycl_host_pipe_registration {\n";
@@ -6155,6 +6179,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
     O << "} // namespace\n";
 
     O << "\n";
+    EmitPragmaDiagnosticPop(O);
   }
 
 
@@ -6169,6 +6194,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
       O << ",";
     O << "\n";
   }
+  // Add a sentinel to avoid warning if the collection is empty
+  // (similar to what we do for kernel_signatures below).
+  O << "  \"\",\n";
   O << "};\n\n";
 
   O << "// array representing signatures of all kernels defined in the\n";
@@ -6720,12 +6748,16 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
     OS << "#include <sycl/detail/device_global_map.hpp>\n";
     DeviceGlobOS.flush();
     OS << "namespace sycl::detail {\n";
+    // Supress the old-style case diagnostic that clang generates
+    // for the construct below in DeviceGlobalsBuf.
+    EmitPragmaDiagnosticPush(OS, "-Wold-style-cast");
     OS << "namespace {\n";
     OS << "__sycl_device_global_registration::__sycl_device_global_"
           "registration() noexcept {\n";
     OS << DeviceGlobalsBuf;
     OS << "}\n";
     OS << "} // namespace (unnamed)\n";
+    EmitPragmaDiagnosticPop(OS);
     OS << "} // namespace sycl::detail\n";
 
     S.getSyclIntegrationHeader().addDeviceGlobalRegistration();
@@ -6735,12 +6767,16 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) {
     OS << "#include <sycl/detail/host_pipe_map.hpp>\n";
     HostPipesOS.flush();
     OS << "namespace sycl::detail {\n";
+    // Supress the old-style case diagnostic that clang generates
+    // for the construct below in HostPipesBuf.
+    EmitPragmaDiagnosticPush(OS, "-Wold-style-cast");
     OS << "namespace {\n";
     OS << "__sycl_host_pipe_registration::__sycl_host_pipe_"
           "registration() noexcept {\n";
     OS << HostPipesBuf;
     OS << "}\n";
     OS << "} // namespace (unnamed)\n";
+    EmitPragmaDiagnosticPop(OS);
     OS << "} // namespace sycl::detail\n";
 
     S.getSyclIntegrationHeader().addHostPipeRegistration();

clang/test/CodeGenSYCL/device_global_int_footer_header.cpp

@@ -16,13 +16,20 @@ int main() {
 // CHECK-HEADER: namespace sycl {
 // CHECK-HEADER-NEXT: inline namespace _V1 {
 // CHECK-HEADER-NEXT: namespace detail {
-// CHECK-HEADER-NEXT: namespace {
+// CHECK-HEADER: #ifdef __clang__
+// CHECK-HEADER-NEXT: #pragma clang diagnostic push
+// CHECK-HEADER-NEXT: #pragma clang diagnostic ignored "-Wreserved-identifier"
+// CHECK-HEADER-NEXT: #endif // defined(__clang__)
+// CHECK-HEADER: namespace {
 // CHECK-HEADER-NEXT: class __sycl_device_global_registration {
 // CHECK-HEADER-NEXT: public:
 // CHECK-HEADER-NEXT:   __sycl_device_global_registration() noexcept;
 // CHECK-HEADER-NEXT: };
 // CHECK-HEADER-NEXT: __sycl_device_global_registration __sycl_device_global_registrar;
 // CHECK-HEADER-NEXT: } // namespace
+// CHECK-HEADER: #ifdef __clang__
+// CHECK-HEADER-NEXT: #pragma clang diagnostic pop
+// CHECK-HEADER-NEXT: #endif // defined(__clang__)
 // CHECK-HEADER: } // namespace detail
 // CHECK-HEADER: } // namespace _V1
 // CHECK-HEADER: } // namespace sycl
@@ -49,7 +56,11 @@ int main() {
 
 // CHECK-FOOTER: #include <sycl/detail/device_global_map.hpp>
 // CHECK-FOOTER: namespace sycl::detail {
-// CHECK-FOOTER-NEXT: namespace {
+// CHECK-FOOTER: #ifdef __clang__
+// CHECK-FOOTER-NEXT: #pragma clang diagnostic push
+// CHECK-FOOTER-NEXT: #pragma clang diagnostic ignored "-Wold-style-cast"
+// CHECK-FOOTER-NEXT: #endif // defined(__clang__)
+// CHECK-FOOTER: namespace {
 // CHECK-FOOTER-NEXT: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {
 
 extern device_global<int> Basic;
@@ -111,3 +122,6 @@ struct HasVarTemplate {
 } // namespace
 const auto x = HasVarTemplate::VarTempl<int>.get();
 // CHECK-FOOTER-NEXT: device_global_map::add((void *)&::__sycl_detail::__shim_[[SHIM1]](), "THE_PREFIX____ZN12_GLOBAL__N_114HasVarTemplate8VarTemplIiEE");
+// CHECK-FOOTER: #ifdef __clang__
+// CHECK-FOOTER-NEXT: #pragma clang diagnostic pop
+// CHECK-FOOTER-NEXT: #endif // defined(__clang__)

clang/test/CodeGenSYCL/device_globals_with_spec_ids.cpp

@@ -249,7 +249,7 @@ struct Wrapper {
 // CHECK: #include <sycl/detail/spec_const_integration.hpp>
 // CHECK-NEXT: #include <sycl/detail/device_global_map.hpp>
 // CHECK-NEXT: namespace sycl::detail {
-// CHECK-NEXT: namespace {
+// CHECK: namespace {
 // CHECK-NEXT: __sycl_device_global_registration::__sycl_device_global_registration() noexcept {
 // CHECK-NEXT: device_global_map::add((void *)&::b, "_Z1b");
 // CHECK-NEXT: device_global_map::add((void *)&::Wrapper::b, "_ZN7Wrapper1bE");
@@ -260,4 +260,4 @@ struct Wrapper {
 // CHECK-NEXT: device_global_map::add((void *)&::outer::__sycl_detail::__shim_[[SHIM11]](), "____ZN5outer12_GLOBAL__N_15inner12_GLOBAL__N_17Wrapper1cE");
 // CHECK-NEXT: }
 // CHECK-NEXT: } // namespace (unnamed)
-// CHECK-NEXT: } // namespace sycl::detail
+// CHECK: } // namespace sycl::detail

clang/test/CodeGenSYCL/free_function_int_header.cpp

@@ -49,6 +49,7 @@ template <> void ff_3<double>(double *ptr, double start, double end) {
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_3IiEvPT_S0_S0_
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_3IfEvPT_S0_S0_
 // CHECK-NEXT:   {{.*}}__sycl_kernel_ff_3IdEvPT_S0_S0_
+// CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
 // CHECK:      const kernel_param_desc_t kernel_signatures[] = {
@@ -201,4 +202,4 @@ template <> void ff_3<double>(double *ptr, double start, double end) {
 // CHECK-NEXT:   kernel_id ext::oneapi::experimental::get_kernel_id<__sycl_shim5()>() {
 // CHECK-NEXT:     return sycl::detail::get_kernel_id_impl(std::string_view{"_Z18__sycl_kernel_ff_3IdEvPT_S0_S0_"});
 // CHECK-NEXT:   }
-// CHECK-NEXT: }
+// CHECK-NEXT: }

clang/test/CodeGenSYCL/host_pipe_int_footer_header.cpp

@@ -23,7 +23,7 @@ void foo() {
 // CHECK-HEADER: namespace sycl {
 // CHECK-HEADER-NEXT: inline namespace _V1 {
 // CHECK-HEADER-NEXT: namespace detail {
-// CHECK-HEADER-NEXT: namespace {
+// CHECK-HEADER: namespace {
 // CHECK-HEADER-NEXT: class __sycl_host_pipe_registration {
 // CHECK-HEADER-NEXT: public:
 // CHECK-HEADER-NEXT: __sycl_host_pipe_registration() noexcept;
@@ -37,7 +37,7 @@ void foo() {
 // CHECK-FOOTER: #include <sycl/detail/defines_elementary.hpp>
 // CHECK-FOOTER: #include <sycl/detail/host_pipe_map.hpp>
 // CHECK-FOOTER-NEXT: namespace sycl::detail {
-// CHECK-FOOTER-NEXT: namespace {
+// CHECK-FOOTER: namespace {
 // CHECK-FOOTER-NEXT: __sycl_host_pipe_registration::__sycl_host_pipe_registration() noexcept {
 
 // CHECK-FOOTER: host_pipe_map::add((void *)&::sycl::ext::intel::experimental::host_pipe<HPInt, int>::__pipe, "THE_PREFIX____ZN4sycl3_V13ext5intel12experimental9host_pipeIZZZ3foovENKUlRNS0_7handlerEE_clES6_ENKUlvE_clEvE5HPIntiE6__pipeE");

clang/test/CodeGenSYCL/int-header-empty-signatures.cpp

@@ -7,6 +7,7 @@
 // CHECK: static constexpr
 // CHECK-NEXT: const char* const kernel_names[] = {
 // CHECK-NEXT:   "_ZTSZ4mainE1K"
+// CHECK-NEXT:   ""
 // CHECK-NEXT: };
 
 // CHECK: static constexpr