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 }