Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -192,6 +192,11 @@
   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]>;
Index: lib/Target/AMDGPU/SIInstrInfo.td
===================================================================
--- lib/Target/AMDGPU/SIInstrInfo.td
+++ lib/Target/AMDGPU/SIInstrInfo.td
@@ -371,6 +371,10 @@
   [{return isUInt<16>(N->getZExtValue());}]
 >;
 
+def SIMM16bit : PatLeaf <(imm),
+  [{return isInt<16>(N->getSExtValue());}]
+>;
+
 def IMM20bit : PatLeaf <(imm),
   [{return isUInt<20>(N->getZExtValue());}]
 >;
Index: lib/Target/AMDGPU/SIInstructions.td
===================================================================
--- lib/Target/AMDGPU/SIInstructions.td
+++ lib/Target/AMDGPU/SIInstructions.td
@@ -501,10 +501,22 @@
 
 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))]
   >;
Index: test/CodeGen/AMDGPU/llvm.amdgcn.s.sleep.ll
===================================================================
--- /dev/null
+++ test/CodeGen/AMDGPU/llvm.amdgcn.s.sleep.ll
@@ -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 }