[SYCL] Move kernel_compiler related information

[steffenlarsen]

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.

[steffenlarsen]

Tag @jopperm @sommerlukas @cperkinsintel @0x12CC - I still need to fix the last couple of regressions and write some tests for sycl::join with bundles created from different paths (interop, regular SYCL offline, and kernel_compiler,) but your thoughts on the approach would be appreciated earlier, if you have time. 😄

[steffenlarsen]

Friendly ping @jopperm & @sommerlukas. I believe the failures should disappear with #17461 and if not then I'll investigate immediately after rebasing. It would be good to have this go in soon though, to avoid having to keep doing heavy rebasing in perpetuum. 😄

[steffenlarsen]

Also ping @cperkinsintel 😉

[jopperm]

Looks great! I really like how you split the somewhat convoluted lifetime management into the ManagedDeviceBinaries and ManagedDeviceGlobalsRegistry classes.

[steffenlarsen]

@jopperm - Thanks a ton for the review! I found out what the issue with the test was: I had broken the dependency resolution logic when I moved away from refetching the binaries from the program manager. It should be fixed with ffbe777 if you want to have another look. 😄

[jopperm]

It should be fixed with ffbe777 if you want to have another look.

I think the code LGTM, but can this be exercised with the functionality in this PR already, or does it need linking as in #17442?

[cperkinsintel]

LGTM.
Besides moving things around and changing interfaces, there are quite a few new functions and methods introduced. Some for kernel_bundle_impl, some for device_image_impl. Are all of these exercised by the tests?

[steffenlarsen]

Besides moving things around and changing interfaces, there are quite a few new functions and methods introduced. Some for kernel_bundle_impl, some for device_image_impl. Are all of these exercised by the tests?

Since the primary changes in this patch is moving kernel compiler related logic from kernel bundle into the device images, the vast majority of this is fairly well covered by existing kernel compiler testing. The main complex new functionality is probably merging of RTC information, which the new test should at least cover at a basic level. I expect the testing we'll be doing for linking in the upcoming patch should improve that coverage significantly however.

[steffenlarsen]

It should be fixed with ffbe777 if you want to have another look.

I think the code LGTM, but can this be exercised with the functionality in this PR already, or does it need linking as in #17442?

I will see if we can make the linking logic resolve these cases, but I am not convinced it will happen in the first iteration. Eventually I would like to drop the "images with dependencies" abstraction in favor of simply populating the bundles with the dependencies as images in the bundle.

[vinser52]

Do we need ownership on Context? I think std::shared_ptr is required when you need to control the lifetime of the object. Otherwise, a raw pointer or reference to the object should be enough.

[steffenlarsen]

Having a reference to the context be a shared pointer was indeed intentional here, though you could argue it is being overly careful. To properly unregister the device globals, the implementation needs the context to be alive, which would currently be guaranteed by the device_image_impl, which owns the KernelCompilerBinaryInfo object that in turn owns this object. However, we have no guarantees that the KernelCompilerBinaryInfo or the shared_ptr to this object could be copied somewhere else or even that the context member in device_image_impl gets moved to after the KernelCompilerBinaryInfo object (order of destruction could release context before KernelCompilerBinaryInfo.)

In summary, we could use a raw pointer, but we risk some very hard-to-find bugs for unobvious changes.

You could argue that we wouldn't necessarily have to release the resources in the context explicitly, so what we could do is use a weak_ptr here instead, then skip the resource release and let the context dtor take care of it. Note though that we would still create a shared_ptr to the context during the dtor of this object when we lock the weak pointer, so at that point we've not saved the reference counting.