From b6e91f42927081637ef4a313a840d8dca37508d4 Mon Sep 17 00:00:00 2001 From: Nico Weber Date: Thu, 25 May 2017 19:19:29 +0000 Subject: [PATCH] Revert r303859, CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll fails on bots. git-svn-id: https://llvm.org/svn/llvm-project/llvm/trunk@303902 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/llvm/IR/IntrinsicsAMDGPU.td | 10 ---------- lib/Target/AMDGPU/SOPInstructions.td | 4 +--- test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll | 15 --------------- 3 files changed, 1 insertion(+), 28 deletions(-) delete mode 100644 test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll diff --git a/include/llvm/IR/IntrinsicsAMDGPU.td b/include/llvm/IR/IntrinsicsAMDGPU.td index e1928546607..d7413fe9e56 100644 --- a/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/include/llvm/IR/IntrinsicsAMDGPU.td @@ -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 values: 0 = P10, 1 = P20, 2 = P0 def int_amdgcn_interp_mov : diff --git a/lib/Target/AMDGPU/SOPInstructions.td b/lib/Target/AMDGPU/SOPInstructions.td index ec29a66c8bb..f2d8b6f7b7a 100644 --- a/lib/Target/AMDGPU/SOPInstructions.td +++ b/lib/Target/AMDGPU/SOPInstructions.td @@ -184,9 +184,7 @@ 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", - [(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 { diff --git a/test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll b/test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll deleted file mode 100644 index c6491e6d1aa..00000000000 --- a/test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll +++ /dev/null @@ -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 } -- 2.11.0