[llvm] r262120 - AMDGPU: Add s_sleep intrinsic

[Matt Arsenault via llvm-commits]

Author: arsenm
Date: Sat Feb 27 02:53:52 2016
New Revision: 262120

URL: http://llvm.org/viewvc/llvm-project?rev=262120&view=rev
Log:
AMDGPU: Add s_sleep intrinsic

Added:
    llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.sleep.ll
Modified:
    llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
    llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td
    llvm/trunk/lib/Target/AMDGPU/SIInstructions.td

Modified: llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td?rev=262120&r1=262119&r2=262120&view=diff
==============================================================================
--- llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td (original)
+++ llvm/trunk/include/llvm/IR/IntrinsicsAMDGPU.td Sat Feb 27 02:53:52 2016
@@ -192,6 +192,11 @@ def int_amdgcn_s_memtime :
   GCCBuiltin<"__builtin_amdgcn_s_memtime">,
   Intrinsic<[llvm_i64_ty], [], []>;
 
+def int_amdgcn_s_sleep :
+  GCCBuiltin<"__builtin_amdgcn_s_sleep">,
+  Intrinsic<[], [llvm_i32_ty], []> {
+}
+
 def int_amdgcn_dispatch_ptr :
   GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">,
   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 2>], [], [IntrNoMem]>;

Modified: llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td?rev=262120&r1=262119&r2=262120&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstrInfo.td Sat Feb 27 02:53:52 2016
@@ -371,6 +371,10 @@ def IMM16bit : PatLeaf <(imm),
   [{return isUInt<16>(N->getZExtValue());}]
 >;
 
+def SIMM16bit : PatLeaf <(imm),
+  [{return isInt<16>(N->getSExtValue());}]
+>;
+
 def IMM20bit : PatLeaf <(imm),
   [{return isUInt<20>(N->getZExtValue());}]
 >;

Modified: llvm/trunk/lib/Target/AMDGPU/SIInstructions.td
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/AMDGPU/SIInstructions.td?rev=262120&r1=262119&r2=262120&view=diff
==============================================================================
--- llvm/trunk/lib/Target/AMDGPU/SIInstructions.td (original)
+++ llvm/trunk/lib/Target/AMDGPU/SIInstructions.td Sat Feb 27 02:53:52 2016
@@ -501,10 +501,22 @@ def S_BARRIER : SOPP <0x0000000a, (ins),
 
 def S_WAITCNT : SOPP <0x0000000c, (ins WAIT_FLAG:$simm16), "s_waitcnt $simm16">;
 def S_SETHALT : SOPP <0x0000000d, (ins i16imm:$simm16), "s_sethalt $simm16">;
-def S_SLEEP : SOPP <0x0000000e, (ins i16imm:$simm16), "s_sleep $simm16">;
+
+// On SI the documentation says sleep for approximately 64 * low 2
+// bits, consistent with the reported maximum of 448. On VI the
+// maximum reported is 960 cycles, so 960 / 64 = 15 max, so is the
+// maximum really 15 on VI?
+def S_SLEEP : SOPP <0x0000000e, (ins i32imm:$simm16),
+  "s_sleep $simm16", [(int_amdgcn_s_sleep SIMM16bit:$simm16)]> {
+  let hasSideEffects = 1;
+  let mayLoad = 1;
+  let mayStore = 1;
+}
+
 def S_SETPRIO : SOPP <0x0000000f, (ins i16imm:$sim16), "s_setprio $sim16">;
 
 let Uses = [EXEC, M0] in {
+  // FIXME: Should this be mayLoad+mayStore?
   def S_SENDMSG : SOPP <0x00000010, (ins SendMsgImm:$simm16), "s_sendmsg $simm16",
       [(AMDGPUsendmsg (i32 imm:$simm16))]
   >;

Added: llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.sleep.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.sleep.ll?rev=262120&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.sleep.ll (added)
+++ llvm/trunk/test/CodeGen/AMDGPU/llvm.amdgcn.s.sleep.ll Sat Feb 27 02:53:52 2016
@@ -0,0 +1,45 @@
+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+
+declare void @llvm.amdgcn.s.sleep(i32) #0
+
+; GCN-LABEL: {{^}}test_s_sleep:
+; GCN: s_sleep 0{{$}}
+; GCN: s_sleep 1{{$}}
+; GCN: s_sleep 2{{$}}
+; GCN: s_sleep 3{{$}}
+; GCN: s_sleep 4{{$}}
+; GCN: s_sleep 5{{$}}
+; GCN: s_sleep 6{{$}}
+; GCN: s_sleep 7{{$}}
+; GCN: s_sleep 8{{$}}
+; GCN: s_sleep 9{{$}}
+; GCN: s_sleep 10{{$}}
+; GCN: s_sleep 11{{$}}
+; GCN: s_sleep 12{{$}}
+; GCN: s_sleep 13{{$}}
+; GCN: s_sleep 14{{$}}
+; GCN: s_sleep 15{{$}}
+define void @test_s_sleep(i32 %x) #0 {
+  call void @llvm.amdgcn.s.sleep(i32 0)
+  call void @llvm.amdgcn.s.sleep(i32 1)
+  call void @llvm.amdgcn.s.sleep(i32 2)
+  call void @llvm.amdgcn.s.sleep(i32 3)
+  call void @llvm.amdgcn.s.sleep(i32 4)
+  call void @llvm.amdgcn.s.sleep(i32 5)
+  call void @llvm.amdgcn.s.sleep(i32 6)
+  call void @llvm.amdgcn.s.sleep(i32 7)
+
+  ; Values that might only work on VI
+  call void @llvm.amdgcn.s.sleep(i32 8)
+  call void @llvm.amdgcn.s.sleep(i32 9)
+  call void @llvm.amdgcn.s.sleep(i32 10)
+  call void @llvm.amdgcn.s.sleep(i32 11)
+  call void @llvm.amdgcn.s.sleep(i32 12)
+  call void @llvm.amdgcn.s.sleep(i32 13)
+  call void @llvm.amdgcn.s.sleep(i32 14)
+  call void @llvm.amdgcn.s.sleep(i32 15)
+  ret void
+}
+
+attributes #0 = { nounwind }