57 lines · plain
1; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 52; RUN: llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | FileCheck %s -check-prefixes=CHECK3; RUN: %if ptxas-sm_100 && ptxas-isa-8.8 %{ llc < %s -march=nvptx64 -mcpu=sm_100 -mattr=+ptx88 | %ptxas-verify -arch=sm_100 %}4 5; Confirm that a masked store with a variable mask is scalarized before lowering6 7define void @global_variable_mask(ptr addrspace(1) %a, ptr addrspace(1) %b, <4 x i1> %mask) {8; CHECK-LABEL: global_variable_mask(9; CHECK: {10; CHECK-NEXT: .reg .pred %p<9>;11; CHECK-NEXT: .reg .b16 %rs<9>;12; CHECK-NEXT: .reg .b64 %rd<7>;13; CHECK-EMPTY:14; CHECK-NEXT: // %bb.0:15; CHECK-NEXT: ld.param.b8 %rs1, [global_variable_mask_param_2+3];16; CHECK-NEXT: ld.param.b8 %rs3, [global_variable_mask_param_2+2];17; CHECK-NEXT: and.b16 %rs4, %rs3, 1;18; CHECK-NEXT: ld.param.b8 %rs5, [global_variable_mask_param_2+1];19; CHECK-NEXT: and.b16 %rs6, %rs5, 1;20; CHECK-NEXT: setp.ne.b16 %p2, %rs6, 0;21; CHECK-NEXT: ld.param.b8 %rs7, [global_variable_mask_param_2];22; CHECK-NEXT: and.b16 %rs8, %rs7, 1;23; CHECK-NEXT: setp.ne.b16 %p1, %rs8, 0;24; CHECK-NEXT: ld.param.b64 %rd5, [global_variable_mask_param_1];25; CHECK-NEXT: ld.param.b64 %rd6, [global_variable_mask_param_0];26; CHECK-NEXT: ld.global.v4.b64 {%rd1, %rd2, %rd3, %rd4}, [%rd6];27; CHECK-NEXT: not.pred %p5, %p1;28; CHECK-NEXT: @%p5 bra $L__BB0_2;29; CHECK-NEXT: // %bb.1: // %cond.store30; CHECK-NEXT: st.global.b64 [%rd5], %rd1;31; CHECK-NEXT: $L__BB0_2: // %else32; CHECK-NEXT: and.b16 %rs2, %rs1, 1;33; CHECK-NEXT: setp.ne.b16 %p3, %rs4, 0;34; CHECK-NEXT: not.pred %p6, %p2;35; CHECK-NEXT: @%p6 bra $L__BB0_4;36; CHECK-NEXT: // %bb.3: // %cond.store137; CHECK-NEXT: st.global.b64 [%rd5+8], %rd2;38; CHECK-NEXT: $L__BB0_4: // %else239; CHECK-NEXT: setp.ne.b16 %p4, %rs2, 0;40; CHECK-NEXT: not.pred %p7, %p3;41; CHECK-NEXT: @%p7 bra $L__BB0_6;42; CHECK-NEXT: // %bb.5: // %cond.store343; CHECK-NEXT: st.global.b64 [%rd5+16], %rd3;44; CHECK-NEXT: $L__BB0_6: // %else445; CHECK-NEXT: not.pred %p8, %p4;46; CHECK-NEXT: @%p8 bra $L__BB0_8;47; CHECK-NEXT: // %bb.7: // %cond.store548; CHECK-NEXT: st.global.b64 [%rd5+24], %rd4;49; CHECK-NEXT: $L__BB0_8: // %else650; CHECK-NEXT: ret;51 %a.load = load <4 x i64>, ptr addrspace(1) %a52 tail call void @llvm.masked.store.v4i64.p1(<4 x i64> %a.load, ptr addrspace(1) align 32 %b, <4 x i1> %mask)53 ret void54}55 56declare void @llvm.masked.store.v4i64.p1(<4 x i64>, ptr addrspace(1), <4 x i1>)57