Skip to content

Commit bab560b

Browse files
abhilash1910abhilash1910
committed
refresh with addrspace
1 parent 0d26914 commit bab560b

File tree

3 files changed

+79
-20
lines changed

3 files changed

+79
-20
lines changed

‎llvm/include/llvm/IR/IntrinsicsNVVM.td‎

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2093,8 +2093,7 @@ let IntrProperties = [IntrArgMemOnly, ReadOnly<ArgIndex<0>>, NoCapture<ArgIndex<
20932093
def int_nvvm_prefetch_local_ # level : DefaultAttrsIntrinsic<[], [llvm_local_ptr_ty]>;
20942094
}
20952095

2096-
def int_nvvm_prefetch_generic_tensormap : DefaultAttrsIntrinsic<[], [llvm_ptr_ty]>;
2097-
def int_nvvm_prefetch_const_tensormap : DefaultAttrsIntrinsic<[], [llvm_constant_ptr_ty]>;
2096+
def int_nvvm_prefetch_tensormap : DefaultAttrsIntrinsic<[], [llvm_anyptr_ty]>;
20982097

20992098
foreach eviction_priority = ["evict_normal", "evict_last"] in
21002099
def int_nvvm_prefetch_global_L2_ # eviction_priority : DefaultAttrsIntrinsic<[], [llvm_global_ptr_ty]>;

‎llvm/lib/Target/NVPTX/NVPTXIntrinsics.td‎

Lines changed: 53 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -39,6 +39,12 @@ def AS_match {
3939
code global = [{
4040
return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_GLOBAL);
4141
}];
42+
code const = [{
43+
return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_CONST);
44+
}];
45+
code param = [{
46+
return ChkMemSDNodeAddressSpace(N, llvm::ADDRESS_SPACE_PARAM);
47+
}];
4248
}
4349

4450
// A node that will be replaced with the current PTX version.
@@ -744,13 +750,57 @@ foreach dim = [1, 2, 3, 4, 5] in {
744750
}
745751
}
746752

747-
//Prefetchu and Prefetch
753+
//Prefetchu and Prefetch
754+
755+
class PREFETCH_CONST_CHK<dag frag>
756+
: PatFrag<!setdagop(frag, ops), frag, [{
757+
auto *Addr = N->getOperand(2).getNode();
758+
auto *MemNode = dyn_cast<MemSDNode>(Addr);
759+
bool result = MemNode->getMemOperand()->getAddrSpace() == llvm::ADDRESS_SPACE_CONST;
760+
return result;
761+
}]>;
762+
763+
764+
class PREFETCH_GENERIC_CHK<dag frag>
765+
: PatFrag<!setdagop(frag, ops), frag, [{
766+
auto *Addr = N->getOperand(2).getNode();
767+
auto *MemNode = dyn_cast<MemSDNode>(Addr);
768+
bool result= MemNode->getMemOperand()->getAddrSpace() == llvm::ADDRESS_SPACE_GENERIC;
769+
return result;
770+
}]>;
771+
772+
773+
class PREFETCH_PARAM_CHK<dag frag>
774+
: PatFrag<!setdagop(frag, ops), frag, [{
775+
auto *Addr = N->getOperand(2).getNode();
776+
auto *MemNode = dyn_cast<MemSDNode>(Addr);
777+
bool result = MemNode->getMemOperand()->getAddrSpace() == llvm::ADDRESS_SPACE_PARAM;
778+
return result;
779+
}]>;
748780

781+
defvar frag_pat = (int_nvvm_prefetch_tensormap node:$addr);
782+
783+
def prefetch_tensormap_const : PREFETCH_CONST_CHK<frag_pat>;
784+
def prefetch_tensormap_gen : PREFETCH_GENERIC_CHK<frag_pat>;
785+
def prefetch_tensormap_param : PREFETCH_PARAM_CHK<frag_pat>;
786+
787+
def PREFETCH_CONST_TENSORMAP : NVPTXInst<(outs), (ins ADDR:$addr),
788+
"prefetch.const.tensormap [$addr];",
789+
[(prefetch_tensormap_const addr:$addr)]>;
790+
791+
def PREFETCH_GENERIC_TENSORMAP : NVPTXInst<(outs), (ins ADDR:$addr),
792+
"prefetch.tensormap [$addr];",
793+
[(prefetch_tensormap_gen addr:$addr)]>;
794+
795+
def PREFETCH_PARAM_TENSORMAP : NVPTXInst<(outs), (ins ADDR:$addr),
796+
"prefetch.param.tensormap [$addr];",
797+
[(prefetch_tensormap_param addr:$addr)]>;
798+
799+
749800
class PREFETCH_INTRS<string InstName, Intrinsic Intr> :
750801
BasicNVPTXInst<(outs), (ins ADDR:$addr),
751802
InstName,
752-
[(Intr addr:$addr)]>,
753-
Requires<[hasPTX<80>, hasSM<90>]>;
803+
[(Intr addr:$addr)]>;
754804

