Skip to content

Commit f4a3cb9

Browse files
[SYCL] Move kernel_compiler related information (#17380)
This commit moves the kernel compiler related information from the kernel bundles to the device images. This separation allows the implementation to properly join and (in the near future) link kernel bundles created from different paths. --------- Signed-off-by: Larsen, Steffen <[email protected]>
1 parent 1f5e50a commit f4a3cb9

File tree

15 files changed

+1747
-973
lines changed

15 files changed

+1747
-973
lines changed

sycl/source/backend.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -299,7 +299,8 @@ make_kernel_bundle(ur_native_handle_t NativeHandle,
299299
// symbols (e.g. when kernel_bundle is supposed to be joined with another).
300300
auto KernelIDs = std::make_shared<std::vector<kernel_id>>();
301301
auto DevImgImpl = std::make_shared<device_image_impl>(
302-
nullptr, TargetContext, Devices, State, KernelIDs, UrProgram);
302+
nullptr, TargetContext, Devices, State, KernelIDs, UrProgram,
303+
ImageOriginInterop);
303304
device_image_plain DevImg{DevImgImpl};
304305

305306
return std::make_shared<kernel_bundle_impl>(TargetContext, Devices, DevImg);

sycl/source/detail/device_image_impl.hpp

Lines changed: 770 additions & 16 deletions
Large diffs are not rendered by default.

sycl/source/detail/graph_impl.cpp

Lines changed: 4 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1534,18 +1534,10 @@ void exec_graph_impl::populateURKernelUpdateStructs(
15341534
std::shared_ptr<sycl::detail::kernel_impl> SyclKernelImpl = nullptr;
15351535
const sycl::detail::KernelArgMask *EliminatedArgMask = nullptr;
15361536

1537-
// Use kernel_bundle if available unless it is interop.
1538-
// Interop bundles can't be used in the first branch, because the kernels
1539-
// in interop kernel bundles (if any) do not have kernel_id
1540-
// and can therefore not be looked up, but since they are self-contained
1541-
// they can simply be launched directly.
1542-
if (KernelBundleImplPtr && !KernelBundleImplPtr->isInterop()) {
1543-
const auto &KernelName = ExecCG.MKernelName;
1544-
kernel_id KernelID =
1545-
sycl::detail::ProgramManager::getInstance().getSYCLKernelID(KernelName);
1546-
kernel SyclKernel =
1547-
KernelBundleImplPtr->get_kernel(KernelID, KernelBundleImplPtr);
1548-
SyclKernelImpl = sycl::detail::getSyclObjImpl(SyclKernel);
1537+
if (auto SyclKernelImpl = KernelBundleImplPtr
1538+
? KernelBundleImplPtr->tryGetKernel(
1539+
ExecCG.MKernelName, KernelBundleImplPtr)
1540+
: std::shared_ptr<kernel_impl>{nullptr}) {
15491541
UrKernel = SyclKernelImpl->getHandleRef();
15501542
EliminatedArgMask = SyclKernelImpl->getKernelArgMask();
15511543
} else if (Kernel != nullptr) {

sycl/source/detail/helpers.cpp

Lines changed: 7 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -72,17 +72,14 @@ retrieveKernelBinary(const QueueImplPtr &Queue, const char *KernelName,
7272

7373
const RTDeviceBinaryImage *DeviceImage = nullptr;
7474
ur_program_handle_t Program = nullptr;
75-
if (KernelCG->getKernelBundle() != nullptr) {
75+
auto KernelBundleImpl = KernelCG->getKernelBundle();
76+
if (auto SyclKernelImpl =
77+
KernelBundleImpl
78+
? KernelBundleImpl->tryGetKernel(KernelName, KernelBundleImpl)
79+
: std::shared_ptr<kernel_impl>{nullptr}) {
7680
// Retrieve the device image from the kernel bundle.
77-
auto KernelBundle = KernelCG->getKernelBundle();
78-
kernel_id KernelID =
79-
detail::ProgramManager::getInstance().getSYCLKernelID(KernelName);
80-
81-
auto SyclKernel = detail::getSyclObjImpl(
82-
KernelBundle->get_kernel(KernelID, KernelBundle));
83-
84-
DeviceImage = SyclKernel->getDeviceImage()->get_bin_image_ref();
85-
Program = SyclKernel->getDeviceImage()->get_ur_program_ref();
81+
DeviceImage = SyclKernelImpl->getDeviceImage()->get_bin_image_ref();
82+
Program = SyclKernelImpl->getDeviceImage()->get_ur_program_ref();
8683
} else if (KernelCG->MSyclKernel != nullptr) {
8784
DeviceImage = KernelCG->MSyclKernel->getDeviceImage()->get_bin_image_ref();
8885
Program = KernelCG->MSyclKernel->getDeviceImage()->get_ur_program_ref();

0 commit comments

Comments
 (0)