diff --git a/sycl/include/sycl/kernel_bundle_enums.hpp b/sycl/include/sycl/kernel_bundle_enums.hpp index 0fbccd917f27..67c7288253c0 100644 --- a/sycl/include/sycl/kernel_bundle_enums.hpp +++ b/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 diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 19f1915943f0..bd5922423f65 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -19,6 +19,14 @@ #include #include +#include +#include +#ifdef _WIN32 +#include +#else +#include // pipe, dup2, read, close +#endif + namespace sycl { inline namespace _V1 { namespace detail { @@ -1173,17 +1181,21 @@ std::vector jit_compiler::compileSYCL( const std::vector &UserArgs, std::string *LogPtr, const std::vector &RegisteredKernelNames) { - // TODO: Handle template instantiation. - if (!RegisteredKernelNames.empty()) { - throw sycl::exception( - sycl::errc::build, - "Property `sycl::ext::oneapi::experimental::registered_kernel_names` " - "is not yet supported for the `sycl_jit` source language"); + // RegisteredKernelNames may contain template specialization that + // we want to make sure are instantiated. So we just put them in main() + // which ensures they are instantiated. + std::ostringstream ss; + ss << "int main() {\n"; + for (const std::string &KernelName : RegisteredKernelNames) { + ss << " (void)" << KernelName << ";\n"; } + ss << " return 0;\n}\n" << std::endl; + + std::string FinalSource = SYCLSource + ss.str(); std::string SYCLFileName = Id + ".cpp"; ::jit_compiler::InMemoryFile SourceFile{SYCLFileName.c_str(), - SYCLSource.c_str()}; + FinalSource.c_str()}; std::vector<::jit_compiler::InMemoryFile> IncludeFilesView; IncludeFilesView.reserve(IncludePairs.size()); @@ -1198,14 +1210,81 @@ std::vector jit_compiler::compileSYCL( std::back_inserter(UserArgsView), [](const auto &Arg) { return Arg.c_str(); }); + // Redirect stderr to a string stream. +#ifdef _WIN32 + HANDLE read_pipe, write_pipe; + SECURITY_ATTRIBUTES sa = {sizeof(SECURITY_ATTRIBUTES), NULL, TRUE}; + if (!CreatePipe(&read_pipe, &write_pipe, &sa, 0)) { + throw sycl::exception(sycl::errc::build, "Failed to create pipe"); + } + + HANDLE saved_stderr = GetStdHandle(STD_ERROR_HANDLE); + HANDLE saved_stdout = GetStdHandle(STD_OUTPUT_HANDLE); + if (!SetStdHandle(STD_ERROR_HANDLE, write_pipe) || + !SetStdHandle(STD_OUTPUT_HANDLE, write_pipe)) { + throw sycl::exception(sycl::errc::build, + "Failed to redirect stderr/stdout"); + } +#else + int pipefd[2]; + if (pipe(pipefd) == -1) { + throw sycl::exception(sycl::errc::build, "Failed to create pipe"); + } + + int saved_stderr = dup(fileno(stderr)); + int saved_stdout = dup(fileno(stdout)); + if (dup2(pipefd[1], fileno(stderr)) == -1 || + dup2(pipefd[1], fileno(stdout)) == -1) { + throw sycl::exception(sycl::errc::build, + "Failed to redirect stderr/stdout"); + } + close(pipefd[1]); +#endif + + std::stringstream error_stream; + + // Compile it! auto Result = CompileSYCLHandle(SourceFile, IncludeFilesView, UserArgsView); - if (Result.failed()) { - throw sycl::exception(sycl::errc::build, Result.getErrorMessage()); + // Restore stderr/stdout. +#ifdef _WIN32 + SetStdHandle(STD_ERROR_HANDLE, saved_stderr); + SetStdHandle(STD_OUTPUT_HANDLE, saved_stdout); + CloseHandle(write_pipe); + + // Read from the pipe + char buffer[1024]; + DWORD count; + while (ReadFile(read_pipe, buffer, sizeof(buffer) - 1, &count, NULL) && + count > 0) { + buffer[count] = '\0'; + error_stream << buffer; + } + CloseHandle(read_pipe); +#else + dup2(saved_stderr, fileno(stderr)); + dup2(saved_stdout, fileno(stdout)); + close(saved_stderr); + close(saved_stdout); + + // Read from the pipe + char buffer[1024]; + ssize_t count; + while ((count = read(pipefd[0], buffer, sizeof(buffer) - 1)) > 0) { + buffer[count] = '\0'; + error_stream << buffer; + } + close(pipefd[0]); +#endif + + if (LogPtr != nullptr) { + LogPtr->append(error_stream.str()); } - // TODO: We currently don't have a meaningful build log. - (void)LogPtr; + if (Result.failed()) { + throw sycl::exception(sycl::errc::build, + Result.getErrorMessage() + error_stream.str()); + } const auto &BI = Result.getKernelInfo().BinaryInfo; assert(BI.Format == ::jit_compiler::BinaryFormat::SPIRV); diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 7ce597171129..711d8e2f80c8 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -494,13 +494,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(this->Source); - return syclex::detail::SYCL_JIT_to_SPIRV(SourceStr, IncludePairs, + return syclex::detail::SYCL_JIT_to_SPIRV(*SourceStrPtr, IncludePairs, BuildOptions, LogPtr, RegisteredKernelNames); } @@ -571,8 +565,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; @@ -595,6 +588,7 @@ class kernel_bundle_impl { "kernel_bundle."); std::string AdjustedName = adjust_kernel_name(Name, Language); + if (!ext_oneapi_has_kernel(Name)) throw sycl::exception(make_error_code(errc::invalid), "kernel '" + AdjustedName + diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 89a8a548a6d8..41c27238ac5b 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -9,322 +9,6 @@ #include "kernel_compiler_sycl.hpp" #include // 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 &UserArgs, std::string *LogPtr, - const std::vector &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 &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 - -#include -#include -#include -#include -#include -#include -#include - -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( - 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 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 &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 &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 &UserArgs, - const std::vector &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; - 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 &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(Spv.data()), Size); - - return Spv; -} - -spirv_vec_t -SYCL_to_SPIRV(const std::string &SYCLSource, include_pairs_t IncludePairs, - const std::vector &UserArgs, std::string *LogPtr, - const std::vector &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 &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 diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp index 948f199c1b3c..73f4f1f058f0 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp @@ -24,13 +24,6 @@ namespace detail { using spirv_vec_t = std::vector; using include_pairs_t = std::vector>; -spirv_vec_t -SYCL_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs, - const std::vector &UserArgs, std::string *LogPtr, - const std::vector &RegisteredKernelNames); - -bool SYCL_Compilation_Available(); - std::string userArgsAsString(const std::vector &UserArguments); spirv_vec_t diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index a5bf9245c6d8..453b776b35b4 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -350,8 +350,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(); } } diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp index ffce6e2ee298..4dd46630d0be 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp @@ -12,14 +12,8 @@ // -- Test the kernel_compiler with SYCL source. // RUN: %{build} -o %t.out -// If clang++ is not on the PATH, or if sycl was compiled with GCC < 8, then -// the kernel_compiler is not available for SYCL language. -// Note: this 'invoking clang++' version for SYCL language support is temporary, -// and will be replaced by the SYCL_JIT version soon. -// DEFINE: %{available} = %t.out available - -// RUN: %if available %{ %{run} %t.out %} -// RUN: %if available %{ %{l0_leak_check} %{run} %t.out %} +// RUN: %{run} %t.out +// RUN: %{l0_leak_check} %{run} %t.out // -- Test again, with caching. // 'reading-from-cache' is just a string we pass to differentiate between the @@ -27,13 +21,13 @@ // DEFINE: %{cache_vars} = %{l0_leak_check} env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=5 SYCL_CACHE_DIR=%t/cache_dir // RUN: rm -rf %t/cache_dir -// RUN: %if available %{ %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE %} -// RUN: %if available %{ %{cache_vars} %t.out reading-from-cache 2>&1 | FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE %} +// RUN: %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE +// RUN: %{cache_vars} %t.out reading-from-cache 2>&1 | FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE // -- Add leak check. // RUN: rm -rf %t/cache_dir -// RUN: %if available %{ %{l0_leak_check} %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE %} -// RUN: %if available %{ %{l0_leak_check} %{cache_vars} %t.out reading-from-cache 2>&1 | FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE %} +// RUN: %{l0_leak_check} %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE +// RUN: %{l0_leak_check} %{cache_vars} %t.out reading-from-cache 2>&1 | FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE // CHECK-WRITTEN-TO-CACHE: [Persistent Cache]: enabled // CHECK-WRITTEN-TO-CACHE-NOT: [kernel_compiler Persistent Cache]: using cached binary @@ -161,6 +155,7 @@ void test_build_and_run(bool readingFromCache) { "kernel bundle extension: " << q.get_device().get_info() << std::endl; + assert(ok); return; } @@ -191,10 +186,11 @@ void test_build_and_run(bool readingFromCache) { // If the kernel was restored from cache, there will not have been // any warning issued by the compilation of the kernel. - if (!readingFromCache) { - assert(log.find("warning: 'this_nd_item<1>' is deprecated") != - std::string::npos); - } + // TODO: get logging working, so this assert can be tested. + // if (!readingFromCache) { + // assert(log.find("warning: 'this_nd_item<1>' is deprecated") != + // std::string::npos); + // } // clang-format off @@ -244,6 +240,41 @@ void test_error() { } } +void test_unsupported_options() { + namespace syclex = sycl::ext::oneapi::experimental; + using source_kb = sycl::kernel_bundle; + + sycl::queue q; + sycl::context ctx = q.get_context(); + + bool ok = + q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl); + if (!ok) { + return; + } + + source_kb kbSrc = syclex::create_kernel_bundle_from_source( + ctx, syclex::source_language::sycl, ""); + std::vector devs = kbSrc.get_devices(); + + auto CheckUnsupported = [&](const std::vector &flags) { + try { + syclex::build(kbSrc, devs, + syclex::properties{syclex::build_options{flags}}); + assert(false && "unsupported option not detected"); + } catch (sycl::exception &e) { + assert(e.code() == sycl::errc::build); + assert(std::string(e.what()).find("Parsing of user arguments failed") != + std::string::npos); + } + }; + + CheckUnsupported({"-fsanitize=address"}); + CheckUnsupported({"-Xsycl-target-frontend", "-fsanitize=address"}); + CheckUnsupported({"-Xsycl-target-frontend=spir64", "-fsanitize=address"}); + CheckUnsupported({"-Xarch_device", "-fsanitize=address"}); +} + void test_esimd() { namespace syclex = sycl::ext::oneapi::experimental; using source_kb = sycl::kernel_bundle; @@ -316,17 +347,16 @@ int main(int argc, char *argv[]) { std::string argument(argv[1]); if (argument == "reading-from-cache") { readingFromCache = true; - } else if (argument == "available") { - sycl::device d; - bool avail = d.ext_oneapi_can_compile(syclex::source_language::sycl); - return avail; } } #ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER test_build_and_run(readingFromCache); test_error(); - test_esimd(); + test_unsupported_options(); + + // TODO: jit_compiler is not supporting ESIMD. + // test_esimd(); #else static_assert(false, "Kernel Compiler feature test macro undefined"); #endif diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp deleted file mode 100644 index 01f25f813b82..000000000000 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ /dev/null @@ -1,194 +0,0 @@ -//==- kernel_compiler_sycl_jit.cpp --- kernel_compiler extension tests -----==// -// -// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. -// See https://llvm.org/LICENSE.txt for license information. -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -// -//===----------------------------------------------------------------------===// - -// REQUIRES: (opencl || level_zero) -// UNSUPPORTED: accelerator - -// RUN: %{build} -o %t.out -// RUN: %{run} %t.out -// RUN: %{l0_leak_check} %{run} %t.out - -// -- Test again, with caching. - -// DEFINE: %{cache_vars} = %{l0_leak_check} env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_TRACE=5 SYCL_CACHE_DIR=%t/cache_dir -// RUN: rm -rf %t/cache_dir -// RUN: %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE -// RUN: %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE - -// -- Add leak check. -// RUN: rm -rf %t/cache_dir -// RUN: %{l0_leak_check} %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-WRITTEN-TO-CACHE -// RUN: %{l0_leak_check} %{cache_vars} %t.out 2>&1 | FileCheck %s --check-prefixes=CHECK-READ-FROM-CACHE - -// CHECK-WRITTEN-TO-CACHE: [Persistent Cache]: enabled -// CHECK-WRITTEN-TO-CACHE-NOT: [kernel_compiler Persistent Cache]: using cached binary -// CHECK-WRITTEN-TO-CACHE: [kernel_compiler Persistent Cache]: binary has been cached - -// CHECK-READ-FROM-CACHE: [Persistent Cache]: enabled -// CHECK-READ-FROM-CACHE-NOT: [kernel_compiler Persistent Cache]: binary has been cached -// CHECK-READ-FROM-CACHE: [kernel_compiler Persistent Cache]: using cached binary - -#include -#include - -auto constexpr AddEmH = R"===( - int AddEm(int a, int b){ - return a + b + 5; - } -)==="; - -auto constexpr PlusEmH = R"===( - int PlusEm(int a, int b){ - return a + b + 6; - } -)==="; - -// TODO: remove SYCL_EXTERNAL once it is no longer needed. -auto constexpr SYCLSource = R"===( -#include -#include "intermediate/AddEm.h" -#include "intermediate/PlusEm.h" - -// use extern "C" to avoid name mangling -extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) -void ff_cp(int *ptr) { - - // intentionally using deprecated routine, as opposed to this_work_item::get_nd_item<1>() - sycl::nd_item<1> Item = sycl::ext::oneapi::experimental::this_nd_item<1>(); - - sycl::id<1> GId = Item.get_global_id(); - ptr[GId.get(0)] = AddEm(GId.get(0), 37); -} -)==="; - -void test_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) { - constexpr int Range = 10; - int *usmPtr = sycl::malloc_shared(Range, Queue); - int start = 3; - - sycl::nd_range<1> R1{{Range}, {1}}; - - bool Passa = true; - - memset(usmPtr, 0, Range * sizeof(int)); - Queue.submit([&](sycl::handler &Handler) { - Handler.set_arg(0, usmPtr); - Handler.parallel_for(R1, Kernel); - }); - Queue.wait(); - - for (int i = 0; i < Range; i++) { - std::cout << usmPtr[i] << "=" << (i + seed) << " "; - assert(usmPtr[i] == i + seed); - } - std::cout << std::endl; - - sycl::free(usmPtr, Queue); -} - -int test_build_and_run() { - namespace syclex = sycl::ext::oneapi::experimental; - using source_kb = sycl::kernel_bundle; - using exe_kb = sycl::kernel_bundle; - - sycl::queue q; - sycl::context ctx = q.get_context(); - - bool ok = - q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl_jit); - if (!ok) { - std::cout << "Apparently this device does not support `sycl_jit` source " - "kernel bundle extension: " - << q.get_device().get_info() - << std::endl; - return -1; - } - - // Create from source. - syclex::include_files incFiles{"intermediate/AddEm.h", AddEmH}; - incFiles.add("intermediate/PlusEm.h", PlusEmH); - source_kb kbSrc = syclex::create_kernel_bundle_from_source( - ctx, syclex::source_language::sycl_jit, SYCLSource, - syclex::properties{incFiles}); - - // Double check kernel_bundle.get_source() / get_backend(). - sycl::context ctxRes = kbSrc.get_context(); - assert(ctxRes == ctx); - sycl::backend beRes = kbSrc.get_backend(); - assert(beRes == ctx.get_backend()); - - // Compilation of empty prop list, no devices. - exe_kb kbExe1 = syclex::build(kbSrc); - - // // Compilation with props and devices - std::string log; - std::vector flags{"-g", "-fno-fast-math", - "-fsycl-instrument-device-code"}; - std::vector devs = kbSrc.get_devices(); - exe_kb kbExe2 = syclex::build( - kbSrc, devs, syclex::properties{syclex::build_options{flags}}); - - // extern "C" was used, so the name "ff_cp" is not mangled. - sycl::kernel k = kbExe2.ext_oneapi_get_kernel("ff_cp"); - - // Test the kernels. - test_1(q, k, 37 + 5); // ff_cp seeds 37. AddEm will add 5 more. - - return 0; -} - -int test_unsupported_options() { - namespace syclex = sycl::ext::oneapi::experimental; - using source_kb = sycl::kernel_bundle; - - sycl::queue q; - sycl::context ctx = q.get_context(); - - bool ok = - q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl_jit); - if (!ok) { - std::cout << "Apparently this device does not support `sycl_jit` source " - "kernel bundle extension: " - << q.get_device().get_info() - << std::endl; - return -1; - } - - source_kb kbSrc = syclex::create_kernel_bundle_from_source( - ctx, syclex::source_language::sycl_jit, ""); - std::vector devs = kbSrc.get_devices(); - - auto CheckUnsupported = [&](const std::vector &flags) { - try { - syclex::build(kbSrc, devs, - syclex::properties{syclex::build_options{flags}}); - assert(false && "unsupported option not detected"); - } catch (sycl::exception &e) { - assert(e.code() == sycl::errc::build); - assert(std::string(e.what()).find("Parsing of user arguments failed") != - std::string::npos); - } - }; - - CheckUnsupported({"-fsanitize=address"}); - CheckUnsupported({"-Xsycl-target-frontend", "-fsanitize=address"}); - CheckUnsupported({"-Xsycl-target-frontend=spir64", "-fsanitize=address"}); - CheckUnsupported({"-Xarch_device", "-fsanitize=address"}); - - return 0; -} - -int main() { - -#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER - return test_build_and_run() || test_unsupported_options(); -#else - static_assert(false, "Kernel Compiler feature test macro undefined"); -#endif - return 0; -} diff --git a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp index ca458439cc4d..a08c92265b15 100644 --- a/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp +++ b/sycl/test/e2e_test_requirements/no-unsupported-without-info.cpp @@ -54,7 +54,7 @@ // tests to match the required format and in that case you should just update // (i.e. reduce) the number and the list below. // -// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 423 +// NUMBER-OF-UNSUPPORTED-WITHOUT-INFO: 422 // // List of improperly UNSUPPORTED tests. // Remove the CHECK once the test has been properly UNSUPPORTED. @@ -231,7 +231,6 @@ // CHECK-NEXT: KernelAndProgram/undefined-symbol.cpp // CHECK-NEXT: KernelCompiler/kernel_compiler_opencl.cpp // CHECK-NEXT: KernelCompiler/kernel_compiler_sycl.cpp -// CHECK-NEXT: KernelCompiler/kernel_compiler_sycl_jit.cpp // CHECK-NEXT: KernelCompiler/multi_device.cpp // CHECK-NEXT: KernelCompiler/sycl_device_flags.cpp // CHECK-NEXT: LLVMIntrinsicLowering/bitreverse.cpp diff --git a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp index 3aaa3e6cb8bf..652169c1082f 100644 --- a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp +++ b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp @@ -6,7 +6,7 @@ // CHECK-DAG: README.md // CHECK-DAG: lit.cfg.py // -// CHECK-NUM-MATCHES: 5 +// CHECK-NUM-MATCHES: 4 // // This test verifies that `` isn't used in E2E tests. Instead, // fine-grained includes should used, see