755805
def PREFETCHU_L1 : PREFETCH_INTRS<"prefetchu.L1", int_nvvm_prefetchu_L1>;
756806
def PREFETCH_L1 : PREFETCH_INTRS<"prefetch.L1", int_nvvm_prefetch_L1>;
@@ -759,10 +809,6 @@ def PREFETCH_GLOBAL_L1 : PREFETCH_INTRS<"prefetch.global.L1", int_nvvm_prefetch_
759809
def PREFETCH_LOCAL_L1 : PREFETCH_INTRS<"prefetch.local.L1", int_nvvm_prefetch_local_L1>;
760810
def PREFETCH_GLOBAL_L2 : PREFETCH_INTRS<"prefetch.global.L2", int_nvvm_prefetch_global_L2>;
761811
def PREFETCH_LOCAL_L2 : PREFETCH_INTRS<"prefetch.local.L2", int_nvvm_prefetch_local_L2>;
762-
def PREFETCH_CONST_TENSORMAP : PREFETCH_INTRS<"prefetch.const.tensormap",
763-
int_nvvm_prefetch_const_tensormap>;
764-
def PREFETCH_GENERIC_TENSORMAP : PREFETCH_INTRS<"prefetch.tensormap",
765-
int_nvvm_prefetch_generic_tensormap>;
766812
def PREFETCH_GLOBAL_L2_EVICT_NORMAL : PREFETCH_INTRS<"prefetch.global.L2::evict_normal",
767813
int_nvvm_prefetch_global_L2_evict_normal>;
768814
def PREFETCH_GLOBAL_L2_EVICT_LAST : PREFETCH_INTRS<"prefetch.global.L2::evict_last",

‎llvm/test/CodeGen/NVPTX/prefetch.ll‎

Lines changed: 25 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -12,8 +12,9 @@ declare void @llvm.nvvm.prefetch.local.L2(ptr addrspace(5) %local_ptr)
1212
declare void @llvm.nvvm.prefetch.L1(ptr %ptr)
1313
declare void @llvm.nvvm.prefetch.L2(ptr %ptr)
1414

15-
declare void @llvm.nvvm.prefetch.generic.tensormap(ptr %ptr)
16-
declare void @llvm.nvvm.prefetch.const.tensormap(ptr addrspace(4) %const_ptr)
15+
declare void @llvm.nvvm.prefetch.tensormap.p0(ptr %ptr)
16+
declare void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
17+
declare void @llvm.nvvm.prefetch.tensormap.p101(ptr addrspace(101) %const_ptr)
1718

1819
declare void @llvm.nvvm.prefetch.global.L2.evict.normal(ptr addrspace(1) %global_ptr)
1920
declare void @llvm.nvvm.prefetch.global.L2.evict.last(ptr addrspace(1) %global_ptr)
@@ -70,17 +71,30 @@ define void @prefetch_(ptr %ptr) {
7071
ret void
7172
}
7273

74+
define void @prefetchu_l1(ptr %ptr) {
75+
; CHECK-PTX64-LABEL: prefetchu_l1(
76+
; CHECK-PTX64: {
77+
; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
78+
; CHECK-PTX64-EMPTY:
79+
; CHECK-PTX64-NEXT: // %bb.0:
80+
; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetchu_l1_param_0];
81+
; CHECK-PTX64-NEXT: prefetchu.L1 [%rd1];
82+
; CHECK-PTX64-NEXT: ret;
83+
tail call void @llvm.nvvm.prefetchu.L1(ptr %ptr)
84+
ret void
85+
}
86+
7387

74-
define void @prefetch_generic_tensormap(ptr %ptr) {
75-
; CHECK-PTX64-LABEL: prefetch_generic_tensormap(
88+
define void @prefetch_tensormap(ptr %ptr) {
89+
; CHECK-PTX64-LABEL: prefetch_tensormap(
7690
; CHECK-PTX64: {
7791
; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
7892
; CHECK-PTX64-EMPTY:
7993
; CHECK-PTX64-NEXT: // %bb.0:
8094
; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_generic_tensormap_param_0];
8195
; CHECK-PTX64-NEXT: prefetch.tensormap [%rd1];
8296
; CHECK-PTX64-NEXT: ret;
83-
tail call void @llvm.nvvm.prefetch.generic.tensormap(ptr %ptr)
97+
tail call void @llvm.nvvm.prefetch.tensormap.p0(ptr addrspace(0) %ptr)
8498
ret void
8599
}
86100

@@ -93,19 +107,19 @@ define void @prefetch_const_tensormap(ptr addrspace(4) %const_ptr) {
93107
; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_const_tensormap_param_0];
94108
; CHECK-PTX64-NEXT: prefetch.const.tensormap [%rd1];
95109
; CHECK-PTX64-NEXT: ret;
96-
tail call void @llvm.nvvm.prefetch.const.tensormap(ptr addrspace(4) %const_ptr)
110+
tail call void @llvm.nvvm.prefetch.tensormap.p4(ptr addrspace(4) %const_ptr)
97111
ret void
98112
}
99113

100-
define void @prefetchu_l1(ptr %ptr) {
101-
; CHECK-PTX64-LABEL: prefetchu_l1(
114+
define void @prefetch_param_tensormap(ptr addrspace(101) %const_ptr) {
115+
; CHECK-PTX64-LABEL: prefetch_param_tensormap(
102116
; CHECK-PTX64: {
103117
; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
104118
; CHECK-PTX64-EMPTY:
105119
; CHECK-PTX64-NEXT: // %bb.0:
106-
; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetchu_l1_param_0];
107-
; CHECK-PTX64-NEXT: prefetchu.L1 [%rd1];
120+
; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_const_tensormap_param_0];
121+
; CHECK-PTX64-NEXT: prefetch.param.tensormap [%rd1];
108122
; CHECK-PTX64-NEXT: ret;
109-
tail call void @llvm.nvvm.prefetchu.L1(ptr %ptr)
123+
tail call void @llvm.nvvm.prefetch.tensormap.p101(ptr addrspace(101) %const_ptr)
110124
ret void
111125
}

0 commit comments

Comments
 (0)