<div dir="ltr"><div class="gmail_default" style="font-family:tahoma,sans-serif">Ref: <a href="https://github.com/AQuickBlackFox/amdgpu-code/blob/master/mul.ll" style="font-family:Arial,Helvetica,sans-serif">https://github.com/AQuickBlackFox/amdgpu-code/blob/master/mul.ll</a></div><div class="gmail_default" style="font-family:tahoma,sans-serif"><br></div><div class="gmail_default" style="font-family:tahoma,sans-serif">Replace "<span class="gmail-pl-c1" style="box-sizing:border-box;color:rgb(0,92,197);font-family:SFMono-Regular,Consolas,"Liberation Mono",Menlo,monospace;font-size:12px;white-space:pre">@llvm.amdgcn.workitem.id.x</span><span style="color:rgb(36,41,46);font-family:SFMono-Regular,Consolas,"Liberation Mono",Menlo,monospace;font-size:12px;white-space:pre">()</span>" with "<span class="gmail-pl-c1" style="box-sizing:border-box;color:rgb(0,92,197);font-family:SFMono-Regular,Consolas,"Liberation Mono",Menlo,monospace;font-size:12px;white-space:pre">@llvm.amdgcn.workgroup.id.x</span><span style="color:rgb(36,41,46);font-family:SFMono-Regular,Consolas,"Liberation Mono",Menlo,monospace;font-size:12px;white-space:pre">()</span>"</div></div><br><div class="gmail_quote"><div dir="ltr" class="gmail_attr">On Mon, Apr 13, 2020 at 10:19 AM Frank Winter via llvm-dev <<a href="mailto:llvm-dev@lists.llvm.org">llvm-dev@lists.llvm.org</a>> wrote:<br></div><blockquote class="gmail_quote" style="margin:0px 0px 0px 0.8ex;border-left:1px solid rgb(204,204,204);padding-left:1ex"><br>
On 4/13/20 11:43 AM, Matt Arsenault wrote:<br>
><br>
>> On Apr 13, 2020, at 11:37, Frank Winter via llvm-dev <<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a>> wrote:<br>
>><br>
>> Hi!<br>
>><br>
>> I'm trying to figure out how to access the workgroup id from within<br>
>> the LLVM IR language when lowering with the AMDGPU backend.<br>
>><br>
>> Looking at the 'llvm/include/llvm/IR/IntrinsicsAMDGPU.td' file there<br>
>> are intrinsics defined to access the workitem index (thread index),<br>
>> but this file lives in 'llvm/include':<br>
>><br>
>> //===----------------------------------------------------------------------===//<br>
>> // ABI Special Intrinsics<br>
>> //===----------------------------------------------------------------------===//<br>
>><br>
>> defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;<br>
>> defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named<br>
>> <"__builtin_amdgcn_workgroup_id">;<br>
>><br>
>> There is no new definition of any intrinsics within the target<br>
>> AMDGPU. I was working before with the NVPTX backend and that target<br>
>> has the special registers are associated with strings for the LLVM IR,<br>
>> e.g., PTX_READ_SREG_R32<"tid.x", int_nvvm_read_ptx_sreg_tid_x>;<br>
>><br>
>> Maybe I'm missing something but does this mean at the moment it is not<br>
>> possible to write a function in the LLVM IR language which accesses<br>
>> the workgroup id?<br>
>><br>
>> Frank<br>
>><br>
>><br>
> 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.<br>
><br>
> -Matt<br>
<br>
<br>
Thanks for the quick response!<br>
<br>
How do these definitions translate into LLVM IR? I tried a very simple<br>
function, e.g. simple_amd.ll:<br>
<br>
define i32 @simple() {<br>
entrypoint:<br>
   ret i32 %__builtin_amdgcn_workgroup_id_x<br>
}<br>
<br>
llc -march=amdgcn -mcpu=gfx906 < simple_amd.ll<br>
<br>
error: use of undefined value '%__builtin_amdgcn_workgroup_id_x'<br>
   ret i32 %__builtin_amdgcn_workgroup_id_x<br>
<br>
llc --version<br>
LLVM (<a href="http://llvm.org/" rel="noreferrer" target="_blank">http://llvm.org/</a>):<br>
   LLVM version 10.0.0<br>
   DEBUG build with assertions.<br>
   Default target: x86_64-unknown-linux-gnu<br>
   Host CPU: skylake-avx512<br>
<br>
   Registered Targets:<br>
     amdgcn - AMD GCN GPUs<br>
     r600   - AMD GPUs HD2XXX-HD6XXX<br>
<br>
<br>
Frank<br>
<br>
<br>
_______________________________________________<br>
LLVM Developers mailing list<br>
<a href="mailto:llvm-dev@lists.llvm.org" target="_blank">llvm-dev@lists.llvm.org</a><br>
<a href="https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev" rel="noreferrer" target="_blank">https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-dev</a><br>
</blockquote></div>