1680 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_90 -mattr=+ptx87 | FileCheck -check-prefixes=CHECK,SM90 %s3; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | %ptxas-verify -arch=sm_90 %}4; RUN: llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | FileCheck %s -check-prefixes=CHECK,SM1005; RUN: %if ptxas-sm_100 && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | %ptxas-verify -arch=sm_100 %}6 7; This test is based on load-store-vectors.ll,8; and contains testing for lowering 256-bit vector loads/stores9 10; Types we are checking: i8, i16, half, bfloat, i32, i64, f32, f6411 12; Address spaces we are checking: generic, global, shared, local13; - Global is the only address space that currently supports 256-bit/v8 loads/stores,14; the other cases will legalize by splitting to smaller vectors15 16; 256-bit vector loads/stores are only legal for blackwell+, so on sm_90, the vectors will be split17 18; Types of loads/stores we are checking: normal, volatile19; - No need to check atomic loads/stores (monotonic and unordered) like load-store-vectors.ll checks,20; because those only work with non-vectors21 22;; generic statespace23 24; generic25 26define void @generic_32xi8(ptr %a, ptr %b) {27; CHECK-LABEL: generic_32xi8(28; CHECK: {29; CHECK-NEXT: .reg .b32 %r<9>;30; CHECK-NEXT: .reg .b64 %rd<3>;31; CHECK-EMPTY:32; CHECK-NEXT: // %bb.0:33; CHECK-NEXT: ld.param.b64 %rd1, [generic_32xi8_param_0];34; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];35; CHECK-NEXT: ld.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];36; CHECK-NEXT: ld.param.b64 %rd2, [generic_32xi8_param_1];37; CHECK-NEXT: st.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};38; CHECK-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};39; CHECK-NEXT: ret;40 %a.load = load <32 x i8>, ptr %a41 store <32 x i8> %a.load, ptr %b42 ret void43}44 45define void @generic_16xi16(ptr %a, ptr %b) {46; CHECK-LABEL: generic_16xi16(47; CHECK: {48; CHECK-NEXT: .reg .b32 %r<9>;49; CHECK-NEXT: .reg .b64 %rd<3>;50; CHECK-EMPTY:51; CHECK-NEXT: // %bb.0:52; CHECK-NEXT: ld.param.b64 %rd1, [generic_16xi16_param_0];53; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];54; CHECK-NEXT: ld.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];55; CHECK-NEXT: ld.param.b64 %rd2, [generic_16xi16_param_1];56; CHECK-NEXT: st.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};57; CHECK-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};58; CHECK-NEXT: ret;59 %a.load = load <16 x i16>, ptr %a60 store <16 x i16> %a.load, ptr %b61 ret void62}63 64define void @generic_16xhalf(ptr %a, ptr %b) {65; CHECK-LABEL: generic_16xhalf(66; CHECK: {67; CHECK-NEXT: .reg .b32 %r<9>;68; CHECK-NEXT: .reg .b64 %rd<3>;69; CHECK-EMPTY:70; CHECK-NEXT: // %bb.0:71; CHECK-NEXT: ld.param.b64 %rd1, [generic_16xhalf_param_0];72; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];73; CHECK-NEXT: ld.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];74; CHECK-NEXT: ld.param.b64 %rd2, [generic_16xhalf_param_1];75; CHECK-NEXT: st.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};76; CHECK-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};77; CHECK-NEXT: ret;78 %a.load = load <16 x half>, ptr %a79 store <16 x half> %a.load, ptr %b80 ret void81}82 83define void @generic_16xbfloat(ptr %a, ptr %b) {84; CHECK-LABEL: generic_16xbfloat(85; CHECK: {86; CHECK-NEXT: .reg .b32 %r<9>;87; CHECK-NEXT: .reg .b64 %rd<3>;88; CHECK-EMPTY:89; CHECK-NEXT: // %bb.0:90; CHECK-NEXT: ld.param.b64 %rd1, [generic_16xbfloat_param_0];91; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];92; CHECK-NEXT: ld.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];93; CHECK-NEXT: ld.param.b64 %rd2, [generic_16xbfloat_param_1];94; CHECK-NEXT: st.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};95; CHECK-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};96; CHECK-NEXT: ret;97 %a.load = load <16 x bfloat>, ptr %a98 store <16 x bfloat> %a.load, ptr %b99 ret void100}101 102define void @generic_8xi32(ptr %a, ptr %b) {103; SM90-LABEL: generic_8xi32(104; SM90: {105; SM90-NEXT: .reg .b32 %r<9>;106; SM90-NEXT: .reg .b64 %rd<3>;107; SM90-EMPTY:108; SM90-NEXT: // %bb.0:109; SM90-NEXT: ld.param.b64 %rd1, [generic_8xi32_param_0];110; SM90-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];111; SM90-NEXT: ld.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];112; SM90-NEXT: ld.param.b64 %rd2, [generic_8xi32_param_1];113; SM90-NEXT: st.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};114; SM90-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};115; SM90-NEXT: ret;116;117; SM100-LABEL: generic_8xi32(118; SM100: {119; SM100-NEXT: .reg .b64 %rd<7>;120; SM100-EMPTY:121; SM100-NEXT: // %bb.0:122; SM100-NEXT: ld.param.b64 %rd1, [generic_8xi32_param_0];123; SM100-NEXT: ld.v2.b64 {%rd2, %rd3}, [%rd1];124; SM100-NEXT: ld.v2.b64 {%rd4, %rd5}, [%rd1+16];125; SM100-NEXT: ld.param.b64 %rd6, [generic_8xi32_param_1];126; SM100-NEXT: st.v2.b64 [%rd6+16], {%rd4, %rd5};127; SM100-NEXT: st.v2.b64 [%rd6], {%rd2, %rd3};128; SM100-NEXT: ret;129 %a.load = load <8 x i32>, ptr %a130 store <8 x i32> %a.load, ptr %b131 ret void132}133 134define void @generic_4xi64(ptr %a, ptr %b) {135; CHECK-LABEL: generic_4xi64(136; CHECK: {137; CHECK-NEXT: .reg .b64 %rd<7>;138; CHECK-EMPTY:139; CHECK-NEXT: // %bb.0:140; CHECK-NEXT: ld.param.b64 %rd1, [generic_4xi64_param_0];141; CHECK-NEXT: ld.v2.b64 {%rd2, %rd3}, [%rd1];142; CHECK-NEXT: ld.v2.b64 {%rd4, %rd5}, [%rd1+16];143; CHECK-NEXT: ld.param.b64 %rd6, [generic_4xi64_param_1];144; CHECK-NEXT: st.v2.b64 [%rd6+16], {%rd4, %rd5};145; CHECK-NEXT: st.v2.b64 [%rd6], {%rd2, %rd3};146; CHECK-NEXT: ret;147 %a.load = load <4 x i64>, ptr %a148 store <4 x i64> %a.load, ptr %b149 ret void150}151 152define void @generic_8xfloat(ptr %a, ptr %b) {153; SM90-LABEL: generic_8xfloat(154; SM90: {155; SM90-NEXT: .reg .b32 %r<9>;156; SM90-NEXT: .reg .b64 %rd<3>;157; SM90-EMPTY:158; SM90-NEXT: // %bb.0:159; SM90-NEXT: ld.param.b64 %rd1, [generic_8xfloat_param_0];160; SM90-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];161; SM90-NEXT: ld.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];162; SM90-NEXT: ld.param.b64 %rd2, [generic_8xfloat_param_1];163; SM90-NEXT: st.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};164; SM90-NEXT: st.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};165; SM90-NEXT: ret;166;167; SM100-LABEL: generic_8xfloat(168; SM100: {169; SM100-NEXT: .reg .b64 %rd<7>;170; SM100-EMPTY:171; SM100-NEXT: // %bb.0:172; SM100-NEXT: ld.param.b64 %rd1, [generic_8xfloat_param_0];173; SM100-NEXT: ld.v2.b64 {%rd2, %rd3}, [%rd1];174; SM100-NEXT: ld.v2.b64 {%rd4, %rd5}, [%rd1+16];175; SM100-NEXT: ld.param.b64 %rd6, [generic_8xfloat_param_1];176; SM100-NEXT: st.v2.b64 [%rd6+16], {%rd4, %rd5};177; SM100-NEXT: st.v2.b64 [%rd6], {%rd2, %rd3};178; SM100-NEXT: ret;179 %a.load = load <8 x float>, ptr %a180 store <8 x float> %a.load, ptr %b181 ret void182}183 184define void @generic_4xdouble(ptr %a, ptr %b) {185; CHECK-LABEL: generic_4xdouble(186; CHECK: {187; CHECK-NEXT: .reg .b64 %rd<7>;188; CHECK-EMPTY:189; CHECK-NEXT: // %bb.0:190; CHECK-NEXT: ld.param.b64 %rd1, [generic_4xdouble_param_0];191; CHECK-NEXT: ld.v2.b64 {%rd2, %rd3}, [%rd1];192; CHECK-NEXT: ld.v2.b64 {%rd4, %rd5}, [%rd1+16];193; CHECK-NEXT: ld.param.b64 %rd6, [generic_4xdouble_param_1];194; CHECK-NEXT: st.v2.b64 [%rd6+16], {%rd4, %rd5};195; CHECK-NEXT: st.v2.b64 [%rd6], {%rd2, %rd3};196; CHECK-NEXT: ret;197 %a.load = load <4 x double>, ptr %a198 store <4 x double> %a.load, ptr %b199 ret void200}201 202; generic_volatile203 204define void @generic_volatile_32xi8(ptr %a, ptr %b) {205; CHECK-LABEL: generic_volatile_32xi8(206; CHECK: {207; CHECK-NEXT: .reg .b32 %r<9>;208; CHECK-NEXT: .reg .b64 %rd<3>;209; CHECK-EMPTY:210; CHECK-NEXT: // %bb.0:211; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_32xi8_param_0];212; CHECK-NEXT: ld.volatile.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];213; CHECK-NEXT: ld.volatile.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];214; CHECK-NEXT: ld.param.b64 %rd2, [generic_volatile_32xi8_param_1];215; CHECK-NEXT: st.volatile.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};216; CHECK-NEXT: st.volatile.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};217; CHECK-NEXT: ret;218 %a.load = load volatile <32 x i8>, ptr %a219 store volatile <32 x i8> %a.load, ptr %b220 ret void221}222 223define void @generic_volatile_16xi16(ptr %a, ptr %b) {224; CHECK-LABEL: generic_volatile_16xi16(225; CHECK: {226; CHECK-NEXT: .reg .b32 %r<9>;227; CHECK-NEXT: .reg .b64 %rd<3>;228; CHECK-EMPTY:229; CHECK-NEXT: // %bb.0:230; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_16xi16_param_0];231; CHECK-NEXT: ld.volatile.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];232; CHECK-NEXT: ld.volatile.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];233; CHECK-NEXT: ld.param.b64 %rd2, [generic_volatile_16xi16_param_1];234; CHECK-NEXT: st.volatile.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};235; CHECK-NEXT: st.volatile.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};236; CHECK-NEXT: ret;237 %a.load = load volatile <16 x i16>, ptr %a238 store volatile <16 x i16> %a.load, ptr %b239 ret void240}241 242define void @generic_volatile_16xhalf(ptr %a, ptr %b) {243; CHECK-LABEL: generic_volatile_16xhalf(244; CHECK: {245; CHECK-NEXT: .reg .b32 %r<9>;246; CHECK-NEXT: .reg .b64 %rd<3>;247; CHECK-EMPTY:248; CHECK-NEXT: // %bb.0:249; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_16xhalf_param_0];250; CHECK-NEXT: ld.volatile.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];251; CHECK-NEXT: ld.volatile.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];252; CHECK-NEXT: ld.param.b64 %rd2, [generic_volatile_16xhalf_param_1];253; CHECK-NEXT: st.volatile.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};254; CHECK-NEXT: st.volatile.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};255; CHECK-NEXT: ret;256 %a.load = load volatile <16 x half>, ptr %a257 store volatile <16 x half> %a.load, ptr %b258 ret void259}260 261define void @generic_volatile_16xbfloat(ptr %a, ptr %b) {262; CHECK-LABEL: generic_volatile_16xbfloat(263; CHECK: {264; CHECK-NEXT: .reg .b32 %r<9>;265; CHECK-NEXT: .reg .b64 %rd<3>;266; CHECK-EMPTY:267; CHECK-NEXT: // %bb.0:268; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_16xbfloat_param_0];269; CHECK-NEXT: ld.volatile.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];270; CHECK-NEXT: ld.volatile.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];271; CHECK-NEXT: ld.param.b64 %rd2, [generic_volatile_16xbfloat_param_1];272; CHECK-NEXT: st.volatile.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};273; CHECK-NEXT: st.volatile.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};274; CHECK-NEXT: ret;275 %a.load = load volatile <16 x bfloat>, ptr %a276 store volatile <16 x bfloat> %a.load, ptr %b277 ret void278}279 280define void @generic_volatile_8xi32(ptr %a, ptr %b) {281; SM90-LABEL: generic_volatile_8xi32(282; SM90: {283; SM90-NEXT: .reg .b32 %r<9>;284; SM90-NEXT: .reg .b64 %rd<3>;285; SM90-EMPTY:286; SM90-NEXT: // %bb.0:287; SM90-NEXT: ld.param.b64 %rd1, [generic_volatile_8xi32_param_0];288; SM90-NEXT: ld.volatile.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];289; SM90-NEXT: ld.volatile.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];290; SM90-NEXT: ld.param.b64 %rd2, [generic_volatile_8xi32_param_1];291; SM90-NEXT: st.volatile.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};292; SM90-NEXT: st.volatile.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};293; SM90-NEXT: ret;294;295; SM100-LABEL: generic_volatile_8xi32(296; SM100: {297; SM100-NEXT: .reg .b64 %rd<7>;298; SM100-EMPTY:299; SM100-NEXT: // %bb.0:300; SM100-NEXT: ld.param.b64 %rd1, [generic_volatile_8xi32_param_0];301; SM100-NEXT: ld.volatile.v2.b64 {%rd2, %rd3}, [%rd1];302; SM100-NEXT: ld.volatile.v2.b64 {%rd4, %rd5}, [%rd1+16];303; SM100-NEXT: ld.param.b64 %rd6, [generic_volatile_8xi32_param_1];304; SM100-NEXT: st.volatile.v2.b64 [%rd6+16], {%rd4, %rd5};305; SM100-NEXT: st.volatile.v2.b64 [%rd6], {%rd2, %rd3};306; SM100-NEXT: ret;307 %a.load = load volatile <8 x i32>, ptr %a308 store volatile <8 x i32> %a.load, ptr %b309 ret void310}311 312define void @generic_volatile_4xi64(ptr %a, ptr %b) {313; CHECK-LABEL: generic_volatile_4xi64(314; CHECK: {315; CHECK-NEXT: .reg .b64 %rd<7>;316; CHECK-EMPTY:317; CHECK-NEXT: // %bb.0:318; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_4xi64_param_0];319; CHECK-NEXT: ld.volatile.v2.b64 {%rd2, %rd3}, [%rd1];320; CHECK-NEXT: ld.volatile.v2.b64 {%rd4, %rd5}, [%rd1+16];321; CHECK-NEXT: ld.param.b64 %rd6, [generic_volatile_4xi64_param_1];322; CHECK-NEXT: st.volatile.v2.b64 [%rd6+16], {%rd4, %rd5};323; CHECK-NEXT: st.volatile.v2.b64 [%rd6], {%rd2, %rd3};324; CHECK-NEXT: ret;325 %a.load = load volatile <4 x i64>, ptr %a326 store volatile <4 x i64> %a.load, ptr %b327 ret void328}329 330define void @generic_volatile_8xfloat(ptr %a, ptr %b) {331; SM90-LABEL: generic_volatile_8xfloat(332; SM90: {333; SM90-NEXT: .reg .b32 %r<9>;334; SM90-NEXT: .reg .b64 %rd<3>;335; SM90-EMPTY:336; SM90-NEXT: // %bb.0:337; SM90-NEXT: ld.param.b64 %rd1, [generic_volatile_8xfloat_param_0];338; SM90-NEXT: ld.volatile.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];339; SM90-NEXT: ld.volatile.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];340; SM90-NEXT: ld.param.b64 %rd2, [generic_volatile_8xfloat_param_1];341; SM90-NEXT: st.volatile.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};342; SM90-NEXT: st.volatile.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};343; SM90-NEXT: ret;344;345; SM100-LABEL: generic_volatile_8xfloat(346; SM100: {347; SM100-NEXT: .reg .b64 %rd<7>;348; SM100-EMPTY:349; SM100-NEXT: // %bb.0:350; SM100-NEXT: ld.param.b64 %rd1, [generic_volatile_8xfloat_param_0];351; SM100-NEXT: ld.volatile.v2.b64 {%rd2, %rd3}, [%rd1];352; SM100-NEXT: ld.volatile.v2.b64 {%rd4, %rd5}, [%rd1+16];353; SM100-NEXT: ld.param.b64 %rd6, [generic_volatile_8xfloat_param_1];354; SM100-NEXT: st.volatile.v2.b64 [%rd6+16], {%rd4, %rd5};355; SM100-NEXT: st.volatile.v2.b64 [%rd6], {%rd2, %rd3};356; SM100-NEXT: ret;357 %a.load = load volatile <8 x float>, ptr %a358 store volatile <8 x float> %a.load, ptr %b359 ret void360}361 362define void @generic_volatile_4xdouble(ptr %a, ptr %b) {363; CHECK-LABEL: generic_volatile_4xdouble(364; CHECK: {365; CHECK-NEXT: .reg .b64 %rd<7>;366; CHECK-EMPTY:367; CHECK-NEXT: // %bb.0:368; CHECK-NEXT: ld.param.b64 %rd1, [generic_volatile_4xdouble_param_0];369; CHECK-NEXT: ld.volatile.v2.b64 {%rd2, %rd3}, [%rd1];370; CHECK-NEXT: ld.volatile.v2.b64 {%rd4, %rd5}, [%rd1+16];371; CHECK-NEXT: ld.param.b64 %rd6, [generic_volatile_4xdouble_param_1];372; CHECK-NEXT: st.volatile.v2.b64 [%rd6+16], {%rd4, %rd5};373; CHECK-NEXT: st.volatile.v2.b64 [%rd6], {%rd2, %rd3};374; CHECK-NEXT: ret;375 %a.load = load volatile <4 x double>, ptr %a376 store volatile <4 x double> %a.load, ptr %b377 ret void378}379 380;; global statespace381 382; global383 384define void @global_32xi8(ptr addrspace(1) %a, ptr addrspace(1) %b) {385; SM90-LABEL: global_32xi8(386; SM90: {387; SM90-NEXT: .reg .b32 %r<9>;388; SM90-NEXT: .reg .b64 %rd<3>;389; SM90-EMPTY:390; SM90-NEXT: // %bb.0:391; SM90-NEXT: ld.param.b64 %rd1, [global_32xi8_param_0];392; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];393; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];394; SM90-NEXT: ld.param.b64 %rd2, [global_32xi8_param_1];395; SM90-NEXT: st.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};396; SM90-NEXT: st.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};397; SM90-NEXT: ret;398;399; SM100-LABEL: global_32xi8(400; SM100: {401; SM100-NEXT: .reg .b32 %r<9>;402; SM100-NEXT: .reg .b64 %rd<3>;403; SM100-EMPTY:404; SM100-NEXT: // %bb.0:405; SM100-NEXT: ld.param.b64 %rd1, [global_32xi8_param_0];406; SM100-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];407; SM100-NEXT: ld.param.b64 %rd2, [global_32xi8_param_1];408; SM100-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};409; SM100-NEXT: ret;410 %a.load = load <32 x i8>, ptr addrspace(1) %a411 store <32 x i8> %a.load, ptr addrspace(1) %b412 ret void413}414 415define void @global_16xi16(ptr addrspace(1) %a, ptr addrspace(1) %b) {416; SM90-LABEL: global_16xi16(417; SM90: {418; SM90-NEXT: .reg .b32 %r<9>;419; SM90-NEXT: .reg .b64 %rd<3>;420; SM90-EMPTY:421; SM90-NEXT: // %bb.0:422; SM90-NEXT: ld.param.b64 %rd1, [global_16xi16_param_0];423; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];424; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];425; SM90-NEXT: ld.param.b64 %rd2, [global_16xi16_param_1];426; SM90-NEXT: st.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};427; SM90-NEXT: st.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};428; SM90-NEXT: ret;429;430; SM100-LABEL: global_16xi16(431; SM100: {432; SM100-NEXT: .reg .b32 %r<9>;433; SM100-NEXT: .reg .b64 %rd<3>;434; SM100-EMPTY:435; SM100-NEXT: // %bb.0:436; SM100-NEXT: ld.param.b64 %rd1, [global_16xi16_param_0];437; SM100-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];438; SM100-NEXT: ld.param.b64 %rd2, [global_16xi16_param_1];439; SM100-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};440; SM100-NEXT: ret;441 %a.load = load <16 x i16>, ptr addrspace(1) %a442 store <16 x i16> %a.load, ptr addrspace(1) %b443 ret void444}445 446define void @global_16xhalf(ptr addrspace(1) %a, ptr addrspace(1) %b) {447; SM90-LABEL: global_16xhalf(448; SM90: {449; SM90-NEXT: .reg .b32 %r<9>;450; SM90-NEXT: .reg .b64 %rd<3>;451; SM90-EMPTY:452; SM90-NEXT: // %bb.0:453; SM90-NEXT: ld.param.b64 %rd1, [global_16xhalf_param_0];454; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];455; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];456; SM90-NEXT: ld.param.b64 %rd2, [global_16xhalf_param_1];457; SM90-NEXT: st.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};458; SM90-NEXT: st.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};459; SM90-NEXT: ret;460;461; SM100-LABEL: global_16xhalf(462; SM100: {463; SM100-NEXT: .reg .b32 %r<9>;464; SM100-NEXT: .reg .b64 %rd<3>;465; SM100-EMPTY:466; SM100-NEXT: // %bb.0:467; SM100-NEXT: ld.param.b64 %rd1, [global_16xhalf_param_0];468; SM100-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];469; SM100-NEXT: ld.param.b64 %rd2, [global_16xhalf_param_1];470; SM100-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};471; SM100-NEXT: ret;472 %a.load = load <16 x half>, ptr addrspace(1) %a473 store <16 x half> %a.load, ptr addrspace(1) %b474 ret void475}476 477define void @global_16xbfloat(ptr addrspace(1) %a, ptr addrspace(1) %b) {478; SM90-LABEL: global_16xbfloat(479; SM90: {480; SM90-NEXT: .reg .b32 %r<9>;481; SM90-NEXT: .reg .b64 %rd<3>;482; SM90-EMPTY:483; SM90-NEXT: // %bb.0:484; SM90-NEXT: ld.param.b64 %rd1, [global_16xbfloat_param_0];485; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];486; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];487; SM90-NEXT: ld.param.b64 %rd2, [global_16xbfloat_param_1];488; SM90-NEXT: st.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};489; SM90-NEXT: st.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};490; SM90-NEXT: ret;491;492; SM100-LABEL: global_16xbfloat(493; SM100: {494; SM100-NEXT: .reg .b32 %r<9>;495; SM100-NEXT: .reg .b64 %rd<3>;496; SM100-EMPTY:497; SM100-NEXT: // %bb.0:498; SM100-NEXT: ld.param.b64 %rd1, [global_16xbfloat_param_0];499; SM100-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];500; SM100-NEXT: ld.param.b64 %rd2, [global_16xbfloat_param_1];501; SM100-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};502; SM100-NEXT: ret;503 %a.load = load <16 x bfloat>, ptr addrspace(1) %a504 store <16 x bfloat> %a.load, ptr addrspace(1) %b505 ret void506}507 508define void @global_8xi32(ptr addrspace(1) %a, ptr addrspace(1) %b) {509; SM90-LABEL: global_8xi32(510; SM90: {511; SM90-NEXT: .reg .b32 %r<9>;512; SM90-NEXT: .reg .b64 %rd<3>;513; SM90-EMPTY:514; SM90-NEXT: // %bb.0:515; SM90-NEXT: ld.param.b64 %rd1, [global_8xi32_param_0];516; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];517; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];518; SM90-NEXT: ld.param.b64 %rd2, [global_8xi32_param_1];519; SM90-NEXT: st.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};520; SM90-NEXT: st.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};521; SM90-NEXT: ret;522;523; SM100-LABEL: global_8xi32(524; SM100: {525; SM100-NEXT: .reg .b64 %rd<7>;526; SM100-EMPTY:527; SM100-NEXT: // %bb.0:528; SM100-NEXT: ld.param.b64 %rd1, [global_8xi32_param_0];529; SM100-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];530; SM100-NEXT: ld.param.b64 %rd6, [global_8xi32_param_1];531; SM100-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};532; SM100-NEXT: ret;533 %a.load = load <8 x i32>, ptr addrspace(1) %a534 store <8 x i32> %a.load, ptr addrspace(1) %b535 ret void536}537 538define void @global_4xi64(ptr addrspace(1) %a, ptr addrspace(1) %b) {539; SM90-LABEL: global_4xi64(540; SM90: {541; SM90-NEXT: .reg .b64 %rd<7>;542; SM90-EMPTY:543; SM90-NEXT: // %bb.0:544; SM90-NEXT: ld.param.b64 %rd1, [global_4xi64_param_0];545; SM90-NEXT: ld.global.v2.b64 {%rd2, %rd3}, [%rd1];546; SM90-NEXT: ld.global.v2.b64 {%rd4, %rd5}, [%rd1+16];547; SM90-NEXT: ld.param.b64 %rd6, [global_4xi64_param_1];548; SM90-NEXT: st.global.v2.b64 [%rd6+16], {%rd4, %rd5};549; SM90-NEXT: st.global.v2.b64 [%rd6], {%rd2, %rd3};550; SM90-NEXT: ret;551;552; SM100-LABEL: global_4xi64(553; SM100: {554; SM100-NEXT: .reg .b64 %rd<7>;555; SM100-EMPTY:556; SM100-NEXT: // %bb.0:557; SM100-NEXT: ld.param.b64 %rd1, [global_4xi64_param_0];558; SM100-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];559; SM100-NEXT: ld.param.b64 %rd6, [global_4xi64_param_1];560; SM100-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};561; SM100-NEXT: ret;562 %a.load = load <4 x i64>, ptr addrspace(1) %a563 store <4 x i64> %a.load, ptr addrspace(1) %b564 ret void565}566 567define void @global_8xfloat(ptr addrspace(1) %a, ptr addrspace(1) %b) {568; SM90-LABEL: global_8xfloat(569; SM90: {570; SM90-NEXT: .reg .b32 %r<9>;571; SM90-NEXT: .reg .b64 %rd<3>;572; SM90-EMPTY:573; SM90-NEXT: // %bb.0:574; SM90-NEXT: ld.param.b64 %rd1, [global_8xfloat_param_0];575; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];576; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];577; SM90-NEXT: ld.param.b64 %rd2, [global_8xfloat_param_1];578; SM90-NEXT: st.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};579; SM90-NEXT: st.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};580; SM90-NEXT: ret;581;582; SM100-LABEL: global_8xfloat(583; SM100: {584; SM100-NEXT: .reg .b64 %rd<7>;585; SM100-EMPTY:586; SM100-NEXT: // %bb.0:587; SM100-NEXT: ld.param.b64 %rd1, [global_8xfloat_param_0];588; SM100-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];589; SM100-NEXT: ld.param.b64 %rd6, [global_8xfloat_param_1];590; SM100-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};591; SM100-NEXT: ret;592 %a.load = load <8 x float>, ptr addrspace(1) %a593 store <8 x float> %a.load, ptr addrspace(1) %b594 ret void595}596 597define void @global_4xdouble(ptr addrspace(1) %a, ptr addrspace(1) %b) {598; SM90-LABEL: global_4xdouble(599; SM90: {600; SM90-NEXT: .reg .b64 %rd<7>;601; SM90-EMPTY:602; SM90-NEXT: // %bb.0:603; SM90-NEXT: ld.param.b64 %rd1, [global_4xdouble_param_0];604; SM90-NEXT: ld.global.v2.b64 {%rd2, %rd3}, [%rd1];605; SM90-NEXT: ld.global.v2.b64 {%rd4, %rd5}, [%rd1+16];606; SM90-NEXT: ld.param.b64 %rd6, [global_4xdouble_param_1];607; SM90-NEXT: st.global.v2.b64 [%rd6+16], {%rd4, %rd5};608; SM90-NEXT: st.global.v2.b64 [%rd6], {%rd2, %rd3};609; SM90-NEXT: ret;610;611; SM100-LABEL: global_4xdouble(612; SM100: {613; SM100-NEXT: .reg .b64 %rd<7>;614; SM100-EMPTY:615; SM100-NEXT: // %bb.0:616; SM100-NEXT: ld.param.b64 %rd1, [global_4xdouble_param_0];617; SM100-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];618; SM100-NEXT: ld.param.b64 %rd6, [global_4xdouble_param_1];619; SM100-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};620; SM100-NEXT: ret;621 %a.load = load <4 x double>, ptr addrspace(1) %a622 store <4 x double> %a.load, ptr addrspace(1) %b623 ret void624}625 626; global_volatile627 628define void @global_volatile_32xi8(ptr addrspace(1) %a, ptr addrspace(1) %b) {629; SM90-LABEL: global_volatile_32xi8(630; SM90: {631; SM90-NEXT: .reg .b32 %r<9>;632; SM90-NEXT: .reg .b64 %rd<3>;633; SM90-EMPTY:634; SM90-NEXT: // %bb.0:635; SM90-NEXT: ld.param.b64 %rd1, [global_volatile_32xi8_param_0];636; SM90-NEXT: ld.volatile.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];637; SM90-NEXT: ld.volatile.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];638; SM90-NEXT: ld.param.b64 %rd2, [global_volatile_32xi8_param_1];639; SM90-NEXT: st.volatile.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};640; SM90-NEXT: st.volatile.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};641; SM90-NEXT: ret;642;643; SM100-LABEL: global_volatile_32xi8(644; SM100: {645; SM100-NEXT: .reg .b32 %r<9>;646; SM100-NEXT: .reg .b64 %rd<3>;647; SM100-EMPTY:648; SM100-NEXT: // %bb.0:649; SM100-NEXT: ld.param.b64 %rd1, [global_volatile_32xi8_param_0];650; SM100-NEXT: ld.volatile.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];651; SM100-NEXT: ld.param.b64 %rd2, [global_volatile_32xi8_param_1];652; SM100-NEXT: st.volatile.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};653; SM100-NEXT: ret;654 %a.load = load volatile <32 x i8>, ptr addrspace(1) %a655 store volatile <32 x i8> %a.load, ptr addrspace(1) %b656 ret void657}658 659define void @global_volatile_16xi16(ptr addrspace(1) %a, ptr addrspace(1) %b) {660; SM90-LABEL: global_volatile_16xi16(661; SM90: {662; SM90-NEXT: .reg .b32 %r<9>;663; SM90-NEXT: .reg .b64 %rd<3>;664; SM90-EMPTY:665; SM90-NEXT: // %bb.0:666; SM90-NEXT: ld.param.b64 %rd1, [global_volatile_16xi16_param_0];667; SM90-NEXT: ld.volatile.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];668; SM90-NEXT: ld.volatile.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];669; SM90-NEXT: ld.param.b64 %rd2, [global_volatile_16xi16_param_1];670; SM90-NEXT: st.volatile.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};671; SM90-NEXT: st.volatile.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};672; SM90-NEXT: ret;673;674; SM100-LABEL: global_volatile_16xi16(675; SM100: {676; SM100-NEXT: .reg .b32 %r<9>;677; SM100-NEXT: .reg .b64 %rd<3>;678; SM100-EMPTY:679; SM100-NEXT: // %bb.0:680; SM100-NEXT: ld.param.b64 %rd1, [global_volatile_16xi16_param_0];681; SM100-NEXT: ld.volatile.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];682; SM100-NEXT: ld.param.b64 %rd2, [global_volatile_16xi16_param_1];683; SM100-NEXT: st.volatile.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};684; SM100-NEXT: ret;685 %a.load = load volatile <16 x i16>, ptr addrspace(1) %a686 store volatile <16 x i16> %a.load, ptr addrspace(1) %b687 ret void688}689 690define void @global_volatile_16xhalf(ptr addrspace(1) %a, ptr addrspace(1) %b) {691; SM90-LABEL: global_volatile_16xhalf(692; SM90: {693; SM90-NEXT: .reg .b32 %r<9>;694; SM90-NEXT: .reg .b64 %rd<3>;695; SM90-EMPTY:696; SM90-NEXT: // %bb.0:697; SM90-NEXT: ld.param.b64 %rd1, [global_volatile_16xhalf_param_0];698; SM90-NEXT: ld.volatile.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];699; SM90-NEXT: ld.volatile.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];700; SM90-NEXT: ld.param.b64 %rd2, [global_volatile_16xhalf_param_1];701; SM90-NEXT: st.volatile.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};702; SM90-NEXT: st.volatile.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};703; SM90-NEXT: ret;704;705; SM100-LABEL: global_volatile_16xhalf(706; SM100: {707; SM100-NEXT: .reg .b32 %r<9>;708; SM100-NEXT: .reg .b64 %rd<3>;709; SM100-EMPTY:710; SM100-NEXT: // %bb.0:711; SM100-NEXT: ld.param.b64 %rd1, [global_volatile_16xhalf_param_0];712; SM100-NEXT: ld.volatile.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];713; SM100-NEXT: ld.param.b64 %rd2, [global_volatile_16xhalf_param_1];714; SM100-NEXT: st.volatile.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};715; SM100-NEXT: ret;716 %a.load = load volatile <16 x half>, ptr addrspace(1) %a717 store volatile <16 x half> %a.load, ptr addrspace(1) %b718 ret void719}720 721define void @global_volatile_16xbfloat(ptr addrspace(1) %a, ptr addrspace(1) %b) {722; SM90-LABEL: global_volatile_16xbfloat(723; SM90: {724; SM90-NEXT: .reg .b32 %r<9>;725; SM90-NEXT: .reg .b64 %rd<3>;726; SM90-EMPTY:727; SM90-NEXT: // %bb.0:728; SM90-NEXT: ld.param.b64 %rd1, [global_volatile_16xbfloat_param_0];729; SM90-NEXT: ld.volatile.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];730; SM90-NEXT: ld.volatile.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];731; SM90-NEXT: ld.param.b64 %rd2, [global_volatile_16xbfloat_param_1];732; SM90-NEXT: st.volatile.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};733; SM90-NEXT: st.volatile.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};734; SM90-NEXT: ret;735;736; SM100-LABEL: global_volatile_16xbfloat(737; SM100: {738; SM100-NEXT: .reg .b32 %r<9>;739; SM100-NEXT: .reg .b64 %rd<3>;740; SM100-EMPTY:741; SM100-NEXT: // %bb.0:742; SM100-NEXT: ld.param.b64 %rd1, [global_volatile_16xbfloat_param_0];743; SM100-NEXT: ld.volatile.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];744; SM100-NEXT: ld.param.b64 %rd2, [global_volatile_16xbfloat_param_1];745; SM100-NEXT: st.volatile.global.v8.b32 [%rd2], {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8};746; SM100-NEXT: ret;747 %a.load = load volatile <16 x bfloat>, ptr addrspace(1) %a748 store volatile <16 x bfloat> %a.load, ptr addrspace(1) %b749 ret void750}751 752define void @global_volatile_8xi32(ptr addrspace(1) %a, ptr addrspace(1) %b) {753; SM90-LABEL: global_volatile_8xi32(754; SM90: {755; SM90-NEXT: .reg .b32 %r<9>;756; SM90-NEXT: .reg .b64 %rd<3>;757; SM90-EMPTY:758; SM90-NEXT: // %bb.0:759; SM90-NEXT: ld.param.b64 %rd1, [global_volatile_8xi32_param_0];760; SM90-NEXT: ld.volatile.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];761; SM90-NEXT: ld.volatile.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];762; SM90-NEXT: ld.param.b64 %rd2, [global_volatile_8xi32_param_1];763; SM90-NEXT: st.volatile.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};764; SM90-NEXT: st.volatile.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};765; SM90-NEXT: ret;766;767; SM100-LABEL: global_volatile_8xi32(768; SM100: {769; SM100-NEXT: .reg .b64 %rd<7>;770; SM100-EMPTY:771; SM100-NEXT: // %bb.0:772; SM100-NEXT: ld.param.b64 %rd1, [global_volatile_8xi32_param_0];773; SM100-NEXT: ld.volatile.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];774; SM100-NEXT: ld.param.b64 %rd6, [global_volatile_8xi32_param_1];775; SM100-NEXT: st.volatile.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};776; SM100-NEXT: ret;777 %a.load = load volatile <8 x i32>, ptr addrspace(1) %a778 store volatile <8 x i32> %a.load, ptr addrspace(1) %b779 ret void780}781 782define void @global_volatile_4xi64(ptr addrspace(1) %a, ptr addrspace(1) %b) {783; SM90-LABEL: global_volatile_4xi64(784; SM90: {785; SM90-NEXT: .reg .b64 %rd<7>;786; SM90-EMPTY:787; SM90-NEXT: // %bb.0:788; SM90-NEXT: ld.param.b64 %rd1, [global_volatile_4xi64_param_0];789; SM90-NEXT: ld.volatile.global.v2.b64 {%rd2, %rd3}, [%rd1];790; SM90-NEXT: ld.volatile.global.v2.b64 {%rd4, %rd5}, [%rd1+16];791; SM90-NEXT: ld.param.b64 %rd6, [global_volatile_4xi64_param_1];792; SM90-NEXT: st.volatile.global.v2.b64 [%rd6+16], {%rd4, %rd5};793; SM90-NEXT: st.volatile.global.v2.b64 [%rd6], {%rd2, %rd3};794; SM90-NEXT: ret;795;796; SM100-LABEL: global_volatile_4xi64(797; SM100: {798; SM100-NEXT: .reg .b64 %rd<7>;799; SM100-EMPTY:800; SM100-NEXT: // %bb.0:801; SM100-NEXT: ld.param.b64 %rd1, [global_volatile_4xi64_param_0];802; SM100-NEXT: ld.volatile.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];803; SM100-NEXT: ld.param.b64 %rd6, [global_volatile_4xi64_param_1];804; SM100-NEXT: st.volatile.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};805; SM100-NEXT: ret;806 %a.load = load volatile <4 x i64>, ptr addrspace(1) %a807 store volatile <4 x i64> %a.load, ptr addrspace(1) %b808 ret void809}810 811define void @global_volatile_8xfloat(ptr addrspace(1) %a, ptr addrspace(1) %b) {812; SM90-LABEL: global_volatile_8xfloat(813; SM90: {814; SM90-NEXT: .reg .b32 %r<9>;815; SM90-NEXT: .reg .b64 %rd<3>;816; SM90-EMPTY:817; SM90-NEXT: // %bb.0:818; SM90-NEXT: ld.param.b64 %rd1, [global_volatile_8xfloat_param_0];819; SM90-NEXT: ld.volatile.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];820; SM90-NEXT: ld.volatile.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];821; SM90-NEXT: ld.param.b64 %rd2, [global_volatile_8xfloat_param_1];822; SM90-NEXT: st.volatile.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};823; SM90-NEXT: st.volatile.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};824; SM90-NEXT: ret;825;826; SM100-LABEL: global_volatile_8xfloat(827; SM100: {828; SM100-NEXT: .reg .b64 %rd<7>;829; SM100-EMPTY:830; SM100-NEXT: // %bb.0:831; SM100-NEXT: ld.param.b64 %rd1, [global_volatile_8xfloat_param_0];832; SM100-NEXT: ld.volatile.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];833; SM100-NEXT: ld.param.b64 %rd6, [global_volatile_8xfloat_param_1];834; SM100-NEXT: st.volatile.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};835; SM100-NEXT: ret;836 %a.load = load volatile <8 x float>, ptr addrspace(1) %a837 store volatile <8 x float> %a.load, ptr addrspace(1) %b838 ret void839}840 841define void @global_volatile_4xdouble(ptr addrspace(1) %a, ptr addrspace(1) %b) {842; SM90-LABEL: global_volatile_4xdouble(843; SM90: {844; SM90-NEXT: .reg .b64 %rd<7>;845; SM90-EMPTY:846; SM90-NEXT: // %bb.0:847; SM90-NEXT: ld.param.b64 %rd1, [global_volatile_4xdouble_param_0];848; SM90-NEXT: ld.volatile.global.v2.b64 {%rd2, %rd3}, [%rd1];849; SM90-NEXT: ld.volatile.global.v2.b64 {%rd4, %rd5}, [%rd1+16];850; SM90-NEXT: ld.param.b64 %rd6, [global_volatile_4xdouble_param_1];851; SM90-NEXT: st.volatile.global.v2.b64 [%rd6+16], {%rd4, %rd5};852; SM90-NEXT: st.volatile.global.v2.b64 [%rd6], {%rd2, %rd3};853; SM90-NEXT: ret;854;855; SM100-LABEL: global_volatile_4xdouble(856; SM100: {857; SM100-NEXT: .reg .b64 %rd<7>;858; SM100-EMPTY:859; SM100-NEXT: // %bb.0:860; SM100-NEXT: ld.param.b64 %rd1, [global_volatile_4xdouble_param_0];861; SM100-NEXT: ld.volatile.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];862; SM100-NEXT: ld.param.b64 %rd6, [global_volatile_4xdouble_param_1];863; SM100-NEXT: st.volatile.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};864; SM100-NEXT: ret;865 %a.load = load volatile <4 x double>, ptr addrspace(1) %a866 store volatile <4 x double> %a.load, ptr addrspace(1) %b867 ret void868}869 870;; shared statespace871 872; shared873 874define void @shared_32xi8(ptr addrspace(3) %a, ptr addrspace(3) %b) {875; CHECK-LABEL: shared_32xi8(876; CHECK: {877; CHECK-NEXT: .reg .b32 %r<9>;878; CHECK-NEXT: .reg .b64 %rd<3>;879; CHECK-EMPTY:880; CHECK-NEXT: // %bb.0:881; CHECK-NEXT: ld.param.b64 %rd1, [shared_32xi8_param_0];882; CHECK-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];883; CHECK-NEXT: ld.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];884; CHECK-NEXT: ld.param.b64 %rd2, [shared_32xi8_param_1];885; CHECK-NEXT: st.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};886; CHECK-NEXT: st.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};887; CHECK-NEXT: ret;888 %a.load = load <32 x i8>, ptr addrspace(3) %a889 store <32 x i8> %a.load, ptr addrspace(3) %b890 ret void891}892 893define void @shared_16xi16(ptr addrspace(3) %a, ptr addrspace(3) %b) {894; CHECK-LABEL: shared_16xi16(895; CHECK: {896; CHECK-NEXT: .reg .b32 %r<9>;897; CHECK-NEXT: .reg .b64 %rd<3>;898; CHECK-EMPTY:899; CHECK-NEXT: // %bb.0:900; CHECK-NEXT: ld.param.b64 %rd1, [shared_16xi16_param_0];901; CHECK-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];902; CHECK-NEXT: ld.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];903; CHECK-NEXT: ld.param.b64 %rd2, [shared_16xi16_param_1];904; CHECK-NEXT: st.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};905; CHECK-NEXT: st.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};906; CHECK-NEXT: ret;907 %a.load = load <16 x i16>, ptr addrspace(3) %a908 store <16 x i16> %a.load, ptr addrspace(3) %b909 ret void910}911 912define void @shared_16xhalf(ptr addrspace(3) %a, ptr addrspace(3) %b) {913; CHECK-LABEL: shared_16xhalf(914; CHECK: {915; CHECK-NEXT: .reg .b32 %r<9>;916; CHECK-NEXT: .reg .b64 %rd<3>;917; CHECK-EMPTY:918; CHECK-NEXT: // %bb.0:919; CHECK-NEXT: ld.param.b64 %rd1, [shared_16xhalf_param_0];920; CHECK-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];921; CHECK-NEXT: ld.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];922; CHECK-NEXT: ld.param.b64 %rd2, [shared_16xhalf_param_1];923; CHECK-NEXT: st.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};924; CHECK-NEXT: st.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};925; CHECK-NEXT: ret;926 %a.load = load <16 x half>, ptr addrspace(3) %a927 store <16 x half> %a.load, ptr addrspace(3) %b928 ret void929}930 931define void @shared_16xbfloat(ptr addrspace(3) %a, ptr addrspace(3) %b) {932; CHECK-LABEL: shared_16xbfloat(933; CHECK: {934; CHECK-NEXT: .reg .b32 %r<9>;935; CHECK-NEXT: .reg .b64 %rd<3>;936; CHECK-EMPTY:937; CHECK-NEXT: // %bb.0:938; CHECK-NEXT: ld.param.b64 %rd1, [shared_16xbfloat_param_0];939; CHECK-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];940; CHECK-NEXT: ld.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];941; CHECK-NEXT: ld.param.b64 %rd2, [shared_16xbfloat_param_1];942; CHECK-NEXT: st.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};943; CHECK-NEXT: st.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};944; CHECK-NEXT: ret;945 %a.load = load <16 x bfloat>, ptr addrspace(3) %a946 store <16 x bfloat> %a.load, ptr addrspace(3) %b947 ret void948}949 950define void @shared_8xi32(ptr addrspace(3) %a, ptr addrspace(3) %b) {951; SM90-LABEL: shared_8xi32(952; SM90: {953; SM90-NEXT: .reg .b32 %r<9>;954; SM90-NEXT: .reg .b64 %rd<3>;955; SM90-EMPTY:956; SM90-NEXT: // %bb.0:957; SM90-NEXT: ld.param.b64 %rd1, [shared_8xi32_param_0];958; SM90-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];959; SM90-NEXT: ld.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];960; SM90-NEXT: ld.param.b64 %rd2, [shared_8xi32_param_1];961; SM90-NEXT: st.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};962; SM90-NEXT: st.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};963; SM90-NEXT: ret;964;965; SM100-LABEL: shared_8xi32(966; SM100: {967; SM100-NEXT: .reg .b64 %rd<7>;968; SM100-EMPTY:969; SM100-NEXT: // %bb.0:970; SM100-NEXT: ld.param.b64 %rd1, [shared_8xi32_param_0];971; SM100-NEXT: ld.shared.v2.b64 {%rd2, %rd3}, [%rd1];972; SM100-NEXT: ld.shared.v2.b64 {%rd4, %rd5}, [%rd1+16];973; SM100-NEXT: ld.param.b64 %rd6, [shared_8xi32_param_1];974; SM100-NEXT: st.shared.v2.b64 [%rd6+16], {%rd4, %rd5};975; SM100-NEXT: st.shared.v2.b64 [%rd6], {%rd2, %rd3};976; SM100-NEXT: ret;977 %a.load = load <8 x i32>, ptr addrspace(3) %a978 store <8 x i32> %a.load, ptr addrspace(3) %b979 ret void980}981 982define void @shared_4xi64(ptr addrspace(3) %a, ptr addrspace(3) %b) {983; CHECK-LABEL: shared_4xi64(984; CHECK: {985; CHECK-NEXT: .reg .b64 %rd<7>;986; CHECK-EMPTY:987; CHECK-NEXT: // %bb.0:988; CHECK-NEXT: ld.param.b64 %rd1, [shared_4xi64_param_0];989; CHECK-NEXT: ld.shared.v2.b64 {%rd2, %rd3}, [%rd1];990; CHECK-NEXT: ld.shared.v2.b64 {%rd4, %rd5}, [%rd1+16];991; CHECK-NEXT: ld.param.b64 %rd6, [shared_4xi64_param_1];992; CHECK-NEXT: st.shared.v2.b64 [%rd6+16], {%rd4, %rd5};993; CHECK-NEXT: st.shared.v2.b64 [%rd6], {%rd2, %rd3};994; CHECK-NEXT: ret;995 %a.load = load <4 x i64>, ptr addrspace(3) %a996 store <4 x i64> %a.load, ptr addrspace(3) %b997 ret void998}999 1000define void @shared_8xfloat(ptr addrspace(3) %a, ptr addrspace(3) %b) {1001; SM90-LABEL: shared_8xfloat(1002; SM90: {1003; SM90-NEXT: .reg .b32 %r<9>;1004; SM90-NEXT: .reg .b64 %rd<3>;1005; SM90-EMPTY:1006; SM90-NEXT: // %bb.0:1007; SM90-NEXT: ld.param.b64 %rd1, [shared_8xfloat_param_0];1008; SM90-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1009; SM90-NEXT: ld.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1010; SM90-NEXT: ld.param.b64 %rd2, [shared_8xfloat_param_1];1011; SM90-NEXT: st.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1012; SM90-NEXT: st.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1013; SM90-NEXT: ret;1014;1015; SM100-LABEL: shared_8xfloat(1016; SM100: {1017; SM100-NEXT: .reg .b64 %rd<7>;1018; SM100-EMPTY:1019; SM100-NEXT: // %bb.0:1020; SM100-NEXT: ld.param.b64 %rd1, [shared_8xfloat_param_0];1021; SM100-NEXT: ld.shared.v2.b64 {%rd2, %rd3}, [%rd1];1022; SM100-NEXT: ld.shared.v2.b64 {%rd4, %rd5}, [%rd1+16];1023; SM100-NEXT: ld.param.b64 %rd6, [shared_8xfloat_param_1];1024; SM100-NEXT: st.shared.v2.b64 [%rd6+16], {%rd4, %rd5};1025; SM100-NEXT: st.shared.v2.b64 [%rd6], {%rd2, %rd3};1026; SM100-NEXT: ret;1027 %a.load = load <8 x float>, ptr addrspace(3) %a1028 store <8 x float> %a.load, ptr addrspace(3) %b1029 ret void1030}1031 1032define void @shared_4xdouble(ptr addrspace(3) %a, ptr addrspace(3) %b) {1033; CHECK-LABEL: shared_4xdouble(1034; CHECK: {1035; CHECK-NEXT: .reg .b64 %rd<7>;1036; CHECK-EMPTY:1037; CHECK-NEXT: // %bb.0:1038; CHECK-NEXT: ld.param.b64 %rd1, [shared_4xdouble_param_0];1039; CHECK-NEXT: ld.shared.v2.b64 {%rd2, %rd3}, [%rd1];1040; CHECK-NEXT: ld.shared.v2.b64 {%rd4, %rd5}, [%rd1+16];1041; CHECK-NEXT: ld.param.b64 %rd6, [shared_4xdouble_param_1];1042; CHECK-NEXT: st.shared.v2.b64 [%rd6+16], {%rd4, %rd5};1043; CHECK-NEXT: st.shared.v2.b64 [%rd6], {%rd2, %rd3};1044; CHECK-NEXT: ret;1045 %a.load = load <4 x double>, ptr addrspace(3) %a1046 store <4 x double> %a.load, ptr addrspace(3) %b1047 ret void1048}1049 1050; shared_volatile1051 1052define void @shared_volatile_32xi8(ptr addrspace(3) %a, ptr addrspace(3) %b) {1053; CHECK-LABEL: shared_volatile_32xi8(1054; CHECK: {1055; CHECK-NEXT: .reg .b32 %r<9>;1056; CHECK-NEXT: .reg .b64 %rd<3>;1057; CHECK-EMPTY:1058; CHECK-NEXT: // %bb.0:1059; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_32xi8_param_0];1060; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1061; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1062; CHECK-NEXT: ld.param.b64 %rd2, [shared_volatile_32xi8_param_1];1063; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1064; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1065; CHECK-NEXT: ret;1066 %a.load = load volatile <32 x i8>, ptr addrspace(3) %a1067 store volatile <32 x i8> %a.load, ptr addrspace(3) %b1068 ret void1069}1070 1071define void @shared_volatile_16xi16(ptr addrspace(3) %a, ptr addrspace(3) %b) {1072; CHECK-LABEL: shared_volatile_16xi16(1073; CHECK: {1074; CHECK-NEXT: .reg .b32 %r<9>;1075; CHECK-NEXT: .reg .b64 %rd<3>;1076; CHECK-EMPTY:1077; CHECK-NEXT: // %bb.0:1078; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_16xi16_param_0];1079; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1080; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1081; CHECK-NEXT: ld.param.b64 %rd2, [shared_volatile_16xi16_param_1];1082; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1083; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1084; CHECK-NEXT: ret;1085 %a.load = load volatile <16 x i16>, ptr addrspace(3) %a1086 store volatile <16 x i16> %a.load, ptr addrspace(3) %b1087 ret void1088}1089 1090define void @shared_volatile_16xhalf(ptr addrspace(3) %a, ptr addrspace(3) %b) {1091; CHECK-LABEL: shared_volatile_16xhalf(1092; CHECK: {1093; CHECK-NEXT: .reg .b32 %r<9>;1094; CHECK-NEXT: .reg .b64 %rd<3>;1095; CHECK-EMPTY:1096; CHECK-NEXT: // %bb.0:1097; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_16xhalf_param_0];1098; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1099; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1100; CHECK-NEXT: ld.param.b64 %rd2, [shared_volatile_16xhalf_param_1];1101; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1102; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1103; CHECK-NEXT: ret;1104 %a.load = load volatile <16 x half>, ptr addrspace(3) %a1105 store volatile <16 x half> %a.load, ptr addrspace(3) %b1106 ret void1107}1108 1109define void @shared_volatile_16xbfloat(ptr addrspace(3) %a, ptr addrspace(3) %b) {1110; CHECK-LABEL: shared_volatile_16xbfloat(1111; CHECK: {1112; CHECK-NEXT: .reg .b32 %r<9>;1113; CHECK-NEXT: .reg .b64 %rd<3>;1114; CHECK-EMPTY:1115; CHECK-NEXT: // %bb.0:1116; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_16xbfloat_param_0];1117; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1118; CHECK-NEXT: ld.volatile.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1119; CHECK-NEXT: ld.param.b64 %rd2, [shared_volatile_16xbfloat_param_1];1120; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1121; CHECK-NEXT: st.volatile.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1122; CHECK-NEXT: ret;1123 %a.load = load volatile <16 x bfloat>, ptr addrspace(3) %a1124 store volatile <16 x bfloat> %a.load, ptr addrspace(3) %b1125 ret void1126}1127 1128define void @shared_volatile_8xi32(ptr addrspace(3) %a, ptr addrspace(3) %b) {1129; SM90-LABEL: shared_volatile_8xi32(1130; SM90: {1131; SM90-NEXT: .reg .b32 %r<9>;1132; SM90-NEXT: .reg .b64 %rd<3>;1133; SM90-EMPTY:1134; SM90-NEXT: // %bb.0:1135; SM90-NEXT: ld.param.b64 %rd1, [shared_volatile_8xi32_param_0];1136; SM90-NEXT: ld.volatile.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1137; SM90-NEXT: ld.volatile.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1138; SM90-NEXT: ld.param.b64 %rd2, [shared_volatile_8xi32_param_1];1139; SM90-NEXT: st.volatile.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1140; SM90-NEXT: st.volatile.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1141; SM90-NEXT: ret;1142;1143; SM100-LABEL: shared_volatile_8xi32(1144; SM100: {1145; SM100-NEXT: .reg .b64 %rd<7>;1146; SM100-EMPTY:1147; SM100-NEXT: // %bb.0:1148; SM100-NEXT: ld.param.b64 %rd1, [shared_volatile_8xi32_param_0];1149; SM100-NEXT: ld.volatile.shared.v2.b64 {%rd2, %rd3}, [%rd1];1150; SM100-NEXT: ld.volatile.shared.v2.b64 {%rd4, %rd5}, [%rd1+16];1151; SM100-NEXT: ld.param.b64 %rd6, [shared_volatile_8xi32_param_1];1152; SM100-NEXT: st.volatile.shared.v2.b64 [%rd6+16], {%rd4, %rd5};1153; SM100-NEXT: st.volatile.shared.v2.b64 [%rd6], {%rd2, %rd3};1154; SM100-NEXT: ret;1155 %a.load = load volatile <8 x i32>, ptr addrspace(3) %a1156 store volatile <8 x i32> %a.load, ptr addrspace(3) %b1157 ret void1158}1159 1160define void @shared_volatile_4xi64(ptr addrspace(3) %a, ptr addrspace(3) %b) {1161; CHECK-LABEL: shared_volatile_4xi64(1162; CHECK: {1163; CHECK-NEXT: .reg .b64 %rd<7>;1164; CHECK-EMPTY:1165; CHECK-NEXT: // %bb.0:1166; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_4xi64_param_0];1167; CHECK-NEXT: ld.volatile.shared.v2.b64 {%rd2, %rd3}, [%rd1];1168; CHECK-NEXT: ld.volatile.shared.v2.b64 {%rd4, %rd5}, [%rd1+16];1169; CHECK-NEXT: ld.param.b64 %rd6, [shared_volatile_4xi64_param_1];1170; CHECK-NEXT: st.volatile.shared.v2.b64 [%rd6+16], {%rd4, %rd5};1171; CHECK-NEXT: st.volatile.shared.v2.b64 [%rd6], {%rd2, %rd3};1172; CHECK-NEXT: ret;1173 %a.load = load volatile <4 x i64>, ptr addrspace(3) %a1174 store volatile <4 x i64> %a.load, ptr addrspace(3) %b1175 ret void1176}1177 1178define void @shared_volatile_8xfloat(ptr addrspace(3) %a, ptr addrspace(3) %b) {1179; SM90-LABEL: shared_volatile_8xfloat(1180; SM90: {1181; SM90-NEXT: .reg .b32 %r<9>;1182; SM90-NEXT: .reg .b64 %rd<3>;1183; SM90-EMPTY:1184; SM90-NEXT: // %bb.0:1185; SM90-NEXT: ld.param.b64 %rd1, [shared_volatile_8xfloat_param_0];1186; SM90-NEXT: ld.volatile.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1187; SM90-NEXT: ld.volatile.shared.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1188; SM90-NEXT: ld.param.b64 %rd2, [shared_volatile_8xfloat_param_1];1189; SM90-NEXT: st.volatile.shared.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1190; SM90-NEXT: st.volatile.shared.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1191; SM90-NEXT: ret;1192;1193; SM100-LABEL: shared_volatile_8xfloat(1194; SM100: {1195; SM100-NEXT: .reg .b64 %rd<7>;1196; SM100-EMPTY:1197; SM100-NEXT: // %bb.0:1198; SM100-NEXT: ld.param.b64 %rd1, [shared_volatile_8xfloat_param_0];1199; SM100-NEXT: ld.volatile.shared.v2.b64 {%rd2, %rd3}, [%rd1];1200; SM100-NEXT: ld.volatile.shared.v2.b64 {%rd4, %rd5}, [%rd1+16];1201; SM100-NEXT: ld.param.b64 %rd6, [shared_volatile_8xfloat_param_1];1202; SM100-NEXT: st.volatile.shared.v2.b64 [%rd6+16], {%rd4, %rd5};1203; SM100-NEXT: st.volatile.shared.v2.b64 [%rd6], {%rd2, %rd3};1204; SM100-NEXT: ret;1205 %a.load = load volatile <8 x float>, ptr addrspace(3) %a1206 store volatile <8 x float> %a.load, ptr addrspace(3) %b1207 ret void1208}1209 1210define void @shared_volatile_4xdouble(ptr addrspace(3) %a, ptr addrspace(3) %b) {1211; CHECK-LABEL: shared_volatile_4xdouble(1212; CHECK: {1213; CHECK-NEXT: .reg .b64 %rd<7>;1214; CHECK-EMPTY:1215; CHECK-NEXT: // %bb.0:1216; CHECK-NEXT: ld.param.b64 %rd1, [shared_volatile_4xdouble_param_0];1217; CHECK-NEXT: ld.volatile.shared.v2.b64 {%rd2, %rd3}, [%rd1];1218; CHECK-NEXT: ld.volatile.shared.v2.b64 {%rd4, %rd5}, [%rd1+16];1219; CHECK-NEXT: ld.param.b64 %rd6, [shared_volatile_4xdouble_param_1];1220; CHECK-NEXT: st.volatile.shared.v2.b64 [%rd6+16], {%rd4, %rd5};1221; CHECK-NEXT: st.volatile.shared.v2.b64 [%rd6], {%rd2, %rd3};1222; CHECK-NEXT: ret;1223 %a.load = load volatile <4 x double>, ptr addrspace(3) %a1224 store volatile <4 x double> %a.load, ptr addrspace(3) %b1225 ret void1226}1227 1228;; local statespace1229 1230; local1231 1232define void @local_32xi8(ptr addrspace(5) %a, ptr addrspace(5) %b) {1233; CHECK-LABEL: local_32xi8(1234; CHECK: {1235; CHECK-NEXT: .reg .b32 %r<9>;1236; CHECK-NEXT: .reg .b64 %rd<3>;1237; CHECK-EMPTY:1238; CHECK-NEXT: // %bb.0:1239; CHECK-NEXT: ld.param.b64 %rd1, [local_32xi8_param_0];1240; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1241; CHECK-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1242; CHECK-NEXT: ld.param.b64 %rd2, [local_32xi8_param_1];1243; CHECK-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1244; CHECK-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1245; CHECK-NEXT: ret;1246 %a.load = load <32 x i8>, ptr addrspace(5) %a1247 store <32 x i8> %a.load, ptr addrspace(5) %b1248 ret void1249}1250 1251define void @local_16xi16(ptr addrspace(5) %a, ptr addrspace(5) %b) {1252; CHECK-LABEL: local_16xi16(1253; CHECK: {1254; CHECK-NEXT: .reg .b32 %r<9>;1255; CHECK-NEXT: .reg .b64 %rd<3>;1256; CHECK-EMPTY:1257; CHECK-NEXT: // %bb.0:1258; CHECK-NEXT: ld.param.b64 %rd1, [local_16xi16_param_0];1259; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1260; CHECK-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1261; CHECK-NEXT: ld.param.b64 %rd2, [local_16xi16_param_1];1262; CHECK-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1263; CHECK-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1264; CHECK-NEXT: ret;1265 %a.load = load <16 x i16>, ptr addrspace(5) %a1266 store <16 x i16> %a.load, ptr addrspace(5) %b1267 ret void1268}1269 1270define void @local_16xhalf(ptr addrspace(5) %a, ptr addrspace(5) %b) {1271; CHECK-LABEL: local_16xhalf(1272; CHECK: {1273; CHECK-NEXT: .reg .b32 %r<9>;1274; CHECK-NEXT: .reg .b64 %rd<3>;1275; CHECK-EMPTY:1276; CHECK-NEXT: // %bb.0:1277; CHECK-NEXT: ld.param.b64 %rd1, [local_16xhalf_param_0];1278; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1279; CHECK-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1280; CHECK-NEXT: ld.param.b64 %rd2, [local_16xhalf_param_1];1281; CHECK-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1282; CHECK-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1283; CHECK-NEXT: ret;1284 %a.load = load <16 x half>, ptr addrspace(5) %a1285 store <16 x half> %a.load, ptr addrspace(5) %b1286 ret void1287}1288 1289define void @local_16xbfloat(ptr addrspace(5) %a, ptr addrspace(5) %b) {1290; CHECK-LABEL: local_16xbfloat(1291; CHECK: {1292; CHECK-NEXT: .reg .b32 %r<9>;1293; CHECK-NEXT: .reg .b64 %rd<3>;1294; CHECK-EMPTY:1295; CHECK-NEXT: // %bb.0:1296; CHECK-NEXT: ld.param.b64 %rd1, [local_16xbfloat_param_0];1297; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1298; CHECK-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1299; CHECK-NEXT: ld.param.b64 %rd2, [local_16xbfloat_param_1];1300; CHECK-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1301; CHECK-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1302; CHECK-NEXT: ret;1303 %a.load = load <16 x bfloat>, ptr addrspace(5) %a1304 store <16 x bfloat> %a.load, ptr addrspace(5) %b1305 ret void1306}1307 1308define void @local_8xi32(ptr addrspace(5) %a, ptr addrspace(5) %b) {1309; SM90-LABEL: local_8xi32(1310; SM90: {1311; SM90-NEXT: .reg .b32 %r<9>;1312; SM90-NEXT: .reg .b64 %rd<3>;1313; SM90-EMPTY:1314; SM90-NEXT: // %bb.0:1315; SM90-NEXT: ld.param.b64 %rd1, [local_8xi32_param_0];1316; SM90-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1317; SM90-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1318; SM90-NEXT: ld.param.b64 %rd2, [local_8xi32_param_1];1319; SM90-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1320; SM90-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1321; SM90-NEXT: ret;1322;1323; SM100-LABEL: local_8xi32(1324; SM100: {1325; SM100-NEXT: .reg .b64 %rd<7>;1326; SM100-EMPTY:1327; SM100-NEXT: // %bb.0:1328; SM100-NEXT: ld.param.b64 %rd1, [local_8xi32_param_0];1329; SM100-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];1330; SM100-NEXT: ld.local.v2.b64 {%rd4, %rd5}, [%rd1+16];1331; SM100-NEXT: ld.param.b64 %rd6, [local_8xi32_param_1];1332; SM100-NEXT: st.local.v2.b64 [%rd6+16], {%rd4, %rd5};1333; SM100-NEXT: st.local.v2.b64 [%rd6], {%rd2, %rd3};1334; SM100-NEXT: ret;1335 %a.load = load <8 x i32>, ptr addrspace(5) %a1336 store <8 x i32> %a.load, ptr addrspace(5) %b1337 ret void1338}1339 1340define void @local_4xi64(ptr addrspace(5) %a, ptr addrspace(5) %b) {1341; CHECK-LABEL: local_4xi64(1342; CHECK: {1343; CHECK-NEXT: .reg .b64 %rd<7>;1344; CHECK-EMPTY:1345; CHECK-NEXT: // %bb.0:1346; CHECK-NEXT: ld.param.b64 %rd1, [local_4xi64_param_0];1347; CHECK-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];1348; CHECK-NEXT: ld.local.v2.b64 {%rd4, %rd5}, [%rd1+16];1349; CHECK-NEXT: ld.param.b64 %rd6, [local_4xi64_param_1];1350; CHECK-NEXT: st.local.v2.b64 [%rd6+16], {%rd4, %rd5};1351; CHECK-NEXT: st.local.v2.b64 [%rd6], {%rd2, %rd3};1352; CHECK-NEXT: ret;1353 %a.load = load <4 x i64>, ptr addrspace(5) %a1354 store <4 x i64> %a.load, ptr addrspace(5) %b1355 ret void1356}1357 1358define void @local_8xfloat(ptr addrspace(5) %a, ptr addrspace(5) %b) {1359; SM90-LABEL: local_8xfloat(1360; SM90: {1361; SM90-NEXT: .reg .b32 %r<9>;1362; SM90-NEXT: .reg .b64 %rd<3>;1363; SM90-EMPTY:1364; SM90-NEXT: // %bb.0:1365; SM90-NEXT: ld.param.b64 %rd1, [local_8xfloat_param_0];1366; SM90-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1367; SM90-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1368; SM90-NEXT: ld.param.b64 %rd2, [local_8xfloat_param_1];1369; SM90-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1370; SM90-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1371; SM90-NEXT: ret;1372;1373; SM100-LABEL: local_8xfloat(1374; SM100: {1375; SM100-NEXT: .reg .b64 %rd<7>;1376; SM100-EMPTY:1377; SM100-NEXT: // %bb.0:1378; SM100-NEXT: ld.param.b64 %rd1, [local_8xfloat_param_0];1379; SM100-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];1380; SM100-NEXT: ld.local.v2.b64 {%rd4, %rd5}, [%rd1+16];1381; SM100-NEXT: ld.param.b64 %rd6, [local_8xfloat_param_1];1382; SM100-NEXT: st.local.v2.b64 [%rd6+16], {%rd4, %rd5};1383; SM100-NEXT: st.local.v2.b64 [%rd6], {%rd2, %rd3};1384; SM100-NEXT: ret;1385 %a.load = load <8 x float>, ptr addrspace(5) %a1386 store <8 x float> %a.load, ptr addrspace(5) %b1387 ret void1388}1389 1390define void @local_4xdouble(ptr addrspace(5) %a, ptr addrspace(5) %b) {1391; CHECK-LABEL: local_4xdouble(1392; CHECK: {1393; CHECK-NEXT: .reg .b64 %rd<7>;1394; CHECK-EMPTY:1395; CHECK-NEXT: // %bb.0:1396; CHECK-NEXT: ld.param.b64 %rd1, [local_4xdouble_param_0];1397; CHECK-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];1398; CHECK-NEXT: ld.local.v2.b64 {%rd4, %rd5}, [%rd1+16];1399; CHECK-NEXT: ld.param.b64 %rd6, [local_4xdouble_param_1];1400; CHECK-NEXT: st.local.v2.b64 [%rd6+16], {%rd4, %rd5};1401; CHECK-NEXT: st.local.v2.b64 [%rd6], {%rd2, %rd3};1402; CHECK-NEXT: ret;1403 %a.load = load <4 x double>, ptr addrspace(5) %a1404 store <4 x double> %a.load, ptr addrspace(5) %b1405 ret void1406}1407 1408; local_volatile1409 1410define void @local_volatile_32xi8(ptr addrspace(5) %a, ptr addrspace(5) %b) {1411; CHECK-LABEL: local_volatile_32xi8(1412; CHECK: {1413; CHECK-NEXT: .reg .b32 %r<9>;1414; CHECK-NEXT: .reg .b64 %rd<3>;1415; CHECK-EMPTY:1416; CHECK-NEXT: // %bb.0:1417; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_32xi8_param_0];1418; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1419; CHECK-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1420; CHECK-NEXT: ld.param.b64 %rd2, [local_volatile_32xi8_param_1];1421; CHECK-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1422; CHECK-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1423; CHECK-NEXT: ret;1424 %a.load = load volatile <32 x i8>, ptr addrspace(5) %a1425 store volatile <32 x i8> %a.load, ptr addrspace(5) %b1426 ret void1427}1428 1429define void @local_volatile_16xi16(ptr addrspace(5) %a, ptr addrspace(5) %b) {1430; CHECK-LABEL: local_volatile_16xi16(1431; CHECK: {1432; CHECK-NEXT: .reg .b32 %r<9>;1433; CHECK-NEXT: .reg .b64 %rd<3>;1434; CHECK-EMPTY:1435; CHECK-NEXT: // %bb.0:1436; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_16xi16_param_0];1437; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1438; CHECK-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1439; CHECK-NEXT: ld.param.b64 %rd2, [local_volatile_16xi16_param_1];1440; CHECK-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1441; CHECK-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1442; CHECK-NEXT: ret;1443 %a.load = load volatile <16 x i16>, ptr addrspace(5) %a1444 store volatile <16 x i16> %a.load, ptr addrspace(5) %b1445 ret void1446}1447 1448define void @local_volatile_16xhalf(ptr addrspace(5) %a, ptr addrspace(5) %b) {1449; CHECK-LABEL: local_volatile_16xhalf(1450; CHECK: {1451; CHECK-NEXT: .reg .b32 %r<9>;1452; CHECK-NEXT: .reg .b64 %rd<3>;1453; CHECK-EMPTY:1454; CHECK-NEXT: // %bb.0:1455; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_16xhalf_param_0];1456; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1457; CHECK-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1458; CHECK-NEXT: ld.param.b64 %rd2, [local_volatile_16xhalf_param_1];1459; CHECK-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1460; CHECK-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1461; CHECK-NEXT: ret;1462 %a.load = load volatile <16 x half>, ptr addrspace(5) %a1463 store volatile <16 x half> %a.load, ptr addrspace(5) %b1464 ret void1465}1466 1467define void @local_volatile_16xbfloat(ptr addrspace(5) %a, ptr addrspace(5) %b) {1468; CHECK-LABEL: local_volatile_16xbfloat(1469; CHECK: {1470; CHECK-NEXT: .reg .b32 %r<9>;1471; CHECK-NEXT: .reg .b64 %rd<3>;1472; CHECK-EMPTY:1473; CHECK-NEXT: // %bb.0:1474; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_16xbfloat_param_0];1475; CHECK-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1476; CHECK-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1477; CHECK-NEXT: ld.param.b64 %rd2, [local_volatile_16xbfloat_param_1];1478; CHECK-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1479; CHECK-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1480; CHECK-NEXT: ret;1481 %a.load = load volatile <16 x bfloat>, ptr addrspace(5) %a1482 store volatile <16 x bfloat> %a.load, ptr addrspace(5) %b1483 ret void1484}1485 1486define void @local_volatile_8xi32(ptr addrspace(5) %a, ptr addrspace(5) %b) {1487; SM90-LABEL: local_volatile_8xi32(1488; SM90: {1489; SM90-NEXT: .reg .b32 %r<9>;1490; SM90-NEXT: .reg .b64 %rd<3>;1491; SM90-EMPTY:1492; SM90-NEXT: // %bb.0:1493; SM90-NEXT: ld.param.b64 %rd1, [local_volatile_8xi32_param_0];1494; SM90-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1495; SM90-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1496; SM90-NEXT: ld.param.b64 %rd2, [local_volatile_8xi32_param_1];1497; SM90-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1498; SM90-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1499; SM90-NEXT: ret;1500;1501; SM100-LABEL: local_volatile_8xi32(1502; SM100: {1503; SM100-NEXT: .reg .b64 %rd<7>;1504; SM100-EMPTY:1505; SM100-NEXT: // %bb.0:1506; SM100-NEXT: ld.param.b64 %rd1, [local_volatile_8xi32_param_0];1507; SM100-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];1508; SM100-NEXT: ld.local.v2.b64 {%rd4, %rd5}, [%rd1+16];1509; SM100-NEXT: ld.param.b64 %rd6, [local_volatile_8xi32_param_1];1510; SM100-NEXT: st.local.v2.b64 [%rd6+16], {%rd4, %rd5};1511; SM100-NEXT: st.local.v2.b64 [%rd6], {%rd2, %rd3};1512; SM100-NEXT: ret;1513 %a.load = load volatile <8 x i32>, ptr addrspace(5) %a1514 store volatile <8 x i32> %a.load, ptr addrspace(5) %b1515 ret void1516}1517 1518define void @local_volatile_4xi64(ptr addrspace(5) %a, ptr addrspace(5) %b) {1519; CHECK-LABEL: local_volatile_4xi64(1520; CHECK: {1521; CHECK-NEXT: .reg .b64 %rd<7>;1522; CHECK-EMPTY:1523; CHECK-NEXT: // %bb.0:1524; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_4xi64_param_0];1525; CHECK-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];1526; CHECK-NEXT: ld.local.v2.b64 {%rd4, %rd5}, [%rd1+16];1527; CHECK-NEXT: ld.param.b64 %rd6, [local_volatile_4xi64_param_1];1528; CHECK-NEXT: st.local.v2.b64 [%rd6+16], {%rd4, %rd5};1529; CHECK-NEXT: st.local.v2.b64 [%rd6], {%rd2, %rd3};1530; CHECK-NEXT: ret;1531 %a.load = load volatile <4 x i64>, ptr addrspace(5) %a1532 store volatile <4 x i64> %a.load, ptr addrspace(5) %b1533 ret void1534}1535 1536define void @local_volatile_8xfloat(ptr addrspace(5) %a, ptr addrspace(5) %b) {1537; SM90-LABEL: local_volatile_8xfloat(1538; SM90: {1539; SM90-NEXT: .reg .b32 %r<9>;1540; SM90-NEXT: .reg .b64 %rd<3>;1541; SM90-EMPTY:1542; SM90-NEXT: // %bb.0:1543; SM90-NEXT: ld.param.b64 %rd1, [local_volatile_8xfloat_param_0];1544; SM90-NEXT: ld.local.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];1545; SM90-NEXT: ld.local.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];1546; SM90-NEXT: ld.param.b64 %rd2, [local_volatile_8xfloat_param_1];1547; SM90-NEXT: st.local.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};1548; SM90-NEXT: st.local.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};1549; SM90-NEXT: ret;1550;1551; SM100-LABEL: local_volatile_8xfloat(1552; SM100: {1553; SM100-NEXT: .reg .b64 %rd<7>;1554; SM100-EMPTY:1555; SM100-NEXT: // %bb.0:1556; SM100-NEXT: ld.param.b64 %rd1, [local_volatile_8xfloat_param_0];1557; SM100-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];1558; SM100-NEXT: ld.local.v2.b64 {%rd4, %rd5}, [%rd1+16];1559; SM100-NEXT: ld.param.b64 %rd6, [local_volatile_8xfloat_param_1];1560; SM100-NEXT: st.local.v2.b64 [%rd6+16], {%rd4, %rd5};1561; SM100-NEXT: st.local.v2.b64 [%rd6], {%rd2, %rd3};1562; SM100-NEXT: ret;1563 %a.load = load volatile <8 x float>, ptr addrspace(5) %a1564 store volatile <8 x float> %a.load, ptr addrspace(5) %b1565 ret void1566}1567 1568define void @local_volatile_4xdouble(ptr addrspace(5) %a, ptr addrspace(5) %b) {1569; CHECK-LABEL: local_volatile_4xdouble(1570; CHECK: {1571; CHECK-NEXT: .reg .b64 %rd<7>;1572; CHECK-EMPTY:1573; CHECK-NEXT: // %bb.0:1574; CHECK-NEXT: ld.param.b64 %rd1, [local_volatile_4xdouble_param_0];1575; CHECK-NEXT: ld.local.v2.b64 {%rd2, %rd3}, [%rd1];1576; CHECK-NEXT: ld.local.v2.b64 {%rd4, %rd5}, [%rd1+16];1577; CHECK-NEXT: ld.param.b64 %rd6, [local_volatile_4xdouble_param_1];1578; CHECK-NEXT: st.local.v2.b64 [%rd6+16], {%rd4, %rd5};1579; CHECK-NEXT: st.local.v2.b64 [%rd6], {%rd2, %rd3};1580; CHECK-NEXT: ret;1581 %a.load = load volatile <4 x double>, ptr addrspace(5) %a1582 store volatile <4 x double> %a.load, ptr addrspace(5) %b1583 ret void1584}1585 1586define void @test_i256_global(ptr addrspace(1) %a, ptr addrspace(1) %b) {1587; SM90-LABEL: test_i256_global(1588; SM90: {1589; SM90-NEXT: .reg .b64 %rd<7>;1590; SM90-EMPTY:1591; SM90-NEXT: // %bb.0:1592; SM90-NEXT: ld.param.b64 %rd1, [test_i256_global_param_0];1593; SM90-NEXT: ld.global.v2.b64 {%rd2, %rd3}, [%rd1];1594; SM90-NEXT: ld.global.v2.b64 {%rd4, %rd5}, [%rd1+16];1595; SM90-NEXT: ld.param.b64 %rd6, [test_i256_global_param_1];1596; SM90-NEXT: st.global.v2.b64 [%rd6+16], {%rd4, %rd5};1597; SM90-NEXT: st.global.v2.b64 [%rd6], {%rd2, %rd3};1598; SM90-NEXT: ret;1599;1600; SM100-LABEL: test_i256_global(1601; SM100: {1602; SM100-NEXT: .reg .b64 %rd<7>;1603; SM100-EMPTY:1604; SM100-NEXT: // %bb.0:1605; SM100-NEXT: ld.param.b64 %rd1, [test_i256_global_param_0];1606; SM100-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];1607; SM100-NEXT: ld.param.b64 %rd6, [test_i256_global_param_1];1608; SM100-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};1609; SM100-NEXT: ret;1610 %a.load = load i256, ptr addrspace(1) %a, align 321611 store i256 %a.load, ptr addrspace(1) %b, align 321612 ret void1613}1614 1615 1616define void @test_i256_global_unaligned(ptr addrspace(1) %a, ptr addrspace(1) %b) {1617; CHECK-LABEL: test_i256_global_unaligned(1618; CHECK: {1619; CHECK-NEXT: .reg .b64 %rd<7>;1620; CHECK-EMPTY:1621; CHECK-NEXT: // %bb.0:1622; CHECK-NEXT: ld.param.b64 %rd1, [test_i256_global_unaligned_param_0];1623; CHECK-NEXT: ld.global.v2.b64 {%rd2, %rd3}, [%rd1];1624; CHECK-NEXT: ld.global.v2.b64 {%rd4, %rd5}, [%rd1+16];1625; CHECK-NEXT: ld.param.b64 %rd6, [test_i256_global_unaligned_param_1];1626; CHECK-NEXT: st.global.v2.b64 [%rd6+16], {%rd4, %rd5};1627; CHECK-NEXT: st.global.v2.b64 [%rd6], {%rd2, %rd3};1628; CHECK-NEXT: ret;1629 %a.load = load i256, ptr addrspace(1) %a, align 161630 store i256 %a.load, ptr addrspace(1) %b, align 161631 ret void1632}1633 1634define void @test_i256_generic(ptr %a, ptr %b) {1635; CHECK-LABEL: test_i256_generic(1636; CHECK: {1637; CHECK-NEXT: .reg .b64 %rd<7>;1638; CHECK-EMPTY:1639; CHECK-NEXT: // %bb.0:1640; CHECK-NEXT: ld.param.b64 %rd1, [test_i256_generic_param_0];1641; CHECK-NEXT: ld.v2.b64 {%rd2, %rd3}, [%rd1];1642; CHECK-NEXT: ld.v2.b64 {%rd4, %rd5}, [%rd1+16];1643; CHECK-NEXT: ld.param.b64 %rd6, [test_i256_generic_param_1];1644; CHECK-NEXT: st.v2.b64 [%rd6+16], {%rd4, %rd5};1645; CHECK-NEXT: st.v2.b64 [%rd6], {%rd2, %rd3};1646; CHECK-NEXT: ret;1647 %a.load = load i256, ptr %a, align 321648 store i256 %a.load, ptr %b, align 321649 ret void1650}1651 1652define void @test_i256_global_volatile(ptr addrspace(1) %a, ptr addrspace(1) %b) {1653; SM90-LABEL: test_i256_global_volatile(1654; SM90: {1655; SM90-NEXT: .reg .b64 %rd<7>;1656; SM90-EMPTY:1657; SM90-NEXT: // %bb.0:1658; SM90-NEXT: ld.param.b64 %rd1, [test_i256_global_volatile_param_0];1659; SM90-NEXT: ld.volatile.global.v2.b64 {%rd2, %rd3}, [%rd1];1660; SM90-NEXT: ld.volatile.global.v2.b64 {%rd4, %rd5}, [%rd1+16];1661; SM90-NEXT: ld.param.b64 %rd6, [test_i256_global_volatile_param_1];1662; SM90-NEXT: st.volatile.global.v2.b64 [%rd6+16], {%rd4, %rd5};1663; SM90-NEXT: st.volatile.global.v2.b64 [%rd6], {%rd2, %rd3};1664; SM90-NEXT: ret;1665;1666; SM100-LABEL: test_i256_global_volatile(1667; SM100: {1668; SM100-NEXT: .reg .b64 %rd<7>;1669; SM100-EMPTY:1670; SM100-NEXT: // %bb.0:1671; SM100-NEXT: ld.param.b64 %rd1, [test_i256_global_volatile_param_0];1672; SM100-NEXT: ld.volatile.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];1673; SM100-NEXT: ld.param.b64 %rd6, [test_i256_global_volatile_param_1];1674; SM100-NEXT: st.volatile.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};1675; SM100-NEXT: ret;1676 %a.load = load volatile i256, ptr addrspace(1) %a, align 321677 store volatile i256 %a.load, ptr addrspace(1) %b, align 321678 ret void1679}1680