llvm dev - Apr 2020 - [EXTERNAL] Re: Are AMDGPU intrinsics available in LLVM IR ?

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

> On Apr 13, 2020, at 11:37, Frank Winter via llvm-dev <llvm-dev at
lists.llvm.org> wrote:
> 
> 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
> 
> 
These are the definitions, and the standard place for targets to define
intrinsics. NVPTX is odd in defining intrinsics in the backend. These are usable
like any other IR intrinsic, through the generated
Intrinsic::amdgcn_workitem_id_* enums.

-Matt

On 4/13/20 11:43 AM, Matt Arsenault wrote:
>
>> On Apr 13, 2020, at 11:37, Frank Winter via llvm-dev <llvm-dev at
lists.llvm.org> wrote:
>>
>> 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
>>
>>
> These are the definitions, and the standard place for targets to define
intrinsics. NVPTX is odd in defining intrinsics in the backend. These are usable
like any other IR intrinsic, through the generated
Intrinsic::amdgcn_workitem_id_* enums.
>
> -Matt

Thanks for the quick response!

How do these definitions translate into LLVM IR? I tried a very simple
function, e.g. simple_amd.ll:

define i32 @simple() {
entrypoint:
   ret i32 %__builtin_amdgcn_workgroup_id_x
}

llc -march=amdgcn -mcpu=gfx906 < simple_amd.ll

error: use of undefined value '%__builtin_amdgcn_workgroup_id_x'
   ret i32 %__builtin_amdgcn_workgroup_id_x

llc --version
LLVM (http://llvm.org/):
   LLVM version 10.0.0
   DEBUG build with assertions.
   Default target: x86_64-unknown-linux-gnu
   Host CPU: skylake-avx512

   Registered Targets:
     amdgcn - AMD GCN GPUs
     r600   - AMD GPUs HD2XXX-HD6XXX


Frank

On Mon, Apr 13, 2020 at 8:43 AM Matt Arsenault via llvm-dev <
llvm-dev at lists.llvm.org> wrote:
> These are the definitions, and the standard place for targets to define
> intrinsics. NVPTX is odd in defining intrinsics in the backend. These are
> usable like any other IR intrinsic, through the generated
> Intrinsic::amdgcn_workitem_id_* enums.
>
This is slightly off topic, but I wanted to point out that intrinsics are
now generated into separate headers (IntrinsicsARM.h, IntrinsicsAMDGPU.h,
...). This opens the door to moving the intrinsics into the backends and
compiling them out when a particular target is disabled. There is obviously
a lot more work to do since there are many mid-level IR transforms that
depend on these headers, but we could sink that logic into targets if we
were motivated to do it.
-------------- next part --------------
An HTML attachment was scrubbed...
URL:
<http://lists.llvm.org/pipermail/llvm-dev/attachments/20200414/88ec0516/attachment.html>