[AMDGPU] add intrinsic for s_getpc

[timcorringham]

Summary

The s_getpc instruction is exposed as intrinsic llvm.amdgcn.s.getpc.

[arsenm]

What are you using this for? I don't see how an unpredictable address like this is useful

[timcorringham]

In D32862#746257, @arsenm wrote:

What are you using this for? I don't see how an unpredictable address like this is useful

The particular use-case this is required for only uses the top part of the address that is returned, so the imprecision isn't a problem. The onus is on the programmer to use barriers if a more precise value is required.

[arsenm]

In D32862#746896, @timcorringham wrote:

In D32862#746257, @arsenm wrote:

What are you using this for? I don't see how an unpredictable address like this is useful

The particular use-case this is required for only uses the top part of the address that is returned, so the imprecision isn't a problem. The onus is on the programmer to use barriers if a more precise value is required.

Even then that's a pretty big assumption relying on the high 32-bits. What are you doing with the address? This might be better served by something more targeted

[timcorringham]

Even then that's a pretty big assumption relying on the high 32-bits. What are you doing with the address? This might be better served by something more targeted

The use-case is for code that follows a convention that ensures the high address is always the value that is required. This intrinsic is unlikely to be widely used, but is probably more useful than an intrinsic that just returns the high address - although I accept that is debatable!

[nhaehnle]

Low & high bits and use cases aside, isn't this something that could be covered by adding support for PC in the read_register intrinsic?

[timcorringham]

In D32862#748486, @nhaehnle wrote:

Low & high bits and use cases aside, isn't this something that could be covered by adding support for PC in the read_register intrinsic?

I'm sure that is possible, but I'm put off by the LLVM LangRef description of read_register which includes the following:

Warning: So far it only works with the stack pointer on selected architectures (ARM, AArch64, PowerPC and x86_64). Significant amount of work is needed to support other registers...

Apart from that warning, the advantage of s_getpc is that it just exposes an instruction which has some documentation and known behaviour, while adding new behaviour to read_register would be more obscure.

[arsenm]

read/write register work for AMDGPU, though there is the hazard you can easily create impossible cases of VGPR->SGPR copies we don't attempt to handle

[arsenm]

read/write register would be inappropriate for AMDGPU because we don't actually have an addressable PC register, and it's accessed through an instruction like this

[timcorringham]

[AMDGPU] add intrinsic for s_getpc

The s_getpc instruction is exposed as intrinsic llvm.amdgcn.s.getpc

Updated comment to clarify that it is provided to allow a specific style
of position independent code to determine the high part of its address when
it is known (through convention) that the code and any data of interest does
not cross a 4Gb address boundary.

[arsenm]

include/llvm/IR/IntrinsicsAMDGPU.td

576

Do you have the patch for the clang component of this?

577

You could also debatably make it speculatable

lib/Target/AMDGPU/SOPInstructions.td

941–943

You should be able to add this directly to the instruction definition without a standalone pattern. The comment is also unnecessary

test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll

16

Not that it matters, but this doesn't match the readnone it was declared with

[timcorringham]

Amendments addressing review comments:

Corresponding Clang changes are in D33276.

I have made the intrinsic speculatable.

I have moved the intrinsic pattern to be next to the instruction def, but I'm too much of a Tablegen novice to be able to roll them in together.

Corrected attributes in test.

[arsenm]

In D32862#757385, @timcorringham wrote:

Amendments addressing review comments:

Corresponding Clang changes are in D33276.

I have made the intrinsic speculatable.

I have moved the intrinsic pattern to be next to the instruction def, but I'm too much of a Tablegen novice to be able to roll them in together.

Corrected attributes in test.

See SI_PS_LIVE, you just need something like [(set i64:$sdst, (int_foo))]

[timcorringham]

Updated s_getpc instruction definition to include intrinsic.

[arsenm]

LGTM

[thakis]

http://lab.llvm.org:8011/builders/llvm-clang-x86_64-expensive-checks-win/builds/2604
http://lab.llvm.org:8011/builders/clang-with-thin-lto-ubuntu/builds/3280

