544 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | FileCheck %s -check-prefixes=PTX3; RUN: %if ptxas-sm_100 && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | %ptxas-verify -arch=sm_100 %}4 5; In this test, we check that all the addressing modes are lowered correctly6; for 256-bit invariant loads, which get lowered to ld.global.nc7; addr can be any of the following:8; - avar : direct address9; - asi: direct address + offset10; - areg_64: 64-bit register11; - ari_64: 64-bit register + offset12; Since this is a blackwell+ feature,13; and support for 32-bit addressing does not exist after sm_90,14; the "areg" and "ari" 32-bit addressing modes are not tested or supported.15 16; For invariant loads, asi is historically not supported,17; and instead it is selected as move into register, add of offset, and loaded as areg6418 19; Checks 8 types: i8, i16, bfloat, half, i32, i64, float, double20 21; Global is the only address space that currently supports 256-bit loads/stores22 23@globalin = external addrspace(1) global ptr24@globalout = external addrspace(1) global ptr25 26define void @avar_i8() {27; PTX-LABEL: avar_i8(28; PTX: {29; PTX-NEXT: .reg .b32 %r<9>;30; PTX-EMPTY:31; PTX-NEXT: // %bb.0:32; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin];33; PTX-NEXT: st.global.v8.b32 [globalout], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};34; PTX-NEXT: ret;35 %load = load <32 x i8>, ptr addrspace(1) @globalin, !invariant.load !036 store <32 x i8> %load, ptr addrspace(1) @globalout37 ret void38}39 40define void @avar_i16() {41; PTX-LABEL: avar_i16(42; PTX: {43; PTX-NEXT: .reg .b32 %r<9>;44; PTX-EMPTY:45; PTX-NEXT: // %bb.0:46; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin];47; PTX-NEXT: st.global.v8.b32 [globalout], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};48; PTX-NEXT: ret;49 %load = load <16 x i16>, ptr addrspace(1) @globalin, !invariant.load !050 store <16 x i16> %load, ptr addrspace(1) @globalout51 ret void52}53 54define void @avar_half() {55; PTX-LABEL: avar_half(56; PTX: {57; PTX-NEXT: .reg .b32 %r<9>;58; PTX-EMPTY:59; PTX-NEXT: // %bb.0:60; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin];61; PTX-NEXT: st.global.v8.b32 [globalout], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};62; PTX-NEXT: ret;63 %load = load <16 x half>, ptr addrspace(1) @globalin, !invariant.load !064 store <16 x half> %load, ptr addrspace(1) @globalout65 ret void66}67 68define void @avar_bfloat() {69; PTX-LABEL: avar_bfloat(70; PTX: {71; PTX-NEXT: .reg .b32 %r<9>;72; PTX-EMPTY:73; PTX-NEXT: // %bb.0:74; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin];75; PTX-NEXT: st.global.v8.b32 [globalout], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};76; PTX-NEXT: ret;77 %load = load <16 x bfloat>, ptr addrspace(1) @globalin, !invariant.load !078 store <16 x bfloat> %load, ptr addrspace(1) @globalout79 ret void80}81 82define void @avar_i32() {83; PTX-LABEL: avar_i32(84; PTX: {85; PTX-NEXT: .reg .b64 %rd<5>;86; PTX-EMPTY:87; PTX-NEXT: // %bb.0:88; PTX-NEXT: ld.global.nc.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin];89; PTX-NEXT: st.global.v4.b64 [globalout], {%rd1, %rd2, %rd3, %rd4};90; PTX-NEXT: ret;91 %load = load <8 x i32>, ptr addrspace(1) @globalin, !invariant.load !092 store <8 x i32> %load, ptr addrspace(1) @globalout93 ret void94}95 96define void @avar_i64() {97; PTX-LABEL: avar_i64(98; PTX: {99; PTX-NEXT: .reg .b64 %rd<5>;100; PTX-EMPTY:101; PTX-NEXT: // %bb.0:102; PTX-NEXT: ld.global.nc.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin];103; PTX-NEXT: st.global.v4.b64 [globalout], {%rd1, %rd2, %rd3, %rd4};104; PTX-NEXT: ret;105 %load = load <4 x i64>, ptr addrspace(1) @globalin, !invariant.load !0106 store <4 x i64> %load, ptr addrspace(1) @globalout107 ret void108}109 110define void @avar_float() {111; PTX-LABEL: avar_float(112; PTX: {113; PTX-NEXT: .reg .b64 %rd<5>;114; PTX-EMPTY:115; PTX-NEXT: // %bb.0:116; PTX-NEXT: ld.global.nc.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin];117; PTX-NEXT: st.global.v4.b64 [globalout], {%rd1, %rd2, %rd3, %rd4};118; PTX-NEXT: ret;119 %load = load <8 x float>, ptr addrspace(1) @globalin, !invariant.load !0120 store <8 x float> %load, ptr addrspace(1) @globalout121 ret void122}123 124define void @avar_double() {125; PTX-LABEL: avar_double(126; PTX: {127; PTX-NEXT: .reg .b64 %rd<5>;128; PTX-EMPTY:129; PTX-NEXT: // %bb.0:130; PTX-NEXT: ld.global.nc.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin];131; PTX-NEXT: st.global.v4.b64 [globalout], {%rd1, %rd2, %rd3, %rd4};132; PTX-NEXT: ret;133 %load = load <4 x double>, ptr addrspace(1) @globalin, !invariant.load !0134 store <4 x double> %load, ptr addrspace(1) @globalout135 ret void136}137 138define void @asi_i8() {139; PTX-LABEL: asi_i8(140; PTX: {141; PTX-NEXT: .reg .b32 %r<9>;142; PTX-EMPTY:143; PTX-NEXT: // %bb.0:144; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin+32];145; PTX-NEXT: st.global.v8.b32 [globalout+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};146; PTX-NEXT: ret;147 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32148 %load = load <32 x i8>, ptr addrspace(1) %in.offset, !invariant.load !0149 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32150 store <32 x i8> %load, ptr addrspace(1) %out.offset151 ret void152}153 154define void @asi_i16() {155; PTX-LABEL: asi_i16(156; PTX: {157; PTX-NEXT: .reg .b32 %r<9>;158; PTX-EMPTY:159; PTX-NEXT: // %bb.0:160; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin+32];161; PTX-NEXT: st.global.v8.b32 [globalout+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};162; PTX-NEXT: ret;163 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32164 %load = load <16 x i16>, ptr addrspace(1) %in.offset, !invariant.load !0165 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32166 store <16 x i16> %load, ptr addrspace(1) %out.offset167 ret void168}169 170define void @asi_half() {171; PTX-LABEL: asi_half(172; PTX: {173; PTX-NEXT: .reg .b32 %r<9>;174; PTX-EMPTY:175; PTX-NEXT: // %bb.0:176; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin+32];177; PTX-NEXT: st.global.v8.b32 [globalout+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};178; PTX-NEXT: ret;179 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32180 %load = load <16 x half>, ptr addrspace(1) %in.offset, !invariant.load !0181 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32182 store <16 x half> %load, ptr addrspace(1) %out.offset183 ret void184}185 186define void @asi_bfloat() {187; PTX-LABEL: asi_bfloat(188; PTX: {189; PTX-NEXT: .reg .b32 %r<9>;190; PTX-EMPTY:191; PTX-NEXT: // %bb.0:192; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin+32];193; PTX-NEXT: st.global.v8.b32 [globalout+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};194; PTX-NEXT: ret;195 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32196 %load = load <16 x bfloat>, ptr addrspace(1) %in.offset, !invariant.load !0197 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32198 store <16 x bfloat> %load, ptr addrspace(1) %out.offset199 ret void200}201 202define void @asi_i32() {203; PTX-LABEL: asi_i32(204; PTX: {205; PTX-NEXT: .reg .b64 %rd<5>;206; PTX-EMPTY:207; PTX-NEXT: // %bb.0:208; PTX-NEXT: ld.global.nc.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin+32];209; PTX-NEXT: st.global.v4.b64 [globalout+32], {%rd1, %rd2, %rd3, %rd4};210; PTX-NEXT: ret;211 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32212 %load = load <8 x i32>, ptr addrspace(1) %in.offset, !invariant.load !0213 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32214 store <8 x i32> %load, ptr addrspace(1) %out.offset215 ret void216}217 218define void @asi_i64() {219; PTX-LABEL: asi_i64(220; PTX: {221; PTX-NEXT: .reg .b64 %rd<5>;222; PTX-EMPTY:223; PTX-NEXT: // %bb.0:224; PTX-NEXT: ld.global.nc.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin+32];225; PTX-NEXT: st.global.v4.b64 [globalout+32], {%rd1, %rd2, %rd3, %rd4};226; PTX-NEXT: ret;227 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32228 %load = load <4 x i64>, ptr addrspace(1) %in.offset, !invariant.load !0229 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32230 store <4 x i64> %load, ptr addrspace(1) %out.offset231 ret void232}233 234define void @asi_float() {235; PTX-LABEL: asi_float(236; PTX: {237; PTX-NEXT: .reg .b64 %rd<5>;238; PTX-EMPTY:239; PTX-NEXT: // %bb.0:240; PTX-NEXT: ld.global.nc.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin+32];241; PTX-NEXT: st.global.v4.b64 [globalout+32], {%rd1, %rd2, %rd3, %rd4};242; PTX-NEXT: ret;243 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32244 %load = load <8 x float>, ptr addrspace(1) %in.offset, !invariant.load !0245 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32246 store <8 x float> %load, ptr addrspace(1) %out.offset247 ret void248}249 250define void @asi_double() {251; PTX-LABEL: asi_double(252; PTX: {253; PTX-NEXT: .reg .b64 %rd<5>;254; PTX-EMPTY:255; PTX-NEXT: // %bb.0:256; PTX-NEXT: ld.global.nc.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin+32];257; PTX-NEXT: st.global.v4.b64 [globalout+32], {%rd1, %rd2, %rd3, %rd4};258; PTX-NEXT: ret;259 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32260 %load = load <4 x double>, ptr addrspace(1) %in.offset, !invariant.load !0261 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32262 store <4 x double> %load, ptr addrspace(1) %out.offset263 ret void264}265 266define void @areg_64_i8(ptr addrspace(1) %in, ptr addrspace(1) %out) {267; PTX-LABEL: areg_64_i8(268; PTX: {269; PTX-NEXT: .reg .b32 %r<9>;270; PTX-NEXT: .reg .b64 %rd<3>;271; PTX-EMPTY:272; PTX-NEXT: // %bb.0:273; PTX-NEXT: ld.param.b64 %rd1, [areg_64_i8_param_0];274; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];275; PTX-NEXT: ld.param.b64 %rd2, [areg_64_i8_param_1];276; PTX-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};277; PTX-NEXT: ret;278 %load = load <32 x i8>, ptr addrspace(1) %in, !invariant.load !0279 store <32 x i8> %load, ptr addrspace(1) %out280 ret void281}282define void @areg_64_i16(ptr addrspace(1) %in, ptr addrspace(1) %out) {283; PTX-LABEL: areg_64_i16(284; PTX: {285; PTX-NEXT: .reg .b32 %r<9>;286; PTX-NEXT: .reg .b64 %rd<3>;287; PTX-EMPTY:288; PTX-NEXT: // %bb.0:289; PTX-NEXT: ld.param.b64 %rd1, [areg_64_i16_param_0];290; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];291; PTX-NEXT: ld.param.b64 %rd2, [areg_64_i16_param_1];292; PTX-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};293; PTX-NEXT: ret;294 %load = load <16 x i16>, ptr addrspace(1) %in, !invariant.load !0295 store <16 x i16> %load, ptr addrspace(1) %out296 ret void297}298define void @areg_64_half(ptr addrspace(1) %in, ptr addrspace(1) %out) {299; PTX-LABEL: areg_64_half(300; PTX: {301; PTX-NEXT: .reg .b32 %r<9>;302; PTX-NEXT: .reg .b64 %rd<3>;303; PTX-EMPTY:304; PTX-NEXT: // %bb.0:305; PTX-NEXT: ld.param.b64 %rd1, [areg_64_half_param_0];306; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];307; PTX-NEXT: ld.param.b64 %rd2, [areg_64_half_param_1];308; PTX-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};309; PTX-NEXT: ret;310 %load = load <16 x half>, ptr addrspace(1) %in, !invariant.load !0311 store <16 x half> %load, ptr addrspace(1) %out312 ret void313}314define void @areg_64_bfloat(ptr addrspace(1) %in, ptr addrspace(1) %out) {315; PTX-LABEL: areg_64_bfloat(316; PTX: {317; PTX-NEXT: .reg .b32 %r<9>;318; PTX-NEXT: .reg .b64 %rd<3>;319; PTX-EMPTY:320; PTX-NEXT: // %bb.0:321; PTX-NEXT: ld.param.b64 %rd1, [areg_64_bfloat_param_0];322; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];323; PTX-NEXT: ld.param.b64 %rd2, [areg_64_bfloat_param_1];324; PTX-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};325; PTX-NEXT: ret;326 %load = load <16 x bfloat>, ptr addrspace(1) %in, !invariant.load !0327 store <16 x bfloat> %load, ptr addrspace(1) %out328 ret void329}330 331define void @areg_64_i32(ptr addrspace(1) %in, ptr addrspace(1) %out) {332; PTX-LABEL: areg_64_i32(333; PTX: {334; PTX-NEXT: .reg .b64 %rd<7>;335; PTX-EMPTY:336; PTX-NEXT: // %bb.0:337; PTX-NEXT: ld.param.b64 %rd1, [areg_64_i32_param_0];338; PTX-NEXT: ld.global.nc.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];339; PTX-NEXT: ld.param.b64 %rd6, [areg_64_i32_param_1];340; PTX-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};341; PTX-NEXT: ret;342 %load = load <8 x i32>, ptr addrspace(1) %in, !invariant.load !0343 store <8 x i32> %load, ptr addrspace(1) %out344 ret void345}346 347define void @areg_64_i64(ptr addrspace(1) %in, ptr addrspace(1) %out) {348; PTX-LABEL: areg_64_i64(349; PTX: {350; PTX-NEXT: .reg .b64 %rd<7>;351; PTX-EMPTY:352; PTX-NEXT: // %bb.0:353; PTX-NEXT: ld.param.b64 %rd1, [areg_64_i64_param_0];354; PTX-NEXT: ld.global.nc.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];355; PTX-NEXT: ld.param.b64 %rd6, [areg_64_i64_param_1];356; PTX-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};357; PTX-NEXT: ret;358 %load = load <4 x i64>, ptr addrspace(1) %in, !invariant.load !0359 store <4 x i64> %load, ptr addrspace(1) %out360 ret void361}362 363define void @areg_64_float(ptr addrspace(1) %in, ptr addrspace(1) %out) {364; PTX-LABEL: areg_64_float(365; PTX: {366; PTX-NEXT: .reg .b64 %rd<7>;367; PTX-EMPTY:368; PTX-NEXT: // %bb.0:369; PTX-NEXT: ld.param.b64 %rd1, [areg_64_float_param_0];370; PTX-NEXT: ld.global.nc.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];371; PTX-NEXT: ld.param.b64 %rd6, [areg_64_float_param_1];372; PTX-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};373; PTX-NEXT: ret;374 %load = load <8 x float>, ptr addrspace(1) %in, !invariant.load !0375 store <8 x float> %load, ptr addrspace(1) %out376 ret void377}378 379define void @areg_64_double(ptr addrspace(1) %in, ptr addrspace(1) %out) {380; PTX-LABEL: areg_64_double(381; PTX: {382; PTX-NEXT: .reg .b64 %rd<7>;383; PTX-EMPTY:384; PTX-NEXT: // %bb.0:385; PTX-NEXT: ld.param.b64 %rd1, [areg_64_double_param_0];386; PTX-NEXT: ld.global.nc.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];387; PTX-NEXT: ld.param.b64 %rd6, [areg_64_double_param_1];388; PTX-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};389; PTX-NEXT: ret;390 %load = load <4 x double>, ptr addrspace(1) %in, !invariant.load !0391 store <4 x double> %load, ptr addrspace(1) %out392 ret void393}394 395define void @ari_64_i8(ptr addrspace(1) %in, ptr addrspace(1) %out) {396; PTX-LABEL: ari_64_i8(397; PTX: {398; PTX-NEXT: .reg .b32 %r<9>;399; PTX-NEXT: .reg .b64 %rd<3>;400; PTX-EMPTY:401; PTX-NEXT: // %bb.0:402; PTX-NEXT: ld.param.b64 %rd1, [ari_64_i8_param_0];403; PTX-NEXT: ld.param.b64 %rd2, [ari_64_i8_param_1];404; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1+32];405; PTX-NEXT: st.global.v8.b32 [%rd2+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};406; PTX-NEXT: ret;407 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32408 %load = load <32 x i8>, ptr addrspace(1) %in.offset, !invariant.load !0409 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32410 store <32 x i8> %load, ptr addrspace(1) %out.offset411 ret void412}413 414define void @ari_64_i16(ptr addrspace(1) %in, ptr addrspace(1) %out) {415; PTX-LABEL: ari_64_i16(416; PTX: {417; PTX-NEXT: .reg .b32 %r<9>;418; PTX-NEXT: .reg .b64 %rd<3>;419; PTX-EMPTY:420; PTX-NEXT: // %bb.0:421; PTX-NEXT: ld.param.b64 %rd1, [ari_64_i16_param_0];422; PTX-NEXT: ld.param.b64 %rd2, [ari_64_i16_param_1];423; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1+32];424; PTX-NEXT: st.global.v8.b32 [%rd2+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};425; PTX-NEXT: ret;426 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32427 %load = load <16 x i16>, ptr addrspace(1) %in.offset, !invariant.load !0428 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32429 store <16 x i16> %load, ptr addrspace(1) %out.offset430 ret void431}432 433define void @ari_64_half(ptr addrspace(1) %in, ptr addrspace(1) %out) {434; PTX-LABEL: ari_64_half(435; PTX: {436; PTX-NEXT: .reg .b32 %r<9>;437; PTX-NEXT: .reg .b64 %rd<3>;438; PTX-EMPTY:439; PTX-NEXT: // %bb.0:440; PTX-NEXT: ld.param.b64 %rd1, [ari_64_half_param_0];441; PTX-NEXT: ld.param.b64 %rd2, [ari_64_half_param_1];442; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1+32];443; PTX-NEXT: st.global.v8.b32 [%rd2+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};444; PTX-NEXT: ret;445 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32446 %load = load <16 x half>, ptr addrspace(1) %in.offset, !invariant.load !0447 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32448 store <16 x half> %load, ptr addrspace(1) %out.offset449 ret void450}451 452define void @ari_64_bfloat(ptr addrspace(1) %in, ptr addrspace(1) %out) {453; PTX-LABEL: ari_64_bfloat(454; PTX: {455; PTX-NEXT: .reg .b32 %r<9>;456; PTX-NEXT: .reg .b64 %rd<3>;457; PTX-EMPTY:458; PTX-NEXT: // %bb.0:459; PTX-NEXT: ld.param.b64 %rd1, [ari_64_bfloat_param_0];460; PTX-NEXT: ld.param.b64 %rd2, [ari_64_bfloat_param_1];461; PTX-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1+32];462; PTX-NEXT: st.global.v8.b32 [%rd2+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};463; PTX-NEXT: ret;464 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32465 %load = load <16 x bfloat>, ptr addrspace(1) %in.offset, !invariant.load !0466 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32467 store <16 x bfloat> %load, ptr addrspace(1) %out.offset468 ret void469}470 471define void @ari_64_i32(ptr addrspace(1) %in, ptr addrspace(1) %out) {472; PTX-LABEL: ari_64_i32(473; PTX: {474; PTX-NEXT: .reg .b64 %rd<7>;475; PTX-EMPTY:476; PTX-NEXT: // %bb.0:477; PTX-NEXT: ld.param.b64 %rd1, [ari_64_i32_param_0];478; PTX-NEXT: ld.param.b64 %rd2, [ari_64_i32_param_1];479; PTX-NEXT: ld.global.nc.v4.b64 {%rd3, %rd4, %rd5, %rd6}, [%rd1+32];480; PTX-NEXT: st.global.v4.b64 [%rd2+32], {%rd3, %rd4, %rd5, %rd6};481; PTX-NEXT: ret;482 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32483 %load = load <8 x i32>, ptr addrspace(1) %in.offset, !invariant.load !0484 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32485 store <8 x i32> %load, ptr addrspace(1) %out.offset486 ret void487}488 489define void @ari_64_i64(ptr addrspace(1) %in, ptr addrspace(1) %out) {490; PTX-LABEL: ari_64_i64(491; PTX: {492; PTX-NEXT: .reg .b64 %rd<7>;493; PTX-EMPTY:494; PTX-NEXT: // %bb.0:495; PTX-NEXT: ld.param.b64 %rd1, [ari_64_i64_param_0];496; PTX-NEXT: ld.param.b64 %rd2, [ari_64_i64_param_1];497; PTX-NEXT: ld.global.nc.v4.b64 {%rd3, %rd4, %rd5, %rd6}, [%rd1+32];498; PTX-NEXT: st.global.v4.b64 [%rd2+32], {%rd3, %rd4, %rd5, %rd6};499; PTX-NEXT: ret;500 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32501 %load = load <4 x i64>, ptr addrspace(1) %in.offset, !invariant.load !0502 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32503 store <4 x i64> %load, ptr addrspace(1) %out.offset504 ret void505}506 507define void @ari_64_float(ptr addrspace(1) %in, ptr addrspace(1) %out) {508; PTX-LABEL: ari_64_float(509; PTX: {510; PTX-NEXT: .reg .b64 %rd<7>;511; PTX-EMPTY:512; PTX-NEXT: // %bb.0:513; PTX-NEXT: ld.param.b64 %rd1, [ari_64_float_param_0];514; PTX-NEXT: ld.param.b64 %rd2, [ari_64_float_param_1];515; PTX-NEXT: ld.global.nc.v4.b64 {%rd3, %rd4, %rd5, %rd6}, [%rd1+32];516; PTX-NEXT: st.global.v4.b64 [%rd2+32], {%rd3, %rd4, %rd5, %rd6};517; PTX-NEXT: ret;518 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32519 %load = load <8 x float>, ptr addrspace(1) %in.offset, !invariant.load !0520 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32521 store <8 x float> %load, ptr addrspace(1) %out.offset522 ret void523}524 525define void @ari_64_double(ptr addrspace(1) %in, ptr addrspace(1) %out) {526; PTX-LABEL: ari_64_double(527; PTX: {528; PTX-NEXT: .reg .b64 %rd<7>;529; PTX-EMPTY:530; PTX-NEXT: // %bb.0:531; PTX-NEXT: ld.param.b64 %rd1, [ari_64_double_param_0];532; PTX-NEXT: ld.param.b64 %rd2, [ari_64_double_param_1];533; PTX-NEXT: ld.global.nc.v4.b64 {%rd3, %rd4, %rd5, %rd6}, [%rd1+32];534; PTX-NEXT: st.global.v4.b64 [%rd2+32], {%rd3, %rd4, %rd5, %rd6};535; PTX-NEXT: ret;536 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32537 %load = load <4 x double>, ptr addrspace(1) %in.offset, !invariant.load !0538 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32539 store <4 x double> %load, ptr addrspace(1) %out.offset540 ret void541}542 543!0 = !{}544