21 namespace experimental {
24 static std::vector<const char *>
26 bool Is64Bit,
const std::string &DeviceStepping,
27 const std::string &UserArgs) {
28 std::vector<const char *> Args = {
"ocloc",
"-q",
"-spv_only",
"-device"};
30 if (DeviceType == sycl::info::device_type::gpu) {
32 case device_arch::gpu_gen9_5:
33 Args.push_back(
"cfl");
36 case device_arch::gpu_gen11:
37 Args.push_back(
"icllp");
41 Args.push_back(
"skl");
46 Args.push_back(
"skl");
49 if (DeviceStepping !=
"") {
50 Args.push_back(
"-revision_id");
51 Args.push_back(DeviceStepping.c_str());
54 Args.push_back(Is64Bit ?
"-64" :
"-32");
57 Args.push_back(
"-options");
58 Args.push_back(UserArgs.c_str());
78 static std::vector<byte>
81 const std::string &DeviceStepping,
void *&CompileToSPIRVHandle,
82 void *&FreeSPIRVOutputsHandle,
83 const std::vector<std::string> &UserArgs) {
85 if (!CompileToSPIRVHandle) {
86 #ifdef __SYCL_RT_OS_WINDOWS
87 static const std::string OclocLibraryName =
"ocloc64.dll";
89 static const std::string OclocLibraryName =
"libocloc.so";
95 void *OclocVersionHandle =
101 if (OclocVersionHandle) {
103 reinterpret_cast<decltype(::
oclocVersion) *
>(OclocVersionHandle);
104 LoadedVersion = OclocVersionFunc();
108 int LoadedVersionMajor = LoadedVersion >> 16;
109 int LoadedVersionMinor = LoadedVersion & 0xffff;
112 if (LoadedVersionMajor != CurrentVersionMajor ||
113 LoadedVersionMinor < CurrentVersionMinor)
115 std::string(
"Found incompatible version of ocloc library: (") +
116 std::to_string(LoadedVersionMajor) +
"." +
117 std::to_string(LoadedVersionMinor) +
118 "). The supported versions are (" +
119 std::to_string(CurrentVersionMajor) +
120 ".N), where (N >= " + std::to_string(CurrentVersionMinor) +
").");
122 CompileToSPIRVHandle =
124 if (!CompileToSPIRVHandle)
127 OclocLibrary,
"oclocFreeOutput");
128 if (!FreeSPIRVOutputsHandle)
132 std::string CombinedUserArgs;
133 for (
auto UserArg : UserArgs) {
136 if (CombinedUserArgs !=
"")
137 CombinedUserArgs = CombinedUserArgs +
" " + UserArg;
139 CombinedUserArgs = UserArg;
142 DeviceType, DeviceArch, Is64Bit, DeviceStepping, CombinedUserArgs);
144 uint32_t NumOutputs = 0;
145 byte **Outputs =
nullptr;
146 size_t *OutputLengths =
nullptr;
147 char **OutputNames =
nullptr;
149 const byte *Sources[] = {
reinterpret_cast<const byte *
>(Source.c_str())};
150 const char *SourceName =
"main.cl";
151 const uint64_t SourceLengths[] = {Source.length() + 1};
153 Args.push_back(
"-file");
154 Args.push_back(SourceName);
157 reinterpret_cast<decltype(::
oclocInvoke) *
>(CompileToSPIRVHandle);
159 OclocInvokeFunc(Args.size(), Args.data(), 1, Sources, SourceLengths,
160 &SourceName, 0,
nullptr,
nullptr,
nullptr, &NumOutputs,
161 &Outputs, &OutputLengths, &OutputNames);
163 std::vector<byte> SpirV;
164 std::string CompileLog;
165 for (uint32_t I = 0; I < NumOutputs; I++) {
166 size_t NameLen = strlen(OutputNames[I]);
167 if (NameLen >= 4 && strstr(OutputNames[I],
".spv") !=
nullptr &&
168 Outputs[I] !=
nullptr) {
169 assert(SpirV.size() == 0 &&
"More than one SPIR-V output found.");
170 SpirV = std::vector<byte>(Outputs[I], Outputs[I] + OutputLengths[I]);
171 }
else if (!strcmp(OutputNames[I],
"stdout.log")) {
172 CompileLog = std::string(
reinterpret_cast<const char *
>(Outputs[I]));
178 reinterpret_cast<decltype(::
oclocFreeOutput) *
>(FreeSPIRVOutputsHandle);
180 OclocFreeOutputFunc(&NumOutputs, &Outputs, &OutputLengths, &OutputNames);
187 "Unexpected output: ocloc did not return SPIR-V");
197 __SYCL_EXPORT std::vector<byte>
199 const std::string &Source,
const std::vector<std::string> &UserArgs) {
201 if (OutputFormatVersion != std::pair<int, int>{0, 0}) {
202 std::string Version = std::to_string(OutputFormatVersion.first) +
", " +
203 std::to_string(OutputFormatVersion.second);
205 Version +
") is not supported yet");
209 DeviceStepping, CompileToSPIRVHandle,
210 FreeSPIRVOutputsHandle, UserArgs);
216 const std::string &Source,
const std::vector<std::string> &UserArgs) {
218 if (OutputFormatVersion != std::pair<int, int>{0, 0}) {
219 std::string Version = std::to_string(OutputFormatVersion.first) +
", " +
220 std::to_string(OutputFormatVersion.second);
222 Version +
") is not supported yet");
225 std::vector<std::string> CMUserArgs = UserArgs;
226 CMUserArgs.push_back(
"-cmc");
228 DeviceStepping, CompileToSPIRVHandle,
229 FreeSPIRVOutputsHandle, CMUserArgs);
237 "use 'ext::intel::experimental'
instead") intel {
238 using namespace ext::intel::experimental;
243 "use 'ext::intel::experimental'
instead") INTEL {
244 using namespace ext::intel::experimental;