Index: include/llvm/IR/IntrinsicsAMDGPU.td =================================================================== --- include/llvm/IR/IntrinsicsAMDGPU.td +++ include/llvm/IR/IntrinsicsAMDGPU.td @@ -104,6 +104,11 @@ // Instruction Intrinsics //===----------------------------------------------------------------------===// +def int_amdgcn_s_sendmsg : GCCBuiltin<"__builtin_amdgcn_s_sendmsg">, + Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], []>; +def int_amdgcn_s_sendmsghalt : GCCBuiltin<"__builtin_amdgcn_s_sendmsghalt">, + Intrinsic <[], [llvm_i32_ty, llvm_i32_ty], []>; + def int_amdgcn_s_barrier : GCCBuiltin<"__builtin_amdgcn_s_barrier">, Intrinsic<[], [], [IntrConvergent]>; Index: lib/Target/AMDGPU/AMDGPUISelLowering.h =================================================================== --- lib/Target/AMDGPU/AMDGPUISelLowering.h +++ lib/Target/AMDGPU/AMDGPUISelLowering.h @@ -307,6 +307,7 @@ /// Pointer to the start of the shader's constant data. CONST_DATA_PTR, SENDMSG, + SENDMSGHALT, INTERP_MOV, INTERP_P1, INTERP_P2, Index: lib/Target/AMDGPU/AMDGPUISelLowering.cpp =================================================================== --- lib/Target/AMDGPU/AMDGPUISelLowering.cpp +++ lib/Target/AMDGPU/AMDGPUISelLowering.cpp @@ -2799,6 +2799,7 @@ NODE_NAME_CASE(KILL) case AMDGPUISD::FIRST_MEM_OPCODE_NUMBER: break; NODE_NAME_CASE(SENDMSG) + NODE_NAME_CASE(SENDMSGHALT) NODE_NAME_CASE(INTERP_MOV) NODE_NAME_CASE(INTERP_P1) NODE_NAME_CASE(INTERP_P2) Index: lib/Target/AMDGPU/AMDGPUInstrInfo.td =================================================================== --- lib/Target/AMDGPU/AMDGPUInstrInfo.td +++ lib/Target/AMDGPU/AMDGPUInstrInfo.td @@ -253,6 +253,10 @@ SDTypeProfile<0, 1, [SDTCisInt<0>]>, [SDNPHasChain, SDNPInGlue]>; +def AMDGPUsendmsghalt : SDNode<"AMDGPUISD::SENDMSGHALT", + SDTypeProfile<0, 1, [SDTCisInt<0>]>, + [SDNPHasChain, SDNPInGlue]>; + def AMDGPUinterp_mov : SDNode<"AMDGPUISD::INTERP_MOV", SDTypeProfile<1, 3, [SDTCisFP<0>]>, [SDNPInGlue]>; Index: lib/Target/AMDGPU/SIISelLowering.cpp =================================================================== --- lib/Target/AMDGPU/SIISelLowering.cpp +++ lib/Target/AMDGPU/SIISelLowering.cpp @@ -2316,12 +2316,19 @@ unsigned IntrinsicID = cast(Op.getOperand(1))->getZExtValue(); switch (IntrinsicID) { - case AMDGPUIntrinsic::SI_sendmsg: { + case AMDGPUIntrinsic::SI_sendmsg: + case Intrinsic::amdgcn_s_sendmsg: { Chain = copyToM0(DAG, Chain, DL, Op.getOperand(3)); SDValue Glue = Chain.getValue(1); return DAG.getNode(AMDGPUISD::SENDMSG, DL, MVT::Other, Chain, Op.getOperand(2), Glue); } + case Intrinsic::amdgcn_s_sendmsghalt: { + Chain = copyToM0(DAG, Chain, DL, Op.getOperand(3)); + SDValue Glue = Chain.getValue(1); + return DAG.getNode(AMDGPUISD::SENDMSGHALT, DL, MVT::Other, Chain, + Op.getOperand(2), Glue); + } case AMDGPUIntrinsic::SI_tbuffer_store: { SDValue Ops[] = { Chain, Index: lib/Target/AMDGPU/SIInsertWaits.cpp =================================================================== --- lib/Target/AMDGPU/SIInsertWaits.cpp +++ lib/Target/AMDGPU/SIInsertWaits.cpp @@ -491,7 +491,7 @@ return; // There must be "S_NOP 0" between an instruction writing M0 and S_SENDMSG. - if (LastInstWritesM0 && I->getOpcode() == AMDGPU::S_SENDMSG) { + if (LastInstWritesM0 && (I->getOpcode() == AMDGPU::S_SENDMSG || I->getOpcode() == AMDGPU::S_SENDMSGHALT)) { BuildMI(MBB, I, DebugLoc(), TII->get(AMDGPU::S_NOP)).addImm(0); LastInstWritesM0 = false; return; @@ -591,7 +591,8 @@ // but we also want to wait for any other outstanding transfers before // signalling other hardware blocks if (I->getOpcode() == AMDGPU::S_BARRIER || - I->getOpcode() == AMDGPU::S_SENDMSG) + I->getOpcode() == AMDGPU::S_SENDMSG || + I->getOpcode() == AMDGPU::S_SENDMSGHALT) Required = LastIssued; else Required = handleOperands(*I); Index: lib/Target/AMDGPU/SOPInstructions.td =================================================================== --- lib/Target/AMDGPU/SOPInstructions.td +++ lib/Target/AMDGPU/SOPInstructions.td @@ -758,9 +758,12 @@ def S_SENDMSG : SOPP <0x00000010, (ins SendMsgImm:$simm16), "s_sendmsg $simm16", [(AMDGPUsendmsg (i32 imm:$simm16))] >; + +def S_SENDMSGHALT : SOPP <0x00000011, (ins SendMsgImm:$simm16), "s_sendmsghalt $simm16", + [(AMDGPUsendmsghalt (i32 imm:$simm16))] +>; } // 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_ICACHE_INV : SOPP <0x00000013, (ins), "s_icache_inv"> { let simm16 = 0; Index: test/CodeGen/AMDGPU/amdgcn.sendmsg-m0.ll =================================================================== --- /dev/null +++ test/CodeGen/AMDGPU/amdgcn.sendmsg-m0.ll @@ -0,0 +1,41 @@ +; RUN: llc -march=amdgcn -mcpu=verde -verify-machineinstrs < %s | FileCheck -check-prefix=SI -check-prefix=GCN %s +; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefix=VI -check-prefix=GCN %s + +; GCN-LABEL: {{^}}main: +; GCN: s_mov_b32 m0, s0 +; VI-NEXT: s_nop 0 +; GCN-NEXT: sendmsg(MSG_GS_DONE, GS_OP_NOP) +; GCN-NEXT: s_endpgm + +define amdgpu_gs void @main(i32 inreg %a) #0 { + call void @llvm.amdgcn.s.sendmsg(i32 3, i32 %a) + ret void +} + +; GCN-LABEL: {{^}}main_halt: +; GCN: s_mov_b32 m0, s0 +; VI-NEXT: s_nop 0 +; GCN-NEXT: s_sendmsghalt sendmsg(MSG_INTERRUPT) +; GCN-NEXT: s_endpgm + +define void @main_halt(i32 inreg %a) #0 { + call void @llvm.amdgcn.s.sendmsghalt(i32 1, i32 %a) + ret void +} + +; GCN-LABEL: {{^}}legacy: +; GCN: s_mov_b32 m0, s0 +; VI-NEXT: s_nop 0 +; GCN-NEXT: sendmsg(MSG_GS_DONE, GS_OP_NOP) +; GCN-NEXT: s_endpgm + +define amdgpu_gs void @legacy(i32 inreg %a) #0 { + call void @llvm.SI.sendmsg(i32 3, i32 %a) + ret void +} + +declare void @llvm.amdgcn.s.sendmsg(i32, i32) #0 +declare void @llvm.amdgcn.s.sendmsghalt(i32, i32) #0 +declare void @llvm.SI.sendmsg(i32, i32) #0 + +attributes #0 = { nounwind } Index: test/CodeGen/AMDGPU/amdgcn.sendmsg.ll =================================================================== --- /dev/null +++ test/CodeGen/AMDGPU/amdgcn.sendmsg.ll @@ -0,0 +1,49 @@ +;RUN: llc < %s -march=amdgcn -mcpu=verde -verify-machineinstrs | FileCheck %s +;RUN: llc < %s -march=amdgcn -mcpu=tonga -verify-machineinstrs | FileCheck %s + +; CHECK-LABEL: {{^}}main: +; CHECK: s_mov_b32 m0, 0 +; CHECK-NOT: s_mov_b32 m0 +; CHECK: s_sendmsg sendmsg(MSG_INTERRUPT) +; CHECK: s_sendmsg sendmsg(MSG_GS, GS_OP_EMIT, 0) +; CHECK: s_sendmsg sendmsg(MSG_GS, GS_OP_CUT, 1) +; CHECK: s_sendmsg sendmsg(MSG_GS, GS_OP_EMIT_CUT, 2) +; CHECK: s_sendmsg sendmsg(MSG_GS_DONE, GS_OP_NOP) +; CHECK: s_sendmsghalt sendmsg(MSG_INTERRUPT) +; CHECK: s_sendmsghalt sendmsg(MSG_GS, GS_OP_EMIT, 0) +; CHECK: s_sendmsghalt sendmsg(MSG_GS, GS_OP_CUT, 1) +; CHECK: s_sendmsghalt sendmsg(MSG_GS, GS_OP_EMIT_CUT, 2) +; CHECK: s_sendmsghalt sendmsg(MSG_GS_DONE, GS_OP_NOP) + +; Legacy +; CHECK: s_sendmsg sendmsg(MSG_GS, GS_OP_EMIT, 0) +; CHECK: s_sendmsg sendmsg(MSG_GS, GS_OP_CUT, 1) +; CHECK: s_sendmsg sendmsg(MSG_GS, GS_OP_EMIT_CUT, 2) +; CHECK: s_sendmsg sendmsg(MSG_GS_DONE, GS_OP_NOP) +define void @main() { +main_body: + call void @llvm.amdgcn.s.sendmsg(i32 1, i32 0); + call void @llvm.amdgcn.s.sendmsg(i32 34, i32 0); + call void @llvm.amdgcn.s.sendmsg(i32 274, i32 0); + call void @llvm.amdgcn.s.sendmsg(i32 562, i32 0); + call void @llvm.amdgcn.s.sendmsg(i32 3, i32 0); + + call void @llvm.amdgcn.s.sendmsghalt(i32 1, i32 0); + call void @llvm.amdgcn.s.sendmsghalt(i32 34, i32 0); + call void @llvm.amdgcn.s.sendmsghalt(i32 274, i32 0); + call void @llvm.amdgcn.s.sendmsghalt(i32 562, i32 0); + call void @llvm.amdgcn.s.sendmsghalt(i32 3, i32 0); + + call void @llvm.SI.sendmsg(i32 34, i32 0); + call void @llvm.SI.sendmsg(i32 274, i32 0); + call void @llvm.SI.sendmsg(i32 562, i32 0); + call void @llvm.SI.sendmsg(i32 3, i32 0); + ret void +} + +; Function Attrs: nounwind +declare void @llvm.amdgcn.s.sendmsg(i32, i32) #0 +declare void @llvm.amdgcn.s.sendmsghalt(i32, i32) #0 +declare void @llvm.SI.sendmsg(i32, i32) #0 + +attributes #0 = { nounwind } Index: test/CodeGen/AMDGPU/llvm.SI.sendmsg-m0.ll =================================================================== --- test/CodeGen/AMDGPU/llvm.SI.sendmsg-m0.ll +++ /dev/null @@ -1,17 +0,0 @@ -; RUN: llc -march=amdgcn -mcpu=verde -verify-machineinstrs < %s | FileCheck -check-prefix=SI -check-prefix=GCN %s -; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck -check-prefix=VI -check-prefix=GCN %s - -; GCN-LABEL: {{^}}main: -; GCN: s_mov_b32 m0, s0 -; VI-NEXT: s_nop 0 -; GCN-NEXT: sendmsg(MSG_GS_DONE, GS_OP_NOP) -; GCN-NEXT: s_endpgm - -define amdgpu_gs void @main(i32 inreg %a) #0 { - call void @llvm.SI.sendmsg(i32 3, i32 %a) - ret void -} - -declare void @llvm.SI.sendmsg(i32, i32) #0 - -attributes #0 = { nounwind } Index: test/CodeGen/AMDGPU/llvm.SI.sendmsg.ll =================================================================== --- test/CodeGen/AMDGPU/llvm.SI.sendmsg.ll +++ /dev/null @@ -1,24 +0,0 @@ -;RUN: llc < %s -march=amdgcn -mcpu=verde -verify-machineinstrs | FileCheck %s -;RUN: llc < %s -march=amdgcn -mcpu=tonga -verify-machineinstrs | FileCheck %s - -; CHECK-LABEL: {{^}}main: -; CHECK: s_mov_b32 m0, 0 -; CHECK-NOT: s_mov_b32 m0 -; CHECK: s_sendmsg sendmsg(MSG_GS, GS_OP_EMIT, 0) -; CHECK: s_sendmsg sendmsg(MSG_GS, GS_OP_CUT, 1) -; CHECK: s_sendmsg sendmsg(MSG_GS, GS_OP_EMIT_CUT, 2) -; CHECK: s_sendmsg sendmsg(MSG_GS_DONE, GS_OP_NOP) - -define void @main() { -main_body: - call void @llvm.SI.sendmsg(i32 34, i32 0); - call void @llvm.SI.sendmsg(i32 274, i32 0); - call void @llvm.SI.sendmsg(i32 562, i32 0); - call void @llvm.SI.sendmsg(i32 3, i32 0); - ret void -} - -; Function Attrs: nounwind -declare void @llvm.SI.sendmsg(i32, i32) #0 - -attributes #0 = { nounwind }