[llvm] r303902 - Revert r303859, CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll fails on bots.

[Nico Weber via llvm-commits]

Author: nico
Date: Thu May 25 14:19:29 2017
New Revision: 303902

URL: http://llvm.org/viewvc/llvm-project?rev=303902&view=rev
Log:
Revert r303859, CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll fails on bots.

Removed:
    llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll
Modified:
    llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
    llvm/trunk/lib/Target/AMDGPU/SOPInstructions.td

Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=303902&r1=303901&r2=303902&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Thu May 25 14:19:29 2017
@@ -566,16 +566,6 @@ def int_amdgcn_s_getreg :
   [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 :

Modified: llvm/trunk/lib/Target/AMDGPU/SOPInstructions.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SOPInstructions.td?rev=303902&r1=303901&r2=303902&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SOPInstructions.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/SOPInstructions.td Thu May 25 14:19:29 2017
@@ -184,9 +184,7 @@ def S_BITSET0_B32 : SOP1_32    <"s_bitse
 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",
-  [(set i64:$sdst, (int_amdgcn_s_getpc))]
->;
+def S_GETPC_B64 : SOP1_64_0  <"s_getpc_b64">;
 
 let isTerminator = 1, isBarrier = 1, SchedRW = [WriteBranch] in {
 

Removed: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll?rev=303901&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll (original)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll (removed)
@@ -1,15 +0,0 @@
-; 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 }