[NVPTX] Update 32-bit NVPTX tests to use 64-bit

Summary:
These used the now-removed `nvptx` target. Because they used
`ptxas-verify` they would then fail since it's been removed. Update
these to 64-bit.

Author: jhuber6

llvm/test/CodeGen/NVPTX/inline-asm-b128-test1.ll

@@ -1,6 +1,6 @@
 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --extra_scrub --version 5
-; RUN: llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
 
 target triple = "nvptx64-nvidia-cuda"
 
@@ -9,47 +9,41 @@ target triple = "nvptx64-nvidia-cuda"
 define void @test_b128_input_from_const() {
 ; CHECK-LABEL: test_b128_input_from_const(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b32 %r<3>;
-; CHECK-NEXT:    .reg .b64 %rd<4>;
+; CHECK-NEXT:    .reg .b64 %rd<5>;
 ; CHECK-NEXT:    .reg .b128 %rq<2>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    mov.u64 %rd2, 0;
 ; CHECK-NEXT:    mov.u64 %rd3, 42;
 ; CHECK-NEXT:    mov.b128 %rq1, {%rd3, %rd2};
-; CHECK-NEXT:    mov.u32 %r1, value;
-; CHECK-NEXT:    cvta.global.u32 %r2, %r1;
-; CHECK-NEXT:    cvt.u64.u32 %rd1, %r2;
+; CHECK-NEXT:    mov.u64 %rd4, value;
+; CHECK-NEXT:    cvta.global.u64 %rd1, %rd4;
 ; CHECK-NEXT:    // begin inline asm
 ; CHECK-NEXT:    { st.b128 [%rd1], %rq1; }
 ; CHECK-NEXT:    // end inline asm
 ; CHECK-NEXT:    ret;
-
   tail call void asm sideeffect "{ st.b128 [$0], $1; }", "l,q"(ptr nonnull addrspacecast (ptr addrspace(1) @value to ptr), i128 42)
   ret void
 }
 
 define void @test_b128_input_from_load(ptr nocapture readonly %data) {
 ; CHECK-LABEL: test_b128_input_from_load(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b32 %r<5>;
-; CHECK-NEXT:    .reg .b64 %rd<4>;
+; CHECK-NEXT:    .reg .b64 %rd<7>;
 ; CHECK-NEXT:    .reg .b128 %rq<2>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
-; CHECK-NEXT:    ld.param.u32 %r1, [test_b128_input_from_load_param_0];
-; CHECK-NEXT:    cvta.to.global.u32 %r2, %r1;
-; CHECK-NEXT:    ld.global.u64 %rd2, [%r2+8];
-; CHECK-NEXT:    ld.global.u64 %rd3, [%r2];
-; CHECK-NEXT:    mov.b128 %rq1, {%rd3, %rd2};
-; CHECK-NEXT:    mov.u32 %r3, value;
-; CHECK-NEXT:    cvta.global.u32 %r4, %r3;
-; CHECK-NEXT:    cvt.u64.u32 %rd1, %r4;
+; CHECK-NEXT:    ld.param.u64 %rd2, [test_b128_input_from_load_param_0];
+; CHECK-NEXT:    cvta.to.global.u64 %rd3, %rd2;
+; CHECK-NEXT:    ld.global.u64 %rd4, [%rd3+8];
+; CHECK-NEXT:    ld.global.u64 %rd5, [%rd3];
+; CHECK-NEXT:    mov.b128 %rq1, {%rd5, %rd4};
+; CHECK-NEXT:    mov.u64 %rd6, value;
+; CHECK-NEXT:    cvta.global.u64 %rd1, %rd6;
 ; CHECK-NEXT:    // begin inline asm
 ; CHECK-NEXT:    { st.b128 [%rd1], %rq1; }
 ; CHECK-NEXT:    // end inline asm
 ; CHECK-NEXT:    ret;
-
   %1 = addrspacecast ptr %data to ptr addrspace(1)
   %2 = load <2 x i64>, ptr addrspace(1) %1, align 16
   %3 = bitcast <2 x i64> %2 to i128
@@ -62,26 +56,23 @@ define void @test_b128_input_from_select(ptr nocapture readonly %flag) {
 ; CHECK:       {
 ; CHECK-NEXT:    .reg .pred %p<2>;
 ; CHECK-NEXT:    .reg .b16 %rs<2>;
-; CHECK-NEXT:    .reg .b32 %r<5>;
-; CHECK-NEXT:    .reg .b64 %rd<4>;
+; CHECK-NEXT:    .reg .b64 %rd<7>;
 ; CHECK-NEXT:    .reg .b128 %rq<2>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
-; CHECK-NEXT:    ld.param.u32 %r1, [test_b128_input_from_select_param_0];
-; CHECK-NEXT:    cvta.to.global.u32 %r2, %r1;
-; CHECK-NEXT:    ld.global.u8 %rs1, [%r2];
+; CHECK-NEXT:    ld.param.u64 %rd2, [test_b128_input_from_select_param_0];
+; CHECK-NEXT:    cvta.to.global.u64 %rd3, %rd2;
+; CHECK-NEXT:    ld.global.u8 %rs1, [%rd3];
 ; CHECK-NEXT:    setp.eq.s16 %p1, %rs1, 0;
-; CHECK-NEXT:    selp.b64 %rd2, 24, 42, %p1;
-; CHECK-NEXT:    mov.u64 %rd3, 0;
-; CHECK-NEXT:    mov.b128 %rq1, {%rd2, %rd3};
-; CHECK-NEXT:    mov.u32 %r3, value;
-; CHECK-NEXT:    cvta.global.u32 %r4, %r3;
-; CHECK-NEXT:    cvt.u64.u32 %rd1, %r4;
+; CHECK-NEXT:    selp.b64 %rd4, 24, 42, %p1;
+; CHECK-NEXT:    mov.u64 %rd5, 0;
+; CHECK-NEXT:    mov.b128 %rq1, {%rd4, %rd5};
+; CHECK-NEXT:    mov.u64 %rd6, value;
+; CHECK-NEXT:    cvta.global.u64 %rd1, %rd6;
 ; CHECK-NEXT:    // begin inline asm
 ; CHECK-NEXT:    { st.b128 [%rd1], %rq1; }
 ; CHECK-NEXT:    // end inline asm
 ; CHECK-NEXT:    ret;
-
   %1 = addrspacecast ptr %flag to ptr addrspace(1)
   %2 = load i8, ptr addrspace(1) %1, align 1
   %3 = icmp eq i8 %2, 0
@@ -106,7 +97,6 @@ define void @test_store_b128_output() {
 ; CHECK-NEXT:    st.global.u64 [value+8], %rd4;
 ; CHECK-NEXT:    st.global.u64 [value], %rd3;
 ; CHECK-NEXT:    ret;
-
   %1 = tail call i128 asm "{ mov.b128 $0, 41; }", "=q"()
   %add = add nsw i128 %1, 1
   %2 = bitcast i128 %add to <2 x i64>
@@ -117,26 +107,24 @@ define void @test_store_b128_output() {
 define void @test_use_of_b128_output(ptr nocapture readonly %data) {
 ; CHECK-LABEL: test_use_of_b128_output(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b32 %r<3>;
-; CHECK-NEXT:    .reg .b64 %rd<7>;
+; CHECK-NEXT:    .reg .b64 %rd<9>;
 ; CHECK-NEXT:    .reg .b128 %rq<3>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
-; CHECK-NEXT:    ld.param.u32 %r1, [test_use_of_b128_output_param_0];
-; CHECK-NEXT:    cvta.to.global.u32 %r2, %r1;
-; CHECK-NEXT:    ld.global.u64 %rd1, [%r2+8];
-; CHECK-NEXT:    ld.global.u64 %rd2, [%r2];
-; CHECK-NEXT:    mov.b128 %rq2, {%rd2, %rd1};
+; CHECK-NEXT:    ld.param.u64 %rd1, [test_use_of_b128_output_param_0];
+; CHECK-NEXT:    cvta.to.global.u64 %rd2, %rd1;
+; CHECK-NEXT:    ld.global.u64 %rd3, [%rd2+8];
+; CHECK-NEXT:    ld.global.u64 %rd4, [%rd2];
+; CHECK-NEXT:    mov.b128 %rq2, {%rd4, %rd3};
 ; CHECK-NEXT:    // begin inline asm
 ; CHECK-NEXT:    { mov.b128 %rq1, %rq2; }
 ; CHECK-NEXT:    // end inline asm
-; CHECK-NEXT:    mov.b128 {%rd3, %rd4}, %rq1;
-; CHECK-NEXT:    add.cc.s64 %rd5, %rd3, 1;
-; CHECK-NEXT:    addc.cc.s64 %rd6, %rd4, 0;
-; CHECK-NEXT:    st.global.u64 [value], %rd5;
-; CHECK-NEXT:    st.global.u64 [value+8], %rd6;
+; CHECK-NEXT:    mov.b128 {%rd5, %rd6}, %rq1;
+; CHECK-NEXT:    add.cc.s64 %rd7, %rd5, 1;
+; CHECK-NEXT:    addc.cc.s64 %rd8, %rd6, 0;
+; CHECK-NEXT:    st.global.u64 [value], %rd7;
+; CHECK-NEXT:    st.global.u64 [value+8], %rd8;
 ; CHECK-NEXT:    ret;
-
   %1 = addrspacecast ptr %data to ptr addrspace(1)
   %2 = load <2 x i64>, ptr addrspace(1) %1, align 16
   %3 = bitcast <2 x i64> %2 to i128

llvm/test/CodeGen/NVPTX/inline-asm-b128-test2.ll

@@ -1,6 +1,6 @@
-; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --extra_scrub --version 5
-; RUN: llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
 
 target triple = "nvptx64-nvidia-cuda"
 
@@ -17,20 +17,16 @@ target triple = "nvptx64-nvidia-cuda"
 define void @test_corner_values() {
 ; CHECK-LABEL: test_corner_values(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b32 %r<20>;
-; CHECK-NEXT:    .reg .b64 %rd<17>;
+; CHECK-NEXT:    .reg .b64 %rd<24>;
 ; CHECK-NEXT:    .reg .b128 %rq<5>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
-; CHECK-NEXT:    ld.global.u32 %r1, [v64];
-; CHECK-NEXT:    add.s32 %r2, %r1, 8;
+; CHECK-NEXT:    ld.global.u64 %rd1, [v64];
+; CHECK-NEXT:    add.s64 %rd2, %rd1, 8;
 ; CHECK-NEXT:    mov.u64 %rd13, -1;
 ; CHECK-NEXT:    mov.b128 %rq1, {%rd13, %rd13};
-; CHECK-NEXT:    cvt.u64.u32 %rd1, %r1;
-; CHECK-NEXT:    cvt.u64.u32 %rd2, %r2;
-; CHECK-NEXT:    mov.u32 %r3, v_u128_max;
-; CHECK-NEXT:    cvta.global.u32 %r4, %r3;
-; CHECK-NEXT:    cvt.u64.u32 %rd3, %r4;
+; CHECK-NEXT:    mov.u64 %rd14, v_u128_max;
+; CHECK-NEXT:    cvta.global.u64 %rd3, %rd14;
 ; CHECK-NEXT:    // begin inline asm
 ; CHECK-NEXT:    {
 ; CHECK-NEXT:    .reg .b64 hi;
@@ -41,16 +37,13 @@ define void @test_corner_values() {
 ; CHECK-NEXT:    st.b128 [%rd3], %rq1;
 ; CHECK-NEXT:    }
 ; CHECK-NEXT:    // end inline asm
-; CHECK-NEXT:    ld.global.u32 %r5, [v64];
-; CHECK-NEXT:    add.s32 %r6, %r5, 16;
-; CHECK-NEXT:    add.s32 %r7, %r5, 24;
-; CHECK-NEXT:    mov.u64 %rd14, 9223372036854775807;
-; CHECK-NEXT:    mov.b128 %rq2, {%rd13, %rd14};
-; CHECK-NEXT:    mov.u32 %r8, v_i128_max;
-; CHECK-NEXT:    cvta.global.u32 %r9, %r8;
-; CHECK-NEXT:    cvt.u64.u32 %rd6, %r9;
-; CHECK-NEXT:    cvt.u64.u32 %rd4, %r6;
-; CHECK-NEXT:    cvt.u64.u32 %rd5, %r7;
+; CHECK-NEXT:    ld.global.u64 %rd15, [v64];
+; CHECK-NEXT:    add.s64 %rd4, %rd15, 16;
+; CHECK-NEXT:    add.s64 %rd5, %rd15, 24;
+; CHECK-NEXT:    mov.u64 %rd16, 9223372036854775807;
+; CHECK-NEXT:    mov.b128 %rq2, {%rd13, %rd16};
+; CHECK-NEXT:    mov.u64 %rd17, v_i128_max;
+; CHECK-NEXT:    cvta.global.u64 %rd6, %rd17;
 ; CHECK-NEXT:    // begin inline asm
 ; CHECK-NEXT:    {
 ; CHECK-NEXT:    .reg .b64 hi;
@@ -61,17 +54,14 @@ define void @test_corner_values() {
 ; CHECK-NEXT:    st.b128 [%rd6], %rq2;
 ; CHECK-NEXT:    }
 ; CHECK-NEXT:    // end inline asm
-; CHECK-NEXT:    ld.global.u32 %r10, [v64];
-; CHECK-NEXT:    add.s32 %r11, %r10, 32;
-; CHECK-NEXT:    add.s32 %r12, %r10, 40;
-; CHECK-NEXT:    mov.u64 %rd15, -9223372036854775808;
-; CHECK-NEXT:    mov.u64 %rd16, 0;
-; CHECK-NEXT:    mov.b128 %rq3, {%rd16, %rd15};
-; CHECK-NEXT:    mov.u32 %r13, v_i128_min;
-; CHECK-NEXT:    cvta.global.u32 %r14, %r13;
-; CHECK-NEXT:    cvt.u64.u32 %rd9, %r14;
-; CHECK-NEXT:    cvt.u64.u32 %rd7, %r11;
-; CHECK-NEXT:    cvt.u64.u32 %rd8, %r12;
+; CHECK-NEXT:    ld.global.u64 %rd18, [v64];
+; CHECK-NEXT:    add.s64 %rd7, %rd18, 32;
+; CHECK-NEXT:    add.s64 %rd8, %rd18, 40;
+; CHECK-NEXT:    mov.u64 %rd19, -9223372036854775808;
+; CHECK-NEXT:    mov.u64 %rd20, 0;
+; CHECK-NEXT:    mov.b128 %rq3, {%rd20, %rd19};
+; CHECK-NEXT:    mov.u64 %rd21, v_i128_min;
+; CHECK-NEXT:    cvta.global.u64 %rd9, %rd21;
 ; CHECK-NEXT:    // begin inline asm
 ; CHECK-NEXT:    {
 ; CHECK-NEXT:    .reg .b64 hi;
@@ -82,15 +72,12 @@ define void @test_corner_values() {
 ; CHECK-NEXT:    st.b128 [%rd9], %rq3;
 ; CHECK-NEXT:    }
 ; CHECK-NEXT:    // end inline asm
-; CHECK-NEXT:    ld.global.u32 %r15, [v64];
-; CHECK-NEXT:    add.s32 %r16, %r15, 48;
-; CHECK-NEXT:    add.s32 %r17, %r15, 56;
-; CHECK-NEXT:    mov.b128 %rq4, {%rd16, %rd16};
-; CHECK-NEXT:    mov.u32 %r18, v_u128_zero;
-; CHECK-NEXT:    cvta.global.u32 %r19, %r18;
-; CHECK-NEXT:    cvt.u64.u32 %rd12, %r19;
-; CHECK-NEXT:    cvt.u64.u32 %rd10, %r16;
-; CHECK-NEXT:    cvt.u64.u32 %rd11, %r17;
+; CHECK-NEXT:    ld.global.u64 %rd22, [v64];
+; CHECK-NEXT:    add.s64 %rd10, %rd22, 48;
+; CHECK-NEXT:    add.s64 %rd11, %rd22, 56;
+; CHECK-NEXT:    mov.b128 %rq4, {%rd20, %rd20};
+; CHECK-NEXT:    mov.u64 %rd23, v_u128_zero;
+; CHECK-NEXT:    cvta.global.u64 %rd12, %rd23;
 ; CHECK-NEXT:    // begin inline asm
 ; CHECK-NEXT:    {
 ; CHECK-NEXT:    .reg .b64 hi;

llvm/test/CodeGen/NVPTX/inline-asm-b128-test3.ll

@@ -1,6 +1,6 @@
 ; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --extra_scrub --version 5
-; RUN: llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
-; RUN: %if ptxas %{ llc < %s -march=nvptx -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx83 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_70 -mattr=+ptx83 | %ptxas-verify -arch=sm_70 %}
 
 target triple = "nvptx64-nvidia-cuda"