multiple use

Author: frasercrmck

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

@@ -6213,6 +6213,11 @@ PerformBUILD_VECTORCombine(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
           Op->getOperand(0).getValueType() == MVT::i32))
       return SDValue();
 
+    // If the truncate has multiple uses, this optimization can increase
+    // register pressure
+    if (!Op->hasOneUse())
+      return SDValue();
+
     *Op = Op->getOperand(0);
 
     // Optionally, fold in a shift-right of the original operand and let permute

llvm/test/CodeGen/NVPTX/i16x2-instructions.ll

@@ -818,6 +818,71 @@ define <2 x i16> @test_trunc_2xi32(<2 x i32> %a) #0 {
   ret <2 x i16> %r
 }
 
+define <2 x i16> @test_trunc_2xi32_muliple_use0(<2 x i32> %a, ptr %p) #0 {
+; I16x2-LABEL: test_trunc_2xi32_muliple_use0(
+; I16x2:       {
+; I16x2-NEXT:    .reg .b32 %r<7>;
+; I16x2-NEXT:    .reg .b64 %rd<2>;
+; I16x2-EMPTY:
+; I16x2-NEXT:  // %bb.0:
+; I16x2-NEXT:    ld.param.v2.u32 {%r1, %r2}, [test_trunc_2xi32_muliple_use0_param_0];
+; I16x2-NEXT:    ld.param.u64 %rd1, [test_trunc_2xi32_muliple_use0_param_1];
+; I16x2-NEXT:    prmt.b32 %r3, %r1, %r2, 0x5410U;
+; I16x2-NEXT:    mov.b32 %r5, 65537;
+; I16x2-NEXT:    add.s16x2 %r6, %r3, %r5;
+; I16x2-NEXT:    st.u32 [%rd1], %r6;
+; I16x2-NEXT:    st.param.b32 [func_retval0], %r3;
+; I16x2-NEXT:    ret;
+;
+; NO-I16x2-LABEL: test_trunc_2xi32_muliple_use0(
+; NO-I16x2:       {
+; NO-I16x2-NEXT:    .reg .b16 %rs<5>;
+; NO-I16x2-NEXT:    .reg .b32 %r<5>;
+; NO-I16x2-NEXT:    .reg .b64 %rd<2>;
+; NO-I16x2-EMPTY:
+; NO-I16x2-NEXT:  // %bb.0:
+; NO-I16x2-NEXT:    ld.param.v2.u32 {%r1, %r2}, [test_trunc_2xi32_muliple_use0_param_0];
+; NO-I16x2-NEXT:    ld.param.u64 %rd1, [test_trunc_2xi32_muliple_use0_param_1];
+; NO-I16x2-NEXT:    cvt.u16.u32 %rs1, %r2;
+; NO-I16x2-NEXT:    cvt.u16.u32 %rs2, %r1;
+; NO-I16x2-NEXT:    mov.b32 %r3, {%rs2, %rs1};
+; NO-I16x2-NEXT:    add.s16 %rs3, %rs1, 1;
+; NO-I16x2-NEXT:    add.s16 %rs4, %rs2, 1;
+; NO-I16x2-NEXT:    mov.b32 %r4, {%rs4, %rs3};
+; NO-I16x2-NEXT:    st.u32 [%rd1], %r4;
+; NO-I16x2-NEXT:    st.param.b32 [func_retval0], %r3;
+; NO-I16x2-NEXT:    ret;
+  %r = trunc <2 x i32> %a to <2 x i16>
+  ; Reuse the truncate - optimizing to PRMT when we don't have i16x2 vectors
+  ; would increase register pressure
+  %s = add <2 x i16> %r, splat (i16 1)
+  store <2 x i16> %s, ptr %p
+  ret <2 x i16> %r
+}
+
+define <2 x i16> @test_trunc_2xi32_muliple_use1(<2 x i32> %a, ptr %p) #0 {
+; COMMON-LABEL: test_trunc_2xi32_muliple_use1(
+; COMMON:       {
+; COMMON-NEXT:    .reg .b32 %r<7>;
+; COMMON-NEXT:    .reg .b64 %rd<2>;
+; COMMON-EMPTY:
+; COMMON-NEXT:  // %bb.0:
+; COMMON-NEXT:    ld.param.v2.u32 {%r1, %r2}, [test_trunc_2xi32_muliple_use1_param_0];
+; COMMON-NEXT:    ld.param.u64 %rd1, [test_trunc_2xi32_muliple_use1_param_1];
+; COMMON-NEXT:    prmt.b32 %r3, %r1, %r2, 0x5410U;
+; COMMON-NEXT:    add.s32 %r5, %r2, 1;
+; COMMON-NEXT:    add.s32 %r6, %r1, 1;
+; COMMON-NEXT:    st.v2.u32 [%rd1], {%r6, %r5};
+; COMMON-NEXT:    st.param.b32 [func_retval0], %r3;
+; COMMON-NEXT:    ret;
+  %r = trunc <2 x i32> %a to <2 x i16>
+  ; Reuse the original value - optimizing to PRMT does not increase register
+  ; pressure
+  %s = add <2 x i32> %a, splat (i32 1)
+  store <2 x i32> %s, ptr %p
+  ret <2 x i16> %r
+}
+
 define <2 x i16> @test_trunc_2xi64(<2 x i64> %a) #0 {
 ; COMMON-LABEL: test_trunc_2xi64(
 ; COMMON:       {