<div dir="ltr">Now <a href="http://lab.llvm.org:8011/builders/clang-x86_64-debian-fast/builds/4222/steps/test/logs/stdio">http://lab.llvm.org:8011/builders/clang-x86_64-debian-fast/builds/4222/steps/test/logs/stdio</a> is not happy about it.<br></div><div class="gmail_extra"><br><div class="gmail_quote">On Thu, May 25, 2017 at 12:19 PM, Nico Weber via llvm-commits <span dir="ltr"><<a href="mailto:llvm-commits@lists.llvm.org" target="_blank">llvm-commits@lists.llvm.org</a>></span> wrote:<br><blockquote class="gmail_quote" style="margin:0 0 0 .8ex;border-left:1px #ccc solid;padding-left:1ex">Author: nico<br>
Date: Thu May 25 14:19:29 2017<br>
New Revision: 303902<br>
<br>
URL: <a href="http://llvm.org/viewvc/llvm-project?rev=303902&view=rev" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-<wbr>project?rev=303902&view=rev</a><br>
Log:<br>
Revert r303859, CodeGen/AMDGPU/llvm.amdgcn.s.<wbr>getpc.ll fails on bots.<br>
<br>
Removed:<br>
    llvm/trunk/test/CodeGen/<wbr>AMDGPU/llvm.amdgcn.s.getpc.ll<br>
Modified:<br>
    llvm/trunk/include/llvm/IR/<wbr>IntrinsicsAMDGPU.td<br>
    llvm/trunk/lib/Target/AMDGPU/<wbr>SOPInstructions.td<br>
<br>
Modified: llvm/trunk/include/llvm/IR/<wbr>IntrinsicsAMDGPU.td<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=303902&r1=303901&r2=303902&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-<wbr>project/llvm/trunk/include/<wbr>llvm/IR/IntrinsicsAMDGPU.td?<wbr>rev=303902&r1=303901&r2=<wbr>303902&view=diff</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- llvm/trunk/include/llvm/IR/<wbr>IntrinsicsAMDGPU.td (original)<br>
+++ llvm/trunk/include/llvm/IR/<wbr>IntrinsicsAMDGPU.td Thu May 25 14:19:29 2017<br>
@@ -566,16 +566,6 @@ def int_amdgcn_s_getreg :<br>
   [IntrReadMem, IntrSpeculatable]<br>
 >;<br>
<br>
-// int_amdgcn_s_getpc is provided to allow a specific style of position<br>
-// independent code to determine the high part of its address when it is<br>
-// known (through convention) that the code and any data of interest does<br>
-// not cross a 4Gb address boundary. Use for any other purpose may not<br>
-// produce the desired results as optimizations may cause code movement,<br>
-// especially as we explicitly use IntrNoMem to allow optimizations.<br>
-def int_amdgcn_s_getpc :<br>
-  GCCBuiltin<"__builtin_amdgcn_<wbr>s_getpc">,<br>
-  Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>;<br>
-<br>
 // __builtin_amdgcn_interp_mov <param>, <attr_chan>, <attr>, <m0><br>
 // param values: 0 = P10, 1 = P20, 2 = P0<br>
 def int_amdgcn_interp_mov :<br>
<br>
Modified: llvm/trunk/lib/Target/AMDGPU/<wbr>SOPInstructions.td<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SOPInstructions.td?rev=303902&r1=303901&r2=303902&view=diff" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-<wbr>project/llvm/trunk/lib/Target/<wbr>AMDGPU/SOPInstructions.td?rev=<wbr>303902&r1=303901&r2=303902&<wbr>view=diff</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- llvm/trunk/lib/Target/AMDGPU/<wbr>SOPInstructions.td (original)<br>
+++ llvm/trunk/lib/Target/AMDGPU/<wbr>SOPInstructions.td Thu May 25 14:19:29 2017<br>
@@ -184,9 +184,7 @@ def S_BITSET0_B32 : SOP1_32    <"s_bitse<br>
 def S_BITSET0_B64 : SOP1_64_32 <"s_bitset0_b64">;<br>
 def S_BITSET1_B32 : SOP1_32    <"s_bitset1_b32">;<br>
 def S_BITSET1_B64 : SOP1_64_32 <"s_bitset1_b64">;<br>
-def S_GETPC_B64 : SOP1_64_0  <"s_getpc_b64",<br>
-  [(set i64:$sdst, (int_amdgcn_s_getpc))]<br>
->;<br>
+def S_GETPC_B64 : SOP1_64_0  <"s_getpc_b64">;<br>
<br>
 let isTerminator = 1, isBarrier = 1, SchedRW = [WriteBranch] in {<br>
<br>
<br>
Removed: llvm/trunk/test/CodeGen/<wbr>AMDGPU/llvm.amdgcn.s.getpc.ll<br>
URL: <a href="http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.getpc.ll?rev=303901&view=auto" rel="noreferrer" target="_blank">http://llvm.org/viewvc/llvm-<wbr>project/llvm/trunk/test/<wbr>CodeGen/AMDGPU/llvm.amdgcn.s.<wbr>getpc.ll?rev=303901&view=auto</a><br>
==============================<wbr>==============================<wbr>==================<br>
--- llvm/trunk/test/CodeGen/<wbr>AMDGPU/llvm.amdgcn.s.getpc.ll (original)<br>
+++ llvm/trunk/test/CodeGen/<wbr>AMDGPU/llvm.amdgcn.s.getpc.ll (removed)<br>
@@ -1,15 +0,0 @@<br>
-; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s<br>
-<br>
-declare i64 @llvm.amdgcn.s.getpc() #0<br>
-<br>
-; GCN-LABEL: {{^}}test_s_getpc:<br>
-; GCN: s_load_dwordx2<br>
-; GCN-DAG: s_getpc_b64 s{{\[[0-9]+:[0-9]+\]}}<br>
-; GCN: buffer_store_dwordx2<br>
-define void @test_s_getpc(i64 addrspace(1)* %out) #0 {<br>
-  %tmp = call i64 @llvm.amdgcn.s.getpc() #1<br>
-  store volatile i64 %tmp, i64 addrspace(1)* %out, align 8<br>
-  ret void<br>
-}<br>
-<br>
-attributes #0 = { nounwind readnone speculatable }<br>
<br>
<br>
______________________________<wbr>_________________<br>
llvm-commits mailing list<br>
<a href="mailto:llvm-commits@lists.llvm.org">llvm-commits@lists.llvm.org</a><br>
<a href="http://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-commits" rel="noreferrer" target="_blank">http://lists.llvm.org/cgi-bin/<wbr>mailman/listinfo/llvm-commits</a><br>
</blockquote></div><br></div>