use fadd v2f32 to keep bitcast pattern in isel

Author: Prince781

llvm/test/CodeGen/NVPTX/f32x2-convert-i32x2.ll

@@ -20,26 +20,49 @@ declare <2 x i32> @return_i32x2(i32 %0)
 
 ; Test with v2i32.
 define ptx_kernel void @store_i32x2(i32 %0, ptr %p) {
-; CHECK-LABEL: store_i32x2(
-; CHECK:       {
-; CHECK-NEXT:    .reg .b32 %r<6>;
-; CHECK-NEXT:    .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT:  // %bb.0:
-; CHECK-NEXT:    ld.param.b64 %rd1, [store_i32x2_param_1];
-; CHECK-NEXT:    ld.param.b32 %r1, [store_i32x2_param_0];
-; CHECK-NEXT:    { // callseq 0, 0
-; CHECK-NEXT:    .param .b32 param0;
-; CHECK-NEXT:    .param .align 8 .b8 retval0[8];
-; CHECK-NEXT:    st.param.b32 [param0], %r1;
-; CHECK-NEXT:    call.uni (retval0), return_i32x2, (param0);
-; CHECK-NEXT:    ld.param.v2.b32 {%r2, %r3}, [retval0];
-; CHECK-NEXT:    } // callseq 0
-; CHECK-NEXT:    st.v2.b32 [%rd1], {%r2, %r3};
-; CHECK-NEXT:    ret;
+; CHECK-SM90A-LABEL: store_i32x2(
+; CHECK-SM90A:       {
+; CHECK-SM90A-NEXT:    .reg .b32 %r<6>;
+; CHECK-SM90A-NEXT:    .reg .b64 %rd<2>;
+; CHECK-SM90A-EMPTY:
+; CHECK-SM90A-NEXT:  // %bb.0:
+; CHECK-SM90A-NEXT:    ld.param.b64 %rd1, [store_i32x2_param_1];
+; CHECK-SM90A-NEXT:    ld.param.b32 %r1, [store_i32x2_param_0];
+; CHECK-SM90A-NEXT:    { // callseq 0, 0
+; CHECK-SM90A-NEXT:    .param .b32 param0;
+; CHECK-SM90A-NEXT:    .param .align 8 .b8 retval0[8];
+; CHECK-SM90A-NEXT:    st.param.b32 [param0], %r1;
+; CHECK-SM90A-NEXT:    call.uni (retval0), return_i32x2, (param0);
+; CHECK-SM90A-NEXT:    ld.param.v2.b32 {%r2, %r3}, [retval0];
+; CHECK-SM90A-NEXT:    } // callseq 0
+; CHECK-SM90A-NEXT:    add.rn.f32 %r4, %r3, %r3;
+; CHECK-SM90A-NEXT:    add.rn.f32 %r5, %r2, %r2;
+; CHECK-SM90A-NEXT:    st.v2.b32 [%rd1], {%r5, %r4};
+; CHECK-SM90A-NEXT:    ret;
+;
+; CHECK-SM100-LABEL: store_i32x2(
+; CHECK-SM100:       {
+; CHECK-SM100-NEXT:    .reg .b32 %r<4>;
+; CHECK-SM100-NEXT:    .reg .b64 %rd<4>;
+; CHECK-SM100-EMPTY:
+; CHECK-SM100-NEXT:  // %bb.0:
+; CHECK-SM100-NEXT:    ld.param.b64 %rd1, [store_i32x2_param_1];
+; CHECK-SM100-NEXT:    ld.param.b32 %r1, [store_i32x2_param_0];
+; CHECK-SM100-NEXT:    { // callseq 0, 0
+; CHECK-SM100-NEXT:    .param .b32 param0;
+; CHECK-SM100-NEXT:    .param .align 8 .b8 retval0[8];
+; CHECK-SM100-NEXT:    st.param.b32 [param0], %r1;
+; CHECK-SM100-NEXT:    call.uni (retval0), return_i32x2, (param0);
+; CHECK-SM100-NEXT:    ld.param.v2.b32 {%r2, %r3}, [retval0];
+; CHECK-SM100-NEXT:    } // callseq 0
+; CHECK-SM100-NEXT:    mov.b64 %rd2, {%r2, %r3};
+; CHECK-SM100-NEXT:    add.rn.f32x2 %rd3, %rd2, %rd2;
+; CHECK-SM100-NEXT:    st.b64 [%rd1], %rd3;
+; CHECK-SM100-NEXT:    ret;
   %v = call <2 x i32> @return_i32x2(i32 %0)
   %v.f32x2 = bitcast <2 x i32> %v to <2 x float>
-  store <2 x float> %v.f32x2, ptr %p, align 8
+  %res = fadd <2 x float> %v.f32x2, %v.f32x2
+  store <2 x float> %res, ptr %p, align 8
   ret void
 }
 
@@ -93,3 +116,5 @@ define ptx_kernel void @inlineasm(ptr %p) {
   store <2 x float> %mul, ptr %p, align 8
   ret void
 }
+;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+; CHECK: {{.*}}