Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -231,6 +231,10 @@
   Intrinsic<[], [llvm_i32_ty], []> {
 }
 
+def int_amdgcn_s_getreg :
+  GCCBuiltin<"__builtin_amdgcn_s_getreg">,
+  Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrReadMem]>;
+
 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
@@ -979,7 +979,7 @@
 
 }
 
-multiclass SOPK_32 <sopk op, string opName, list<dag> pattern> {
+multiclass SOPK_16 <sopk op, string opName, list<dag> pattern> {
   def "" : SOPK_Pseudo <opName, (outs SReg_32:$sdst), (ins u16imm:$simm16),
     pattern>;
 
@@ -990,6 +990,17 @@
     opName#" $sdst, $simm16">;
 }
 
+multiclass SOPK_32 <sopk op, string opName, list<dag> pattern> {
+  def "" : SOPK_Pseudo <opName, (outs SReg_32:$sdst), (ins i32imm:$simm16),
+    pattern>;
+
+  def _si : SOPK_Real_si <op, opName, (outs SReg_32:$sdst), (ins i32imm:$simm16),
+    opName#" $sdst, $simm16">;
+
+  def _vi : SOPK_Real_vi <op, opName, (outs SReg_32:$sdst), (ins i32imm:$simm16),
+    opName#" $sdst, $simm16">;
+}
+
 multiclass SOPK_SCC <sopk op, string opName, list<dag> pattern> {
   def "" : SOPK_Pseudo <opName, (outs),
     (ins SReg_32:$src0, u16imm:$src1), pattern> {
@@ -1008,7 +1019,7 @@
   }
 }
 
-multiclass SOPK_32TIE <sopk op, string opName, list<dag> pattern> : SOPK_m <
+multiclass SOPK_16TIE <sopk op, string opName, list<dag> pattern> : SOPK_m <
   op, opName, (outs SReg_32:$sdst), (ins SReg_32:$src0, u16imm:$simm16),
   " $sdst, $simm16"
 >;
Index: lib/Target/AMDGPU/SIInstructions.td
===================================================================
--- lib/Target/AMDGPU/SIInstructions.td
+++ lib/Target/AMDGPU/SIInstructions.td
@@ -367,10 +367,10 @@
 //===----------------------------------------------------------------------===//
 
 let isReMaterializable = 1, isMoveImm = 1 in {
-defm S_MOVK_I32 : SOPK_32 <sopk<0x00>, "s_movk_i32", []>;
+defm S_MOVK_I32 : SOPK_16 <sopk<0x00>, "s_movk_i32", []>;
 } // End isReMaterializable = 1
 let Uses = [SCC] in {
-  defm S_CMOVK_I32 : SOPK_32 <sopk<0x02, 0x01>, "s_cmovk_i32", []>;
+  defm S_CMOVK_I32 : SOPK_16 <sopk<0x02, 0x01>, "s_cmovk_i32", []>;
 }
 
 let isCompare = 1 in {
@@ -408,15 +408,20 @@
 
 let Defs = [SCC], isCommutable = 1, DisableEncoding = "$src0",
     Constraints = "$sdst = $src0" in {
-  defm S_ADDK_I32 : SOPK_32TIE <sopk<0x0f, 0x0e>, "s_addk_i32", []>;
-  defm S_MULK_I32 : SOPK_32TIE <sopk<0x10, 0x0f>, "s_mulk_i32", []>;
+  defm S_ADDK_I32 : SOPK_16TIE <sopk<0x0f, 0x0e>, "s_addk_i32", []>;
+  defm S_MULK_I32 : SOPK_16TIE <sopk<0x10, 0x0f>, "s_mulk_i32", []>;
 }
 
 defm S_CBRANCH_I_FORK : SOPK_m <
   sopk<0x11, 0x10>, "s_cbranch_i_fork", (outs),
   (ins SReg_64:$sdst, u16imm:$simm16), " $sdst, $simm16"
 >;
-defm S_GETREG_B32 : SOPK_32 <sopk<0x12, 0x11>, "s_getreg_b32", []>;
+
+let mayLoad = 1 in {
+defm S_GETREG_B32 : SOPK_32 <sopk<0x12, 0x11>, "s_getreg_b32",
+  [(set SReg_32:$sdst, (int_amdgcn_s_getreg imm:$simm16))]>;
+}
+
 defm S_SETREG_B32 : SOPK_m <
   sopk<0x13, 0x12>, "s_setreg_b32", (outs),
   (ins SReg_32:$sdst, u16imm:$simm16), " $sdst, $simm16"
Index: test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll
===================================================================
--- /dev/null
+++ test/CodeGen/AMDGPU/llvm.amdgcn.s.getreg.ll
@@ -0,0 +1,16 @@
+; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=kaveri -verify-machineinstrs < %s | FileCheck %s
+; RUN: llc -mtriple=amdgcn--amdhsa -mcpu=fiji -verify-machineinstrs < %s | FileCheck %s
+
+; FUNC-LABEL: {{^}}s_getreg_test:
+; CHECK: s_getreg_b32 s{{[0-9]+}}, 0xb206
+define void @s_getreg_test(i32 addrspace(1)* %out) { ; simm16=45574 for lds size.
+  %lds_size_64dwords = call i32 @llvm.amdgcn.s.getreg(i32 45574) #0
+  %lds_size_bytes = shl i32 %lds_size_64dwords, 8
+  store i32 %lds_size_bytes, i32 addrspace(1)* %out
+  ret void
+}
+
+declare i32 @llvm.amdgcn.s.getreg(i32) #0
+
+attributes #0 = { nounwind readonly}