249 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 | FileCheck %s --check-prefixes IR,IRC3; RUN: opt < %s -S -nvptx-lower-args --mtriple nvptx64-nvidia-nvcl | FileCheck %s --check-prefixes IR,IRO4; RUN: llc < %s -mcpu=sm_20 --mtriple nvptx64-nvidia-cuda | FileCheck %s --check-prefixes PTX,PTXC5; RUN: llc < %s -mcpu=sm_20 --mtriple nvptx64-nvidia-nvcl| FileCheck %s --check-prefixes PTX,PTXO6; RUN: %if ptxas %{ llc < %s -mcpu=sm_20 | %ptxas-verify %}7 8target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"9target triple = "nvptx64-nvidia-cuda"10 11%class.outer = type <{ %class.inner, i32, [4 x i8] }>12%class.inner = type { ptr, ptr }13%class.padded = type { i8, i32 }14 15; Check that nvptx-lower-args preserves arg alignment16define void @load_alignment(ptr nocapture readonly byval(%class.outer) align 8 %arg) {17; IR-LABEL: define void @load_alignment(18; IR-SAME: ptr readonly byval([[CLASS_OUTER:%.*]]) align 8 captures(none) [[ARG:%.*]]) {19; IR-NEXT: [[ENTRY:.*:]]20; IR-NEXT: [[ARG_IDX_VAL:%.*]] = load ptr, ptr [[ARG]], align 821; IR-NEXT: [[ARG_IDX1:%.*]] = getelementptr [[CLASS_OUTER]], ptr [[ARG]], i64 0, i32 0, i32 122; IR-NEXT: [[ARG_IDX1_VAL:%.*]] = load ptr, ptr [[ARG_IDX1]], align 823; IR-NEXT: [[ARG_IDX2:%.*]] = getelementptr [[CLASS_OUTER]], ptr [[ARG]], i64 0, i32 124; IR-NEXT: [[ARG_IDX2_VAL:%.*]] = load i32, ptr [[ARG_IDX2]], align 825; IR-NEXT: [[ARG_IDX_VAL_VAL:%.*]] = load i32, ptr [[ARG_IDX_VAL]], align 426; IR-NEXT: [[ADD_I:%.*]] = add nsw i32 [[ARG_IDX_VAL_VAL]], [[ARG_IDX2_VAL]]27; IR-NEXT: store i32 [[ADD_I]], ptr [[ARG_IDX1_VAL]], align 428; IR-NEXT: [[TMP:%.*]] = call ptr @escape(ptr nonnull [[ARG_IDX2]])29; IR-NEXT: ret void30;31; PTX-LABEL: load_alignment(32; PTX: {33; PTX-NEXT: .reg .b32 %r<4>;34; PTX-NEXT: .reg .b64 %rd<6>;35; PTX-EMPTY:36; PTX-NEXT: // %bb.0: // %entry37; PTX-NEXT: mov.b64 %rd1, load_alignment_param_0;38; PTX-NEXT: ld.local.b64 %rd2, [%rd1];39; PTX-NEXT: ld.local.b64 %rd3, [%rd1+8];40; PTX-NEXT: add.s64 %rd4, %rd1, 16;41; PTX-NEXT: cvta.local.u64 %rd5, %rd4;42; PTX-NEXT: ld.local.b32 %r1, [%rd1+16];43; PTX-NEXT: ld.b32 %r2, [%rd2];44; PTX-NEXT: add.s32 %r3, %r2, %r1;45; PTX-NEXT: st.b32 [%rd3], %r3;46; PTX-NEXT: { // callseq 0, 047; PTX-NEXT: .param .b64 param0;48; PTX-NEXT: .param .b64 retval0;49; PTX-NEXT: st.param.b64 [param0], %rd5;50; PTX-NEXT: call.uni (retval0), escape, (param0);51; PTX-NEXT: } // callseq 052; PTX-NEXT: ret;53entry:54 %arg.idx.val = load ptr, ptr %arg, align 855 %arg.idx1 = getelementptr %class.outer, ptr %arg, i64 0, i32 0, i32 156 %arg.idx1.val = load ptr, ptr %arg.idx1, align 857 %arg.idx2 = getelementptr %class.outer, ptr %arg, i64 0, i32 158 %arg.idx2.val = load i32, ptr %arg.idx2, align 859 %arg.idx.val.val = load i32, ptr %arg.idx.val, align 460 %add.i = add nsw i32 %arg.idx.val.val, %arg.idx2.val61 store i32 %add.i, ptr %arg.idx1.val, align 462 63 ; let the pointer escape so we still create a local copy this test uses to64 ; check the load alignment.65 %tmp = call ptr @escape(ptr nonnull %arg.idx2)66 ret void67}68 69; Check that nvptx-lower-args copies padding as the struct may have been a union70define void @load_padding(ptr nocapture readonly byval(%class.padded) %arg) {71; IR-LABEL: define void @load_padding(72; IR-SAME: ptr readonly byval([[CLASS_PADDED:%.*]]) align 4 captures(none) [[ARG:%.*]]) {73; IR-NEXT: [[TMP:%.*]] = call ptr @escape(ptr nonnull align 16 [[ARG]])74; IR-NEXT: ret void75;76; PTX-LABEL: load_padding(77; PTX: {78; PTX-NEXT: .reg .b64 %rd<3>;79; PTX-EMPTY:80; PTX-NEXT: // %bb.0:81; PTX-NEXT: mov.b64 %rd1, load_padding_param_0;82; PTX-NEXT: cvta.local.u64 %rd2, %rd1;83; PTX-NEXT: { // callseq 1, 084; PTX-NEXT: .param .b64 param0;85; PTX-NEXT: .param .b64 retval0;86; PTX-NEXT: st.param.b64 [param0], %rd2;87; PTX-NEXT: call.uni (retval0), escape, (param0);88; PTX-NEXT: } // callseq 189; PTX-NEXT: ret;90 %tmp = call ptr @escape(ptr nonnull align 16 %arg)91 ret void92}93 94; OpenCL can't make assumptions about incoming pointer, so we should generate95; generic pointers load/store.96define ptx_kernel void @ptr_generic(ptr %out, ptr %in) {97; IR-LABEL: define ptx_kernel void @ptr_generic(98; IR-SAME: ptr [[OUT:%.*]], ptr [[IN:%.*]]) {99; IR-NEXT: [[V:%.*]] = load i32, ptr [[IN]], align 4100; IR-NEXT: store i32 [[V]], ptr [[OUT]], align 4101; IR-NEXT: ret void102;103; PTXC-LABEL: ptr_generic(104; PTXC: {105; PTXC-NEXT: .reg .b32 %r<2>;106; PTXC-NEXT: .reg .b64 %rd<5>;107; PTXC-EMPTY:108; PTXC-NEXT: // %bb.0:109; PTXC-NEXT: ld.param.b64 %rd1, [ptr_generic_param_0];110; PTXC-NEXT: ld.param.b64 %rd2, [ptr_generic_param_1];111; PTXC-NEXT: cvta.to.global.u64 %rd3, %rd2;112; PTXC-NEXT: cvta.to.global.u64 %rd4, %rd1;113; PTXC-NEXT: ld.global.b32 %r1, [%rd3];114; PTXC-NEXT: st.global.b32 [%rd4], %r1;115; PTXC-NEXT: ret;116;117; PTXO-LABEL: ptr_generic(118; PTXO: {119; PTXO-NEXT: .reg .b32 %r<2>;120; PTXO-NEXT: .reg .b64 %rd<3>;121; PTXO-EMPTY:122; PTXO-NEXT: // %bb.0:123; PTXO-NEXT: ld.param.b64 %rd1, [ptr_generic_param_0];124; PTXO-NEXT: ld.param.b64 %rd2, [ptr_generic_param_1];125; PTXO-NEXT: ld.b32 %r1, [%rd2];126; PTXO-NEXT: st.b32 [%rd1], %r1;127; PTXO-NEXT: ret;128 %v = load i32, ptr %in, align 4129 store i32 %v, ptr %out, align 4130 ret void131}132 133define ptx_kernel void @ptr_nongeneric(ptr addrspace(1) %out, ptr addrspace(3) %in) {134; IR-LABEL: define ptx_kernel void @ptr_nongeneric(135; IR-SAME: ptr addrspace(1) [[OUT:%.*]], ptr addrspace(3) [[IN:%.*]]) {136; IR-NEXT: [[V:%.*]] = load i32, ptr addrspace(3) [[IN]], align 4137; IR-NEXT: store i32 [[V]], ptr addrspace(1) [[OUT]], align 4138; IR-NEXT: ret void139;140; PTX-LABEL: ptr_nongeneric(141; PTX: {142; PTX-NEXT: .reg .b32 %r<2>;143; PTX-NEXT: .reg .b64 %rd<3>;144; PTX-EMPTY:145; PTX-NEXT: // %bb.0:146; PTX-NEXT: ld.param.b64 %rd1, [ptr_nongeneric_param_0];147; PTX-NEXT: ld.param.b64 %rd2, [ptr_nongeneric_param_1];148; PTX-NEXT: ld.shared.b32 %r1, [%rd2];149; PTX-NEXT: st.global.b32 [%rd1], %r1;150; PTX-NEXT: ret;151 %v = load i32, ptr addrspace(3) %in, align 4152 store i32 %v, ptr addrspace(1) %out, align 4153 ret void154}155 156define ptx_kernel void @ptr_as_int(i64 noundef %i, i32 noundef %v) {157; IRC-LABEL: define ptx_kernel void @ptr_as_int(158; IRC-SAME: i64 noundef [[I:%.*]], i32 noundef [[V:%.*]]) {159; IRC-NEXT: [[P:%.*]] = inttoptr i64 [[I]] to ptr160; IRC-NEXT: [[P1:%.*]] = addrspacecast ptr [[P]] to ptr addrspace(1)161; IRC-NEXT: [[P2:%.*]] = addrspacecast ptr addrspace(1) [[P1]] to ptr162; IRC-NEXT: store i32 [[V]], ptr [[P2]], align 4163; IRC-NEXT: ret void164;165; IRO-LABEL: define ptx_kernel void @ptr_as_int(166; IRO-SAME: i64 noundef [[I:%.*]], i32 noundef [[V:%.*]]) {167; IRO-NEXT: [[P:%.*]] = inttoptr i64 [[I]] to ptr168; IRO-NEXT: store i32 [[V]], ptr [[P]], align 4169; IRO-NEXT: ret void170;171; PTXC-LABEL: ptr_as_int(172; PTXC: {173; PTXC-NEXT: .reg .b32 %r<2>;174; PTXC-NEXT: .reg .b64 %rd<3>;175; PTXC-EMPTY:176; PTXC-NEXT: // %bb.0:177; PTXC-NEXT: ld.param.b64 %rd1, [ptr_as_int_param_0];178; PTXC-NEXT: ld.param.b32 %r1, [ptr_as_int_param_1];179; PTXC-NEXT: cvta.to.global.u64 %rd2, %rd1;180; PTXC-NEXT: st.global.b32 [%rd2], %r1;181; PTXC-NEXT: ret;182;183; PTXO-LABEL: ptr_as_int(184; PTXO: {185; PTXO-NEXT: .reg .b32 %r<2>;186; PTXO-NEXT: .reg .b64 %rd<2>;187; PTXO-EMPTY:188; PTXO-NEXT: // %bb.0:189; PTXO-NEXT: ld.param.b64 %rd1, [ptr_as_int_param_0];190; PTXO-NEXT: ld.param.b32 %r1, [ptr_as_int_param_1];191; PTXO-NEXT: st.b32 [%rd1], %r1;192; PTXO-NEXT: ret;193 %p = inttoptr i64 %i to ptr194 store i32 %v, ptr %p, align 4195 ret void196}197 198%struct.S = type { i64 }199 200define ptx_kernel void @ptr_as_int_aggr(ptr nocapture noundef readonly byval(%struct.S) align 8 %s, i32 noundef %v) {201; IRC-LABEL: define ptx_kernel void @ptr_as_int_aggr(202; IRC-SAME: ptr noundef readonly byval([[STRUCT_S:%.*]]) align 8 captures(none) [[S:%.*]], i32 noundef [[V:%.*]]) {203; IRC-NEXT: [[S3:%.*]] = call align 8 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])204; IRC-NEXT: [[I:%.*]] = load i64, ptr addrspace(101) [[S3]], align 8205; IRC-NEXT: [[P:%.*]] = inttoptr i64 [[I]] to ptr206; IRC-NEXT: [[P1:%.*]] = addrspacecast ptr [[P]] to ptr addrspace(1)207; IRC-NEXT: [[P2:%.*]] = addrspacecast ptr addrspace(1) [[P1]] to ptr208; IRC-NEXT: store i32 [[V]], ptr [[P2]], align 4209; IRC-NEXT: ret void210;211; IRO-LABEL: define ptx_kernel void @ptr_as_int_aggr(212; IRO-SAME: ptr noundef readonly byval([[STRUCT_S:%.*]]) align 8 captures(none) [[S:%.*]], i32 noundef [[V:%.*]]) {213; IRO-NEXT: [[S1:%.*]] = call align 8 ptr addrspace(101) @llvm.nvvm.internal.addrspace.wrap.p101.p0(ptr [[S]])214; IRO-NEXT: [[I:%.*]] = load i64, ptr addrspace(101) [[S1]], align 8215; IRO-NEXT: [[P:%.*]] = inttoptr i64 [[I]] to ptr216; IRO-NEXT: store i32 [[V]], ptr [[P]], align 4217; IRO-NEXT: ret void218;219; PTXC-LABEL: ptr_as_int_aggr(220; PTXC: {221; PTXC-NEXT: .reg .b32 %r<2>;222; PTXC-NEXT: .reg .b64 %rd<3>;223; PTXC-EMPTY:224; PTXC-NEXT: // %bb.0:225; PTXC-NEXT: ld.param.b32 %r1, [ptr_as_int_aggr_param_1];226; PTXC-NEXT: ld.param.b64 %rd1, [ptr_as_int_aggr_param_0];227; PTXC-NEXT: cvta.to.global.u64 %rd2, %rd1;228; PTXC-NEXT: st.global.b32 [%rd2], %r1;229; PTXC-NEXT: ret;230;231; PTXO-LABEL: ptr_as_int_aggr(232; PTXO: {233; PTXO-NEXT: .reg .b32 %r<2>;234; PTXO-NEXT: .reg .b64 %rd<2>;235; PTXO-EMPTY:236; PTXO-NEXT: // %bb.0:237; PTXO-NEXT: ld.param.b32 %r1, [ptr_as_int_aggr_param_1];238; PTXO-NEXT: ld.param.b64 %rd1, [ptr_as_int_aggr_param_0];239; PTXO-NEXT: st.b32 [%rd1], %r1;240; PTXO-NEXT: ret;241 %i = load i64, ptr %s, align 8242 %p = inttoptr i64 %i to ptr243 store i32 %v, ptr %p, align 4244 ret void245}246 247; Function Attrs: convergent nounwind248declare dso_local ptr @escape(ptr) local_unnamed_addr249