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
loadOsLibrary(const std::string & PluginPath)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
getOsLibraryFuncAddress(void * Library,const std::string & FunctionName)38 void *getOsLibraryFuncAddress(void *Library, const std::string &FunctionName) { return dlsym(Library, FunctionName.c_str()); }
39
prepareOclocArgs(sycl::info::device_type DeviceType,device_arch DeviceArch,bool Is64Bit,const std::string & DeviceStepping,const std::string & UserArgs)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.
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<std::string> & UserArgs)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 <>
compile(const std::string & Source,const std::vector<std::string> & UserArgs)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 <>
compile(const std::string & Source,const std::vector<std::string> & UserArgs)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