AMDGPU: Add s_sleep intrinsic

[• tstellarAMD]

LGTM.

[arsenm]

r262120

include/llvm/IR/IntrinsicsAMDGPU.td

 def int_amdgcn_s_dcache_inv :
   GCCBuiltin<"__builtin_amdgcn_s_dcache_inv">,
   Intrinsic<[], [], []>;
 
 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]>;
 
 // __builtin_amdgcn_interp_p1 <i>, <attr_chan>, <attr>, <m0>
 def int_amdgcn_interp_p1 :
   GCCBuiltin<"__builtin_amdgcn_interp_p1">,
   Intrinsic<[llvm_float_ty],

lib/Target/AMDGPU/SIInstrInfo.td

 def IMM12bit : PatLeaf <(imm),
   [{return isUInt<12>(N->getZExtValue());}]
 >;
 
 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());}]
 >;
 
 def IMM32bit : PatLeaf <(imm),
   [{return isUInt<32>(N->getZExtValue());}]
 >;
 

lib/Target/AMDGPU/SIInstructions.td

> {
   let simm16 = 0;
   let mayLoad = 1;
   let mayStore = 1;
   let isConvergent = 1;
 }
 
 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))]
   >;
 } // End Uses = [EXEC, M0]
 
 def S_SENDMSGHALT : SOPP <0x00000011, (ins i16imm:$simm16), "s_sendmsghalt $simm16">;
 def S_TRAP : SOPP <0x00000012, (ins i16imm:$simm16), "s_trap $simm16">;
 def S_ICACHE_INV : SOPP <0x00000013, (ins), "s_icache_inv"> {

test/CodeGen/AMDGPU/llvm.amdgcn.s.sleep.ll

+; 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 }