[NFC] Fix conflict resolution in b6bd010

The incoming commit use int reg instead of float reg,
we should only update accoridly, don't blindly accept all the upstream
code.

Author: jsji

llvm/test/CodeGen/NVPTX/surf-read-cuda.ll

@@ -14,14 +14,15 @@ define void @foo(i64 %img, ptr %red, i32 %idx) {
 ; CHECK-LABEL: foo(
 ; CHECK:       {
 ; CHECK-NEXT:    .reg .b32 %r<4>;
-; CHECK-NEXT:    .reg .b64 %rd<4>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b64 %rd1, [foo_param_0];
 ; CHECK-NEXT:    ld.param.b32 %r1, [foo_param_2];
 ; CHECK-NEXT:    suld.b.1d.b32.trap {%r2}, [%rd1, {%r1}];
+; CHECK-NEXT:    ld.param.b64 %rd2, [foo_param_1];
 ; CHECK-NEXT:    cvt.rn.f32.s32 %r3, %r2;
-; CHECK-NEXT:    st.global.b32 [%rd3], %r3;
+; CHECK-NEXT:    st.b32 [%rd2], %r3;
 ; CHECK-NEXT:    ret;
   %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %img, i32 %idx)
   %ret = sitofp i32 %val to float
@@ -35,14 +36,14 @@ define void @bar(ptr %red, i32 %idx) {
 ; CHECK-LABEL: bar(
 ; CHECK:       {
 ; CHECK-NEXT:    .reg .b32 %r<4>;
-; CHECK-NEXT:    .reg .b64 %rd<4>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b64 %rd1, [bar_param_0];
 ; CHECK-NEXT:    ld.param.b32 %r1, [bar_param_1];
 ; CHECK-NEXT:    suld.b.1d.b32.trap {%r2}, [surf0, {%r1}];
 ; CHECK-NEXT:    cvt.rn.f32.s32 %r3, %r2;
-; CHECK-NEXT:    st.global.b32 [%rd2], %r3;
+; CHECK-NEXT:    st.b32 [%rd1], %r3;
 ; CHECK-NEXT:    ret;
   %surfHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @surf0)
   %val = tail call i32 @llvm.nvvm.suld.1d.i32.trap(i64 %surfHandle, i32 %idx)

llvm/test/CodeGen/NVPTX/tex-read-cuda.ll

@@ -14,13 +14,14 @@ define void @foo(i64 %img, ptr %red, i32 %idx) {
 ; CHECK-LABEL: foo(
 ; CHECK:       {
 ; CHECK-NEXT:    .reg .b32 %r<6>;
-; CHECK-NEXT:    .reg .b64 %rd<4>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b64 %rd1, [foo_param_0];
 ; CHECK-NEXT:    ld.param.b32 %r1, [foo_param_2];
 ; CHECK-NEXT:    tex.1d.v4.f32.s32 {%r2, %r3, %r4, %r5}, [%rd1, {%r1}];
-; CHECK-NEXT:    st.global.b32 [%rd3], %r2;
+; CHECK-NEXT:    ld.param.b64 %rd2, [foo_param_1];
+; CHECK-NEXT:    st.b32 [%rd2], %r2;
 ; CHECK-NEXT:    ret;
   %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %img, i32 %idx)
   %ret = extractvalue { float, float, float, float } %val, 0
@@ -35,13 +36,13 @@ define void @bar(ptr %red, i32 %idx) {
 ; CHECK-LABEL: bar(
 ; CHECK:       {
 ; CHECK-NEXT:    .reg .b32 %r<6>;
-; CHECK-NEXT:    .reg .b64 %rd<4>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b64 %rd1, [bar_param_0];
 ; CHECK-NEXT:    ld.param.b32 %r1, [bar_param_1];
 ; CHECK-NEXT:    tex.1d.v4.f32.s32 {%r2, %r3, %r4, %r5}, [tex0, {%r1}];
-; CHECK-NEXT:    st.global.b32 [%rd2], %r2;
+; CHECK-NEXT:    st.b32 [%rd1], %r2;
 ; CHECK-NEXT:    ret;
   %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0)
   %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx)
@@ -56,12 +57,12 @@ define void @baz(ptr %red, i32 %idx) {
 ; CHECK-LABEL: baz(
 ; CHECK:       {
 ; CHECK-NEXT:    .reg .b32 %r<9>;
-; CHECK-NEXT:    .reg .b64 %rd<4>;
+; CHECK-NEXT:    .reg .b64 %rd<3>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b64 %rd1, [baz_param_0];
 ; CHECK-NEXT:    ld.param.b32 %r1, [baz_param_1];
-; CHECK-NEXT:    mov.u64 %rd3, tex0;
+; CHECK-NEXT:    mov.u64 %rd2, tex0;
 ; CHECK-NEXT:    tex.1d.v4.f32.s32 {%r2, %r3, %r4, %r5}, [tex0, {%r1}];
 ; CHECK-NEXT:    { // callseq 0, 0
 ; CHECK-NEXT:    .param .b64 param0;
@@ -75,7 +76,7 @@ define void @baz(ptr %red, i32 %idx) {
 ; CHECK-NEXT:    ld.param.b32 %r6, [retval0];
 ; CHECK-NEXT:    } // callseq 0
 ; CHECK-NEXT:    add.rn.f32 %r8, %r2, %r6;
-; CHECK-NEXT:    st.global.b32 [%rd2], %r8;
+; CHECK-NEXT:    st.b32 [%rd1], %r8;
 ; CHECK-NEXT:    ret;
   %texHandle = tail call i64 @llvm.nvvm.texsurf.handle.internal.p1(ptr addrspace(1) @tex0)
   %val = tail call { float, float, float, float } @llvm.nvvm.tex.unified.1d.v4f32.s32(i64 %texHandle, i32 %idx)