19 namespace ext::intel::experimental {
22 static std::vector<const char *>
24 bool Is64Bit,
const std::string &DeviceStepping,
25 const std::string &UserArgs) {
26 std::vector<const char *> Args = {
"ocloc",
"-q",
"-spv_only",
"-device"};
28 if (DeviceType == sycl::info::device_type::gpu) {
30 case device_arch::gpu_gen9:
31 Args.push_back(
"skl");
34 case device_arch::gpu_gen9_5:
35 Args.push_back(
"cfl");
38 case device_arch::gpu_gen11:
39 Args.push_back(
"icllp");
42 case device_arch::gpu_gen12:
43 Args.push_back(
"tgllp");
47 Args.push_back(
"tgllp");
52 Args.push_back(
"tgllp");
55 if (DeviceStepping !=
"") {
56 Args.push_back(
"-revision_id");
57 Args.push_back(DeviceStepping.c_str());
60 Args.push_back(Is64Bit ?
"-64" :
"-32");
63 Args.push_back(
"-options");
64 Args.push_back(UserArgs.c_str());
84 static std::vector<byte>
87 const std::string &DeviceStepping,
void *&CompileToSPIRVHandle,
88 void *&FreeSPIRVOutputsHandle,
89 const std::vector<std::string> &UserArgs) {
91 if (!CompileToSPIRVHandle) {
92 #ifdef __SYCL_RT_OS_WINDOWS
93 static const std::string OclocLibraryName =
"ocloc64.dll";
95 static const std::string OclocLibraryName =
"libocloc.so";
101 void *OclocVersionHandle =
107 if (OclocVersionHandle) {
109 reinterpret_cast<decltype(::
oclocVersion) *
>(OclocVersionHandle);
110 LoadedVersion = OclocVersionFunc();
114 int LoadedVersionMajor = LoadedVersion >> 16;
115 int LoadedVersionMinor = LoadedVersion & 0xffff;
118 if (LoadedVersionMajor != CurrentVersionMajor ||
119 LoadedVersionMinor < CurrentVersionMinor)
121 std::string(
"Found incompatible version of ocloc library: (") +
122 std::to_string(LoadedVersionMajor) +
"." +
123 std::to_string(LoadedVersionMinor) +
124 "). The supported versions are (" +
125 std::to_string(CurrentVersionMajor) +
126 ".N), where (N >= " + std::to_string(CurrentVersionMinor) +
").");
128 CompileToSPIRVHandle =
130 if (!CompileToSPIRVHandle)
133 OclocLibrary,
"oclocFreeOutput");
134 if (!FreeSPIRVOutputsHandle)
138 std::string CombinedUserArgs;
139 for (
auto UserArg : UserArgs) {
142 if (CombinedUserArgs !=
"")
143 CombinedUserArgs = CombinedUserArgs +
" " + UserArg;
145 CombinedUserArgs = UserArg;
148 DeviceType, DeviceArch, Is64Bit, DeviceStepping, CombinedUserArgs);
150 uint32_t NumOutputs = 0;
151 byte **Outputs =
nullptr;
152 uint64_t *OutputLengths =
nullptr;
153 char **OutputNames =
nullptr;
155 const byte *Sources[] = {
reinterpret_cast<const byte *
>(Source.c_str())};
156 const char *SourceName =
"main.cl";
157 const uint64_t SourceLengths[] = {Source.length() + 1};
159 Args.push_back(
"-file");
160 Args.push_back(SourceName);
163 reinterpret_cast<decltype(::
oclocInvoke) *
>(CompileToSPIRVHandle);
165 OclocInvokeFunc(Args.size(), Args.data(), 1, Sources, SourceLengths,
166 &SourceName, 0,
nullptr,
nullptr,
nullptr, &NumOutputs,
167 &Outputs, &OutputLengths, &OutputNames);
169 std::vector<byte> SpirV;
170 std::string CompileLog;
171 for (uint32_t I = 0; I < NumOutputs; I++) {
172 size_t NameLen = strlen(OutputNames[I]);
173 if (NameLen >= 4 && strstr(OutputNames[I],
".spv") !=
nullptr &&
174 Outputs[I] !=
nullptr) {
175 assert(SpirV.size() == 0 &&
"More than one SPIR-V output found.");
176 SpirV = std::vector<byte>(Outputs[I], Outputs[I] + OutputLengths[I]);
177 }
else if (!strcmp(OutputNames[I],
"stdout.log")) {
178 CompileLog = std::string(
reinterpret_cast<const char *
>(Outputs[I]));
184 reinterpret_cast<decltype(::
oclocFreeOutput) *
>(FreeSPIRVOutputsHandle);
186 OclocFreeOutputFunc(&NumOutputs, &Outputs, &OutputLengths, &OutputNames);
193 "Unexpected output: ocloc did not return SPIR-V");
203 __SYCL_EXPORT std::vector<byte>
205 const std::string &Source,
const std::vector<std::string> &UserArgs) {
207 if (OutputFormatVersion != std::pair<int, int>{0, 0}) {
208 std::string Version = std::to_string(OutputFormatVersion.first) +
", " +
209 std::to_string(OutputFormatVersion.second);
211 Version +
") is not supported yet");
215 DeviceStepping, CompileToSPIRVHandle,
216 FreeSPIRVOutputsHandle, UserArgs);
222 const std::string &Source,
const std::vector<std::string> &UserArgs) {
224 if (OutputFormatVersion != std::pair<int, int>{0, 0}) {
225 std::string Version = std::to_string(OutputFormatVersion.first) +
", " +
226 std::to_string(OutputFormatVersion.second);
228 Version +
") is not supported yet");
231 std::vector<std::string> CMUserArgs = UserArgs;
232 CMUserArgs.push_back(
"-cmc");
234 DeviceStepping, CompileToSPIRVHandle,
235 FreeSPIRVOutputsHandle, CMUserArgs);
241 "
use 'ext::intel::experimental'
instead") intel {
242 using namespace ext::intel::experimental;
247 "
use 'ext::intel::experimental'
instead") INTEL {
248 using namespace ext::intel::experimental;