319 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; This test is based on load-store-vectors.ll,8; and contains testing for lowering 256-bit masked vector stores9 10; Types we are checking: i32, i64, f32, f6411 12; Address spaces we are checking: generic, global13; - Global is the only address space that currently supports masked stores.14; - The generic stores will get legalized before the backend via scalarization,15; this file tests that even though we won't be generating them in the LSV.16 17; 256-bit vector loads/stores are only legal for blackwell+, so on sm_90, the vectors will be split18 19; generic address space20 21define void @generic_8xi32(ptr %a, ptr %b) {22; CHECK-LABEL: generic_8xi32(23; CHECK: {24; CHECK-NEXT: .reg .b32 %r<9>;25; CHECK-NEXT: .reg .b64 %rd<3>;26; CHECK-EMPTY:27; CHECK-NEXT: // %bb.0:28; CHECK-NEXT: ld.param.b64 %rd1, [generic_8xi32_param_0];29; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];30; CHECK-NEXT: ld.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];31; CHECK-NEXT: ld.param.b64 %rd2, [generic_8xi32_param_1];32; CHECK-NEXT: st.b32 [%rd2], %r5;33; CHECK-NEXT: st.b32 [%rd2+8], %r7;34; CHECK-NEXT: st.b32 [%rd2+28], %r4;35; CHECK-NEXT: ret;36 %a.load = load <8 x i32>, ptr %a37 tail call void @llvm.masked.store.v8i32.p0(<8 x i32> %a.load, ptr align 32 %b, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 false, i1 false, i1 false, i1 true>)38 ret void39}40 41define void @generic_4xi64(ptr %a, ptr %b) {42; CHECK-LABEL: generic_4xi64(43; CHECK: {44; CHECK-NEXT: .reg .b64 %rd<7>;45; CHECK-EMPTY:46; CHECK-NEXT: // %bb.0:47; CHECK-NEXT: ld.param.b64 %rd1, [generic_4xi64_param_0];48; CHECK-NEXT: ld.v2.b64 {%rd2, %rd3}, [%rd1+16];49; CHECK-NEXT: ld.v2.b64 {%rd4, %rd5}, [%rd1];50; CHECK-NEXT: ld.param.b64 %rd6, [generic_4xi64_param_1];51; CHECK-NEXT: st.b64 [%rd6], %rd4;52; CHECK-NEXT: st.b64 [%rd6+16], %rd2;53; CHECK-NEXT: ret;54 %a.load = load <4 x i64>, ptr %a55 tail call void @llvm.masked.store.v4i64.p0(<4 x i64> %a.load, ptr align 32 %b, <4 x i1> <i1 true, i1 false, i1 true, i1 false>)56 ret void57}58 59define void @generic_8xfloat(ptr %a, ptr %b) {60; CHECK-LABEL: generic_8xfloat(61; CHECK: {62; CHECK-NEXT: .reg .b32 %r<9>;63; CHECK-NEXT: .reg .b64 %rd<3>;64; CHECK-EMPTY:65; CHECK-NEXT: // %bb.0:66; CHECK-NEXT: ld.param.b64 %rd1, [generic_8xfloat_param_0];67; CHECK-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];68; CHECK-NEXT: ld.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];69; CHECK-NEXT: ld.param.b64 %rd2, [generic_8xfloat_param_1];70; CHECK-NEXT: st.b32 [%rd2], %r5;71; CHECK-NEXT: st.b32 [%rd2+8], %r7;72; CHECK-NEXT: st.b32 [%rd2+28], %r4;73; CHECK-NEXT: ret;74 %a.load = load <8 x float>, ptr %a75 tail call void @llvm.masked.store.v8f32.p0(<8 x float> %a.load, ptr align 32 %b, <8 x i1> <i1 true, i1 false, i1 true, i1 false, i1 false, i1 false, i1 false, i1 true>)76 ret void77}78 79define void @generic_4xdouble(ptr %a, ptr %b) {80; CHECK-LABEL: generic_4xdouble(81; CHECK: {82; CHECK-NEXT: .reg .b64 %rd<7>;83; CHECK-EMPTY:84; CHECK-NEXT: // %bb.0:85; CHECK-NEXT: ld.param.b64 %rd1, [generic_4xdouble_param_0];86; CHECK-NEXT: ld.v2.b64 {%rd2, %rd3}, [%rd1+16];87; CHECK-NEXT: ld.v2.b64 {%rd4, %rd5}, [%rd1];88; CHECK-NEXT: ld.param.b64 %rd6, [generic_4xdouble_param_1];89; CHECK-NEXT: st.b64 [%rd6], %rd4;90; CHECK-NEXT: st.b64 [%rd6+16], %rd2;91; CHECK-NEXT: ret;92 %a.load = load <4 x double>, ptr %a93 tail call void @llvm.masked.store.v4f64.p0(<4 x double> %a.load, ptr align 32 %b, <4 x i1> <i1 true, i1 false, i1 true, i1 false>)94 ret void95}96 97; global address space98 99define void @global_8xi32(ptr addrspace(1) %a, ptr addrspace(1) %b) {100; SM90-LABEL: global_8xi32(101; SM90: {102; SM90-NEXT: .reg .b32 %r<9>;103; SM90-NEXT: .reg .b64 %rd<3>;104; SM90-EMPTY:105; SM90-NEXT: // %bb.0:106; SM90-NEXT: ld.param.b64 %rd1, [global_8xi32_param_0];107; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];108; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];109; SM90-NEXT: ld.param.b64 %rd2, [global_8xi32_param_1];110; SM90-NEXT: st.global.b32 [%rd2], %r5;111; SM90-NEXT: st.global.b32 [%rd2+8], %r7;112; SM90-NEXT: st.global.b32 [%rd2+28], %r4;113; SM90-NEXT: ret;114;115; SM100-LABEL: global_8xi32(116; SM100: {117; SM100-NEXT: .reg .b32 %r<9>;118; SM100-NEXT: .reg .b64 %rd<3>;119; SM100-EMPTY:120; SM100-NEXT: // %bb.0:121; SM100-NEXT: ld.param.b64 %rd1, [global_8xi32_param_0];122; SM100-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];123; SM100-NEXT: ld.param.b64 %rd2, [global_8xi32_param_1];124; SM100-NEXT: st.global.v8.b32 [%rd2], {%r1, _, %r3, _, _, _, _, %r8};125; SM100-NEXT: ret;126 %a.load = load <8 x i32>, ptr addrspace(1) %a127 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>)128 ret void129}130 131define void @global_4xi64(ptr addrspace(1) %a, ptr addrspace(1) %b) {132; SM90-LABEL: global_4xi64(133; SM90: {134; SM90-NEXT: .reg .b64 %rd<7>;135; SM90-EMPTY:136; SM90-NEXT: // %bb.0:137; SM90-NEXT: ld.param.b64 %rd1, [global_4xi64_param_0];138; SM90-NEXT: ld.global.v2.b64 {%rd2, %rd3}, [%rd1+16];139; SM90-NEXT: ld.global.v2.b64 {%rd4, %rd5}, [%rd1];140; SM90-NEXT: ld.param.b64 %rd6, [global_4xi64_param_1];141; SM90-NEXT: st.global.b64 [%rd6], %rd4;142; SM90-NEXT: st.global.b64 [%rd6+16], %rd2;143; SM90-NEXT: ret;144;145; SM100-LABEL: global_4xi64(146; SM100: {147; SM100-NEXT: .reg .b64 %rd<7>;148; SM100-EMPTY:149; SM100-NEXT: // %bb.0:150; SM100-NEXT: ld.param.b64 %rd1, [global_4xi64_param_0];151; SM100-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];152; SM100-NEXT: ld.param.b64 %rd6, [global_4xi64_param_1];153; SM100-NEXT: st.global.v4.b64 [%rd6], {%rd2, _, %rd4, _};154; SM100-NEXT: ret;155 %a.load = load <4 x i64>, ptr addrspace(1) %a156 tail call void @llvm.masked.store.v4i64.p1(<4 x i64> %a.load, ptr addrspace(1) align 32 %b, <4 x i1> <i1 true, i1 false, i1 true, i1 false>)157 ret void158}159 160define void @global_8xfloat(ptr addrspace(1) %a, ptr addrspace(1) %b) {161; SM90-LABEL: global_8xfloat(162; SM90: {163; SM90-NEXT: .reg .b32 %r<9>;164; SM90-NEXT: .reg .b64 %rd<3>;165; SM90-EMPTY:166; SM90-NEXT: // %bb.0:167; SM90-NEXT: ld.param.b64 %rd1, [global_8xfloat_param_0];168; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];169; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];170; SM90-NEXT: ld.param.b64 %rd2, [global_8xfloat_param_1];171; SM90-NEXT: st.global.b32 [%rd2], %r5;172; SM90-NEXT: st.global.b32 [%rd2+8], %r7;173; SM90-NEXT: st.global.b32 [%rd2+28], %r4;174; SM90-NEXT: ret;175;176; SM100-LABEL: global_8xfloat(177; SM100: {178; SM100-NEXT: .reg .b32 %r<9>;179; SM100-NEXT: .reg .b64 %rd<3>;180; SM100-EMPTY:181; SM100-NEXT: // %bb.0:182; SM100-NEXT: ld.param.b64 %rd1, [global_8xfloat_param_0];183; SM100-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];184; SM100-NEXT: ld.param.b64 %rd2, [global_8xfloat_param_1];185; SM100-NEXT: st.global.v8.b32 [%rd2], {%r1, _, %r3, _, _, _, _, %r8};186; SM100-NEXT: ret;187 %a.load = load <8 x float>, ptr addrspace(1) %a188 tail call void @llvm.masked.store.v8f32.p1(<8 x float> %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>)189 ret void190}191 192define void @global_4xdouble(ptr addrspace(1) %a, ptr addrspace(1) %b) {193; SM90-LABEL: global_4xdouble(194; SM90: {195; SM90-NEXT: .reg .b64 %rd<7>;196; SM90-EMPTY:197; SM90-NEXT: // %bb.0:198; SM90-NEXT: ld.param.b64 %rd1, [global_4xdouble_param_0];199; SM90-NEXT: ld.global.v2.b64 {%rd2, %rd3}, [%rd1+16];200; SM90-NEXT: ld.global.v2.b64 {%rd4, %rd5}, [%rd1];201; SM90-NEXT: ld.param.b64 %rd6, [global_4xdouble_param_1];202; SM90-NEXT: st.global.b64 [%rd6], %rd4;203; SM90-NEXT: st.global.b64 [%rd6+16], %rd2;204; SM90-NEXT: ret;205;206; SM100-LABEL: global_4xdouble(207; SM100: {208; SM100-NEXT: .reg .b64 %rd<7>;209; SM100-EMPTY:210; SM100-NEXT: // %bb.0:211; SM100-NEXT: ld.param.b64 %rd1, [global_4xdouble_param_0];212; SM100-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];213; SM100-NEXT: ld.param.b64 %rd6, [global_4xdouble_param_1];214; SM100-NEXT: st.global.v4.b64 [%rd6], {%rd2, _, %rd4, _};215; SM100-NEXT: ret;216 %a.load = load <4 x double>, ptr addrspace(1) %a217 tail call void @llvm.masked.store.v4f64.p1(<4 x double> %a.load, ptr addrspace(1) align 32 %b, <4 x i1> <i1 true, i1 false, i1 true, i1 false>)218 ret void219}220 221; edge cases222define void @global_8xi32_all_mask_on(ptr addrspace(1) %a, ptr addrspace(1) %b) {223; SM90-LABEL: global_8xi32_all_mask_on(224; SM90: {225; SM90-NEXT: .reg .b32 %r<9>;226; SM90-NEXT: .reg .b64 %rd<3>;227; SM90-EMPTY:228; SM90-NEXT: // %bb.0:229; SM90-NEXT: ld.param.b64 %rd1, [global_8xi32_all_mask_on_param_0];230; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];231; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1+16];232; SM90-NEXT: ld.param.b64 %rd2, [global_8xi32_all_mask_on_param_1];233; SM90-NEXT: st.global.v4.b32 [%rd2+16], {%r5, %r6, %r7, %r8};234; SM90-NEXT: st.global.v4.b32 [%rd2], {%r1, %r2, %r3, %r4};235; SM90-NEXT: ret;236;237; SM100-LABEL: global_8xi32_all_mask_on(238; SM100: {239; SM100-NEXT: .reg .b64 %rd<7>;240; SM100-EMPTY:241; SM100-NEXT: // %bb.0:242; SM100-NEXT: ld.param.b64 %rd1, [global_8xi32_all_mask_on_param_0];243; SM100-NEXT: ld.global.v4.b64 {%rd2, %rd3, %rd4, %rd5}, [%rd1];244; SM100-NEXT: ld.param.b64 %rd6, [global_8xi32_all_mask_on_param_1];245; SM100-NEXT: st.global.v4.b64 [%rd6], {%rd2, %rd3, %rd4, %rd5};246; SM100-NEXT: ret;247 %a.load = load <8 x i32>, ptr addrspace(1) %a248 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 true, i1 true, i1 true, i1 true, i1 true, i1 true, i1 true>)249 ret void250}251 252define void @global_8xi32_all_mask_off(ptr addrspace(1) %a, ptr addrspace(1) %b) {253; CHECK-LABEL: global_8xi32_all_mask_off(254; CHECK: {255; CHECK-EMPTY:256; CHECK-EMPTY:257; CHECK-NEXT: // %bb.0:258; CHECK-NEXT: ret;259 %a.load = load <8 x i32>, ptr addrspace(1) %a260 tail call void @llvm.masked.store.v8i32.p1(<8 x i32> %a.load, ptr addrspace(1) align 32 %b, <8 x i1> <i1 false, i1 false, i1 false, i1 false, i1 false, i1 false, i1 false, i1 false>)261 ret void262}263 264; This is an example pattern for the LSV's output of these masked stores265define void @vectorizerOutput(ptr addrspace(1) %in, ptr addrspace(1) %out) {266; SM90-LABEL: vectorizerOutput(267; SM90: {268; SM90-NEXT: .reg .b32 %r<9>;269; SM90-NEXT: .reg .b64 %rd<3>;270; SM90-EMPTY:271; SM90-NEXT: // %bb.0:272; SM90-NEXT: ld.param.b64 %rd1, [vectorizerOutput_param_0];273; SM90-NEXT: ld.global.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];274; SM90-NEXT: ld.global.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];275; SM90-NEXT: ld.param.b64 %rd2, [vectorizerOutput_param_1];276; SM90-NEXT: st.global.b32 [%rd2], %r5;277; SM90-NEXT: st.global.b32 [%rd2+4], %r6;278; SM90-NEXT: st.global.b32 [%rd2+12], %r8;279; SM90-NEXT: st.global.b32 [%rd2+16], %r1;280; SM90-NEXT: ret;281;282; SM100-LABEL: vectorizerOutput(283; SM100: {284; SM100-NEXT: .reg .b32 %r<9>;285; SM100-NEXT: .reg .b64 %rd<3>;286; SM100-EMPTY:287; SM100-NEXT: // %bb.0:288; SM100-NEXT: ld.param.b64 %rd1, [vectorizerOutput_param_0];289; SM100-NEXT: ld.global.v8.b32 {%r1, %r2, %r3, %r4, %r5, %r6, %r7, %r8}, [%rd1];290; SM100-NEXT: ld.param.b64 %rd2, [vectorizerOutput_param_1];291; SM100-NEXT: st.global.v8.b32 [%rd2], {%r1, %r2, _, %r4, %r5, _, _, _};292; SM100-NEXT: ret;293 %1 = load <8 x i32>, ptr addrspace(1) %in, align 32294 %load05 = extractelement <8 x i32> %1, i32 0295 %load16 = extractelement <8 x i32> %1, i32 1296 %load38 = extractelement <8 x i32> %1, i32 3297 %load49 = extractelement <8 x i32> %1, i32 4298 %2 = insertelement <8 x i32> poison, i32 %load05, i32 0299 %3 = insertelement <8 x i32> %2, i32 %load16, i32 1300 %4 = insertelement <8 x i32> %3, i32 poison, i32 2301 %5 = insertelement <8 x i32> %4, i32 %load38, i32 3302 %6 = insertelement <8 x i32> %5, i32 %load49, i32 4303 %7 = insertelement <8 x i32> %6, i32 poison, i32 5304 %8 = insertelement <8 x i32> %7, i32 poison, i32 6305 %9 = insertelement <8 x i32> %8, i32 poison, i32 7306 call void @llvm.masked.store.v8i32.p1(<8 x i32> %9, ptr addrspace(1) align 32 %out, <8 x i1> <i1 true, i1 true, i1 false, i1 true, i1 true, i1 false, i1 false, i1 false>)307 ret void308}309 310declare void @llvm.masked.store.v8i32.p0(<8 x i32>, ptr, <8 x i1>)311declare void @llvm.masked.store.v4i64.p0(<4 x i64>, ptr, <4 x i1>)312declare void @llvm.masked.store.v8f32.p0(<8 x float>, ptr, <8 x i1>)313declare void @llvm.masked.store.v4f64.p0(<4 x double>, ptr, <4 x i1>)314 315declare void @llvm.masked.store.v8i32.p1(<8 x i32>, ptr addrspace(1), <8 x i1>)316declare void @llvm.masked.store.v4i64.p1(<4 x i64>, ptr addrspace(1), <4 x i1>)317declare void @llvm.masked.store.v8f32.p1(<8 x float>, ptr addrspace(1), <8 x i1>)318declare void @llvm.masked.store.v4f64.p1(<4 x double>, ptr addrspace(1), <4 x i1>)319