1 //==----------- online_compiler.cpp ----------------------------------------==// 2 // 3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. 4 // See https://llvm.org/LICENSE.txt for license information. 5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception 6 // 7 //===----------------------------------------------------------------------===// 8 #include <cstring> 9 #include <string> 10 #include <dlfcn.h> 11 12 #include <sycl/sycl.hpp> 13 #include "ocloc_api.h" 14 #include "online_compiler.hpp" 15 16 namespace sycl { 17 namespace ext::libceed { 18 19 void *loadOsLibrary(const std::string &PluginPath) { 20 // TODO: Check if the option RTLD_NOW is correct. Explore using 21 // RTLD_DEEPBIND option when there are multiple plugins. 22 void *so = dlopen(PluginPath.c_str(), RTLD_NOW); 23 if (!so) { 24 char *Error = dlerror(); 25 std::cerr << "dlopen(" << PluginPath << ") failed with <" << (Error ? Error : "unknown error") << ">" << std::endl; 26 } 27 return so; 28 } 29 30 // int unloadOsPluginLibrary(void *Library) { 31 // // The mock plugin does not have an associated library, so we allow nullptr 32 // // here to avoid it trying to free a non-existent library. 33 // if (!Library) 34 // return 0; 35 // return dlclose(Library); 36 // } 37 38 void *getOsLibraryFuncAddress(void *Library, const std::string &FunctionName) { return dlsym(Library, FunctionName.c_str()); } 39 40 static std::vector<const char *> prepareOclocArgs(sycl::info::device_type DeviceType, device_arch DeviceArch, bool Is64Bit, 41 const std::string &DeviceStepping, const std::string &UserArgs) { 42 std::vector<const char *> Args = {"ocloc", "-q", "-spv_only", "-device"}; 43 44 if (DeviceType == sycl::info::device_type::gpu) { 45 switch (DeviceArch) { 46 case device_arch::gpu_gen9: 47 Args.push_back("skl"); 48 break; 49 50 case device_arch::gpu_gen9_5: 51 Args.push_back("cfl"); 52 break; 53 54 case device_arch::gpu_gen11: 55 Args.push_back("icllp"); 56 break; 57 58 case device_arch::gpu_gen12: 59 Args.push_back("tgllp"); 60 break; 61 62 default: 63 Args.push_back("pvc"); 64 } 65 } else { 66 // TODO: change that to generic device when ocloc adds support for it. 67 // For now "tgllp" is used as the option supported on all known GPU RT. 68 Args.push_back("pvc"); 69 } 70 71 if (DeviceStepping != "") { 72 Args.push_back("-revision_id"); 73 Args.push_back(DeviceStepping.c_str()); 74 } 75 76 Args.push_back(Is64Bit ? "-64" : "-32"); 77 78 if (UserArgs != "") { 79 Args.push_back("-options"); 80 Args.push_back(UserArgs.c_str()); 81 } 82 83 return Args; 84 } 85 86 /// Compiles the given source \p Source to SPIR-V IL and returns IL as a vector 87 /// of bytes. 88 /// @param Source - Either OpenCL or CM source code. 89 /// @param DeviceType - SYCL device type, e.g. cpu, gpu, accelerator, etc. 90 /// @param DeviceArch - More detailed info on the target device architecture. 91 /// @param Is64Bit - If set to true, specifies the 64-bit architecture. 92 /// Otherwise, 32-bit is assumed. 93 /// @param DeviceStepping - implementation specific target device stepping. 94 /// @param CompileToSPIRVHandle - Output parameter. It is set to the address 95 /// of the library function doing the compilation. 96 /// @param FreeSPIRVOutputsHandle - Output parameter. It is set to the address 97 /// of the library function freeing memory 98 /// allocated during the compilation. 99 /// @param UserArgs - User's options to ocloc compiler. 100 static std::vector<byte> compileToSPIRV(const std::string &Source, sycl::info::device_type DeviceType, device_arch DeviceArch, bool Is64Bit, 101 const std::string &DeviceStepping, void *&CompileToSPIRVHandle, void *&FreeSPIRVOutputsHandle, 102 const std::vector<std::string> &UserArgs) { 103 if (!CompileToSPIRVHandle) { 104 #ifdef __SYCL_RT_OS_WINDOWS 105 static const std::string OclocLibraryName = "ocloc64.dll"; 106 #else 107 static const std::string OclocLibraryName = "libocloc.so"; 108 #endif 109 void *OclocLibrary = loadOsLibrary(OclocLibraryName); 110 if (!OclocLibrary) throw online_compile_error("Cannot load ocloc library: " + OclocLibraryName); 111 void *OclocVersionHandle = getOsLibraryFuncAddress(OclocLibrary, "oclocVersion"); 112 // The initial versions of ocloc library did not have the oclocVersion() 113 // function. Those versions had the same API as the first version of ocloc 114 // library having that oclocVersion() function. 115 int LoadedVersion = ocloc_version_t::OCLOC_VERSION_1_0; 116 if (OclocVersionHandle) { 117 decltype(::oclocVersion) *OclocVersionFunc = reinterpret_cast<decltype(::oclocVersion) *>(OclocVersionHandle); 118 LoadedVersion = OclocVersionFunc(); 119 } 120 // The loaded library with version (A.B) is compatible with expected API/ABI 121 // version (X.Y) used here if A == B and B >= Y. 122 int LoadedVersionMajor = LoadedVersion >> 16; 123 int LoadedVersionMinor = LoadedVersion & 0xffff; 124 int CurrentVersionMajor = ocloc_version_t::OCLOC_VERSION_CURRENT >> 16; 125 int CurrentVersionMinor = ocloc_version_t::OCLOC_VERSION_CURRENT & 0xffff; 126 if (LoadedVersionMajor != CurrentVersionMajor || LoadedVersionMinor < CurrentVersionMinor) 127 throw online_compile_error(std::string("Found incompatible version of ocloc library: (") + std::to_string(LoadedVersionMajor) + "." + 128 std::to_string(LoadedVersionMinor) + "). The supported versions are (" + std::to_string(CurrentVersionMajor) + 129 ".N), where (N >= " + std::to_string(CurrentVersionMinor) + ")."); 130 131 CompileToSPIRVHandle = getOsLibraryFuncAddress(OclocLibrary, "oclocInvoke"); 132 if (!CompileToSPIRVHandle) throw online_compile_error("Cannot load oclocInvoke() function"); 133 FreeSPIRVOutputsHandle = getOsLibraryFuncAddress(OclocLibrary, "oclocFreeOutput"); 134 if (!FreeSPIRVOutputsHandle) throw online_compile_error("Cannot load oclocFreeOutput() function"); 135 } 136 137 std::string CombinedUserArgs; 138 for (auto UserArg : UserArgs) { 139 if (UserArg == "") continue; 140 if (CombinedUserArgs != "") CombinedUserArgs = CombinedUserArgs + " " + UserArg; 141 else CombinedUserArgs = UserArg; 142 } 143 std::vector<const char *> Args = prepareOclocArgs(DeviceType, DeviceArch, Is64Bit, DeviceStepping, CombinedUserArgs); 144 145 uint32_t NumOutputs = 0; 146 byte **Outputs = nullptr; 147 uint64_t *OutputLengths = nullptr; 148 char **OutputNames = nullptr; 149 150 const byte *Sources[] = {reinterpret_cast<const byte *>(Source.c_str())}; 151 const char *SourceName = "main.cl"; 152 const uint64_t SourceLengths[] = {Source.length() + 1}; 153 154 Args.push_back("-file"); 155 Args.push_back(SourceName); 156 157 decltype(::oclocInvoke) *OclocInvokeFunc = reinterpret_cast<decltype(::oclocInvoke) *>(CompileToSPIRVHandle); 158 int CompileError = OclocInvokeFunc(Args.size(), Args.data(), 1, Sources, SourceLengths, &SourceName, 0, nullptr, nullptr, nullptr, &NumOutputs, 159 &Outputs, &OutputLengths, &OutputNames); 160 161 std::vector<byte> SpirV; 162 std::string CompileLog; 163 for (uint32_t I = 0; I < NumOutputs; I++) { 164 size_t NameLen = strlen(OutputNames[I]); 165 if (NameLen >= 4 && strstr(OutputNames[I], ".spv") != nullptr && Outputs[I] != nullptr) { 166 assert(SpirV.size() == 0 && "More than one SPIR-V output found."); 167 SpirV = std::vector<byte>(Outputs[I], Outputs[I] + OutputLengths[I]); 168 } else if (!strcmp(OutputNames[I], "stdout.log")) { 169 CompileLog = std::string(reinterpret_cast<const char *>(Outputs[I])); 170 } 171 } 172 173 // Try to free memory before reporting possible error. 174 decltype(::oclocFreeOutput) *OclocFreeOutputFunc = reinterpret_cast<decltype(::oclocFreeOutput) *>(FreeSPIRVOutputsHandle); 175 int MemFreeError = OclocFreeOutputFunc(&NumOutputs, &Outputs, &OutputLengths, &OutputNames); 176 177 if (CompileError) throw online_compile_error("ocloc reported compilation errors: {\n" + CompileLog + "\n}"); 178 if (SpirV.empty()) throw online_compile_error("Unexpected output: ocloc did not return SPIR-V"); 179 if (MemFreeError) throw online_compile_error("ocloc cannot safely free resources"); 180 181 return SpirV; 182 } 183 184 template <> 185 template <> 186 std::vector<byte> online_compiler<source_language::opencl_c>::compile(const std::string &Source, const std::vector<std::string> &UserArgs) { 187 if (OutputFormatVersion != std::pair<int, int>{0, 0}) { 188 std::string Version = std::to_string(OutputFormatVersion.first) + ", " + std::to_string(OutputFormatVersion.second); 189 throw online_compile_error(std::string("The output format version (") + Version + ") is not supported yet"); 190 } 191 192 return compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit, DeviceStepping, CompileToSPIRVHandle, FreeSPIRVOutputsHandle, UserArgs); 193 } 194 195 template <> 196 template <> 197 std::vector<byte> online_compiler<source_language::cm>::compile(const std::string &Source, const std::vector<std::string> &UserArgs) { 198 if (OutputFormatVersion != std::pair<int, int>{0, 0}) { 199 std::string Version = std::to_string(OutputFormatVersion.first) + ", " + std::to_string(OutputFormatVersion.second); 200 throw online_compile_error(std::string("The output format version (") + Version + ") is not supported yet"); 201 } 202 203 std::vector<std::string> CMUserArgs = UserArgs; 204 CMUserArgs.push_back("-cmc"); 205 return compileToSPIRV(Source, DeviceType, DeviceArch, Is64Bit, DeviceStepping, CompileToSPIRVHandle, FreeSPIRVOutputsHandle, CMUserArgs); 206 } 207 208 } // namespace ext::libceed 209 } // namespace sycl 210