$ svn commit -m 'Revert r303859, CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll fails on bots.'
Sending include/llvm/IR/IntrinsicsAMDGPU.td
Sending lib/Target/AMDGPU/SOPInstructions.td
Deleting test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll
Transmitting file data ..
Committed revision 303902.

[kzhuravl]

llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll

9 ↗

(On Diff #100242)

@timcorringham, Add amdgpu_kernel convention.

include/llvm/IR/IntrinsicsAMDGPU.td

 }
 
 def int_amdgcn_s_getreg :
   GCCBuiltin<"__builtin_amdgcn_s_getreg">,
   Intrinsic<[llvm_i32_ty], [llvm_i32_ty],
   [IntrReadMem, IntrSpeculatable]
 >;
 
+// int_amdgcn_s_getpc is provided to allow a specific style of position
+// independent code to determine the high part of its address when it is
+// known (through convention) that the code and any data of interest does
+// not cross a 4Gb address boundary. Use for any other purpose may not
+// produce the desired results as optimizations may cause code movement,
+// especially as we explicitly use IntrNoMem to allow optimizations.
+def int_amdgcn_s_getpc :
+  GCCBuiltin<"__builtin_amdgcn_s_getpc">,
+  Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>;
+
 // __builtin_amdgcn_interp_mov <param>, <attr_chan>, <attr>, <m0>
 // param values: 0 = P10, 1 = P20, 2 = P0
 def int_amdgcn_interp_mov :
   GCCBuiltin<"__builtin_amdgcn_interp_mov">,
   Intrinsic<[llvm_float_ty],
             [llvm_i32_ty, llvm_i32_ty, llvm_i32_ty, llvm_i32_ty],
             [IntrNoMem, IntrSpeculatable]>;
 

lib/Target/AMDGPU/SOPInstructions.td

def S_SEXT_I32_I16 : SOP1_32 <"s_sext_i32_i16",
   [(set i32:$sdst, (sext_inreg i32:$src0, i16))]
 >;
 
 def S_BITSET0_B32 : SOP1_32    <"s_bitset0_b32">;
 def S_BITSET0_B64 : SOP1_64_32 <"s_bitset0_b64">;
 def S_BITSET1_B32 : SOP1_32    <"s_bitset1_b32">;
 def S_BITSET1_B64 : SOP1_64_32 <"s_bitset1_b64">;
 def S_GETPC_B64 : SOP1_64_0  <"s_getpc_b64">;
+def : Pat <
+  (int_amdgcn_s_getpc),
+  (S_GETPC_B64)
+>;
+
 
 let isTerminator = 1, isBarrier = 1,
     isBranch = 1, isIndirectBranch = 1 in {
 def S_SETPC_B64 : SOP1_1  <"s_setpc_b64">;
 }
 def S_SWAPPC_B64 : SOP1_64 <"s_swappc_b64">;
 def S_RFE_B64 : SOP1_1  <"s_rfe_b64">;
 

 }
 }
 
 let Predicates = [isGCN] in {
 
 //===----------------------------------------------------------------------===//
 // S_GETREG_B32 Intrinsic Pattern.
 //===----------------------------------------------------------------------===//
 def : Pat <
   (int_amdgcn_s_getreg imm:$simm16),
   (S_GETREG_B32 (as_i16imm $simm16))
 >;
 
 //===----------------------------------------------------------------------===//
 // SOP1 Patterns
 //===----------------------------------------------------------------------===//
 
 def : Pat <
   (i64 (ctpop i64:$src)),

test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll

+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+
+declare i64 @llvm.amdgcn.s.getpc() #0
+
+; GCN-LABEL: {{^}}test_s_getpc:
+; GCN: s_load_dwordx2
+; GCN-DAG: s_getpc_b64 s{{\[[0-9]+:[0-9]+\]}}
+; GCN: buffer_store_dwordx2
+define void @test_s_getpc(i64 addrspace(1)* %out) #0 {
+  %tmp = call i64 @llvm.amdgcn.s.getpc() #1
+  store volatile i64 %tmp, i64 addrspace(1)* %out, align 8
+  ret void
+}
+
+attributes #0 = { nounwind readnone speculatable }