[llvm] [AMDGPU] New intrinsic void llvm.amdgcn.s.nop(i16) (PR #65757)
David Stuttard via llvm-commits
llvm-commits at lists.llvm.org
Fri Sep 8 06:49:11 PDT 2023
https://github.com/dstutt created https://github.com/llvm/llvm-project/pull/65757:
This allows front ends to insert s_nops - this is most often when a delay less
than s_sleep 1 is required.
>From 6945398024d9048812a49a66aa1b41ba6cf3b208 Mon Sep 17 00:00:00 2001
From: David Stuttard <david.stuttard at amd.com>
Date: Thu, 24 Aug 2023 13:57:24 +0100
Subject: [PATCH] [AMDGPU] New intrinsic void llvm.amdgcn.s.nop(i16)
This allows front ends to insert s_nops - this is most often when a delay less
than s_sleep 1 is required.
---
llvm/include/llvm/IR/IntrinsicsAMDGPU.td | 5 ++++
llvm/lib/Target/AMDGPU/SOPInstructions.td | 3 +-
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.nop.ll | 30 +++++++++++++++++++
3 files changed, 37 insertions(+), 1 deletion(-)
create mode 100644 llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.nop.ll
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index e9b13c3adcbaa54..124f22c1a9b27c7 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -1677,6 +1677,11 @@ def int_amdgcn_s_sleep :
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,
diff --git a/llvm/lib/Target/AMDGPU/SOPInstructions.td b/llvm/lib/Target/AMDGPU/SOPInstructions.td
index 38fa90bdc937276..175045a8a893e92 100644
--- a/llvm/lib/Target/AMDGPU/SOPInstructions.td
+++ b/llvm/lib/Target/AMDGPU/SOPInstructions.td
@@ -1161,7 +1161,8 @@ multiclass SOPP_With_Relaxation <string opName, dag ins,
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"> {
+def S_NOP : SOPP_Pseudo<"s_nop" , (ins i16imm:$simm16), "$simm16",
+ [(int_amdgcn_s_nop timm:$simm16)]> {
let hasSideEffects = 1;
}
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.nop.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.nop.ll
new file mode 100644
index 000000000000000..a625f973c0b8f55
--- /dev/null
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.s.nop.ll
@@ -0,0 +1,30 @@
+; 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)
More information about the llvm-commits
mailing list