Skip to content

Commit a2a79a7

Browse files
Prince781Prince781
committed
[NVPTX] update tests for mov.b32 canonicalization
1 parent 43f0f49 commit a2a79a7

File tree

1 file changed

+36
-36
lines changed

1 file changed

+36
-36
lines changed

‎llvm/test/CodeGen/NVPTX/f32x2-instructions.ll‎

Lines changed: 36 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -18,8 +18,8 @@ define <2 x float> @test_ret_const() #0 {
1818
; CHECK-NEXT: .reg .f32 %f<3>;
1919
; CHECK-EMPTY:
2020
; CHECK-NEXT: // %bb.0:
21-
; CHECK-NEXT: mov.f32 %f1, 0f40000000;
22-
; CHECK-NEXT: mov.f32 %f2, 0f3F800000;
21+
; CHECK-NEXT: mov.b32 %f1, 0f40000000;
22+
; CHECK-NEXT: mov.b32 %f2, 0f3F800000;
2323
; CHECK-NEXT: st.param.v2.f32 [func_retval0], {%f2, %f1};
2424
; CHECK-NEXT: ret;
2525
ret <2 x float> <float 1.0, float 2.0>
@@ -86,8 +86,8 @@ define <2 x float> @test_fadd_imm_0(<2 x float> %a) #0 {
8686
; CHECK-EMPTY:
8787
; CHECK-NEXT: // %bb.0:
8888
; CHECK-NEXT: ld.param.b64 %rd1, [test_fadd_imm_0_param_0];
89-
; CHECK-NEXT: mov.f32 %f1, 0f40000000;
90-
; CHECK-NEXT: mov.f32 %f2, 0f3F800000;
89+
; CHECK-NEXT: mov.b32 %f1, 0f40000000;
90+
; CHECK-NEXT: mov.b32 %f2, 0f3F800000;
9191
; CHECK-NEXT: mov.b64 %rd2, {%f2, %f1};
9292
; CHECK-NEXT: add.rn.f32x2 %rd3, %rd1, %rd2;
9393
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
@@ -104,8 +104,8 @@ define <2 x float> @test_fadd_imm_1(<2 x float> %a) #0 {
104104
; CHECK-EMPTY:
105105
; CHECK-NEXT: // %bb.0:
106106
; CHECK-NEXT: ld.param.b64 %rd1, [test_fadd_imm_1_param_0];
107-
; CHECK-NEXT: mov.f32 %f1, 0f40000000;
108-
; CHECK-NEXT: mov.f32 %f2, 0f3F800000;
107+
; CHECK-NEXT: mov.b32 %f1, 0f40000000;
108+
; CHECK-NEXT: mov.b32 %f2, 0f3F800000;
109109
; CHECK-NEXT: mov.b64 %rd2, {%f2, %f1};
110110
; CHECK-NEXT: add.rn.f32x2 %rd3, %rd1, %rd2;
111111
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
@@ -138,12 +138,12 @@ define <4 x float> @test_fadd_imm_0_v4(<4 x float> %a) #0 {
138138
; CHECK-EMPTY:
139139
; CHECK-NEXT: // %bb.0:
140140
; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_0_v4_param_0];
141-
; CHECK-NEXT: mov.f32 %f1, 0f40800000;
142-
; CHECK-NEXT: mov.f32 %f2, 0f40400000;
141+
; CHECK-NEXT: mov.b32 %f1, 0f40800000;
142+
; CHECK-NEXT: mov.b32 %f2, 0f40400000;
143143
; CHECK-NEXT: mov.b64 %rd3, {%f2, %f1};
144144
; CHECK-NEXT: add.rn.f32x2 %rd4, %rd2, %rd3;
145-
; CHECK-NEXT: mov.f32 %f3, 0f40000000;
146-
; CHECK-NEXT: mov.f32 %f4, 0f3F800000;
145+
; CHECK-NEXT: mov.b32 %f3, 0f40000000;
146+
; CHECK-NEXT: mov.b32 %f4, 0f3F800000;
147147
; CHECK-NEXT: mov.b64 %rd5, {%f4, %f3};
148148
; CHECK-NEXT: add.rn.f32x2 %rd6, %rd1, %rd5;
149149
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd4};
@@ -160,12 +160,12 @@ define <4 x float> @test_fadd_imm_1_v4(<4 x float> %a) #0 {
160160
; CHECK-EMPTY:
161161
; CHECK-NEXT: // %bb.0:
162162
; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_1_v4_param_0];
163-
; CHECK-NEXT: mov.f32 %f1, 0f40800000;
164-
; CHECK-NEXT: mov.f32 %f2, 0f40400000;
163+
; CHECK-NEXT: mov.b32 %f1, 0f40800000;
164+
; CHECK-NEXT: mov.b32 %f2, 0f40400000;
165165
; CHECK-NEXT: mov.b64 %rd3, {%f2, %f1};
166166
; CHECK-NEXT: add.rn.f32x2 %rd4, %rd2, %rd3;
167-
; CHECK-NEXT: mov.f32 %f3, 0f40000000;
168-
; CHECK-NEXT: mov.f32 %f4, 0f3F800000;
167+
; CHECK-NEXT: mov.b32 %f3, 0f40000000;
168+
; CHECK-NEXT: mov.b32 %f4, 0f3F800000;
169169
; CHECK-NEXT: mov.b64 %rd5, {%f4, %f3};
170170
; CHECK-NEXT: add.rn.f32x2 %rd6, %rd1, %rd5;
171171
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd4};
@@ -197,7 +197,7 @@ define <2 x float> @test_fneg(<2 x float> %a) #0 {
197197
; CHECK-EMPTY:
198198
; CHECK-NEXT: // %bb.0:
199199
; CHECK-NEXT: ld.param.b64 %rd1, [test_fneg_param_0];
200-
; CHECK-NEXT: mov.f32 %f1, 0f00000000;
200+
; CHECK-NEXT: mov.b32 %f1, 0f00000000;
201201
; CHECK-NEXT: mov.b64 %rd2, {%f1, %f1};
202202
; CHECK-NEXT: sub.rn.f32x2 %rd3, %rd2, %rd1;
203203
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
@@ -270,14 +270,14 @@ define <2 x float> @test_frem(<2 x float> %a, <2 x float> %b) #0 {
270270
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
271271
; CHECK-NEXT: div.rn.f32 %f5, %f4, %f2;
272272
; CHECK-NEXT: cvt.rzi.f32.f32 %f6, %f5;
273-
; CHECK-NEXT: mul.f32 %f7, %f6, %f2;
274-
; CHECK-NEXT: sub.f32 %f8, %f4, %f7;
273+
; CHECK-NEXT: neg.f32 %f7, %f6;
274+
; CHECK-NEXT: fma.rn.f32 %f8, %f7, %f2, %f4;
275275
; CHECK-NEXT: testp.infinite.f32 %p1, %f2;
276276
; CHECK-NEXT: selp.f32 %f9, %f4, %f8, %p1;
277277
; CHECK-NEXT: div.rn.f32 %f10, %f3, %f1;
278278
; CHECK-NEXT: cvt.rzi.f32.f32 %f11, %f10;
279-
; CHECK-NEXT: mul.f32 %f12, %f11, %f1;
280-
; CHECK-NEXT: sub.f32 %f13, %f3, %f12;
279+
; CHECK-NEXT: neg.f32 %f12, %f11;
280+
; CHECK-NEXT: fma.rn.f32 %f13, %f12, %f1, %f3;
281281
; CHECK-NEXT: testp.infinite.f32 %p2, %f1;
282282
; CHECK-NEXT: selp.f32 %f14, %f3, %f13, %p2;
283283
; CHECK-NEXT: st.param.v2.f32 [func_retval0], {%f14, %f9};
@@ -309,8 +309,8 @@ define <2 x float> @test_fadd_imm_0_ftz(<2 x float> %a) #2 {
309309
; CHECK-EMPTY:
310310
; CHECK-NEXT: // %bb.0:
311311
; CHECK-NEXT: ld.param.b64 %rd1, [test_fadd_imm_0_ftz_param_0];
312-
; CHECK-NEXT: mov.f32 %f1, 0f40000000;
313-
; CHECK-NEXT: mov.f32 %f2, 0f3F800000;
312+
; CHECK-NEXT: mov.b32 %f1, 0f40000000;
313+
; CHECK-NEXT: mov.b32 %f2, 0f3F800000;
314314
; CHECK-NEXT: mov.b64 %rd2, {%f2, %f1};
315315
; CHECK-NEXT: add.rn.ftz.f32x2 %rd3, %rd1, %rd2;
316316
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
@@ -327,8 +327,8 @@ define <2 x float> @test_fadd_imm_1_ftz(<2 x float> %a) #2 {
327327
; CHECK-EMPTY:
328328
; CHECK-NEXT: // %bb.0:
329329
; CHECK-NEXT: ld.param.b64 %rd1, [test_fadd_imm_1_ftz_param_0];
330-
; CHECK-NEXT: mov.f32 %f1, 0f40000000;
331-
; CHECK-NEXT: mov.f32 %f2, 0f3F800000;
330+
; CHECK-NEXT: mov.b32 %f1, 0f40000000;
331+
; CHECK-NEXT: mov.b32 %f2, 0f3F800000;
332332
; CHECK-NEXT: mov.b64 %rd2, {%f2, %f1};
333333
; CHECK-NEXT: add.rn.ftz.f32x2 %rd3, %rd1, %rd2;
334334
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
@@ -361,12 +361,12 @@ define <4 x float> @test_fadd_imm_0_v4_ftz(<4 x float> %a) #2 {
361361
; CHECK-EMPTY:
362362
; CHECK-NEXT: // %bb.0:
363363
; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_0_v4_ftz_param_0];
364-
; CHECK-NEXT: mov.f32 %f1, 0f40800000;
365-
; CHECK-NEXT: mov.f32 %f2, 0f40400000;
364+
; CHECK-NEXT: mov.b32 %f1, 0f40800000;
365+
; CHECK-NEXT: mov.b32 %f2, 0f40400000;
366366
; CHECK-NEXT: mov.b64 %rd3, {%f2, %f1};
367367
; CHECK-NEXT: add.rn.ftz.f32x2 %rd4, %rd2, %rd3;
368-
; CHECK-NEXT: mov.f32 %f3, 0f40000000;
369-
; CHECK-NEXT: mov.f32 %f4, 0f3F800000;
368+
; CHECK-NEXT: mov.b32 %f3, 0f40000000;
369+
; CHECK-NEXT: mov.b32 %f4, 0f3F800000;
370370
; CHECK-NEXT: mov.b64 %rd5, {%f4, %f3};
371371
; CHECK-NEXT: add.rn.ftz.f32x2 %rd6, %rd1, %rd5;
372372
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd4};
@@ -383,12 +383,12 @@ define <4 x float> @test_fadd_imm_1_v4_ftz(<4 x float> %a) #2 {
383383
; CHECK-EMPTY:
384384
; CHECK-NEXT: // %bb.0:
385385
; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_1_v4_ftz_param_0];
386-
; CHECK-NEXT: mov.f32 %f1, 0f40800000;
387-
; CHECK-NEXT: mov.f32 %f2, 0f40400000;
386+
; CHECK-NEXT: mov.b32 %f1, 0f40800000;
387+
; CHECK-NEXT: mov.b32 %f2, 0f40400000;
388388
; CHECK-NEXT: mov.b64 %rd3, {%f2, %f1};
389389
; CHECK-NEXT: add.rn.ftz.f32x2 %rd4, %rd2, %rd3;
390-
; CHECK-NEXT: mov.f32 %f3, 0f40000000;
391-
; CHECK-NEXT: mov.f32 %f4, 0f3F800000;
390+
; CHECK-NEXT: mov.b32 %f3, 0f40000000;
391+
; CHECK-NEXT: mov.b32 %f4, 0f3F800000;
392392
; CHECK-NEXT: mov.b64 %rd5, {%f4, %f3};
393393
; CHECK-NEXT: add.rn.ftz.f32x2 %rd6, %rd1, %rd5;
394394
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd6, %rd4};
@@ -420,7 +420,7 @@ define <2 x float> @test_fneg_ftz(<2 x float> %a) #2 {
420420
; CHECK-EMPTY:
421421
; CHECK-NEXT: // %bb.0:
422422
; CHECK-NEXT: ld.param.b64 %rd1, [test_fneg_ftz_param_0];
423-
; CHECK-NEXT: mov.f32 %f1, 0f00000000;
423+
; CHECK-NEXT: mov.b32 %f1, 0f00000000;
424424
; CHECK-NEXT: mov.b64 %rd2, {%f1, %f1};
425425
; CHECK-NEXT: sub.rn.ftz.f32x2 %rd3, %rd2, %rd1;
426426
; CHECK-NEXT: st.param.b64 [func_retval0], %rd3;
@@ -493,14 +493,14 @@ define <2 x float> @test_frem_ftz(<2 x float> %a, <2 x float> %b) #2 {
493493
; CHECK-NEXT: mov.b64 {%f3, %f4}, %rd1;
494494
; CHECK-NEXT: div.rn.ftz.f32 %f5, %f4, %f2;
495495
; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %f6, %f5;
496-
; CHECK-NEXT: mul.ftz.f32 %f7, %f6, %f2;
497-
; CHECK-NEXT: sub.ftz.f32 %f8, %f4, %f7;
496+
; CHECK-NEXT: neg.ftz.f32 %f7, %f6;
497+
; CHECK-NEXT: fma.rn.ftz.f32 %f8, %f7, %f2, %f4;
498498
; CHECK-NEXT: testp.infinite.f32 %p1, %f2;
499499
; CHECK-NEXT: selp.f32 %f9, %f4, %f8, %p1;
500500
; CHECK-NEXT: div.rn.ftz.f32 %f10, %f3, %f1;
501501
; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %f11, %f10;
502-
; CHECK-NEXT: mul.ftz.f32 %f12, %f11, %f1;
503-
; CHECK-NEXT: sub.ftz.f32 %f13, %f3, %f12;
502+
; CHECK-NEXT: neg.ftz.f32 %f12, %f11;
503+
; CHECK-NEXT: fma.rn.ftz.f32 %f13, %f12, %f1, %f3;
504504
; CHECK-NEXT: testp.infinite.f32 %p2, %f1;
505505
; CHECK-NEXT: selp.f32 %f14, %f3, %f13, %p2;
506506
; CHECK-NEXT: st.param.v2.f32 [func_retval0], {%f14, %f9};

0 commit comments

Comments
 (0)