Skip to content

Commit b1eeac3

Browse files
joppermjopperm
committed
WIP, UR program is created via program manager
Signed-off-by: Julian Oppermann <[email protected]>
1 parent a726365 commit b1eeac3

File tree

7 files changed

+111
-61
lines changed

7 files changed

+111
-61
lines changed

‎sycl/source/detail/jit_compiler.cpp‎

Lines changed: 45 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1119,6 +1119,49 @@ sycl_device_binaries jit_compiler::createPIDeviceBinary(
11191119
return JITDeviceBinaries.back().getPIDeviceStruct();
11201120
}
11211121

1122+
const RTDeviceBinaryImage &jit_compiler::createDeviceBinaryImage(
1123+
const ::jit_compiler::RTCBundleInfo &BundleInfo) {
1124+
DeviceBinaryContainer Binary;
1125+
for (const auto &Symbol : BundleInfo.SymbolTable) {
1126+
// Create an offload entry each kernel.
1127+
// It seems to be OK to set zero for most of the information here, at least
1128+
// that is the case for compiled SPIR-V binaries.
1129+
OffloadEntryContainer Entry{Symbol.c_str(), nullptr, 0, 0, 0};
1130+
Binary.addOffloadEntry(std::move(Entry));
1131+
}
1132+
1133+
for (const auto &FPS : BundleInfo.Properties) {
1134+
PropertySetContainer PropSet{FPS.Name.c_str()};
1135+
for (const auto &FPV : FPS.Values) {
1136+
if (FPV.IsUIntValue) {
1137+
PropSet.addProperty(PropertyContainer{FPV.Name.c_str(), FPV.UIntValue});
1138+
} else {
1139+
PropSet.addProperty(PropertyContainer{
1140+
FPV.Name.c_str(), FPV.Bytes.begin(), FPV.Bytes.size(),
1141+
sycl_property_type::SYCL_PROPERTY_TYPE_BYTE_ARRAY});
1142+
}
1143+
}
1144+
Binary.addProperty(std::move(PropSet));
1145+
}
1146+
1147+
DeviceBinariesCollection Collection;
1148+
Collection.addDeviceBinary(std::move(Binary),
1149+
BundleInfo.BinaryInfo.BinaryStart,
1150+
BundleInfo.BinaryInfo.BinarySize,
1151+
(BundleInfo.BinaryInfo.AddressBits == 64)
1152+
? __SYCL_DEVICE_BINARY_TARGET_SPIRV64
1153+
: __SYCL_DEVICE_BINARY_TARGET_SPIRV32,
1154+
SYCL_DEVICE_BINARY_TYPE_SPIRV);
1155+
JITDeviceBinaries.push_back(std::move(Collection));
1156+
// TODO: If we want to handle multiple device binary images, we should instead
1157+
// return `sycl_device_binaries`, to be passed to
1158+
// `program_manager::addImages`. The program manager then creates and
1159+
// owns the `RTDeviceBinaryImage` instances.
1160+
RTCDeviceBinaryImages.emplace_back(
1161+
&JITDeviceBinaries.back().getPIDeviceStruct()->DeviceBinaries[0]);
1162+
return RTCDeviceBinaryImages.back();
1163+
}
1164+
11221165
std::vector<uint8_t> jit_compiler::encodeArgUsageMask(
11231166
const ::jit_compiler::ArgUsageMask &Mask) const {
11241167
// This must match the decoding logic in program_manager.cpp.
@@ -1167,7 +1210,7 @@ std::vector<uint8_t> jit_compiler::encodeReqdWorkGroupSize(
11671210
return Encoded;
11681211
}
11691212

1170-
std::vector<uint8_t> jit_compiler::compileSYCL(
1213+
const RTDeviceBinaryImage &jit_compiler::compileSYCL(
11711214
const std::string &Id, const std::string &SYCLSource,
11721215
const std::vector<std::pair<std::string, std::string>> &IncludePairs,
11731216
const std::vector<std::string> &UserArgs, std::string *LogPtr,
@@ -1207,10 +1250,7 @@ std::vector<uint8_t> jit_compiler::compileSYCL(
12071250
// TODO: We currently don't have a meaningful build log.
12081251
(void)LogPtr;
12091252

1210-
const auto &BI = Result.getBundleInfo().BinaryInfo;
1211-
assert(BI.Format == ::jit_compiler::BinaryFormat::SPIRV);
1212-
std::vector<uint8_t> SPV(BI.BinaryStart, BI.BinaryStart + BI.BinarySize);
1213-
return SPV;
1253+
return createDeviceBinaryImage(Result.getBundleInfo());
12141254
}
12151255

12161256
} // namespace detail

‎sycl/source/detail/jit_compiler.hpp‎

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -44,7 +44,7 @@ class jit_compiler {
4444
const std::string &KernelName,
4545
const std::vector<unsigned char> &SpecConstBlob);
4646

47-
std::vector<uint8_t> compileSYCL(
47+
const RTDeviceBinaryImage &compileSYCL(
4848
const std::string &Id, const std::string &SYCLSource,
4949
const std::vector<std::pair<std::string, std::string>> &IncludePairs,
5050
const std::vector<std::string> &UserArgs, std::string *LogPtr,
@@ -69,6 +69,9 @@ class jit_compiler {
6969
createPIDeviceBinary(const ::jit_compiler::SYCLKernelInfo &FusedKernelInfo,
7070
::jit_compiler::BinaryFormat Format);
7171

72+
const RTDeviceBinaryImage &
73+
createDeviceBinaryImage(const ::jit_compiler::RTCBundleInfo &BundleInfo);
74+
7275
std::vector<uint8_t>
7376
encodeArgUsageMask(const ::jit_compiler::ArgUsageMask &Mask) const;
7477

@@ -81,6 +84,9 @@ class jit_compiler {
8184
// Manages the lifetime of the UR structs for device binaries.
8285
std::vector<DeviceBinariesCollection> JITDeviceBinaries;
8386

87+
// Manages the lifetime of the runtime wrappers for device binary images.
88+
std::vector<RTDeviceBinaryImage> RTCDeviceBinaryImages;
89+
8490
#if SYCL_EXT_JIT_ENABLE
8591
// Handles to the entry points of the lazily loaded JIT library.
8692
using FuseKernelsFuncT = decltype(::jit_compiler::fuseKernels) *;

‎sycl/source/detail/jit_device_binaries.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -27,7 +27,7 @@ _sycl_offload_entry_struct OffloadEntryContainer::getPIOffloadEntry() {
2727
EntryFlags, EntryReserved};
2828
}
2929

30-
PropertyContainer::PropertyContainer(const std::string &Name, void *Data,
30+
PropertyContainer::PropertyContainer(const std::string &Name, const void *Data,
3131
size_t Size, uint32_t Type)
3232
: PropName{new char[Name.length() + 1]}, Value{new unsigned char[Size]},
3333
ValueSize{Size}, PropType{Type} {

‎sycl/source/detail/jit_device_binaries.hpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -50,7 +50,7 @@ class OffloadEntryContainer {
5050
class PropertyContainer {
5151

5252
public:
53-
PropertyContainer(const std::string &Name, void *Data, size_t Size,
53+
PropertyContainer(const std::string &Name, const void *Data, size_t Size,
5454
uint32_t Type);
5555
// Set a UR_PROPERTY_TYPE_UINT32 property
5656
PropertyContainer(const std::string &Name, uint32_t Data);

‎sycl/source/detail/kernel_bundle_impl.hpp‎

Lines changed: 53 additions & 51 deletions
Original file line numberDiff line numberDiff line change
@@ -468,57 +468,59 @@ class kernel_bundle_impl {
468468
}
469469

470470
if (!FetchedFromCache) {
471-
const auto spirv = [&]() -> std::vector<uint8_t> {
472-
if (Language == syclex::source_language::opencl) {
473-
// if successful, the log is empty. if failed, throws an error with
474-
// the compilation log.
475-
std::vector<uint32_t> IPVersionVec(Devices.size());
476-
std::transform(DeviceVec.begin(), DeviceVec.end(),
477-
IPVersionVec.begin(), [&](ur_device_handle_t d) {
478-
uint32_t ipVersion = 0;
479-
Adapter->call<UrApiKind::urDeviceGetInfo>(
480-
d, UR_DEVICE_INFO_IP_VERSION, sizeof(uint32_t),
481-
&ipVersion, nullptr);
482-
return ipVersion;
483-
});
484-
return syclex::detail::OpenCLC_to_SPIRV(*SourceStrPtr, IPVersionVec,
485-
BuildOptions, LogPtr);
486-
}
487-
if (Language == syclex::source_language::spirv) {
488-
const auto &SourceBytes =
489-
std::get<std::vector<std::byte>>(this->Source);
490-
std::vector<uint8_t> Result(SourceBytes.size());
491-
std::transform(SourceBytes.cbegin(), SourceBytes.cend(),
492-
Result.begin(),
493-
[](std::byte B) { return static_cast<uint8_t>(B); });
494-
return Result;
495-
}
496-
if (Language == syclex::source_language::sycl) {
497-
return syclex::detail::SYCL_to_SPIRV(*SourceStrPtr, IncludePairs,
498-
BuildOptions, LogPtr,
499-
RegisteredKernelNames);
500-
}
501-
if (Language == syclex::source_language::sycl_jit) {
502-
const auto &SourceStr = std::get<std::string>(this->Source);
503-
return syclex::detail::SYCL_JIT_to_SPIRV(SourceStr, IncludePairs,
504-
BuildOptions, LogPtr,
505-
RegisteredKernelNames);
506-
}
507-
throw sycl::exception(
508-
make_error_code(errc::invalid),
509-
"SYCL C++, OpenCL C and SPIR-V are the only supported "
510-
"languages at this time");
511-
}();
512-
513-
Adapter->call<UrApiKind::urProgramCreateWithIL>(
514-
ContextImpl->getHandleRef(), spirv.data(), spirv.size(), nullptr,
515-
&UrProgram);
516-
// program created by urProgramCreateWithIL is implicitly retained.
517-
if (UrProgram == nullptr)
518-
throw sycl::exception(
519-
sycl::make_error_code(errc::invalid),
520-
"urProgramCreateWithIL resulted in a null program handle.");
521-
471+
if (Language == syclex::source_language::sycl_jit) {
472+
const auto &SourceStr = std::get<std::string>(this->Source);
473+
const auto &Img = syclex::detail::SYCL_JIT_to_SPIRV(
474+
SourceStr, IncludePairs, BuildOptions, LogPtr,
475+
RegisteredKernelNames);
476+
UrProgram = ProgramManager::getInstance().createURProgram(Img, MContext,
477+
MDevices);
478+
} else {
479+
const auto spirv = [&]() -> std::vector<uint8_t> {
480+
if (Language == syclex::source_language::opencl) {
481+
// if successful, the log is empty. if failed, throws an error with
482+
// the compilation log.
483+
std::vector<uint32_t> IPVersionVec(Devices.size());
484+
std::transform(DeviceVec.begin(), DeviceVec.end(),
485+
IPVersionVec.begin(), [&](ur_device_handle_t d) {
486+
uint32_t ipVersion = 0;
487+
Adapter->call<UrApiKind::urDeviceGetInfo>(
488+
d, UR_DEVICE_INFO_IP_VERSION, sizeof(uint32_t),
489+
&ipVersion, nullptr);
490+
return ipVersion;
491+
});
492+
return syclex::detail::OpenCLC_to_SPIRV(*SourceStrPtr, IPVersionVec,
493+
BuildOptions, LogPtr);
494+
}
495+
if (Language == syclex::source_language::spirv) {
496+
const auto &SourceBytes =
497+
std::get<std::vector<std::byte>>(this->Source);
498+
std::vector<uint8_t> Result(SourceBytes.size());
499+
std::transform(SourceBytes.cbegin(), SourceBytes.cend(),
500+
Result.begin(),
501+
[](std::byte B) { return static_cast<uint8_t>(B); });
502+
return Result;
503+
}
504+
if (Language == syclex::source_language::sycl) {
505+
return syclex::detail::SYCL_to_SPIRV(*SourceStrPtr, IncludePairs,
506+
BuildOptions, LogPtr,
507+
RegisteredKernelNames);
508+
}
509+
throw sycl::exception(
510+
make_error_code(errc::invalid),
511+
"SYCL C++, OpenCL C and SPIR-V are the only supported "
512+
"languages at this time");
513+
}();
514+
515+
Adapter->call<UrApiKind::urProgramCreateWithIL>(
516+
ContextImpl->getHandleRef(), spirv.data(), spirv.size(), nullptr,
517+
&UrProgram);
518+
// program created by urProgramCreateWithIL is implicitly retained.
519+
if (UrProgram == nullptr)
520+
throw sycl::exception(
521+
sycl::make_error_code(errc::invalid),
522+
"urProgramCreateWithIL resulted in a null program handle.");
523+
}
522524
} // if(!FetchedFromCache)
523525

524526
std::string XsFlags = extractXsFlags(BuildOptions);

‎sycl/source/detail/kernel_compiler/kernel_compiler_sycl.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -342,7 +342,7 @@ bool SYCL_JIT_Compilation_Available() {
342342
#endif
343343
}
344344

345-
spirv_vec_t SYCL_JIT_to_SPIRV(
345+
const sycl::detail::RTDeviceBinaryImage &SYCL_JIT_to_SPIRV(
346346
[[maybe_unused]] const std::string &SYCLSource,
347347
[[maybe_unused]] include_pairs_t IncludePairs,
348348
[[maybe_unused]] const std::vector<std::string> &UserArgs,

‎sycl/source/detail/kernel_compiler/kernel_compiler_sycl.hpp‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -12,6 +12,8 @@
1212
#include <sycl/detail/export.hpp> // __SYCL_EXPORT
1313
#include <sycl/device.hpp>
1414

15+
#include <detail/device_binary_image.hpp>
16+
1517
#include <numeric> // std::accumulate
1618
#include <string>
1719
#include <vector>
@@ -33,7 +35,7 @@ bool SYCL_Compilation_Available();
3335

3436
std::string userArgsAsString(const std::vector<std::string> &UserArguments);
3537

36-
spirv_vec_t
38+
const sycl::detail::RTDeviceBinaryImage &
3739
SYCL_JIT_to_SPIRV(const std::string &Source, include_pairs_t IncludePairs,
3840
const std::vector<std::string> &UserArgs, std::string *LogPtr,
3941
const std::vector<std::string> &RegisteredKernelNames);

0 commit comments

Comments
 (0)