[llvm] [NVPTX] Add prefetch tensormap variant (PR #146203)
Abhilash Majumder via llvm-commits
llvm-commits at lists.llvm.org
Tue Jul 1 01:10:08 PDT 2025
https://github.com/abhilash1910 updated https://github.com/llvm/llvm-project/pull/146203
>From 7e3815bafdcba79ad0f0f66b29b6d15ddf796d13 Mon Sep 17 00:00:00 2001
From: Abhilash Majumder <abmajumder at nvidia.com>
Date: Sat, 28 Jun 2025 13:55:04 +0530
Subject: [PATCH 1/5] add prefetch tensormap variant
---
llvm/docs/NVPTXUsage.rst | 8 ++++++-
llvm/include/llvm/IR/IntrinsicsNVVM.td | 4 ++++
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 9 +++++++
llvm/test/CodeGen/NVPTX/prefetch.ll | 30 ++++++++++++++++++++++++
4 files changed, 50 insertions(+), 1 deletion(-)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 11017fe4e01b4..ca951811b73dd 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -971,6 +971,9 @@ Syntax:
declare void @llvm.nvvm.prefetch.L1(ptr %ptr)
declare void @llvm.nvvm.prefetch.L2(ptr %ptr)
+ declare void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
+ declare void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
+
declare void @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
declare void @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)
@@ -983,7 +986,10 @@ The '``@llvm.nvvm.prefetch.*``' and '``@llvm.nvvm.prefetchu.*``' intrinsic
correspond to the '``prefetch.*``;' and '``prefetchu.*``' family of PTX instructions.
The '``prefetch.*``' instructions bring the cache line containing the
specified address in global or local memory address space into the
-specified cache level (L1 or L2). The '`prefetchu.*``' instruction brings the cache line
+specified cache level (L1 or L2). If the '``.tensormap``' qualifier is specified then the
+prefetch instruction brings the cache line containing the specified address in the
+'``.const``' or '``.param memory``' state space for subsequent use by the '``cp.async.bulk.tensor``'
+instruction.The '`prefetchu.*``' instruction brings the cache line
containing the specified generic address into the specified uniform cache level.
If no address space is specified, it is assumed to be generic address. The intrinsic
uses and eviction priority which can be accessed by the '``.level::eviction_priority``' modifier.
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 0375f29ad8906..0678bba51e4a3 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -137,6 +137,7 @@
def llvm_global_ptr_ty : LLVMQualPointerType<1>; // (global)ptr
def llvm_shared_ptr_ty : LLVMQualPointerType<3>; // (shared)ptr
+def llvm_constant_ptr_ty: LLVMQualPointerType<4>; // (const)ptr
def llvm_local_ptr_ty : LLVMQualPointerType<5>; // (local)ptr
def llvm_tmem_ptr_ty : LLVMQualPointerType<6>; // (tensor memory)ptr
def llvm_shared_cluster_ptr_ty : LLVMQualPointerType<7>; // (shared_cluster)ptr
@@ -2092,6 +2093,9 @@ let IntrProperties = [IntrArgMemOnly, ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<
def int_nvvm_prefetch_local_ # level : Intrinsic<[], [llvm_local_ptr_ty]>;
}
+ def int_nvvm_prefetch_tensormap_p0: DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
+ def int_nvvm_prefetch_tensormap_p4 : DefaultAttrsIntrinsic<[], [llvm_constant_ptr_ty]>;
+
foreach eviction_priority = ["evict_normal", "evict_last"] in
def int_nvvm_prefetch_global_L2_ # eviction_priority : Intrinsic<[], [llvm_global_ptr_ty]>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index cc1fd027d8515..8afc7063c363a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -760,6 +760,15 @@ def PREFETCH_GLOBAL_L1 : PREFETCH_INTRS<"prefetch.global.L1">;
def PREFETCH_LOCAL_L1 : PREFETCH_INTRS<"prefetch.local.L1">;
def PREFETCH_GLOBAL_L2 : PREFETCH_INTRS<"prefetch.global.L2">;
def PREFETCH_LOCAL_L2 : PREFETCH_INTRS<"prefetch.local.L2">;
+def PREFETCH_CONST_TENSORMAP : BasicNVPTXInst<(outs), (ins ADDR:$addr),
+ "prefetch.const.tensormap",
+ [(int_nvvm_prefetch_tensormap_p4 addr:$addr)]>,
+ Requires<[hasPTX<80>, hasSM<90>]>;
+
+def PREFETCH_GENERIC_TENSORMAP : BasicNVPTXInst<(outs), (ins ADDR:$addr),
+ "prefetch.tensormap",
+ [(int_nvvm_prefetch_tensormap_p0 addr:$addr)]>,
+ Requires<[hasPTX<80>, hasSM<90>]>;
def PREFETCH_GLOBAL_L2_EVICT_NORMAL : BasicNVPTXInst<(outs), (ins ADDR:$addr),
"prefetch.global.L2::evict_normal",
diff --git a/llvm/test/CodeGen/NVPTX/prefetch.ll b/llvm/test/CodeGen/NVPTX/prefetch.ll
index a64e4fe7a508e..b63155ff49185 100644
--- a/llvm/test/CodeGen/NVPTX/prefetch.ll
+++ b/llvm/test/CodeGen/NVPTX/prefetch.ll
@@ -12,6 +12,9 @@ declare void @llvm.nvvm.prefetch.local.L2(ptr addrspace(5) %local_ptr)
declare void @llvm.nvvm.prefetch.L1(ptr %ptr)
declare void @llvm.nvvm.prefetch.L2(ptr %ptr)
+declare void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
+declare void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
+
declare void @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
declare void @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)
@@ -67,6 +70,33 @@ define void @prefetch_(ptr %ptr) {
ret void
}
+
+define void @prefetch_generic_tensormap(ptr %ptr) {
+; CHECK-PTX64-LABEL: prefetch_generic_tensormap(
+; CHECK-PTX64: {
+; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT: // %bb.0:
+; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_generic_tensormap_param_0];
+; CHECK-PTX64-NEXT: prefetch.tensormap [%rd1];
+; CHECK-PTX64-NEXT: ret;
+ tail call void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
+ ret void
+}
+
+define void @prefetch_const_tensormap(ptr addrspace(4) %const_ptr) {
+; CHECK-PTX64-LABEL: prefetch_const_tensormap(
+; CHECK-PTX64: {
+; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
+; CHECK-PTX64-EMPTY:
+; CHECK-PTX64-NEXT: // %bb.0:
+; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_const_tensormap_param_0];
+; CHECK-PTX64-NEXT: prefetch.const.tensormap [%rd1];
+; CHECK-PTX64-NEXT: ret;
+ tail call void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
+ ret void
+}
+
define void @prefetchu_l1(ptr %ptr) {
; CHECK-PTX64-LABEL: prefetchu_l1(
; CHECK-PTX64: {
>From 775daa341731a301e9cfbe544685962e5c7122fc Mon Sep 17 00:00:00 2001
From: Abhilash Majumder <abmajumder at nvidia.com>
Date: Tue, 1 Jul 2025 12:04:37 +0530
Subject: [PATCH 2/5] use generic and const names
---
llvm/docs/NVPTXUsage.rst | 4 ++--
llvm/include/llvm/IR/IntrinsicsNVVM.td | 4 ++--
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 8 ++------
llvm/test/CodeGen/NVPTX/prefetch.ll | 8 ++++----
4 files changed, 10 insertions(+), 14 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index ca951811b73dd..33d36b9411c1a 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -971,8 +971,8 @@ Syntax:
declare void @llvm.nvvm.prefetch.L1(ptr %ptr)
declare void @llvm.nvvm.prefetch.L2(ptr %ptr)
- declare void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
- declare void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
+ declare void @llvm.nvvm.prefetch.generic.tensormap(ptr %ptr)
+ declare void @llvm.nvvm.prefetch.const.tensormap(ptr addrspace(4) %const_ptr)
declare void @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
declare void @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index 0678bba51e4a3..c8df95994011b 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2093,8 +2093,8 @@ let IntrProperties = [IntrArgMemOnly, ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<
def int_nvvm_prefetch_local_ # level : Intrinsic<[], [llvm_local_ptr_ty]>;
}
- def int_nvvm_prefetch_tensormap_p0: DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
- def int_nvvm_prefetch_tensormap_p4 : DefaultAttrsIntrinsic<[], [llvm_constant_ptr_ty]>;
+ def int_nvvm_prefetch_generic_tensormap : DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
+ def int_nvvm_prefetch_const_tensormap: DefaultAttrsIntrinsic<[], [llvm_constant_ptr_ty]>;
foreach eviction_priority = ["evict_normal", "evict_last"] in
def int_nvvm_prefetch_global_L2_ # eviction_priority : Intrinsic<[], [llvm_global_ptr_ty]>;
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index 8afc7063c363a..d8446b4b4dbe6 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -760,14 +760,10 @@ def PREFETCH_GLOBAL_L1 : PREFETCH_INTRS<"prefetch.global.L1">;
def PREFETCH_LOCAL_L1 : PREFETCH_INTRS<"prefetch.local.L1">;
def PREFETCH_GLOBAL_L2 : PREFETCH_INTRS<"prefetch.global.L2">;
def PREFETCH_LOCAL_L2 : PREFETCH_INTRS<"prefetch.local.L2">;
-def PREFETCH_CONST_TENSORMAP : BasicNVPTXInst<(outs), (ins ADDR:$addr),
- "prefetch.const.tensormap",
- [(int_nvvm_prefetch_tensormap_p4 addr:$addr)]>,
- Requires<[hasPTX<80>, hasSM<90>]>;
-
+def PREFETCH_CONST_TENSORMAP : PREFETCH_INTRS<"prefetch.const.tensormap">;
def PREFETCH_GENERIC_TENSORMAP : BasicNVPTXInst<(outs), (ins ADDR:$addr),
"prefetch.tensormap",
- [(int_nvvm_prefetch_tensormap_p0 addr:$addr)]>,
+ [(int_nvvm_prefetch_generic_tensormap addr:$addr)]>,
Requires<[hasPTX<80>, hasSM<90>]>;
def PREFETCH_GLOBAL_L2_EVICT_NORMAL : BasicNVPTXInst<(outs), (ins ADDR:$addr),
diff --git a/llvm/test/CodeGen/NVPTX/prefetch.ll b/llvm/test/CodeGen/NVPTX/prefetch.ll
index b63155ff49185..d9b4e48167310 100644
--- a/llvm/test/CodeGen/NVPTX/prefetch.ll
+++ b/llvm/test/CodeGen/NVPTX/prefetch.ll
@@ -12,8 +12,8 @@ declare void @llvm.nvvm.prefetch.local.L2(ptr addrspace(5) %local_ptr)
declare void @llvm.nvvm.prefetch.L1(ptr %ptr)
declare void @llvm.nvvm.prefetch.L2(ptr %ptr)
-declare void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
-declare void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
+declare void @llvm.nvvm.prefetch.generic.tensormap(ptr %ptr)
+declare void @llvm.nvvm.prefetch.const.tensormap(ptr addrspace(4) %const_ptr)
declare void @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
declare void @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)
@@ -80,7 +80,7 @@ define void @prefetch_generic_tensormap(ptr %ptr) {
; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_generic_tensormap_param_0];
; CHECK-PTX64-NEXT: prefetch.tensormap [%rd1];
; CHECK-PTX64-NEXT: ret;
- tail call void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
+ tail call void @llvm.nvvm.prefetch.generic.tensormap(ptr %ptr)
ret void
}
@@ -93,7 +93,7 @@ define void @prefetch_const_tensormap(ptr addrspace(4) %const_ptr) {
; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_const_tensormap_param_0];
; CHECK-PTX64-NEXT: prefetch.const.tensormap [%rd1];
; CHECK-PTX64-NEXT: ret;
- tail call void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
+ tail call void @llvm.nvvm.prefetch.const.tensormap(ptr addrspace(4) %const_ptr)
ret void
}
>From aa8e4d019341f2bde555e2cc5de636c865ecea22 Mon Sep 17 00:00:00 2001
From: Abhilash Majumder <abmajumder at nvidia.com>
Date: Tue, 1 Jul 2025 12:06:26 +0530
Subject: [PATCH 3/5] format
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index c8df95994011b..a26a35bb0d947 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2093,8 +2093,8 @@ let IntrProperties = [IntrArgMemOnly, ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<
def int_nvvm_prefetch_local_ # level : Intrinsic<[], [llvm_local_ptr_ty]>;
}
- def int_nvvm_prefetch_generic_tensormap : DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
- def int_nvvm_prefetch_const_tensormap: DefaultAttrsIntrinsic<[], [llvm_constant_ptr_ty]>;
+ def int_nvvm_prefetch_generic_tensormap : DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
+ def int_nvvm_prefetch_const_tensormap : DefaultAttrsIntrinsic<[], [llvm_constant_ptr_ty]>;
foreach eviction_priority = ["evict_normal", "evict_last"] in
def int_nvvm_prefetch_global_L2_ # eviction_priority : Intrinsic<[], [llvm_global_ptr_ty]>;
>From 6931c80342ccc5acfc2101d93b61c10cf79051d0 Mon Sep 17 00:00:00 2001
From: Abhilash Majumder <abmajumder at nvidia.com>
Date: Tue, 1 Jul 2025 13:23:17 +0530
Subject: [PATCH 4/5] refresh
---
llvm/docs/NVPTXUsage.rst | 2 +-
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 10 +++++-----
2 files changed, 6 insertions(+), 6 deletions(-)
diff --git a/llvm/docs/NVPTXUsage.rst b/llvm/docs/NVPTXUsage.rst
index 33d36b9411c1a..2c7a531f34a8f 100644
--- a/llvm/docs/NVPTXUsage.rst
+++ b/llvm/docs/NVPTXUsage.rst
@@ -989,7 +989,7 @@ specified address in global or local memory address space into the
specified cache level (L1 or L2). If the '``.tensormap``' qualifier is specified then the
prefetch instruction brings the cache line containing the specified address in the
'``.const``' or '``.param memory``' state space for subsequent use by the '``cp.async.bulk.tensor``'
-instruction.The '`prefetchu.*``' instruction brings the cache line
+instruction. The '`prefetchu.*``' instruction brings the cache line
containing the specified generic address into the specified uniform cache level.
If no address space is specified, it is assumed to be generic address. The intrinsic
uses and eviction priority which can be accessed by the '``.level::eviction_priority``' modifier.
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index d8446b4b4dbe6..a4ee24aaf4ce3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -760,11 +760,11 @@ def PREFETCH_GLOBAL_L1 : PREFETCH_INTRS<"prefetch.global.L1">;
def PREFETCH_LOCAL_L1 : PREFETCH_INTRS<"prefetch.local.L1">;
def PREFETCH_GLOBAL_L2 : PREFETCH_INTRS<"prefetch.global.L2">;
def PREFETCH_LOCAL_L2 : PREFETCH_INTRS<"prefetch.local.L2">;
-def PREFETCH_CONST_TENSORMAP : PREFETCH_INTRS<"prefetch.const.tensormap">;
-def PREFETCH_GENERIC_TENSORMAP : BasicNVPTXInst<(outs), (ins ADDR:$addr),
- "prefetch.tensormap",
- [(int_nvvm_prefetch_generic_tensormap addr:$addr)]>,
- Requires<[hasPTX<80>, hasSM<90>]>;
+def PREFETCH_CONST_TENSORMAP : PREFETCH_INTRS<"prefetch.const.tensormap">;
+def PREFETCH_GENERIC_TENSORMAP : BasicNVPTXInst<(outs), (ins ADDR:$addr),
+ "prefetch.tensormap",
+ [(int_nvvm_prefetch_generic_tensormap addr:$addr)]>,
+ Requires<[hasPTX<80>, hasSM<90>]>;
def PREFETCH_GLOBAL_L2_EVICT_NORMAL : BasicNVPTXInst<(outs), (ins ADDR:$addr),
"prefetch.global.L2::evict_normal",
>From 887e139b515d60936c834e6149c6ebded8ba860c Mon Sep 17 00:00:00 2001
From: Abhilash Majumder <abmajumder at nvidia.com>
Date: Tue, 1 Jul 2025 13:39:41 +0530
Subject: [PATCH 5/5] refactor and refresh
---
llvm/include/llvm/IR/IntrinsicsNVVM.td | 10 +++---
llvm/lib/Target/NVPTX/NVPTXIntrinsics.td | 41 ++++++++----------------
2 files changed, 19 insertions(+), 32 deletions(-)
diff --git a/llvm/include/llvm/IR/IntrinsicsNVVM.td b/llvm/include/llvm/IR/IntrinsicsNVVM.td
index a26a35bb0d947..1be79bb5525ae 100644
--- a/llvm/include/llvm/IR/IntrinsicsNVVM.td
+++ b/llvm/include/llvm/IR/IntrinsicsNVVM.td
@@ -2088,18 +2088,18 @@ foreach dim = 1...5 in {
// Intrinsics for Prefetch and Prefetchu
let IntrProperties = [IntrArgMemOnly, ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>] in {
foreach level = ["L1", "L2"] in {
- def int_nvvm_prefetch_ # level : Intrinsic<[], [llvm_ptr_ty]>;
- def int_nvvm_prefetch_global_ # level : Intrinsic<[], [llvm_global_ptr_ty]>;
- def int_nvvm_prefetch_local_ # level : Intrinsic<[], [llvm_local_ptr_ty]>;
+ def int_nvvm_prefetch_ # level : DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
+ def int_nvvm_prefetch_global_ # level : DefaultAttrsIntrinsic<[], [llvm_global_ptr_ty]>;
+ def int_nvvm_prefetch_local_ # level : DefaultAttrsIntrinsic<[], [llvm_local_ptr_ty]>;
}
def int_nvvm_prefetch_generic_tensormap : DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
def int_nvvm_prefetch_const_tensormap : DefaultAttrsIntrinsic<[], [llvm_constant_ptr_ty]>;
foreach eviction_priority = ["evict_normal", "evict_last"] in
- def int_nvvm_prefetch_global_L2_ # eviction_priority : Intrinsic<[], [llvm_global_ptr_ty]>;
+ def int_nvvm_prefetch_global_L2_ # eviction_priority : DefaultAttrsIntrinsic<[], [llvm_global_ptr_ty]>;
- def int_nvvm_prefetchu_L1 : Intrinsic<[], [llvm_ptr_ty]>;
+ def int_nvvm_prefetchu_L1 : DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
}
// applypriority
diff --git a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
index a4ee24aaf4ce3..58990bfc1f1a1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
+++ b/llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
@@ -746,38 +746,25 @@ foreach dim = [1, 2, 3, 4, 5] in {
//Prefetch and Prefetchu
-class PREFETCH_INTRS<string InstName> :
+class PREFETCH_INTRS<string InstName, string IntrName> :
BasicNVPTXInst<(outs), (ins ADDR:$addr),
InstName,
- [(!cast<Intrinsic>(!strconcat("int_nvvm_",
- !subst(".", "_", InstName))) addr:$addr)]>,
+ [(!cast<Intrinsic>(IntrName) addr:$addr)]>,
Requires<[hasPTX<80>, hasSM<90>]>;
-def PREFETCH_L1 : PREFETCH_INTRS<"prefetch.L1">;
-def PREFETCH_L2 : PREFETCH_INTRS<"prefetch.L2">;
-def PREFETCH_GLOBAL_L1 : PREFETCH_INTRS<"prefetch.global.L1">;
-def PREFETCH_LOCAL_L1 : PREFETCH_INTRS<"prefetch.local.L1">;
-def PREFETCH_GLOBAL_L2 : PREFETCH_INTRS<"prefetch.global.L2">;
-def PREFETCH_LOCAL_L2 : PREFETCH_INTRS<"prefetch.local.L2">;
-def PREFETCH_CONST_TENSORMAP : PREFETCH_INTRS<"prefetch.const.tensormap">;
-def PREFETCH_GENERIC_TENSORMAP : BasicNVPTXInst<(outs), (ins ADDR:$addr),
- "prefetch.tensormap",
- [(int_nvvm_prefetch_generic_tensormap addr:$addr)]>,
- Requires<[hasPTX<80>, hasSM<90>]>;
-
-def PREFETCH_GLOBAL_L2_EVICT_NORMAL : BasicNVPTXInst<(outs), (ins ADDR:$addr),
- "prefetch.global.L2::evict_normal",
- [(int_nvvm_prefetch_global_L2_evict_normal addr:$addr)]>,
- Requires<[hasPTX<80>, hasSM<90>]>;
-
-def PREFETCH_GLOBAL_L2_EVICT_LAST : BasicNVPTXInst<(outs), (ins ADDR:$addr),
- "prefetch.global.L2::evict_last",
- [(int_nvvm_prefetch_global_L2_evict_last addr:$addr)]>,
- Requires<[hasPTX<80>, hasSM<90>]>;
-
-
-def PREFETCHU_L1 : PREFETCH_INTRS<"prefetchu.L1">;
+def PREFETCH_L1 : PREFETCH_INTRS<"prefetch.L1", "int_nvvm_prefetch_L1">;
+def PREFETCH_L2 : PREFETCH_INTRS<"prefetch.L2", "int_nvvm_prefetch_L2">;
+def PREFETCH_GLOBAL_L1 : PREFETCH_INTRS<"prefetch.global.L1", "int_nvvm_prefetch_global_L1">;
+def PREFETCH_LOCAL_L1 : PREFETCH_INTRS<"prefetch.local.L1", "int_nvvm_prefetch_local_L1">;
+def PREFETCH_GLOBAL_L2 : PREFETCH_INTRS<"prefetch.global.L2", "int_nvvm_prefetch_global_L2">;
+def PREFETCH_LOCAL_L2 : PREFETCH_INTRS<"prefetch.local.L2", "int_nvvm_prefetch_local_L2">;
+def PREFETCH_CONST_TENSORMAP : PREFETCH_INTRS<"prefetch.const.tensormap", "int_nvvm_prefetch_const_tensormap">;
+def PREFETCH_GENERIC_TENSORMAP : PREFETCH_INTRS<"prefetch.tensormap", "int_nvvm_prefetch_generic_tensormap">;
+def PREFETCH_GLOBAL_L2_EVICT_NORMAL : PREFETCH_INTRS<"prefetch.global.L2::evict_normal", "int_nvvm_prefetch_global_L2_evict_normal">;
+def PREFETCH_GLOBAL_L2_EVICT_LAST : PREFETCH_INTRS<"prefetch.global.L2::evict_last", "int_nvvm_prefetch_global_L2_evict_last">;
+
+def PREFETCHU_L1 : PREFETCH_INTRS<"prefetchu.L1", "int_nvvm_prefetchu_L1">;
//Applypriority intrinsics
class APPLYPRIORITY_L2_INTRS<string addrspace> :
More information about the llvm-commits
mailing list