[llvm-dev] Are AMDGPU intrinsics available in LLVM IR ?

[Frank Winter via llvm-dev]

Hi!

I'm trying to figure out how to access the workgroup id from within
the LLVM IR language when lowering with the AMDGPU backend.

Looking at the 'llvm/include/llvm/IR/IntrinsicsAMDGPU.td' file there
are intrinsics defined to access the workitem index (thread index),
but this file lives in 'llvm/include':

//===----------------------------------------------------------------------===//
// ABI Special Intrinsics
//===----------------------------------------------------------------------===//

defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;
defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
<"__builtin_amdgcn_workgroup_id">;

There is no new definition of any intrinsics within the target
AMDGPU. I was working before with the NVPTX backend and that target
has the special registers are associated with strings for the LLVM IR,
e.g., PTX_READ_SREG_R32<"tid.x", int_nvvm_read_ptx_sreg_tid_x>;

Maybe I'm missing something but does this mean at the moment it is not
possible to write a function in the LLVM IR language which accesses
the workgroup id?

Frank