Skip to content

Commit 6028618

Browse files
tomnatan30tomnatan30
authored
Revert "[NVPTX] Add folding for cvt.rn.bf16x2.f32 (#116109)"
This reverts commit 90cbd4a.
1 parent 10b048c commit 6028618

File tree

5 files changed

+145
-247
lines changed

5 files changed

+145
-247
lines changed

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

Lines changed: 0 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -739,20 +739,6 @@ let hasSideEffects = false in {
739739
def CVT_f16x2_e5m2x2 : CVT_f16x2_fp8<"e5m2">;
740740
}
741741

742-
def fpround_oneuse : PatFrag<(ops node:$a), (fpround node:$a), [{
743-
return N->hasOneUse();
744-
}]>;
745-
746-
def : Pat<(v2bf16 (build_vector (bf16 (fpround_oneuse Float32Regs:$a)),
747-
(bf16 (fpround_oneuse Float32Regs:$b)))),
748-
(CVT_bf16x2_f32 Float32Regs:$a, Float32Regs:$b, CvtRN)>,
749-
Requires<[hasPTX<70>, hasSM<80>, hasBF16Math]>;
750-
751-
def : Pat<(v2f16 (build_vector (f16 (fpround_oneuse Float32Regs:$a)),
752-
(f16 (fpround_oneuse Float32Regs:$b)))),
753-
(CVT_f16x2_f32 Float32Regs:$a, Float32Regs:$b, CvtRN)>,
754-
Requires<[hasPTX<70>, hasSM<80>, useFP16Math]>;
755-
756742
//-----------------------------------
757743
// Selection instructions (selp)
758744
//-----------------------------------

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

