diff --git a/sycl/include/CL/sycl/INTEL/online_compiler.hpp b/sycl/include/CL/sycl/INTEL/online_compiler.hpp index a2e8c6ba5b41a..f8bb63bc79dc5 100644 --- a/sycl/include/CL/sycl/INTEL/online_compiler.hpp +++ b/sycl/include/CL/sycl/INTEL/online_compiler.hpp @@ -8,11 +8,11 @@ #pragma once -#include #include // for __SYCL_INLINE_NAMESPACE +#include // for __SYCL_EXPORT #include -#include +#include #include __SYCL_INLINE_NAMESPACE(cl) { @@ -59,7 +59,9 @@ class device_arch { /// Represents an error happend during online compilation. class online_compile_error : public sycl::exception { - // TBD +public: + online_compile_error() = default; + online_compile_error(const string_class &Msg) : sycl::exception(Msg) {} }; /// Designates a source language for the online compiler. @@ -70,25 +72,32 @@ enum class source_language { opencl_c, cm }; template class online_compiler { public: /// Constructs online compiler which can target any device and produces - /// given compiled code format. Produces device code is 64-bit. + /// given compiled code format. Produces 64-bit device code. /// The created compiler is "optimistic" - it assumes all applicable SYCL /// device capabilities are supported by the target device(s). online_compiler(compiled_code_format fmt = compiled_code_format::spir_v) : OutputFormat(fmt), OutputFormatVersion({0, 0}), - DeviceArch(device_arch::any), Is64Bit(true), DeviceStepping("") {} + DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), + Is64Bit(true), DeviceStepping("") {} /// Constructs online compiler which targets given architecture and produces - /// given compiled code format. Produces device code is 64-bit. + /// given compiled code format. Produces 64-bit device code. /// Throws online_compile_error if values of constructor arguments are /// contradictory or not supported - e.g. if the source language is not /// supported for given device type. online_compiler(sycl::info::device_type dev_type, device_arch arch, compiled_code_format fmt = compiled_code_format::spir_v) - : OutputFormat(fmt), OutputFormatVersion({0, 0}), DeviceArch(arch), - Is64Bit(true), DeviceStepping("") {} + : OutputFormat(fmt), OutputFormatVersion({0, 0}), DeviceType(dev_type), + DeviceArch(arch), Is64Bit(true), DeviceStepping("") {} /// Constructs online compiler for the target specified by given SYCL device. - online_compiler(const sycl::device &dev); + // TODO: the initial version generates the generic code (SKL now), need + // to do additional device::info calls to determine the device by it's + // features. + online_compiler(const sycl::device &dev) + : OutputFormat(compiled_code_format::spir_v), OutputFormatVersion({0, 0}), + DeviceType(sycl::info::device_type::all), DeviceArch(device_arch::any), + Is64Bit(true), DeviceStepping("") {} /// Compiles given in-memory \c Lang source to a binary blob. Blob format, /// other parameters are set in the constructor by the compilation target @@ -100,31 +109,50 @@ template class online_compiler { std::vector compile(const std::string &src, const Tys &... args); /// Sets the compiled code format of the compilation target and returns *this. - online_compiler &setOutputFormat(compiled_code_format fmt); + online_compiler &setOutputFormat(compiled_code_format fmt) { + OutputFormat = fmt; + return *this; + } /// Sets the compiled code format version of the compilation target and /// returns *this. - online_compiler &setOutputFormatVersion(int major, int minor); + online_compiler &setOutputFormatVersion(int major, int minor) { + OutputFormatVersion = {major, minor}; + return *this; + } /// Sets the device type of the compilation target and returns *this. - online_compiler &setTargetDeviceType(sycl::info::device_type type); + online_compiler &setTargetDeviceType(sycl::info::device_type type) { + DeviceType = type; + return *this; + } /// Sets the device architecture of the compilation target and returns *this. - online_compiler &setTargetDeviceArch(device_arch arch); + online_compiler &setTargetDeviceArch(device_arch arch) { + DeviceArch = arch; + return *this; + } /// Makes the compilation target 32-bit and returns *this. - online_compiler &set32bitTarget(); + online_compiler &set32bitTarget() { + Is64Bit = false; + return *this; + }; /// Makes the compilation target 64-bit and returns *this. - online_compiler &set64bitTarget(); + online_compiler &set64bitTarget() { + Is64Bit = true; + return *this; + }; /// Sets implementation-defined target device stepping of the compilation /// target and returns *this. - online_compiler &setTargetDeviceStepping(const std::string &id); + online_compiler &setTargetDeviceStepping(const std::string &id) { + DeviceStepping = id; + return *this; + } private: - // Compilation target specification fields: { - /// Compiled code format. compiled_code_format OutputFormat; @@ -142,51 +170,48 @@ template class online_compiler { /// Target device stepping (implementation defined) std::string DeviceStepping; - // } + + /// Handles to helper functions used by the implementation. + void *CompileToSPIRVHandle = nullptr; + void *FreeSPIRVOutputsHandle = nullptr; }; // Specializations of the online_compiler class and 'compile' function for // particular languages and parameter types. -/// Compiles given OpenCL source. May throw \c online_compile_error. -/// @param src - contents of the source +/// Compiles the given OpenCL source. May throw \c online_compile_error. +/// @param src - contents of the source. +/// @param options - compilation options (implementation defined); standard +/// OpenCL JIT compiler options must be supported. +template <> +template <> +__SYCL_EXPORT std::vector +online_compiler::compile( + const std::string &src, const std::vector &options); + +/// Compiles the given OpenCL source. May throw \c online_compile_error. +/// @param src - contents of the source. template <> template <> std::vector online_compiler::compile(const std::string &src) { - // real implementation will call some non-templated impl function here - return std::vector{}; + return compile(src, std::vector{}); } -/// Compiles given OpenCL source. May throw \c online_compile_error. -/// @param src - contents of the source -/// @param options - compilation options (implementation defined); standard -/// OpenCL JIT compiler options must be supported +/// Compiles the given CM source \p src. +/// @param src - contents of the source. +/// @param options - compilation options (implementation defined). template <> template <> -std::vector online_compiler::compile( - const std::string &src, const std::vector &options) { - // real implementation will call some non-templated impl function here - return std::vector{}; -} +__SYCL_EXPORT std::vector online_compiler::compile( + const std::string &src, const std::vector &options); -/// Compiles given CM source. +/// Compiles the given CM source \p src. template <> template <> std::vector online_compiler::compile(const std::string &src) { - // real implementation will call some non-templated impl function here - return std::vector{}; -} - -/// Compiles given CM source. -/// @param options - compilation options (implementation defined) -template <> -template <> -std::vector online_compiler::compile( - const std::string &src, const std::vector &options) { - // real implementation will call some non-templated impl function here - return std::vector{}; + return compile(src, std::vector{}); } } // namespace INTEL diff --git a/sycl/include/CL/sycl/exception.hpp b/sycl/include/CL/sycl/exception.hpp index 038ddeb94fc73..10c0d55a47d0b 100644 --- a/sycl/include/CL/sycl/exception.hpp +++ b/sycl/include/CL/sycl/exception.hpp @@ -52,6 +52,8 @@ class __SYCL_EXPORT exception : public std::exception { shared_ptr_class Context = nullptr) : MMsg(Msg + " " + detail::codeToString(CLErr)), MCLErr(CLErr), MContext(Context) {} + + exception(const string_class &Msg) : MMsg(Msg), MContext(nullptr) {} }; class runtime_error : public exception { diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index b1eeba9f8b2eb..b02fcb00b74a4 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -132,6 +132,7 @@ set(SYCL_SOURCES "detail/program_impl.cpp" "detail/program_manager/program_manager.cpp" "detail/queue_impl.cpp" + "detail/online_compiler/online_compiler.cpp" "detail/os_util.cpp" "detail/platform_util.cpp" "detail/reduction.cpp" diff --git a/sycl/source/detail/online_compiler/ocloc_api.h b/sycl/source/detail/online_compiler/ocloc_api.h new file mode 100644 index 0000000000000..f073c5c524f7a --- /dev/null +++ b/sycl/source/detail/online_compiler/ocloc_api.h @@ -0,0 +1,110 @@ +//===------- ocloc_api.h --------------------------------------------------===// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// This file is copied from +// https://github.com/intel/compute-runtime/blob/master/shared/offline_compiler/source/ocloc_api.h + +#include + +#ifndef OCLOC_MAKE_VERSION +/// Generates ocloc API versions +#define OCLOC_MAKE_VERSION(_major, _minor) \ + ((_major << 16) | (_minor & 0x0000ffff)) +#endif // OCLOC_MAKE_VERSION + +typedef enum _ocloc_version_t { + OCLOC_VERSION_1_0 = OCLOC_MAKE_VERSION(1, 0), ///< version 1.0 + OCLOC_VERSION_CURRENT = OCLOC_MAKE_VERSION(1, 0), ///< latest known version + OCLOC_VERSION_FORCE_UINT32 = 0x7fffffff +} ocloc_version_t; + +#ifdef _WIN32 +#define SIGNATURE __declspec(dllexport) int __cdecl +#else +#define SIGNATURE int +#endif + +extern "C" { +/// Invokes ocloc API using C interface. Supported commands match +/// the functionality of ocloc executable (check ocloc's "help" +/// for reference : shared/offline_compiler/source/ocloc_api.cpp) +/// at https://github.com/intel/compute-runtime. +/// +/// NumArgs and argv params represent the command line. +/// Remaining params represent I/O. +/// Output params should be freed using oclocFreeOutput() when +/// no longer needed. +/// List and names of outputs match outputs of ocloc executable. +/// +/// \param NumArgs is the number of arguments to pass to ocloc. +/// +/// \param Argv is an array of arguments to be passed to ocloc. +/// +/// \param NumSources is the number of in-memory representations +/// of source files to be passed to ocloc. +/// +/// \param DataSources is an array of in-memory representations +/// of source files to be passed to ocloc. +/// +/// \param LenSources is an array of sizes of in-memory representations +/// of source files passed to ocloc as DataSources. +/// +/// \param NameSources is an array of names of in-memory representations +/// of source files passed to ocloc as DataSources. +/// +/// \param NumInputHeaders is the number of in-memory representations +/// of header files to be passed to ocloc. +/// +/// \param DataInputHeaders is an array of in-memory representations +/// of header files to be passed to ocloc. +/// +/// \param LenInputHeaders is an array of sizes of in-memory representations +/// of header files passed to ocloc as DataInputHeaders. +/// +/// \param NameInputHeaders is an array of names of in-memory representations +/// of header files passed to ocloc as DataInputHeaders. +/// +/// \param NumOutputs returns the number of outputs. +/// +/// \param DataOutputs returns an array of in-memory representations +/// of output files. +/// +/// \param LenOutputs returns an array of sizes of in-memory representations +/// of output files. +/// +/// \param NameOutputs returns an array of names of in-memory representations +/// of output files. Special name stdout.log describes output that contains +/// messages generated by ocloc (e.g. compiler errors/warnings). +/// +/// \returns 0 on succes. Returns non-0 in case of failure. +SIGNATURE oclocInvoke(uint32_t NumArgs, const char *Argv[], uint32_t NumSources, + const uint8_t **DataSources, const uint64_t *LenSources, + const char **NameSources, uint32_t NumInputHeaders, + const uint8_t **DataInputHeaders, + const uint64_t *LenInputHeaders, + const char **NameInputHeaders, uint32_t *NumOutputs, + uint8_t ***DataOutputs, uint64_t **LenOutputs, + char ***NameOutputs); + +/// Frees results of oclocInvoke +/// +/// \param NumOutputs is number of outputs as returned by oclocInvoke(). +/// +/// \param DataOutputs is array of outputs as returned by oclocInvoke(). +/// +/// \param LenOutputs is array of sizes of outputs as returned by oclocInvoke(). +/// +/// \param NameOutputs is array of names of outputs as returned by oclocInvoke() +/// +/// \returns 0 on succes. Returns non-0 in case of failure. +SIGNATURE oclocFreeOutput(uint32_t *NumOutputs, uint8_t ***DataOutputs, + uint64_t **LenOutputs, char ***NameOutputs); + +/// Returns the current version of ocloc. +SIGNATURE oclocVersion(); +} diff --git a/sycl/source/detail/online_compiler/online_compiler.cpp b/sycl/source/detail/online_compiler/online_compiler.cpp new file mode 100644 index 0000000000000..ca352c387b8a9 --- /dev/null +++ b/sycl/source/detail/online_compiler/online_compiler.cpp @@ -0,0 +1,234 @@ +//==----------- online_compiler.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 +#include +#include + +#include + +#include "ocloc_api.h" + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace INTEL { +namespace detail { + +static std::vector +prepareOclocArgs(sycl::info::device_type DeviceType, device_arch DeviceArch, + bool Is64Bit, const std::string &DeviceStepping, + const std::string &UserArgs) { + std::vector Args = {"ocloc", "-q", "-spv_only", "-device"}; + + if (DeviceType == sycl::info::device_type::gpu) { + switch (DeviceArch) { + case device_arch::gpu_gen9_5: + Args.push_back("cfl"); + break; + + case device_arch::gpu_gen11: + Args.push_back("icllp"); + break; + + default: + Args.push_back("skl"); + } + } else { + // TODO: change that to generic device when ocloc adds support for it. + // For now "skl" is used as the lowest arch with GEN9 arch. + Args.push_back("skl"); + } + + if (DeviceStepping != "") { + Args.push_back("-revision_id"); + Args.push_back(DeviceStepping.c_str()); + } + + Args.push_back(Is64Bit ? "-64" : "-32"); + + if (UserArgs != "") { + Args.push_back("-options"); + Args.push_back(UserArgs.c_str()); + } + + return Args; +} + +/// Compiles the given source \p Source to SPIR-V IL and returns IL as a vector +/// of bytes. +/// @param Source - Either OpenCL or CM source code. +/// @param DeviceType - SYCL device type, e.g. cpu, gpu, accelerator, etc. +/// @param DeviceArch - More detailed info on the target device architecture. +/// @param Is64Bit - If set to true, specifies the 64-bit architecture. +/// Otherwise, 32-bit is assumed. +/// @param DeviceStepping - implementation specific target device stepping. +/// @param CompileToSPIRVHandle - Output parameter. It is set to the address +/// of the library function doing the compilation. +/// @param FreeSPIRVOutputsHandle - Output parameter. It is set to the address +/// of the library function freeing memory +/// allocated during the compilation. +/// @param UserArgs - User's options to ocloc compiler. +static std::vector +compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, + device_arch DeviceArch, bool Is64Bit, + const std::string &DeviceStepping, void *&CompileToSPIRVHandle, + void *&FreeSPIRVOutputsHandle, + const std::vector &UserArgs) { + + if (!CompileToSPIRVHandle) { +#ifdef __SYCL_RT_OS_WINDOWS + static const std::string OclocLibraryName = "ocloc64.dll"; +#else + static const std::string OclocLibraryName = "libocloc.so"; +#endif + void *OclocLibrary = sycl::detail::pi::loadOsLibrary(OclocLibraryName); + if (!OclocLibrary) + throw online_compile_error("Cannot load ocloc library: " + + OclocLibraryName); + void *OclocVersionHandle = + sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocVersion"); + // The initial versions of ocloc library did not have the oclocVersion() + // function. Those versions had the same API as the first version of ocloc + // library having that oclocVersion() function. + int LoadedVersion = ocloc_version_t::OCLOC_VERSION_1_0; + if (OclocVersionHandle) { + decltype(::oclocVersion) *OclocVersionFunc = + reinterpret_cast(OclocVersionHandle); + LoadedVersion = OclocVersionFunc(); + } + // The loaded library with version (A.B) is compatible with expected API/ABI + // version (X.Y) used here if A == B and B >= Y. + int LoadedVersionMajor = LoadedVersion >> 16; + int LoadedVersionMinor = LoadedVersion & 0xffff; + int CurrentVersionMajor = ocloc_version_t::OCLOC_VERSION_CURRENT >> 16; + int CurrentVersionMinor = ocloc_version_t::OCLOC_VERSION_CURRENT & 0xffff; + if (LoadedVersionMajor != CurrentVersionMajor || + LoadedVersionMinor < CurrentVersionMinor) + throw online_compile_error( + std::string("Found incompatible version of ocloc library: (") + + std::to_string(LoadedVersionMajor) + "." + + std::to_string(LoadedVersionMinor) + + "). The supported versions are (" + + std::to_string(CurrentVersionMajor) + + ".N), where (N >= " + std::to_string(CurrentVersionMinor) + ")."); + + CompileToSPIRVHandle = + sycl::detail::pi::getOsLibraryFuncAddress(OclocLibrary, "oclocInvoke"); + if (!CompileToSPIRVHandle) + throw online_compile_error("Cannot load oclocInvoke() function"); + FreeSPIRVOutputsHandle = sycl::detail::pi::getOsLibraryFuncAddress( + OclocLibrary, "oclocFreeOutput"); + if (!FreeSPIRVOutputsHandle) + throw online_compile_error("Cannot load oclocFreeOutput() function"); + } + + std::string CombinedUserArgs; + for (auto UserArg : UserArgs) { + if (UserArg == "") + continue; + if (CombinedUserArgs != "") + CombinedUserArgs = CombinedUserArgs + " " + UserArg; + else + CombinedUserArgs = UserArg; + } + std::vector Args = detail::prepareOclocArgs( + DeviceType, DeviceArch, Is64Bit, DeviceStepping, CombinedUserArgs); + + uint32_t NumOutputs = 0; + byte **Outputs = nullptr; + size_t *OutputLengths = nullptr; + char **OutputNames = nullptr; + + const byte *Sources[] = {reinterpret_cast(Source.c_str())}; + const char *SourceName = "main.cl"; + const uint64_t SourceLengths[] = {Source.length() + 1}; + + Args.push_back("-file"); + Args.push_back(SourceName); + + decltype(::oclocInvoke) *OclocInvokeFunc = + reinterpret_cast(CompileToSPIRVHandle); + int CompileError = + OclocInvokeFunc(Args.size(), Args.data(), 1, Sources, SourceLengths, + &SourceName, 0, nullptr, nullptr, nullptr, &NumOutputs, + &Outputs, &OutputLengths, &OutputNames); + + byte *SpirV = nullptr; + std::string CompileLog; + size_t SpirVSize = 0; + for (uint32_t I = 0; I < NumOutputs; I++) { + size_t NameLen = strlen(OutputNames[I]); + if (NameLen >= 4 && strstr(OutputNames[I], ".spv") != nullptr && + Outputs[I] != nullptr) { + SpirVSize = OutputLengths[I]; + SpirV = new byte[SpirVSize]; + std::memcpy(SpirV, Outputs[I], SpirVSize); + } else if (!strcmp(OutputNames[I], "stdout.log")) { + CompileLog = std::string(reinterpret_cast(Outputs[I])); + } + } + + // Try to free memory before reporting possible error. + decltype(::oclocFreeOutput) *OclocFreeOutputFunc = + reinterpret_cast(FreeSPIRVOutputsHandle); + int MemFreeError = + OclocFreeOutputFunc(&NumOutputs, &Outputs, &OutputLengths, &OutputNames); + + if (CompileError) + throw online_compile_error("ocloc reported compilation errors: {\n" + + CompileLog + "\n}"); + if (!SpirV) + throw online_compile_error( + "Unexpected output: ocloc did not return SPIR-V"); + if (MemFreeError) + throw online_compile_error("ocloc cannot safely free resources"); + + return std::vector(SpirV, SpirV + SpirVSize); +} +} // namespace detail + +template <> +template <> +__SYCL_EXPORT std::vector +online_compiler::compile( + const std::string &Source, const std::vector &UserArgs) { + + if (OutputFormatVersion != std::pair{0, 0}) { + std::string Version = std::to_string(OutputFormatVersion.first) + ", " + + std::to_string(OutputFormatVersion.second); + throw online_compile_error(std::string("The output format version (") + + Version + ") is not supported yet"); + } + + return detail::compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit, + DeviceStepping, CompileToSPIRVHandle, + FreeSPIRVOutputsHandle, UserArgs); +} + +template <> +template <> +__SYCL_EXPORT std::vector online_compiler::compile( + const std::string &Source, const std::vector &UserArgs) { + + if (OutputFormatVersion != std::pair{0, 0}) { + std::string Version = std::to_string(OutputFormatVersion.first) + ", " + + std::to_string(OutputFormatVersion.second); + throw online_compile_error(std::string("The output format version (") + + Version + ") is not supported yet"); + } + + std::vector CMUserArgs = UserArgs; + CMUserArgs.push_back("-cmc"); + return detail::compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit, + DeviceStepping, CompileToSPIRVHandle, + FreeSPIRVOutputsHandle, CMUserArgs); +} + +} // namespace INTEL +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 921652e70037b..0bac17e92b320 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3626,6 +3626,8 @@ _ZN2cl4sycl5eventC1Ev _ZN2cl4sycl5eventC2EP9_cl_eventRKNS0_7contextE _ZN2cl4sycl5eventC2ESt10shared_ptrINS0_6detail10event_implEE _ZN2cl4sycl5eventC2Ev +_ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE0EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISC_EEEEES6_IhSaIhEERKSC_DpRKT_ +_ZN2cl4sycl5INTEL15online_compilerILNS1_15source_languageE1EE7compileIJSt6vectorINSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEESaISC_EEEEES6_IhSaIhEERKSC_DpRKT_ _ZN2cl4sycl5queue10mem_adviseEPKvm14_pi_mem_advice _ZN2cl4sycl5queue10wait_proxyERKNS0_6detail13code_locationE _ZN2cl4sycl5queue11submit_implESt8functionIFvRNS0_7handlerEEERKNS0_6detail13code_locationE diff --git a/sycl/test/on-device/online_compiler/online_compiler_L0.cpp b/sycl/test/on-device/online_compiler/online_compiler_L0.cpp new file mode 100644 index 0000000000000..5fd556f81a4d7 --- /dev/null +++ b/sycl/test/on-device/online_compiler/online_compiler_L0.cpp @@ -0,0 +1,60 @@ +// REQUIRES: level_zero + +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -I %sycl_source_dir -DRUN_KERNELS -lze_loader %s -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -I %sycl_source_dir -lze_loader %s -o %th.out +// RUN: %RUN_ON_HOST %th.out + +// This test checks INTEL feature class online_compiler for Level-Zero. +// All Level-Zero specific code is kept here and the common part that can be +// re-used by other backends is kept in online_compiler_common.hpp file. + +#include +#include + +#include + +// TODO: The testing of CM compilation is temporarily turned OFF on Linux +// due to problems with dependencies on libclangFEWrapper.so which is not +// currently included into NEO package on Linux. +#ifdef _WIN32 +#define COMPILE_CM_KERNEL 1 +#endif + +// clang-format off +#include +#include +// clang-format on + +using byte = unsigned char; + +#ifdef RUN_KERNELS +sycl::kernel getSYCLKernelWithIL(sycl::context &Context, + const std::vector &IL) { + + ze_module_desc_t ZeModuleDesc = {}; + ZeModuleDesc.format = ZE_MODULE_FORMAT_IL_SPIRV; + ZeModuleDesc.inputSize = IL.size(); + ZeModuleDesc.pInputModule = IL.data(); + ZeModuleDesc.pBuildFlags = ""; + ZeModuleDesc.pConstants = nullptr; + + assert(Context.get_devices().size() == 1 && "Expected to have only 1 device"); + sycl::device Device = Context.get_devices()[0]; + auto ZeDevice = Device.get_native(); + auto ZeContext = Context.get_native(); + + ze_module_build_log_handle_t ZeBuildLog; + ze_module_handle_t ZeModule; + ze_result_t ZeResult = zeModuleCreate(ZeContext, ZeDevice, &ZeModuleDesc, + &ZeModule, &ZeBuildLog); + if (ZeResult != ZE_RESULT_SUCCESS) + throw sycl::runtime_error(); + sycl::program SyclProgram = + sycl::level_zero::make(Context, ZeModule); + return SyclProgram.get_kernel("my_kernel"); +} +#endif // RUN_KERNELS + +#include "online_compiler_common.hpp" diff --git a/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp b/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp new file mode 100644 index 0000000000000..086b2da992f16 --- /dev/null +++ b/sycl/test/on-device/online_compiler/online_compiler_OpenCL.cpp @@ -0,0 +1,48 @@ +// REQUIRES: opencl + +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -DRUN_KERNELS -lOpenCL -o %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -lOpenCL -o %th.out +// RUN: %RUN_ON_HOST %th.out + +// This test checks INTEL feature class online_compiler for OpenCL. +// All OpenCL specific code is kept here and the common part that can be +// re-used by other backends is kept in online_compiler_common.hpp file. + +// TODO: The testing of CM compilation is temporarily turned OFF on Linux +// due to problems with dependencies on libclangFEWrapper.so which is not +// currently included into NEO package on Linux. +#ifdef _WIN32 +#define COMPILE_CM_KERNEL 1 +#endif + +#include +#include + +#include + +using byte = unsigned char; + +#ifdef RUN_KERNELS +sycl::kernel getSYCLKernelWithIL(sycl::context &Context, + const std::vector &IL) { + cl_int Err; + cl_program ClProgram = + clCreateProgramWithIL(Context.get(), IL.data(), IL.size(), &Err); + if (Err != CL_SUCCESS) + throw sycl::runtime_error("clCreateProgramWithIL() failed", Err); + + Err = clBuildProgram(ClProgram, 0, nullptr, nullptr, nullptr, nullptr); + if (Err != CL_SUCCESS) + throw sycl::runtime_error("clBuildProgram() failed", Err); + + cl_kernel ClKernel = clCreateKernel(ClProgram, "my_kernel", &Err); + if (Err != CL_SUCCESS) + throw sycl::runtime_error("clCreateKernel() failed", Err); + + return sycl::kernel(ClKernel, Context); +} +#endif // RUN_KERNELS + +#include "online_compiler_common.hpp" diff --git a/sycl/test/on-device/online_compiler/online_compiler_common.hpp b/sycl/test/on-device/online_compiler/online_compiler_common.hpp new file mode 100644 index 0000000000000..9ef97c45020bb --- /dev/null +++ b/sycl/test/on-device/online_compiler/online_compiler_common.hpp @@ -0,0 +1,177 @@ +#include +#include + +#include +#include + +auto constexpr CLSource = R"===( +__kernel void my_kernel(__global int *in, __global int *out) { + size_t i = get_global_id(0); + out[i] = in[i]*2 + 100; +} +)==="; + +auto constexpr CLSourceSyntaxError = R"===( +__kernel void my_kernel(__global int *in, __global int *out) { + syntax error here + size_t i = get_global_id(0); + out[i] = in[i]*2 + 100; +} +)==="; + +auto constexpr CMSource = R"===( +extern "C" +void cm_kernel() { +} +)==="; + +using namespace sycl::INTEL; + +#ifdef RUN_KERNELS +void testSyclKernel(sycl::queue &Q, sycl::kernel Kernel) { + std::cout << "Run the kernel now:\n"; + constexpr int N = 4; + int InputArray[N] = {0, 1, 2, 3}; + int OutputArray[N] = {}; + + sycl::buffer InputBuf(InputArray, sycl::range<1>(N)); + sycl::buffer OutputBuf(OutputArray, sycl::range<1>(N)); + + Q.submit([&](sycl::handler &CGH) { + CGH.set_arg(0, InputBuf.get_access(CGH)); + CGH.set_arg(1, OutputBuf.get_access(CGH)); + CGH.parallel_for(sycl::range<1>{N}, Kernel); + }); + + auto Out = OutputBuf.get_access(); + for (int I = 0; I < N; I++) + std::cout << I << "*2 + 100 = " << Out[I] << "\n"; +} +#endif // RUN_KERNELS + +int main(int argc, char **argv) { + cl::sycl::queue Q; + cl::sycl::context Context = Q.get_context(); + cl::sycl::device Device = Q.get_device(); + + { // Compile and run a trivial OpenCL kernel. + std::cout << "Test case1\n"; + online_compiler Compiler; + std::vector IL; + try { + IL = Compiler.compile( + CLSource, + // Pass two options to check that more than one is accepted. + std::vector{"-cl-fast-relaxed-math", + "-cl-finite-math-only"}); + std::cout << "IL size = " << IL.size() << "\n"; + assert(IL.size() > 0 && "Unexpected IL size"); + } catch (sycl::exception &e) { + std::cout << "Compilation to IL failed: " << e.what() << "\n"; + return 1; + } +#ifdef RUN_KERNELS + testSyclKernel(Q, getSYCLKernelWithIL(Context, IL)); +#endif // RUN_KERNELS + } + + { // Compile and run a trivial OpenCL kernel using online_compiler() + // constructor accepting SYCL device. + std::cout << "Test case2\n"; + online_compiler Compiler(Device); + std::vector IL; + try { + IL = Compiler.compile(CLSource); + std::cout << "IL size = " << IL.size() << "\n"; + assert(IL.size() > 0 && "Unexpected IL size"); + } catch (sycl::exception &e) { + std::cout << "Compilation to IL failed: " << e.what() << "\n"; + return 1; + } +#ifdef RUN_KERNELS + testSyclKernel(Q, getSYCLKernelWithIL(Context, IL)); +#endif // RUN_KERNELS + } + +#ifdef COMPILE_CM_KERNEL + // TODO: this test is temporarily turned off because CI buildbots do not set + // PATHs to clangFEWrapper library properly. + { // Compile a trivial CM kernel. + std::cout << "Test case3\n"; + online_compiler Compiler; + try { + std::vector IL = Compiler.compile(CMSource); + + std::cout << "IL size = " << IL.size() << "\n"; + assert(IL.size() > 0 && "Unexpected IL size"); + } catch (sycl::exception &e) { + std::cout << "Compilation to IL failed: " << e.what() << "\n"; + return 1; + } + } +#endif // COMPILE_CM_KERNEL + + { // Compile a source with syntax errors. + std::cout << "Test case4\n"; + online_compiler Compiler; + std::vector IL; + bool TestPassed = false; + try { + IL = Compiler.compile(CLSourceSyntaxError); + } catch (sycl::exception &e) { + std::string Msg = e.what(); + if (Msg.find("syntax error here") != std::string::npos) + TestPassed = true; + else + std::cerr << "Unexpected exception: " << Msg << "\n"; + } + assert(TestPassed && "Failed to throw an exception for syntax error"); + if (!TestPassed) + return 1; + } + + { // Compile a good CL source using unrecognized compilation options. + std::cout << "Test case5\n"; + online_compiler Compiler; + std::vector IL; + bool TestPassed = false; + try { + IL = Compiler.compile(CLSource, + // Intentionally use incorrect option. + std::vector{"WRONG_OPTION"}); + } catch (sycl::exception &e) { + std::string Msg = e.what(); + if (Msg.find("WRONG_OPTION") != std::string::npos) + TestPassed = true; + else + std::cerr << "Unexpected exception: " << Msg << "\n"; + } + assert(TestPassed && + "Failed to throw an exception for unrecognized option"); + if (!TestPassed) + return 1; + } + + { // Try compiling CM source with OpenCL compiler. + std::cout << "Test case6\n"; + online_compiler Compiler; + std::vector IL; + bool TestPassed = false; + try { + // Intentionally pass CMSource instead of CLSource. + IL = Compiler.compile(CMSource); + } catch (sycl::exception &e) { + std::string Msg = e.what(); + if (Msg.find("error: expected identifier or '('") != std::string::npos) + TestPassed = true; + else + std::cerr << "Unexpected exception: " << Msg << "\n"; + } + assert(TestPassed && "Failed to throw an exception for wrong program"); + if (!TestPassed) + return 1; + } + + std::cout << "\nAll test cases passed.\n"; + return 0; +}