Index: include/llvm/IR/IntrinsicsAMDGPU.td
===================================================================
--- include/llvm/IR/IntrinsicsAMDGPU.td
+++ include/llvm/IR/IntrinsicsAMDGPU.td
@@ -458,6 +458,10 @@
   Intrinsic<[], [llvm_i32_ty], []> {
 }
 
+def int_amdgcn_s_trap :
+  GCCBuiltin<"__builtin_amdgcn_s_trap">,
+  Intrinsic<[], [llvm_i32_ty], []>;
+
 def int_amdgcn_s_incperflevel :
   GCCBuiltin<"__builtin_amdgcn_s_incperflevel">,
   Intrinsic<[], [llvm_i32_ty], []> {
Index: lib/Target/AMDGPU/SOPInstructions.td
===================================================================
--- lib/Target/AMDGPU/SOPInstructions.td
+++ lib/Target/AMDGPU/SOPInstructions.td
@@ -827,7 +827,13 @@
 } // End Uses = [EXEC, M0]
 
 def S_SENDMSGHALT : SOPP <0x00000011, (ins SendMsgImm:$simm16), "s_sendmsghalt $simm16">;
-def S_TRAP : SOPP <0x00000012, (ins i16imm:$simm16), "s_trap $simm16">;
+def S_TRAP : SOPP <0x00000012, (ins i32imm:$simm16), "s_trap $simm16", 
+  [(int_amdgcn_s_trap (SIMM16bit:$simm16))]> {
+  let hasSideEffects = 1;
+  let mayLoad = 1;
+  let mayStore = 1;    
+}
+
 def S_ICACHE_INV : SOPP <0x00000013, (ins), "s_icache_inv"> {
   let simm16 = 0;
 }
Index: test/CodeGen/AMDGPU/llvm.amdgcn.s.trap.ll
===================================================================
--- /dev/null
+++ test/CodeGen/AMDGPU/llvm.amdgcn.s.trap.ll
@@ -0,0 +1,25 @@
+; RUN: llc -march=amdgcn -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -march=amdgcn -mcpu=fiji -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+
+declare void @llvm.amdgcn.s.trap(i32) #0
+
+; GCN-LABEL: {{^}}test_s_trap:
+; GCN: s_trap 0{{$}}
+; GCN: s_trap 1{{$}}
+; GCN: s_trap 2{$}}
+; GCN: s_trap 3{{$}}
+; GCN: s_trap 4{{$}}
+; GCN: s_trap 5{{$}}
+; GCN: s_trap 6{{$}}
+define void @test_s_trap(i32 %x) #0 {
+  call void @llvm.amdgcn.s.trap(i32 0)
+  call void @llvm.amdgcn.s.trap(i32 1)
+  call void @llvm.amdgcn.s.trap(i32 2)
+  call void @llvm.amdgcn.s.trap(i32 3)
+  call void @llvm.amdgcn.s.trap(i32 4)
+  call void @llvm.amdgcn.s.trap(i32 5)
+  call void @llvm.amdgcn.s.trap(i32 6)
+  ret void
+}
+
+attributes #0 = { nounwind }