918 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: opt < %s -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -nvptx-lower-args -S | FileCheck %s --check-prefixes=COMMON,LOWER-ARGS,SM_603; RUN: opt < %s -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -nvptx-lower-args -S | FileCheck %s --check-prefixes=COMMON,LOWER-ARGS,SM_704; RUN: opt < %s -mtriple nvptx64 -mcpu=sm_60 -mattr=ptx77 -passes=nvptx-lower-args -S | FileCheck %s --check-prefixes=COMMON,LOWER-ARGS,SM_605; RUN: opt < %s -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-lower-args -S | FileCheck %s --check-prefixes=COMMON,LOWER-ARGS,SM_706; RUN: opt < %s -mtriple nvptx64 -mcpu=sm_70 -mattr=ptx77 -passes=nvptx-copy-byval-args -S | FileCheck %s --check-prefixes=COMMON,COPY7; RUN: llc < %s -mcpu=sm_60 -mattr=ptx77 | FileCheck %s --check-prefixes=PTX,PTX_608; RUN: llc < %s -mcpu=sm_70 -mattr=ptx77 | FileCheck %s --check-prefixes=PTX,PTX_709source_filename = "<stdin>"10target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"11target triple = "nvptx64-nvidia-cuda"12 13%struct.S = type { i32, i32 }14 15; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)16declare dso_local void @_Z6escapePv(ptr noundef) local_unnamed_addr #017 18; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)19declare dso_local void @_Z6escapei(i32 noundef) local_unnamed_addr #020 21; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite)22declare void @llvm.memcpy.p0.p0.i64(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg) #123 24; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: readwrite)25declare void @llvm.memmove.p0.p0.i64(ptr nocapture writeonly, ptr nocapture readonly, i64, i1 immarg) #126 27; Function Attrs: nocallback nofree nounwind willreturn memory(argmem: write)28declare void @llvm.memset.p0.i64(ptr nocapture writeonly, i8, i64, i1 immarg) #229 30; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)31define dso_local ptx_kernel void @read_only(ptr nocapture noundef writeonly %out, ptr nocapture noundef readonly byval(%struct.S) align 4 %s) local_unnamed_addr #0 {32; LOWER-ARGS-LABEL: define dso_local ptx_kernel void @read_only(33; LOWER-ARGS-SAME: ptr noundef writeonly captures(none) [[OUT:%.*]], ptr noundef readonly byval([[STRUCT_S:%.*]]) align 4 captures(none) [[S:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {34; LOWER-ARGS-NEXT: [[ENTRY:.*:]]35; LOWER-ARGS-NEXT: [[S3:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])36; LOWER-ARGS-NEXT: [[I:%.*]] = load i32, ptr addrspace(101) [[S3]], align 437; LOWER-ARGS-NEXT: store i32 [[I]], ptr [[OUT]], align 438; LOWER-ARGS-NEXT: ret void39;40; COPY-LABEL: define dso_local ptx_kernel void @read_only(41; COPY-SAME: ptr noundef writeonly captures(none) [[OUT:%.*]], ptr noundef readonly byval([[STRUCT_S:%.*]]) align 4 captures(none) [[S:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {42; COPY-NEXT: [[ENTRY:.*:]]43; COPY-NEXT: [[I:%.*]] = load i32, ptr [[S]], align 444; COPY-NEXT: store i32 [[I]], ptr [[OUT]], align 445; COPY-NEXT: ret void46;47; PTX-LABEL: read_only(48; PTX: {49; PTX-NEXT: .reg .b32 %r<2>;50; PTX-NEXT: .reg .b64 %rd<3>;51; PTX-EMPTY:52; PTX-NEXT: // %bb.0: // %entry53; PTX-NEXT: ld.param.b64 %rd1, [read_only_param_0];54; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;55; PTX-NEXT: ld.param.b32 %r1, [read_only_param_1];56; PTX-NEXT: st.global.b32 [%rd2], %r1;57; PTX-NEXT: ret;58entry:59 %i = load i32, ptr %s, align 460 store i32 %i, ptr %out, align 461 ret void62}63 64; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)65define dso_local ptx_kernel void @read_only_gep(ptr nocapture noundef writeonly %out, ptr nocapture noundef readonly byval(%struct.S) align 4 %s) local_unnamed_addr #0 {66; LOWER-ARGS-LABEL: define dso_local ptx_kernel void @read_only_gep(67; LOWER-ARGS-SAME: ptr noundef writeonly captures(none) [[OUT:%.*]], ptr noundef readonly byval([[STRUCT_S:%.*]]) align 4 captures(none) [[S:%.*]]) local_unnamed_addr #[[ATTR0]] {68; LOWER-ARGS-NEXT: [[ENTRY:.*:]]69; LOWER-ARGS-NEXT: [[S3:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])70; LOWER-ARGS-NEXT: [[B4:%.*]] = getelementptr inbounds i8, ptr addrspace(101) [[S3]], i64 471; LOWER-ARGS-NEXT: [[I:%.*]] = load i32, ptr addrspace(101) [[B4]], align 472; LOWER-ARGS-NEXT: store i32 [[I]], ptr [[OUT]], align 473; LOWER-ARGS-NEXT: ret void74;75; COPY-LABEL: define dso_local ptx_kernel void @read_only_gep(76; COPY-SAME: ptr noundef writeonly captures(none) [[OUT:%.*]], ptr noundef readonly byval([[STRUCT_S:%.*]]) align 4 captures(none) [[S:%.*]]) local_unnamed_addr #[[ATTR0]] {77; COPY-NEXT: [[ENTRY:.*:]]78; COPY-NEXT: [[B:%.*]] = getelementptr inbounds nuw i8, ptr [[S]], i64 479; COPY-NEXT: [[I:%.*]] = load i32, ptr [[B]], align 480; COPY-NEXT: store i32 [[I]], ptr [[OUT]], align 481; COPY-NEXT: ret void82;83; PTX-LABEL: read_only_gep(84; PTX: {85; PTX-NEXT: .reg .b32 %r<2>;86; PTX-NEXT: .reg .b64 %rd<3>;87; PTX-EMPTY:88; PTX-NEXT: // %bb.0: // %entry89; PTX-NEXT: ld.param.b64 %rd1, [read_only_gep_param_0];90; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;91; PTX-NEXT: ld.param.b32 %r1, [read_only_gep_param_1+4];92; PTX-NEXT: st.global.b32 [%rd2], %r1;93; PTX-NEXT: ret;94entry:95 %b = getelementptr inbounds nuw i8, ptr %s, i64 496 %i = load i32, ptr %b, align 497 store i32 %i, ptr %out, align 498 ret void99}100 101;; TODO: This test has been disabled because the addrspacecast is not legal on102;; sm_60, and not well supported within nvptx-lower-args. We should determine103;; in what cases it is safe to make assumptions about the address of a byval104;; parameter and improve our handling of addrspacecast in nvptx-lower-args.105 106; define dso_local ptx_kernel void @read_only_gep_asc(ptr nocapture noundef writeonly %out, ptr nocapture noundef readonly byval(%struct.S) align 4 %s) local_unnamed_addr #0 {107; entry:108; %b = getelementptr inbounds nuw i8, ptr %s, i64 4109; %asc = addrspacecast ptr %b to ptr addrspace(101)110; %i = load i32, ptr addrspace(101) %asc, align 4111; store i32 %i, ptr %out, align 4112; ret void113; }114;115; define dso_local ptx_kernel void @read_only_gep_asc0(ptr nocapture noundef writeonly %out, ptr nocapture noundef readonly byval(%struct.S) align 4 %s) local_unnamed_addr #0 {116; entry:117; %b = getelementptr inbounds nuw i8, ptr %s, i64 4118; %asc = addrspacecast ptr %b to ptr addrspace(101)119; %asc0 = addrspacecast ptr addrspace(101) %asc to ptr120; %i = load i32, ptr %asc0, align 4121; store i32 %i, ptr %out, align 4122; ret void123; }124 125; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)126define dso_local ptx_kernel void @escape_ptr(ptr nocapture noundef readnone %out, ptr noundef byval(%struct.S) align 4 %s) local_unnamed_addr #0 {127; COMMON-LABEL: define dso_local ptx_kernel void @escape_ptr(128; COMMON-SAME: ptr noundef readnone captures(none) [[OUT:%.*]], ptr noundef byval([[STRUCT_S:%.*]]) align 4 [[S:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {129; COMMON-NEXT: [[ENTRY:.*:]]130; COMMON-NEXT: [[S1:%.*]] = alloca [[STRUCT_S]], align 4131; COMMON-NEXT: [[S2:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])132; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[S1]], ptr addrspace(101) align 4 [[S2]], i64 8, i1 false)133; COMMON-NEXT: call void @_Z6escapePv(ptr noundef nonnull [[S1]]) #[[ATTR6:[0-9]+]]134; COMMON-NEXT: ret void135;136; PTX-LABEL: escape_ptr(137; PTX: {138; PTX-NEXT: .local .align 4 .b8 __local_depot2[8];139; PTX-NEXT: .reg .b64 %SP;140; PTX-NEXT: .reg .b64 %SPL;141; PTX-NEXT: .reg .b32 %r<3>;142; PTX-NEXT: .reg .b64 %rd<3>;143; PTX-EMPTY:144; PTX-NEXT: // %bb.0: // %entry145; PTX-NEXT: mov.b64 %SPL, __local_depot2;146; PTX-NEXT: cvta.local.u64 %SP, %SPL;147; PTX-NEXT: add.u64 %rd1, %SP, 0;148; PTX-NEXT: add.u64 %rd2, %SPL, 0;149; PTX-NEXT: ld.param.b32 %r1, [escape_ptr_param_1+4];150; PTX-NEXT: st.local.b32 [%rd2+4], %r1;151; PTX-NEXT: ld.param.b32 %r2, [escape_ptr_param_1];152; PTX-NEXT: st.local.b32 [%rd2], %r2;153; PTX-NEXT: { // callseq 0, 0154; PTX-NEXT: .param .b64 param0;155; PTX-NEXT: st.param.b64 [param0], %rd1;156; PTX-NEXT: call.uni _Z6escapePv, (param0);157; PTX-NEXT: } // callseq 0158; PTX-NEXT: ret;159entry:160 call void @_Z6escapePv(ptr noundef nonnull %s) #0161 ret void162}163 164; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)165define dso_local ptx_kernel void @escape_ptr_gep(ptr nocapture noundef readnone %out, ptr noundef byval(%struct.S) align 4 %s) local_unnamed_addr #0 {166; COMMON-LABEL: define dso_local ptx_kernel void @escape_ptr_gep(167; COMMON-SAME: ptr noundef readnone captures(none) [[OUT:%.*]], ptr noundef byval([[STRUCT_S:%.*]]) align 4 [[S:%.*]]) local_unnamed_addr #[[ATTR0]] {168; COMMON-NEXT: [[ENTRY:.*:]]169; COMMON-NEXT: [[S1:%.*]] = alloca [[STRUCT_S]], align 4170; COMMON-NEXT: [[S2:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])171; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[S1]], ptr addrspace(101) align 4 [[S2]], i64 8, i1 false)172; COMMON-NEXT: [[B:%.*]] = getelementptr inbounds nuw i8, ptr [[S1]], i64 4173; COMMON-NEXT: call void @_Z6escapePv(ptr noundef nonnull [[B]]) #[[ATTR6]]174; COMMON-NEXT: ret void175;176; PTX-LABEL: escape_ptr_gep(177; PTX: {178; PTX-NEXT: .local .align 4 .b8 __local_depot3[8];179; PTX-NEXT: .reg .b64 %SP;180; PTX-NEXT: .reg .b64 %SPL;181; PTX-NEXT: .reg .b32 %r<3>;182; PTX-NEXT: .reg .b64 %rd<4>;183; PTX-EMPTY:184; PTX-NEXT: // %bb.0: // %entry185; PTX-NEXT: mov.b64 %SPL, __local_depot3;186; PTX-NEXT: cvta.local.u64 %SP, %SPL;187; PTX-NEXT: add.u64 %rd1, %SP, 0;188; PTX-NEXT: add.u64 %rd2, %SPL, 0;189; PTX-NEXT: ld.param.b32 %r1, [escape_ptr_gep_param_1+4];190; PTX-NEXT: st.local.b32 [%rd2+4], %r1;191; PTX-NEXT: ld.param.b32 %r2, [escape_ptr_gep_param_1];192; PTX-NEXT: st.local.b32 [%rd2], %r2;193; PTX-NEXT: add.s64 %rd3, %rd1, 4;194; PTX-NEXT: { // callseq 1, 0195; PTX-NEXT: .param .b64 param0;196; PTX-NEXT: st.param.b64 [param0], %rd3;197; PTX-NEXT: call.uni _Z6escapePv, (param0);198; PTX-NEXT: } // callseq 1199; PTX-NEXT: ret;200entry:201 %b = getelementptr inbounds nuw i8, ptr %s, i64 4202 call void @_Z6escapePv(ptr noundef nonnull %b) #0203 ret void204}205 206; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)207define dso_local ptx_kernel void @escape_ptr_store(ptr nocapture noundef writeonly %out, ptr noundef byval(%struct.S) align 4 %s) local_unnamed_addr #0 {208; COMMON-LABEL: define dso_local ptx_kernel void @escape_ptr_store(209; COMMON-SAME: ptr noundef writeonly captures(none) [[OUT:%.*]], ptr noundef byval([[STRUCT_S:%.*]]) align 4 [[S:%.*]]) local_unnamed_addr #[[ATTR0]] {210; COMMON-NEXT: [[ENTRY:.*:]]211; COMMON-NEXT: [[S1:%.*]] = alloca [[STRUCT_S]], align 4212; COMMON-NEXT: [[S2:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])213; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[S1]], ptr addrspace(101) align 4 [[S2]], i64 8, i1 false)214; COMMON-NEXT: store ptr [[S1]], ptr [[OUT]], align 8215; COMMON-NEXT: ret void216;217; PTX-LABEL: escape_ptr_store(218; PTX: {219; PTX-NEXT: .local .align 4 .b8 __local_depot4[8];220; PTX-NEXT: .reg .b64 %SP;221; PTX-NEXT: .reg .b64 %SPL;222; PTX-NEXT: .reg .b32 %r<3>;223; PTX-NEXT: .reg .b64 %rd<5>;224; PTX-EMPTY:225; PTX-NEXT: // %bb.0: // %entry226; PTX-NEXT: mov.b64 %SPL, __local_depot4;227; PTX-NEXT: cvta.local.u64 %SP, %SPL;228; PTX-NEXT: ld.param.b64 %rd1, [escape_ptr_store_param_0];229; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;230; PTX-NEXT: add.u64 %rd3, %SP, 0;231; PTX-NEXT: add.u64 %rd4, %SPL, 0;232; PTX-NEXT: ld.param.b32 %r1, [escape_ptr_store_param_1+4];233; PTX-NEXT: st.local.b32 [%rd4+4], %r1;234; PTX-NEXT: ld.param.b32 %r2, [escape_ptr_store_param_1];235; PTX-NEXT: st.local.b32 [%rd4], %r2;236; PTX-NEXT: st.global.b64 [%rd2], %rd3;237; PTX-NEXT: ret;238entry:239 store ptr %s, ptr %out, align 8240 ret void241}242 243; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)244define dso_local ptx_kernel void @escape_ptr_gep_store(ptr nocapture noundef writeonly %out, ptr noundef byval(%struct.S) align 4 %s) local_unnamed_addr #0 {245; COMMON-LABEL: define dso_local ptx_kernel void @escape_ptr_gep_store(246; COMMON-SAME: ptr noundef writeonly captures(none) [[OUT:%.*]], ptr noundef byval([[STRUCT_S:%.*]]) align 4 [[S:%.*]]) local_unnamed_addr #[[ATTR0]] {247; COMMON-NEXT: [[ENTRY:.*:]]248; COMMON-NEXT: [[S1:%.*]] = alloca [[STRUCT_S]], align 4249; COMMON-NEXT: [[S2:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])250; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[S1]], ptr addrspace(101) align 4 [[S2]], i64 8, i1 false)251; COMMON-NEXT: [[B:%.*]] = getelementptr inbounds nuw i8, ptr [[S1]], i64 4252; COMMON-NEXT: store ptr [[B]], ptr [[OUT]], align 8253; COMMON-NEXT: ret void254;255; PTX-LABEL: escape_ptr_gep_store(256; PTX: {257; PTX-NEXT: .local .align 4 .b8 __local_depot5[8];258; PTX-NEXT: .reg .b64 %SP;259; PTX-NEXT: .reg .b64 %SPL;260; PTX-NEXT: .reg .b32 %r<3>;261; PTX-NEXT: .reg .b64 %rd<6>;262; PTX-EMPTY:263; PTX-NEXT: // %bb.0: // %entry264; PTX-NEXT: mov.b64 %SPL, __local_depot5;265; PTX-NEXT: cvta.local.u64 %SP, %SPL;266; PTX-NEXT: ld.param.b64 %rd1, [escape_ptr_gep_store_param_0];267; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;268; PTX-NEXT: add.u64 %rd3, %SP, 0;269; PTX-NEXT: add.u64 %rd4, %SPL, 0;270; PTX-NEXT: ld.param.b32 %r1, [escape_ptr_gep_store_param_1+4];271; PTX-NEXT: st.local.b32 [%rd4+4], %r1;272; PTX-NEXT: ld.param.b32 %r2, [escape_ptr_gep_store_param_1];273; PTX-NEXT: st.local.b32 [%rd4], %r2;274; PTX-NEXT: add.s64 %rd5, %rd3, 4;275; PTX-NEXT: st.global.b64 [%rd2], %rd5;276; PTX-NEXT: ret;277entry:278 %b = getelementptr inbounds nuw i8, ptr %s, i64 4279 store ptr %b, ptr %out, align 8280 ret void281}282 283; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)284define dso_local ptx_kernel void @escape_ptrtoint(ptr nocapture noundef writeonly %out, ptr noundef byval(%struct.S) align 4 %s) local_unnamed_addr #0 {285; COMMON-LABEL: define dso_local ptx_kernel void @escape_ptrtoint(286; COMMON-SAME: ptr noundef writeonly captures(none) [[OUT:%.*]], ptr noundef byval([[STRUCT_S:%.*]]) align 4 [[S:%.*]]) local_unnamed_addr #[[ATTR0]] {287; COMMON-NEXT: [[ENTRY:.*:]]288; COMMON-NEXT: [[S1:%.*]] = alloca [[STRUCT_S]], align 4289; COMMON-NEXT: [[S2:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])290; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[S1]], ptr addrspace(101) align 4 [[S2]], i64 8, i1 false)291; COMMON-NEXT: [[I:%.*]] = ptrtoint ptr [[S1]] to i64292; COMMON-NEXT: store i64 [[I]], ptr [[OUT]], align 8293; COMMON-NEXT: ret void294;295; PTX-LABEL: escape_ptrtoint(296; PTX: {297; PTX-NEXT: .local .align 4 .b8 __local_depot6[8];298; PTX-NEXT: .reg .b64 %SP;299; PTX-NEXT: .reg .b64 %SPL;300; PTX-NEXT: .reg .b32 %r<3>;301; PTX-NEXT: .reg .b64 %rd<5>;302; PTX-EMPTY:303; PTX-NEXT: // %bb.0: // %entry304; PTX-NEXT: mov.b64 %SPL, __local_depot6;305; PTX-NEXT: cvta.local.u64 %SP, %SPL;306; PTX-NEXT: ld.param.b64 %rd1, [escape_ptrtoint_param_0];307; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;308; PTX-NEXT: add.u64 %rd3, %SP, 0;309; PTX-NEXT: add.u64 %rd4, %SPL, 0;310; PTX-NEXT: ld.param.b32 %r1, [escape_ptrtoint_param_1+4];311; PTX-NEXT: st.local.b32 [%rd4+4], %r1;312; PTX-NEXT: ld.param.b32 %r2, [escape_ptrtoint_param_1];313; PTX-NEXT: st.local.b32 [%rd4], %r2;314; PTX-NEXT: st.global.b64 [%rd2], %rd3;315; PTX-NEXT: ret;316entry:317 %i = ptrtoint ptr %s to i64318 store i64 %i, ptr %out, align 8319 ret void320}321 322; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)323define dso_local ptx_kernel void @memcpy_from_param(ptr nocapture noundef writeonly %out, ptr nocapture noundef readonly byval(%struct.S) align 4 %s) local_unnamed_addr #0 {324; LOWER-ARGS-LABEL: define dso_local ptx_kernel void @memcpy_from_param(325; LOWER-ARGS-SAME: ptr noundef writeonly captures(none) [[OUT:%.*]], ptr noundef readonly byval([[STRUCT_S:%.*]]) align 4 captures(none) [[S:%.*]]) local_unnamed_addr #[[ATTR0]] {326; LOWER-ARGS-NEXT: [[ENTRY:.*:]]327; LOWER-ARGS-NEXT: [[S3:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])328; LOWER-ARGS-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr [[OUT]], ptr addrspace(101) [[S3]], i64 16, i1 true)329; LOWER-ARGS-NEXT: ret void330;331; COPY-LABEL: define dso_local ptx_kernel void @memcpy_from_param(332; COPY-SAME: ptr noundef writeonly captures(none) [[OUT:%.*]], ptr noundef readonly byval([[STRUCT_S:%.*]]) align 4 captures(none) [[S:%.*]]) local_unnamed_addr #[[ATTR0]] {333; COPY-NEXT: [[ENTRY:.*:]]334; COPY-NEXT: tail call void @llvm.memcpy.p0.p0.i64(ptr [[OUT]], ptr [[S]], i64 16, i1 true)335; COPY-NEXT: ret void336;337; PTX-LABEL: memcpy_from_param(338; PTX: {339; PTX-NEXT: .reg .b16 %rs<17>;340; PTX-NEXT: .reg .b64 %rd<2>;341; PTX-EMPTY:342; PTX-NEXT: // %bb.0: // %entry343; PTX-NEXT: ld.param.b64 %rd1, [memcpy_from_param_param_0];344; PTX-NEXT: ld.param.b8 %rs1, [memcpy_from_param_param_1+15];345; PTX-NEXT: st.volatile.b8 [%rd1+15], %rs1;346; PTX-NEXT: ld.param.b8 %rs2, [memcpy_from_param_param_1+14];347; PTX-NEXT: st.volatile.b8 [%rd1+14], %rs2;348; PTX-NEXT: ld.param.b8 %rs3, [memcpy_from_param_param_1+13];349; PTX-NEXT: st.volatile.b8 [%rd1+13], %rs3;350; PTX-NEXT: ld.param.b8 %rs4, [memcpy_from_param_param_1+12];351; PTX-NEXT: st.volatile.b8 [%rd1+12], %rs4;352; PTX-NEXT: ld.param.b8 %rs5, [memcpy_from_param_param_1+11];353; PTX-NEXT: st.volatile.b8 [%rd1+11], %rs5;354; PTX-NEXT: ld.param.b8 %rs6, [memcpy_from_param_param_1+10];355; PTX-NEXT: st.volatile.b8 [%rd1+10], %rs6;356; PTX-NEXT: ld.param.b8 %rs7, [memcpy_from_param_param_1+9];357; PTX-NEXT: st.volatile.b8 [%rd1+9], %rs7;358; PTX-NEXT: ld.param.b8 %rs8, [memcpy_from_param_param_1+8];359; PTX-NEXT: st.volatile.b8 [%rd1+8], %rs8;360; PTX-NEXT: ld.param.b8 %rs9, [memcpy_from_param_param_1+7];361; PTX-NEXT: st.volatile.b8 [%rd1+7], %rs9;362; PTX-NEXT: ld.param.b8 %rs10, [memcpy_from_param_param_1+6];363; PTX-NEXT: st.volatile.b8 [%rd1+6], %rs10;364; PTX-NEXT: ld.param.b8 %rs11, [memcpy_from_param_param_1+5];365; PTX-NEXT: st.volatile.b8 [%rd1+5], %rs11;366; PTX-NEXT: ld.param.b8 %rs12, [memcpy_from_param_param_1+4];367; PTX-NEXT: st.volatile.b8 [%rd1+4], %rs12;368; PTX-NEXT: ld.param.b8 %rs13, [memcpy_from_param_param_1+3];369; PTX-NEXT: st.volatile.b8 [%rd1+3], %rs13;370; PTX-NEXT: ld.param.b8 %rs14, [memcpy_from_param_param_1+2];371; PTX-NEXT: st.volatile.b8 [%rd1+2], %rs14;372; PTX-NEXT: ld.param.b8 %rs15, [memcpy_from_param_param_1+1];373; PTX-NEXT: st.volatile.b8 [%rd1+1], %rs15;374; PTX-NEXT: ld.param.b8 %rs16, [memcpy_from_param_param_1];375; PTX-NEXT: st.volatile.b8 [%rd1], %rs16;376; PTX-NEXT: ret;377entry:378 tail call void @llvm.memcpy.p0.p0.i64(ptr %out, ptr %s, i64 16, i1 true)379 ret void380}381 382; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)383define dso_local ptx_kernel void @memcpy_from_param_noalign (ptr nocapture noundef writeonly %out, ptr nocapture noundef readonly byval(%struct.S) %s) local_unnamed_addr #0 {384; LOWER-ARGS-LABEL: define dso_local ptx_kernel void @memcpy_from_param_noalign(385; LOWER-ARGS-SAME: ptr noundef writeonly captures(none) [[OUT:%.*]], ptr noundef readonly byval([[STRUCT_S:%.*]]) align 4 captures(none) [[S:%.*]]) local_unnamed_addr #[[ATTR0]] {386; LOWER-ARGS-NEXT: [[ENTRY:.*:]]387; LOWER-ARGS-NEXT: [[S3:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])388; LOWER-ARGS-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr [[OUT]], ptr addrspace(101) [[S3]], i64 16, i1 true)389; LOWER-ARGS-NEXT: ret void390;391; COPY-LABEL: define dso_local ptx_kernel void @memcpy_from_param_noalign(392; COPY-SAME: ptr noundef writeonly captures(none) [[OUT:%.*]], ptr noundef readonly byval([[STRUCT_S:%.*]]) captures(none) [[S:%.*]]) local_unnamed_addr #[[ATTR0]] {393; COPY-NEXT: [[ENTRY:.*:]]394; COPY-NEXT: tail call void @llvm.memcpy.p0.p0.i64(ptr [[OUT]], ptr [[S]], i64 16, i1 true)395; COPY-NEXT: ret void396;397; PTX-LABEL: memcpy_from_param_noalign(398; PTX: {399; PTX-NEXT: .reg .b16 %rs<17>;400; PTX-NEXT: .reg .b64 %rd<2>;401; PTX-EMPTY:402; PTX-NEXT: // %bb.0: // %entry403; PTX-NEXT: ld.param.b64 %rd1, [memcpy_from_param_noalign_param_0];404; PTX-NEXT: ld.param.b8 %rs1, [memcpy_from_param_noalign_param_1+15];405; PTX-NEXT: st.volatile.b8 [%rd1+15], %rs1;406; PTX-NEXT: ld.param.b8 %rs2, [memcpy_from_param_noalign_param_1+14];407; PTX-NEXT: st.volatile.b8 [%rd1+14], %rs2;408; PTX-NEXT: ld.param.b8 %rs3, [memcpy_from_param_noalign_param_1+13];409; PTX-NEXT: st.volatile.b8 [%rd1+13], %rs3;410; PTX-NEXT: ld.param.b8 %rs4, [memcpy_from_param_noalign_param_1+12];411; PTX-NEXT: st.volatile.b8 [%rd1+12], %rs4;412; PTX-NEXT: ld.param.b8 %rs5, [memcpy_from_param_noalign_param_1+11];413; PTX-NEXT: st.volatile.b8 [%rd1+11], %rs5;414; PTX-NEXT: ld.param.b8 %rs6, [memcpy_from_param_noalign_param_1+10];415; PTX-NEXT: st.volatile.b8 [%rd1+10], %rs6;416; PTX-NEXT: ld.param.b8 %rs7, [memcpy_from_param_noalign_param_1+9];417; PTX-NEXT: st.volatile.b8 [%rd1+9], %rs7;418; PTX-NEXT: ld.param.b8 %rs8, [memcpy_from_param_noalign_param_1+8];419; PTX-NEXT: st.volatile.b8 [%rd1+8], %rs8;420; PTX-NEXT: ld.param.b8 %rs9, [memcpy_from_param_noalign_param_1+7];421; PTX-NEXT: st.volatile.b8 [%rd1+7], %rs9;422; PTX-NEXT: ld.param.b8 %rs10, [memcpy_from_param_noalign_param_1+6];423; PTX-NEXT: st.volatile.b8 [%rd1+6], %rs10;424; PTX-NEXT: ld.param.b8 %rs11, [memcpy_from_param_noalign_param_1+5];425; PTX-NEXT: st.volatile.b8 [%rd1+5], %rs11;426; PTX-NEXT: ld.param.b8 %rs12, [memcpy_from_param_noalign_param_1+4];427; PTX-NEXT: st.volatile.b8 [%rd1+4], %rs12;428; PTX-NEXT: ld.param.b8 %rs13, [memcpy_from_param_noalign_param_1+3];429; PTX-NEXT: st.volatile.b8 [%rd1+3], %rs13;430; PTX-NEXT: ld.param.b8 %rs14, [memcpy_from_param_noalign_param_1+2];431; PTX-NEXT: st.volatile.b8 [%rd1+2], %rs14;432; PTX-NEXT: ld.param.b8 %rs15, [memcpy_from_param_noalign_param_1+1];433; PTX-NEXT: st.volatile.b8 [%rd1+1], %rs15;434; PTX-NEXT: ld.param.b8 %rs16, [memcpy_from_param_noalign_param_1];435; PTX-NEXT: st.volatile.b8 [%rd1], %rs16;436; PTX-NEXT: ret;437entry:438 tail call void @llvm.memcpy.p0.p0.i64(ptr %out, ptr %s, i64 16, i1 true)439 ret void440}441 442; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)443define dso_local ptx_kernel void @memcpy_to_param(ptr nocapture noundef readonly %in, ptr nocapture noundef byval(%struct.S) align 4 %s) local_unnamed_addr #0 {444; COMMON-LABEL: define dso_local ptx_kernel void @memcpy_to_param(445; COMMON-SAME: ptr noundef readonly captures(none) [[IN:%.*]], ptr noundef byval([[STRUCT_S:%.*]]) align 4 captures(none) [[S:%.*]]) local_unnamed_addr #[[ATTR0]] {446; COMMON-NEXT: [[ENTRY:.*:]]447; COMMON-NEXT: [[S1:%.*]] = alloca [[STRUCT_S]], align 4448; COMMON-NEXT: [[S2:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])449; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[S1]], ptr addrspace(101) align 4 [[S2]], i64 8, i1 false)450; COMMON-NEXT: tail call void @llvm.memcpy.p0.p0.i64(ptr [[S1]], ptr [[IN]], i64 16, i1 true)451; COMMON-NEXT: ret void452;453; PTX-LABEL: memcpy_to_param(454; PTX: {455; PTX-NEXT: .local .align 8 .b8 __local_depot9[8];456; PTX-NEXT: .reg .b64 %SP;457; PTX-NEXT: .reg .b64 %SPL;458; PTX-NEXT: .reg .b32 %r<3>;459; PTX-NEXT: .reg .b64 %rd<47>;460; PTX-EMPTY:461; PTX-NEXT: // %bb.0: // %entry462; PTX-NEXT: mov.b64 %SPL, __local_depot9;463; PTX-NEXT: cvta.local.u64 %SP, %SPL;464; PTX-NEXT: ld.param.b64 %rd1, [memcpy_to_param_param_0];465; PTX-NEXT: add.u64 %rd2, %SPL, 0;466; PTX-NEXT: ld.param.b32 %r1, [memcpy_to_param_param_1+4];467; PTX-NEXT: st.local.b32 [%rd2+4], %r1;468; PTX-NEXT: ld.param.b32 %r2, [memcpy_to_param_param_1];469; PTX-NEXT: st.local.b32 [%rd2], %r2;470; PTX-NEXT: ld.volatile.b8 %rd3, [%rd1];471; PTX-NEXT: ld.volatile.b8 %rd4, [%rd1+1];472; PTX-NEXT: shl.b64 %rd5, %rd4, 8;473; PTX-NEXT: or.b64 %rd6, %rd5, %rd3;474; PTX-NEXT: ld.volatile.b8 %rd7, [%rd1+2];475; PTX-NEXT: shl.b64 %rd8, %rd7, 16;476; PTX-NEXT: ld.volatile.b8 %rd9, [%rd1+3];477; PTX-NEXT: shl.b64 %rd10, %rd9, 24;478; PTX-NEXT: or.b64 %rd11, %rd10, %rd8;479; PTX-NEXT: or.b64 %rd12, %rd11, %rd6;480; PTX-NEXT: ld.volatile.b8 %rd13, [%rd1+4];481; PTX-NEXT: ld.volatile.b8 %rd14, [%rd1+5];482; PTX-NEXT: shl.b64 %rd15, %rd14, 8;483; PTX-NEXT: or.b64 %rd16, %rd15, %rd13;484; PTX-NEXT: ld.volatile.b8 %rd17, [%rd1+6];485; PTX-NEXT: shl.b64 %rd18, %rd17, 16;486; PTX-NEXT: ld.volatile.b8 %rd19, [%rd1+7];487; PTX-NEXT: shl.b64 %rd20, %rd19, 24;488; PTX-NEXT: or.b64 %rd21, %rd20, %rd18;489; PTX-NEXT: or.b64 %rd22, %rd21, %rd16;490; PTX-NEXT: shl.b64 %rd23, %rd22, 32;491; PTX-NEXT: or.b64 %rd24, %rd23, %rd12;492; PTX-NEXT: st.volatile.b64 [%SP], %rd24;493; PTX-NEXT: ld.volatile.b8 %rd25, [%rd1+8];494; PTX-NEXT: ld.volatile.b8 %rd26, [%rd1+9];495; PTX-NEXT: shl.b64 %rd27, %rd26, 8;496; PTX-NEXT: or.b64 %rd28, %rd27, %rd25;497; PTX-NEXT: ld.volatile.b8 %rd29, [%rd1+10];498; PTX-NEXT: shl.b64 %rd30, %rd29, 16;499; PTX-NEXT: ld.volatile.b8 %rd31, [%rd1+11];500; PTX-NEXT: shl.b64 %rd32, %rd31, 24;501; PTX-NEXT: or.b64 %rd33, %rd32, %rd30;502; PTX-NEXT: or.b64 %rd34, %rd33, %rd28;503; PTX-NEXT: ld.volatile.b8 %rd35, [%rd1+12];504; PTX-NEXT: ld.volatile.b8 %rd36, [%rd1+13];505; PTX-NEXT: shl.b64 %rd37, %rd36, 8;506; PTX-NEXT: or.b64 %rd38, %rd37, %rd35;507; PTX-NEXT: ld.volatile.b8 %rd39, [%rd1+14];508; PTX-NEXT: shl.b64 %rd40, %rd39, 16;509; PTX-NEXT: ld.volatile.b8 %rd41, [%rd1+15];510; PTX-NEXT: shl.b64 %rd42, %rd41, 24;511; PTX-NEXT: or.b64 %rd43, %rd42, %rd40;512; PTX-NEXT: or.b64 %rd44, %rd43, %rd38;513; PTX-NEXT: shl.b64 %rd45, %rd44, 32;514; PTX-NEXT: or.b64 %rd46, %rd45, %rd34;515; PTX-NEXT: st.volatile.b64 [%SP+8], %rd46;516; PTX-NEXT: ret;517entry:518 tail call void @llvm.memcpy.p0.p0.i64(ptr %s, ptr %in, i64 16, i1 true)519 ret void520}521 522; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)523define dso_local ptx_kernel void @copy_on_store(ptr nocapture noundef readonly %in, ptr nocapture noundef byval(%struct.S) align 4 %s, i1 noundef zeroext %b) local_unnamed_addr #0 {524; COMMON-LABEL: define dso_local ptx_kernel void @copy_on_store(525; COMMON-SAME: ptr noundef readonly captures(none) [[IN:%.*]], ptr noundef byval([[STRUCT_S:%.*]]) align 4 captures(none) [[S:%.*]], i1 noundef zeroext [[B:%.*]]) local_unnamed_addr #[[ATTR0]] {526; COMMON-NEXT: [[BB:.*:]]527; COMMON-NEXT: [[S1:%.*]] = alloca [[STRUCT_S]], align 4528; COMMON-NEXT: [[S2:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])529; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[S1]], ptr addrspace(101) align 4 [[S2]], i64 8, i1 false)530; COMMON-NEXT: [[I:%.*]] = load i32, ptr [[IN]], align 4531; COMMON-NEXT: store i32 [[I]], ptr [[S1]], align 4532; COMMON-NEXT: ret void533;534; PTX-LABEL: copy_on_store(535; PTX: {536; PTX-EMPTY:537; PTX-EMPTY:538; PTX-NEXT: // %bb.0: // %bb539; PTX-NEXT: ret;540bb:541 %i = load i32, ptr %in, align 4542 store i32 %i, ptr %s, align 4543 ret void544}545 546define ptx_kernel void @test_select(ptr byval(i32) align 4 %input1, ptr byval(i32) %input2, ptr %out, i1 %cond) {547; SM_60-LABEL: define ptx_kernel void @test_select(548; SM_60-SAME: ptr byval(i32) align 4 [[INPUT1:%.*]], ptr byval(i32) [[INPUT2:%.*]], ptr [[OUT:%.*]], i1 [[COND:%.*]]) #[[ATTR3:[0-9]+]] {549; SM_60-NEXT: [[BB:.*:]]550; SM_60-NEXT: [[INPUT24:%.*]] = alloca i32, align 4551; SM_60-NEXT: [[INPUT25:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT2]])552; SM_60-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[INPUT24]], ptr addrspace(101) align 4 [[INPUT25]], i64 4, i1 false)553; SM_60-NEXT: [[INPUT11:%.*]] = alloca i32, align 4554; SM_60-NEXT: [[INPUT12:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])555; SM_60-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[INPUT11]], ptr addrspace(101) align 4 [[INPUT12]], i64 4, i1 false)556; SM_60-NEXT: [[PTRNEW:%.*]] = select i1 [[COND]], ptr [[INPUT11]], ptr [[INPUT24]]557; SM_60-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4558; SM_60-NEXT: store i32 [[VALLOADED]], ptr [[OUT]], align 4559; SM_60-NEXT: ret void560;561; SM_70-LABEL: define ptx_kernel void @test_select(562; SM_70-SAME: ptr byval(i32) align 4 [[INPUT1:%.*]], ptr byval(i32) [[INPUT2:%.*]], ptr [[OUT:%.*]], i1 [[COND:%.*]]) #[[ATTR3:[0-9]+]] {563; SM_70-NEXT: [[BB:.*:]]564; SM_70-NEXT: [[TMP0:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT2]])565; SM_70-NEXT: [[INPUT2_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP0]] to ptr566; SM_70-NEXT: [[TMP1:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])567; SM_70-NEXT: [[INPUT1_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP1]] to ptr568; SM_70-NEXT: [[PTRNEW:%.*]] = select i1 [[COND]], ptr [[INPUT1_PARAM_GEN]], ptr [[INPUT2_PARAM_GEN]]569; SM_70-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4570; SM_70-NEXT: store i32 [[VALLOADED]], ptr [[OUT]], align 4571; SM_70-NEXT: ret void572;573; COPY-LABEL: define ptx_kernel void @test_select(574; COPY-SAME: ptr byval(i32) align 4 [[INPUT1:%.*]], ptr byval(i32) [[INPUT2:%.*]], ptr [[OUT:%.*]], i1 [[COND:%.*]]) #[[ATTR3:[0-9]+]] {575; COPY-NEXT: [[BB:.*:]]576; COPY-NEXT: [[INPUT23:%.*]] = alloca i32, align 4577; COPY-NEXT: [[INPUT24:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT2]])578; COPY-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[INPUT23]], ptr addrspace(101) align 4 [[INPUT24]], i64 4, i1 false)579; COPY-NEXT: [[INPUT11:%.*]] = alloca i32, align 4580; COPY-NEXT: [[INPUT12:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])581; COPY-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[INPUT11]], ptr addrspace(101) align 4 [[INPUT12]], i64 4, i1 false)582; COPY-NEXT: [[PTRNEW:%.*]] = select i1 [[COND]], ptr [[INPUT11]], ptr [[INPUT23]]583; COPY-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4584; COPY-NEXT: store i32 [[VALLOADED]], ptr [[OUT]], align 4585; COPY-NEXT: ret void586;587; PTX-LABEL: test_select(588; PTX: {589; PTX-NEXT: .reg .pred %p<2>;590; PTX-NEXT: .reg .b16 %rs<3>;591; PTX-NEXT: .reg .b32 %r<2>;592; PTX-NEXT: .reg .b64 %rd<6>;593; PTX-EMPTY:594; PTX-NEXT: // %bb.0: // %bb595; PTX-NEXT: ld.param.b8 %rs1, [test_select_param_3];596; PTX-NEXT: and.b16 %rs2, %rs1, 1;597; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;598; PTX-NEXT: mov.b64 %rd1, test_select_param_0;599; PTX-NEXT: ld.param.b64 %rd2, [test_select_param_2];600; PTX-NEXT: cvta.to.global.u64 %rd3, %rd2;601; PTX-NEXT: mov.b64 %rd4, test_select_param_1;602; PTX-NEXT: selp.b64 %rd5, %rd1, %rd4, %p1;603; PTX-NEXT: ld.param.b32 %r1, [%rd5];604; PTX-NEXT: st.global.b32 [%rd3], %r1;605; PTX-NEXT: ret;606bb:607 %ptrnew = select i1 %cond, ptr %input1, ptr %input2608 %valloaded = load i32, ptr %ptrnew, align 4609 store i32 %valloaded, ptr %out, align 4610 ret void611}612 613define ptx_kernel void @test_select_write(ptr byval(i32) align 4 %input1, ptr byval(i32) %input2, ptr %out, i1 %cond) {614; COMMON-LABEL: define ptx_kernel void @test_select_write(615; COMMON-SAME: ptr byval(i32) align 4 [[INPUT1:%.*]], ptr byval(i32) [[INPUT2:%.*]], ptr [[OUT:%.*]], i1 [[COND:%.*]]) #[[ATTR3:[0-9]+]] {616; COMMON-NEXT: [[BB:.*:]]617; COMMON-NEXT: [[INPUT23:%.*]] = alloca i32, align 4618; COMMON-NEXT: [[INPUT24:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT2]])619; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[INPUT23]], ptr addrspace(101) align 4 [[INPUT24]], i64 4, i1 false)620; COMMON-NEXT: [[INPUT11:%.*]] = alloca i32, align 4621; COMMON-NEXT: [[INPUT12:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])622; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[INPUT11]], ptr addrspace(101) align 4 [[INPUT12]], i64 4, i1 false)623; COMMON-NEXT: [[PTRNEW:%.*]] = select i1 [[COND]], ptr [[INPUT11]], ptr [[INPUT23]]624; COMMON-NEXT: store i32 1, ptr [[PTRNEW]], align 4625; COMMON-NEXT: ret void626;627; PTX-LABEL: test_select_write(628; PTX: {629; PTX-NEXT: .local .align 4 .b8 __local_depot12[8];630; PTX-NEXT: .reg .b64 %SP;631; PTX-NEXT: .reg .b64 %SPL;632; PTX-NEXT: .reg .pred %p<2>;633; PTX-NEXT: .reg .b16 %rs<3>;634; PTX-NEXT: .reg .b32 %r<3>;635; PTX-NEXT: .reg .b64 %rd<4>;636; PTX-EMPTY:637; PTX-NEXT: // %bb.0: // %bb638; PTX-NEXT: mov.b64 %SPL, __local_depot12;639; PTX-NEXT: cvta.local.u64 %SP, %SPL;640; PTX-NEXT: ld.param.b8 %rs1, [test_select_write_param_3];641; PTX-NEXT: and.b16 %rs2, %rs1, 1;642; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;643; PTX-NEXT: ld.param.b32 %r1, [test_select_write_param_1];644; PTX-NEXT: st.b32 [%SP], %r1;645; PTX-NEXT: ld.param.b32 %r2, [test_select_write_param_0];646; PTX-NEXT: st.b32 [%SP+4], %r2;647; PTX-NEXT: add.u64 %rd1, %SPL, 4;648; PTX-NEXT: add.u64 %rd2, %SPL, 0;649; PTX-NEXT: selp.b64 %rd3, %rd1, %rd2, %p1;650; PTX-NEXT: st.local.b32 [%rd3], 1;651; PTX-NEXT: ret;652bb:653 %ptrnew = select i1 %cond, ptr %input1, ptr %input2654 store i32 1, ptr %ptrnew, align 4655 ret void656}657 658define ptx_kernel void @test_phi(ptr byval(%struct.S) align 4 %input1, ptr byval(%struct.S) %input2, ptr %inout, i1 %cond) {659; SM_60-LABEL: define ptx_kernel void @test_phi(660; SM_60-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 [[INPUT1:%.*]], ptr byval([[STRUCT_S]]) [[INPUT2:%.*]], ptr [[INOUT:%.*]], i1 [[COND:%.*]]) #[[ATTR3]] {661; SM_60-NEXT: [[BB:.*:]]662; SM_60-NEXT: [[INPUT24:%.*]] = alloca [[STRUCT_S]], align 8663; SM_60-NEXT: [[INPUT25:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT2]])664; SM_60-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 8 [[INPUT24]], ptr addrspace(101) align 8 [[INPUT25]], i64 8, i1 false)665; SM_60-NEXT: [[INPUT11:%.*]] = alloca [[STRUCT_S]], align 4666; SM_60-NEXT: [[INPUT12:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])667; SM_60-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[INPUT11]], ptr addrspace(101) align 4 [[INPUT12]], i64 8, i1 false)668; SM_60-NEXT: br i1 [[COND]], label %[[FIRST:.*]], label %[[SECOND:.*]]669; SM_60: [[FIRST]]:670; SM_60-NEXT: [[PTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT11]], i32 0, i32 0671; SM_60-NEXT: br label %[[MERGE:.*]]672; SM_60: [[SECOND]]:673; SM_60-NEXT: [[PTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT24]], i32 0, i32 1674; SM_60-NEXT: br label %[[MERGE]]675; SM_60: [[MERGE]]:676; SM_60-NEXT: [[PTRNEW:%.*]] = phi ptr [ [[PTR1]], %[[FIRST]] ], [ [[PTR2]], %[[SECOND]] ]677; SM_60-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4678; SM_60-NEXT: store i32 [[VALLOADED]], ptr [[INOUT]], align 4679; SM_60-NEXT: ret void680;681; SM_70-LABEL: define ptx_kernel void @test_phi(682; SM_70-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 [[INPUT1:%.*]], ptr byval([[STRUCT_S]]) [[INPUT2:%.*]], ptr [[INOUT:%.*]], i1 [[COND:%.*]]) #[[ATTR3]] {683; SM_70-NEXT: [[BB:.*:]]684; SM_70-NEXT: [[TMP0:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT2]])685; SM_70-NEXT: [[INPUT2_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP0]] to ptr686; SM_70-NEXT: [[TMP1:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])687; SM_70-NEXT: [[INPUT1_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP1]] to ptr688; SM_70-NEXT: br i1 [[COND]], label %[[FIRST:.*]], label %[[SECOND:.*]]689; SM_70: [[FIRST]]:690; SM_70-NEXT: [[PTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1_PARAM_GEN]], i32 0, i32 0691; SM_70-NEXT: br label %[[MERGE:.*]]692; SM_70: [[SECOND]]:693; SM_70-NEXT: [[PTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT2_PARAM_GEN]], i32 0, i32 1694; SM_70-NEXT: br label %[[MERGE]]695; SM_70: [[MERGE]]:696; SM_70-NEXT: [[PTRNEW:%.*]] = phi ptr [ [[PTR1]], %[[FIRST]] ], [ [[PTR2]], %[[SECOND]] ]697; SM_70-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4698; SM_70-NEXT: store i32 [[VALLOADED]], ptr [[INOUT]], align 4699; SM_70-NEXT: ret void700;701; COPY-LABEL: define ptx_kernel void @test_phi(702; COPY-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 [[INPUT1:%.*]], ptr byval([[STRUCT_S]]) [[INPUT2:%.*]], ptr [[INOUT:%.*]], i1 [[COND:%.*]]) #[[ATTR3]] {703; COPY-NEXT: [[BB:.*:]]704; COPY-NEXT: [[INPUT23:%.*]] = alloca [[STRUCT_S]], align 8705; COPY-NEXT: [[INPUT24:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT2]])706; COPY-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 8 [[INPUT23]], ptr addrspace(101) align 8 [[INPUT24]], i64 8, i1 false)707; COPY-NEXT: [[INPUT11:%.*]] = alloca [[STRUCT_S]], align 4708; COPY-NEXT: [[INPUT12:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])709; COPY-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[INPUT11]], ptr addrspace(101) align 4 [[INPUT12]], i64 8, i1 false)710; COPY-NEXT: br i1 [[COND]], label %[[FIRST:.*]], label %[[SECOND:.*]]711; COPY: [[FIRST]]:712; COPY-NEXT: [[PTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT11]], i32 0, i32 0713; COPY-NEXT: br label %[[MERGE:.*]]714; COPY: [[SECOND]]:715; COPY-NEXT: [[PTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT23]], i32 0, i32 1716; COPY-NEXT: br label %[[MERGE]]717; COPY: [[MERGE]]:718; COPY-NEXT: [[PTRNEW:%.*]] = phi ptr [ [[PTR1]], %[[FIRST]] ], [ [[PTR2]], %[[SECOND]] ]719; COPY-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4720; COPY-NEXT: store i32 [[VALLOADED]], ptr [[INOUT]], align 4721; COPY-NEXT: ret void722;723; PTX_60-LABEL: test_phi(724; PTX_60: {725; PTX_60-NEXT: .reg .pred %p<2>;726; PTX_60-NEXT: .reg .b16 %rs<3>;727; PTX_60-NEXT: .reg .b32 %r<2>;728; PTX_60-NEXT: .reg .b64 %rd<3>;729; PTX_60-EMPTY:730; PTX_60-NEXT: // %bb.0: // %bb731; PTX_60-NEXT: ld.param.b8 %rs1, [test_phi_param_3];732; PTX_60-NEXT: and.b16 %rs2, %rs1, 1;733; PTX_60-NEXT: setp.ne.b16 %p1, %rs2, 0;734; PTX_60-NEXT: ld.param.b64 %rd2, [test_phi_param_2];735; PTX_60-NEXT: cvta.to.global.u64 %rd1, %rd2;736; PTX_60-NEXT: ld.param.b32 %r1, [test_phi_param_0];737; PTX_60-NEXT: @%p1 bra $L__BB13_2;738; PTX_60-NEXT: // %bb.1: // %second739; PTX_60-NEXT: ld.param.b32 %r1, [test_phi_param_1+4];740; PTX_60-NEXT: $L__BB13_2: // %merge741; PTX_60-NEXT: st.global.b32 [%rd1], %r1;742; PTX_60-NEXT: ret;743;744; PTX_70-LABEL: test_phi(745; PTX_70: {746; PTX_70-NEXT: .reg .pred %p<2>;747; PTX_70-NEXT: .reg .b16 %rs<3>;748; PTX_70-NEXT: .reg .b32 %r<2>;749; PTX_70-NEXT: .reg .b64 %rd<5>;750; PTX_70-EMPTY:751; PTX_70-NEXT: // %bb.0: // %bb752; PTX_70-NEXT: ld.param.b8 %rs1, [test_phi_param_3];753; PTX_70-NEXT: and.b16 %rs2, %rs1, 1;754; PTX_70-NEXT: setp.ne.b16 %p1, %rs2, 0;755; PTX_70-NEXT: mov.b64 %rd4, test_phi_param_0;756; PTX_70-NEXT: ld.param.b64 %rd3, [test_phi_param_2];757; PTX_70-NEXT: cvta.to.global.u64 %rd1, %rd3;758; PTX_70-NEXT: @%p1 bra $L__BB13_2;759; PTX_70-NEXT: // %bb.1: // %second760; PTX_70-NEXT: mov.b64 %rd2, test_phi_param_1;761; PTX_70-NEXT: add.s64 %rd4, %rd2, 4;762; PTX_70-NEXT: $L__BB13_2: // %merge763; PTX_70-NEXT: ld.param.b32 %r1, [%rd4];764; PTX_70-NEXT: st.global.b32 [%rd1], %r1;765; PTX_70-NEXT: ret;766bb:767 br i1 %cond, label %first, label %second768 769first: ; preds = %bb770 %ptr1 = getelementptr inbounds %struct.S, ptr %input1, i32 0, i32 0771 br label %merge772 773second: ; preds = %bb774 %ptr2 = getelementptr inbounds %struct.S, ptr %input2, i32 0, i32 1775 br label %merge776 777merge: ; preds = %second, %first778 %ptrnew = phi ptr [ %ptr1, %first ], [ %ptr2, %second ]779 %valloaded = load i32, ptr %ptrnew, align 4780 store i32 %valloaded, ptr %inout, align 4781 ret void782}783 784define ptx_kernel void @test_phi_write(ptr byval(%struct.S) align 4 %input1, ptr byval(%struct.S) %input2, i1 %cond) {785; COMMON-LABEL: define ptx_kernel void @test_phi_write(786; COMMON-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 [[INPUT1:%.*]], ptr byval([[STRUCT_S]]) [[INPUT2:%.*]], i1 [[COND:%.*]]) #[[ATTR3]] {787; COMMON-NEXT: [[BB:.*:]]788; COMMON-NEXT: [[INPUT24:%.*]] = alloca [[STRUCT_S]], align 8789; COMMON-NEXT: [[INPUT25:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT2]])790; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 8 [[INPUT24]], ptr addrspace(101) align 8 [[INPUT25]], i64 8, i1 false)791; COMMON-NEXT: [[INPUT11:%.*]] = alloca [[STRUCT_S]], align 4792; COMMON-NEXT: [[INPUT12:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])793; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[INPUT11]], ptr addrspace(101) align 4 [[INPUT12]], i64 8, i1 false)794; COMMON-NEXT: br i1 [[COND]], label %[[FIRST:.*]], label %[[SECOND:.*]]795; COMMON: [[FIRST]]:796; COMMON-NEXT: [[PTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT11]], i32 0, i32 0797; COMMON-NEXT: br label %[[MERGE:.*]]798; COMMON: [[SECOND]]:799; COMMON-NEXT: [[PTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT24]], i32 0, i32 1800; COMMON-NEXT: br label %[[MERGE]]801; COMMON: [[MERGE]]:802; COMMON-NEXT: [[PTRNEW:%.*]] = phi ptr [ [[PTR1]], %[[FIRST]] ], [ [[PTR2]], %[[SECOND]] ]803; COMMON-NEXT: store i32 1, ptr [[PTRNEW]], align 4804; COMMON-NEXT: ret void805;806; PTX-LABEL: test_phi_write(807; PTX: {808; PTX-NEXT: .local .align 4 .b8 __local_depot14[8];809; PTX-NEXT: .reg .b64 %SP;810; PTX-NEXT: .reg .b64 %SPL;811; PTX-NEXT: .reg .pred %p<2>;812; PTX-NEXT: .reg .b16 %rs<3>;813; PTX-NEXT: .reg .b32 %r<3>;814; PTX-NEXT: .reg .b64 %rd<3>;815; PTX-EMPTY:816; PTX-NEXT: // %bb.0: // %bb817; PTX-NEXT: mov.b64 %SPL, __local_depot14;818; PTX-NEXT: cvta.local.u64 %SP, %SPL;819; PTX-NEXT: ld.param.b8 %rs1, [test_phi_write_param_2];820; PTX-NEXT: and.b16 %rs2, %rs1, 1;821; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;822; PTX-NEXT: add.u64 %rd1, %SPL, 0;823; PTX-NEXT: ld.param.b32 %r1, [test_phi_write_param_1+4];824; PTX-NEXT: st.b32 [%SP], %r1;825; PTX-NEXT: add.u64 %rd2, %SPL, 4;826; PTX-NEXT: ld.param.b32 %r2, [test_phi_write_param_0];827; PTX-NEXT: st.b32 [%SP+4], %r2;828; PTX-NEXT: @%p1 bra $L__BB14_2;829; PTX-NEXT: // %bb.1: // %second830; PTX-NEXT: mov.b64 %rd2, %rd1;831; PTX-NEXT: $L__BB14_2: // %merge832; PTX-NEXT: st.local.b32 [%rd2], 1;833; PTX-NEXT: ret;834bb:835 br i1 %cond, label %first, label %second836 837first: ; preds = %bb838 %ptr1 = getelementptr inbounds %struct.S, ptr %input1, i32 0, i32 0839 br label %merge840 841second: ; preds = %bb842 %ptr2 = getelementptr inbounds %struct.S, ptr %input2, i32 0, i32 1843 br label %merge844 845merge: ; preds = %second, %first846 %ptrnew = phi ptr [ %ptr1, %first ], [ %ptr2, %second ]847 store i32 1, ptr %ptrnew, align 4848 ret void849}850 851define ptx_kernel void @test_forward_byval_arg(ptr byval(i32) align 4 %input) {852; COMMON-LABEL: define ptx_kernel void @test_forward_byval_arg(853; COMMON-SAME: ptr byval(i32) align 4 [[INPUT:%.*]]) #[[ATTR3]] {854; COMMON-NEXT: [[INPUT1:%.*]] = alloca i32, align 4855; COMMON-NEXT: [[INPUT2:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT]])856; COMMON-NEXT: call void @llvm.memcpy.p0.p101.i64(ptr align 4 [[INPUT1]], ptr addrspace(101) align 4 [[INPUT2]], i64 4, i1 false)857; COMMON-NEXT: call void @device_func(ptr byval(i32) align 4 [[INPUT1]])858; COMMON-NEXT: ret void859;860; PTX-LABEL: test_forward_byval_arg(861; PTX: {862; PTX-NEXT: .local .align 4 .b8 __local_depot15[4];863; PTX-NEXT: .reg .b64 %SP;864; PTX-NEXT: .reg .b64 %SPL;865; PTX-NEXT: .reg .b32 %r<2>;866; PTX-NEXT: .reg .b64 %rd<2>;867; PTX-EMPTY:868; PTX-NEXT: // %bb.0:869; PTX-NEXT: mov.b64 %SPL, __local_depot15;870; PTX-NEXT: add.u64 %rd1, %SPL, 0;871; PTX-NEXT: ld.param.b32 %r1, [test_forward_byval_arg_param_0];872; PTX-NEXT: st.local.b32 [%rd1], %r1;873; PTX-NEXT: { // callseq 2, 0874; PTX-NEXT: .param .align 4 .b8 param0[4];875; PTX-NEXT: st.param.b32 [param0], %r1;876; PTX-NEXT: call.uni device_func, (param0);877; PTX-NEXT: } // callseq 2878; PTX-NEXT: ret;879 call void @device_func(ptr byval(i32) align 4 %input)880 ret void881}882 883define void @device_func(ptr byval(i32) align 4 %input) {884; COMMON-LABEL: define void @device_func(885; COMMON-SAME: ptr byval(i32) align 4 [[INPUT:%.*]]) #[[ATTR3]] {886; COMMON-NEXT: call void @device_func(ptr byval(i32) align 4 [[INPUT]])887; COMMON-NEXT: ret void888;889; PTX-LABEL: device_func(890; PTX: {891; PTX-NEXT: .reg .b32 %r<2>;892; PTX-EMPTY:893; PTX-NEXT: // %bb.0:894; PTX-NEXT: { // callseq 3, 0895; PTX-NEXT: .param .align 4 .b8 param0[4];896; PTX-NEXT: ld.param.b32 %r1, [device_func_param_0];897; PTX-NEXT: st.param.b32 [param0], %r1;898; PTX-NEXT: call.uni device_func, (param0);899; PTX-NEXT: } // callseq 3900; PTX-NEXT: ret;901 call void @device_func(ptr byval(i32) align 4 %input)902 ret void903}904 905attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite) "no-trapping-math"="true" "target-cpu"="sm_60" "target-features"="+ptx78,+sm_60" "uniform-work-group-size"="true" }906attributes #1 = { nocallback nofree nounwind willreturn memory(argmem: readwrite) }907attributes #2 = { nocallback nofree nounwind willreturn memory(argmem: write) }908 909!llvm.module.flags = !{!0, !1, !2, !3}910!llvm.ident = !{!20, !21}911 912!0 = !{i32 2, !"SDK Version", [2 x i32] [i32 11, i32 8]}913!1 = !{i32 1, !"wchar_size", i32 4}914!2 = !{i32 4, !"nvvm-reflect-ftz", i32 0}915!3 = !{i32 7, !"frame-pointer", i32 2}916!20 = !{!"clang version 20.0.0git"}917!21 = !{!"clang version 3.8.0 (tags/RELEASE_380/final)"}918