@@ -14,7 +14,7 @@ declare void @llvm.nvvm.prefetch.L2(ptr %ptr)
14
14
15
15
declare void @llvm.nvvm.prefetch.tensormap.p0 (ptr %ptr )
16
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 )
17
+ declare void @llvm.nvvm.prefetch.tensormap.p101 (ptr addrspace (101 ) %param_ptr )
18
18
19
19
declare void @llvm.nvvm.prefetch.global.L2.evict.normal (ptr addrspace (1 ) %global_ptr )
20
20
declare void @llvm.nvvm.prefetch.global.L2.evict.last (ptr addrspace (1 ) %global_ptr )
@@ -84,17 +84,16 @@ define void @prefetchu_l1(ptr %ptr) {
84
84
ret void
85
85
}
86
86
87
-
88
87
define void @prefetch_tensormap (ptr %ptr ) {
89
88
; CHECK-PTX64-LABEL: prefetch_tensormap(
90
89
; CHECK-PTX64: {
91
90
; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
92
91
; CHECK-PTX64-EMPTY:
93
92
; CHECK-PTX64-NEXT: // %bb.0:
94
- ; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_generic_tensormap_param_0 ];
93
+ ; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_tensormap_param_0 ];
95
94
; CHECK-PTX64-NEXT: prefetch.tensormap [%rd1];
96
95
; CHECK-PTX64-NEXT: ret;
97
- tail call void @llvm.nvvm.prefetch.tensormap.p0 (ptr addrspace ( 0 ) %ptr )
96
+ tail call void @llvm.nvvm.prefetch.tensormap.p0 (ptr %ptr )
98
97
ret void
99
98
}
100
99
@@ -107,19 +106,19 @@ define void @prefetch_const_tensormap(ptr addrspace(4) %const_ptr) {
107
106
; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_const_tensormap_param_0];
108
107
; CHECK-PTX64-NEXT: prefetch.const.tensormap [%rd1];
109
108
; CHECK-PTX64-NEXT: ret;
110
- tail call void @llvm.nvvm.prefetch.tensormap.p4 (ptr addrspace (4 ) %const_ptr )
109
+ tail call void @llvm.nvvm.prefetch.tensormap.p4 (ptr addrspace (4 ) %const_ptr )
111
110
ret void
112
111
}
113
112
114
- define void @prefetch_param_tensormap (ptr addrspace (101 ) %const_ptr ) {
113
+ define void @prefetch_param_tensormap (ptr addrspace (101 ) %param_ptr ) {
115
114
; CHECK-PTX64-LABEL: prefetch_param_tensormap(
116
115
; CHECK-PTX64: {
117
116
; CHECK-PTX64-NEXT: .reg .b64 %rd<2>;
118
117
; CHECK-PTX64-EMPTY:
119
118
; CHECK-PTX64-NEXT: // %bb.0:
120
- ; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_const_tensormap_param_0 ];
119
+ ; CHECK-PTX64-NEXT: ld.param.b64 %rd1, [prefetch_param_tensormap_param_0 ];
121
120
; CHECK-PTX64-NEXT: prefetch.param.tensormap [%rd1];
122
121
; CHECK-PTX64-NEXT: ret;
123
- tail call void @llvm.nvvm.prefetch.tensormap.p101 (ptr addrspace (101 ) %const_ptr )
122
+ tail call void @llvm.nvvm.prefetch.tensormap.p101 (ptr addrspace (101 ) %param_ptr )
124
123
ret void
125
124
}
0 commit comments