Lines changed: 36 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -204,7 +204,7 @@ define <2 x bfloat> @test_faddx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
204204
;
205205
; SM80-LABEL: test_faddx2(
206206
; SM80: {
207-
; SM80-NEXT: .reg .b16 %rs<5>;
207+
; SM80-NEXT: .reg .b16 %rs<7>;
208208
; SM80-NEXT: .reg .b32 %r<4>;
209209
; SM80-NEXT: .reg .f32 %f<7>;
210210
; SM80-EMPTY:
@@ -216,16 +216,18 @@ define <2 x bfloat> @test_faddx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
216216
; SM80-NEXT: mov.b32 {%rs3, %rs4}, %r1;
217217
; SM80-NEXT: cvt.f32.bf16 %f2, %rs4;
218218
; SM80-NEXT: add.rn.f32 %f3, %f2, %f1;
219+
; SM80-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
219220
; SM80-NEXT: cvt.f32.bf16 %f4, %rs1;
220221
; SM80-NEXT: cvt.f32.bf16 %f5, %rs3;
221222
; SM80-NEXT: add.rn.f32 %f6, %f5, %f4;
222-
; SM80-NEXT: cvt.rn.bf16x2.f32 %r3, %f6, %f3;
223+
; SM80-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
224+
; SM80-NEXT: mov.b32 %r3, {%rs6, %rs5};
223225
; SM80-NEXT: st.param.b32 [func_retval0], %r3;
224226
; SM80-NEXT: ret;
225227
;
226228
; SM80-FTZ-LABEL: test_faddx2(
227229
; SM80-FTZ: {
228-
; SM80-FTZ-NEXT: .reg .b16 %rs<5>;
230+
; SM80-FTZ-NEXT: .reg .b16 %rs<7>;
229231
; SM80-FTZ-NEXT: .reg .b32 %r<4>;
230232
; SM80-FTZ-NEXT: .reg .f32 %f<7>;
231233
; SM80-FTZ-EMPTY:
@@ -237,10 +239,12 @@ define <2 x bfloat> @test_faddx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
237239
; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r1;
238240
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs4;
239241
; SM80-FTZ-NEXT: add.rn.ftz.f32 %f3, %f2, %f1;
242+
; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
240243
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f4, %rs1;
241244
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f5, %rs3;
242245
; SM80-FTZ-NEXT: add.rn.ftz.f32 %f6, %f5, %f4;
243-
; SM80-FTZ-NEXT: cvt.rn.bf16x2.f32 %r3, %f6, %f3;
246+
; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
247+
; SM80-FTZ-NEXT: mov.b32 %r3, {%rs6, %rs5};
244248
; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r3;
245249
; SM80-FTZ-NEXT: ret;
246250
;
@@ -307,7 +311,7 @@ define <2 x bfloat> @test_fsubx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
307311
;
308312
; SM80-LABEL: test_fsubx2(
309313
; SM80: {
310-
; SM80-NEXT: .reg .b16 %rs<5>;
314+
; SM80-NEXT: .reg .b16 %rs<7>;
311315
; SM80-NEXT: .reg .b32 %r<4>;
312316
; SM80-NEXT: .reg .f32 %f<7>;
313317
; SM80-EMPTY:
@@ -319,16 +323,18 @@ define <2 x bfloat> @test_fsubx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
319323
; SM80-NEXT: mov.b32 {%rs3, %rs4}, %r1;
320324
; SM80-NEXT: cvt.f32.bf16 %f2, %rs4;
321325
; SM80-NEXT: sub.rn.f32 %f3, %f2, %f1;
326+
; SM80-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
322327
; SM80-NEXT: cvt.f32.bf16 %f4, %rs1;
323328
; SM80-NEXT: cvt.f32.bf16 %f5, %rs3;
324329
; SM80-NEXT: sub.rn.f32 %f6, %f5, %f4;
325-
; SM80-NEXT: cvt.rn.bf16x2.f32 %r3, %f6, %f3;
330+
; SM80-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
331+
; SM80-NEXT: mov.b32 %r3, {%rs6, %rs5};
326332
; SM80-NEXT: st.param.b32 [func_retval0], %r3;
327333
; SM80-NEXT: ret;
328334
;
329335
; SM80-FTZ-LABEL: test_fsubx2(
330336
; SM80-FTZ: {
331-
; SM80-FTZ-NEXT: .reg .b16 %rs<5>;
337+
; SM80-FTZ-NEXT: .reg .b16 %rs<7>;
332338
; SM80-FTZ-NEXT: .reg .b32 %r<4>;
333339
; SM80-FTZ-NEXT: .reg .f32 %f<7>;
334340
; SM80-FTZ-EMPTY:
@@ -340,10 +346,12 @@ define <2 x bfloat> @test_fsubx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
340346
; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r1;
341347
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs4;
342348
; SM80-FTZ-NEXT: sub.rn.ftz.f32 %f3, %f2, %f1;
349+
; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
343350
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f4, %rs1;
344351
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f5, %rs3;
345352
; SM80-FTZ-NEXT: sub.rn.ftz.f32 %f6, %f5, %f4;
346-
; SM80-FTZ-NEXT: cvt.rn.bf16x2.f32 %r3, %f6, %f3;
353+
; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
354+
; SM80-FTZ-NEXT: mov.b32 %r3, {%rs6, %rs5};
347355
; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r3;
348356
; SM80-FTZ-NEXT: ret;
349357
;
@@ -410,7 +418,7 @@ define <2 x bfloat> @test_fmulx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
410418
;
411419
; SM80-LABEL: test_fmulx2(
412420
; SM80: {
413-
; SM80-NEXT: .reg .b16 %rs<5>;
421+
; SM80-NEXT: .reg .b16 %rs<7>;
414422
; SM80-NEXT: .reg .b32 %r<4>;
415423
; SM80-NEXT: .reg .f32 %f<7>;
416424
; SM80-EMPTY:
@@ -422,16 +430,18 @@ define <2 x bfloat> @test_fmulx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
422430
; SM80-NEXT: mov.b32 {%rs3, %rs4}, %r1;
423431
; SM80-NEXT: cvt.f32.bf16 %f2, %rs4;
424432
; SM80-NEXT: mul.rn.f32 %f3, %f2, %f1;
433+
; SM80-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
425434
; SM80-NEXT: cvt.f32.bf16 %f4, %rs1;
426435
; SM80-NEXT: cvt.f32.bf16 %f5, %rs3;
427436
; SM80-NEXT: mul.rn.f32 %f6, %f5, %f4;
428-
; SM80-NEXT: cvt.rn.bf16x2.f32 %r3, %f6, %f3;
437+
; SM80-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
438+
; SM80-NEXT: mov.b32 %r3, {%rs6, %rs5};
429439
; SM80-NEXT: st.param.b32 [func_retval0], %r3;
430440
; SM80-NEXT: ret;
431441
;
432442
; SM80-FTZ-LABEL: test_fmulx2(
433443
; SM80-FTZ: {
434-
; SM80-FTZ-NEXT: .reg .b16 %rs<5>;
444+
; SM80-FTZ-NEXT: .reg .b16 %rs<7>;
435445
; SM80-FTZ-NEXT: .reg .b32 %r<4>;
436446
; SM80-FTZ-NEXT: .reg .f32 %f<7>;
437447
; SM80-FTZ-EMPTY:
@@ -443,10 +453,12 @@ define <2 x bfloat> @test_fmulx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
443453
; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r1;
444454
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs4;
445455
; SM80-FTZ-NEXT: mul.rn.ftz.f32 %f3, %f2, %f1;
456+
; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
446457
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f4, %rs1;
447458
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f5, %rs3;
448459
; SM80-FTZ-NEXT: mul.rn.ftz.f32 %f6, %f5, %f4;
449-
; SM80-FTZ-NEXT: cvt.rn.bf16x2.f32 %r3, %f6, %f3;
460+
; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
461+
; SM80-FTZ-NEXT: mov.b32 %r3, {%rs6, %rs5};
450462
; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r3;
451463
; SM80-FTZ-NEXT: ret;
452464
;
@@ -513,7 +525,7 @@ define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
513525
;
514526
; SM80-LABEL: test_fdiv(
515527
; SM80: {
516-
; SM80-NEXT: .reg .b16 %rs<5>;
528+
; SM80-NEXT: .reg .b16 %rs<7>;
517529
; SM80-NEXT: .reg .b32 %r<4>;
518530
; SM80-NEXT: .reg .f32 %f<7>;
519531
; SM80-EMPTY:
@@ -525,16 +537,18 @@ define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
525537
; SM80-NEXT: mov.b32 {%rs3, %rs4}, %r1;
526538
; SM80-NEXT: cvt.f32.bf16 %f2, %rs4;
527539
; SM80-NEXT: div.rn.f32 %f3, %f2, %f1;
540+
; SM80-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
528541
; SM80-NEXT: cvt.f32.bf16 %f4, %rs1;
529542
; SM80-NEXT: cvt.f32.bf16 %f5, %rs3;
530543
; SM80-NEXT: div.rn.f32 %f6, %f5, %f4;
531-
; SM80-NEXT: cvt.rn.bf16x2.f32 %r3, %f6, %f3;
544+
; SM80-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
545+
; SM80-NEXT: mov.b32 %r3, {%rs6, %rs5};
532546
; SM80-NEXT: st.param.b32 [func_retval0], %r3;
533547
; SM80-NEXT: ret;
534548
;
535549
; SM80-FTZ-LABEL: test_fdiv(
536550
; SM80-FTZ: {
537-
; SM80-FTZ-NEXT: .reg .b16 %rs<5>;
551+
; SM80-FTZ-NEXT: .reg .b16 %rs<7>;
538552
; SM80-FTZ-NEXT: .reg .b32 %r<4>;
539553
; SM80-FTZ-NEXT: .reg .f32 %f<7>;
540554
; SM80-FTZ-EMPTY:
@@ -546,16 +560,18 @@ define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
546560
; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r1;
547561
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs4;
548562
; SM80-FTZ-NEXT: div.rn.ftz.f32 %f3, %f2, %f1;
563+
; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
549564
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f4, %rs1;
550565
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f5, %rs3;
551566
; SM80-FTZ-NEXT: div.rn.ftz.f32 %f6, %f5, %f4;
552-
; SM80-FTZ-NEXT: cvt.rn.bf16x2.f32 %r3, %f6, %f3;
567+
; SM80-FTZ-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
568+
; SM80-FTZ-NEXT: mov.b32 %r3, {%rs6, %rs5};
553569
; SM80-FTZ-NEXT: st.param.b32 [func_retval0], %r3;
554570
; SM80-FTZ-NEXT: ret;
555571
;
556572
; SM90-LABEL: test_fdiv(
557573
; SM90: {
558-
; SM90-NEXT: .reg .b16 %rs<5>;
574+
; SM90-NEXT: .reg .b16 %rs<7>;
559575
; SM90-NEXT: .reg .b32 %r<4>;
560576
; SM90-NEXT: .reg .f32 %f<7>;
561577
; SM90-EMPTY:
@@ -567,10 +583,12 @@ define <2 x bfloat> @test_fdiv(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
567583
; SM90-NEXT: mov.b32 {%rs3, %rs4}, %r1;
568584
; SM90-NEXT: cvt.f32.bf16 %f2, %rs4;
569585
; SM90-NEXT: div.rn.f32 %f3, %f2, %f1;
586+
; SM90-NEXT: cvt.rn.bf16.f32 %rs5, %f3;
570587
; SM90-NEXT: cvt.f32.bf16 %f4, %rs1;
571588
; SM90-NEXT: cvt.f32.bf16 %f5, %rs3;
572589
; SM90-NEXT: div.rn.f32 %f6, %f5, %f4;
573-
; SM90-NEXT: cvt.rn.bf16x2.f32 %r3, %f6, %f3;
590+
; SM90-NEXT: cvt.rn.bf16.f32 %rs6, %f6;
591+
; SM90-NEXT: mov.b32 %r3, {%rs6, %rs5};
574592
; SM90-NEXT: st.param.b32 [func_retval0], %r3;
575593
; SM90-NEXT: ret;
576594
%r = fdiv <2 x bfloat> %a, %b

‎llvm/test/CodeGen/NVPTX/bf16x2-instructions-approx.ll

Lines changed: 6 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -13,7 +13,9 @@ declare <2 x bfloat> @llvm.cos.f16(<2 x bfloat> %a) #0
1313
; CHECK-DAG: cvt.f32.bf16 [[AF1:%f[0-9]+]], [[A1]];
1414
; CHECK-DAG: sin.approx.f32 [[RF0:%f[0-9]+]], [[AF0]];
1515
; CHECK-DAG: sin.approx.f32 [[RF1:%f[0-9]+]], [[AF1]];
16-
; CHECK: cvt.rn.bf16x2.f32 [[R:%r[0-9]+]], [[RF0]], [[RF1]]
16+
; CHECK-DAG: cvt.rn.bf16.f32 [[R0:%rs[0-9]+]], [[RF0]];
17+
; CHECK-DAG: cvt.rn.bf16.f32 [[R1:%rs[0-9]+]], [[RF1]];
18+
; CHECK: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
1719
; CHECK: st.param.b32 [func_retval0], [[R]];
1820
; CHECK: ret;
1921
define <2 x bfloat> @test_sin(<2 x bfloat> %a) #0 #1 {
@@ -28,7 +30,9 @@ define <2 x bfloat> @test_sin(<2 x bfloat> %a) #0 #1 {
2830
; CHECK-DAG: cvt.f32.bf16 [[AF1:%f[0-9]+]], [[A1]];
2931
; CHECK-DAG: cos.approx.f32 [[RF0:%f[0-9]+]], [[AF0]];
3032
; CHECK-DAG: cos.approx.f32 [[RF1:%f[0-9]+]], [[AF1]];
31-
; CHECK: cvt.rn.bf16x2.f32 [[R:%r[0-9]+]], [[RF0]], [[RF1]]
33+
; CHECK-DAG: cvt.rn.bf16.f32 [[R0:%rs[0-9]+]], [[RF0]];
34+
; CHECK-DAG: cvt.rn.bf16.f32 [[R1:%rs[0-9]+]], [[RF1]];
35+
; CHECK: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
3236
; CHECK: st.param.b32 [func_retval0], [[R]];
3337
; CHECK: ret;
3438
define <2 x bfloat> @test_cos(<2 x bfloat> %a) #0 #1 {

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

Lines changed: 26 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -26,7 +26,9 @@ define <2 x bfloat> @test_ret_const() #0 {
2626
; SM80-DAG: cvt.f32.bf16 [[FA1:%f[0-9]+]], [[A1]]
2727
; SM80-DAG: add.rn.f32 [[FR0:%f[0-9]+]], [[FA0]], 0f3F800000;
2828
; SM80-DAG: add.rn.f32 [[FR1:%f[0-9]+]], [[FA1]], 0f40000000;
29-
; SM80-DAG: cvt.rn.bf16x2.f32 [[R:%r[0-9]+]], [[FR0]], [[FR1]];
29+
; SM80-DAG: cvt.rn.bf16.f32 [[R0:%rs[0-9]+]], [[FR0]]
30+
; SM80-DAG: cvt.rn.bf16.f32 [[R1:%rs[0-9]+]], [[FR1]]
31+
; SM80-DAG: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
3032
;
3133
; CHECK-NEXT: st.param.b32 [func_retval0], [[R]];
3234
; CHECK-NEXT: ret;
@@ -66,7 +68,9 @@ define bfloat @test_fadd_imm_1(bfloat %a) #0 {
6668
; SM80-DAG: cvt.f32.bf16 [[FB1:%f[0-9]+]], [[B1]];
6769
; SM80-DAG: sub.rn.f32 [[FR0:%f[0-9]+]], [[FA0]], [[FB0]];
6870
; SM80-DAG: sub.rn.f32 [[FR1:%f[0-9]+]], [[FA1]], [[FB1]];
69-
; SM80-DAG: cvt.rn.bf16x2.f32 [[R:%r[0-9]+]], [[FR0]], [[FR1]];
71+
; SM80-DAG: cvt.rn.bf16.f32 [[R0:%rs[0-9]+]], [[FR0]];
72+
; SM80-DAG: cvt.rn.bf16.f32 [[R1:%rs[0-9]+]], [[FR1]];
73+
; SM80: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]};
7074

7175
; CHECK: st.param.b32 [func_retval0], [[R]];
7276
; CHECK: ret;
@@ -89,7 +93,9 @@ define <2 x bfloat> @test_fsubx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
8993
; SM80-DAG: cvt.f32.bf16 [[FB1:%f[0-9]+]], [[B1]];
9094
; SM80-DAG: mul.rn.f32 [[FR0:%f[0-9]+]], [[FA0]], [[FB0]];
9195
; SM80-DAG: mul.rn.f32 [[FR1:%f[0-9]+]], [[FA1]], [[FB1]];
92-
; SM80-DAG: cvt.rn.bf16x2.f32 [[R:%r[0-9]+]], [[FR0]], [[FR1]];
96+
; SM80-DAG: cvt.rn.bf16.f32 [[R0:%rs[0-9]+]], [[FR0]];
97+
; SM80-DAG: cvt.rn.bf16.f32 [[R1:%rs[0-9]+]], [[FR1]];
98+
; SM80: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]};
9399

94100
; CHECK: st.param.b32 [func_retval0], [[R]];
95101
; CHECK: ret;
@@ -110,7 +116,9 @@ define <2 x bfloat> @test_fmulx2(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
110116
; CHECK-DAG: cvt.f32.bf16 [[FB1:%f[0-9]+]], [[B1]];
111117
; CHECK-DAG: div.rn.f32 [[FR0:%f[0-9]+]], [[FA0]], [[FB0]];
112118
; CHECK-DAG: div.rn.f32 [[FR1:%f[0-9]+]], [[FA1]], [[FB1]];
113-
; CHECK: cvt.rn.bf16x2.f32 [[R:%r[0-9]+]], [[FR0]], [[FR1]];
119+
; CHECK-DAG: cvt.rn.bf16.f32 [[R0:%rs[0-9]+]], [[FR0]];
120+
; CHECK-DAG: cvt.rn.bf16.f32 [[R1:%rs[0-9]+]], [[FR1]];
121+
; CHECK-NEXT: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
114122
; CHECK-NEXT: st.param.b32 [func_retval0], [[R]];
115123
; CHECK-NEXT: ret;
116124

@@ -279,7 +287,9 @@ define <2 x bfloat> @test_select_cc_bf16_f32(<2 x bfloat> %a, <2 x bfloat> %b,
279287

280288
; CHECK-LABEL: test_fptrunc_2xfloat(
281289
; CHECK: ld.param.v2.f32 {[[A0:%f[0-9]+]], [[A1:%f[0-9]+]]}, [test_fptrunc_2xfloat_param_0];
282-
; CHECK: cvt.rn.bf16x2.f32 [[R:%r[0-9]+]], [[A0]], [[A1]];
290+
; CHECK-DAG: cvt.rn.bf16.f32 [[R0:%rs[0-9]+]], [[A0]];
291+
; CHECK-DAG: cvt.rn.bf16.f32 [[R1:%rs[0-9]+]], [[A1]];
292+
; CHECK: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
283293
; CHECK: st.param.b32 [func_retval0], [[R]];
284294
; CHECK: ret;
285295
define <2 x bfloat> @test_fptrunc_2xfloat(<2 x float> %a) #0 {
@@ -349,7 +359,9 @@ declare <2 x bfloat> @llvm.fmuladd.f16(<2 x bfloat> %a, <2 x bfloat> %b, <2 x bf
349359
; CHECK-DAG: cvt.f32.bf16 [[AF1:%f[0-9]+]], [[A1]];
350360
; CHECK-DAG: sqrt.rn.f32 [[RF0:%f[0-9]+]], [[AF0]];
351361
; CHECK-DAG: sqrt.rn.f32 [[RF1:%f[0-9]+]], [[AF1]];
352-
; CHECK-DAG: cvt.rn.bf16x2.f32 [[R:%r[0-9]+]], [[RF0]], [[RF1]];
362+
; CHECK-DAG: cvt.rn.bf16.f32 [[R0:%rs[0-9]+]], [[RF0]];
363+
; CHECK-DAG: cvt.rn.bf16.f32 [[R1:%rs[0-9]+]], [[RF1]];
364+
; CHECK: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
353365
; CHECK: st.param.b32 [func_retval0], [[R]];
354366
; CHECK: ret;
355367
define <2 x bfloat> @test_sqrt(<2 x bfloat> %a) #0 {
@@ -424,7 +436,9 @@ define <2 x bfloat> @test_maxnum(<2 x bfloat> %a, <2 x bfloat> %b) #0 {
424436
; SM80-DAG: cvt.f32.bf16 [[FA1:%f[0-9]+]], [[A1]];
425437
; SM80-DAG: cvt.rmi.f32.f32 [[RF0:%f[0-9]+]], [[FA0]];
426438
; SM80-DAG: cvt.rmi.f32.f32 [[RF1:%f[0-9]+]], [[FA1]];
427-
; SM80: cvt.rn.bf16x2.f32 [[R:%r[0-9]+]], [[RF0]], [[RF1]];
439+
; SM80-DAG: cvt.rn.bf16.f32 [[R0:%rs[0-9]+]], [[RF0]];
440+
; SM80-DAG: cvt.rn.bf16.f32 [[R1:%rs[0-9]+]], [[RF1]];
441+
; CHECK: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
428442
; CHECK: st.param.b32 [func_retval0], [[R]];
429443
; CHECK: ret;
430444
define <2 x bfloat> @test_floor(<2 x bfloat> %a) #0 {
@@ -441,7 +455,9 @@ define <2 x bfloat> @test_floor(<2 x bfloat> %a) #0 {
441455
; SM80-DAG: cvt.f32.bf16 [[FA1:%f[0-9]+]], [[A1]];
442456
; SM80-DAG: cvt.rpi.f32.f32 [[RF0:%f[0-9]+]], [[FA0]];
443457
; SM80-DAG: cvt.rpi.f32.f32 [[RF1:%f[0-9]+]], [[FA1]];
444-
; SM80: cvt.rn.bf16x2.f32 [[R:%r[0-9]+]], [[RF0]], [[RF1]];
458+
; SM80-DAG: cvt.rn.bf16.f32 [[R0:%rs[0-9]+]], [[RF0]];
459+
; SM80-DAG: cvt.rn.bf16.f32 [[R1:%rs[0-9]+]], [[RF1]];
460+
; CHECK: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
445461
; CHECK: st.param.b32 [func_retval0], [[R]];
446462
; CHECK: ret;
447463
define <2 x bfloat> @test_ceil(<2 x bfloat> %a) #0 {
@@ -454,7 +470,7 @@ define <2 x bfloat> @test_ceil(<2 x bfloat> %a) #0 {
454470
; CHECK-DAG: mov.b32 {[[A0:%rs[0-9]+]], [[A1:%rs[0-9]+]]}, [[A]];
455471
; SM90: cvt.rzi.bf16.bf16 [[R1:%rs[0-9]+]], [[A1]];
456472
; SM90: cvt.rzi.bf16.bf16 [[R0:%rs[0-9]+]], [[A0]];
457-
; SM90: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
473+
; CHECK: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
458474
; CHECK: st.param.b32 [func_retval0], [[R]];
459475
; CHECK: ret;
460476
define <2 x bfloat> @test_trunc(<2 x bfloat> %a) #0 {
@@ -467,7 +483,7 @@ define <2 x bfloat> @test_trunc(<2 x bfloat> %a) #0 {
467483
; CHECK-DAG: mov.b32 {[[A0:%rs[0-9]+]], [[A1:%rs[0-9]+]]}, [[A]];
468484
; SM90: cvt.rni.bf16.bf16 [[R1:%rs[0-9]+]], [[A1]];
469485
; SM90: cvt.rni.bf16.bf16 [[R0:%rs[0-9]+]], [[A0]];
470-
; SM90: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
486+
; CHECK: mov.b32 [[R:%r[0-9]+]], {[[R0]], [[R1]]}
471487
; CHECK: st.param.b32 [func_retval0], [[R]];
472488
; CHECK: ret;
473489
define <2 x bfloat> @test_rint(<2 x bfloat> %a) #0 {

0 commit comments

Comments
 (0)