Skip to content

[NVPTX] Lower i1 select with logical ops in the general case #135868

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 3 commits into from
Apr 17, 2025
Merged

[NVPTX] Lower i1 select with logical ops in the general case #135868

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
b521b6b
pre-commit tests
AlexMaclean Apr 11, 2025
abc1479
[NVPTX] Lower i1 select with logical ops in the general case
AlexMaclean Apr 15, 2025
b09cb17
update tests
AlexMaclean Apr 16, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
52 changes: 35 additions & 17 deletions llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2850,6 +2850,40 @@ static SDValue lowerFREM(SDValue Op, SelectionDAG &DAG,
return DAG.getSelect(DL, Ty, IsInf, X, Sub);
}

static SDValue lowerSELECT(SDValue Op, SelectionDAG &DAG) {
assert(Op.getValueType() == MVT::i1 && "Custom lowering enabled only for i1");

SDValue Cond = Op->getOperand(0);
SDValue TrueVal = Op->getOperand(1);
SDValue FalseVal = Op->getOperand(2);
SDLoc DL(Op);

// If both operands are truncated, we push the select through the truncates.
if (TrueVal.getOpcode() == ISD::TRUNCATE &&
FalseVal.getOpcode() == ISD::TRUNCATE) {
TrueVal = TrueVal.getOperand(0);
FalseVal = FalseVal.getOperand(0);

EVT VT = TrueVal.getSimpleValueType().bitsLE(FalseVal.getSimpleValueType())
? TrueVal.getValueType()
: FalseVal.getValueType();
TrueVal = DAG.getAnyExtOrTrunc(TrueVal, DL, VT);
FalseVal = DAG.getAnyExtOrTrunc(FalseVal, DL, VT);
SDValue Select = DAG.getSelect(DL, VT, Cond, TrueVal, FalseVal);
return DAG.getNode(ISD::TRUNCATE, DL, MVT::i1, Select);
}

// Otherwise, expand the select into a series of logical operations. These
// often can be folded into other operations either by us or ptxas.
TrueVal = DAG.getFreeze(TrueVal);
FalseVal = DAG.getFreeze(FalseVal);
SDValue And1 = DAG.getNode(ISD::AND, DL, MVT::i1, Cond, TrueVal);
SDValue NotCond = DAG.getNOT(DL, Cond, MVT::i1);
SDValue And2 = DAG.getNode(ISD::AND, DL, MVT::i1, NotCond, FalseVal);
SDValue Or = DAG.getNode(ISD::OR, DL, MVT::i1, And1, And2);
return Or;
}

