DPC++ HIP nd_range kernel launch fails with hipErrorIllegalState on gfx1200

### Describe the bug

The HIP backend built from intel/llvm commit
eca4d070277a1e62b196a5fddefe72bc7f98ee24 against Ubuntu 26.04 ROCm 7.1
compiles the attached one-file program for gfx1200, but the first nd_range kernel
fails in urEnqueueKernelLaunch -> hipModuleLaunchKernel with HIP error 401
(hipErrorIllegalState). A single_task kernel built by the same compiler works.

This does not appear to be a general ROCm or hardware failure: the equivalent
nd_range program builds and runs on the same RX 9060 XT with AdaptiveCpp 25.10
and the same distro ROCm 7.1 packages.

### To reproduce

Reproduce in one command after installing the DPC++ HIP toolchain:

/opt/dpcpp/bin/clang++ -O3 -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1200 --rocm-device-lib-path=/opt/rocm-device-libs/amdgcn/bitcode repro.cpp -o repro && env LD_LIBRARY_PATH=/opt/dpcpp/lib 'ONEAPI_DEVICE_SELECTOR=hip:*' ./repro

Expected: prints 42.
Actual: throws sycl::exception; the HIP adapter reports error 401 while launching.

[repro.zip](https://github.com/user-attachments/files/29932599/repro.zip) contains only the minimized source.

### Environment

- Ubuntu 26.04 LTS, Linux 7.0.0-27-generic, x86_64
- AMD Radeon RX 9060 XT (gfx1200), amdgpu kernel driver
- ROCm/HIP 7.1 from Ubuntu packages
- intel/llvm nightly-2026-07-11, commit
eca4d070277a1e62b196a5fddefe72bc7f98ee24, Release build with --hip
- AOT target amdgcn-amd-amdhsa / gfx1200; ONEAPI_DEVICE_SELECTOR=hip:*

### Additional context

Investigation update / workaround (verified on the toolchain above):

add: -Xoffload-linker=amdgcn-amd-amdhsa '--lto-newpm-passes=globaloffset,lto'

This makes GlobalOffsetPass resolve __spirv_BuiltInGlobalOffset before the
AMDGPU full-LTO pipeline. Without it, AMDGPUAttributor sees that unresolved
callee and emits hostcall/default-queue/completion/heap hidden arguments; HIP
rejects the resulting launch. The workaround prints 42 and changes the two
kernel argument segments from 264/280 bytes to 8/20 bytes. A memory-bandwidth
control measured 299.4 GB/s versus 306.0 GB/s for native HIP (97.8%). The
regression appears after intel/llvm c54fde47 / PR #21385.

Checked before preparing this report:
- Intel LLVM and Unified Runtime issue searches contained no report with this
exact hipModuleLaunchKernel/401/gfx1200 failure.
- Code object versions 4, 5, and 6 fail identically.
- Work-group sizes 1 and 64 fail identically.
- Disabling SYCL dead-argument optimization does not help.
- A PCI reset and a complete amdgpu unbind/rebind do not help.
- AdaptiveCpp's HIP backend runs the equivalent nd_range correctly.

good-bad-dumps.zip contains the working single_task control, failing nd_range source, logs, and emitted code objects.

[good-bad-dumps.zip](https://github.com/user-attachments/files/29932602/good-bad-dumps.zip)

Start by running the attached repro.cpp with the one-command DPC++ HIP build and compare the failing nd_range launch with the working single_task control in good-bad-dumps.zip. Trace urEnqueueKernelLaunch into hipModuleLaunchKernel, then inspect GlobalOffsetPass and AMDGPUAttributor around the regression after PR #21385. Done means the nd_range program prints 42 without the linker workaround.

Written by the indexing model from the issue text.