345 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}4 5; Even though general vector types are not supported in PTX, we can still6; optimize loads/stores with pseudo-vector instructions of the form:7;8; ld.v2.f32 {%r0, %r1}, [%r0]9;10; which will load two floats at once into scalar registers.11 12define void @foo(ptr %a) {13; CHECK-LABEL: foo(14; CHECK: {15; CHECK-NEXT: .reg .b32 %r<5>;16; CHECK-NEXT: .reg .b64 %rd<2>;17; CHECK-EMPTY:18; CHECK-NEXT: // %bb.0:19; CHECK-NEXT: ld.param.b64 %rd1, [foo_param_0];20; CHECK-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];21; CHECK-NEXT: mul.rn.f32 %r3, %r2, %r2;22; CHECK-NEXT: mul.rn.f32 %r4, %r1, %r1;23; CHECK-NEXT: st.v2.b32 [%rd1], {%r4, %r3};24; CHECK-NEXT: ret;25 %t1 = load <2 x float>, ptr %a26 %t2 = fmul <2 x float> %t1, %t127 store <2 x float> %t2, ptr %a28 ret void29}30 31define void @foo2(ptr %a) {32; CHECK-LABEL: foo2(33; CHECK: {34; CHECK-NEXT: .reg .b32 %r<9>;35; CHECK-NEXT: .reg .b64 %rd<2>;36; CHECK-EMPTY:37; CHECK-NEXT: // %bb.0:38; CHECK-NEXT: ld.param.b64 %rd1, [foo2_param_0];39; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];40; CHECK-NEXT: mul.rn.f32 %r5, %r4, %r4;41; CHECK-NEXT: mul.rn.f32 %r6, %r3, %r3;42; CHECK-NEXT: mul.rn.f32 %r7, %r2, %r2;43; CHECK-NEXT: mul.rn.f32 %r8, %r1, %r1;44; CHECK-NEXT: st.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};45; CHECK-NEXT: ret;46 %t1 = load <4 x float>, ptr %a47 %t2 = fmul <4 x float> %t1, %t148 store <4 x float> %t2, ptr %a49 ret void50}51 52define void @foo3(ptr %a) {53; CHECK-LABEL: foo3(54; CHECK: {55; CHECK-NEXT: .reg .b32 %r<17>;56; CHECK-NEXT: .reg .b64 %rd<2>;57; CHECK-EMPTY:58; CHECK-NEXT: // %bb.0:59; CHECK-NEXT: ld.param.b64 %rd1, [foo3_param_0];60; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];61; CHECK-NEXT: ld.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];62; CHECK-NEXT: mul.rn.f32 %r9, %r8, %r8;63; CHECK-NEXT: mul.rn.f32 %r10, %r7, %r7;64; CHECK-NEXT: mul.rn.f32 %r11, %r6, %r6;65; CHECK-NEXT: mul.rn.f32 %r12, %r5, %r5;66; CHECK-NEXT: mul.rn.f32 %r13, %r4, %r4;67; CHECK-NEXT: mul.rn.f32 %r14, %r3, %r3;68; CHECK-NEXT: mul.rn.f32 %r15, %r2, %r2;69; CHECK-NEXT: mul.rn.f32 %r16, %r1, %r1;70; CHECK-NEXT: st.v4.b32 [%rd1+16], {%r16, %r15, %r14, %r13};71; CHECK-NEXT: st.v4.b32 [%rd1], {%r12, %r11, %r10, %r9};72; CHECK-NEXT: ret;73 %t1 = load <8 x float>, ptr %a74 %t2 = fmul <8 x float> %t1, %t175 store <8 x float> %t2, ptr %a76 ret void77}78 79 80 81define void @foo4(ptr %a) {82; CHECK-LABEL: foo4(83; CHECK: {84; CHECK-NEXT: .reg .b32 %r<5>;85; CHECK-NEXT: .reg .b64 %rd<2>;86; CHECK-EMPTY:87; CHECK-NEXT: // %bb.0:88; CHECK-NEXT: ld.param.b64 %rd1, [foo4_param_0];89; CHECK-NEXT: ld.v2.b32 {%r1, %r2}, [%rd1];90; CHECK-NEXT: mul.lo.s32 %r3, %r2, %r2;91; CHECK-NEXT: mul.lo.s32 %r4, %r1, %r1;92; CHECK-NEXT: st.v2.b32 [%rd1], {%r4, %r3};93; CHECK-NEXT: ret;94 %t1 = load <2 x i32>, ptr %a95 %t2 = mul <2 x i32> %t1, %t196 store <2 x i32> %t2, ptr %a97 ret void98}99 100define void @foo5(ptr %a) {101; CHECK-LABEL: foo5(102; CHECK: {103; CHECK-NEXT: .reg .b32 %r<9>;104; CHECK-NEXT: .reg .b64 %rd<2>;105; CHECK-EMPTY:106; CHECK-NEXT: // %bb.0:107; CHECK-NEXT: ld.param.b64 %rd1, [foo5_param_0];108; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];109; CHECK-NEXT: mul.lo.s32 %r5, %r4, %r4;110; CHECK-NEXT: mul.lo.s32 %r6, %r3, %r3;111; CHECK-NEXT: mul.lo.s32 %r7, %r2, %r2;112; CHECK-NEXT: mul.lo.s32 %r8, %r1, %r1;113; CHECK-NEXT: st.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};114; CHECK-NEXT: ret;115 %t1 = load <4 x i32>, ptr %a116 %t2 = mul <4 x i32> %t1, %t1117 store <4 x i32> %t2, ptr %a118 ret void119}120 121define void @foo6(ptr %a) {122; CHECK-LABEL: foo6(123; CHECK: {124; CHECK-NEXT: .reg .b32 %r<17>;125; CHECK-NEXT: .reg .b64 %rd<2>;126; CHECK-EMPTY:127; CHECK-NEXT: // %bb.0:128; CHECK-NEXT: ld.param.b64 %rd1, [foo6_param_0];129; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];130; CHECK-NEXT: ld.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];131; CHECK-NEXT: mul.lo.s32 %r9, %r8, %r8;132; CHECK-NEXT: mul.lo.s32 %r10, %r7, %r7;133; CHECK-NEXT: mul.lo.s32 %r11, %r6, %r6;134; CHECK-NEXT: mul.lo.s32 %r12, %r5, %r5;135; CHECK-NEXT: mul.lo.s32 %r13, %r4, %r4;136; CHECK-NEXT: mul.lo.s32 %r14, %r3, %r3;137; CHECK-NEXT: mul.lo.s32 %r15, %r2, %r2;138; CHECK-NEXT: mul.lo.s32 %r16, %r1, %r1;139; CHECK-NEXT: st.v4.b32 [%rd1+16], {%r16, %r15, %r14, %r13};140; CHECK-NEXT: st.v4.b32 [%rd1], {%r12, %r11, %r10, %r9};141; CHECK-NEXT: ret;142 %t1 = load <8 x i32>, ptr %a143 %t2 = mul <8 x i32> %t1, %t1144 store <8 x i32> %t2, ptr %a145 ret void146}147 148; The following test wasn't passing previously as the address149; computation was still too complex when LSV was called.150declare i32 @llvm.nvvm.read.ptx.sreg.ctaid.x() #0151declare i32 @llvm.nvvm.read.ptx.sreg.tid.x() #0152define void @foo_complex(ptr nocapture readonly align 16 dereferenceable(134217728) %alloc0) {153; CHECK-LABEL: foo_complex(154; CHECK: {155; CHECK-NEXT: .reg .b16 %rs<4>;156; CHECK-NEXT: .reg .b32 %r<8>;157; CHECK-NEXT: .reg .b64 %rd<6>;158; CHECK-EMPTY:159; CHECK-NEXT: // %bb.0:160; CHECK-NEXT: ld.param.b64 %rd1, [foo_complex_param_0];161; CHECK-NEXT: mov.u32 %r1, %tid.x;162; CHECK-NEXT: mov.u32 %r2, %ctaid.x;163; CHECK-NEXT: shr.u32 %r3, %r2, 8;164; CHECK-NEXT: shl.b32 %r4, %r2, 9;165; CHECK-NEXT: and.b32 %r5, %r4, 130560;166; CHECK-NEXT: shl.b32 %r6, %r1, 1;167; CHECK-NEXT: or.b32 %r7, %r5, %r6;168; CHECK-NEXT: cvt.u64.u32 %rd2, %r7;169; CHECK-NEXT: mul.wide.u32 %rd3, %r3, 131072;170; CHECK-NEXT: add.s64 %rd4, %rd1, %rd3;171; CHECK-NEXT: add.s64 %rd5, %rd4, %rd2;172; CHECK-NEXT: ld.v2.b8 {%rs1, %rs2}, [%rd5+128];173; CHECK-NEXT: max.u16 %rs3, %rs1, %rs2;174; CHECK-NEXT: st.b8 [%rd5+129], %rs3;175; CHECK-NEXT: ret;176 %t0 = tail call i32 @llvm.nvvm.read.ptx.sreg.tid.x(), !range !1177 %t1 = tail call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()178 %t2 = lshr i32 %t1, 8179 %t3 = shl nuw nsw i32 %t1, 9180 %ttile_origin.2 = and i32 %t3, 130560181 %tstart_offset_x_mul = shl nuw nsw i32 %t0, 1182 %t4 = or disjoint i32 %ttile_origin.2, %tstart_offset_x_mul183 %t6 = or disjoint i32 %t4, 1184 %t8 = or disjoint i32 %t4, 128185 %t9 = zext i32 %t8 to i64186 %t10 = or disjoint i32 %t4, 129187 %t11 = zext i32 %t10 to i64188 %t20 = zext i32 %t2 to i64189 %t27 = getelementptr inbounds [1024 x [131072 x i8]], ptr %alloc0, i64 0, i64 %t20, i64 %t9190 %t28 = load i8, ptr %t27, align 2191 %t31 = getelementptr inbounds [1024 x [131072 x i8]], ptr %alloc0, i64 0, i64 %t20, i64 %t11192 %t32 = load i8, ptr %t31, align 1193 %t33 = icmp ult i8 %t28, %t32194 %t34 = select i1 %t33, i8 %t32, i8 %t28195 store i8 %t34, ptr %t31196 ret void197}198 199define void @extv8f16_global_a16(ptr addrspace(1) noalias readonly align 16 %dst, ptr addrspace(1) noalias readonly align 16 %src) #0 {200; CHECK-LABEL: extv8f16_global_a16(201; CHECK: {202; CHECK-NEXT: .reg .b16 %rs<9>;203; CHECK-NEXT: .reg .b32 %r<13>;204; CHECK-NEXT: .reg .b64 %rd<3>;205; CHECK-EMPTY:206; CHECK-NEXT: // %bb.0:207; CHECK-NEXT: ld.param.b64 %rd1, [extv8f16_global_a16_param_0];208; CHECK-NEXT: ld.param.b64 %rd2, [extv8f16_global_a16_param_1];209; CHECK-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd2];210; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r3;211; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r4;212; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;213; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r2;214; CHECK-NEXT: cvt.f32.f16 %r5, %rs8;215; CHECK-NEXT: cvt.f32.f16 %r6, %rs7;216; CHECK-NEXT: cvt.f32.f16 %r7, %rs6;217; CHECK-NEXT: cvt.f32.f16 %r8, %rs5;218; CHECK-NEXT: cvt.f32.f16 %r9, %rs4;219; CHECK-NEXT: cvt.f32.f16 %r10, %rs3;220; CHECK-NEXT: cvt.f32.f16 %r11, %rs2;221; CHECK-NEXT: cvt.f32.f16 %r12, %rs1;222; CHECK-NEXT: st.global.v4.b32 [%rd1+16], {%r12, %r11, %r10, %r9};223; CHECK-NEXT: st.global.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};224; CHECK-NEXT: ret;225 %v = load <8 x half>, ptr addrspace(1) %src, align 16226 %ext = fpext <8 x half> %v to <8 x float>227 store <8 x float> %ext, ptr addrspace(1) %dst, align 16228 ret void229}230 231define void @extv8f16_global_a4(ptr addrspace(1) noalias readonly align 16 %dst, ptr addrspace(1) noalias readonly align 16 %src) #0 {232; CHECK-LABEL: extv8f16_global_a4(233; CHECK: {234; CHECK-NEXT: .reg .b16 %rs<9>;235; CHECK-NEXT: .reg .b32 %r<9>;236; CHECK-NEXT: .reg .b64 %rd<3>;237; CHECK-EMPTY:238; CHECK-NEXT: // %bb.0:239; CHECK-NEXT: ld.param.b64 %rd1, [extv8f16_global_a4_param_0];240; CHECK-NEXT: ld.param.b64 %rd2, [extv8f16_global_a4_param_1];241; CHECK-NEXT: ld.global.v2.b16 {%rs1, %rs2}, [%rd2+8];242; CHECK-NEXT: ld.global.v2.b16 {%rs3, %rs4}, [%rd2+12];243; CHECK-NEXT: ld.global.v2.b16 {%rs5, %rs6}, [%rd2];244; CHECK-NEXT: ld.global.v2.b16 {%rs7, %rs8}, [%rd2+4];245; CHECK-NEXT: cvt.f32.f16 %r1, %rs8;246; CHECK-NEXT: cvt.f32.f16 %r2, %rs7;247; CHECK-NEXT: cvt.f32.f16 %r3, %rs6;248; CHECK-NEXT: cvt.f32.f16 %r4, %rs5;249; CHECK-NEXT: cvt.f32.f16 %r5, %rs4;250; CHECK-NEXT: cvt.f32.f16 %r6, %rs3;251; CHECK-NEXT: cvt.f32.f16 %r7, %rs2;252; CHECK-NEXT: cvt.f32.f16 %r8, %rs1;253; CHECK-NEXT: st.global.v4.b32 [%rd1+16], {%r8, %r7, %r6, %r5};254; CHECK-NEXT: st.global.v4.b32 [%rd1], {%r4, %r3, %r2, %r1};255; CHECK-NEXT: ret;256 %v = load <8 x half>, ptr addrspace(1) %src, align 4257 %ext = fpext <8 x half> %v to <8 x float>258 store <8 x float> %ext, ptr addrspace(1) %dst, align 16259 ret void260}261 262 263define void @extv8f16_generic_a16(ptr noalias readonly align 16 %dst, ptr noalias readonly align 16 %src) #0 {264; CHECK-LABEL: extv8f16_generic_a16(265; CHECK: {266; CHECK-NEXT: .reg .b16 %rs<9>;267; CHECK-NEXT: .reg .b32 %r<13>;268; CHECK-NEXT: .reg .b64 %rd<3>;269; CHECK-EMPTY:270; CHECK-NEXT: // %bb.0:271; CHECK-NEXT: ld.param.b64 %rd1, [extv8f16_generic_a16_param_0];272; CHECK-NEXT: ld.param.b64 %rd2, [extv8f16_generic_a16_param_1];273; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd2];274; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r3;275; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r4;276; CHECK-NEXT: mov.b32 {%rs5, %rs6}, %r1;277; CHECK-NEXT: mov.b32 {%rs7, %rs8}, %r2;278; CHECK-NEXT: cvt.f32.f16 %r5, %rs8;279; CHECK-NEXT: cvt.f32.f16 %r6, %rs7;280; CHECK-NEXT: cvt.f32.f16 %r7, %rs6;281; CHECK-NEXT: cvt.f32.f16 %r8, %rs5;282; CHECK-NEXT: cvt.f32.f16 %r9, %rs4;283; CHECK-NEXT: cvt.f32.f16 %r10, %rs3;284; CHECK-NEXT: cvt.f32.f16 %r11, %rs2;285; CHECK-NEXT: cvt.f32.f16 %r12, %rs1;286; CHECK-NEXT: st.v4.b32 [%rd1+16], {%r12, %r11, %r10, %r9};287; CHECK-NEXT: st.v4.b32 [%rd1], {%r8, %r7, %r6, %r5};288; CHECK-NEXT: ret;289 %v = load <8 x half>, ptr %src, align 16290 %ext = fpext <8 x half> %v to <8 x float>291 store <8 x float> %ext, ptr %dst, align 16292 ret void293}294 295define void @extv8f16_generic_a4(ptr noalias readonly align 16 %dst, ptr noalias readonly align 16 %src) #0 {296; CHECK-LABEL: extv8f16_generic_a4(297; CHECK: {298; CHECK-NEXT: .reg .b16 %rs<9>;299; CHECK-NEXT: .reg .b32 %r<9>;300; CHECK-NEXT: .reg .b64 %rd<3>;301; CHECK-EMPTY:302; CHECK-NEXT: // %bb.0:303; CHECK-NEXT: ld.param.b64 %rd1, [extv8f16_generic_a4_param_0];304; CHECK-NEXT: ld.param.b64 %rd2, [extv8f16_generic_a4_param_1];305; CHECK-NEXT: ld.v2.b16 {%rs1, %rs2}, [%rd2+8];306; CHECK-NEXT: ld.v2.b16 {%rs3, %rs4}, [%rd2+12];307; CHECK-NEXT: ld.v2.b16 {%rs5, %rs6}, [%rd2];308; CHECK-NEXT: ld.v2.b16 {%rs7, %rs8}, [%rd2+4];309; CHECK-NEXT: cvt.f32.f16 %r1, %rs8;310; CHECK-NEXT: cvt.f32.f16 %r2, %rs7;311; CHECK-NEXT: cvt.f32.f16 %r3, %rs6;312; CHECK-NEXT: cvt.f32.f16 %r4, %rs5;313; CHECK-NEXT: cvt.f32.f16 %r5, %rs4;314; CHECK-NEXT: cvt.f32.f16 %r6, %rs3;315; CHECK-NEXT: cvt.f32.f16 %r7, %rs2;316; CHECK-NEXT: cvt.f32.f16 %r8, %rs1;317; CHECK-NEXT: st.v4.b32 [%rd1+16], {%r8, %r7, %r6, %r5};318; CHECK-NEXT: st.v4.b32 [%rd1], {%r4, %r3, %r2, %r1};319; CHECK-NEXT: ret;320 %v = load <8 x half>, ptr %src, align 4321 %ext = fpext <8 x half> %v to <8 x float>322 store <8 x float> %ext, ptr %dst, align 16323 ret void324}325 326 327!1 = !{i32 0, i32 64}328 329define dso_local void @bf16_v4_align_load_store(ptr noundef %0, ptr noundef %1) #0 {330; CHECK-LABEL: bf16_v4_align_load_store(331; CHECK: {332; CHECK-NEXT: .reg .b32 %r<3>;333; CHECK-NEXT: .reg .b64 %rd<3>;334; CHECK-EMPTY:335; CHECK-NEXT: // %bb.0:336; CHECK-NEXT: ld.param.b64 %rd1, [bf16_v4_align_load_store_param_0];337; CHECK-NEXT: ld.param.b64 %rd2, [bf16_v4_align_load_store_param_1];338; CHECK-NEXT: ld.v2.b32 {%r1, %r2}, [%rd2];339; CHECK-NEXT: st.v2.b32 [%rd1], {%r1, %r2};340; CHECK-NEXT: ret;341 %3 = load <4 x bfloat>, ptr %1, align 8342 store <4 x bfloat> %3, ptr %0, align 8343 ret void344}345