538 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 correctly,6; addr can be any of the following:7; - avar : direct address8; - asi: direct address + offset9; - areg_64: 64-bit register10; - ari_64: 64-bit register + offset11; Since this is a blackwell+ feature,12; and support for 32-bit addressing does not exist after sm_90,13; the "areg" and "ari" 32-bit addressing modes are not tested or supported.14 15; Checks 8 types: i8, i16, bfloat, half, i32, i64, float, double16 17; Global is the only address space that currently supports 256-bit loads/stores18 19@globalin = external addrspace(1) global ptr20@globalout = external addrspace(1) global ptr21 22define void @avar_i8() {23; PTX-LABEL: avar_i8(24; PTX: {25; PTX-NEXT: .reg .b32 %r<9>;26; PTX-EMPTY:27; PTX-NEXT: // %bb.0:28; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin];29; PTX-NEXT: st.global.v8.b32 [globalout], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};30; PTX-NEXT: ret;31 %load = load <32 x i8>, ptr addrspace(1) @globalin32 store <32 x i8> %load, ptr addrspace(1) @globalout33 ret void34}35 36define void @avar_i16() {37; PTX-LABEL: avar_i16(38; PTX: {39; PTX-NEXT: .reg .b32 %r<9>;40; PTX-EMPTY:41; PTX-NEXT: // %bb.0:42; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin];43; PTX-NEXT: st.global.v8.b32 [globalout], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};44; PTX-NEXT: ret;45 %load = load <16 x i16>, ptr addrspace(1) @globalin46 store <16 x i16> %load, ptr addrspace(1) @globalout47 ret void48}49 50define void @avar_half() {51; PTX-LABEL: avar_half(52; PTX: {53; PTX-NEXT: .reg .b32 %r<9>;54; PTX-EMPTY:55; PTX-NEXT: // %bb.0:56; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin];57; PTX-NEXT: st.global.v8.b32 [globalout], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};58; PTX-NEXT: ret;59 %load = load <16 x half>, ptr addrspace(1) @globalin60 store <16 x half> %load, ptr addrspace(1) @globalout61 ret void62}63 64define void @avar_bfloat() {65; PTX-LABEL: avar_bfloat(66; PTX: {67; PTX-NEXT: .reg .b32 %r<9>;68; PTX-EMPTY:69; PTX-NEXT: // %bb.0:70; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin];71; PTX-NEXT: st.global.v8.b32 [globalout], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};72; PTX-NEXT: ret;73 %load = load <16 x bfloat>, ptr addrspace(1) @globalin74 store <16 x bfloat> %load, ptr addrspace(1) @globalout75 ret void76}77 78define void @avar_i32() {79; PTX-LABEL: avar_i32(80; PTX: {81; PTX-NEXT: .reg .b64 %rd<5>;82; PTX-EMPTY:83; PTX-NEXT: // %bb.0:84; PTX-NEXT: ld.global.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin];85; PTX-NEXT: st.global.v4.b64 [globalout], {%rd1, %rd2, %rd3, %rd4};86; PTX-NEXT: ret;87 %load = load <8 x i32>, ptr addrspace(1) @globalin88 store <8 x i32> %load, ptr addrspace(1) @globalout89 ret void90}91 92define void @avar_i64() {93; PTX-LABEL: avar_i64(94; PTX: {95; PTX-NEXT: .reg .b64 %rd<5>;96; PTX-EMPTY:97; PTX-NEXT: // %bb.0:98; PTX-NEXT: ld.global.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin];99; PTX-NEXT: st.global.v4.b64 [globalout], {%rd1, %rd2, %rd3, %rd4};100; PTX-NEXT: ret;101 %load = load <4 x i64>, ptr addrspace(1) @globalin102 store <4 x i64> %load, ptr addrspace(1) @globalout103 ret void104}105 106define void @avar_float() {107; PTX-LABEL: avar_float(108; PTX: {109; PTX-NEXT: .reg .b64 %rd<5>;110; PTX-EMPTY:111; PTX-NEXT: // %bb.0:112; PTX-NEXT: ld.global.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin];113; PTX-NEXT: st.global.v4.b64 [globalout], {%rd1, %rd2, %rd3, %rd4};114; PTX-NEXT: ret;115 %load = load <8 x float>, ptr addrspace(1) @globalin116 store <8 x float> %load, ptr addrspace(1) @globalout117 ret void118}119 120define void @avar_double() {121; PTX-LABEL: avar_double(122; PTX: {123; PTX-NEXT: .reg .b64 %rd<5>;124; PTX-EMPTY:125; PTX-NEXT: // %bb.0:126; PTX-NEXT: ld.global.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin];127; PTX-NEXT: st.global.v4.b64 [globalout], {%rd1, %rd2, %rd3, %rd4};128; PTX-NEXT: ret;129 %load = load <4 x double>, ptr addrspace(1) @globalin130 store <4 x double> %load, ptr addrspace(1) @globalout131 ret void132}133 134define void @asi_i8() {135; PTX-LABEL: asi_i8(136; PTX: {137; PTX-NEXT: .reg .b32 %r<9>;138; PTX-EMPTY:139; PTX-NEXT: // %bb.0:140; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin+32];141; PTX-NEXT: st.global.v8.b32 [globalout+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};142; PTX-NEXT: ret;143 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32144 %load = load <32 x i8>, ptr addrspace(1) %in.offset145 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32146 store <32 x i8> %load, ptr addrspace(1) %out.offset147 ret void148}149 150define void @asi_i16() {151; PTX-LABEL: asi_i16(152; PTX: {153; PTX-NEXT: .reg .b32 %r<9>;154; PTX-EMPTY:155; PTX-NEXT: // %bb.0:156; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin+32];157; PTX-NEXT: st.global.v8.b32 [globalout+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};158; PTX-NEXT: ret;159 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32160 %load = load <16 x i16>, ptr addrspace(1) %in.offset161 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32162 store <16 x i16> %load, ptr addrspace(1) %out.offset163 ret void164}165 166define void @asi_half() {167; PTX-LABEL: asi_half(168; PTX: {169; PTX-NEXT: .reg .b32 %r<9>;170; PTX-EMPTY:171; PTX-NEXT: // %bb.0:172; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin+32];173; PTX-NEXT: st.global.v8.b32 [globalout+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};174; PTX-NEXT: ret;175 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32176 %load = load <16 x half>, ptr addrspace(1) %in.offset177 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32178 store <16 x half> %load, ptr addrspace(1) %out.offset179 ret void180}181 182define void @asi_bfloat() {183; PTX-LABEL: asi_bfloat(184; PTX: {185; PTX-NEXT: .reg .b32 %r<9>;186; PTX-EMPTY:187; PTX-NEXT: // %bb.0:188; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [globalin+32];189; PTX-NEXT: st.global.v8.b32 [globalout+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};190; PTX-NEXT: ret;191 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32192 %load = load <16 x bfloat>, ptr addrspace(1) %in.offset193 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32194 store <16 x bfloat> %load, ptr addrspace(1) %out.offset195 ret void196}197 198define void @asi_i32() {199; PTX-LABEL: asi_i32(200; PTX: {201; PTX-NEXT: .reg .b64 %rd<5>;202; PTX-EMPTY:203; PTX-NEXT: // %bb.0:204; PTX-NEXT: ld.global.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin+32];205; PTX-NEXT: st.global.v4.b64 [globalout+32], {%rd1, %rd2, %rd3, %rd4};206; PTX-NEXT: ret;207 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32208 %load = load <8 x i32>, ptr addrspace(1) %in.offset209 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32210 store <8 x i32> %load, ptr addrspace(1) %out.offset211 ret void212}213 214define void @asi_i64() {215; PTX-LABEL: asi_i64(216; PTX: {217; PTX-NEXT: .reg .b64 %rd<5>;218; PTX-EMPTY:219; PTX-NEXT: // %bb.0:220; PTX-NEXT: ld.global.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin+32];221; PTX-NEXT: st.global.v4.b64 [globalout+32], {%rd1, %rd2, %rd3, %rd4};222; PTX-NEXT: ret;223 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32224 %load = load <4 x i64>, ptr addrspace(1) %in.offset225 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32226 store <4 x i64> %load, ptr addrspace(1) %out.offset227 ret void228}229 230define void @asi_float() {231; PTX-LABEL: asi_float(232; PTX: {233; PTX-NEXT: .reg .b64 %rd<5>;234; PTX-EMPTY:235; PTX-NEXT: // %bb.0:236; PTX-NEXT: ld.global.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin+32];237; PTX-NEXT: st.global.v4.b64 [globalout+32], {%rd1, %rd2, %rd3, %rd4};238; PTX-NEXT: ret;239 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32240 %load = load <8 x float>, ptr addrspace(1) %in.offset241 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32242 store <8 x float> %load, ptr addrspace(1) %out.offset243 ret void244}245 246define void @asi_double() {247; PTX-LABEL: asi_double(248; PTX: {249; PTX-NEXT: .reg .b64 %rd<5>;250; PTX-EMPTY:251; PTX-NEXT: // %bb.0:252; PTX-NEXT: ld.global.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [globalin+32];253; PTX-NEXT: st.global.v4.b64 [globalout+32], {%rd1, %rd2, %rd3, %rd4};254; PTX-NEXT: ret;255 %in.offset = getelementptr inbounds i8, ptr addrspace(1) @globalin, i32 32256 %load = load <4 x double>, ptr addrspace(1) %in.offset257 %out.offset = getelementptr inbounds i8, ptr addrspace(1) @globalout, i32 32258 store <4 x double> %load, ptr addrspace(1) %out.offset259 ret void260}261 262define void @areg_64_i8(ptr addrspace(1) %in, ptr addrspace(1) %out) {263; PTX-LABEL: areg_64_i8(264; PTX: {265; PTX-NEXT: .reg .b32 %r<9>;266; PTX-NEXT: .reg .b64 %rd<3>;267; PTX-EMPTY:268; PTX-NEXT: // %bb.0:269; PTX-NEXT: ld.param.b64 %rd1, [areg_64_i8_param_0];270; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];271; PTX-NEXT: ld.param.b64 %rd2, [areg_64_i8_param_1];272; PTX-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};273; PTX-NEXT: ret;274 %load = load <32 x i8>, ptr addrspace(1) %in275 store <32 x i8> %load, ptr addrspace(1) %out276 ret void277}278define void @areg_64_i16(ptr addrspace(1) %in, ptr addrspace(1) %out) {279; PTX-LABEL: areg_64_i16(280; PTX: {281; PTX-NEXT: .reg .b32 %r<9>;282; PTX-NEXT: .reg .b64 %rd<3>;283; PTX-EMPTY:284; PTX-NEXT: // %bb.0:285; PTX-NEXT: ld.param.b64 %rd1, [areg_64_i16_param_0];286; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];287; PTX-NEXT: ld.param.b64 %rd2, [areg_64_i16_param_1];288; PTX-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};289; PTX-NEXT: ret;290 %load = load <16 x i16>, ptr addrspace(1) %in291 store <16 x i16> %load, ptr addrspace(1) %out292 ret void293}294define void @areg_64_half(ptr addrspace(1) %in, ptr addrspace(1) %out) {295; PTX-LABEL: areg_64_half(296; PTX: {297; PTX-NEXT: .reg .b32 %r<9>;298; PTX-NEXT: .reg .b64 %rd<3>;299; PTX-EMPTY:300; PTX-NEXT: // %bb.0:301; PTX-NEXT: ld.param.b64 %rd1, [areg_64_half_param_0];302; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];303; PTX-NEXT: ld.param.b64 %rd2, [areg_64_half_param_1];304; PTX-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};305; PTX-NEXT: ret;306 %load = load <16 x half>, ptr addrspace(1) %in307 store <16 x half> %load, ptr addrspace(1) %out308 ret void309}310define void @areg_64_bfloat(ptr addrspace(1) %in, ptr addrspace(1) %out) {311; PTX-LABEL: areg_64_bfloat(312; PTX: {313; PTX-NEXT: .reg .b32 %r<9>;314; PTX-NEXT: .reg .b64 %rd<3>;315; PTX-EMPTY:316; PTX-NEXT: // %bb.0:317; PTX-NEXT: ld.param.b64 %rd1, [areg_64_bfloat_param_0];318; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];319; PTX-NEXT: ld.param.b64 %rd2, [areg_64_bfloat_param_1];320; PTX-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};321; PTX-NEXT: ret;322 %load = load <16 x bfloat>, ptr addrspace(1) %in323 store <16 x bfloat> %load, ptr addrspace(1) %out324 ret void325}326 327define void @areg_64_i32(ptr addrspace(1) %in, ptr addrspace(1) %out) {328; PTX-LABEL: areg_64_i32(329; PTX: {330; PTX-NEXT: .reg .b64 %rd<7>;331; PTX-EMPTY:332; PTX-NEXT: // %bb.0:333; PTX-NEXT: ld.param.b64 %rd1, [areg_64_i32_param_0];334; PTX-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];335; PTX-NEXT: ld.param.b64 %rd6, [areg_64_i32_param_1];336; PTX-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};337; PTX-NEXT: ret;338 %load = load <8 x i32>, ptr addrspace(1) %in339 store <8 x i32> %load, ptr addrspace(1) %out340 ret void341}342 343define void @areg_64_i64(ptr addrspace(1) %in, ptr addrspace(1) %out) {344; PTX-LABEL: areg_64_i64(345; PTX: {346; PTX-NEXT: .reg .b64 %rd<7>;347; PTX-EMPTY:348; PTX-NEXT: // %bb.0:349; PTX-NEXT: ld.param.b64 %rd1, [areg_64_i64_param_0];350; PTX-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];351; PTX-NEXT: ld.param.b64 %rd6, [areg_64_i64_param_1];352; PTX-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};353; PTX-NEXT: ret;354 %load = load <4 x i64>, ptr addrspace(1) %in355 store <4 x i64> %load, ptr addrspace(1) %out356 ret void357}358 359define void @areg_64_float(ptr addrspace(1) %in, ptr addrspace(1) %out) {360; PTX-LABEL: areg_64_float(361; PTX: {362; PTX-NEXT: .reg .b64 %rd<7>;363; PTX-EMPTY:364; PTX-NEXT: // %bb.0:365; PTX-NEXT: ld.param.b64 %rd1, [areg_64_float_param_0];366; PTX-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];367; PTX-NEXT: ld.param.b64 %rd6, [areg_64_float_param_1];368; PTX-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};369; PTX-NEXT: ret;370 %load = load <8 x float>, ptr addrspace(1) %in371 store <8 x float> %load, ptr addrspace(1) %out372 ret void373}374 375define void @areg_64_double(ptr addrspace(1) %in, ptr addrspace(1) %out) {376; PTX-LABEL: areg_64_double(377; PTX: {378; PTX-NEXT: .reg .b64 %rd<7>;379; PTX-EMPTY:380; PTX-NEXT: // %bb.0:381; PTX-NEXT: ld.param.b64 %rd1, [areg_64_double_param_0];382; PTX-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];383; PTX-NEXT: ld.param.b64 %rd6, [areg_64_double_param_1];384; PTX-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};385; PTX-NEXT: ret;386 %load = load <4 x double>, ptr addrspace(1) %in387 store <4 x double> %load, ptr addrspace(1) %out388 ret void389}390 391define void @ari_64_i8(ptr addrspace(1) %in, ptr addrspace(1) %out) {392; PTX-LABEL: ari_64_i8(393; PTX: {394; PTX-NEXT: .reg .b32 %r<9>;395; PTX-NEXT: .reg .b64 %rd<3>;396; PTX-EMPTY:397; PTX-NEXT: // %bb.0:398; PTX-NEXT: ld.param.b64 %rd1, [ari_64_i8_param_0];399; PTX-NEXT: ld.param.b64 %rd2, [ari_64_i8_param_1];400; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1+32];401; PTX-NEXT: st.global.v8.b32 [%rd2+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};402; PTX-NEXT: ret;403 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32404 %load = load <32 x i8>, ptr addrspace(1) %in.offset405 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32406 store <32 x i8> %load, ptr addrspace(1) %out.offset407 ret void408}409 410define void @ari_64_i16(ptr addrspace(1) %in, ptr addrspace(1) %out) {411; PTX-LABEL: ari_64_i16(412; PTX: {413; PTX-NEXT: .reg .b32 %r<9>;414; PTX-NEXT: .reg .b64 %rd<3>;415; PTX-EMPTY:416; PTX-NEXT: // %bb.0:417; PTX-NEXT: ld.param.b64 %rd1, [ari_64_i16_param_0];418; PTX-NEXT: ld.param.b64 %rd2, [ari_64_i16_param_1];419; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1+32];420; PTX-NEXT: st.global.v8.b32 [%rd2+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};421; PTX-NEXT: ret;422 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32423 %load = load <16 x i16>, ptr addrspace(1) %in.offset424 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32425 store <16 x i16> %load, ptr addrspace(1) %out.offset426 ret void427}428 429define void @ari_64_half(ptr addrspace(1) %in, ptr addrspace(1) %out) {430; PTX-LABEL: ari_64_half(431; PTX: {432; PTX-NEXT: .reg .b32 %r<9>;433; PTX-NEXT: .reg .b64 %rd<3>;434; PTX-EMPTY:435; PTX-NEXT: // %bb.0:436; PTX-NEXT: ld.param.b64 %rd1, [ari_64_half_param_0];437; PTX-NEXT: ld.param.b64 %rd2, [ari_64_half_param_1];438; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1+32];439; PTX-NEXT: st.global.v8.b32 [%rd2+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};440; PTX-NEXT: ret;441 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32442 %load = load <16 x half>, ptr addrspace(1) %in.offset443 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32444 store <16 x half> %load, ptr addrspace(1) %out.offset445 ret void446}447 448define void @ari_64_bfloat(ptr addrspace(1) %in, ptr addrspace(1) %out) {449; PTX-LABEL: ari_64_bfloat(450; PTX: {451; PTX-NEXT: .reg .b32 %r<9>;452; PTX-NEXT: .reg .b64 %rd<3>;453; PTX-EMPTY:454; PTX-NEXT: // %bb.0:455; PTX-NEXT: ld.param.b64 %rd1, [ari_64_bfloat_param_0];456; PTX-NEXT: ld.param.b64 %rd2, [ari_64_bfloat_param_1];457; PTX-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1+32];458; PTX-NEXT: st.global.v8.b32 [%rd2+32], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};459; PTX-NEXT: ret;460 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32461 %load = load <16 x bfloat>, ptr addrspace(1) %in.offset462 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32463 store <16 x bfloat> %load, ptr addrspace(1) %out.offset464 ret void465}466 467define void @ari_64_i32(ptr addrspace(1) %in, ptr addrspace(1) %out) {468; PTX-LABEL: ari_64_i32(469; PTX: {470; PTX-NEXT: .reg .b64 %rd<7>;471; PTX-EMPTY:472; PTX-NEXT: // %bb.0:473; PTX-NEXT: ld.param.b64 %rd1, [ari_64_i32_param_0];474; PTX-NEXT: ld.param.b64 %rd2, [ari_64_i32_param_1];475; PTX-NEXT: ld.global.v4.b64 {%rd3, %rd4, %rd5, %rd6}, [%rd1+32];476; PTX-NEXT: st.global.v4.b64 [%rd2+32], {%rd3, %rd4, %rd5, %rd6};477; PTX-NEXT: ret;478 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32479 %load = load <8 x i32>, ptr addrspace(1) %in.offset480 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32481 store <8 x i32> %load, ptr addrspace(1) %out.offset482 ret void483}484 485define void @ari_64_i64(ptr addrspace(1) %in, ptr addrspace(1) %out) {486; PTX-LABEL: ari_64_i64(487; PTX: {488; PTX-NEXT: .reg .b64 %rd<7>;489; PTX-EMPTY:490; PTX-NEXT: // %bb.0:491; PTX-NEXT: ld.param.b64 %rd1, [ari_64_i64_param_0];492; PTX-NEXT: ld.param.b64 %rd2, [ari_64_i64_param_1];493; PTX-NEXT: ld.global.v4.b64 {%rd3, %rd4, %rd5, %rd6}, [%rd1+32];494; PTX-NEXT: st.global.v4.b64 [%rd2+32], {%rd3, %rd4, %rd5, %rd6};495; PTX-NEXT: ret;496 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32497 %load = load <4 x i64>, ptr addrspace(1) %in.offset498 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32499 store <4 x i64> %load, ptr addrspace(1) %out.offset500 ret void501}502 503define void @ari_64_float(ptr addrspace(1) %in, ptr addrspace(1) %out) {504; PTX-LABEL: ari_64_float(505; PTX: {506; PTX-NEXT: .reg .b64 %rd<7>;507; PTX-EMPTY:508; PTX-NEXT: // %bb.0:509; PTX-NEXT: ld.param.b64 %rd1, [ari_64_float_param_0];510; PTX-NEXT: ld.param.b64 %rd2, [ari_64_float_param_1];511; PTX-NEXT: ld.global.v4.b64 {%rd3, %rd4, %rd5, %rd6}, [%rd1+32];512; PTX-NEXT: st.global.v4.b64 [%rd2+32], {%rd3, %rd4, %rd5, %rd6};513; PTX-NEXT: ret;514 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32515 %load = load <8 x float>, ptr addrspace(1) %in.offset516 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32517 store <8 x float> %load, ptr addrspace(1) %out.offset518 ret void519}520 521define void @ari_64_double(ptr addrspace(1) %in, ptr addrspace(1) %out) {522; PTX-LABEL: ari_64_double(523; PTX: {524; PTX-NEXT: .reg .b64 %rd<7>;525; PTX-EMPTY:526; PTX-NEXT: // %bb.0:527; PTX-NEXT: ld.param.b64 %rd1, [ari_64_double_param_0];528; PTX-NEXT: ld.param.b64 %rd2, [ari_64_double_param_1];529; PTX-NEXT: ld.global.v4.b64 {%rd3, %rd4, %rd5, %rd6}, [%rd1+32];530; PTX-NEXT: st.global.v4.b64 [%rd2+32], {%rd3, %rd4, %rd5, %rd6};531; PTX-NEXT: ret;532 %in.offset = getelementptr inbounds i8, ptr addrspace(1) %in, i32 32533 %load = load <4 x double>, ptr addrspace(1) %in.offset534 %out.offset = getelementptr inbounds i8, ptr addrspace(1) %out, i32 32535 store <4 x double> %load, ptr addrspace(1) %out.offset536 ret void537}538