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

[MoneroOcean]

Describe the bug

The HIP backend built from intel/llvm commit
eca4d07 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 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
  eca4d07, 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 c54fde4 / 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