[SYCL][Driver] Make -fno-sycl-rdc objects self-contained

[dvrogozh]

This commit is entirely authored by Claude. I provide it as a draft PR in the hope it can be used as a fix or will help to narrow the root cause of the issue. It requires a review from compiler team expert. See #23097

An object compiled with -fsycl -fno-sycl-rdc -c could not be linked into a program by anything but the SYCL-aware driver, even though non-RDC means the device code of a translation unit is final at compile time. Linking such an object with a plain host linker produced a binary that aborted at the first kernel launch:

  $ clang++ -fsycl -fno-gpu-rdc -fPIE -c test.cpp -o test.o
  $ g++ test.o $(pkg-config --libs sycl-dpcpp-7) -o test.gcc
  $ ./test.gcc
  Assertion `It != m_DeviceKernelInfoMap.end()' failed.

Two things were in the way.

First, the compile-time device link was gated on Triple::isSPIRAOT(), so for the default JIT spir64 target -fno-sycl-rdc -c had no effect at all: the object carried plain device IR in a __CLANG_OFFLOAD_BUNDLE__ section and needed the final link to finish the device pipeline. Gate it on SYCL::hasFinalDeviceImage() instead, which accepts every SPIR-V target. NVPTX, AMDGCN and NativeCPU still need the device link to run as part of the host link, so they keep the old behaviour.

Second, even for AOT the finalized image was hidden in an offload bundle section. Bundles are created readonly,exclude (SHF_EXCLUDE), which is exactly what makes a host linker drop them, so no device image was registered and the kernel info map stayed empty. Merge the wrapped device object into the host object with a partial link (<linker> -r) instead of bundling it, via a new PartialLinkJobAction and a matching tools::PartialLink. The result is an ordinary relocatable object whose __sycl_register_lib constructor registers the image on its own, so any host linker will do. MSVC provides no -r, so Windows keeps bundling.

The wrapper names the section holding the image data after the offload bundle convention, so with the image now living in the host object the bundler would mistake such an object for a fat one. Require SHF_EXCLUDE in ObjectFileHandler::IsOffloadSection to tell the two apart.

Link-time handling of the _image suffixed bundle target name is kept - Windows objects and objects from older compilers still use it - and the list of architecture names it looks for is now shared with the driver through SYCL::getFinalDeviceImageArchNames().

Finally, replace the two asserts in ProgramManager::getDeviceKernelInfo() with a sycl::exception: a kernel missing from the map means no device image provided it, which is a build problem to be reported rather than an internal invariant to abort on (and It->second on end() is undefined behavior in an NDEBUG build).

Co-Authored-By: Claude Opus 5 noreply@anthropic.com

CC: @tahonermann, @YuriPlyakhin, @srividya-sundaram