llvm
d93715de - [SYCL][SYCLBIN] Fix native-AOT object-state cross-library link (#22509)

Commit
14 hours ago
[SYCL][SYCLBIN] Fix native-AOT object-state cross-library link (#22509) ## Fix native-AOT object-state cross-library link Loading an AOT-only object-state SYCLBIN and linking it via `sycl::link` did not work end to end. #22196 enabled the importer-side object load (and shipped a load-only e2e), but a full cross-library link of native AOT objects hit three defects, none covered by that load-only test. ### Bugs fixed 1. **Export-only AOT library not loadable as `object`.** `SYCLBINBinaries::getBestCompatibleImages` gated the native-image fallback on `getBinImageState(img) == State`. An export-only library imports nothing, so it classifies as `executable` and was dropped for an `object` request . 2. **Assertion on loaded image state.** The SYCLBIN loader stamped the image with its intrinsic state (`executable`), tripping `DevImageState == bundle_state::object` in `bringSYCLDeviceImageToState`. The created image is now clamped to the requested (already-validated) declared state. 3. **Segfault on AOT-only link.** `kernel_bundle_impl::link` always invoked `ProgramManager::link` on the JIT image set; for an AOT-only link that set is empty, so `urProgramLinkExp` was called with zero programs and segfaulted in the L0 adapter. The JIT link is now skipped when there are no JIT images; AOT images are resolved via `dynamicLink`. ### Tests - **e2e** `SYCLBIN/link_object_aot.cpp` — cross-library link of two AOT-only object SYCLBINs on Level Zero (spir64_gen). Exercises the AOT partition, `ProgramManager::build` with `AllowUnresolvedSymbols`, and `dynamicLink`. Restricted to L0 (`REQUIRES: level_zero, ocloc`); - **unit** `AOTBinaryTarget.cpp` — `isAOTBinaryTarget` and `getBinImageState` classifiers. - **unit** `DynamicLinking.cpp` — `build` with `AllowUnresolvedSymbols` (happy path + program-cache-key bool), and `feature_not_supported` guards on the build and link entry points when the adapter lacks the Exp variants. - Fixed `RuntimeLinkingCommon::clear()` to reset the CreateWithBinary counter (latent; exposed under `--gtest_shuffle`).
Author
Parents
Loading