@@ -72,21 +72,24 @@ define void @grid_const_escape(ptr byval(%struct.s) align 4 %input) {
72
72
; PTX-LABEL: grid_const_escape(
73
73
; PTX: {
74
74
; PTX-NEXT: .reg .b32 %r<3>;
75
- ; PTX-NEXT: .reg .b64 %rd<4 >;
75
+ ; PTX-NEXT: .reg .b64 %rd<5 >;
76
76
; PTX-EMPTY:
77
77
; PTX-NEXT: // %bb.0:
78
- ; PTX-NEXT: mov.b64 %rd1, grid_const_escape_param_0;
79
- ; PTX-NEXT: mov.u64 %rd2, %rd1;
80
- ; PTX-NEXT: cvta.param.u64 %rd3, %rd2;
78
+ ; PTX-NEXT: mov.b64 %rd2, grid_const_escape_param_0;
79
+ ; PTX-NEXT: mov.u64 %rd3, %rd2;
80
+ ; PTX-NEXT: cvta.param.u64 %rd4, %rd3;
81
+ ; PTX-NEXT: mov.u64 %rd1, escape;
81
82
; PTX-NEXT: { // callseq 0, 0
82
83
; PTX-NEXT: .param .b64 param0;
83
- ; PTX-NEXT: st.param.b64 [param0+0], %rd3 ;
84
+ ; PTX-NEXT: st.param.b64 [param0+0], %rd4 ;
84
85
; PTX-NEXT: .param .b32 retval0;
85
- ; PTX-NEXT: call.uni (retval0),
86
- ; PTX-NEXT: escape,
86
+ ; PTX-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);
87
+ ; PTX-NEXT: call (retval0),
88
+ ; PTX-NEXT: %rd1,
87
89
; PTX-NEXT: (
88
90
; PTX-NEXT: param0
89
- ; PTX-NEXT: );
91
+ ; PTX-NEXT: )
92
+ ; PTX-NEXT: , prototype_0;
90
93
; PTX-NEXT: ld.param.b32 %r1, [retval0+0];
91
94
; PTX-NEXT: } // callseq 0
92
95
; PTX-NEXT: ret;
@@ -107,36 +110,39 @@ define void @multiple_grid_const_escape(ptr byval(%struct.s) align 4 %input, i32
107
110
; PTX-NEXT: .reg .b64 %SP;
108
111
; PTX-NEXT: .reg .b64 %SPL;
109
112
; PTX-NEXT: .reg .b32 %r<4>;
110
- ; PTX-NEXT: .reg .b64 %rd<9 >;
113
+ ; PTX-NEXT: .reg .b64 %rd<10 >;
111
114
; PTX-EMPTY:
112
115
; PTX-NEXT: // %bb.0:
113
116
; PTX-NEXT: mov.u64 %SPL, __local_depot3;
114
117
; PTX-NEXT: cvta.local.u64 %SP, %SPL;
115
- ; PTX-NEXT: mov.b64 %rd1 , multiple_grid_const_escape_param_0;
116
- ; PTX-NEXT: mov.b64 %rd2 , multiple_grid_const_escape_param_2;
117
- ; PTX-NEXT: mov.u64 %rd3 , %rd2 ;
118
+ ; PTX-NEXT: mov.b64 %rd2 , multiple_grid_const_escape_param_0;
119
+ ; PTX-NEXT: mov.b64 %rd3 , multiple_grid_const_escape_param_2;
120
+ ; PTX-NEXT: mov.u64 %rd4 , %rd3 ;
118
121
; PTX-NEXT: ld.param.u32 %r1, [multiple_grid_const_escape_param_1];
119
- ; PTX-NEXT: cvta.param.u64 %rd4, %rd3;
120
- ; PTX-NEXT: mov.u64 %rd5, %rd1;
121
- ; PTX-NEXT: cvta.param.u64 %rd6, %rd5;
122
- ; PTX-NEXT: add.u64 %rd7, %SP, 0;
123
- ; PTX-NEXT: add.u64 %rd8, %SPL, 0;
124
- ; PTX-NEXT: st.local.u32 [%rd8], %r1;
122
+ ; PTX-NEXT: cvta.param.u64 %rd5, %rd4;
123
+ ; PTX-NEXT: mov.u64 %rd6, %rd2;
124
+ ; PTX-NEXT: cvta.param.u64 %rd7, %rd6;
125
+ ; PTX-NEXT: add.u64 %rd8, %SP, 0;
126
+ ; PTX-NEXT: add.u64 %rd9, %SPL, 0;
127
+ ; PTX-NEXT: st.local.u32 [%rd9], %r1;
128
+ ; PTX-NEXT: mov.u64 %rd1, escape3;
125
129
; PTX-NEXT: { // callseq 1, 0
126
130
; PTX-NEXT: .param .b64 param0;
127
- ; PTX-NEXT: st.param.b64 [param0+0], %rd6 ;
131
+ ; PTX-NEXT: st.param.b64 [param0+0], %rd7 ;
128
132
; PTX-NEXT: .param .b64 param1;
129
- ; PTX-NEXT: st.param.b64 [param1+0], %rd7 ;
133
+ ; PTX-NEXT: st.param.b64 [param1+0], %rd8 ;
130
134
; PTX-NEXT: .param .b64 param2;
131
- ; PTX-NEXT: st.param.b64 [param2+0], %rd4 ;
135
+ ; PTX-NEXT: st.param.b64 [param2+0], %rd5 ;
132
136
; PTX-NEXT: .param .b32 retval0;
133
- ; PTX-NEXT: call.uni (retval0),
134
- ; PTX-NEXT: escape3,
137
+ ; PTX-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);
138
+ ; PTX-NEXT: call (retval0),
139
+ ; PTX-NEXT: %rd1,
135
140
; PTX-NEXT: (
136
141
; PTX-NEXT: param0,
137
142
; PTX-NEXT: param1,
138
143
; PTX-NEXT: param2
139
- ; PTX-NEXT: );
144
+ ; PTX-NEXT: )
145
+ ; PTX-NEXT: , prototype_1;
140
146
; PTX-NEXT: ld.param.b32 %r2, [retval0+0];
141
147
; PTX-NEXT: } // callseq 1
142
148
; PTX-NEXT: ret;
@@ -221,26 +227,29 @@ define void @grid_const_partial_escape(ptr byval(i32) %input, ptr %output) {
221
227
; PTX-LABEL: grid_const_partial_escape(
222
228
; PTX: {
223
229
; PTX-NEXT: .reg .b32 %r<5>;
224
- ; PTX-NEXT: .reg .b64 %rd<6 >;
230
+ ; PTX-NEXT: .reg .b64 %rd<7 >;
225
231
; PTX-EMPTY:
226
232
; PTX-NEXT: // %bb.0:
227
- ; PTX-NEXT: mov.b64 %rd1 , grid_const_partial_escape_param_0;
228
- ; PTX-NEXT: ld.param.u64 %rd2 , [grid_const_partial_escape_param_1];
229
- ; PTX-NEXT: cvta.to.global.u64 %rd3 , %rd2 ;
230
- ; PTX-NEXT: mov.u64 %rd4 , %rd1 ;
231
- ; PTX-NEXT: cvta.param.u64 %rd5 , %rd4 ;
232
- ; PTX-NEXT: ld.u32 %r1, [%rd5 ];
233
+ ; PTX-NEXT: mov.b64 %rd2 , grid_const_partial_escape_param_0;
234
+ ; PTX-NEXT: ld.param.u64 %rd3 , [grid_const_partial_escape_param_1];
235
+ ; PTX-NEXT: cvta.to.global.u64 %rd4 , %rd3 ;
236
+ ; PTX-NEXT: mov.u64 %rd5 , %rd2 ;
237
+ ; PTX-NEXT: cvta.param.u64 %rd6 , %rd5 ;
238
+ ; PTX-NEXT: ld.u32 %r1, [%rd6 ];
233
239
; PTX-NEXT: add.s32 %r2, %r1, %r1;
234
- ; PTX-NEXT: st.global.u32 [%rd3], %r2;
240
+ ; PTX-NEXT: st.global.u32 [%rd4], %r2;
241
+ ; PTX-NEXT: mov.u64 %rd1, escape;
235
242
; PTX-NEXT: { // callseq 2, 0
236
243
; PTX-NEXT: .param .b64 param0;
237
- ; PTX-NEXT: st.param.b64 [param0+0], %rd5 ;
244
+ ; PTX-NEXT: st.param.b64 [param0+0], %rd6 ;
238
245
; PTX-NEXT: .param .b32 retval0;
239
- ; PTX-NEXT: call.uni (retval0),
240
- ; PTX-NEXT: escape,
246
+ ; PTX-NEXT: prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);
247
+ ; PTX-NEXT: call (retval0),
248
+ ; PTX-NEXT: %rd1,
241
249
; PTX-NEXT: (
242
250
; PTX-NEXT: param0
243
- ; PTX-NEXT: );
251
+ ; PTX-NEXT: )
252
+ ; PTX-NEXT: , prototype_2;
244
253
; PTX-NEXT: ld.param.b32 %r3, [retval0+0];
245
254
; PTX-NEXT: } // callseq 2
246
255
; PTX-NEXT: ret;
@@ -266,27 +275,30 @@ define i32 @grid_const_partial_escapemem(ptr byval(%struct.s) %input, ptr %outpu
266
275
; PTX-LABEL: grid_const_partial_escapemem(
267
276
; PTX: {
268
277
; PTX-NEXT: .reg .b32 %r<6>;
269
- ; PTX-NEXT: .reg .b64 %rd<6 >;
278
+ ; PTX-NEXT: .reg .b64 %rd<7 >;
270
279
; PTX-EMPTY:
271
280
; PTX-NEXT: // %bb.0:
272
- ; PTX-NEXT: mov.b64 %rd1 , grid_const_partial_escapemem_param_0;
273
- ; PTX-NEXT: ld.param.u64 %rd2 , [grid_const_partial_escapemem_param_1];
274
- ; PTX-NEXT: cvta.to.global.u64 %rd3 , %rd2 ;
275
- ; PTX-NEXT: mov.u64 %rd4 , %rd1 ;
276
- ; PTX-NEXT: cvta.param.u64 %rd5 , %rd4 ;
277
- ; PTX-NEXT: ld.u32 %r1, [%rd5 ];
278
- ; PTX-NEXT: ld.u32 %r2, [%rd5 +4];
279
- ; PTX-NEXT: st.global.u64 [%rd3 ], %rd5 ;
281
+ ; PTX-NEXT: mov.b64 %rd2 , grid_const_partial_escapemem_param_0;
282
+ ; PTX-NEXT: ld.param.u64 %rd3 , [grid_const_partial_escapemem_param_1];
283
+ ; PTX-NEXT: cvta.to.global.u64 %rd4 , %rd3 ;
284
+ ; PTX-NEXT: mov.u64 %rd5 , %rd2 ;
285
+ ; PTX-NEXT: cvta.param.u64 %rd6 , %rd5 ;
286
+ ; PTX-NEXT: ld.u32 %r1, [%rd6 ];
287
+ ; PTX-NEXT: ld.u32 %r2, [%rd6 +4];
288
+ ; PTX-NEXT: st.global.u64 [%rd4 ], %rd6 ;
280
289
; PTX-NEXT: add.s32 %r3, %r1, %r2;
290
+ ; PTX-NEXT: mov.u64 %rd1, escape;
281
291
; PTX-NEXT: { // callseq 3, 0
282
292
; PTX-NEXT: .param .b64 param0;
283
- ; PTX-NEXT: st.param.b64 [param0+0], %rd5 ;
293
+ ; PTX-NEXT: st.param.b64 [param0+0], %rd6 ;
284
294
; PTX-NEXT: .param .b32 retval0;
285
- ; PTX-NEXT: call.uni (retval0),
286
- ; PTX-NEXT: escape,
295
+ ; PTX-NEXT: prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);
296
+ ; PTX-NEXT: call (retval0),
297
+ ; PTX-NEXT: %rd1,
287
298
; PTX-NEXT: (
288
299
; PTX-NEXT: param0
289
- ; PTX-NEXT: );
300
+ ; PTX-NEXT: )
301
+ ; PTX-NEXT: , prototype_3;
290
302
; PTX-NEXT: ld.param.b32 %r4, [retval0+0];
291
303
; PTX-NEXT: } // callseq 3
292
304
; PTX-NEXT: st.param.b32 [func_retval0+0], %r3;
0 commit comments