SDValue
NVPTXTargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const {
switch (Op.getOpcode()) {
Expand Down Expand Up @@ -2889,7 +2923,7 @@ NVPTXTargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const {
case ISD::SRL_PARTS:
return LowerShiftRightParts(Op, DAG);
case ISD::SELECT:
return LowerSelect(Op, DAG);
return lowerSELECT(Op, DAG);
case ISD::FROUND:
return LowerFROUND(Op, DAG);
case ISD::FCOPYSIGN:
Expand Down Expand Up @@ -3056,22 +3090,6 @@ SDValue NVPTXTargetLowering::LowerVASTART(SDValue Op, SelectionDAG &DAG) const {
MachinePointerInfo(SV));
}

SDValue NVPTXTargetLowering::LowerSelect(SDValue Op, SelectionDAG &DAG) const {
SDValue Op0 = Op->getOperand(0);
SDValue Op1 = Op->getOperand(1);
SDValue Op2 = Op->getOperand(2);
SDLoc DL(Op.getNode());

assert(Op.getValueType() == MVT::i1 && "Custom lowering enabled only for i1");

Op1 = DAG.getNode(ISD::ANY_EXTEND, DL, MVT::i32, Op1);
Op2 = DAG.getNode(ISD::ANY_EXTEND, DL, MVT::i32, Op2);
SDValue Select = DAG.getNode(ISD::SELECT, DL, MVT::i32, Op0, Op1, Op2);
SDValue Trunc = DAG.getNode(ISD::TRUNCATE, DL, MVT::i1, Select);

return Trunc;
}

SDValue NVPTXTargetLowering::LowerLOAD(SDValue Op, SelectionDAG &DAG) const {
if (Op.getValueType() == MVT::i1)
return LowerLOADi1(Op, DAG);
Expand Down
2 changes: 0 additions & 2 deletions llvm/lib/Target/NVPTX/NVPTXISelLowering.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -324,8 +324,6 @@ class NVPTXTargetLowering : public TargetLowering {
SDValue LowerShiftRightParts(SDValue Op, SelectionDAG &DAG) const;
SDValue LowerShiftLeftParts(SDValue Op, SelectionDAG &DAG) const;

SDValue LowerSelect(SDValue Op, SelectionDAG &DAG) const;

SDValue LowerBR_JT(SDValue Op, SelectionDAG &DAG) const;

SDValue LowerVAARG(SDValue Op, SelectionDAG &DAG) const;
Expand Down
20 changes: 18 additions & 2 deletions llvm/test/CodeGen/NVPTX/bug22246.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,13 +1,29 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}

target datalayout = "e-i64:64-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

; CHECK-LABEL: _Z3foobbbPb
define void @_Z3foobbbPb(i1 zeroext %p1, i1 zeroext %p2, i1 zeroext %p3, ptr nocapture %output) {
; CHECK-LABEL: _Z3foobbbPb(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<2>;
; CHECK-NEXT: .reg .b16 %rs<7>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0: // %entry
; CHECK-NEXT: ld.param.u8 %rs1, [_Z3foobbbPb_param_0];
; CHECK-NEXT: and.b16 %rs2, %rs1, 1;
; CHECK-NEXT: setp.ne.b16 %p1, %rs2, 0;
; CHECK-NEXT: ld.param.u8 %rs3, [_Z3foobbbPb_param_1];
; CHECK-NEXT: ld.param.u8 %rs4, [_Z3foobbbPb_param_2];
; CHECK-NEXT: selp.b16 %rs5, %rs3, %rs4, %p1;
; CHECK-NEXT: and.b16 %rs6, %rs5, 1;
; CHECK-NEXT: ld.param.u64 %rd1, [_Z3foobbbPb_param_3];
; CHECK-NEXT: st.u8 [%rd1], %rs6;
; CHECK-NEXT: ret;
entry:
; CHECK: selp.b32 %r{{[0-9]+}}, %r{{[0-9]+}}, %r{{[0-9]+}}, %p{{[0-9]+}}
%.sink.v = select i1 %p1, i1 %p2, i1 %p3
%frombool5 = zext i1 %.sink.v to i8
store i8 %frombool5, ptr %output, align 1
Expand Down
128 changes: 128 additions & 0 deletions llvm/test/CodeGen/NVPTX/i1-select.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,128 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_20 | FileCheck %s
; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_20 | %ptxas-verify %}

target triple = "nvptx-nvidia-cuda"

define i32 @test_select_i1_trunc(i32 %a, i32 %b, i32 %c, i32 %true, i32 %false) {
; CHECK-LABEL: test_select_i1_trunc(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<3>;
; CHECK-NEXT: .reg .b32 %r<10>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u32 %r1, [test_select_i1_trunc_param_0];
; CHECK-NEXT: and.b32 %r2, %r1, 1;
; CHECK-NEXT: setp.ne.b32 %p1, %r2, 0;
; CHECK-NEXT: ld.param.u32 %r3, [test_select_i1_trunc_param_1];
; CHECK-NEXT: ld.param.u32 %r4, [test_select_i1_trunc_param_2];
; CHECK-NEXT: ld.param.u32 %r5, [test_select_i1_trunc_param_3];
; CHECK-NEXT: selp.b32 %r6, %r3, %r4, %p1;
; CHECK-NEXT: and.b32 %r7, %r6, 1;
; CHECK-NEXT: setp.ne.b32 %p2, %r7, 0;
; CHECK-NEXT: ld.param.u32 %r8, [test_select_i1_trunc_param_4];
; CHECK-NEXT: selp.b32 %r9, %r5, %r8, %p2;
; CHECK-NEXT: st.param.b32 [func_retval0], %r9;
; CHECK-NEXT: ret;
%a_trunc = trunc i32 %a to i1
%b_trunc = trunc i32 %b to i1
%c_trunc = trunc i32 %c to i1
%select_i1 = select i1 %a_trunc, i1 %b_trunc, i1 %c_trunc
%select_ret = select i1 %select_i1, i32 %true, i32 %false
ret i32 %select_ret
}

define i32 @test_select_i1_trunc_2(i64 %a, i16 %b, i32 %c, i32 %true, i32 %false) {
; CHECK-LABEL: test_select_i1_trunc_2(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<3>;
; CHECK-NEXT: .reg .b16 %rs<5>;
; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u64 %rd1, [test_select_i1_trunc_2_param_0];
; CHECK-NEXT: and.b64 %rd2, %rd1, 1;
; CHECK-NEXT: setp.ne.b64 %p1, %rd2, 0;
; CHECK-NEXT: ld.param.u16 %rs1, [test_select_i1_trunc_2_param_1];
; CHECK-NEXT: ld.param.u16 %rs2, [test_select_i1_trunc_2_param_2];
; CHECK-NEXT: ld.param.u32 %r1, [test_select_i1_trunc_2_param_3];
; CHECK-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;
; CHECK-NEXT: and.b16 %rs4, %rs3, 1;
; CHECK-NEXT: setp.ne.b16 %p2, %rs4, 0;
; CHECK-NEXT: ld.param.u32 %r2, [test_select_i1_trunc_2_param_4];
; CHECK-NEXT: selp.b32 %r3, %r1, %r2, %p2;
; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
; CHECK-NEXT: ret;
%a_trunc = trunc i64 %a to i1
%b_trunc = trunc i16 %b to i1
%c_trunc = trunc i32 %c to i1
%select_i1 = select i1 %a_trunc, i1 %b_trunc, i1 %c_trunc
%select_ret = select i1 %select_i1, i32 %true, i32 %false
ret i32 %select_ret
}

define i32 @test_select_i1_basic(i32 %v1, i32 %v2, i32 %v3, i32 %true, i32 %false) {
; CHECK-LABEL: test_select_i1_basic(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<4>;
; CHECK-NEXT: .reg .b32 %r<12>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u32 %r1, [test_select_i1_basic_param_0];
; CHECK-NEXT: ld.param.u32 %r2, [test_select_i1_basic_param_1];
; CHECK-NEXT: or.b32 %r4, %r1, %r2;
; CHECK-NEXT: setp.ne.s32 %p1, %r1, 0;
; CHECK-NEXT: ld.param.u32 %r5, [test_select_i1_basic_param_2];
; CHECK-NEXT: setp.eq.s32 %p2, %r5, 0;
; CHECK-NEXT: ld.param.u32 %r7, [test_select_i1_basic_param_3];
; CHECK-NEXT: setp.eq.s32 %p3, %r4, 0;
; CHECK-NEXT: ld.param.u32 %r8, [test_select_i1_basic_param_4];
; CHECK-NEXT: selp.b32 %r9, %r7, %r8, %p2;
; CHECK-NEXT: selp.b32 %r10, %r9, %r8, %p1;
; CHECK-NEXT: selp.b32 %r11, %r7, %r10, %p3;
; CHECK-NEXT: st.param.b32 [func_retval0], %r11;
; CHECK-NEXT: ret;
%b1 = icmp eq i32 %v1, 0
%b2 = icmp eq i32 %v2, 0
%b3 = icmp eq i32 %v3, 0
%select_i1 = select i1 %b1, i1 %b2, i1 %b3
%select_ret = select i1 %select_i1, i32 %true, i32 %false
ret i32 %select_ret
}

define i32 @test_select_i1_basic_folding(i32 %v1, i32 %v2, i32 %v3, i32 %true, i32 %false) {
; CHECK-LABEL: test_select_i1_basic_folding(
; CHECK: {
; CHECK-NEXT: .reg .pred %p<13>;
; CHECK-NEXT: .reg .b32 %r<7>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.u32 %r1, [test_select_i1_basic_folding_param_0];
; CHECK-NEXT: setp.eq.s32 %p1, %r1, 0;
; CHECK-NEXT: ld.param.u32 %r2, [test_select_i1_basic_folding_param_1];
; CHECK-NEXT: setp.ne.s32 %p2, %r2, 0;
; CHECK-NEXT: setp.eq.s32 %p3, %r2, 0;
; CHECK-NEXT: ld.param.u32 %r3, [test_select_i1_basic_folding_param_2];
; CHECK-NEXT: setp.eq.s32 %p4, %r3, 0;
; CHECK-NEXT: ld.param.u32 %r4, [test_select_i1_basic_folding_param_3];
; CHECK-NEXT: xor.pred %p6, %p1, %p3;
; CHECK-NEXT: ld.param.u32 %r5, [test_select_i1_basic_folding_param_4];
; CHECK-NEXT: and.pred %p7, %p6, %p4;
; CHECK-NEXT: and.pred %p9, %p2, %p4;
; CHECK-NEXT: and.pred %p10, %p3, %p7;
; CHECK-NEXT: or.pred %p11, %p10, %p9;
; CHECK-NEXT: xor.pred %p12, %p11, %p3;
; CHECK-NEXT: selp.b32 %r6, %r4, %r5, %p12;
; CHECK-NEXT: st.param.b32 [func_retval0], %r6;
; CHECK-NEXT: ret;
%b1 = icmp eq i32 %v1, 0
%b2 = icmp eq i32 %v2, 0
%b3 = icmp eq i32 %v3, 0
%b4 = xor i1 %b1, %b2
%b5 = and i1 %b4, %b3
%select_i1 = select i1 %b2, i1 %b5, i1 %b3
%b6 = xor i1 %select_i1, %b2
%select_ret = select i1 %b6, i32 %true, i32 %false
ret i32 %select_ret
}
Loading
Loading