[SYCL] remove invoking kernel_compiler

sycl/include/sycl/kernel_bundle_enums.hpp

@@ -25,7 +25,6 @@ enum class source_language : int {
   spirv = 1,
   sycl = 2,
   /* cuda */
-  sycl_jit = 99 /* temporary, alternative implementation for SYCL */
 };
 
 // opencl versions

sycl/source/detail/kernel_bundle_impl.hpp

@@ -499,13 +499,7 @@ class kernel_bundle_impl {
           return Result;
         }
         if (Language == syclex::source_language::sycl) {
-          return syclex::detail::SYCL_to_SPIRV(*SourceStrPtr, IncludePairs,
-                                               BuildOptions, LogPtr,
-                                               RegisteredKernelNames);
-        }
-        if (Language == syclex::source_language::sycl_jit) {
-          const auto &SourceStr = std::get<std::string>(this->Source);
-          return syclex::detail::SYCL_JIT_to_SPIRV(SourceStr, IncludePairs,
+          return syclex::detail::SYCL_JIT_to_SPIRV(*SourceStrPtr, IncludePairs,
                                                    BuildOptions, LogPtr,
                                                    RegisteredKernelNames);
         }
@@ -578,8 +572,7 @@ class kernel_bundle_impl {
   std::string adjust_kernel_name(const std::string &Name,
                                  syclex::source_language Lang) {
     // Once name demangling support is in, we won't need this.
-    if (Lang != syclex::source_language::sycl &&
-        Lang != syclex::source_language::sycl_jit)
+    if (Lang != syclex::source_language::sycl)
       return Name;
 
     bool isMangled = Name.find("__sycl_kernel_") != std::string::npos;

sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp

@@ -9,322 +9,6 @@
 #include "kernel_compiler_sycl.hpp"
 #include <sycl/exception.hpp> // make_error_code
 
-#if __GNUC__ && __GNUC__ < 8
-
-// std::filesystem is not availalbe for GCC < 8
-// and much of the  cross-platform file handling code depends upon it.
-// Given that this extension is experimental and that the file
-// handling aspects are most likely temporary, it makes sense to
-// simply not support GCC<8.
-
-namespace sycl {
-inline namespace _V1 {
-namespace ext::oneapi::experimental {
-namespace detail {
-
-bool SYCL_Compilation_Available() { return false; }
-
-spirv_vec_t
-SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs,
-              const std::vector<std::string> &UserArgs, std::string *LogPtr,
-              const std::vector<std::string> &RegisteredKernelNames) {
-  (void)SYCLSource;
-  (void)IncludePairs;
-  (void)UserArgs;
-  (void)LogPtr;
-  (void)RegisteredKernelNames;
-  throw sycl::exception(sycl::errc::build,
-                        "kernel_compiler does not support GCC<8");
-}
-
-std::string userArgsAsString(const std::vector<std::string> &UserArguments) {
-  return std::accumulate(UserArguments.begin(), UserArguments.end(),
-                         std::string(""),
-                         [](const std::string &A, const std::string &B) {
-                           return A.empty() ? B : A + " " + B;
-                         });
-}
-
-} // namespace detail
-} // namespace ext::oneapi::experimental
-} // namespace _V1
-} // namespace sycl
-
-#else
-
-#include <sycl/detail/os_util.hpp>
-
-#include <ctime>
-#include <filesystem>
-#include <fstream>
-#include <random>
-#include <regex>
-#include <sstream>
-#include <stdio.h>
-
-namespace sycl {
-inline namespace _V1 {
-namespace ext::oneapi::experimental {
-namespace detail {
-
-std::string generateSemiUniqueId() {
-  auto Now = std::chrono::high_resolution_clock::now();
-  auto Milliseconds = std::chrono::duration_cast<std::chrono::milliseconds>(
-      Now.time_since_epoch());
-
-  // Generate random number between 10'000 and 99'900.
-  std::random_device RD;
-  std::mt19937 Gen(RD());
-  std::uniform_int_distribution<int> Distrib(10'000, 99'999);
-  int RandomNumber = Distrib(Gen);
-
-  // Combine time and random number into a string.
-  std::stringstream Ss;
-  Ss << Milliseconds.count() << "_" << std::setfill('0') << std::setw(5)
-     << RandomNumber;
-
-  return Ss.str();
-}
-
-std::filesystem::path prepareWS(const std::string &Id) {
-  namespace fs = std::filesystem;
-  const fs::path TmpDirectoryPath = fs::temp_directory_path();
-  fs::path NewDirectoryPath = TmpDirectoryPath / Id;
-
-  try {
-    fs::create_directories(NewDirectoryPath);
-    fs::permissions(NewDirectoryPath, fs::perms::owner_read |
-                                          fs::perms::owner_write |
-                                          fs::perms::owner_exec); // 0700
-
-  } catch (const fs::filesystem_error &E) {
-    throw sycl::exception(sycl::errc::build, E.what());
-  }
-
-  return NewDirectoryPath;
-}
-
-void deleteWS(const std::filesystem::path &ParentDir) {
-  try {
-    std::filesystem::remove_all(ParentDir);
-  } catch (const std::filesystem::filesystem_error &E) {
-    // We could simply suppress this, since deleting the directory afterwards
-    // is not critical. But if there are problems, seems good to know.
-    throw sycl::exception(sycl::errc::build, E.what());
-  }
-}
-
-std::string userArgsAsString(const std::vector<std::string> &UserArguments) {
-  return std::accumulate(UserArguments.begin(), UserArguments.end(),
-                         std::string(""),
-                         [](const std::string &A, const std::string &B) {
-                           return A.empty() ? B : A + " " + B;
-                         });
-}
-
-void outputPreamble(std::ofstream &Os, const std::filesystem::path &FilePath,
-                    const std::string &Id,
-                    const std::vector<std::string> &UserArgs) {
-
-  Os << "/*\n";
-  Os << "  clang++ -fsycl -o " << Id << ".bin ";
-  Os << userArgsAsString(UserArgs);
-  Os << " -fno-sycl-dead-args-optimization -fsycl-dump-device-code=./ " << Id;
-  Os << ".cpp \n */" << std::endl;
-}
-
-std::filesystem::path
-outputCpp(const std::filesystem::path &ParentDir, const std::string &Id,
-          std::string RawCodeString, const std::vector<std::string> &UserArgs,
-          const std::vector<std::string> &RegisteredKernelNames) {
-  std::filesystem::path FilePath = ParentDir / (Id + ".cpp");
-  std::ofstream Outfile(FilePath, std::ios::out | std::ios::trunc);
-
-  if (Outfile.is_open()) {
-    outputPreamble(Outfile, FilePath, Id, UserArgs);
-    Outfile << RawCodeString << std::endl;
-
-    // Temporarily needed until -c works with -fsycl-dump-spirv.
-    Outfile << "int main() {\n";
-    for (const std::string &KernelName : RegisteredKernelNames) {
-      Outfile << "  " << KernelName << ";\n";
-    }
-    Outfile << "  return 0;\n}\n" << std::endl;
-
-    Outfile.close();
-  } else {
-    throw sycl::exception(sycl::errc::build,
-                          "Failed to open .cpp file for write: " +
-                              FilePath.string());
-  }
-
-  return FilePath;
-}
-
-void outputIncludeFiles(const std::filesystem::path &Dirpath,
-                        include_pairs_t IncludePairs) {
-  using pairStrings = std::pair<std::string, std::string>;
-  for (pairStrings p : IncludePairs) {
-    std::filesystem::path FilePath = Dirpath / p.first;
-    std::filesystem::create_directories(FilePath.parent_path());
-    std::ofstream outfile(FilePath, std::ios::out | std::ios::trunc);
-    if (outfile.is_open()) {
-      outfile << p.second << std::endl;
-
-      outfile.close();
-    } else {
-      throw sycl::exception(sycl::errc::build,
-                            "Failed to open include file for write: " +
-                                FilePath.string());
-    }
-  }
-}
-
-std::string getCompilerName() {
-#ifdef _WIN32
-  std::string Compiler = "clang++.exe";
-#else
-  std::string Compiler = "clang++";
-#endif
-  return Compiler;
-}
-
-// We are assuming that the compiler is in /bin and the shared lib in
-// the adjacent /lib.
-std::filesystem::path getCompilerPath() {
-  std::string Compiler = getCompilerName();
-  const std::string LibSYCLDir = sycl::detail::OSUtil::getCurrentDSODir();
-  std::filesystem::path CompilerPath =
-      std::filesystem::path(LibSYCLDir) / ".." / "bin" / Compiler;
-  return CompilerPath;
-}
-
-int invokeCommand(const std::string &command, std::string &output) {
-#ifdef _WIN32
-  FILE *pipe = _popen(command.c_str(), "r");
-#else
-  FILE *pipe = popen(command.c_str(), "r");
-#endif
-  if (!pipe) {
-    return -1;
-  }
-
-  char buffer[1024];
-  while (!feof(pipe)) {
-    if (fgets(buffer, sizeof(buffer), pipe) != NULL) {
-      output += buffer;
-    }
-  }
-
-#ifdef _WIN32
-  _pclose(pipe);
-#else
-  pclose(pipe);
-#endif
-
-  return 0;
-}
-
-std::string invokeCompiler(const std::filesystem::path &FPath,
-                           const std::filesystem::path &DPath,
-                           const std::string &Id,
-                           const std::vector<std::string> &UserArgs,
-                           std::string *LogPtr) {
-
-  std::filesystem::path FilePath(FPath);
-  std::filesystem::path ParentDir(DPath);
-  std::filesystem::path TargetPath = ParentDir / (Id + ".bin");
-  std::filesystem::path LogPath = ParentDir / "compilation_log.txt";
-  std::string Compiler = getCompilerPath().make_preferred().string();
-
-  std::string Command =
-      Compiler + " -fsycl -o " + TargetPath.make_preferred().string() + " " +
-      userArgsAsString(UserArgs) +
-      " -fno-sycl-dead-args-optimization -fsycl-dump-device-code=" +
-      ParentDir.make_preferred().string() + " " +
-      FilePath.make_preferred().string() + " 2>&1";
-
-  std::string CompileLog;
-  int Result = invokeCommand(Command, CompileLog);
-
-  if (LogPtr != nullptr) {
-    LogPtr->append(CompileLog);
-  }
-
-  // There is little chance of Result being non-zero.
-  // Actual compilation failure is not detected by error code,
-  // but by missing .spv files.
-  if (Result != 0) {
-    throw sycl::exception(sycl::errc::build,
-                          "Compile failure: " + std::to_string(Result) + " " +
-                              CompileLog);
-  }
-  return CompileLog;
-}
-
-std::filesystem::path findSpv(const std::filesystem::path &ParentDir,
-                              const std::string &Id, std::string &CompileLog) {
-  std::regex PatternRegex(Id + R"(.*\.spv)");
-
-  // Iterate through all files in the directory matching the pattern.
-  for (const auto &Entry : std::filesystem::directory_iterator(ParentDir)) {
-    if (Entry.is_regular_file() &&
-        std::regex_match(Entry.path().filename().string(), PatternRegex)) {
-      return Entry.path(); // Return the path if it matches the SPV pattern.
-    }
-  }
-
-  // Missing .spv file indicates there was a compilation failure.
-  throw sycl::exception(sycl::errc::build, "Compile failure: " + CompileLog);
-}
-
-spirv_vec_t loadSpvFromFile(const std::filesystem::path &FileName) {
-  std::ifstream SpvStream(FileName, std::ios::binary);
-  SpvStream.seekg(0, std::ios::end);
-  size_t Size = SpvStream.tellg();
-  SpvStream.seekg(0);
-  spirv_vec_t Spv(Size);
-  SpvStream.read(reinterpret_cast<char *>(Spv.data()), Size);
-
-  return Spv;
-}
-
-spirv_vec_t
-SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs,
-              const std::vector<std::string> &UserArgs, std::string *LogPtr,
-              const std::vector<std::string> &RegisteredKernelNames) {
-  // clang-format off
-  const std::string id                   = generateSemiUniqueId();
-  const std::filesystem::path ParentDir  = prepareWS(id);
-  std::filesystem::path FilePath         = outputCpp(ParentDir, id, SYCLSource, UserArgs, RegisteredKernelNames);
-                                           outputIncludeFiles(ParentDir, IncludePairs);
-  std::string CompileLog                 = invokeCompiler(FilePath, ParentDir, id, UserArgs, LogPtr);
-  std::filesystem::path SpvPath          = findSpv(ParentDir, id, CompileLog);
-  spirv_vec_t Spv                        = loadSpvFromFile(SpvPath);
-                                           deleteWS(ParentDir);
-                                           return Spv;
-  // clang-format on
-}
-
-bool SYCL_Compilation_Available() {
-  // Is compiler on $PATH ? We try to invoke it.
-  std::string id = generateSemiUniqueId();
-  const std::filesystem::path tmp = std::filesystem::temp_directory_path();
-  std::filesystem::path DumpPath = tmp / (id + "_version.txt");
-  std::string Compiler = getCompilerPath().make_preferred().string();
-  std::string TestCommand =
-      Compiler + " --version > " + DumpPath.make_preferred().string();
-  int result = std::system(TestCommand.c_str());
-
-  return (result == 0);
-}
-
-} // namespace detail
-} // namespace ext::oneapi::experimental
-} // namespace _V1
-} // namespace sycl
-#endif
-
 #if SYCL_EXT_JIT_ENABLE
 #include "../jit_compiler.hpp"
 #endif
@@ -357,6 +41,14 @@ spirv_vec_t SYCL_JIT_to_SPIRV(
 #endif
 }
 
+std::string userArgsAsString(const std::vector<std::string> &UserArguments) {
+  return std::accumulate(UserArguments.begin(), UserArguments.end(),
+                         std::string(""),
+                         [](const std::string &A, const std::string &B) {
+                           return A.empty() ? B : A + " " + B;
+                         });
+}
+
 } // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1

sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp

@@ -24,13 +24,6 @@ namespace detail {
 using spirv_vec_t = std::vector<uint8_t>;
 using include_pairs_t = std::vector<std::pair<std::string, std::string>>;
 
-spirv_vec_t
-SYCL_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs,
-              const std::vector<std::string> &UserArgs, std::string *LogPtr,
-              const std::vector<std::string> &RegisteredKernelNames);
-
-bool SYCL_Compilation_Available();
-
 std::string userArgsAsString(const std::vector<std::string> &UserArguments);
 
 spirv_vec_t

sycl/source/kernel_bundle.cpp

@@ -379,8 +379,6 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) {
     } else if (Language == source_language::spirv) {
       return true;
     } else if (Language == source_language::sycl) {
-      return detail::SYCL_Compilation_Available();
-    } else if (Language == source_language::sycl_jit) {
       return detail::SYCL_JIT_Compilation_Available();
     }
   }