[SYCL][SYCLBIN] Fix kernel lookup after link of SYCLBIN + RTC bundles#22218
Open
koparasy wants to merge 8 commits into
Open
[SYCL][SYCLBIN] Fix kernel lookup after link of SYCLBIN + RTC bundles#22218koparasy wants to merge 8 commits into
koparasy wants to merge 8 commits into
Conversation
againull
approved these changes
Jun 5, 2026
`sycl::link({RTC_obj, SYCLBIN_obj})` collapses both inputs into a single
device_image_impl whose origin mask is KernelCompiler|SYCLBIN and whose
MRTCBinInfo is set. tryGetExtensionKernel always took the RTC branch and
looked up SYCLBIN-origin kernels via ProgramManager::tryGetSYCLKernelID,
which only knows RTC-registered kernels, returning nullptr without ever
falling through to the SYCLBIN urKernelCreate path. Any kernel that came
from the SYCLBIN object was unreachable after the link, while RTC-origin
kernels kept working.
Fall through to the SYCLBIN urKernelCreate path on RTC lookup miss when
the merged image carries ImageOriginSYCLBIN and the requested name is in
the image's kernel-name set.
Adds e2e tests link_rtc_bidir_{object,input}.cpp exercising the merged
kernel-name lookup for both origins post-link.
Documents -fsycl-allow-device-image-dependencies in UsersManual.md,
SYCLBINDesign.md (driver flag table + new "Linking SYCLBIN files"
section) and the sycl_ext_oneapi_syclbin extension (Examples section).
…ject
When the user requests a linkable SYCLBIN artifact via
-fsyclbin=input or -fsyclbin=object, the driver now implies:
* -fsycl-allow-device-image-dependencies, so sycl-post-link emits
the SYCL/exported and SYCL/imported symbol property sets that the
runtime link graph needs.
* -ftarget-export-symbols (-> -library-compilation) for AOT
targets, so IGC keeps symbols externally visible in the resulting
zebin for resolution via zeModuleDynamicLink.
-fsyclbin=executable is unchanged: closed-world codegen path
preserved, no flag implied.
Diagnostics added for incompatible combinations:
* error on -fsyclbin=input/object combined with
-fno-sycl-allow-device-image-dependencies (a contradiction; the
message redirects users to -fsyclbin=executable).
* warning on -fsyclbin=executable combined with
-fsycl-allow-device-image-dependencies (harmless but pointless).
Motivation: producing a SYCLBIN that participates in a runtime
sycl::link previously required passing two extra flags in lockstep
with -fsyclbin=. Missing either left users with a SYCLBIN that
loaded but failed at runtime ("No exported symbol \"X\"" at
sycl::link or "Unresolved Symbol \"X\"" inside
zeModuleDynamicLink). The state argument to -fsyclbin= already
declares user intent, so the additional flags are now derived from
it.
koparasy
force-pushed
the
jira/CMPLRLLVM-74442
branch
from
7596220 to
9a2feac
Compare
…or syclexp::compile
The user-visible API contract for syclexp::compile on an
ext_oneapi_source kernel_bundle is to produce a bundle in the
object state, which is intended to be passed to sycl::link for
cross-image symbol resolution. That requires the SYCL/exported
and SYCL/imported symbol property sets emitted by the RTC
sycl-post-link pass, which today is gated on the user passing
build_options{"-fsycl-allow-device-image-dependencies"} to
compile. Missing the flag leaves users with an object bundle that
silently drops cross-image metadata, and a subsequent sycl::link
fails with "No exported symbol \"X\" found in linked images." or
"Unresolved Symbol \"X\"" depending on the backend.
Imply the flag from the operation. createSYCLImages adds the flag
to the user options when the target state is not executable
(i.e. compile path). The implication is suppressed when the user
explicitly opts out via -fno-sycl-allow-device-image-dependencies
in build_options. syclex::build (executable target) is unchanged;
fully-linked images do not need the metadata.
Mirrors the producer-side coupling shipped for -fsyclbin=input/object
in the clang driver, so the same implication holds at both layers.
koparasy
force-pushed
the
jira/CMPLRLLVM-74442
branch
from
9a2feac to
258ee27
Compare
Contributor
Author
|
@dpcpp-specification-reviewers @dpcpp-doc-reviewers @dpcpp-clang-driver-reviewers can you please review this? |
srividya-sundaram
approved these changes
Jun 8, 2026
srividya-sundaram
left a comment
srividya-sundaram
left a comment
Contributor
There was a problem hiding this comment.
Driver changes LGTM, just some suggestions.
Co-authored-by: Srividya Sundaram <srividya.sundaram@intel.com>
mdtoguchi
reviewed
Jun 8, 2026
Co-authored-by: Michael Toguchi <michael.d.toguchi@intel.com>
Contributor
Author
|
@mdtoguchi any other comments? Or this good to go? |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Closes
CMPLRLLVM-74442. Addresses driver/runtime UX surface ofCMPLRLLVM-75983(runtime AOT fixes already landed upstream as #22196).
Three logical commits:
[SYCL][SYCLBIN] Fix kernel lookup after link of SYCLBIN + RTC bundles
sycl::link({RTC_obj, SYCLBIN_obj})merged image had origin bitsOR'd to (KernelCompiler|SYCLBIN). tryGetExtensionKernel routed all
lookups through RTC ProgramManager and missed SYCLBIN-origin
kernels. Add fall-through to SYCLBIN urKernelCreate path.
the link semantics, 3 supported configurations, and one
end-to-end example.
[SYCL][Driver] Imply cross-image symbol flags from -fsyclbin=input/object
spir64_gen, -library-compilation, when -fsyclbin=input/object.
is a hard error.
emits warning (new diag warn_drv_argument_has_no_effect_with):
"'-fsycl-allow-device-image-dependencies' has no effect with
'-fsyclbin=executable' [-Woption-ignored]"
[SYCL][KernelCompiler] Imply
-fsycl-allow-device-image-dependenciesfor syclexp::compile
(object state). syclex::build (executable target) unchanged.
E2E suite results: