diff --git a/sycl-jit/common/include/Kernel.h b/sycl-jit/common/include/Kernel.h index 2959b2e1b9871..bbb9a802995ff 100644 --- a/sycl-jit/common/include/Kernel.h +++ b/sycl-jit/common/include/Kernel.h @@ -350,6 +350,11 @@ struct SYCLKernelInfo { : Name{KernelName}, Args{NumArgs}, Attributes{}, NDR{}, BinaryInfo{} {} }; +struct InMemoryFile { + const char *Path; + const char *Contents; +}; + } // namespace jit_compiler #endif // SYCL_FUSION_COMMON_KERNEL_H diff --git a/sycl-jit/jit-compiler/CMakeLists.txt b/sycl-jit/jit-compiler/CMakeLists.txt index f49833b5ac7f8..09af2de6853ae 100644 --- a/sycl-jit/jit-compiler/CMakeLists.txt +++ b/sycl-jit/jit-compiler/CMakeLists.txt @@ -7,6 +7,7 @@ add_llvm_library(sycl-jit lib/fusion/FusionHelper.cpp lib/fusion/JITContext.cpp lib/fusion/ModuleHelper.cpp + lib/rtc/DeviceCompilation.cpp lib/helper/ConfigHelper.cpp SHARED @@ -29,6 +30,14 @@ add_llvm_library(sycl-jit TargetParser MC ${LLVM_TARGETS_TO_BUILD} + + LINK_LIBS + clangBasic + clangDriver + clangFrontend + clangCodeGen + clangTooling + clangSerialization ) target_compile_options(sycl-jit PRIVATE ${SYCL_JIT_WARNING_FLAGS}) @@ -40,6 +49,8 @@ target_include_directories(sycl-jit SYSTEM PRIVATE ${LLVM_MAIN_INCLUDE_DIR} ${LLVM_SPIRV_INCLUDE_DIRS} + ${LLVM_EXTERNAL_CLANG_SOURCE_DIR}/include + ${CMAKE_BINARY_DIR}/tools/clang/include ) target_include_directories(sycl-jit PUBLIC diff --git a/sycl-jit/jit-compiler/include/KernelFusion.h b/sycl-jit/jit-compiler/include/KernelFusion.h index 51a66e684afbe..f149e05692627 100644 --- a/sycl-jit/jit-compiler/include/KernelFusion.h +++ b/sycl-jit/jit-compiler/include/KernelFusion.h @@ -66,6 +66,9 @@ JITResult materializeSpecConstants(const char *KernelName, jit_compiler::SYCLKernelBinaryInfo &BinInfo, View SpecConstBlob); +JITResult compileSYCL(InMemoryFile SourceFile, View IncludeFiles, + View UserArgs); + /// Clear all previously set options. void resetJITConfiguration(); diff --git a/sycl-jit/jit-compiler/ld-version-script.txt b/sycl-jit/jit-compiler/ld-version-script.txt index eb7892fdfec9c..c12256659ce30 100644 --- a/sycl-jit/jit-compiler/ld-version-script.txt +++ b/sycl-jit/jit-compiler/ld-version-script.txt @@ -3,6 +3,7 @@ /* Export the library entry points */ fuseKernels; materializeSpecConstants; + compileSYCL; resetJITConfiguration; addToJITConfiguration; diff --git a/sycl-jit/jit-compiler/lib/KernelFusion.cpp b/sycl-jit/jit-compiler/lib/KernelFusion.cpp index 0a03c90946cde..81037438061ae 100644 --- a/sycl-jit/jit-compiler/lib/KernelFusion.cpp +++ b/sycl-jit/jit-compiler/lib/KernelFusion.cpp @@ -14,6 +14,7 @@ #include "fusion/FusionPipeline.h" #include "helper/ConfigHelper.h" #include "helper/ErrorHandling.h" +#include "rtc/DeviceCompilation.h" #include "translation/KernelTranslation.h" #include "translation/SPIRVLLVMTranslation.h" #include @@ -235,6 +236,31 @@ extern "C" JITResult fuseKernels(View KernelInformation, return JITResult{FusedKernelInfo}; } +extern "C" JITResult compileSYCL(InMemoryFile SourceFile, + View IncludeFiles, + View UserArgs) { + auto ModuleOrErr = compileDeviceCode(SourceFile, IncludeFiles, UserArgs); + if (!ModuleOrErr) { + return errorToFusionResult(ModuleOrErr.takeError(), + "Device compilation failed"); + } + std::unique_ptr Module = std::move(*ModuleOrErr); + + SYCLKernelInfo Kernel; + auto Error = translation::KernelTranslator::translateKernel( + Kernel, *Module, JITContext::getInstance(), BinaryFormat::SPIRV); + + auto *LLVMCtx = &Module->getContext(); + Module.reset(); + delete LLVMCtx; + + if (Error) { + return errorToFusionResult(std::move(Error), "SPIR-V translation failed"); + } + + return JITResult{Kernel}; +} + extern "C" void resetJITConfiguration() { ConfigHelper::reset(); } extern "C" void addToJITConfiguration(OptionStorage &&Opt) { diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp new file mode 100644 index 0000000000000..55b8714ce5ac2 --- /dev/null +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -0,0 +1,147 @@ +//==---------------------- DeviceCompilation.cpp ---------------------------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#include "DeviceCompilation.h" + +#include +#include +#include +#include +#include +#include + +#ifdef _GNU_SOURCE +#include +static char X; // Dummy symbol, used as an anchor for `dlinfo` below. +#endif + +static constexpr auto InvalidDPCPPRoot = ""; +static constexpr auto JITLibraryPathSuffix = "/lib/libsycl-jit.so"; + +static const std::string &getDPCPPRoot() { + thread_local std::string DPCPPRoot; + + if (!DPCPPRoot.empty()) { + return DPCPPRoot; + } + DPCPPRoot = InvalidDPCPPRoot; + +#ifdef _GNU_SOURCE + Dl_info Info; + if (dladdr(&X, &Info)) { + std::string LoadedLibraryPath = Info.dli_fname; + auto Pos = LoadedLibraryPath.rfind(JITLibraryPathSuffix); + if (Pos != std::string::npos) { + DPCPPRoot = LoadedLibraryPath.substr(0, Pos); + } + } +#endif // _GNU_SOURCE + + // TODO: Implemenent other means of determining the DPCPP root, e.g. + // evaluating the `CMPLR_ROOT` env. + + return DPCPPRoot; +} + +namespace { +using namespace clang; +using namespace clang::tooling; +using namespace clang::driver; + +struct GetLLVMModuleAction : public ToolAction { + // Code adapted from `FrontendActionFactory::runInvocation`. + bool runInvocation(std::shared_ptr Invocation, + FileManager *Files, + std::shared_ptr PCHContainerOps, + DiagnosticConsumer *DiagConsumer) override { + assert(!Module && "Action should only be invoked on a single file"); + + // Create a compiler instance to handle the actual work. + CompilerInstance Compiler(std::move(PCHContainerOps)); + Compiler.setInvocation(std::move(Invocation)); + Compiler.setFileManager(Files); + + // Create the compiler's actual diagnostics engine. + Compiler.createDiagnostics(DiagConsumer, /*ShouldOwnClient=*/false); + if (!Compiler.hasDiagnostics()) { + return false; + } + + Compiler.createSourceManager(*Files); + + // Ignore `Compiler.getFrontendOpts().ProgramAction` (would be `EmitBC`) and + // create/execute an `EmitLLVMOnlyAction` (= codegen to LLVM module without + // emitting anything) instead. + EmitLLVMOnlyAction ELOA; + const bool Success = Compiler.ExecuteAction(ELOA); + Files->clearStatCache(); + if (!Success) { + return false; + } + + // Take the module and its context to extend the objects' lifetime. + Module = ELOA.takeModule(); + ELOA.takeLLVMContext(); + + return true; + } + + std::unique_ptr Module; +}; + +} // anonymous namespace + +llvm::Expected> +jit_compiler::compileDeviceCode(InMemoryFile SourceFile, + View IncludeFiles, + View UserArgs) { + const std::string &DPCPPRoot = getDPCPPRoot(); + if (DPCPPRoot == InvalidDPCPPRoot) { + return llvm::createStringError("Could not locate DPCPP root directory"); + } + + SmallVector CommandLine = {"-fsycl-device-only"}; + // TODO: Allow instrumentation again when device library linking is + // implemented. + CommandLine.push_back("-fno-sycl-instrument-device-code"); + CommandLine.append(UserArgs.begin(), UserArgs.end()); + clang::tooling::FixedCompilationDatabase DB{".", CommandLine}; + + clang::tooling::ClangTool Tool{DB, {SourceFile.Path}}; + + // Set up in-memory filesystem. + Tool.mapVirtualFile(SourceFile.Path, SourceFile.Contents); + for (const auto &IF : IncludeFiles) { + Tool.mapVirtualFile(IF.Path, IF.Contents); + } + + // Reset argument adjusters to drop the `-fsyntax-only` flag which is added by + // default by this API. + Tool.clearArgumentsAdjusters(); + // Then, modify argv[0] and set the resource directory so that the driver + // picks up the correct SYCL environment. + Tool.appendArgumentsAdjuster( + [&DPCPPRoot](const CommandLineArguments &Args, + StringRef Filename) -> CommandLineArguments { + (void)Filename; + CommandLineArguments NewArgs = Args; + NewArgs[0] = (Twine(DPCPPRoot) + "/bin/clang++").str(); + NewArgs.push_back((Twine("-resource-dir=") + DPCPPRoot + "/lib/clang/" + + Twine(CLANG_VERSION_MAJOR)) + .str()); + return NewArgs; + }); + + GetLLVMModuleAction Action; + if (!Tool.run(&Action)) { + return std::move(Action.Module); + } + + // TODO: Capture compiler errors from the ClangTool. + return llvm::createStringError("Unable to obtain LLVM module"); +} diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h new file mode 100644 index 0000000000000..2373457c5847f --- /dev/null +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.h @@ -0,0 +1,28 @@ +//==---- DeviceCompilation.h - Compile SYCL device code with libtooling ----==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#ifndef SYCL_JIT_COMPILER_RTC_DEVICE_COMPILATION_H +#define SYCL_JIT_COMPILER_RTC_DEVICE_COMPILATION_H + +#include "Kernel.h" +#include "View.h" + +#include +#include + +#include + +namespace jit_compiler { + +llvm::Expected> +compileDeviceCode(InMemoryFile SourceFile, View IncludeFiles, + View UserArgs); + +} // namespace jit_compiler + +#endif // SYCL_JIT_COMPILER_RTC_DEVICE_COMPILATION_H diff --git a/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp b/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp index 4092f9dd96fc8..16e7af8d8b08c 100644 --- a/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp +++ b/sycl-jit/jit-compiler/lib/translation/SPIRVLLVMTranslation.cpp @@ -41,6 +41,9 @@ SPIRV::TranslatorOpts &SPIRVLLVMTranslator::translatorOpts() { // there's currently no obvious way to iterate the // array of extensions in KernelInfo. TransOpt.enableAllExtensions(); + // TODO: Remove this workaround. + TransOpt.setAllowedToUseExtension( + SPIRV::ExtensionID::SPV_KHR_untyped_pointers, false); TransOpt.setDesiredBIsRepresentation( SPIRV::BIsRepresentation::SPIRVFriendlyIR); // TODO: We need to take care of specialization constants, either by diff --git a/sycl/include/sycl/kernel_bundle_enums.hpp b/sycl/include/sycl/kernel_bundle_enums.hpp index fd53f8cd3a740..0fbccd917f27f 100644 --- a/sycl/include/sycl/kernel_bundle_enums.hpp +++ b/sycl/include/sycl/kernel_bundle_enums.hpp @@ -20,7 +20,13 @@ enum class bundle_state : char { namespace ext::oneapi::experimental { -enum class source_language : int { opencl = 0, spirv = 1, sycl = 2 /* cuda */ }; +enum class source_language : int { + opencl = 0, + spirv = 1, + sycl = 2, + /* cuda */ + sycl_jit = 99 /* temporary, alternative implementation for SYCL */ +}; // opencl versions struct cl_version { diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index f4332143aa617..c0c22954822b7 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -74,6 +74,14 @@ jit_compiler::jit_compiler() { return false; } + this->CompileSYCLHandle = reinterpret_cast( + sycl::detail::ur::getOsLibraryFuncAddress(LibraryPtr, "compileSYCL")); + if (!this->CompileSYCLHandle) { + printPerformanceWarning( + "Cannot resolve JIT library function entry point"); + return false; + } + return true; }; Available = checkJITLibrary(); @@ -1145,6 +1153,52 @@ std::vector jit_compiler::encodeReqdWorkGroupSize( return Encoded; } +std::vector jit_compiler::compileSYCL( + const std::string &Id, const std::string &SYCLSource, + const std::vector> &IncludePairs, + 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"); + } + + std::string SYCLFileName = Id + ".cpp"; + ::jit_compiler::InMemoryFile SourceFile{SYCLFileName.c_str(), + SYCLSource.c_str()}; + + std::vector<::jit_compiler::InMemoryFile> IncludeFilesView; + IncludeFilesView.reserve(IncludePairs.size()); + std::transform(IncludePairs.begin(), IncludePairs.end(), + std::back_inserter(IncludeFilesView), [](const auto &Pair) { + return ::jit_compiler::InMemoryFile{Pair.first.c_str(), + Pair.second.c_str()}; + }); + std::vector UserArgsView; + UserArgsView.reserve(UserArgs.size()); + std::transform(UserArgs.begin(), UserArgs.end(), + std::back_inserter(UserArgsView), + [](const auto &Arg) { return Arg.c_str(); }); + + auto Result = CompileSYCLHandle(SourceFile, IncludeFilesView, UserArgsView); + + if (Result.failed()) { + throw sycl::exception(sycl::errc::build, Result.getErrorMessage()); + } + + // TODO: We currently don't have a meaningful build log. + (void)LogPtr; + + const auto &BI = Result.getKernelInfo().BinaryInfo; + assert(BI.Format == ::jit_compiler::BinaryFormat::SPIRV); + std::vector SPV(BI.BinaryStart, BI.BinaryStart + BI.BinarySize); + return SPV; +} + } // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/source/detail/jit_compiler.hpp b/sycl/source/detail/jit_compiler.hpp index b2bdc091dfbd9..1908defa42e77 100644 --- a/sycl/source/detail/jit_compiler.hpp +++ b/sycl/source/detail/jit_compiler.hpp @@ -44,6 +44,12 @@ class jit_compiler { const std::string &KernelName, const std::vector &SpecConstBlob); + std::vector compileSYCL( + const std::string &Id, const std::string &SYCLSource, + const std::vector> &IncludePairs, + const std::vector &UserArgs, std::string *LogPtr, + const std::vector &RegisteredKernelNames); + bool isAvailable() { return Available; } static jit_compiler &get_instance() { @@ -80,10 +86,12 @@ class jit_compiler { using FuseKernelsFuncT = decltype(::jit_compiler::fuseKernels) *; using MaterializeSpecConstFuncT = decltype(::jit_compiler::materializeSpecConstants) *; + using CompileSYCLFuncT = decltype(::jit_compiler::compileSYCL) *; using ResetConfigFuncT = decltype(::jit_compiler::resetJITConfiguration) *; using AddToConfigFuncT = decltype(::jit_compiler::addToJITConfiguration) *; FuseKernelsFuncT FuseKernelsHandle = nullptr; MaterializeSpecConstFuncT MaterializeSpecConstHandle = nullptr; + CompileSYCLFuncT CompileSYCLHandle = nullptr; ResetConfigFuncT ResetConfigHandle = nullptr; AddToConfigFuncT AddToConfigHandle = nullptr; #endif // SYCL_EXT_JIT_ENABLE diff --git a/sycl/source/detail/kernel_bundle_impl.hpp b/sycl/source/detail/kernel_bundle_impl.hpp index 38b6bb1deb920..d6930633c2cd1 100644 --- a/sycl/source/detail/kernel_bundle_impl.hpp +++ b/sycl/source/detail/kernel_bundle_impl.hpp @@ -446,6 +446,12 @@ class kernel_bundle_impl { 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, + BuildOptions, LogPtr, + RegisteredKernelNames); + } throw sycl::exception( make_error_code(errc::invalid), "OpenCL C and SPIR-V are the only supported languages at this time"); @@ -502,7 +508,8 @@ 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) + if (Lang != syclex::source_language::sycl && + Lang != syclex::source_language::sycl_jit) return Name; bool isMangled = Name.find("__sycl_kernel_") != std::string::npos; diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp index 1e3106255b9c5..195ed3f1987d8 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp @@ -315,3 +315,40 @@ bool SYCL_Compilation_Available() { } // namespace _V1 } // namespace sycl #endif + +#if SYCL_EXT_JIT_ENABLE +#include "../jit_compiler.hpp" +#endif + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { +namespace detail { + +bool SYCL_JIT_Compilation_Available() { +#if SYCL_EXT_JIT_ENABLE + return sycl::detail::jit_compiler::get_instance().isAvailable(); +#else + return false; +#endif +} + +spirv_vec_t SYCL_JIT_to_SPIRV( + [[maybe_unused]] const std::string &SYCLSource, + [[maybe_unused]] include_pairs_t IncludePairs, + [[maybe_unused]] const std::vector &UserArgs, + [[maybe_unused]] std::string *LogPtr, + [[maybe_unused]] const std::vector &RegisteredKernelNames) { +#if SYCL_EXT_JIT_ENABLE + return sycl::detail::jit_compiler::get_instance().compileSYCL( + "rtc", SYCLSource, IncludePairs, UserArgs, LogPtr, RegisteredKernelNames); +#else + throw sycl::exception(sycl::errc::build, + "kernel_compiler via sycl-jit is not available"); +#endif +} + +} // namespace detail +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl diff --git a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp index dfff9ac839e84..2d591cfb0913a 100644 --- a/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp +++ b/sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp @@ -31,6 +31,13 @@ SYCL_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs, bool SYCL_Compilation_Available(); +spirv_vec_t +SYCL_JIT_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs, + const std::vector &UserArgs, std::string *LogPtr, + const std::vector &RegisteredKernelNames); + +bool SYCL_JIT_Compilation_Available(); + } // namespace detail } // namespace ext::oneapi::experimental diff --git a/sycl/source/kernel_bundle.cpp b/sycl/source/kernel_bundle.cpp index 12ca87b6604f5..5e81084c9aaaa 100644 --- a/sycl/source/kernel_bundle.cpp +++ b/sycl/source/kernel_bundle.cpp @@ -380,6 +380,8 @@ bool is_source_kernel_bundle_supported(backend BE, source_language Language) { 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_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp new file mode 100644 index 0000000000000..1491483781834 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -0,0 +1,129 @@ +//==- 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 + +#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); +} + +void 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; + } + + // 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"}; + 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. +} + +int main() { + +#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER + test_build_and_run(); +#else + static_assert(false, "Kernel Compiler feature test macro undefined"); +#endif + return 0; +} 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 652169c1082f1..3aaa3e6cb8bf4 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: 4 +// CHECK-NUM-MATCHES: 5 // // This test verifies that `` isn't used in E2E tests. Instead, // fine-grained includes should used, see