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

[Aditya Atluri via llvm-dev]

Ref: https://github.com/AQuickBlackFox/amdgpu-code/blob/master/mul.ll

Replace "@llvm.amdgcn.workitem.id.x()" with "@llvm.amdgcn.workgroup.id.x()"

On Mon, Apr 13, 2020 at 10:19 AM Frank Winter via llvm-dev <
llvm-dev at lists.llvm.org> wrote:

>
> 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
>
>
> _______________________________________________
> LLVM Developers mailing list
> llvm-dev at lists.llvm.org
> https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/llvm-dev/attachments/20200413/7c94600d/attachment.html>