548 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: opt < %s -S -nvptx-lower-args --mtriple nvptx64-nvidia-cuda -mcpu=sm_70 -mattr=+ptx77 | FileCheck %s --check-prefixes OPT3; RUN: llc < %s --mtriple nvptx64-nvidia-cuda -mcpu=sm_70 -mattr=+ptx77 | FileCheck %s --check-prefixes PTX4 5%struct.uint4 = type { i32, i32, i32, i32 }6 7@gi = dso_local addrspace(1) externally_initialized global %struct.uint4 { i32 50462976, i32 117835012, i32 185207048, i32 252579084 }, align 168 9; Function Attrs: mustprogress nofree noinline norecurse nosync nounwind willreturn memory(read, inaccessiblemem: none)10; Regular functions mus still make a copy. `cvta.param` does not always work there.11define dso_local noundef i32 @non_kernel_function(ptr nocapture noundef readonly byval(%struct.uint4) align 16 %a, i1 noundef zeroext %b, i32 noundef %c) local_unnamed_addr #0 {12; OPT-LABEL: define dso_local noundef i32 @non_kernel_function(13; OPT-SAME: ptr noundef readonly byval([[STRUCT_UINT4:%.*]]) align 16 captures(none) [[A:%.*]], i1 noundef zeroext [[B:%.*]], i32 noundef [[C:%.*]]) local_unnamed_addr #[[ATTR0:[0-9]+]] {14; OPT-NEXT: [[ENTRY:.*:]]15; OPT-NEXT: [[A_:%.*]] = select i1 [[B]], ptr [[A]], ptr addrspacecast (ptr addrspace(1) @gi to ptr)16; OPT-NEXT: [[IDX_EXT:%.*]] = sext i32 [[C]] to i6417; OPT-NEXT: [[ADD_PTR:%.*]] = getelementptr inbounds i8, ptr [[A_]], i64 [[IDX_EXT]]18; OPT-NEXT: [[TMP0:%.*]] = load i32, ptr [[ADD_PTR]], align 119; OPT-NEXT: ret i32 [[TMP0]]20;21; PTX-LABEL: non_kernel_function(22; PTX: {23; PTX-NEXT: .reg .pred %p<2>;24; PTX-NEXT: .reg .b16 %rs<3>;25; PTX-NEXT: .reg .b32 %r<11>;26; PTX-NEXT: .reg .b64 %rd<8>;27; PTX-EMPTY:28; PTX-NEXT: // %bb.0: // %entry29; PTX-NEXT: mov.b64 %rd1, non_kernel_function_param_0;30; PTX-NEXT: cvta.local.u64 %rd2, %rd1;31; PTX-NEXT: ld.param.b8 %rs1, [non_kernel_function_param_1];32; PTX-NEXT: and.b16 %rs2, %rs1, 1;33; PTX-NEXT: setp.ne.b16 %p1, %rs2, 0;34; PTX-NEXT: mov.b64 %rd3, gi;35; PTX-NEXT: cvta.global.u64 %rd4, %rd3;36; PTX-NEXT: selp.b64 %rd5, %rd2, %rd4, %p1;37; PTX-NEXT: ld.param.s32 %rd6, [non_kernel_function_param_2];38; PTX-NEXT: add.s64 %rd7, %rd5, %rd6;39; PTX-NEXT: ld.b8 %r1, [%rd7];40; PTX-NEXT: ld.b8 %r2, [%rd7+1];41; PTX-NEXT: shl.b32 %r3, %r2, 8;42; PTX-NEXT: or.b32 %r4, %r3, %r1;43; PTX-NEXT: ld.b8 %r5, [%rd7+2];44; PTX-NEXT: shl.b32 %r6, %r5, 16;45; PTX-NEXT: ld.b8 %r7, [%rd7+3];46; PTX-NEXT: shl.b32 %r8, %r7, 24;47; PTX-NEXT: or.b32 %r9, %r8, %r6;48; PTX-NEXT: or.b32 %r10, %r9, %r4;49; PTX-NEXT: st.param.b32 [func_retval0], %r10;50; PTX-NEXT: ret;51entry:52 %a. = select i1 %b, ptr %a, ptr addrspacecast (ptr addrspace(1) @gi to ptr)53 %idx.ext = sext i32 %c to i6454 %add.ptr = getelementptr inbounds i8, ptr %a., i64 %idx.ext55 %0 = load i32, ptr %add.ptr, align 156 ret i32 %057}58 59define ptx_kernel void @grid_const_int(ptr byval(i32) align 4 "nvvm.grid_constant" %input1, i32 %input2, ptr %out, i32 %n) {60; PTX-LABEL: grid_const_int(61; PTX: {62; PTX-NEXT: .reg .b32 %r<4>;63; PTX-NEXT: .reg .b64 %rd<3>;64; PTX-EMPTY:65; PTX-NEXT: // %bb.0:66; PTX-NEXT: ld.param.b64 %rd1, [grid_const_int_param_2];67; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;68; PTX-NEXT: ld.param.b32 %r1, [grid_const_int_param_1];69; PTX-NEXT: ld.param.b32 %r2, [grid_const_int_param_0];70; PTX-NEXT: add.s32 %r3, %r2, %r1;71; PTX-NEXT: st.global.b32 [%rd2], %r3;72; PTX-NEXT: ret;73; OPT-LABEL: define ptx_kernel void @grid_const_int(74; OPT-SAME: ptr byval(i32) align 4 "nvvm.grid_constant" [[INPUT1:%.*]], i32 [[INPUT2:%.*]], ptr [[OUT:%.*]], i32 [[N:%.*]]) #[[ATTR0]] {75; OPT-NEXT: [[INPUT11:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])76; OPT-NEXT: [[TMP:%.*]] = load i32, ptr addrspace(101) [[INPUT11]], align 477; OPT-NEXT: [[ADD:%.*]] = add i32 [[TMP]], [[INPUT2]]78; OPT-NEXT: store i32 [[ADD]], ptr [[OUT]], align 479; OPT-NEXT: ret void80 %tmp = load i32, ptr %input1, align 481 %add = add i32 %tmp, %input282 store i32 %add, ptr %out83 ret void84}85 86%struct.s = type { i32, i32 }87 88define ptx_kernel void @grid_const_struct(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input, ptr %out){89; PTX-LABEL: grid_const_struct(90; PTX: {91; PTX-NEXT: .reg .b32 %r<4>;92; PTX-NEXT: .reg .b64 %rd<3>;93; PTX-EMPTY:94; PTX-NEXT: // %bb.0:95; PTX-NEXT: ld.param.b64 %rd1, [grid_const_struct_param_1];96; PTX-NEXT: cvta.to.global.u64 %rd2, %rd1;97; PTX-NEXT: ld.param.b32 %r1, [grid_const_struct_param_0];98; PTX-NEXT: ld.param.b32 %r2, [grid_const_struct_param_0+4];99; PTX-NEXT: add.s32 %r3, %r1, %r2;100; PTX-NEXT: st.global.b32 [%rd2], %r3;101; PTX-NEXT: ret;102; OPT-LABEL: define ptx_kernel void @grid_const_struct(103; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT:%.*]], ptr [[OUT:%.*]]) #[[ATTR0]] {104; OPT-NEXT: [[INPUT1:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT]])105; OPT-NEXT: [[GEP13:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr addrspace(101) [[INPUT1]], i32 0, i32 0106; OPT-NEXT: [[GEP22:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr addrspace(101) [[INPUT1]], i32 0, i32 1107; OPT-NEXT: [[TMP1:%.*]] = load i32, ptr addrspace(101) [[GEP13]], align 4108; OPT-NEXT: [[TMP2:%.*]] = load i32, ptr addrspace(101) [[GEP22]], align 4109; OPT-NEXT: [[ADD:%.*]] = add i32 [[TMP1]], [[TMP2]]110; OPT-NEXT: store i32 [[ADD]], ptr [[OUT]], align 4111; OPT-NEXT: ret void112 %gep1 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 0113 %gep2 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 1114 %int1 = load i32, ptr %gep1115 %int2 = load i32, ptr %gep2116 %add = add i32 %int1, %int2117 store i32 %add, ptr %out118 ret void119}120 121define ptx_kernel void @grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input) {122; PTX-LABEL: grid_const_escape(123; PTX: {124; PTX-NEXT: .reg .b64 %rd<4>;125; PTX-EMPTY:126; PTX-NEXT: // %bb.0:127; PTX-NEXT: mov.b64 %rd1, grid_const_escape_param_0;128; PTX-NEXT: cvta.param.u64 %rd2, %rd1;129; PTX-NEXT: { // callseq 0, 0130; PTX-NEXT: .param .b64 param0;131; PTX-NEXT: .param .b32 retval0;132; PTX-NEXT: st.param.b64 [param0], %rd2;133; PTX-NEXT: prototype_0 : .callprototype (.param .b32 _) _ (.param .b64 _);134; PTX-NEXT: mov.b64 %rd3, escape;135; PTX-NEXT: call (retval0), %rd3, (param0), prototype_0;136; PTX-NEXT: } // callseq 0137; PTX-NEXT: ret;138; OPT-LABEL: define ptx_kernel void @grid_const_escape(139; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT:%.*]]) #[[ATTR0]] {140; OPT-NEXT: [[TMP1:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT]])141; OPT-NEXT: [[INPUT_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP1]] to ptr142; OPT-NEXT: [[CALL:%.*]] = call i32 @escape(ptr [[INPUT_PARAM_GEN]])143; OPT-NEXT: ret void144 %call = call i32 @escape(ptr %input)145 ret void146}147 148define ptx_kernel void @multiple_grid_const_escape(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input, i32 %a, ptr byval(i32) align 4 "nvvm.grid_constant" %b) {149; PTX-LABEL: multiple_grid_const_escape(150; PTX: {151; PTX-NEXT: .local .align 4 .b8 __local_depot4[4];152; PTX-NEXT: .reg .b64 %SP;153; PTX-NEXT: .reg .b64 %SPL;154; PTX-NEXT: .reg .b32 %r<2>;155; PTX-NEXT: .reg .b64 %rd<8>;156; PTX-EMPTY:157; PTX-NEXT: // %bb.0:158; PTX-NEXT: mov.b64 %SPL, __local_depot4;159; PTX-NEXT: cvta.local.u64 %SP, %SPL;160; PTX-NEXT: mov.b64 %rd1, multiple_grid_const_escape_param_0;161; PTX-NEXT: ld.param.b32 %r1, [multiple_grid_const_escape_param_1];162; PTX-NEXT: mov.b64 %rd2, multiple_grid_const_escape_param_2;163; PTX-NEXT: cvta.param.u64 %rd3, %rd2;164; PTX-NEXT: cvta.param.u64 %rd4, %rd1;165; PTX-NEXT: add.u64 %rd5, %SP, 0;166; PTX-NEXT: add.u64 %rd6, %SPL, 0;167; PTX-NEXT: st.local.b32 [%rd6], %r1;168; PTX-NEXT: { // callseq 1, 0169; PTX-NEXT: .param .b64 param0;170; PTX-NEXT: .param .b64 param1;171; PTX-NEXT: .param .b64 param2;172; PTX-NEXT: .param .b32 retval0;173; PTX-NEXT: st.param.b64 [param2], %rd3;174; PTX-NEXT: st.param.b64 [param1], %rd5;175; PTX-NEXT: st.param.b64 [param0], %rd4;176; PTX-NEXT: prototype_1 : .callprototype (.param .b32 _) _ (.param .b64 _, .param .b64 _, .param .b64 _);177; PTX-NEXT: mov.b64 %rd7, escape3;178; PTX-NEXT: call (retval0), %rd7, (param0, param1, param2), prototype_1;179; PTX-NEXT: } // callseq 1180; PTX-NEXT: ret;181; OPT-LABEL: define ptx_kernel void @multiple_grid_const_escape(182; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT:%.*]], i32 [[A:%.*]], ptr byval(i32) align 4 "nvvm.grid_constant" [[B:%.*]]) #[[ATTR0]] {183; OPT-NEXT: [[TMP1:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[B]])184; OPT-NEXT: [[B_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP1]] to ptr185; OPT-NEXT: [[TMP2:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT]])186; OPT-NEXT: [[INPUT_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP2]] to ptr187; OPT-NEXT: [[A_ADDR:%.*]] = alloca i32, align 4188; OPT-NEXT: store i32 [[A]], ptr [[A_ADDR]], align 4189; OPT-NEXT: [[CALL:%.*]] = call i32 @escape3(ptr [[INPUT_PARAM_GEN]], ptr [[A_ADDR]], ptr [[B_PARAM_GEN]])190; OPT-NEXT: ret void191 %a.addr = alloca i32, align 4192 store i32 %a, ptr %a.addr, align 4193 %call = call i32 @escape3(ptr %input, ptr %a.addr, ptr %b)194 ret void195}196 197define ptx_kernel void @grid_const_memory_escape(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input, ptr %addr) {198; PTX-LABEL: grid_const_memory_escape(199; PTX: {200; PTX-NEXT: .reg .b64 %rd<5>;201; PTX-EMPTY:202; PTX-NEXT: // %bb.0:203; PTX-NEXT: mov.b64 %rd1, grid_const_memory_escape_param_0;204; PTX-NEXT: ld.param.b64 %rd2, [grid_const_memory_escape_param_1];205; PTX-NEXT: cvta.to.global.u64 %rd3, %rd2;206; PTX-NEXT: cvta.param.u64 %rd4, %rd1;207; PTX-NEXT: st.global.b64 [%rd3], %rd4;208; PTX-NEXT: ret;209; OPT-LABEL: define ptx_kernel void @grid_const_memory_escape(210; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT:%.*]], ptr [[ADDR:%.*]]) #[[ATTR0]] {211; OPT-NEXT: [[TMP1:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT]])212; OPT-NEXT: [[INPUT1:%.*]] = addrspacecast ptr addrspace(101) [[TMP1]] to ptr213; OPT-NEXT: store ptr [[INPUT1]], ptr [[ADDR]], align 8214; OPT-NEXT: ret void215 store ptr %input, ptr %addr, align 8216 ret void217}218 219define ptx_kernel void @grid_const_inlineasm_escape(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input, ptr %result) {220; PTX-LABEL: grid_const_inlineasm_escape(221; PTX: {222; PTX-NEXT: .reg .b64 %rd<7>;223; PTX-EMPTY:224; PTX-NEXT: // %bb.0:225; PTX-NEXT: mov.b64 %rd4, grid_const_inlineasm_escape_param_0;226; PTX-NEXT: ld.param.b64 %rd5, [grid_const_inlineasm_escape_param_1];227; PTX-NEXT: cvta.to.global.u64 %rd6, %rd5;228; PTX-NEXT: cvta.param.u64 %rd2, %rd4;229; PTX-NEXT: add.s64 %rd3, %rd2, 4;230; PTX-NEXT: // begin inline asm231; PTX-NEXT: add.s64 %rd1, %rd2, %rd3;232; PTX-NEXT: // end inline asm233; PTX-NEXT: st.global.b64 [%rd6], %rd1;234; PTX-NEXT: ret;235; PTX-NOT .local236; OPT-LABEL: define ptx_kernel void @grid_const_inlineasm_escape(237; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT:%.*]], ptr [[RESULT:%.*]]) #[[ATTR0]] {238; OPT-NEXT: [[TMP1:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT]])239; OPT-NEXT: [[INPUT1:%.*]] = addrspacecast ptr addrspace(101) [[TMP1]] to ptr240; OPT-NEXT: [[TMPPTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1]], i32 0, i32 0241; OPT-NEXT: [[TMPPTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1]], i32 0, i32 1242; OPT-NEXT: [[TMP2:%.*]] = call i64 asm "add.s64 $0, $1, $2243; OPT-NEXT: store i64 [[TMP2]], ptr [[RESULT]], align 8244; OPT-NEXT: ret void245 %tmpptr1 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 0246 %tmpptr2 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 1247 %1 = call i64 asm "add.s64 $0, $1, $2;", "=l,l,l"(ptr %tmpptr1, ptr %tmpptr2) #1248 store i64 %1, ptr %result, align 8249 ret void250}251 252define ptx_kernel void @grid_const_partial_escape(ptr byval(i32) "nvvm.grid_constant" %input, ptr %output) {253; PTX-LABEL: grid_const_partial_escape(254; PTX: {255; PTX-NEXT: .reg .b32 %r<3>;256; PTX-NEXT: .reg .b64 %rd<6>;257; PTX-EMPTY:258; PTX-NEXT: // %bb.0:259; PTX-NEXT: mov.b64 %rd1, grid_const_partial_escape_param_0;260; PTX-NEXT: ld.param.b64 %rd2, [grid_const_partial_escape_param_1];261; PTX-NEXT: cvta.to.global.u64 %rd3, %rd2;262; PTX-NEXT: cvta.param.u64 %rd4, %rd1;263; PTX-NEXT: ld.param.b32 %r1, [grid_const_partial_escape_param_0];264; PTX-NEXT: add.s32 %r2, %r1, %r1;265; PTX-NEXT: st.global.b32 [%rd3], %r2;266; PTX-NEXT: { // callseq 2, 0267; PTX-NEXT: .param .b64 param0;268; PTX-NEXT: .param .b32 retval0;269; PTX-NEXT: st.param.b64 [param0], %rd4;270; PTX-NEXT: prototype_2 : .callprototype (.param .b32 _) _ (.param .b64 _);271; PTX-NEXT: mov.b64 %rd5, escape;272; PTX-NEXT: call (retval0), %rd5, (param0), prototype_2;273; PTX-NEXT: } // callseq 2274; PTX-NEXT: ret;275; OPT-LABEL: define ptx_kernel void @grid_const_partial_escape(276; OPT-SAME: ptr byval(i32) "nvvm.grid_constant" [[INPUT:%.*]], ptr [[OUTPUT:%.*]]) #[[ATTR0]] {277; OPT-NEXT: [[TMP1:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT]])278; OPT-NEXT: [[INPUT1_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP1]] to ptr279; OPT-NEXT: [[VAL1:%.*]] = load i32, ptr [[INPUT1_GEN]], align 4280; OPT-NEXT: [[TWICE:%.*]] = add i32 [[VAL1]], [[VAL1]]281; OPT-NEXT: store i32 [[TWICE]], ptr [[OUTPUT]], align 4282; OPT-NEXT: [[CALL:%.*]] = call i32 @escape(ptr [[INPUT1_GEN]])283; OPT-NEXT: ret void284 %val = load i32, ptr %input285 %twice = add i32 %val, %val286 store i32 %twice, ptr %output287 %call = call i32 @escape(ptr %input)288 ret void289}290 291define ptx_kernel i32 @grid_const_partial_escapemem(ptr byval(%struct.s) "nvvm.grid_constant" %input, ptr %output) {292; PTX-LABEL: grid_const_partial_escapemem(293; PTX: {294; PTX-NEXT: .reg .b32 %r<4>;295; PTX-NEXT: .reg .b64 %rd<6>;296; PTX-EMPTY:297; PTX-NEXT: // %bb.0:298; PTX-NEXT: mov.b64 %rd1, grid_const_partial_escapemem_param_0;299; PTX-NEXT: ld.param.b64 %rd2, [grid_const_partial_escapemem_param_1];300; PTX-NEXT: cvta.to.global.u64 %rd3, %rd2;301; PTX-NEXT: cvta.param.u64 %rd4, %rd1;302; PTX-NEXT: ld.param.b32 %r1, [grid_const_partial_escapemem_param_0];303; PTX-NEXT: ld.param.b32 %r2, [grid_const_partial_escapemem_param_0+4];304; PTX-NEXT: st.global.b64 [%rd3], %rd4;305; PTX-NEXT: add.s32 %r3, %r1, %r2;306; PTX-NEXT: { // callseq 3, 0307; PTX-NEXT: .param .b64 param0;308; PTX-NEXT: .param .b32 retval0;309; PTX-NEXT: st.param.b64 [param0], %rd4;310; PTX-NEXT: prototype_3 : .callprototype (.param .b32 _) _ (.param .b64 _);311; PTX-NEXT: mov.b64 %rd5, escape;312; PTX-NEXT: call (retval0), %rd5, (param0), prototype_3;313; PTX-NEXT: } // callseq 3314; PTX-NEXT: st.param.b32 [func_retval0], %r3;315; PTX-NEXT: ret;316; OPT-LABEL: define ptx_kernel i32 @grid_const_partial_escapemem(317; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) "nvvm.grid_constant" [[INPUT:%.*]], ptr [[OUTPUT:%.*]]) #[[ATTR0]] {318; OPT-NEXT: [[TMP1:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT]])319; OPT-NEXT: [[INPUT1:%.*]] = addrspacecast ptr addrspace(101) [[TMP1]] to ptr320; OPT-NEXT: [[PTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1]], i32 0, i32 0321; OPT-NEXT: [[VAL1:%.*]] = load i32, ptr [[PTR1]], align 4322; OPT-NEXT: [[PTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1]], i32 0, i32 1323; OPT-NEXT: [[VAL2:%.*]] = load i32, ptr [[PTR2]], align 4324; OPT-NEXT: store ptr [[INPUT1]], ptr [[OUTPUT]], align 8325; OPT-NEXT: [[ADD:%.*]] = add i32 [[VAL1]], [[VAL2]]326; OPT-NEXT: [[CALL2:%.*]] = call i32 @escape(ptr [[PTR1]])327; OPT-NEXT: ret i32 [[ADD]]328 %ptr1 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 0329 %val1 = load i32, ptr %ptr1330 %ptr2 = getelementptr inbounds %struct.s, ptr %input, i32 0, i32 1331 %val2 = load i32, ptr %ptr2332 store ptr %input, ptr %output333 %add = add i32 %val1, %val2334 %call2 = call i32 @escape(ptr %ptr1)335 ret i32 %add336}337 338define ptx_kernel void @grid_const_phi(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input1, ptr %inout) {339; PTX-LABEL: grid_const_phi(340; PTX: {341; PTX-NEXT: .reg .pred %p<2>;342; PTX-NEXT: .reg .b32 %r<3>;343; PTX-NEXT: .reg .b64 %rd<4>;344; PTX-EMPTY:345; PTX-NEXT: // %bb.0:346; PTX-NEXT: mov.b64 %rd3, grid_const_phi_param_0;347; PTX-NEXT: ld.param.b64 %rd2, [grid_const_phi_param_1];348; PTX-NEXT: cvta.to.global.u64 %rd1, %rd2;349; PTX-NEXT: ld.global.b32 %r1, [%rd1];350; PTX-NEXT: setp.lt.s32 %p1, %r1, 0;351; PTX-NEXT: @%p1 bra $L__BB9_2;352; PTX-NEXT: // %bb.1: // %second353; PTX-NEXT: add.s64 %rd3, %rd3, 4;354; PTX-NEXT: $L__BB9_2: // %merge355; PTX-NEXT: ld.param.b32 %r2, [%rd3];356; PTX-NEXT: st.global.b32 [%rd1], %r2;357; PTX-NEXT: ret;358; OPT-LABEL: define ptx_kernel void @grid_const_phi(359; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT1:%.*]], ptr [[INOUT:%.*]]) #[[ATTR0]] {360; OPT-NEXT: [[TMP1:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])361; OPT-NEXT: [[INPUT1_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP1]] to ptr362; OPT-NEXT: [[VAL:%.*]] = load i32, ptr [[INOUT]], align 4363; OPT-NEXT: [[LESS:%.*]] = icmp slt i32 [[VAL]], 0364; OPT-NEXT: br i1 [[LESS]], label %[[FIRST:.*]], label %[[SECOND:.*]]365; OPT: [[FIRST]]:366; OPT-NEXT: [[PTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1_PARAM_GEN]], i32 0, i32 0367; OPT-NEXT: br label %[[MERGE:.*]]368; OPT: [[SECOND]]:369; OPT-NEXT: [[PTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1_PARAM_GEN]], i32 0, i32 1370; OPT-NEXT: br label %[[MERGE]]371; OPT: [[MERGE]]:372; OPT-NEXT: [[PTRNEW:%.*]] = phi ptr [ [[PTR1]], %[[FIRST]] ], [ [[PTR2]], %[[SECOND]] ]373; OPT-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4374; OPT-NEXT: store i32 [[VALLOADED]], ptr [[INOUT]], align 4375; OPT-NEXT: ret void376 377 %val = load i32, ptr %inout378 %less = icmp slt i32 %val, 0379 br i1 %less, label %first, label %second380first:381 %ptr1 = getelementptr inbounds %struct.s, ptr %input1, i32 0, i32 0382 br label %merge383second:384 %ptr2 = getelementptr inbounds %struct.s, ptr %input1, i32 0, i32 1385 br label %merge386merge:387 %ptrnew = phi ptr [%ptr1, %first], [%ptr2, %second]388 %valloaded = load i32, ptr %ptrnew389 store i32 %valloaded, ptr %inout390 ret void391}392 393; NOTE: %input2 is *not* grid_constant394define ptx_kernel void @grid_const_phi_ngc(ptr byval(%struct.s) align 4 "nvvm.grid_constant" %input1, ptr byval(%struct.s) %input2, ptr %inout) {395; PTX-LABEL: grid_const_phi_ngc(396; PTX: {397; PTX-NEXT: .reg .pred %p<2>;398; PTX-NEXT: .reg .b32 %r<3>;399; PTX-NEXT: .reg .b64 %rd<5>;400; PTX-EMPTY:401; PTX-NEXT: // %bb.0:402; PTX-NEXT: mov.b64 %rd4, grid_const_phi_ngc_param_0;403; PTX-NEXT: ld.param.b64 %rd3, [grid_const_phi_ngc_param_2];404; PTX-NEXT: cvta.to.global.u64 %rd1, %rd3;405; PTX-NEXT: ld.global.b32 %r1, [%rd1];406; PTX-NEXT: setp.lt.s32 %p1, %r1, 0;407; PTX-NEXT: @%p1 bra $L__BB10_2;408; PTX-NEXT: // %bb.1: // %second409; PTX-NEXT: mov.b64 %rd2, grid_const_phi_ngc_param_1;410; PTX-NEXT: add.s64 %rd4, %rd2, 4;411; PTX-NEXT: $L__BB10_2: // %merge412; PTX-NEXT: ld.param.b32 %r2, [%rd4];413; PTX-NEXT: st.global.b32 [%rd1], %r2;414; PTX-NEXT: ret;415; OPT-LABEL: define ptx_kernel void @grid_const_phi_ngc(416; OPT-SAME: ptr byval([[STRUCT_S:%.*]]) align 4 "nvvm.grid_constant" [[INPUT1:%.*]], ptr byval([[STRUCT_S]]) [[INPUT2:%.*]], ptr [[INOUT:%.*]]) #[[ATTR0]] {417; OPT-NEXT: [[TMP1:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT2]])418; OPT-NEXT: [[INPUT2_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP1]] to ptr419; OPT-NEXT: [[TMP2:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])420; OPT-NEXT: [[INPUT1_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP2]] to ptr421; OPT-NEXT: [[VAL:%.*]] = load i32, ptr [[INOUT]], align 4422; OPT-NEXT: [[LESS:%.*]] = icmp slt i32 [[VAL]], 0423; OPT-NEXT: br i1 [[LESS]], label %[[FIRST:.*]], label %[[SECOND:.*]]424; OPT: [[FIRST]]:425; OPT-NEXT: [[PTR1:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT1_PARAM_GEN]], i32 0, i32 0426; OPT-NEXT: br label %[[MERGE:.*]]427; OPT: [[SECOND]]:428; OPT-NEXT: [[PTR2:%.*]] = getelementptr inbounds [[STRUCT_S]], ptr [[INPUT2_PARAM_GEN]], i32 0, i32 1429; OPT-NEXT: br label %[[MERGE]]430; OPT: [[MERGE]]:431; OPT-NEXT: [[PTRNEW:%.*]] = phi ptr [ [[PTR1]], %[[FIRST]] ], [ [[PTR2]], %[[SECOND]] ]432; OPT-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4433; OPT-NEXT: store i32 [[VALLOADED]], ptr [[INOUT]], align 4434; OPT-NEXT: ret void435 %val = load i32, ptr %inout436 %less = icmp slt i32 %val, 0437 br i1 %less, label %first, label %second438first:439 %ptr1 = getelementptr inbounds %struct.s, ptr %input1, i32 0, i32 0440 br label %merge441second:442 %ptr2 = getelementptr inbounds %struct.s, ptr %input2, i32 0, i32 1443 br label %merge444merge:445 %ptrnew = phi ptr [%ptr1, %first], [%ptr2, %second]446 %valloaded = load i32, ptr %ptrnew447 store i32 %valloaded, ptr %inout448 ret void449}450 451; NOTE: %input2 is *not* grid_constant452define ptx_kernel void @grid_const_select(ptr byval(i32) align 4 "nvvm.grid_constant" %input1, ptr byval(i32) %input2, ptr %inout) {453; PTX-LABEL: grid_const_select(454; PTX: {455; PTX-NEXT: .reg .pred %p<2>;456; PTX-NEXT: .reg .b32 %r<3>;457; PTX-NEXT: .reg .b64 %rd<6>;458; PTX-EMPTY:459; PTX-NEXT: // %bb.0:460; PTX-NEXT: mov.b64 %rd1, grid_const_select_param_0;461; PTX-NEXT: ld.param.b64 %rd2, [grid_const_select_param_2];462; PTX-NEXT: cvta.to.global.u64 %rd3, %rd2;463; PTX-NEXT: mov.b64 %rd4, grid_const_select_param_1;464; PTX-NEXT: ld.global.b32 %r1, [%rd3];465; PTX-NEXT: setp.lt.s32 %p1, %r1, 0;466; PTX-NEXT: selp.b64 %rd5, %rd1, %rd4, %p1;467; PTX-NEXT: ld.param.b32 %r2, [%rd5];468; PTX-NEXT: st.global.b32 [%rd3], %r2;469; PTX-NEXT: ret;470; OPT-LABEL: define ptx_kernel void @grid_const_select(471; OPT-SAME: ptr byval(i32) align 4 "nvvm.grid_constant" [[INPUT1:%.*]], ptr byval(i32) [[INPUT2:%.*]], ptr [[INOUT:%.*]]) #[[ATTR0]] {472; OPT-NEXT: [[TMP1:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT2]])473; OPT-NEXT: [[INPUT2_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP1]] to ptr474; OPT-NEXT: [[TMP2:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT1]])475; OPT-NEXT: [[INPUT1_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[TMP2]] to ptr476; OPT-NEXT: [[VAL:%.*]] = load i32, ptr [[INOUT]], align 4477; OPT-NEXT: [[LESS:%.*]] = icmp slt i32 [[VAL]], 0478; OPT-NEXT: [[PTRNEW:%.*]] = select i1 [[LESS]], ptr [[INPUT1_PARAM_GEN]], ptr [[INPUT2_PARAM_GEN]]479; OPT-NEXT: [[VALLOADED:%.*]] = load i32, ptr [[PTRNEW]], align 4480; OPT-NEXT: store i32 [[VALLOADED]], ptr [[INOUT]], align 4481; OPT-NEXT: ret void482 %val = load i32, ptr %inout483 %less = icmp slt i32 %val, 0484 %ptrnew = select i1 %less, ptr %input1, ptr %input2485 %valloaded = load i32, ptr %ptrnew486 store i32 %valloaded, ptr %inout487 ret void488}489 490define ptx_kernel i32 @grid_const_ptrtoint(ptr byval(i32) "nvvm.grid_constant" %input) {491; PTX-LABEL: grid_const_ptrtoint(492; PTX: {493; PTX-NEXT: .reg .b32 %r<4>;494; PTX-NEXT: .reg .b64 %rd<3>;495; PTX-EMPTY:496; PTX-NEXT: // %bb.0:497; PTX-NEXT: mov.b64 %rd1, grid_const_ptrtoint_param_0;498; PTX-NEXT: ld.param.b32 %r1, [grid_const_ptrtoint_param_0];499; PTX-NEXT: cvta.param.u64 %rd2, %rd1;500; PTX-NEXT: cvt.u32.u64 %r2, %rd2;501; PTX-NEXT: add.s32 %r3, %r1, %r2;502; PTX-NEXT: st.param.b32 [func_retval0], %r3;503; PTX-NEXT: ret;504; OPT-LABEL: define ptx_kernel i32 @grid_const_ptrtoint(505; OPT-SAME: ptr byval(i32) align 4 "nvvm.grid_constant" [[INPUT:%.*]]) #[[ATTR0]] {506; OPT-NEXT: [[INPUT2:%.*]] = call ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT]])507; OPT-NEXT: [[INPUT3:%.*]] = load i32, ptr addrspace(101) [[INPUT2]], align 4508; OPT-NEXT: [[INPUT1:%.*]] = addrspacecast ptr addrspace(101) [[INPUT2]] to ptr509; OPT-NEXT: [[PTRVAL:%.*]] = ptrtoint ptr [[INPUT1]] to i32510; OPT-NEXT: [[KEEPALIVE:%.*]] = add i32 [[INPUT3]], [[PTRVAL]]511; OPT-NEXT: ret i32 [[KEEPALIVE]]512 %val = load i32, ptr %input513 %ptrval = ptrtoint ptr %input to i32514 %keepalive = add i32 %val, %ptrval515 ret i32 %keepalive516}517 518declare void @device_func(ptr byval(i32) align 4)519 520define ptx_kernel void @test_forward_byval_arg(ptr byval(i32) align 4 "nvvm.grid_constant" %input) {521; OPT-LABEL: define ptx_kernel void @test_forward_byval_arg(522; OPT-SAME: ptr byval(i32) align 4 "nvvm.grid_constant" [[INPUT:%.*]]) #[[ATTR0]] {523; OPT-NEXT: [[INPUT_PARAM:%.*]] = call align 4 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[INPUT]])524; OPT-NEXT: [[INPUT_PARAM_GEN:%.*]] = addrspacecast ptr addrspace(101) [[INPUT_PARAM]] to ptr525; OPT-NEXT: call void @device_func(ptr byval(i32) align 4 [[INPUT_PARAM_GEN]])526; OPT-NEXT: ret void527;528; PTX-LABEL: test_forward_byval_arg(529; PTX: {530; PTX-NEXT: .reg .b32 %r<2>;531; PTX-EMPTY:532; PTX-NEXT: // %bb.0:533; PTX-NEXT: { // callseq 4, 0534; PTX-NEXT: .param .align 4 .b8 param0[4];535; PTX-NEXT: ld.param.b32 %r1, [test_forward_byval_arg_param_0];536; PTX-NEXT: st.param.b32 [param0], %r1;537; PTX-NEXT: call.uni device_func, (param0);538; PTX-NEXT: } // callseq 4539; PTX-NEXT: ret;540 call void @device_func(ptr byval(i32) align 4 %input)541 ret void542}543 544 545declare dso_local void @dummy() local_unnamed_addr546declare dso_local ptr @escape(ptr) local_unnamed_addr547declare dso_local ptr @escape3(ptr, ptr, ptr) local_unnamed_addr548