New intrinsic void llvm.amdgcn.s.nop(i16)

[arsenm]

What’s the point?

[dstuttard]

In D158750#4614294, @arsenm wrote:

What’s the point?

It can be used by a front-end to insert a delay that's less than s_sleep 1

[arsenm]

llvm/lib/Target/AMDGPU/SOPInstructions.td

1166

The regular one should be marked with side effects

[dstuttard]

llvm/lib/Target/AMDGPU/SOPInstructions.td

1166

My initial implementation did that - even though I was unsure if a nop should be tagged as having side-effects.
Also, there were failures in some of the lit tests (mir variants) which made me even more reluctant to tag s_nop as having side-effects since it would no longer be a NFC for vanilla s_nops.
I got failures (mainly (all?) out of order matches) for:

LLVM :: CodeGen/AMDGPU/copy-vgpr-clobber-spill-vgpr.mir
LLVM :: CodeGen/AMDGPU/fold-immediate-operand-shrink.mir
LLVM :: CodeGen/AMDGPU/sched-barrier-pre-RA.mir
LLVM :: CodeGen/AMDGPU/sched-group-barrier-pipeline-solver.mir
LLVM :: CodeGen/AMDGPU/sched-group-barrier-pre-RA.mir
LLVM :: CodeGen/AMDGPU/spill-empty-live-interval.mir

[arsenm]

llvm/lib/Target/AMDGPU/SOPInstructions.td

1166

First I've noticed this odd bug that it doesn't have side effects before. That should be fixed independently.

Also, pseudos that modify any real instruction should use a lowercase suffix

[dstuttard]

Rename pseudo instruction (lowercase suffix)

[foad]

Depending on the use case, you might also have to mark it as a scheduling barrier, not just having side effects.

[dstuttard]

In D158750#4637244, @foad wrote:

Depending on the use case, you might also have to mark it as a scheduling barrier, not just having side effects.

It is possible that might be required, but the behaviour is modelled on s_sleep which only tags "hasSideEffects" and doesn't attempt to make it a scheduling barrier as well.
I propose to leave the implementation as-is and it can be re-visited if a scheduling barrier needs to be enforced.

[foad]

In D158750#4637940, @dstuttard wrote:

In D158750#4637244, @foad wrote:

Depending on the use case, you might also have to mark it as a scheduling barrier, not just having side effects.

It is possible that might be required, but the behaviour is modelled on s_sleep which only tags "hasSideEffects" and doesn't attempt to make it a scheduling barrier as well.
I propose to leave the implementation as-is and it can be re-visited if a scheduling barrier needs to be enforced.

That's fine with me.

[dstuttard]

@arsenm are you happy with this?

[foad]

llvm/lib/Target/AMDGPU/SOPInstructions.td

1166

The regular one should be marked with side effects

https://github.com/llvm/llvm-project/pull/65745

[dstuttard]

Rebased on @foad 's change and submitted as PR here: https://github.com/llvm/llvm-project/pull/65757

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

def int_amdgcn_s_memtime :
   DefaultAttrsIntrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrHasSideEffects]>;
 
 def int_amdgcn_s_sleep :
   ClangBuiltin<"__builtin_amdgcn_s_sleep">,
   DefaultAttrsIntrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
                                 IntrHasSideEffects]> {
 }
 
+def int_amdgcn_s_nop :
+  DefaultAttrsIntrinsic<[], [llvm_i16_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
+                                IntrHasSideEffects]> {
+}
+
 def int_amdgcn_s_incperflevel :
   ClangBuiltin<"__builtin_amdgcn_s_incperflevel">,
   DefaultAttrsIntrinsic<[], [llvm_i32_ty], [ImmArg<ArgIndex<0>>, IntrNoMem,
                                 IntrHasSideEffects]> {
 }
 
 def int_amdgcn_s_decperflevel :
   ClangBuiltin<"__builtin_amdgcn_s_decperflevel">,

llvm/lib/Target/AMDGPU/AMDGPUMCInstLower.cpp

else if (Opcode == AMDGPU::SI_CALL) {
     lowerOperand(MI->getOperand(1), Src);
     OutMI.addOperand(Dest);
     OutMI.addOperand(Src);
     return;
   } else if (Opcode == AMDGPU::SI_TCRETURN ||
              Opcode == AMDGPU::SI_TCRETURN_GFX) {
     // TODO: How to use branch immediate and avoid register+add?
     Opcode = AMDGPU::S_SETPC_B64;
+  } else if (Opcode == AMDGPU::S_NOP_se) {
+    Opcode = AMDGPU::S_NOP;
   }
 
   int MCOpcode = TII->pseudoToMCOpcode(Opcode);
   if (MCOpcode == -1) {
     LLVMContext &C = MI->getParent()->getParent()->getFunction().getContext();
     C.emitError("AMDGPUMCInstLower::lower - Pseudo instruction doesn't have "
                 "a target-specific version: " + Twine(MI->getOpcode()));
   }

llvm/lib/Target/AMDGPU/SOPInstructions.td

 multiclass SOPP_With_Relaxation <string opName, dag ins,
                   string asmOps, list<dag> pattern=[]> {
   def "" : SOPP_Pseudo <opName, ins, asmOps, pattern>;
   def _pad_s_nop : SOPP_Pseudo <opName # "_pad_s_nop", ins, asmOps, pattern, " ", opName>;
 }
 
 def S_NOP : SOPP_Pseudo<"s_nop" , (ins i16imm:$simm16), "$simm16">;
 
+// Define variant marked as having side effects for use with intrinsic
+def S_NOP_se : SOPP_Pseudo<"" , (ins i16imm:$simm16), "$simm16",
+  [(int_amdgcn_s_nop timm:$simm16)]> {
+  let hasSideEffects = 1;
+}
+
 let isTerminator = 1 in {
 def S_ENDPGM : SOPP_Pseudo<"s_endpgm", (ins Endpgm:$simm16), "$simm16", [], ""> {
   let isBarrier = 1;
   let isReturn = 1;
   let hasSideEffects = 1;
 }
 
 def S_ENDPGM_SAVED : SOPP_Pseudo<"s_endpgm_saved", (ins)> {

llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.nop.ll

+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -march=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+; RUN: llc -global-isel -march=amdgcn -mcpu=gfx900 -verify-machineinstrs < %s | FileCheck -check-prefix=GCN %s
+
+define amdgpu_kernel void @test_s_nop() {
+; GCN-LABEL: test_s_nop:
+; GCN:       ; %bb.0:
+; GCN-NEXT:    s_nop 0
+; GCN-NEXT:    s_nop 1
+; GCN-NEXT:    s_nop 2
+; GCN-NEXT:    s_nop 3
+; GCN-NEXT:    s_nop 4
+; GCN-NEXT:    s_nop 5
+; GCN-NEXT:    s_nop 6
+; GCN-NEXT:    s_nop 7
+; GCN-NEXT:    s_nop 63
+; GCN-NEXT:    s_endpgm
+  call void @llvm.amdgcn.s.nop(i16 0)
+  call void @llvm.amdgcn.s.nop(i16 1)
+  call void @llvm.amdgcn.s.nop(i16 2)
+  call void @llvm.amdgcn.s.nop(i16 3)
+  call void @llvm.amdgcn.s.nop(i16 4)
+  call void @llvm.amdgcn.s.nop(i16 5)
+  call void @llvm.amdgcn.s.nop(i16 6)
+  call void @llvm.amdgcn.s.nop(i16 7)
+  call void @llvm.amdgcn.s.nop(i16 63)
+  ret void
+}
+
+declare void @llvm.amdgcn.s.nop(i16)