ptx test update

Author: YonahGoldberg

llvm/test/CodeGen/NVPTX/lower-byval-args.ll

@@ -455,64 +455,42 @@ define dso_local ptx_kernel void @memcpy_to_param(ptr nocapture noundef readonly
 ; PTX-NEXT:    .local .align 8 .b8 __local_depot9[8];
 ; PTX-NEXT:    .reg .b64 %SP;
 ; PTX-NEXT:    .reg .b64 %SPL;
-; PTX-NEXT:    .reg .b32 %r<3>;
-; PTX-NEXT:    .reg .b64 %rd<47>;
+; PTX-NEXT:    .reg .b64 %rd<30>;
 ; PTX-EMPTY:
 ; PTX-NEXT:  // %bb.0: // %entry
 ; PTX-NEXT:    mov.b64 %SPL, __local_depot9;
 ; PTX-NEXT:    cvta.local.u64 %SP, %SPL;
 ; PTX-NEXT:    ld.param.b64 %rd1, [memcpy_to_param_param_0];
-; PTX-NEXT:    add.u64 %rd2, %SPL, 0;
-; PTX-NEXT:    ld.param.b32 %r1, [memcpy_to_param_param_1+4];
-; PTX-NEXT:    st.local.b32 [%rd2+4], %r1;
-; PTX-NEXT:    ld.param.b32 %r2, [memcpy_to_param_param_1];
-; PTX-NEXT:    st.local.b32 [%rd2], %r2;
-; PTX-NEXT:    ld.volatile.b8 %rd3, [%rd1];
-; PTX-NEXT:    ld.volatile.b8 %rd4, [%rd1+1];
-; PTX-NEXT:    shl.b64 %rd5, %rd4, 8;
-; PTX-NEXT:    or.b64 %rd6, %rd5, %rd3;
-; PTX-NEXT:    ld.volatile.b8 %rd7, [%rd1+2];
-; PTX-NEXT:    shl.b64 %rd8, %rd7, 16;
-; PTX-NEXT:    ld.volatile.b8 %rd9, [%rd1+3];
-; PTX-NEXT:    shl.b64 %rd10, %rd9, 24;
-; PTX-NEXT:    or.b64 %rd11, %rd10, %rd8;
-; PTX-NEXT:    or.b64 %rd12, %rd11, %rd6;
-; PTX-NEXT:    ld.volatile.b8 %rd13, [%rd1+4];
-; PTX-NEXT:    ld.volatile.b8 %rd14, [%rd1+5];
-; PTX-NEXT:    shl.b64 %rd15, %rd14, 8;
-; PTX-NEXT:    or.b64 %rd16, %rd15, %rd13;
-; PTX-NEXT:    ld.volatile.b8 %rd17, [%rd1+6];
-; PTX-NEXT:    shl.b64 %rd18, %rd17, 16;
-; PTX-NEXT:    ld.volatile.b8 %rd19, [%rd1+7];
-; PTX-NEXT:    shl.b64 %rd20, %rd19, 24;
-; PTX-NEXT:    or.b64 %rd21, %rd20, %rd18;
-; PTX-NEXT:    or.b64 %rd22, %rd21, %rd16;
-; PTX-NEXT:    shl.b64 %rd23, %rd22, 32;
-; PTX-NEXT:    or.b64 %rd24, %rd23, %rd12;
-; PTX-NEXT:    st.volatile.b64 [%SP], %rd24;
-; PTX-NEXT:    ld.volatile.b8 %rd25, [%rd1+8];
-; PTX-NEXT:    ld.volatile.b8 %rd26, [%rd1+9];
-; PTX-NEXT:    shl.b64 %rd27, %rd26, 8;
-; PTX-NEXT:    or.b64 %rd28, %rd27, %rd25;
-; PTX-NEXT:    ld.volatile.b8 %rd29, [%rd1+10];
-; PTX-NEXT:    shl.b64 %rd30, %rd29, 16;
-; PTX-NEXT:    ld.volatile.b8 %rd31, [%rd1+11];
-; PTX-NEXT:    shl.b64 %rd32, %rd31, 24;
-; PTX-NEXT:    or.b64 %rd33, %rd32, %rd30;
-; PTX-NEXT:    or.b64 %rd34, %rd33, %rd28;
-; PTX-NEXT:    ld.volatile.b8 %rd35, [%rd1+12];
-; PTX-NEXT:    ld.volatile.b8 %rd36, [%rd1+13];
-; PTX-NEXT:    shl.b64 %rd37, %rd36, 8;
-; PTX-NEXT:    or.b64 %rd38, %rd37, %rd35;
-; PTX-NEXT:    ld.volatile.b8 %rd39, [%rd1+14];
-; PTX-NEXT:    shl.b64 %rd40, %rd39, 16;
-; PTX-NEXT:    ld.volatile.b8 %rd41, [%rd1+15];
-; PTX-NEXT:    shl.b64 %rd42, %rd41, 24;
-; PTX-NEXT:    or.b64 %rd43, %rd42, %rd40;
-; PTX-NEXT:    or.b64 %rd44, %rd43, %rd38;
-; PTX-NEXT:    shl.b64 %rd45, %rd44, 32;
-; PTX-NEXT:    or.b64 %rd46, %rd45, %rd34;
-; PTX-NEXT:    st.volatile.b64 [%SP+8], %rd46;
+; PTX-NEXT:    cvta.to.global.u64 %rd2, %rd1;
+; PTX-NEXT:    ld.param.b32 %rd3, [memcpy_to_param_param_1+4];
+; PTX-NEXT:    shl.b64 %rd4, %rd3, 32;
+; PTX-NEXT:    ld.param.b32 %rd5, [memcpy_to_param_param_1];
+; PTX-NEXT:    or.b64 %rd6, %rd4, %rd5;
+; PTX-NEXT:    st.b64 [%SP], %rd6;
+; PTX-NEXT:    ld.volatile.global.b8 %rd7, [%rd2];
+; PTX-NEXT:    ld.volatile.global.b8 %rd8, [%rd2+1];
+; PTX-NEXT:    shl.b64 %rd9, %rd8, 8;
+; PTX-NEXT:    or.b64 %rd10, %rd9, %rd7;
+; PTX-NEXT:    ld.volatile.global.b8 %rd11, [%rd2+2];
+; PTX-NEXT:    shl.b64 %rd12, %rd11, 16;
+; PTX-NEXT:    ld.volatile.global.b8 %rd13, [%rd2+3];
+; PTX-NEXT:    shl.b64 %rd14, %rd13, 24;
+; PTX-NEXT:    or.b64 %rd15, %rd14, %rd12;
+; PTX-NEXT:    or.b64 %rd16, %rd15, %rd10;
+; PTX-NEXT:    ld.volatile.global.b8 %rd17, [%rd2+4];
+; PTX-NEXT:    ld.volatile.global.b8 %rd18, [%rd2+5];
+; PTX-NEXT:    shl.b64 %rd19, %rd18, 8;
+; PTX-NEXT:    or.b64 %rd20, %rd19, %rd17;
+; PTX-NEXT:    ld.volatile.global.b8 %rd21, [%rd2+6];
+; PTX-NEXT:    shl.b64 %rd22, %rd21, 16;
+; PTX-NEXT:    ld.volatile.global.b8 %rd23, [%rd2+7];
+; PTX-NEXT:    shl.b64 %rd24, %rd23, 24;
+; PTX-NEXT:    or.b64 %rd25, %rd24, %rd22;
+; PTX-NEXT:    or.b64 %rd26, %rd25, %rd20;
+; PTX-NEXT:    shl.b64 %rd27, %rd26, 32;
+; PTX-NEXT:    or.b64 %rd28, %rd27, %rd16;
+; PTX-NEXT:    add.u64 %rd29, %SPL, 0;
+; PTX-NEXT:    st.local.b64 [%rd29], %rd28;
 ; PTX-NEXT:    ret;
 entry:
   tail call void @llvm.memcpy.p0.p0.i64(ptr %s, ptr %in, i64 16, i1 true)