292 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_35 -verify-machineinstrs | FileCheck %s3; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_35 | %ptxas-verify %}4 5; Check that invariant loads from the global addrspace are lowered to6; ld.global.nc.7 8define i32 @ld_global(ptr addrspace(1) %ptr) {9; CHECK-LABEL: ld_global(10; CHECK: {11; CHECK-NEXT: .reg .b32 %r<2>;12; CHECK-NEXT: .reg .b64 %rd<2>;13; CHECK-EMPTY:14; CHECK-NEXT: // %bb.0:15; CHECK-NEXT: ld.param.b64 %rd1, [ld_global_param_0];16; CHECK-NEXT: ld.global.nc.b32 %r1, [%rd1];17; CHECK-NEXT: st.param.b32 [func_retval0], %r1;18; CHECK-NEXT: ret;19 %a = load i32, ptr addrspace(1) %ptr, !invariant.load !020 ret i32 %a21}22 23define half @ld_global_v2f16(ptr addrspace(1) %ptr) {24; Load of v2f16 is weird. We consider it to be a legal type, which happens to be25; loaded/stored as a 32-bit scalar.26; CHECK-LABEL: ld_global_v2f16(27; CHECK: {28; CHECK-NEXT: .reg .b16 %rs<4>;29; CHECK-NEXT: .reg .b32 %r<4>;30; CHECK-NEXT: .reg .b64 %rd<2>;31; CHECK-EMPTY:32; CHECK-NEXT: // %bb.0:33; CHECK-NEXT: ld.param.b64 %rd1, [ld_global_v2f16_param_0];34; CHECK-NEXT: ld.global.nc.v2.b16 {%rs1, %rs2}, [%rd1];35; CHECK-NEXT: cvt.f32.f16 %r1, %rs2;36; CHECK-NEXT: cvt.f32.f16 %r2, %rs1;37; CHECK-NEXT: add.rn.f32 %r3, %r2, %r1;38; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r3;39; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;40; CHECK-NEXT: ret;41 %a = load <2 x half>, ptr addrspace(1) %ptr, !invariant.load !042 %v1 = extractelement <2 x half> %a, i32 043 %v2 = extractelement <2 x half> %a, i32 144 %sum = fadd half %v1, %v245 ret half %sum46}47 48define half @ld_global_v4f16(ptr addrspace(1) %ptr) {49; Larger f16 vectors may be split into individual f16 elements and multiple50; loads/stores may be vectorized using f16 element type. Practically it's51; limited to v4 variant only.52; CHECK-LABEL: ld_global_v4f16(53; CHECK: {54; CHECK-NEXT: .reg .b16 %rs<8>;55; CHECK-NEXT: .reg .b32 %r<10>;56; CHECK-NEXT: .reg .b64 %rd<2>;57; CHECK-EMPTY:58; CHECK-NEXT: // %bb.0:59; CHECK-NEXT: ld.param.b64 %rd1, [ld_global_v4f16_param_0];60; CHECK-NEXT: ld.global.nc.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];61; CHECK-NEXT: cvt.f32.f16 %r1, %rs2;62; CHECK-NEXT: cvt.f32.f16 %r2, %rs1;63; CHECK-NEXT: add.rn.f32 %r3, %r2, %r1;64; CHECK-NEXT: cvt.rn.f16.f32 %rs5, %r3;65; CHECK-NEXT: cvt.f32.f16 %r4, %rs4;66; CHECK-NEXT: cvt.f32.f16 %r5, %rs3;67; CHECK-NEXT: add.rn.f32 %r6, %r5, %r4;68; CHECK-NEXT: cvt.rn.f16.f32 %rs6, %r6;69; CHECK-NEXT: cvt.f32.f16 %r7, %rs6;70; CHECK-NEXT: cvt.f32.f16 %r8, %rs5;71; CHECK-NEXT: add.rn.f32 %r9, %r8, %r7;72; CHECK-NEXT: cvt.rn.f16.f32 %rs7, %r9;73; CHECK-NEXT: st.param.b16 [func_retval0], %rs7;74; CHECK-NEXT: ret;75 %a = load <4 x half>, ptr addrspace(1) %ptr, !invariant.load !076 %v1 = extractelement <4 x half> %a, i32 077 %v2 = extractelement <4 x half> %a, i32 178 %v3 = extractelement <4 x half> %a, i32 279 %v4 = extractelement <4 x half> %a, i32 380 %sum1 = fadd half %v1, %v281 %sum2 = fadd half %v3, %v482 %sum = fadd half %sum1, %sum283 ret half %sum84}85 86define half @ld_global_v8f16(ptr addrspace(1) %ptr) {87; Larger vectors are, again, loaded as v4i32. PTX has no v8 variants of loads/stores,88; so load/store vectorizer has to convert v8f16 -> v4 x v2f16.89; CHECK-LABEL: ld_global_v8f16(90; CHECK: {91; CHECK-NEXT: .reg .b16 %rs<8>;92; CHECK-NEXT: .reg .b32 %r<14>;93; CHECK-NEXT: .reg .b64 %rd<2>;94; CHECK-EMPTY:95; CHECK-NEXT: // %bb.0:96; CHECK-NEXT: ld.param.b64 %rd1, [ld_global_v8f16_param_0];97; CHECK-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];98; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs1, tmp}, %r3; }99; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs2, tmp}, %r4; }100; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs3, tmp}, %r1; }101; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs4, tmp}, %r2; }102; CHECK-NEXT: cvt.f32.f16 %r5, %rs4;103; CHECK-NEXT: cvt.f32.f16 %r6, %rs3;104; CHECK-NEXT: add.rn.f32 %r7, %r6, %r5;105; CHECK-NEXT: cvt.rn.f16.f32 %rs5, %r7;106; CHECK-NEXT: cvt.f32.f16 %r8, %rs2;107; CHECK-NEXT: cvt.f32.f16 %r9, %rs1;108; CHECK-NEXT: add.rn.f32 %r10, %r9, %r8;109; CHECK-NEXT: cvt.rn.f16.f32 %rs6, %r10;110; CHECK-NEXT: cvt.f32.f16 %r11, %rs6;111; CHECK-NEXT: cvt.f32.f16 %r12, %rs5;112; CHECK-NEXT: add.rn.f32 %r13, %r12, %r11;113; CHECK-NEXT: cvt.rn.f16.f32 %rs7, %r13;114; CHECK-NEXT: st.param.b16 [func_retval0], %rs7;115; CHECK-NEXT: ret;116 %a = load <8 x half>, ptr addrspace(1) %ptr, !invariant.load !0117 %v1 = extractelement <8 x half> %a, i32 0118 %v2 = extractelement <8 x half> %a, i32 2119 %v3 = extractelement <8 x half> %a, i32 4120 %v4 = extractelement <8 x half> %a, i32 6121 %sum1 = fadd half %v1, %v2122 %sum2 = fadd half %v3, %v4123 %sum = fadd half %sum1, %sum2124 ret half %sum125}126 127define i8 @ld_global_v8i8(ptr addrspace(1) %ptr) {128; CHECK-LABEL: ld_global_v8i8(129; CHECK: {130; CHECK-NEXT: .reg .b16 %rs<8>;131; CHECK-NEXT: .reg .b32 %r<8>;132; CHECK-NEXT: .reg .b64 %rd<2>;133; CHECK-EMPTY:134; CHECK-NEXT: // %bb.0:135; CHECK-NEXT: ld.param.b64 %rd1, [ld_global_v8i8_param_0];136; CHECK-NEXT: ld.global.nc.v2.b32 {%r1, %r2}, [%rd1];137; CHECK-NEXT: prmt.b32 %r3, %r2, 0, 0x7772U;138; CHECK-NEXT: cvt.u16.u32 %rs1, %r3;139; CHECK-NEXT: prmt.b32 %r4, %r2, 0, 0x7770U;140; CHECK-NEXT: cvt.u16.u32 %rs2, %r4;141; CHECK-NEXT: prmt.b32 %r5, %r1, 0, 0x7772U;142; CHECK-NEXT: cvt.u16.u32 %rs3, %r5;143; CHECK-NEXT: prmt.b32 %r6, %r1, 0, 0x7770U;144; CHECK-NEXT: cvt.u16.u32 %rs4, %r6;145; CHECK-NEXT: add.s16 %rs5, %rs4, %rs3;146; CHECK-NEXT: add.s16 %rs6, %rs2, %rs1;147; CHECK-NEXT: add.s16 %rs7, %rs5, %rs6;148; CHECK-NEXT: cvt.u32.u16 %r7, %rs7;149; CHECK-NEXT: st.param.b32 [func_retval0], %r7;150; CHECK-NEXT: ret;151 %a = load <8 x i8>, ptr addrspace(1) %ptr, !invariant.load !0152 %v1 = extractelement <8 x i8> %a, i32 0153 %v2 = extractelement <8 x i8> %a, i32 2154 %v3 = extractelement <8 x i8> %a, i32 4155 %v4 = extractelement <8 x i8> %a, i32 6156 %sum1 = add i8 %v1, %v2157 %sum2 = add i8 %v3, %v4158 %sum = add i8 %sum1, %sum2159 ret i8 %sum160}161 162define i8 @ld_global_v16i8(ptr addrspace(1) %ptr) {163; CHECK-LABEL: ld_global_v16i8(164; CHECK: {165; CHECK-NEXT: .reg .b16 %rs<16>;166; CHECK-NEXT: .reg .b32 %r<14>;167; CHECK-NEXT: .reg .b64 %rd<2>;168; CHECK-EMPTY:169; CHECK-NEXT: // %bb.0:170; CHECK-NEXT: ld.param.b64 %rd1, [ld_global_v16i8_param_0];171; CHECK-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];172; CHECK-NEXT: prmt.b32 %r5, %r4, 0, 0x7772U;173; CHECK-NEXT: cvt.u16.u32 %rs1, %r5;174; CHECK-NEXT: prmt.b32 %r6, %r4, 0, 0x7770U;175; CHECK-NEXT: cvt.u16.u32 %rs2, %r6;176; CHECK-NEXT: prmt.b32 %r7, %r3, 0, 0x7772U;177; CHECK-NEXT: cvt.u16.u32 %rs3, %r7;178; CHECK-NEXT: prmt.b32 %r8, %r3, 0, 0x7770U;179; CHECK-NEXT: cvt.u16.u32 %rs4, %r8;180; CHECK-NEXT: prmt.b32 %r9, %r2, 0, 0x7772U;181; CHECK-NEXT: cvt.u16.u32 %rs5, %r9;182; CHECK-NEXT: prmt.b32 %r10, %r2, 0, 0x7770U;183; CHECK-NEXT: cvt.u16.u32 %rs6, %r10;184; CHECK-NEXT: prmt.b32 %r11, %r1, 0, 0x7772U;185; CHECK-NEXT: cvt.u16.u32 %rs7, %r11;186; CHECK-NEXT: prmt.b32 %r12, %r1, 0, 0x7770U;187; CHECK-NEXT: cvt.u16.u32 %rs8, %r12;188; CHECK-NEXT: add.s16 %rs9, %rs8, %rs7;189; CHECK-NEXT: add.s16 %rs10, %rs6, %rs5;190; CHECK-NEXT: add.s16 %rs11, %rs4, %rs3;191; CHECK-NEXT: add.s16 %rs12, %rs2, %rs1;192; CHECK-NEXT: add.s16 %rs13, %rs9, %rs10;193; CHECK-NEXT: add.s16 %rs14, %rs11, %rs12;194; CHECK-NEXT: add.s16 %rs15, %rs13, %rs14;195; CHECK-NEXT: cvt.u32.u16 %r13, %rs15;196; CHECK-NEXT: st.param.b32 [func_retval0], %r13;197; CHECK-NEXT: ret;198 %a = load <16 x i8>, ptr addrspace(1) %ptr, !invariant.load !0199 %v1 = extractelement <16 x i8> %a, i32 0200 %v2 = extractelement <16 x i8> %a, i32 2201 %v3 = extractelement <16 x i8> %a, i32 4202 %v4 = extractelement <16 x i8> %a, i32 6203 %v5 = extractelement <16 x i8> %a, i32 8204 %v6 = extractelement <16 x i8> %a, i32 10205 %v7 = extractelement <16 x i8> %a, i32 12206 %v8 = extractelement <16 x i8> %a, i32 14207 %sum1 = add i8 %v1, %v2208 %sum2 = add i8 %v3, %v4209 %sum3 = add i8 %v5, %v6210 %sum4 = add i8 %v7, %v8211 %sum5 = add i8 %sum1, %sum2212 %sum6 = add i8 %sum3, %sum4213 %sum7 = add i8 %sum5, %sum6214 ret i8 %sum7215}216 217define i32 @ld_global_v2i32(ptr addrspace(1) %ptr) {218; CHECK-LABEL: ld_global_v2i32(219; CHECK: {220; CHECK-NEXT: .reg .b32 %r<4>;221; CHECK-NEXT: .reg .b64 %rd<2>;222; CHECK-EMPTY:223; CHECK-NEXT: // %bb.0:224; CHECK-NEXT: ld.param.b64 %rd1, [ld_global_v2i32_param_0];225; CHECK-NEXT: ld.global.nc.v2.b32 {%r1, %r2}, [%rd1];226; CHECK-NEXT: add.s32 %r3, %r1, %r2;227; CHECK-NEXT: st.param.b32 [func_retval0], %r3;228; CHECK-NEXT: ret;229 %a = load <2 x i32>, ptr addrspace(1) %ptr, !invariant.load !0230 %v1 = extractelement <2 x i32> %a, i32 0231 %v2 = extractelement <2 x i32> %a, i32 1232 %sum = add i32 %v1, %v2233 ret i32 %sum234}235 236define i32 @ld_global_v4i32(ptr addrspace(1) %ptr) {237; CHECK-LABEL: ld_global_v4i32(238; CHECK: {239; CHECK-NEXT: .reg .b32 %r<8>;240; CHECK-NEXT: .reg .b64 %rd<2>;241; CHECK-EMPTY:242; CHECK-NEXT: // %bb.0:243; CHECK-NEXT: ld.param.b64 %rd1, [ld_global_v4i32_param_0];244; CHECK-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];245; CHECK-NEXT: add.s32 %r5, %r1, %r2;246; CHECK-NEXT: add.s32 %r6, %r3, %r4;247; CHECK-NEXT: add.s32 %r7, %r5, %r6;248; CHECK-NEXT: st.param.b32 [func_retval0], %r7;249; CHECK-NEXT: ret;250 %a = load <4 x i32>, ptr addrspace(1) %ptr, !invariant.load !0251 %v1 = extractelement <4 x i32> %a, i32 0252 %v2 = extractelement <4 x i32> %a, i32 1253 %v3 = extractelement <4 x i32> %a, i32 2254 %v4 = extractelement <4 x i32> %a, i32 3255 %sum1 = add i32 %v1, %v2256 %sum2 = add i32 %v3, %v4257 %sum3 = add i32 %sum1, %sum2258 ret i32 %sum3259}260 261define i32 @ld_not_invariant(ptr addrspace(1) %ptr) {262; CHECK-LABEL: ld_not_invariant(263; CHECK: {264; CHECK-NEXT: .reg .b32 %r<2>;265; CHECK-NEXT: .reg .b64 %rd<2>;266; CHECK-EMPTY:267; CHECK-NEXT: // %bb.0:268; CHECK-NEXT: ld.param.b64 %rd1, [ld_not_invariant_param_0];269; CHECK-NEXT: ld.global.b32 %r1, [%rd1];270; CHECK-NEXT: st.param.b32 [func_retval0], %r1;271; CHECK-NEXT: ret;272 %a = load i32, ptr addrspace(1) %ptr273 ret i32 %a274}275 276define i32 @ld_not_global_addrspace(ptr addrspace(0) %ptr) {277; CHECK-LABEL: ld_not_global_addrspace(278; CHECK: {279; CHECK-NEXT: .reg .b32 %r<2>;280; CHECK-NEXT: .reg .b64 %rd<2>;281; CHECK-EMPTY:282; CHECK-NEXT: // %bb.0:283; CHECK-NEXT: ld.param.b64 %rd1, [ld_not_global_addrspace_param_0];284; CHECK-NEXT: ld.b32 %r1, [%rd1];285; CHECK-NEXT: st.param.b32 [func_retval0], %r1;286; CHECK-NEXT: ret;287 %a = load i32, ptr addrspace(0) %ptr288 ret i32 %a289}290 291!0 = !{}292