367 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=+ptx88 | FileCheck %s -check-prefixes=CHECK,SM903; RUN: %if ptxas-sm_90 && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx88 | %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 8; Different architectures are tested in this file for the following reasons:9; - SM90 does not have 256-bit load/store instructions10; - SM90 does not have masked store instructions11; - SM90 does not support packed f32x2 instructions12 13define void @global_8xi32(ptr addrspace(1) %a, ptr addrspace(1) %b) {14; SM90-LABEL: global_8xi32(15; SM90: {16; SM90-NEXT: .reg .b32 %r<9>;17; SM90-NEXT: .reg .b64 %rd<3>;18; SM90-EMPTY:19; SM90-NEXT: // %bb.0:20; SM90-NEXT: ld.param.b64 %rd1, [global_8xi32_param_0];21; SM90-NEXT: .pragma "used_bytes_mask 0xf000";22; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];23; SM90-NEXT: .pragma "used_bytes_mask 0xf0f";24; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];25; SM90-NEXT: ld.param.b64 %rd2, [global_8xi32_param_1];26; SM90-NEXT: st.global.b32 [%rd2], %r5;27; SM90-NEXT: st.global.b32 [%rd2+8], %r7;28; SM90-NEXT: st.global.b32 [%rd2+28], %r4;29; SM90-NEXT: ret;30;31; SM100-LABEL: global_8xi32(32; SM100: {33; SM100-NEXT: .reg .b32 %r<9>;34; SM100-NEXT: .reg .b64 %rd<3>;35; SM100-EMPTY:36; SM100-NEXT: // %bb.0:37; SM100-NEXT: ld.param.b64 %rd1, [global_8xi32_param_0];38; SM100-NEXT: .pragma "used_bytes_mask 0xf0000f0f";39; SM100-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];40; SM100-NEXT: ld.param.b64 %rd2, [global_8xi32_param_1];41; SM100-NEXT: st.global.v8.b32 [%rd2], {%r1, _, %r3, _, _, _, _, %r8};42; SM100-NEXT: ret;43 %a.load = tail call <8 x i32> @llvm.masked.load.v8i32.p1(ptr addrspace(1) align 32 %a, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 false, i1 false, i1 false, i1 true>, <8 x i32> poison)44 tail call void @llvm.masked.store.v8i32.p1(<8 x i32> %a.load, ptr addrspace(1) align 32 %b, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 false, i1 false, i1 false, i1 true>)45 ret void46}47 48; Masked stores are only supported for 32-bit element types,49; while masked loads are supported for all element types.50define void @global_16xi16(ptr addrspace(1) %a, ptr addrspace(1) %b) {51; SM90-LABEL: global_16xi16(52; SM90: {53; SM90-NEXT: .reg .b16 %rs<7>;54; SM90-NEXT: .reg .b32 %r<9>;55; SM90-NEXT: .reg .b64 %rd<3>;56; SM90-EMPTY:57; SM90-NEXT: // %bb.0:58; SM90-NEXT: ld.param.b64 %rd1, [global_16xi16_param_0];59; SM90-NEXT: .pragma "used_bytes_mask 0xf000";60; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];61; SM90-NEXT: mov.b32 {%rs1, %rs2}, %r4;62; SM90-NEXT: .pragma "used_bytes_mask 0xf0f";63; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];64; SM90-NEXT: mov.b32 {%rs3, %rs4}, %r7;65; SM90-NEXT: mov.b32 {%rs5, %rs6}, %r5;66; SM90-NEXT: ld.param.b64 %rd2, [global_16xi16_param_1];67; SM90-NEXT: st.global.b16 [%rd2], %rs5;68; SM90-NEXT: st.global.b16 [%rd2+2], %rs6;69; SM90-NEXT: st.global.b16 [%rd2+8], %rs3;70; SM90-NEXT: st.global.b16 [%rd2+10], %rs4;71; SM90-NEXT: st.global.b16 [%rd2+28], %rs1;72; SM90-NEXT: st.global.b16 [%rd2+30], %rs2;73; SM90-NEXT: ret;74;75; SM100-LABEL: global_16xi16(76; SM100: {77; SM100-NEXT: .reg .b16 %rs<7>;78; SM100-NEXT: .reg .b32 %r<9>;79; SM100-NEXT: .reg .b64 %rd<3>;80; SM100-EMPTY:81; SM100-NEXT: // %bb.0:82; SM100-NEXT: ld.param.b64 %rd1, [global_16xi16_param_0];83; SM100-NEXT: .pragma "used_bytes_mask 0xf0000f0f";84; SM100-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];85; SM100-NEXT: mov.b32 {%rs1, %rs2}, %r8;86; SM100-NEXT: mov.b32 {%rs3, %rs4}, %r3;87; SM100-NEXT: mov.b32 {%rs5, %rs6}, %r1;88; SM100-NEXT: ld.param.b64 %rd2, [global_16xi16_param_1];89; SM100-NEXT: st.global.b16 [%rd2], %rs5;90; SM100-NEXT: st.global.b16 [%rd2+2], %rs6;91; SM100-NEXT: st.global.b16 [%rd2+8], %rs3;92; SM100-NEXT: st.global.b16 [%rd2+10], %rs4;93; SM100-NEXT: st.global.b16 [%rd2+28], %rs1;94; SM100-NEXT: st.global.b16 [%rd2+30], %rs2;95; SM100-NEXT: ret;96 %a.load = tail call <16 x i16> @llvm.masked.load.v16i16.p1(ptr addrspace(1) align 32 %a, <16 x i1> <i1 true, i1 true, i1 false, i1 false, i1 true, i1 true, i1 false, i1 false, i1 false, i1 false, i1 false, i1 false, i1 false, i1 false, i1 true, i1 true>, <16 x i16> poison)97 tail call void @llvm.masked.store.v16i16.p1(<16 x i16> %a.load, ptr addrspace(1) align 32 %b, <16 x i1> <i1 true, i1 true, i1 false, i1 false, i1 true, i1 true, i1 false, i1 false, i1 false, i1 false, i1 false, i1 false, i1 false, i1 false, i1 true, i1 true>)98 ret void99}100 101define void @global_8xi32_no_align(ptr addrspace(1) %a, ptr addrspace(1) %b) {102; CHECK-LABEL: global_8xi32_no_align(103; CHECK: {104; CHECK-NEXT: .reg .b32 %r<4>;105; CHECK-NEXT: .reg .b64 %rd<3>;106; CHECK-EMPTY:107; CHECK-NEXT: // %bb.0:108; CHECK-NEXT: ld.param.b64 %rd1, [global_8xi32_no_align_param_0];109; CHECK-NEXT: ld.global.b32 %r1, [%rd1];110; CHECK-NEXT: ld.param.b64 %rd2, [global_8xi32_no_align_param_1];111; CHECK-NEXT: ld.global.b32 %r2, [%rd1+8];112; CHECK-NEXT: ld.global.b32 %r3, [%rd1+28];113; CHECK-NEXT: st.global.b32 [%rd2], %r1;114; CHECK-NEXT: st.global.b32 [%rd2+8], %r2;115; CHECK-NEXT: st.global.b32 [%rd2+28], %r3;116; CHECK-NEXT: ret;117 %a.load = tail call <8 x i32> @llvm.masked.load.v8i32.p1(ptr addrspace(1) align 16 %a, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 false, i1 false, i1 false, i1 true>, <8 x i32> poison)118 tail call void @llvm.masked.store.v8i32.p1(<8 x i32> %a.load, ptr addrspace(1) align 16 %b, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 false, i1 false, i1 false, i1 true>)119 ret void120}121 122 123define void @global_8xi32_invariant(ptr addrspace(1) %a, ptr addrspace(1) %b) {124; SM90-LABEL: global_8xi32_invariant(125; SM90: {126; SM90-NEXT: .reg .b32 %r<9>;127; SM90-NEXT: .reg .b64 %rd<3>;128; SM90-EMPTY:129; SM90-NEXT: // %bb.0:130; SM90-NEXT: ld.param.b64 %rd1, [global_8xi32_invariant_param_0];131; SM90-NEXT: .pragma "used_bytes_mask 0xf000";132; SM90-NEXT: ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];133; SM90-NEXT: .pragma "used_bytes_mask 0xf0f";134; SM90-NEXT: ld.global.nc.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];135; SM90-NEXT: ld.param.b64 %rd2, [global_8xi32_invariant_param_1];136; SM90-NEXT: st.global.b32 [%rd2], %r5;137; SM90-NEXT: st.global.b32 [%rd2+8], %r7;138; SM90-NEXT: st.global.b32 [%rd2+28], %r4;139; SM90-NEXT: ret;140;141; SM100-LABEL: global_8xi32_invariant(142; SM100: {143; SM100-NEXT: .reg .b32 %r<9>;144; SM100-NEXT: .reg .b64 %rd<3>;145; SM100-EMPTY:146; SM100-NEXT: // %bb.0:147; SM100-NEXT: ld.param.b64 %rd1, [global_8xi32_invariant_param_0];148; SM100-NEXT: .pragma "used_bytes_mask 0xf0000f0f";149; SM100-NEXT: ld.global.nc.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];150; SM100-NEXT: ld.param.b64 %rd2, [global_8xi32_invariant_param_1];151; SM100-NEXT: st.global.v8.b32 [%rd2], {%r1, _, %r3, _, _, _, _, %r8};152; SM100-NEXT: ret;153 %a.load = tail call <8 x i32> @llvm.masked.load.v8i32.p1(ptr addrspace(1) align 32 %a, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 false, i1 false, i1 false, i1 true>, <8 x i32> poison), !invariant.load !0154 tail call void @llvm.masked.store.v8i32.p1(<8 x i32> %a.load, ptr addrspace(1) align 32 %b, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 false, i1 false, i1 false, i1 true>)155 ret void156}157 158define void @global_2xi16(ptr addrspace(1) %a, ptr addrspace(1) %b) {159; CHECK-LABEL: global_2xi16(160; CHECK: {161; CHECK-NEXT: .reg .b16 %rs<2>;162; CHECK-NEXT: .reg .b32 %r<2>;163; CHECK-NEXT: .reg .b64 %rd<3>;164; CHECK-EMPTY:165; CHECK-NEXT: // %bb.0:166; CHECK-NEXT: ld.param.b64 %rd1, [global_2xi16_param_0];167; CHECK-NEXT: .pragma "used_bytes_mask 0x3";168; CHECK-NEXT: ld.global.b32 %r1, [%rd1];169; CHECK-NEXT: ld.param.b64 %rd2, [global_2xi16_param_1];170; CHECK-NEXT: mov.b32 {%rs1, _}, %r1;171; CHECK-NEXT: st.global.b16 [%rd2], %rs1;172; CHECK-NEXT: ret;173 %a.load = tail call <2 x i16> @llvm.masked.load.v2i16.p1(ptr addrspace(1) align 4 %a, <2 x i1> <i1 true, i1 false>, <2 x i16> poison)174 tail call void @llvm.masked.store.v2i16.p1(<2 x i16> %a.load, ptr addrspace(1) align 4 %b, <2 x i1> <i1 true, i1 false>)175 ret void176}177 178define void @global_2xi16_invariant(ptr addrspace(1) %a, ptr addrspace(1) %b) {179; CHECK-LABEL: global_2xi16_invariant(180; CHECK: {181; CHECK-NEXT: .reg .b16 %rs<2>;182; CHECK-NEXT: .reg .b32 %r<2>;183; CHECK-NEXT: .reg .b64 %rd<3>;184; CHECK-EMPTY:185; CHECK-NEXT: // %bb.0:186; CHECK-NEXT: ld.param.b64 %rd1, [global_2xi16_invariant_param_0];187; CHECK-NEXT: .pragma "used_bytes_mask 0x3";188; CHECK-NEXT: ld.global.nc.b32 %r1, [%rd1];189; CHECK-NEXT: ld.param.b64 %rd2, [global_2xi16_invariant_param_1];190; CHECK-NEXT: mov.b32 {%rs1, _}, %r1;191; CHECK-NEXT: st.global.b16 [%rd2], %rs1;192; CHECK-NEXT: ret;193 %a.load = tail call <2 x i16> @llvm.masked.load.v2i16.p1(ptr addrspace(1) align 4 %a, <2 x i1> <i1 true, i1 false>, <2 x i16> poison), !invariant.load !0194 tail call void @llvm.masked.store.v2i16.p1(<2 x i16> %a.load, ptr addrspace(1) align 4 %b, <2 x i1> <i1 true, i1 false>)195 ret void196}197 198define void @global_2xi16_no_align(ptr addrspace(1) %a, ptr addrspace(1) %b) {199; CHECK-LABEL: global_2xi16_no_align(200; CHECK: {201; CHECK-NEXT: .reg .b16 %rs<2>;202; CHECK-NEXT: .reg .b64 %rd<3>;203; CHECK-EMPTY:204; CHECK-NEXT: // %bb.0:205; CHECK-NEXT: ld.param.b64 %rd1, [global_2xi16_no_align_param_0];206; CHECK-NEXT: ld.global.b16 %rs1, [%rd1];207; CHECK-NEXT: ld.param.b64 %rd2, [global_2xi16_no_align_param_1];208; CHECK-NEXT: st.global.b16 [%rd2], %rs1;209; CHECK-NEXT: ret;210 %a.load = tail call <2 x i16> @llvm.masked.load.v2i16.p1(ptr addrspace(1) align 2 %a, <2 x i1> <i1 true, i1 false>, <2 x i16> poison)211 tail call void @llvm.masked.store.v2i16.p1(<2 x i16> %a.load, ptr addrspace(1) align 4 %b, <2 x i1> <i1 true, i1 false>)212 ret void213}214 215define void @global_4xi8(ptr addrspace(1) %a, ptr addrspace(1) %b) {216; CHECK-LABEL: global_4xi8(217; CHECK: {218; CHECK-NEXT: .reg .b32 %r<3>;219; CHECK-NEXT: .reg .b64 %rd<3>;220; CHECK-EMPTY:221; CHECK-NEXT: // %bb.0:222; CHECK-NEXT: ld.param.b64 %rd1, [global_4xi8_param_0];223; CHECK-NEXT: .pragma "used_bytes_mask 0x5";224; CHECK-NEXT: ld.global.b32 %r1, [%rd1];225; CHECK-NEXT: ld.param.b64 %rd2, [global_4xi8_param_1];226; CHECK-NEXT: st.global.b8 [%rd2], %r1;227; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U;228; CHECK-NEXT: st.global.b8 [%rd2+2], %r2;229; CHECK-NEXT: ret;230 %a.load = tail call <4 x i8> @llvm.masked.load.v4i8.p1(ptr addrspace(1) align 4 %a, <4 x i1> <i1 true, i1 false, i1 true, i1 false>, <4 x i8> poison)231 tail call void @llvm.masked.store.v4i8.p1(<4 x i8> %a.load, ptr addrspace(1) align 4 %b, <4 x i1> <i1 true, i1 false, i1 true, i1 false>)232 ret void233}234 235define void @global_4xi8_invariant(ptr addrspace(1) %a, ptr addrspace(1) %b) {236; CHECK-LABEL: global_4xi8_invariant(237; CHECK: {238; CHECK-NEXT: .reg .b32 %r<3>;239; CHECK-NEXT: .reg .b64 %rd<3>;240; CHECK-EMPTY:241; CHECK-NEXT: // %bb.0:242; CHECK-NEXT: ld.param.b64 %rd1, [global_4xi8_invariant_param_0];243; CHECK-NEXT: .pragma "used_bytes_mask 0x5";244; CHECK-NEXT: ld.global.nc.b32 %r1, [%rd1];245; CHECK-NEXT: ld.param.b64 %rd2, [global_4xi8_invariant_param_1];246; CHECK-NEXT: st.global.b8 [%rd2], %r1;247; CHECK-NEXT: prmt.b32 %r2, %r1, 0, 0x7772U;248; CHECK-NEXT: st.global.b8 [%rd2+2], %r2;249; CHECK-NEXT: ret;250 %a.load = tail call <4 x i8> @llvm.masked.load.v4i8.p1(ptr addrspace(1) align 4 %a, <4 x i1> <i1 true, i1 false, i1 true, i1 false>, <4 x i8> poison), !invariant.load !0251 tail call void @llvm.masked.store.v4i8.p1(<4 x i8> %a.load, ptr addrspace(1) align 4 %b, <4 x i1> <i1 true, i1 false, i1 true, i1 false>)252 ret void253}254 255define void @global_4xi8_no_align(ptr addrspace(1) %a, ptr addrspace(1) %b) {256; CHECK-LABEL: global_4xi8_no_align(257; CHECK: {258; CHECK-NEXT: .reg .b16 %rs<3>;259; CHECK-NEXT: .reg .b64 %rd<3>;260; CHECK-EMPTY:261; CHECK-NEXT: // %bb.0:262; CHECK-NEXT: ld.param.b64 %rd1, [global_4xi8_no_align_param_0];263; CHECK-NEXT: ld.global.b8 %rs1, [%rd1];264; CHECK-NEXT: ld.param.b64 %rd2, [global_4xi8_no_align_param_1];265; CHECK-NEXT: ld.global.b8 %rs2, [%rd1+2];266; CHECK-NEXT: st.global.b8 [%rd2], %rs1;267; CHECK-NEXT: st.global.b8 [%rd2+2], %rs2;268; CHECK-NEXT: ret;269 %a.load = tail call <4 x i8> @llvm.masked.load.v4i8.p1(ptr addrspace(1) align 2 %a, <4 x i1> <i1 true, i1 false, i1 true, i1 false>, <4 x i8> poison)270 tail call void @llvm.masked.store.v4i8.p1(<4 x i8> %a.load, ptr addrspace(1) align 4 %b, <4 x i1> <i1 true, i1 false, i1 true, i1 false>)271 ret void272}273 274; In sm100+, we pack 2xf32 loads into a single b64 load while lowering275define void @global_2xf32(ptr addrspace(1) %a, ptr addrspace(1) %b) {276; SM90-LABEL: global_2xf32(277; SM90: {278; SM90-NEXT: .reg .b32 %r<3>;279; SM90-NEXT: .reg .b64 %rd<3>;280; SM90-EMPTY:281; SM90-NEXT: // %bb.0:282; SM90-NEXT: ld.param.b64 %rd1, [global_2xf32_param_0];283; SM90-NEXT: .pragma "used_bytes_mask 0xf";284; SM90-NEXT: ld.global.v2.b32 {%r1, %r2}, [%rd1];285; SM90-NEXT: ld.param.b64 %rd2, [global_2xf32_param_1];286; SM90-NEXT: st.global.b32 [%rd2], %r1;287; SM90-NEXT: ret;288;289; SM100-LABEL: global_2xf32(290; SM100: {291; SM100-NEXT: .reg .b32 %r<2>;292; SM100-NEXT: .reg .b64 %rd<4>;293; SM100-EMPTY:294; SM100-NEXT: // %bb.0:295; SM100-NEXT: ld.param.b64 %rd1, [global_2xf32_param_0];296; SM100-NEXT: .pragma "used_bytes_mask 0xf";297; SM100-NEXT: ld.global.b64 %rd2, [%rd1];298; SM100-NEXT: ld.param.b64 %rd3, [global_2xf32_param_1];299; SM100-NEXT: mov.b64 {%r1, _}, %rd2;300; SM100-NEXT: st.global.b32 [%rd3], %r1;301; SM100-NEXT: ret;302 %a.load = tail call <2 x float> @llvm.masked.load.v2f32.p1(ptr addrspace(1) align 8 %a, <2 x i1> <i1 true, i1 false>, <2 x float> poison)303 tail call void @llvm.masked.store.v2f32.p1(<2 x float> %a.load, ptr addrspace(1) align 8 %b, <2 x i1> <i1 true, i1 false>)304 ret void305}306 307define void @global_2xf32_invariant(ptr addrspace(1) %a, ptr addrspace(1) %b) {308; SM90-LABEL: global_2xf32_invariant(309; SM90: {310; SM90-NEXT: .reg .b32 %r<3>;311; SM90-NEXT: .reg .b64 %rd<3>;312; SM90-EMPTY:313; SM90-NEXT: // %bb.0:314; SM90-NEXT: ld.param.b64 %rd1, [global_2xf32_invariant_param_0];315; SM90-NEXT: .pragma "used_bytes_mask 0xf";316; SM90-NEXT: ld.global.nc.v2.b32 {%r1, %r2}, [%rd1];317; SM90-NEXT: ld.param.b64 %rd2, [global_2xf32_invariant_param_1];318; SM90-NEXT: st.global.b32 [%rd2], %r1;319; SM90-NEXT: ret;320;321; SM100-LABEL: global_2xf32_invariant(322; SM100: {323; SM100-NEXT: .reg .b32 %r<2>;324; SM100-NEXT: .reg .b64 %rd<4>;325; SM100-EMPTY:326; SM100-NEXT: // %bb.0:327; SM100-NEXT: ld.param.b64 %rd1, [global_2xf32_invariant_param_0];328; SM100-NEXT: .pragma "used_bytes_mask 0xf";329; SM100-NEXT: ld.global.nc.b64 %rd2, [%rd1];330; SM100-NEXT: ld.param.b64 %rd3, [global_2xf32_invariant_param_1];331; SM100-NEXT: mov.b64 {%r1, _}, %rd2;332; SM100-NEXT: st.global.b32 [%rd3], %r1;333; SM100-NEXT: ret;334 %a.load = tail call <2 x float> @llvm.masked.load.v2f32.p1(ptr addrspace(1) align 8 %a, <2 x i1> <i1 true, i1 false>, <2 x float> poison), !invariant.load !0335 tail call void @llvm.masked.store.v2f32.p1(<2 x float> %a.load, ptr addrspace(1) align 8 %b, <2 x i1> <i1 true, i1 false>)336 ret void337}338 339define void @global_2xf32_no_align(ptr addrspace(1) %a, ptr addrspace(1) %b) {340; CHECK-LABEL: global_2xf32_no_align(341; CHECK: {342; CHECK-NEXT: .reg .b32 %r<2>;343; CHECK-NEXT: .reg .b64 %rd<3>;344; CHECK-EMPTY:345; CHECK-NEXT: // %bb.0:346; CHECK-NEXT: ld.param.b64 %rd1, [global_2xf32_no_align_param_0];347; CHECK-NEXT: ld.global.b32 %r1, [%rd1];348; CHECK-NEXT: ld.param.b64 %rd2, [global_2xf32_no_align_param_1];349; CHECK-NEXT: st.global.b32 [%rd2], %r1;350; CHECK-NEXT: ret;351 %a.load = tail call <2 x float> @llvm.masked.load.v2f32.p1(ptr addrspace(1) align 4 %a, <2 x i1> <i1 true, i1 false>, <2 x float> poison)352 tail call void @llvm.masked.store.v2f32.p1(<2 x float> %a.load, ptr addrspace(1) align 8 %b, <2 x i1> <i1 true, i1 false>)353 ret void354}355 356declare <8 x i32> @llvm.masked.load.v8i32.p1(ptr addrspace(1), <8 x i1>, <8 x i32>)357declare void @llvm.masked.store.v8i32.p1(<8 x i32>, ptr addrspace(1), <8 x i1>)358declare <16 x i16> @llvm.masked.load.v16i16.p1(ptr addrspace(1), <16 x i1>, <16 x i16>)359declare void @llvm.masked.store.v16i16.p1(<16 x i16>, ptr addrspace(1), <16 x i1>)360declare <2 x i16> @llvm.masked.load.v2i16.p1(ptr addrspace(1), <2 x i1>, <2 x i16>)361declare void @llvm.masked.store.v2i16.p1(<2 x i16>, ptr addrspace(1), <2 x i1>)362declare <4 x i8> @llvm.masked.load.v4i8.p1(ptr addrspace(1), <4 x i1>, <4 x i8>)363declare void @llvm.masked.store.v4i8.p1(<4 x i8>, ptr addrspace(1), <4 x i1>)364declare <2 x float> @llvm.masked.load.v2f32.p1(ptr addrspace(1), <2 x i1>, <2 x float>)365declare void @llvm.masked.store.v2f32.p1(<2 x float>, ptr addrspace(1), <2 x i1>)366!0 = !{}367