diff options
| -rw-r--r-- | llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 4 | ||||
| -rw-r--r-- | llvm/lib/Target/AMDGPU/SIInstructions.td | 12 | ||||
| -rw-r--r-- | llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll | 16 |
3 files changed, 32 insertions, 0 deletions
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index b8479a58e6d..35a86aceb55 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -231,6 +231,10 @@ def int_amdgcn_s_sleep : Intrinsic<[], [llvm_i32_ty], []> { } +def int_amdgcn_s_getreg : + GCCBuiltin<"__builtin_amdgcn_s_getreg">, + Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrReadMem]>; + def int_amdgcn_dispatch_ptr : GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">, Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>; diff --git a/llvm/lib/Target/AMDGPU/SIInstructions.td b/llvm/lib/Target/AMDGPU/SIInstructions.td index 7602498c274..edc60925a3e 100644 --- a/llvm/lib/Target/AMDGPU/SIInstructions.td +++ b/llvm/lib/Target/AMDGPU/SIInstructions.td @@ -416,7 +416,11 @@ defm S_CBRANCH_I_FORK : SOPK_m < sopk<0x11, 0x10>, "s_cbranch_i_fork", (outs), (ins SReg_64:$sdst, u16imm:$simm16), " $sdst, $simm16" >; + +let mayLoad = 1 in { defm S_GETREG_B32 : SOPK_32 <sopk<0x12, 0x11>, "s_getreg_b32", []>; +} + defm S_SETREG_B32 : SOPK_m < sopk<0x13, 0x12>, "s_setreg_b32", (outs), (ins SReg_32:$sdst, u16imm:$simm16), " $sdst, $simm16" @@ -2090,6 +2094,14 @@ def : Pat < >; //===----------------------------------------------------------------------===// +// S_GETREG_B32 Intrinsic Pattern. +//===----------------------------------------------------------------------===// +def : Pat < + (int_amdgcn_s_getreg imm:$simm16), + (S_GETREG_B32 (as_i16imm $simm16)) +>; + +//===----------------------------------------------------------------------===// // SMRD Patterns //===----------------------------------------------------------------------===// diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll new file mode 100644 index 00000000000..d903b00e393 --- /dev/null +++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll @@ -0,0 +1,16 @@ +; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck %s +; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=kaveri -verify-machineinstrs < %s | FileCheck %s +; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck %s + +; FUNC-LABEL: {{^}}s_getreg_test: +; CHECK: s_getreg_b32 s{{[0-9]+}}, 0xb206 +define void @s_getreg_test(i32 addrspace(1)* %out) { ; simm16=45574 for lds size. + %lds_size_64dwords = call i32 @llvm.amdgcn.s.getreg(i32 45574) #0 + %lds_size_bytes = shl i32 %lds_size_64dwords, 8 + store i32 %lds_size_bytes, i32 addrspace(1)* %out + ret void +} + +declare i32 @llvm.amdgcn.s.getreg(i32) #0 + +attributes #0 = { nounwind readonly} |
