381 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64-- -mtriple=nvptx64 -mcpu=sm_52 -mattr=+ptx64 < %s | FileCheck %s --check-prefix=CHECK-PTX3; RUN: %if ptxas-isa-6.4 %{ llc < %s -mtriple=nvptx64-- -mtriple=nvptx64 -mcpu=sm_52 -mattr=+ptx64 | %ptxas-verify %}4 5%struct.S1 = type { i32, i8, i64 }6%struct.S2 = type { i64, i64 }7 8@__const.bar.s1 = private unnamed_addr constant %struct.S1 { i32 1, i8 1, i64 1 }, align 89@__const.qux.s = private unnamed_addr constant %struct.S2 { i64 1, i64 1 }, align 810 11define dso_local i32 @variadics1(i32 noundef %first, ...) {12; CHECK-PTX-LABEL: variadics1(13; CHECK-PTX: {14; CHECK-PTX-NEXT: .reg .b32 %r<11>;15; CHECK-PTX-NEXT: .reg .b64 %rd<17>;16; CHECK-PTX-EMPTY:17; CHECK-PTX-NEXT: // %bb.0: // %entry18; CHECK-PTX-NEXT: ld.param.b32 %r1, [variadics1_param_0];19; CHECK-PTX-NEXT: ld.param.b64 %rd1, [variadics1_param_1];20; CHECK-PTX-NEXT: ld.b32 %r2, [%rd1];21; CHECK-PTX-NEXT: add.s32 %r3, %r1, %r2;22; CHECK-PTX-NEXT: ld.b32 %r4, [%rd1+4];23; CHECK-PTX-NEXT: add.s32 %r5, %r3, %r4;24; CHECK-PTX-NEXT: ld.b32 %r6, [%rd1+8];25; CHECK-PTX-NEXT: add.s32 %r7, %r5, %r6;26; CHECK-PTX-NEXT: add.s64 %rd2, %rd1, 19;27; CHECK-PTX-NEXT: and.b64 %rd3, %rd2, -8;28; CHECK-PTX-NEXT: ld.b64 %rd4, [%rd3];29; CHECK-PTX-NEXT: cvt.u64.u32 %rd5, %r7;30; CHECK-PTX-NEXT: add.s64 %rd6, %rd5, %rd4;31; CHECK-PTX-NEXT: cvt.u32.u64 %r8, %rd6;32; CHECK-PTX-NEXT: add.s64 %rd7, %rd3, 15;33; CHECK-PTX-NEXT: and.b64 %rd8, %rd7, -8;34; CHECK-PTX-NEXT: ld.b64 %rd9, [%rd8];35; CHECK-PTX-NEXT: cvt.rn.f64.s32 %rd10, %r8;36; CHECK-PTX-NEXT: add.rn.f64 %rd11, %rd10, %rd9;37; CHECK-PTX-NEXT: cvt.rzi.s32.f64 %r9, %rd11;38; CHECK-PTX-NEXT: add.s64 %rd12, %rd8, 15;39; CHECK-PTX-NEXT: and.b64 %rd13, %rd12, -8;40; CHECK-PTX-NEXT: ld.b64 %rd14, [%rd13];41; CHECK-PTX-NEXT: cvt.rn.f64.s32 %rd15, %r9;42; CHECK-PTX-NEXT: add.rn.f64 %rd16, %rd15, %rd14;43; CHECK-PTX-NEXT: cvt.rzi.s32.f64 %r10, %rd16;44; CHECK-PTX-NEXT: st.param.b32 [func_retval0], %r10;45; CHECK-PTX-NEXT: ret;46entry:47 %vlist = alloca ptr, align 848 call void @llvm.va_start.p0(ptr %vlist)49 %argp.cur = load ptr, ptr %vlist, align 850 %argp.next = getelementptr inbounds i8, ptr %argp.cur, i64 451 store ptr %argp.next, ptr %vlist, align 852 %0 = load i32, ptr %argp.cur, align 453 %add = add nsw i32 %first, %054 %argp.cur1 = load ptr, ptr %vlist, align 855 %argp.next2 = getelementptr inbounds i8, ptr %argp.cur1, i64 456 store ptr %argp.next2, ptr %vlist, align 857 %1 = load i32, ptr %argp.cur1, align 458 %add3 = add nsw i32 %add, %159 %argp.cur4 = load ptr, ptr %vlist, align 860 %argp.next5 = getelementptr inbounds i8, ptr %argp.cur4, i64 461 store ptr %argp.next5, ptr %vlist, align 862 %2 = load i32, ptr %argp.cur4, align 463 %add6 = add nsw i32 %add3, %264 %argp.cur7 = load ptr, ptr %vlist, align 865 %3 = getelementptr inbounds i8, ptr %argp.cur7, i32 766 %argp.cur7.aligned = call ptr @llvm.ptrmask.p0.i64(ptr %3, i64 -8)67 %argp.next8 = getelementptr inbounds i8, ptr %argp.cur7.aligned, i64 868 store ptr %argp.next8, ptr %vlist, align 869 %4 = load i64, ptr %argp.cur7.aligned, align 870 %conv = sext i32 %add6 to i6471 %add9 = add nsw i64 %conv, %472 %conv10 = trunc i64 %add9 to i3273 %argp.cur11 = load ptr, ptr %vlist, align 874 %5 = getelementptr inbounds i8, ptr %argp.cur11, i32 775 %argp.cur11.aligned = call ptr @llvm.ptrmask.p0.i64(ptr %5, i64 -8)76 %argp.next12 = getelementptr inbounds i8, ptr %argp.cur11.aligned, i64 877 store ptr %argp.next12, ptr %vlist, align 878 %6 = load double, ptr %argp.cur11.aligned, align 879 %conv13 = sitofp i32 %conv10 to double80 %add14 = fadd double %conv13, %681 %conv15 = fptosi double %add14 to i3282 %argp.cur16 = load ptr, ptr %vlist, align 883 %7 = getelementptr inbounds i8, ptr %argp.cur16, i32 784 %argp.cur16.aligned = call ptr @llvm.ptrmask.p0.i64(ptr %7, i64 -8)85 %argp.next17 = getelementptr inbounds i8, ptr %argp.cur16.aligned, i64 886 store ptr %argp.next17, ptr %vlist, align 887 %8 = load double, ptr %argp.cur16.aligned, align 888 %conv18 = sitofp i32 %conv15 to double89 %add19 = fadd double %conv18, %890 %conv20 = fptosi double %add19 to i3291 call void @llvm.va_end.p0(ptr %vlist)92 ret i32 %conv2093}94 95declare void @llvm.va_start.p0(ptr)96 97declare ptr @llvm.ptrmask.p0.i64(ptr, i64)98 99declare void @llvm.va_end.p0(ptr)100 101define dso_local i32 @foo() {102; CHECK-PTX-LABEL: foo(103; CHECK-PTX: {104; CHECK-PTX-NEXT: .local .align 8 .b8 __local_depot1[40];105; CHECK-PTX-NEXT: .reg .b64 %SP;106; CHECK-PTX-NEXT: .reg .b64 %SPL;107; CHECK-PTX-NEXT: .reg .b32 %r<2>;108; CHECK-PTX-NEXT: .reg .b64 %rd<2>;109; CHECK-PTX-EMPTY:110; CHECK-PTX-NEXT: // %bb.0: // %entry111; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot1;112; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;113; CHECK-PTX-NEXT: st.b64 [%SP], 4294967297;114; CHECK-PTX-NEXT: st.b32 [%SP+8], 1;115; CHECK-PTX-NEXT: st.b64 [%SP+16], 1;116; CHECK-PTX-NEXT: st.b64 [%SP+24], 4607182418800017408;117; CHECK-PTX-NEXT: st.b64 [%SP+32], 4607182418800017408;118; CHECK-PTX-NEXT: { // callseq 0, 0119; CHECK-PTX-NEXT: .param .b32 param0;120; CHECK-PTX-NEXT: .param .b64 param1;121; CHECK-PTX-NEXT: .param .b32 retval0;122; CHECK-PTX-NEXT: add.u64 %rd1, %SP, 0;123; CHECK-PTX-NEXT: st.param.b64 [param1], %rd1;124; CHECK-PTX-NEXT: st.param.b32 [param0], 1;125; CHECK-PTX-NEXT: call.uni (retval0), variadics1, (param0, param1);126; CHECK-PTX-NEXT: ld.param.b32 %r1, [retval0];127; CHECK-PTX-NEXT: } // callseq 0128; CHECK-PTX-NEXT: st.param.b32 [func_retval0], %r1;129; CHECK-PTX-NEXT: ret;130entry:131 %conv = sext i8 1 to i32132 %conv1 = sext i16 1 to i32133 %conv2 = fpext float 1.000000e+00 to double134 %call = call i32 (i32, ...) @variadics1(i32 noundef 1, i32 noundef %conv, i32 noundef %conv1, i32 noundef 1, i64 noundef 1, double noundef %conv2, double noundef 1.000000e+00)135 ret i32 %call136}137 138define dso_local i32 @variadics2(i32 noundef %first, ...) {139; CHECK-PTX-LABEL: variadics2(140; CHECK-PTX: {141; CHECK-PTX-NEXT: .local .align 1 .b8 __local_depot2[3];142; CHECK-PTX-NEXT: .reg .b64 %SP;143; CHECK-PTX-NEXT: .reg .b64 %SPL;144; CHECK-PTX-NEXT: .reg .b16 %rs<4>;145; CHECK-PTX-NEXT: .reg .b32 %r<6>;146; CHECK-PTX-NEXT: .reg .b64 %rd<8>;147; CHECK-PTX-EMPTY:148; CHECK-PTX-NEXT: // %bb.0: // %entry149; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot2;150; CHECK-PTX-NEXT: ld.param.b32 %r1, [variadics2_param_0];151; CHECK-PTX-NEXT: ld.param.b64 %rd1, [variadics2_param_1];152; CHECK-PTX-NEXT: add.u64 %rd2, %SPL, 0;153; CHECK-PTX-NEXT: add.s64 %rd3, %rd1, 7;154; CHECK-PTX-NEXT: and.b64 %rd4, %rd3, -8;155; CHECK-PTX-NEXT: ld.b32 %r2, [%rd4];156; CHECK-PTX-NEXT: ld.s8 %r3, [%rd4+4];157; CHECK-PTX-NEXT: ld.b8 %rs1, [%rd4+7];158; CHECK-PTX-NEXT: st.local.b8 [%rd2+2], %rs1;159; CHECK-PTX-NEXT: ld.b8 %rs2, [%rd4+6];160; CHECK-PTX-NEXT: st.local.b8 [%rd2+1], %rs2;161; CHECK-PTX-NEXT: ld.b8 %rs3, [%rd4+5];162; CHECK-PTX-NEXT: st.local.b8 [%rd2], %rs3;163; CHECK-PTX-NEXT: ld.b64 %rd5, [%rd4+8];164; CHECK-PTX-NEXT: add.s32 %r4, %r1, %r2;165; CHECK-PTX-NEXT: add.s32 %r5, %r4, %r3;166; CHECK-PTX-NEXT: cvt.u64.u32 %rd6, %r5;167; CHECK-PTX-NEXT: add.s64 %rd7, %rd6, %rd5;168; CHECK-PTX-NEXT: st.param.b32 [func_retval0], %rd7;169; CHECK-PTX-NEXT: ret;170entry:171 %vlist = alloca ptr, align 8172 %s1.sroa.3 = alloca [3 x i8], align 1173 call void @llvm.va_start.p0(ptr %vlist)174 %argp.cur = load ptr, ptr %vlist, align 8175 %0 = getelementptr inbounds i8, ptr %argp.cur, i32 7176 %argp.cur.aligned = call ptr @llvm.ptrmask.p0.i64(ptr %0, i64 -8)177 %argp.next = getelementptr inbounds i8, ptr %argp.cur.aligned, i64 16178 store ptr %argp.next, ptr %vlist, align 8179 %s1.sroa.0.0.copyload = load i32, ptr %argp.cur.aligned, align 8180 %s1.sroa.2.0.argp.cur.aligned.sroa_idx = getelementptr inbounds i8, ptr %argp.cur.aligned, i64 4181 %s1.sroa.2.0.copyload = load i8, ptr %s1.sroa.2.0.argp.cur.aligned.sroa_idx, align 4182 %s1.sroa.3.0.argp.cur.aligned.sroa_idx = getelementptr inbounds i8, ptr %argp.cur.aligned, i64 5183 call void @llvm.memcpy.p0.p0.i64(ptr align 1 %s1.sroa.3, ptr align 1 %s1.sroa.3.0.argp.cur.aligned.sroa_idx, i64 3, i1 false)184 %s1.sroa.31.0.argp.cur.aligned.sroa_idx = getelementptr inbounds i8, ptr %argp.cur.aligned, i64 8185 %s1.sroa.31.0.copyload = load i64, ptr %s1.sroa.31.0.argp.cur.aligned.sroa_idx, align 8186 %add = add nsw i32 %first, %s1.sroa.0.0.copyload187 %conv = sext i8 %s1.sroa.2.0.copyload to i32188 %add1 = add nsw i32 %add, %conv189 %conv2 = sext i32 %add1 to i64190 %add3 = add nsw i64 %conv2, %s1.sroa.31.0.copyload191 %conv4 = trunc i64 %add3 to i32192 call void @llvm.va_end.p0(ptr %vlist)193 ret i32 %conv4194}195 196declare void @llvm.memcpy.p0.p0.i64(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg)197 198define dso_local i32 @bar() {199; CHECK-PTX-LABEL: bar(200; CHECK-PTX: {201; CHECK-PTX-NEXT: .local .align 8 .b8 __local_depot3[24];202; CHECK-PTX-NEXT: .reg .b64 %SP;203; CHECK-PTX-NEXT: .reg .b64 %SPL;204; CHECK-PTX-NEXT: .reg .b16 %rs<4>;205; CHECK-PTX-NEXT: .reg .b32 %r<2>;206; CHECK-PTX-NEXT: .reg .b64 %rd<3>;207; CHECK-PTX-EMPTY:208; CHECK-PTX-NEXT: // %bb.0: // %entry209; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot3;210; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;211; CHECK-PTX-NEXT: add.u64 %rd1, %SPL, 0;212; CHECK-PTX-NEXT: ld.global.nc.b8 %rs1, [__const_$_bar_$_s1+7];213; CHECK-PTX-NEXT: st.local.b8 [%rd1+2], %rs1;214; CHECK-PTX-NEXT: ld.global.nc.b8 %rs2, [__const_$_bar_$_s1+6];215; CHECK-PTX-NEXT: st.local.b8 [%rd1+1], %rs2;216; CHECK-PTX-NEXT: ld.global.nc.b8 %rs3, [__const_$_bar_$_s1+5];217; CHECK-PTX-NEXT: st.local.b8 [%rd1], %rs3;218; CHECK-PTX-NEXT: st.b32 [%SP+8], 1;219; CHECK-PTX-NEXT: st.b8 [%SP+12], 1;220; CHECK-PTX-NEXT: st.b64 [%SP+16], 1;221; CHECK-PTX-NEXT: { // callseq 1, 0222; CHECK-PTX-NEXT: .param .b32 param0;223; CHECK-PTX-NEXT: .param .b64 param1;224; CHECK-PTX-NEXT: .param .b32 retval0;225; CHECK-PTX-NEXT: add.u64 %rd2, %SP, 8;226; CHECK-PTX-NEXT: st.param.b64 [param1], %rd2;227; CHECK-PTX-NEXT: st.param.b32 [param0], 1;228; CHECK-PTX-NEXT: call.uni (retval0), variadics2, (param0, param1);229; CHECK-PTX-NEXT: ld.param.b32 %r1, [retval0];230; CHECK-PTX-NEXT: } // callseq 1231; CHECK-PTX-NEXT: st.param.b32 [func_retval0], %r1;232; CHECK-PTX-NEXT: ret;233entry:234 %s1.sroa.3 = alloca [3 x i8], align 1235 %s1.sroa.0.0.copyload = load i32, ptr @__const.bar.s1, align 8236 %s1.sroa.2.0.copyload = load i8, ptr getelementptr inbounds (i8, ptr @__const.bar.s1, i64 4), align 4237 call void @llvm.memcpy.p0.p0.i64(ptr align 1 %s1.sroa.3, ptr align 1 getelementptr inbounds (i8, ptr @__const.bar.s1, i64 5), i64 3, i1 false)238 %s1.sroa.31.0.copyload = load i64, ptr getelementptr inbounds (i8, ptr @__const.bar.s1, i64 8), align 8239 %call = call i32 (i32, ...) @variadics2(i32 noundef 1, i32 %s1.sroa.0.0.copyload, i8 %s1.sroa.2.0.copyload, i64 %s1.sroa.31.0.copyload)240 ret i32 %call241}242 243define dso_local i32 @variadics3(i32 noundef %first, ...) {244; CHECK-PTX-LABEL: variadics3(245; CHECK-PTX: {246; CHECK-PTX-NEXT: .reg .b32 %r<8>;247; CHECK-PTX-NEXT: .reg .b64 %rd<4>;248; CHECK-PTX-EMPTY:249; CHECK-PTX-NEXT: // %bb.0: // %entry250; CHECK-PTX-NEXT: ld.param.b64 %rd1, [variadics3_param_1];251; CHECK-PTX-NEXT: add.s64 %rd2, %rd1, 15;252; CHECK-PTX-NEXT: and.b64 %rd3, %rd2, -16;253; CHECK-PTX-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd3];254; CHECK-PTX-NEXT: add.s32 %r5, %r1, %r2;255; CHECK-PTX-NEXT: add.s32 %r6, %r5, %r3;256; CHECK-PTX-NEXT: add.s32 %r7, %r6, %r4;257; CHECK-PTX-NEXT: st.param.b32 [func_retval0], %r7;258; CHECK-PTX-NEXT: ret;259entry:260 %vlist = alloca ptr, align 8261 call void @llvm.va_start.p0(ptr %vlist)262 %argp.cur = load ptr, ptr %vlist, align 8263 %0 = getelementptr inbounds i8, ptr %argp.cur, i32 15264 %argp.cur.aligned = call ptr @llvm.ptrmask.p0.i64(ptr %0, i64 -16)265 %argp.next = getelementptr inbounds i8, ptr %argp.cur.aligned, i64 16266 store ptr %argp.next, ptr %vlist, align 8267 %1 = load <4 x i32>, ptr %argp.cur.aligned, align 16268 call void @llvm.va_end.p0(ptr %vlist)269 %2 = extractelement <4 x i32> %1, i64 0270 %3 = extractelement <4 x i32> %1, i64 1271 %add = add nsw i32 %2, %3272 %4 = extractelement <4 x i32> %1, i64 2273 %add1 = add nsw i32 %add, %4274 %5 = extractelement <4 x i32> %1, i64 3275 %add2 = add nsw i32 %add1, %5276 ret i32 %add2277}278 279define dso_local i32 @baz() {280; CHECK-PTX-LABEL: baz(281; CHECK-PTX: {282; CHECK-PTX-NEXT: .local .align 16 .b8 __local_depot5[16];283; CHECK-PTX-NEXT: .reg .b64 %SP;284; CHECK-PTX-NEXT: .reg .b64 %SPL;285; CHECK-PTX-NEXT: .reg .b32 %r<2>;286; CHECK-PTX-NEXT: .reg .b64 %rd<2>;287; CHECK-PTX-EMPTY:288; CHECK-PTX-NEXT: // %bb.0: // %entry289; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot5;290; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;291; CHECK-PTX-NEXT: st.v4.b32 [%SP], {1, 1, 1, 1};292; CHECK-PTX-NEXT: { // callseq 2, 0293; CHECK-PTX-NEXT: .param .b32 param0;294; CHECK-PTX-NEXT: .param .b64 param1;295; CHECK-PTX-NEXT: .param .b32 retval0;296; CHECK-PTX-NEXT: add.u64 %rd1, %SP, 0;297; CHECK-PTX-NEXT: st.param.b64 [param1], %rd1;298; CHECK-PTX-NEXT: st.param.b32 [param0], 1;299; CHECK-PTX-NEXT: call.uni (retval0), variadics3, (param0, param1);300; CHECK-PTX-NEXT: ld.param.b32 %r1, [retval0];301; CHECK-PTX-NEXT: } // callseq 2302; CHECK-PTX-NEXT: st.param.b32 [func_retval0], %r1;303; CHECK-PTX-NEXT: ret;304entry:305 %call = call i32 (i32, ...) @variadics3(i32 noundef 1, <4 x i32> noundef <i32 1, i32 1, i32 1, i32 1>)306 ret i32 %call307}308 309define dso_local i32 @variadics4(ptr noundef byval(%struct.S2) align 8 %first, ...) {310; CHECK-PTX-LABEL: variadics4(311; CHECK-PTX: {312; CHECK-PTX-NEXT: .reg .b64 %rd<9>;313; CHECK-PTX-EMPTY:314; CHECK-PTX-NEXT: // %bb.0: // %entry315; CHECK-PTX-NEXT: ld.param.b64 %rd1, [variadics4_param_1];316; CHECK-PTX-NEXT: add.s64 %rd2, %rd1, 7;317; CHECK-PTX-NEXT: and.b64 %rd3, %rd2, -8;318; CHECK-PTX-NEXT: ld.b64 %rd4, [%rd3];319; CHECK-PTX-NEXT: ld.param.b64 %rd5, [variadics4_param_0];320; CHECK-PTX-NEXT: ld.param.b64 %rd6, [variadics4_param_0+8];321; CHECK-PTX-NEXT: add.s64 %rd7, %rd5, %rd6;322; CHECK-PTX-NEXT: add.s64 %rd8, %rd7, %rd4;323; CHECK-PTX-NEXT: st.param.b32 [func_retval0], %rd8;324; CHECK-PTX-NEXT: ret;325entry:326 %vlist = alloca ptr, align 8327 call void @llvm.va_start.p0(ptr %vlist)328 %argp.cur = load ptr, ptr %vlist, align 8329 %0 = getelementptr inbounds i8, ptr %argp.cur, i32 7330 %argp.cur.aligned = call ptr @llvm.ptrmask.p0.i64(ptr %0, i64 -8)331 %argp.next = getelementptr inbounds i8, ptr %argp.cur.aligned, i64 8332 store ptr %argp.next, ptr %vlist, align 8333 %1 = load i64, ptr %argp.cur.aligned, align 8334 %x1 = getelementptr inbounds %struct.S2, ptr %first, i32 0, i32 0335 %2 = load i64, ptr %x1, align 8336 %y = getelementptr inbounds %struct.S2, ptr %first, i32 0, i32 1337 %3 = load i64, ptr %y, align 8338 %add = add nsw i64 %2, %3339 %add2 = add nsw i64 %add, %1340 %conv = trunc i64 %add2 to i32341 call void @llvm.va_end.p0(ptr %vlist)342 ret i32 %conv343}344 345define dso_local void @qux() {346; CHECK-PTX-LABEL: qux(347; CHECK-PTX: {348; CHECK-PTX-NEXT: .local .align 8 .b8 __local_depot7[24];349; CHECK-PTX-NEXT: .reg .b64 %SP;350; CHECK-PTX-NEXT: .reg .b64 %SPL;351; CHECK-PTX-NEXT: .reg .b64 %rd<7>;352; CHECK-PTX-EMPTY:353; CHECK-PTX-NEXT: // %bb.0: // %entry354; CHECK-PTX-NEXT: mov.b64 %SPL, __local_depot7;355; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;356; CHECK-PTX-NEXT: add.u64 %rd1, %SPL, 0;357; CHECK-PTX-NEXT: ld.global.nc.b64 %rd2, [__const_$_qux_$_s+8];358; CHECK-PTX-NEXT: st.local.b64 [%rd1+8], %rd2;359; CHECK-PTX-NEXT: ld.global.nc.b64 %rd3, [__const_$_qux_$_s];360; CHECK-PTX-NEXT: st.local.b64 [%rd1], %rd3;361; CHECK-PTX-NEXT: st.b64 [%SP+16], 1;362; CHECK-PTX-NEXT: { // callseq 3, 0363; CHECK-PTX-NEXT: .param .align 8 .b8 param0[16];364; CHECK-PTX-NEXT: .param .b64 param1;365; CHECK-PTX-NEXT: .param .b32 retval0;366; CHECK-PTX-NEXT: add.u64 %rd4, %SP, 16;367; CHECK-PTX-NEXT: st.param.b64 [param1], %rd4;368; CHECK-PTX-NEXT: ld.local.b64 %rd5, [%rd1+8];369; CHECK-PTX-NEXT: st.param.b64 [param0+8], %rd5;370; CHECK-PTX-NEXT: ld.local.b64 %rd6, [%rd1];371; CHECK-PTX-NEXT: st.param.b64 [param0], %rd6;372; CHECK-PTX-NEXT: call.uni (retval0), variadics4, (param0, param1);373; CHECK-PTX-NEXT: } // callseq 3374; CHECK-PTX-NEXT: ret;375entry:376 %s = alloca %struct.S2, align 8377 call void @llvm.memcpy.p0.p0.i64(ptr align 8 %s, ptr align 8 @__const.qux.s, i64 16, i1 false)378 %call = call i32 (ptr, ...) @variadics4(ptr noundef byval(%struct.S2) align 8 %s, i64 noundef 1)379 ret void380}381