Skip to content

Commit 9fe8fdb

Browse files
justinfargnolijustinfargnoli
committed
Fix selection value for PRMT__32
1 parent e908645 commit 9fe8fdb

File tree

3 files changed

+44
-46
lines changed

3 files changed

+44
-46
lines changed

‎llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

Lines changed: 19 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -2319,38 +2319,36 @@ NVPTXTargetLowering::LowerCONCAT_VECTORS(SDValue Op, SelectionDAG &DAG) const {
23192319
// mov.b32 %r2, 0x40003C00
23202320
SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op,
23212321
SelectionDAG &DAG) const {
2322+
SDLoc DL(Op);
23222323
EVT VT = Op->getValueType(0);
23232324
if (!(Isv2x16VT(VT) || VT == MVT::v4i8))
23242325
return Op;
23252326

2326-
SDLoc DL(Op);
2327-
23282327
if (!llvm::all_of(Op->ops(), [](SDValue Operand) {
23292328
return Operand->isUndef() || isa<ConstantSDNode>(Operand) ||
23302329
isa<ConstantFPSDNode>(Operand);
23312330
})) {
2331+
if (VT != MVT::v4i8)
2332+
return Op;
23322333
// Lower non-const v4i8 vector as byte-wise constructed i32, which allows us
23332334
// to optimize calculation of constant parts.
2334-
if (VT == MVT::v4i8) {
2335-
SDValue PRMT__10 = DAG.getNode(
2336-
NVPTXISD::PRMT, DL, MVT::v4i8,
2337-
{DAG.getAnyExtOrTrunc(Op->getOperand(0), DL, MVT::i32),
2338-
DAG.getAnyExtOrTrunc(Op->getOperand(1), DL, MVT::i32),
2339-
DAG.getConstant(0x3340, DL, MVT::i32),
2340-
DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)});
2341-
SDValue PRMT32__ = DAG.getNode(
2342-
NVPTXISD::PRMT, DL, MVT::v4i8,
2343-
{DAG.getAnyExtOrTrunc(Op->getOperand(2), DL, MVT::i32),
2344-
DAG.getAnyExtOrTrunc(Op->getOperand(3), DL, MVT::i32),
2345-
DAG.getConstant(0x4033, DL, MVT::i32),
2346-
DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)});
2347-
SDValue PRMT3210 = DAG.getNode(
2348-
NVPTXISD::PRMT, DL, MVT::v4i8,
2349-
{PRMT__10, PRMT32__, DAG.getConstant(0x5410, DL, MVT::i32),
2335+
auto GetPRMT = [&](const SDValue Left, const SDValue Right, bool Cast,
2336+
uint64_t SelectionValue) -> SDValue {
2337+
SDValue L = Left;
2338+
SDValue R = Right;
2339+
if (Cast) {
2340+
L = DAG.getAnyExtOrTrunc(L, DL, MVT::i32);
2341+
R = DAG.getAnyExtOrTrunc(R, DL, MVT::i32);
2342+
}
2343+
return DAG.getNode(
2344+
NVPTXISD::PRMT, DL, MVT::v4i8, {L, R,
2345+
DAG.getConstant(SelectionValue, DL, MVT::i32),
23502346
DAG.getConstant(NVPTX::PTXPrmtMode::NONE, DL, MVT::i32)});
2351-
return DAG.getNode(ISD::BITCAST, DL, VT, PRMT3210);
2352-
}
2353-
return Op;
2347+
};
2348+
auto PRMT__10 = GetPRMT(Op->getOperand(0), Op->getOperand(1), true, 0x3340);
2349+
auto PRMT__32 = GetPRMT(Op->getOperand(2), Op->getOperand(3), true, 0x3340);
2350+
auto PRMT3210 = GetPRMT(PRMT__10, PRMT__32, false, 0x5410);
2351+
return DAG.getNode(ISD::BITCAST, DL, VT, PRMT3210);
23542352
}
23552353

23562354
// Get value or the Nth operand as an APInt(32). Undef values treated as 0.

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

Lines changed: 24 additions & 24 deletions
Original file line numberDiff line numberDiff line change
@@ -118,7 +118,7 @@ define <4 x i8> @test_add(<4 x i8> %a, <4 x i8> %b) #0 {
118118
; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;
119119
; CHECK-NEXT: add.s16 %rs6, %rs5, %rs4;
120120
; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;
121-
; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435;
121+
; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120;
122122
; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8;
123123
; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;
124124
; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8;
@@ -155,7 +155,7 @@ define <4 x i8> @test_add_imm_0(<4 x i8> %a) #0 {
155155
; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;
156156
; CHECK-NEXT: add.s16 %rs4, %rs3, 3;
157157
; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;
158-
; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 16435;
158+
; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 13120;
159159
; CHECK-NEXT: bfe.u32 %r7, %r1, 8, 8;
160160
; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;
161161
; CHECK-NEXT: add.s16 %rs6, %rs5, 2;
@@ -188,7 +188,7 @@ define <4 x i8> @test_add_imm_1(<4 x i8> %a) #0 {
188188
; CHECK-NEXT: cvt.u16.u32 %rs3, %r4;
189189
; CHECK-NEXT: add.s16 %rs4, %rs3, 3;
190190
; CHECK-NEXT: cvt.u32.u16 %r5, %rs4;
191-
; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 16435;
191+
; CHECK-NEXT: prmt.b32 %r6, %r5, %r3, 13120;
192192
; CHECK-NEXT: bfe.u32 %r7, %r1, 8, 8;
193193
; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;
194194
; CHECK-NEXT: add.s16 %rs6, %rs5, 2;
@@ -226,7 +226,7 @@ define <4 x i8> @test_sub(<4 x i8> %a, <4 x i8> %b) #0 {
226226
; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;
227227
; CHECK-NEXT: sub.s16 %rs6, %rs5, %rs4;
228228
; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;
229-
; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435;
229+
; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120;
230230
; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8;
231231
; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;
232232
; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8;
@@ -276,7 +276,7 @@ define <4 x i8> @test_smax(<4 x i8> %a, <4 x i8> %b) #0 {
276276
; CHECK-NEXT: selp.b32 %r16, %r14, %r15, %p4;
277277
; CHECK-NEXT: bfe.u32 %r17, %r2, 16, 8;
278278
; CHECK-NEXT: selp.b32 %r18, %r13, %r17, %p3;
279-
; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 16435;
279+
; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 13120;
280280
; CHECK-NEXT: bfe.u32 %r20, %r2, 8, 8;
281281
; CHECK-NEXT: selp.b32 %r21, %r12, %r20, %p2;
282282
; CHECK-NEXT: bfe.u32 %r22, %r2, 0, 8;
@@ -313,7 +313,7 @@ define <4 x i8> @test_umax(<4 x i8> %a, <4 x i8> %b) #0 {
313313
; CHECK-NEXT: setp.hi.u32 %p4, %r10, %r9;
314314
; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4;
315315
; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3;
316-
; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 16435;
316+
; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120;
317317
; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2;
318318
; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1;
319319
; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 13120;
@@ -354,7 +354,7 @@ define <4 x i8> @test_smin(<4 x i8> %a, <4 x i8> %b) #0 {
354354
; CHECK-NEXT: selp.b32 %r16, %r14, %r15, %p4;
355355
; CHECK-NEXT: bfe.u32 %r17, %r2, 16, 8;
356356
; CHECK-NEXT: selp.b32 %r18, %r13, %r17, %p3;
357-
; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 16435;
357+
; CHECK-NEXT: prmt.b32 %r19, %r18, %r16, 13120;
358358
; CHECK-NEXT: bfe.u32 %r20, %r2, 8, 8;
359359
; CHECK-NEXT: selp.b32 %r21, %r12, %r20, %p2;
360360
; CHECK-NEXT: bfe.u32 %r22, %r2, 0, 8;
@@ -391,7 +391,7 @@ define <4 x i8> @test_umin(<4 x i8> %a, <4 x i8> %b) #0 {
391391
; CHECK-NEXT: setp.ls.u32 %p4, %r10, %r9;
392392
; CHECK-NEXT: selp.b32 %r11, %r10, %r9, %p4;
393393
; CHECK-NEXT: selp.b32 %r12, %r8, %r7, %p3;
394-
; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 16435;
394+
; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120;
395395
; CHECK-NEXT: selp.b32 %r14, %r6, %r5, %p2;
396396
; CHECK-NEXT: selp.b32 %r15, %r4, %r3, %p1;
397397
; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 13120;
@@ -429,7 +429,7 @@ define <4 x i8> @test_eq(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 {
429429
; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4;
430430
; CHECK-NEXT: bfe.u32 %r14, %r3, 16, 8;
431431
; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3;
432-
; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 16435;
432+
; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 13120;
433433
; CHECK-NEXT: bfe.u32 %r17, %r3, 8, 8;
434434
; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2;
435435
; CHECK-NEXT: bfe.u32 %r19, %r3, 0, 8;
@@ -469,7 +469,7 @@ define <4 x i8> @test_ne(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c) #0 {
469469
; CHECK-NEXT: selp.b32 %r13, %r11, %r12, %p4;
470470
; CHECK-NEXT: bfe.u32 %r14, %r3, 16, 8;
471471
; CHECK-NEXT: selp.b32 %r15, %r9, %r14, %p3;
472-
; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 16435;
472+
; CHECK-NEXT: prmt.b32 %r16, %r15, %r13, 13120;
473473
; CHECK-NEXT: bfe.u32 %r17, %r3, 8, 8;
474474
; CHECK-NEXT: selp.b32 %r18, %r7, %r17, %p2;
475475
; CHECK-NEXT: bfe.u32 %r19, %r3, 0, 8;
@@ -504,7 +504,7 @@ define <4 x i8> @test_mul(<4 x i8> %a, <4 x i8> %b) #0 {
504504
; CHECK-NEXT: cvt.u16.u32 %rs5, %r7;
505505
; CHECK-NEXT: mul.lo.s16 %rs6, %rs5, %rs4;
506506
; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;
507-
; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435;
507+
; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120;
508508
; CHECK-NEXT: bfe.u32 %r10, %r2, 8, 8;
509509
; CHECK-NEXT: cvt.u16.u32 %rs7, %r10;
510510
; CHECK-NEXT: bfe.u32 %r11, %r1, 8, 8;
@@ -549,7 +549,7 @@ define <4 x i8> @test_or_computed(i8 %a) {
549549
; CHECK-NEXT: // %bb.0:
550550
; CHECK-NEXT: ld.param.u8 %rs1, [test_or_computed_param_0];
551551
; CHECK-NEXT: mov.b32 %r1, 0;
552-
; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 16435;
552+
; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120;
553553
; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
554554
; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 13120;
555555
; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520;
@@ -615,7 +615,7 @@ define <4 x i8> @test_xor_computed(i8 %a) {
615615
; CHECK-NEXT: // %bb.0:
616616
; CHECK-NEXT: ld.param.u8 %rs1, [test_xor_computed_param_0];
617617
; CHECK-NEXT: mov.b32 %r1, 0;
618-
; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 16435;
618+
; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120;
619619
; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
620620
; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 13120;
621621
; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520;
@@ -681,7 +681,7 @@ define <4 x i8> @test_and_computed(i8 %a) {
681681
; CHECK-NEXT: // %bb.0:
682682
; CHECK-NEXT: ld.param.u8 %rs1, [test_and_computed_param_0];
683683
; CHECK-NEXT: mov.b32 %r1, 0;
684-
; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 16435;
684+
; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 13120;
685685
; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
686686
; CHECK-NEXT: prmt.b32 %r4, %r3, 0, 13120;
687687
; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520;
@@ -954,7 +954,7 @@ define <4 x i8> @test_select_cc(<4 x i8> %a, <4 x i8> %b, <4 x i8> %c, <4 x i8>
954954
; CHECK-NEXT: bfe.u32 %r16, %r2, 16, 8;
955955
; CHECK-NEXT: bfe.u32 %r17, %r1, 16, 8;
956956
; CHECK-NEXT: selp.b32 %r18, %r17, %r16, %p3;
957-
; CHECK-NEXT: prmt.b32 %r19, %r18, %r15, 16435;
957+
; CHECK-NEXT: prmt.b32 %r19, %r18, %r15, 13120;
958958
; CHECK-NEXT: bfe.u32 %r20, %r2, 8, 8;
959959
; CHECK-NEXT: bfe.u32 %r21, %r1, 8, 8;
960960
; CHECK-NEXT: selp.b32 %r22, %r21, %r20, %p2;
@@ -1026,7 +1026,7 @@ define <4 x i8> @test_select_cc_i8_i32(<4 x i8> %a, <4 x i8> %b,
10261026
; CHECK-NEXT: bfe.u32 %r14, %r2, 16, 8;
10271027
; CHECK-NEXT: bfe.u32 %r15, %r1, 16, 8;
10281028
; CHECK-NEXT: selp.b32 %r16, %r15, %r14, %p3;
1029-
; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 16435;
1029+
; CHECK-NEXT: prmt.b32 %r17, %r16, %r13, 13120;
10301030
; CHECK-NEXT: bfe.u32 %r18, %r2, 8, 8;
10311031
; CHECK-NEXT: bfe.u32 %r19, %r1, 8, 8;
10321032
; CHECK-NEXT: selp.b32 %r20, %r19, %r18, %p2;
@@ -1051,7 +1051,7 @@ define <4 x i8> @test_trunc_2xi32(<4 x i32> %a) #0 {
10511051
; CHECK-EMPTY:
10521052
; CHECK-NEXT: // %bb.0:
10531053
; CHECK-NEXT: ld.param.v4.u32 {%r1, %r2, %r3, %r4}, [test_trunc_2xi32_param_0];
1054-
; CHECK-NEXT: prmt.b32 %r5, %r3, %r4, 16435;
1054+
; CHECK-NEXT: prmt.b32 %r5, %r3, %r4, 13120;
10551055
; CHECK-NEXT: prmt.b32 %r6, %r1, %r2, 13120;
10561056
; CHECK-NEXT: prmt.b32 %r7, %r6, %r5, 21520;
10571057
; CHECK-NEXT: st.param.b32 [func_retval0+0], %r7;
@@ -1071,7 +1071,7 @@ define <4 x i8> @test_trunc_2xi64(<4 x i64> %a) #0 {
10711071
; CHECK-NEXT: ld.param.v2.u64 {%rd1, %rd2}, [test_trunc_2xi64_param_0];
10721072
; CHECK-NEXT: cvt.u32.u64 %r1, %rd4;
10731073
; CHECK-NEXT: cvt.u32.u64 %r2, %rd3;
1074-
; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 16435;
1074+
; CHECK-NEXT: prmt.b32 %r3, %r2, %r1, 13120;
10751075
; CHECK-NEXT: cvt.u32.u64 %r4, %rd2;
10761076
; CHECK-NEXT: cvt.u32.u64 %r5, %rd1;
10771077
; CHECK-NEXT: prmt.b32 %r6, %r5, %r4, 13120;
@@ -1192,7 +1192,7 @@ define <2 x half> @test_bitcast_4xi8_to_2xhalf(i8 %a) #0 {
11921192
; CHECK-NEXT: // %bb.0:
11931193
; CHECK-NEXT: ld.param.u8 %rs1, [test_bitcast_4xi8_to_2xhalf_param_0];
11941194
; CHECK-NEXT: mov.b32 %r1, 6;
1195-
; CHECK-NEXT: prmt.b32 %r2, %r1, 7, 16435;
1195+
; CHECK-NEXT: prmt.b32 %r2, %r1, 7, 13120;
11961196
; CHECK-NEXT: cvt.u32.u16 %r3, %rs1;
11971197
; CHECK-NEXT: prmt.b32 %r4, %r3, 5, 13120;
11981198
; CHECK-NEXT: prmt.b32 %r5, %r4, %r2, 21520;
@@ -1270,7 +1270,7 @@ define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 {
12701270
; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r5;
12711271
; CHECK-NEXT: cvt.u32.u16 %r6, %rs6;
12721272
; CHECK-NEXT: cvt.u32.u16 %r7, %rs5;
1273-
; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 16435;
1273+
; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120;
12741274
; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r3;
12751275
; CHECK-NEXT: cvt.rzi.s16.f16 %rs9, %rs8;
12761276
; CHECK-NEXT: cvt.rzi.s16.f16 %rs10, %rs7;
@@ -1301,7 +1301,7 @@ define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 {
13011301
; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r5;
13021302
; CHECK-NEXT: cvt.u32.u16 %r6, %rs6;
13031303
; CHECK-NEXT: cvt.u32.u16 %r7, %rs5;
1304-
; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 16435;
1304+
; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120;
13051305
; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r3;
13061306
; CHECK-NEXT: cvt.rzi.u16.f16 %rs9, %rs8;
13071307
; CHECK-NEXT: cvt.rzi.u16.f16 %rs10, %rs7;
@@ -1342,7 +1342,7 @@ define void @test_srem_v4i8(ptr %a, ptr %b, ptr %c) {
13421342
; CHECK-NEXT: cvt.s8.s32 %rs5, %r7;
13431343
; CHECK-NEXT: rem.s16 %rs6, %rs5, %rs4;
13441344
; CHECK-NEXT: cvt.u32.u16 %r8, %rs6;
1345-
; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 16435;
1345+
; CHECK-NEXT: prmt.b32 %r9, %r8, %r5, 13120;
13461346
; CHECK-NEXT: bfe.s32 %r10, %r2, 8, 8;
13471347
; CHECK-NEXT: cvt.s8.s32 %rs7, %r10;
13481348
; CHECK-NEXT: bfe.s32 %r11, %r1, 8, 8;
@@ -1411,7 +1411,7 @@ define void @test_srem_v3i8(ptr %a, ptr %b, ptr %c) {
14111411
; CHECK-NEXT: prmt.b32 %r11, %r10, %r7, 13120;
14121412
; CHECK-NEXT: // implicit-def: %r13
14131413
; CHECK-NEXT: // implicit-def: %r14
1414-
; CHECK-NEXT: prmt.b32 %r12, %r13, %r14, 16435;
1414+
; CHECK-NEXT: prmt.b32 %r12, %r13, %r14, 13120;
14151415
; CHECK-NEXT: prmt.b32 %r15, %r11, %r12, 21520;
14161416
; CHECK-NEXT: rem.s16 %rs17, %rs5, %rs10;
14171417
; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs18, tmp}, %r15; }
@@ -1455,7 +1455,7 @@ define void @test_sext_v4i1_to_v4i8(ptr %a, ptr %b, ptr %c) {
14551455
; CHECK-NEXT: setp.hi.u32 %p4, %r10, %r9;
14561456
; CHECK-NEXT: selp.s32 %r11, -1, 0, %p4;
14571457
; CHECK-NEXT: selp.s32 %r12, -1, 0, %p3;
1458-
; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 16435;
1458+
; CHECK-NEXT: prmt.b32 %r13, %r12, %r11, 13120;
14591459
; CHECK-NEXT: selp.s32 %r14, -1, 0, %p2;
14601460
; CHECK-NEXT: selp.s32 %r15, -1, 0, %p1;
14611461
; CHECK-NEXT: prmt.b32 %r16, %r15, %r14, 13120;

‎llvm/test/CodeGen/NVPTX/sext-setcc.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -57,7 +57,7 @@ define <4 x i8> @sext_setcc_v4i1_to_v4i8(ptr %p) {
5757
; CHECK-NEXT: setp.eq.s16 %p4, %rs8, 0;
5858
; CHECK-NEXT: selp.s32 %r6, -1, 0, %p4;
5959
; CHECK-NEXT: selp.s32 %r7, -1, 0, %p3;
60-
; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 16435;
60+
; CHECK-NEXT: prmt.b32 %r8, %r7, %r6, 13120;
6161
; CHECK-NEXT: selp.s32 %r9, -1, 0, %p2;
6262
; CHECK-NEXT: selp.s32 %r10, -1, 0, %p1;
6363
; CHECK-NEXT: prmt.b32 %r11, %r10, %r9, 13120;

0 commit comments

Comments
 (0)