Reverse translation drops addrspace(4) on @llvm.amdgcn.implicitarg.ptr → invalid IR / codegen crash (amdgcnspirv blockDim)

[lamb-j]

Summary

When reverse-translating (-r, SPIR-V → LLVM) a module that reads a HIP block dimension
(blockDim.x, which the HIP headers now lower to __builtin_amdgcn_workgroup_size_x()),
amd-llvm-spirv emits the AMDGPU intrinsic @llvm.amdgcn.implicitarg.ptr with return type
ptr instead of the required ptr addrspace(4) (and the derived GEP/load chain in the wrong
address space). The result is invalid LLVM IR:

Fails to verify module: intrinsic return type expected ptr addrspace(4), but got ptr
declare ptr @llvm.amdgcn.implicitarg.ptr()

The verifier rejects it; and when the verifier is bypassed (as Comgr's in-process JIT does),
the AMDGPU backend segfaults during codegen.

Impact

This breaks the amdgcnspirv offload path for essentially every HIP kernel that reads a
block dimension (blockDim.{x,y,z}) — i.e. nearly all real compute kernels. At the first
hipLaunchKernel, HIP asks Comgr to JIT the SPIR-V code object
(AMD_COMGR_ACTION_COMPILE_SPIRV_TO_RELOCATABLE); the invalid IR crashes codegen and the
process dies with SIGSEGV.

Concretely, this took the rocm-examples job of the SPIRV Compiler CI (on both
ROCm/SPIRV-LLVM-Translator and ROCm/llvm-project amd-staging) red — 16 of 23 HIP-Basic
examples SIGSEGV (saxpy, matrix_multiplication, shared_memory, …), while kernels that never
read a block dim (e.g. hello_world) pass.

Reproducer

Minimal HIP kernel (any recent ROCm whose HIP headers lower blockDim via
__builtin_amdgcn_workgroup_size_x):

// bd.hip
#include <hip/hip_runtime.h>
__global__ void k(int* p){ p[0] = blockDim.x; }

clang++ -x hip --offload-arch=amdgcnspirv --rocm-path=$ROCM -c bd.hip -o bd.o
# extract the embedded SPIR-V code object (e.g. via clang-offload-bundler, or grab
# hip_code_object.spv from a run with AMD_COMGR_SAVE_TEMPS=1), then:
amd-llvm-spirv -r --spirv-target-env=CL2.0 hip_code_object.spv -o out.bc

Observed (straight from amd-llvm-spirv):

Fails to verify module: intrinsic return type expected ptr addrspace(4), but got ptr
declare ptr @llvm.amdgcn.implicitarg.ptr()

The offending function in the reverse-translated IR:

declare ptr @llvm.amdgcn.implicitarg.ptr()          ; <-- should be: ptr addrspace(4)

define internal i32 @_ZL21__hip_get_block_dim_xv() {
entry:
  %0 = call ptr @llvm.amdgcn.implicitarg.ptr()
  %1 = getelementptr inbounds i8, ptr %0, i64 12
  %2 = load i16, ptr %1, align 2
  %3 = zext i16 %2 to i32
  ret i32 %3
}

Suggested fix

Reverse translation should declare and use the intrinsic in the constant address space:

declare ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
...
  %0 = call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
  %1 = getelementptr inbounds i8, ptr addrspace(4) %0, i64 12
  %2 = load i16, ptr addrspace(4) %1, align 2

I verified that hand-patching the reverse-translated IR to addrspace(4) (as above) makes the
AMDGPU backend compile the module to a relocatable cleanly — so this addrspace is the whole fix.

Where it likely lives

The defect is in the reader / reverse-translation path (lib/SPIRV/SPIRVReader.cpp) that
reconstructs the AMDGPU intrinsic signature. This looks like the mirror of the forward-side fix
in #246 ("Fix return type and argument types for AMD intrinsics", lib/SPIRV/SPIRVWriter.cpp) —
the return-type address space needs the same treatment on the reader side. The existing
test/llvm-intrinsics/amdgcn-intrinsic-addrspace.ll /
test/llvm-intrinsics/amdgcn-intrinsic-ptr-arg-addrspace.ll tests look like the right place to
add coverage for the implicitarg.ptr return type.

Context / trigger

The translator has mishandled this signature independent of any recent translator change (the
CI's translator SHA was identical on the last-good and first-bad builds). What newly exercised
it was a HIP header change in ROCm/rocm-systems (SWDEV-548892) that switched
__hip_get_block_dim_x/y/z from __ockl_get_local_size(n) to
__builtin_amdgcn_workgroup_size_{x,y,z}(); the builtin lowers to implicitarg.ptr, which the
round-trip then mistranslates. (Grid dims still use __ockl_get_num_groups, so they are
unaffected.) The header change is a legitimate simplification — the fix belongs here, in the
translator.

Environment

• amd-llvm-spirv / translator: amd-staging
• clang: 24.0.0git (ROCm/llvm-project amd-staging @ 632dcda68437)
• Comgr 3.5.0, HIP runtime 7.16, ROCr 1.21
• Reproduced on gfx1030 locally and gfx942 in CI (address-space defect is arch-independent)

The full reverse-translated IR and the SPIR-V input for the minimal blockDim kernel are
available on request.