32 inline namespace _V1 {
58 using SpecConstMapT = std::map<std::string, std::vector<SpecConstDescT>>;
62 std::shared_ptr<std::vector<kernel_id>> KernelIDs,
64 : MBinImage(BinImage), MContext(
std::move(Context)),
65 MDevices(
std::move(Devices)), MState(State), MProgram(Program),
66 MKernelIDs(
std::move(KernelIDs)),
67 MSpecConstsDefValBlob(getSpecConstsDefValBlob()) {
68 updateSpecConstSymMap();
73 std::shared_ptr<std::vector<kernel_id>> KernelIDs,
76 const std::vector<unsigned char> &SpecConstsBlob)
77 : MBinImage(BinImage), MContext(
std::move(Context)),
78 MDevices(
std::move(Devices)), MState(State), MProgram(Program),
79 MKernelIDs(
std::move(KernelIDs)), MSpecConstsBlob(SpecConstsBlob),
80 MSpecConstsDefValBlob(getSpecConstsDefValBlob()),
81 MSpecConstSymMap(SpecConstMap) {}
84 return std::binary_search(MKernelIDs->begin(), MKernelIDs->end(),
92 for (
const device &Device : MDevices)
93 if (Device == DeviceCand)
100 DeviceCand.get_info<info::device::parent_device>());
113 const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
114 return !MSpecConstSymMap.empty();
127 "native_specialization_constant() called for unimplemented case");
129 auto IsJITSPIRVTarget = [](
const char *Target) {
142 const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
143 return MSpecConstSymMap.count(SpecName) != 0;
151 const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
153 if (MSpecConstSymMap.count(
std::string{SpecName}) == 0)
156 std::vector<SpecConstDescT> &Descs =
163 if (MSpecConstsDefValBlob.
size() &&
164 (std::memcmp(MSpecConstsDefValBlob.
begin() + Desc.BlobOffset,
165 static_cast<const char *
>(Value) + Desc.CompositeOffset,
175 std::memcpy(MSpecConstsBlob.data() + Desc.BlobOffset,
176 static_cast<const char *
>(Value) + Desc.CompositeOffset,
187 const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
188 assert(IsSet || MSpecConstsDefValBlob.
size());
190 const std::vector<SpecConstDescT> &Descs =
194 IsSet ? MSpecConstsBlob.data() : MSpecConstsDefValBlob.
begin();
195 std::memcpy(
static_cast<char *
>(ValueRet) + Desc.CompositeOffset,
196 Blob + Desc.BlobOffset, Desc.Size);
204 const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
205 if (MSpecConstSymMap.count(
std::string{SpecName}) == 0)
208 const std::vector<SpecConstDescT> &Descs =
210 return Descs.front().IsSet;
217 const std::lock_guard<std::mutex> SpecConstLock(MSpecConstAccessMtx);
218 for (
auto &SpecConst : MSpecConstSymMap) {
219 for (
auto &Desc : SpecConst.second) {
230 MBinImage->
getProperty(
"specConstsReplacedWithDefault");
242 MDevices.begin(), MDevices.end(),
243 [&Dev](
const device &DevCand) { return Dev == DevCand; });
259 return MSpecConstsBlob;
263 std::lock_guard<std::mutex> Lock{MSpecConstAccessMtx};
264 if (
nullptr == MSpecConstsBuffer && !MSpecConstsBlob.empty()) {
274 MSpecConstsBlob.size(), MSpecConstsBlob.data(),
275 &MSpecConstsBuffer,
nullptr);
277 return MSpecConstsBuffer;
281 return MSpecConstSymMap;
285 return MSpecConstAccessMtx;
299 return NativeProgram;
308 if (MSpecConstsBuffer) {
309 std::lock_guard<std::mutex> Lock{MSpecConstAccessMtx};
317 ByteArray getSpecConstsDefValBlob()
const {
324 if (!SCDefValRange.
size())
331 return DefValDescriptors;
334 void updateSpecConstSymMap() {
342 unsigned BlobOffset = 0;
343 for (SCItTy SCIt : SCRange) {
344 const char *SCName = (*SCIt)->Name;
346 ByteArray Descriptors = DeviceBinaryProperty(*SCIt).asByteArray();
348 Descriptors.dropBytes(8);
355 unsigned LocalOffset = 0;
356 while (!Descriptors.empty()) {
357 auto [Id, CompositeOffset, Size] =
358 Descriptors.consume<uint32_t, uint32_t, uint32_t>();
361 const unsigned OffsetFromLast = CompositeOffset - LocalOffset;
362 BlobOffset += OffsetFromLast;
370 SpecConstDescT{Id, CompositeOffset, Size, BlobOffset});
372 LocalOffset += OffsetFromLast + Size;
376 MSpecConstsBlob.resize(BlobOffset);
378 if (MSpecConstsDefValBlob.
size()) {
379 assert(MSpecConstsDefValBlob.
size() == MSpecConstsBlob.size() &&
380 "Specialization constant default value blob do not have the "
382 std::uninitialized_copy(MSpecConstsDefValBlob.
begin(),
383 MSpecConstsDefValBlob.
begin() +
384 MSpecConstsBlob.size(),
385 MSpecConstsBlob.data());
390 const RTDeviceBinaryImage *MBinImage =
nullptr;
392 std::vector<device> MDevices;
398 std::shared_ptr<std::vector<kernel_id>> MKernelIDs;
402 mutable std::mutex MSpecConstAccessMtx;
405 std::vector<unsigned char> MSpecConstsBlob;
408 const ByteArray MSpecConstsDefValBlob;
415 std::map<std::string, std::vector<SpecConstDescT>> MSpecConstSymMap;
The context class represents a SYCL context on which kernel functions may be executed.
backend get_backend() const noexcept
Returns the backend associated with this context.
ConstIterator begin() const
void dropBytes(std::size_t Bytes)
ByteArray asByteArray() const
pi_uint32 asUint32() const
ConstIterator begin() const
const pi_device_binary_struct & getRawData() const
pi_device_binary_property getProperty(const char *PropName) const
Returns a single property from SYCL_MISC_PROP category.
const PropertyRange & getSpecConstants() const
Gets the iterator range over specialization constants in this binary image.
const PropertyRange & getSpecConstantsDefaultValues() const
const std::vector< device > & get_devices() const noexcept
std::shared_ptr< std::vector< kernel_id > > & get_kernel_ids_ptr() noexcept
bool has_kernel(const kernel_id &KernelIDCand) const noexcept
std::vector< unsigned char > & get_spec_const_blob_ref() noexcept
bool has_kernel(const kernel_id &KernelIDCand, const device &DeviceCand) const noexcept
device_image_impl(const RTDeviceBinaryImage *BinImage, context Context, std::vector< device > Devices, bundle_state State, std::shared_ptr< std::vector< kernel_id >> KernelIDs, sycl::detail::pi::PiProgram Program, const SpecConstMapT &SpecConstMap, const std::vector< unsigned char > &SpecConstsBlob)
bool is_specialization_constant_set(const char *SpecName) const noexcept
const SpecConstMapT & get_spec_const_data_ref() const noexcept
void set_specialization_constant_raw_value(const char *SpecName, const void *Value) noexcept
sycl::detail::pi::PiMem & get_spec_const_buffer_ref() noexcept
std::map< std::string, std::vector< SpecConstDescT > > SpecConstMapT
bool specialization_constants_replaced_with_default() const noexcept
device_image_impl(const RTDeviceBinaryImage *BinImage, context Context, std::vector< device > Devices, bundle_state State, std::shared_ptr< std::vector< kernel_id >> KernelIDs, sycl::detail::pi::PiProgram Program)
bool has_specialization_constants() const noexcept
const context & get_context() const noexcept
const std::vector< kernel_id > & get_kernel_ids() const noexcept
pi_native_handle getNative() const
std::mutex & get_spec_const_data_lock() noexcept
bundle_state get_state() const noexcept
bool all_specialization_constant_native() const noexcept
bool compatible_with_device(const device &Dev) const
const sycl::detail::pi::PiProgram & get_program_ref() const noexcept
void get_specialization_constant_raw_value(const char *SpecName, void *ValueRet) const noexcept
bool has_specialization_constant(const char *SpecName) const noexcept
bool is_any_specialization_constant_set() const noexcept
const RTDeviceBinaryImage *& get_bin_image_ref() noexcept
void set_state(bundle_state NewState) noexcept
The SYCL device class encapsulates a single SYCL device on which kernels may be executed.
Objects of the class identify kernel is some kernel_bundle related APIs.
void memBufferCreateHelper(const PluginPtr &Plugin, pi_context Ctx, pi_mem_flags Flags, size_t Size, void *HostPtr, pi_mem *RetMem, const pi_mem_properties *Props=nullptr)
std::shared_ptr< sycl::detail::context_impl > ContextImplPtr
decltype(Obj::impl) getSyclObjImpl(const Obj &SyclObject)
void memReleaseHelper(const PluginPtr &Plugin, pi_mem Mem)
std::shared_ptr< plugin > PluginPtr
uintptr_t pi_native_handle
pi_result piProgramRetain(pi_program program)
pi_result piProgramRelease(pi_program program)
constexpr pi_mem_flags PI_MEM_FLAGS_HOST_PTR_COPY
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64
SPIR-V 64-bit image <-> "spir64", 64-bit OpenCL device.
constexpr pi_mem_flags PI_MEM_FLAGS_ACCESS_RW
#define __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV32
SPIR-V 32-bit image <-> "spir", 32-bit OpenCL device.
pi_result piextProgramGetNativeHandle(pi_program program, pi_native_handle *nativeHandle)
Gets the native handle of a PI program object.
C++ wrapper of extern "C" PI interfaces.
bool any_of(const simd_mask< _Tp, _Abi > &) noexcept
_Abi const simd< _Tp, _Abi > & noexcept
const char * DeviceTargetSpec
null-terminated string representation of the device's target architecture which holds one of: __SYCL_...
bool operator()(const T &LHS, const T &RHS) const
unsigned int CompositeOffset