517 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_90 -mattr=+ptx87 -verify-machineinstrs | FileCheck %s -check-prefixes=SM903; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | %ptxas-verify -arch=sm_90 %}4; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx88 -verify-machineinstrs | FileCheck %s -check-prefixes=SM1005; RUN: %if ptxas-sm_100 && ptxas-isa-8.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | %ptxas-verify -arch=sm_100 %}6 7; For 256-bit vectors, check that invariant loads from the8; global addrspace are lowered to ld.global.nc.9 10define i8 @ld_global_v32i8(ptr addrspace(1) %ptr) {11; SM90-LABEL: ld_global_v32i8(12; SM90: {13; SM90-NEXT: .reg .b16 %rs<16>;14; SM90-NEXT: .reg .b32 %r<18>;15; SM90-NEXT: .reg .b64 %rd<2>;16; SM90-EMPTY:17; SM90-NEXT: // %bb.0:18; SM90-NEXT: ld.param.b64 %rd1, [ld_global_v32i8_param_0];19; SM90-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];20; SM90-NEXT: prmt.b32 %r5, %r4, 0, 0x7770U;21; SM90-NEXT: cvt.u16.u32 %rs1, %r5;22; SM90-NEXT: prmt.b32 %r6, %r3, 0, 0x7770U;23; SM90-NEXT: cvt.u16.u32 %rs2, %r6;24; SM90-NEXT: prmt.b32 %r7, %r2, 0, 0x7770U;25; SM90-NEXT: cvt.u16.u32 %rs3, %r7;26; SM90-NEXT: prmt.b32 %r8, %r1, 0, 0x7770U;27; SM90-NEXT: cvt.u16.u32 %rs4, %r8;28; SM90-NEXT: ld.global.nc.v4.b32 {%r9, %r10, %r11, %r12}, [%rd1];29; SM90-NEXT: prmt.b32 %r13, %r12, 0, 0x7770U;30; SM90-NEXT: cvt.u16.u32 %rs5, %r13;31; SM90-NEXT: prmt.b32 %r14, %r11, 0, 0x7770U;32; SM90-NEXT: cvt.u16.u32 %rs6, %r14;33; SM90-NEXT: prmt.b32 %r15, %r10, 0, 0x7770U;34; SM90-NEXT: cvt.u16.u32 %rs7, %r15;35; SM90-NEXT: prmt.b32 %r16, %r9, 0, 0x7770U;36; SM90-NEXT: cvt.u16.u32 %rs8, %r16;37; SM90-NEXT: add.s16 %rs9, %rs8, %rs7;38; SM90-NEXT: add.s16 %rs10, %rs6, %rs5;39; SM90-NEXT: add.s16 %rs11, %rs4, %rs3;40; SM90-NEXT: add.s16 %rs12, %rs2, %rs1;41; SM90-NEXT: add.s16 %rs13, %rs9, %rs10;42; SM90-NEXT: add.s16 %rs14, %rs11, %rs12;43; SM90-NEXT: add.s16 %rs15, %rs13, %rs14;44; SM90-NEXT: cvt.u32.u16 %r17, %rs15;45; SM90-NEXT: st.param.b32 [func_retval0], %r17;46; SM90-NEXT: ret;47;48; SM100-LABEL: ld_global_v32i8(49; SM100: {50; SM100-NEXT: .reg .b16 %rs<16>;51; SM100-NEXT: .reg .b32 %r<18>;52; SM100-NEXT: .reg .b64 %rd<2>;53; SM100-EMPTY:54; SM100-NEXT: // %bb.0:55; SM100-NEXT: ld.param.b64 %rd1, [ld_global_v32i8_param_0];56; SM100-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];57; SM100-NEXT: prmt.b32 %r9, %r8, 0, 0x7770U;58; SM100-NEXT: cvt.u16.u32 %rs1, %r9;59; SM100-NEXT: prmt.b32 %r10, %r7, 0, 0x7770U;60; SM100-NEXT: cvt.u16.u32 %rs2, %r10;61; SM100-NEXT: prmt.b32 %r11, %r6, 0, 0x7770U;62; SM100-NEXT: cvt.u16.u32 %rs3, %r11;63; SM100-NEXT: prmt.b32 %r12, %r5, 0, 0x7770U;64; SM100-NEXT: cvt.u16.u32 %rs4, %r12;65; SM100-NEXT: prmt.b32 %r13, %r4, 0, 0x7770U;66; SM100-NEXT: cvt.u16.u32 %rs5, %r13;67; SM100-NEXT: prmt.b32 %r14, %r3, 0, 0x7770U;68; SM100-NEXT: cvt.u16.u32 %rs6, %r14;69; SM100-NEXT: prmt.b32 %r15, %r2, 0, 0x7770U;70; SM100-NEXT: cvt.u16.u32 %rs7, %r15;71; SM100-NEXT: prmt.b32 %r16, %r1, 0, 0x7770U;72; SM100-NEXT: cvt.u16.u32 %rs8, %r16;73; SM100-NEXT: add.s16 %rs9, %rs8, %rs7;74; SM100-NEXT: add.s16 %rs10, %rs6, %rs5;75; SM100-NEXT: add.s16 %rs11, %rs4, %rs3;76; SM100-NEXT: add.s16 %rs12, %rs2, %rs1;77; SM100-NEXT: add.s16 %rs13, %rs9, %rs10;78; SM100-NEXT: add.s16 %rs14, %rs11, %rs12;79; SM100-NEXT: add.s16 %rs15, %rs13, %rs14;80; SM100-NEXT: cvt.u32.u16 %r17, %rs15;81; SM100-NEXT: st.param.b32 [func_retval0], %r17;82; SM100-NEXT: ret;83 %a = load <32 x i8>, ptr addrspace(1) %ptr, !invariant.load !084 %v1 = extractelement <32 x i8> %a, i32 085 %v2 = extractelement <32 x i8> %a, i32 486 %v3 = extractelement <32 x i8> %a, i32 887 %v4 = extractelement <32 x i8> %a, i32 1288 %v5 = extractelement <32 x i8> %a, i32 1689 %v6 = extractelement <32 x i8> %a, i32 2090 %v7 = extractelement <32 x i8> %a, i32 2491 %v8 = extractelement <32 x i8> %a, i32 2892 %sum1 = add i8 %v1, %v293 %sum2 = add i8 %v3, %v494 %sum3 = add i8 %v5, %v695 %sum4 = add i8 %v7, %v896 %sum5 = add i8 %sum1, %sum297 %sum6 = add i8 %sum3, %sum498 %sum7 = add i8 %sum5, %sum699 ret i8 %sum7100}101 102define i16 @ld_global_v16i16(ptr addrspace(1) %ptr) {103; SM90-LABEL: ld_global_v16i16(104; SM90: {105; SM90-NEXT: .reg .b16 %rs<16>;106; SM90-NEXT: .reg .b32 %r<10>;107; SM90-NEXT: .reg .b64 %rd<2>;108; SM90-EMPTY:109; SM90-NEXT: // %bb.0:110; SM90-NEXT: ld.param.b64 %rd1, [ld_global_v16i16_param_0];111; SM90-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];112; SM90-NEXT: mov.b32 {%rs1, _}, %r4;113; SM90-NEXT: mov.b32 {%rs2, _}, %r3;114; SM90-NEXT: mov.b32 {%rs3, _}, %r2;115; SM90-NEXT: mov.b32 {%rs4, _}, %r1;116; SM90-NEXT: ld.global.nc.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];117; SM90-NEXT: mov.b32 {%rs5, _}, %r8;118; SM90-NEXT: mov.b32 {%rs6, _}, %r7;119; SM90-NEXT: mov.b32 {%rs7, _}, %r6;120; SM90-NEXT: mov.b32 {%rs8, _}, %r5;121; SM90-NEXT: add.s16 %rs9, %rs8, %rs7;122; SM90-NEXT: add.s16 %rs10, %rs6, %rs5;123; SM90-NEXT: add.s16 %rs11, %rs4, %rs3;124; SM90-NEXT: add.s16 %rs12, %rs2, %rs1;125; SM90-NEXT: add.s16 %rs13, %rs9, %rs10;126; SM90-NEXT: add.s16 %rs14, %rs11, %rs12;127; SM90-NEXT: add.s16 %rs15, %rs13, %rs14;128; SM90-NEXT: cvt.u32.u16 %r9, %rs15;129; SM90-NEXT: st.param.b32 [func_retval0], %r9;130; SM90-NEXT: ret;131;132; SM100-LABEL: ld_global_v16i16(133; SM100: {134; SM100-NEXT: .reg .b16 %rs<16>;135; SM100-NEXT: .reg .b32 %r<10>;136; SM100-NEXT: .reg .b64 %rd<2>;137; SM100-EMPTY:138; SM100-NEXT: // %bb.0:139; SM100-NEXT: ld.param.b64 %rd1, [ld_global_v16i16_param_0];140; SM100-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];141; SM100-NEXT: mov.b32 {%rs1, _}, %r8;142; SM100-NEXT: mov.b32 {%rs2, _}, %r7;143; SM100-NEXT: mov.b32 {%rs3, _}, %r6;144; SM100-NEXT: mov.b32 {%rs4, _}, %r5;145; SM100-NEXT: mov.b32 {%rs5, _}, %r4;146; SM100-NEXT: mov.b32 {%rs6, _}, %r3;147; SM100-NEXT: mov.b32 {%rs7, _}, %r2;148; SM100-NEXT: mov.b32 {%rs8, _}, %r1;149; SM100-NEXT: add.s16 %rs9, %rs8, %rs7;150; SM100-NEXT: add.s16 %rs10, %rs6, %rs5;151; SM100-NEXT: add.s16 %rs11, %rs4, %rs3;152; SM100-NEXT: add.s16 %rs12, %rs2, %rs1;153; SM100-NEXT: add.s16 %rs13, %rs9, %rs10;154; SM100-NEXT: add.s16 %rs14, %rs11, %rs12;155; SM100-NEXT: add.s16 %rs15, %rs13, %rs14;156; SM100-NEXT: cvt.u32.u16 %r9, %rs15;157; SM100-NEXT: st.param.b32 [func_retval0], %r9;158; SM100-NEXT: ret;159 %a = load <16 x i16>, ptr addrspace(1) %ptr, !invariant.load !0160 %v1 = extractelement <16 x i16> %a, i32 0161 %v2 = extractelement <16 x i16> %a, i32 2162 %v3 = extractelement <16 x i16> %a, i32 4163 %v4 = extractelement <16 x i16> %a, i32 6164 %v5 = extractelement <16 x i16> %a, i32 8165 %v6 = extractelement <16 x i16> %a, i32 10166 %v7 = extractelement <16 x i16> %a, i32 12167 %v8 = extractelement <16 x i16> %a, i32 14168 %sum1 = add i16 %v1, %v2169 %sum2 = add i16 %v3, %v4170 %sum3 = add i16 %v5, %v6171 %sum4 = add i16 %v7, %v8172 %sum5 = add i16 %sum1, %sum2173 %sum6 = add i16 %sum3, %sum4174 %sum7 = add i16 %sum5, %sum6175 ret i16 %sum7176}177 178define half @ld_global_v16f16(ptr addrspace(1) %ptr) {179; SM90-LABEL: ld_global_v16f16(180; SM90: {181; SM90-NEXT: .reg .b16 %rs<16>;182; SM90-NEXT: .reg .b32 %r<9>;183; SM90-NEXT: .reg .b64 %rd<2>;184; SM90-EMPTY:185; SM90-NEXT: // %bb.0:186; SM90-NEXT: ld.param.b64 %rd1, [ld_global_v16f16_param_0];187; SM90-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];188; SM90-NEXT: mov.b32 {%rs1, _}, %r4;189; SM90-NEXT: mov.b32 {%rs2, _}, %r3;190; SM90-NEXT: mov.b32 {%rs3, _}, %r2;191; SM90-NEXT: mov.b32 {%rs4, _}, %r1;192; SM90-NEXT: ld.global.nc.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];193; SM90-NEXT: mov.b32 {%rs5, _}, %r8;194; SM90-NEXT: mov.b32 {%rs6, _}, %r7;195; SM90-NEXT: mov.b32 {%rs7, _}, %r6;196; SM90-NEXT: mov.b32 {%rs8, _}, %r5;197; SM90-NEXT: add.rn.f16 %rs9, %rs8, %rs7;198; SM90-NEXT: add.rn.f16 %rs10, %rs6, %rs5;199; SM90-NEXT: add.rn.f16 %rs11, %rs4, %rs3;200; SM90-NEXT: add.rn.f16 %rs12, %rs2, %rs1;201; SM90-NEXT: add.rn.f16 %rs13, %rs9, %rs10;202; SM90-NEXT: add.rn.f16 %rs14, %rs11, %rs12;203; SM90-NEXT: add.rn.f16 %rs15, %rs13, %rs14;204; SM90-NEXT: st.param.b16 [func_retval0], %rs15;205; SM90-NEXT: ret;206;207; SM100-LABEL: ld_global_v16f16(208; SM100: {209; SM100-NEXT: .reg .b16 %rs<16>;210; SM100-NEXT: .reg .b32 %r<9>;211; SM100-NEXT: .reg .b64 %rd<2>;212; SM100-EMPTY:213; SM100-NEXT: // %bb.0:214; SM100-NEXT: ld.param.b64 %rd1, [ld_global_v16f16_param_0];215; SM100-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];216; SM100-NEXT: mov.b32 {%rs1, _}, %r8;217; SM100-NEXT: mov.b32 {%rs2, _}, %r7;218; SM100-NEXT: mov.b32 {%rs3, _}, %r6;219; SM100-NEXT: mov.b32 {%rs4, _}, %r5;220; SM100-NEXT: mov.b32 {%rs5, _}, %r4;221; SM100-NEXT: mov.b32 {%rs6, _}, %r3;222; SM100-NEXT: mov.b32 {%rs7, _}, %r2;223; SM100-NEXT: mov.b32 {%rs8, _}, %r1;224; SM100-NEXT: add.rn.f16 %rs9, %rs8, %rs7;225; SM100-NEXT: add.rn.f16 %rs10, %rs6, %rs5;226; SM100-NEXT: add.rn.f16 %rs11, %rs4, %rs3;227; SM100-NEXT: add.rn.f16 %rs12, %rs2, %rs1;228; SM100-NEXT: add.rn.f16 %rs13, %rs9, %rs10;229; SM100-NEXT: add.rn.f16 %rs14, %rs11, %rs12;230; SM100-NEXT: add.rn.f16 %rs15, %rs13, %rs14;231; SM100-NEXT: st.param.b16 [func_retval0], %rs15;232; SM100-NEXT: ret;233 %a = load <16 x half>, ptr addrspace(1) %ptr, !invariant.load !0234 %v1 = extractelement <16 x half> %a, i32 0235 %v2 = extractelement <16 x half> %a, i32 2236 %v3 = extractelement <16 x half> %a, i32 4237 %v4 = extractelement <16 x half> %a, i32 6238 %v5 = extractelement <16 x half> %a, i32 8239 %v6 = extractelement <16 x half> %a, i32 10240 %v7 = extractelement <16 x half> %a, i32 12241 %v8 = extractelement <16 x half> %a, i32 14242 %sum1 = fadd half %v1, %v2243 %sum2 = fadd half %v3, %v4244 %sum3 = fadd half %v5, %v6245 %sum4 = fadd half %v7, %v8246 %sum5 = fadd half %sum1, %sum2247 %sum6 = fadd half %sum3, %sum4248 %sum7 = fadd half %sum5, %sum6249 ret half %sum7250}251 252define bfloat @ld_global_v16bf16(ptr addrspace(1) %ptr) {253; SM90-LABEL: ld_global_v16bf16(254; SM90: {255; SM90-NEXT: .reg .b16 %rs<16>;256; SM90-NEXT: .reg .b32 %r<9>;257; SM90-NEXT: .reg .b64 %rd<2>;258; SM90-EMPTY:259; SM90-NEXT: // %bb.0:260; SM90-NEXT: ld.param.b64 %rd1, [ld_global_v16bf16_param_0];261; SM90-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];262; SM90-NEXT: mov.b32 {%rs1, _}, %r4;263; SM90-NEXT: mov.b32 {%rs2, _}, %r3;264; SM90-NEXT: mov.b32 {%rs3, _}, %r2;265; SM90-NEXT: mov.b32 {%rs4, _}, %r1;266; SM90-NEXT: ld.global.nc.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];267; SM90-NEXT: mov.b32 {%rs5, _}, %r8;268; SM90-NEXT: mov.b32 {%rs6, _}, %r7;269; SM90-NEXT: mov.b32 {%rs7, _}, %r6;270; SM90-NEXT: mov.b32 {%rs8, _}, %r5;271; SM90-NEXT: add.rn.bf16 %rs9, %rs8, %rs7;272; SM90-NEXT: add.rn.bf16 %rs10, %rs6, %rs5;273; SM90-NEXT: add.rn.bf16 %rs11, %rs4, %rs3;274; SM90-NEXT: add.rn.bf16 %rs12, %rs2, %rs1;275; SM90-NEXT: add.rn.bf16 %rs13, %rs9, %rs10;276; SM90-NEXT: add.rn.bf16 %rs14, %rs11, %rs12;277; SM90-NEXT: add.rn.bf16 %rs15, %rs13, %rs14;278; SM90-NEXT: st.param.b16 [func_retval0], %rs15;279; SM90-NEXT: ret;280;281; SM100-LABEL: ld_global_v16bf16(282; SM100: {283; SM100-NEXT: .reg .b16 %rs<16>;284; SM100-NEXT: .reg .b32 %r<9>;285; SM100-NEXT: .reg .b64 %rd<2>;286; SM100-EMPTY:287; SM100-NEXT: // %bb.0:288; SM100-NEXT: ld.param.b64 %rd1, [ld_global_v16bf16_param_0];289; SM100-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];290; SM100-NEXT: mov.b32 {%rs1, _}, %r8;291; SM100-NEXT: mov.b32 {%rs2, _}, %r7;292; SM100-NEXT: mov.b32 {%rs3, _}, %r6;293; SM100-NEXT: mov.b32 {%rs4, _}, %r5;294; SM100-NEXT: mov.b32 {%rs5, _}, %r4;295; SM100-NEXT: mov.b32 {%rs6, _}, %r3;296; SM100-NEXT: mov.b32 {%rs7, _}, %r2;297; SM100-NEXT: mov.b32 {%rs8, _}, %r1;298; SM100-NEXT: add.rn.bf16 %rs9, %rs8, %rs7;299; SM100-NEXT: add.rn.bf16 %rs10, %rs6, %rs5;300; SM100-NEXT: add.rn.bf16 %rs11, %rs4, %rs3;301; SM100-NEXT: add.rn.bf16 %rs12, %rs2, %rs1;302; SM100-NEXT: add.rn.bf16 %rs13, %rs9, %rs10;303; SM100-NEXT: add.rn.bf16 %rs14, %rs11, %rs12;304; SM100-NEXT: add.rn.bf16 %rs15, %rs13, %rs14;305; SM100-NEXT: st.param.b16 [func_retval0], %rs15;306; SM100-NEXT: ret;307 %a = load <16 x bfloat>, ptr addrspace(1) %ptr, !invariant.load !0308 %v1 = extractelement <16 x bfloat> %a, i32 0309 %v2 = extractelement <16 x bfloat> %a, i32 2310 %v3 = extractelement <16 x bfloat> %a, i32 4311 %v4 = extractelement <16 x bfloat> %a, i32 6312 %v5 = extractelement <16 x bfloat> %a, i32 8313 %v6 = extractelement <16 x bfloat> %a, i32 10314 %v7 = extractelement <16 x bfloat> %a, i32 12315 %v8 = extractelement <16 x bfloat> %a, i32 14316 %sum1 = fadd bfloat %v1, %v2317 %sum2 = fadd bfloat %v3, %v4318 %sum3 = fadd bfloat %v5, %v6319 %sum4 = fadd bfloat %v7, %v8320 %sum5 = fadd bfloat %sum1, %sum2321 %sum6 = fadd bfloat %sum3, %sum4322 %sum7 = fadd bfloat %sum5, %sum6323 ret bfloat %sum7324}325 326define i32 @ld_global_v8i32(ptr addrspace(1) %ptr) {327; SM90-LABEL: ld_global_v8i32(328; SM90: {329; SM90-NEXT: .reg .b32 %r<16>;330; SM90-NEXT: .reg .b64 %rd<2>;331; SM90-EMPTY:332; SM90-NEXT: // %bb.0:333; SM90-NEXT: ld.param.b64 %rd1, [ld_global_v8i32_param_0];334; SM90-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];335; SM90-NEXT: ld.global.nc.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];336; SM90-NEXT: add.s32 %r9, %r5, %r6;337; SM90-NEXT: add.s32 %r10, %r7, %r8;338; SM90-NEXT: add.s32 %r11, %r1, %r2;339; SM90-NEXT: add.s32 %r12, %r3, %r4;340; SM90-NEXT: add.s32 %r13, %r9, %r10;341; SM90-NEXT: add.s32 %r14, %r11, %r12;342; SM90-NEXT: add.s32 %r15, %r13, %r14;343; SM90-NEXT: st.param.b32 [func_retval0], %r15;344; SM90-NEXT: ret;345;346; SM100-LABEL: ld_global_v8i32(347; SM100: {348; SM100-NEXT: .reg .b32 %r<16>;349; SM100-NEXT: .reg .b64 %rd<2>;350; SM100-EMPTY:351; SM100-NEXT: // %bb.0:352; SM100-NEXT: ld.param.b64 %rd1, [ld_global_v8i32_param_0];353; SM100-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];354; SM100-NEXT: add.s32 %r9, %r1, %r2;355; SM100-NEXT: add.s32 %r10, %r3, %r4;356; SM100-NEXT: add.s32 %r11, %r5, %r6;357; SM100-NEXT: add.s32 %r12, %r7, %r8;358; SM100-NEXT: add.s32 %r13, %r9, %r10;359; SM100-NEXT: add.s32 %r14, %r11, %r12;360; SM100-NEXT: add.s32 %r15, %r13, %r14;361; SM100-NEXT: st.param.b32 [func_retval0], %r15;362; SM100-NEXT: ret;363 %a = load <8 x i32>, ptr addrspace(1) %ptr, !invariant.load !0364 %v1 = extractelement <8 x i32> %a, i32 0365 %v2 = extractelement <8 x i32> %a, i32 1366 %v3 = extractelement <8 x i32> %a, i32 2367 %v4 = extractelement <8 x i32> %a, i32 3368 %v5 = extractelement <8 x i32> %a, i32 4369 %v6 = extractelement <8 x i32> %a, i32 5370 %v7 = extractelement <8 x i32> %a, i32 6371 %v8 = extractelement <8 x i32> %a, i32 7372 %sum1 = add i32 %v1, %v2373 %sum2 = add i32 %v3, %v4374 %sum3 = add i32 %v5, %v6375 %sum4 = add i32 %v7, %v8376 %sum5 = add i32 %sum1, %sum2377 %sum6 = add i32 %sum3, %sum4378 %sum7 = add i32 %sum5, %sum6379 380 ret i32 %sum7381}382 383define float @ld_global_v8f32(ptr addrspace(1) %ptr) {384; SM90-LABEL: ld_global_v8f32(385; SM90: {386; SM90-NEXT: .reg .b32 %r<16>;387; SM90-NEXT: .reg .b64 %rd<2>;388; SM90-EMPTY:389; SM90-NEXT: // %bb.0:390; SM90-NEXT: ld.param.b64 %rd1, [ld_global_v8f32_param_0];391; SM90-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];392; SM90-NEXT: ld.global.nc.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];393; SM90-NEXT: add.rn.f32 %r9, %r5, %r6;394; SM90-NEXT: add.rn.f32 %r10, %r7, %r8;395; SM90-NEXT: add.rn.f32 %r11, %r1, %r2;396; SM90-NEXT: add.rn.f32 %r12, %r3, %r4;397; SM90-NEXT: add.rn.f32 %r13, %r9, %r10;398; SM90-NEXT: add.rn.f32 %r14, %r11, %r12;399; SM90-NEXT: add.rn.f32 %r15, %r13, %r14;400; SM90-NEXT: st.param.b32 [func_retval0], %r15;401; SM90-NEXT: ret;402;403; SM100-LABEL: ld_global_v8f32(404; SM100: {405; SM100-NEXT: .reg .b32 %r<16>;406; SM100-NEXT: .reg .b64 %rd<2>;407; SM100-EMPTY:408; SM100-NEXT: // %bb.0:409; SM100-NEXT: ld.param.b64 %rd1, [ld_global_v8f32_param_0];410; SM100-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];411; SM100-NEXT: add.rn.f32 %r9, %r1, %r2;412; SM100-NEXT: add.rn.f32 %r10, %r3, %r4;413; SM100-NEXT: add.rn.f32 %r11, %r5, %r6;414; SM100-NEXT: add.rn.f32 %r12, %r7, %r8;415; SM100-NEXT: add.rn.f32 %r13, %r9, %r10;416; SM100-NEXT: add.rn.f32 %r14, %r11, %r12;417; SM100-NEXT: add.rn.f32 %r15, %r13, %r14;418; SM100-NEXT: st.param.b32 [func_retval0], %r15;419; SM100-NEXT: ret;420 %a = load <8 x float>, ptr addrspace(1) %ptr, !invariant.load !0421 %v1 = extractelement <8 x float> %a, i32 0422 %v2 = extractelement <8 x float> %a, i32 1423 %v3 = extractelement <8 x float> %a, i32 2424 %v4 = extractelement <8 x float> %a, i32 3425 %v5 = extractelement <8 x float> %a, i32 4426 %v6 = extractelement <8 x float> %a, i32 5427 %v7 = extractelement <8 x float> %a, i32 6428 %v8 = extractelement <8 x float> %a, i32 7429 %sum1 = fadd float %v1, %v2430 %sum2 = fadd float %v3, %v4431 %sum3 = fadd float %v5, %v6432 %sum4 = fadd float %v7, %v8433 %sum5 = fadd float %sum1, %sum2434 %sum6 = fadd float %sum3, %sum4435 %sum7 = fadd float %sum5, %sum6436 437 ret float %sum7438}439 440define i64 @ld_global_v4i64(ptr addrspace(1) %ptr) {441; SM90-LABEL: ld_global_v4i64(442; SM90: {443; SM90-NEXT: .reg .b64 %rd<9>;444; SM90-EMPTY:445; SM90-NEXT: // %bb.0:446; SM90-NEXT: ld.param.b64 %rd1, [ld_global_v4i64_param_0];447; SM90-NEXT: ld.global.nc.v2.b64 {%rd2, %rd3}, [%rd1+16];448; SM90-NEXT: ld.global.nc.v2.b64 {%rd4, %rd5}, [%rd1];449; SM90-NEXT: add.s64 %rd6, %rd4, %rd5;450; SM90-NEXT: add.s64 %rd7, %rd2, %rd3;451; SM90-NEXT: add.s64 %rd8, %rd6, %rd7;452; SM90-NEXT: st.param.b64 [func_retval0], %rd8;453; SM90-NEXT: ret;454;455; SM100-LABEL: ld_global_v4i64(456; SM100: {457; SM100-NEXT: .reg .b64 %rd<9>;458; SM100-EMPTY:459; SM100-NEXT: // %bb.0:460; SM100-NEXT: ld.param.b64 %rd1, [ld_global_v4i64_param_0];461; SM100-NEXT: ld.global.nc.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];462; SM100-NEXT: add.s64 %rd6, %rd2, %rd3;463; SM100-NEXT: add.s64 %rd7, %rd4, %rd5;464; SM100-NEXT: add.s64 %rd8, %rd6, %rd7;465; SM100-NEXT: st.param.b64 [func_retval0], %rd8;466; SM100-NEXT: ret;467 %a = load <4 x i64>, ptr addrspace(1) %ptr, !invariant.load !0468 %v1 = extractelement <4 x i64> %a, i32 0469 %v2 = extractelement <4 x i64> %a, i32 1470 %v3 = extractelement <4 x i64> %a, i32 2471 %v4 = extractelement <4 x i64> %a, i32 3472 %sum1 = add i64 %v1, %v2473 %sum2 = add i64 %v3, %v4474 %sum3 = add i64 %sum1, %sum2475 ret i64 %sum3476}477 478define double @ld_global_v4f64(ptr addrspace(1) %ptr) {479; SM90-LABEL: ld_global_v4f64(480; SM90: {481; SM90-NEXT: .reg .b64 %rd<9>;482; SM90-EMPTY:483; SM90-NEXT: // %bb.0:484; SM90-NEXT: ld.param.b64 %rd1, [ld_global_v4f64_param_0];485; SM90-NEXT: ld.global.nc.v2.b64 {%rd2, %rd3}, [%rd1+16];486; SM90-NEXT: ld.global.nc.v2.b64 {%rd4, %rd5}, [%rd1];487; SM90-NEXT: add.rn.f64 %rd6, %rd4, %rd5;488; SM90-NEXT: add.rn.f64 %rd7, %rd2, %rd3;489; SM90-NEXT: add.rn.f64 %rd8, %rd6, %rd7;490; SM90-NEXT: st.param.b64 [func_retval0], %rd8;491; SM90-NEXT: ret;492;493; SM100-LABEL: ld_global_v4f64(494; SM100: {495; SM100-NEXT: .reg .b64 %rd<9>;496; SM100-EMPTY:497; SM100-NEXT: // %bb.0:498; SM100-NEXT: ld.param.b64 %rd1, [ld_global_v4f64_param_0];499; SM100-NEXT: ld.global.nc.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];500; SM100-NEXT: add.rn.f64 %rd6, %rd2, %rd3;501; SM100-NEXT: add.rn.f64 %rd7, %rd4, %rd5;502; SM100-NEXT: add.rn.f64 %rd8, %rd6, %rd7;503; SM100-NEXT: st.param.b64 [func_retval0], %rd8;504; SM100-NEXT: ret;505 %a = load <4 x double>, ptr addrspace(1) %ptr, !invariant.load !0506 %v1 = extractelement <4 x double> %a, i32 0507 %v2 = extractelement <4 x double> %a, i32 1508 %v3 = extractelement <4 x double> %a, i32 2509 %v4 = extractelement <4 x double> %a, i32 3510 %sum1 = fadd double %v1, %v2511 %sum2 = fadd double %v3, %v4512 %sum3 = fadd double %sum1, %sum2513 ret double %sum3514}515 516!0 = !{}517