[llvm] 6018820 - [NVPTX] Fix lowering of i1 SETCC (#115035)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Dec 5 12:54:27 PST 2024
Author: Alex MacLean
Date: 2024-12-05T12:54:24-08:00
New Revision: 6018820c48322ee9db4971efa9b048f16521d753
URL: https://github.com/llvm/llvm-project/commit/6018820c48322ee9db4971efa9b048f16521d753
DIFF: https://github.com/llvm/llvm-project/commit/6018820c48322ee9db4971efa9b048f16521d753.diff
LOG: [NVPTX] Fix lowering of i1 SETCC (#115035)
Add DAG legalization support for expanding i1 SETCC nodes using
appropriate logical operations to simulate integer comparisons. Use
these expansions to handle i1 SETCC in NVPTX.
fixes #58428 and #57405
Added:
llvm/test/CodeGen/NVPTX/i1-icmp.ll
Modified:
llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
Removed:
################################################################################
diff --git a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
index 48018ac29bd089..f5334e8c63964a 100644
--- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
@@ -18763,9 +18763,13 @@ SDValue DAGCombiner::rebuildSetCC(SDValue N) {
EVT SetCCVT = N.getValueType();
if (LegalTypes)
SetCCVT = getSetCCResultType(SetCCVT);
- // Replace the uses of XOR with SETCC
- return DAG.getSetCC(SDLoc(N), SetCCVT, Op0, Op1,
- Equal ? ISD::SETEQ : ISD::SETNE);
+ // Replace the uses of XOR with SETCC. Note, avoid this transformation if
+ // it would introduce illegal operations post-legalization as this can
+ // result in infinite looping between converting xor->setcc here, and
+ // expanding setcc->xor in LegalizeSetCCCondCode if requested.
+ const ISD::CondCode CC = Equal ? ISD::SETEQ : ISD::SETNE;
+ if (!LegalOperations || TLI.isCondCodeLegal(CC, Op0.getSimpleValueType()))
+ return DAG.getSetCC(SDLoc(N), SetCCVT, Op0, Op1, CC);
}
}
diff --git a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
index eeba4b7d20f9c1..e87d809f88eb97 100644
--- a/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/TargetLowering.cpp
@@ -11889,6 +11889,47 @@ bool TargetLowering::LegalizeSetCCCondCode(SelectionDAG &DAG, EVT VT,
return true;
}
+ // Special case: expand i1 comparisons using logical operations.
+ if (OpVT == MVT::i1) {
+ SDValue Ret;
+ switch (CCCode) {
+ default:
+ llvm_unreachable("Unknown integer setcc!");
+ case ISD::SETEQ: // X == Y --> ~(X ^ Y)
+ Ret = DAG.getNOT(dl, DAG.getNode(ISD::XOR, dl, MVT::i1, LHS, RHS),
+ MVT::i1);
+ break;
+ case ISD::SETNE: // X != Y --> (X ^ Y)
+ Ret = DAG.getNode(ISD::XOR, dl, MVT::i1, LHS, RHS);
+ break;
+ case ISD::SETGT: // X >s Y --> X == 0 & Y == 1 --> ~X & Y
+ case ISD::SETULT: // X <u Y --> X == 0 & Y == 1 --> ~X & Y
+ Ret = DAG.getNode(ISD::AND, dl, MVT::i1, RHS,
+ DAG.getNOT(dl, LHS, MVT::i1));
+ break;
+ case ISD::SETLT: // X <s Y --> X == 1 & Y == 0 --> ~Y & X
+ case ISD::SETUGT: // X >u Y --> X == 1 & Y == 0 --> ~Y & X
+ Ret = DAG.getNode(ISD::AND, dl, MVT::i1, LHS,
+ DAG.getNOT(dl, RHS, MVT::i1));
+ break;
+ case ISD::SETULE: // X <=u Y --> X == 0 | Y == 1 --> ~X | Y
+ case ISD::SETGE: // X >=s Y --> X == 0 | Y == 1 --> ~X | Y
+ Ret = DAG.getNode(ISD::OR, dl, MVT::i1, RHS,
+ DAG.getNOT(dl, LHS, MVT::i1));
+ break;
+ case ISD::SETUGE: // X >=u Y --> X == 1 | Y == 0 --> ~Y | X
+ case ISD::SETLE: // X <=s Y --> X == 1 | Y == 0 --> ~Y | X
+ Ret = DAG.getNode(ISD::OR, dl, MVT::i1, LHS,
+ DAG.getNOT(dl, RHS, MVT::i1));
+ break;
+ }
+
+ LHS = DAG.getZExtOrTrunc(Ret, dl, VT);
+ RHS = SDValue();
+ CC = SDValue();
+ return true;
+ }
+
ISD::CondCode CC1 = ISD::SETCC_INVALID, CC2 = ISD::SETCC_INVALID;
unsigned Opc = 0;
switch (CCCode) {
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 62647b31285188..ce94dded815b8f 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -667,6 +667,11 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
setTruncStoreAction(VT, MVT::i1, Expand);
}
+ setCondCodeAction({ISD::SETNE, ISD::SETEQ, ISD::SETUGE, ISD::SETULE,
+ ISD::SETUGT, ISD::SETULT, ISD::SETGT, ISD::SETLT,
+ ISD::SETGE, ISD::SETLE},
+ MVT::i1, Expand);
+
// expand extload of vector of integers.
setLoadExtAction({ISD::EXTLOAD, ISD::SEXTLOAD, ISD::ZEXTLOAD}, MVT::v2i16,
MVT::v2i8, Expand);
diff --git a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
index 64f437fd4e4d2d..b8a55851e8d3e1 100644
--- a/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
+++ b/llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
@@ -2120,17 +2120,6 @@ defm : ISET_FORMAT_UNSIGNED<setule, CmpLE>;
defm : ISET_FORMAT_UNSIGNED<setueq, CmpEQ>;
defm : ISET_FORMAT_UNSIGNED<setune, CmpNE>;
-// i1 compares
-def : Pat<(setne Int1Regs:$a, Int1Regs:$b),
- (XORb1rr Int1Regs:$a, Int1Regs:$b)>;
-def : Pat<(setune Int1Regs:$a, Int1Regs:$b),
- (XORb1rr Int1Regs:$a, Int1Regs:$b)>;
-
-def : Pat<(seteq Int1Regs:$a, Int1Regs:$b),
- (NOT1 (XORb1rr Int1Regs:$a, Int1Regs:$b))>;
-def : Pat<(setueq Int1Regs:$a, Int1Regs:$b),
- (NOT1 (XORb1rr Int1Regs:$a, Int1Regs:$b))>;
-
// comparisons of i8 extracted with BFE as i32
// It's faster to do comparison directly on i32 extracted by BFE,
// instead of the long conversion and sign extending.
diff --git a/llvm/test/CodeGen/NVPTX/i1-icmp.ll b/llvm/test/CodeGen/NVPTX/i1-icmp.ll
new file mode 100644
index 00000000000000..b189f5ff8bca82
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/i1-icmp.ll
@@ -0,0 +1,317 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -march=nvptx64 -mcpu=sm_20 | %ptxas-verify %}
+
+target triple = "nvptx-nvidia-cuda"
+
+define i32 @icmp_i1_eq(i32 %a, i32 %b) {
+; CHECK-LABEL: icmp_i1_eq(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<4>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_eq_param_0];
+; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;
+; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_eq_param_1];
+; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1;
+; CHECK-NEXT: xor.pred %p3, %p1, %p2;
+; CHECK-NEXT: @%p3 bra $L__BB0_2;
+; CHECK-NEXT: // %bb.1: // %bb1
+; CHECK-NEXT: mov.b32 %r4, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+; CHECK-NEXT: $L__BB0_2: // %bb2
+; CHECK-NEXT: mov.b32 %r3, 127;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %p1 = icmp sgt i32 %a, 1
+ %p2 = icmp sgt i32 %b, 1
+ %c = icmp eq i1 %p1, %p2
+ br i1 %c, label %bb1, label %bb2
+bb1:
+ ret i32 1
+bb2:
+ ret i32 127
+}
+
+define i32 @icmp_i1_ne(i32 %a, i32 %b) {
+; CHECK-LABEL: icmp_i1_ne(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<5>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_ne_param_0];
+; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;
+; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_ne_param_1];
+; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1;
+; CHECK-NEXT: xor.pred %p3, %p1, %p2;
+; CHECK-NEXT: not.pred %p4, %p3;
+; CHECK-NEXT: @%p4 bra $L__BB1_2;
+; CHECK-NEXT: // %bb.1: // %bb1
+; CHECK-NEXT: mov.b32 %r4, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+; CHECK-NEXT: $L__BB1_2: // %bb2
+; CHECK-NEXT: mov.b32 %r3, 127;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %p1 = icmp sgt i32 %a, 1
+ %p2 = icmp sgt i32 %b, 1
+ %c = icmp ne i1 %p1, %p2
+ br i1 %c, label %bb1, label %bb2
+bb1:
+ ret i32 1
+bb2:
+ ret i32 127
+}
+
+define i32 @icmp_i1_sgt(i32 %a, i32 %b) {
+; CHECK-LABEL: icmp_i1_sgt(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<4>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_sgt_param_0];
+; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;
+; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_sgt_param_1];
+; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2;
+; CHECK-NEXT: or.pred %p3, %p1, %p2;
+; CHECK-NEXT: @%p3 bra $L__BB2_2;
+; CHECK-NEXT: // %bb.1: // %bb1
+; CHECK-NEXT: mov.b32 %r4, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+; CHECK-NEXT: $L__BB2_2: // %bb2
+; CHECK-NEXT: mov.b32 %r3, 127;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %p1 = icmp sgt i32 %a, 1
+ %p2 = icmp sgt i32 %b, 1
+ %c = icmp sgt i1 %p1, %p2
+ br i1 %c, label %bb1, label %bb2
+bb1:
+ ret i32 1
+bb2:
+ ret i32 127
+}
+
+define i32 @icmp_i1_slt(i32 %a, i32 %b) {
+; CHECK-LABEL: icmp_i1_slt(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<4>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_slt_param_0];
+; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2;
+; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_slt_param_1];
+; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1;
+; CHECK-NEXT: or.pred %p3, %p2, %p1;
+; CHECK-NEXT: @%p3 bra $L__BB3_2;
+; CHECK-NEXT: // %bb.1: // %bb1
+; CHECK-NEXT: mov.b32 %r4, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+; CHECK-NEXT: $L__BB3_2: // %bb2
+; CHECK-NEXT: mov.b32 %r3, 127;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %p1 = icmp sgt i32 %a, 1
+ %p2 = icmp sgt i32 %b, 1
+ %c = icmp slt i1 %p1, %p2
+ br i1 %c, label %bb1, label %bb2
+bb1:
+ ret i32 1
+bb2:
+ ret i32 127
+}
+
+define i32 @icmp_i1_sge(i32 %a, i32 %b) {
+; CHECK-LABEL: icmp_i1_sge(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<4>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_sge_param_0];
+; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;
+; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_sge_param_1];
+; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2;
+; CHECK-NEXT: and.pred %p3, %p1, %p2;
+; CHECK-NEXT: @%p3 bra $L__BB4_2;
+; CHECK-NEXT: // %bb.1: // %bb1
+; CHECK-NEXT: mov.b32 %r4, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+; CHECK-NEXT: $L__BB4_2: // %bb2
+; CHECK-NEXT: mov.b32 %r3, 127;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %p1 = icmp sgt i32 %a, 1
+ %p2 = icmp sgt i32 %b, 1
+ %c = icmp sge i1 %p1, %p2
+ br i1 %c, label %bb1, label %bb2
+bb1:
+ ret i32 1
+bb2:
+ ret i32 127
+}
+
+define i32 @icmp_i1_sle(i32 %a, i32 %b) {
+; CHECK-LABEL: icmp_i1_sle(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<4>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_sle_param_0];
+; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2;
+; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_sle_param_1];
+; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1;
+; CHECK-NEXT: and.pred %p3, %p2, %p1;
+; CHECK-NEXT: @%p3 bra $L__BB5_2;
+; CHECK-NEXT: // %bb.1: // %bb1
+; CHECK-NEXT: mov.b32 %r4, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+; CHECK-NEXT: $L__BB5_2: // %bb2
+; CHECK-NEXT: mov.b32 %r3, 127;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %p1 = icmp sgt i32 %a, 1
+ %p2 = icmp sgt i32 %b, 1
+ %c = icmp sle i1 %p1, %p2
+ br i1 %c, label %bb1, label %bb2
+bb1:
+ ret i32 1
+bb2:
+ ret i32 127
+}
+
+define i32 @icmp_i1_uge(i32 %a, i32 %b) {
+; CHECK-LABEL: icmp_i1_uge(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<4>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_uge_param_0];
+; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2;
+; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_uge_param_1];
+; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1;
+; CHECK-NEXT: and.pred %p3, %p2, %p1;
+; CHECK-NEXT: @%p3 bra $L__BB6_2;
+; CHECK-NEXT: // %bb.1: // %bb1
+; CHECK-NEXT: mov.b32 %r4, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+; CHECK-NEXT: $L__BB6_2: // %bb2
+; CHECK-NEXT: mov.b32 %r3, 127;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %p1 = icmp sgt i32 %a, 1
+ %p2 = icmp sgt i32 %b, 1
+ %c = icmp uge i1 %p1, %p2
+ br i1 %c, label %bb1, label %bb2
+bb1:
+ ret i32 1
+bb2:
+ ret i32 127
+}
+
+define i32 @icmp_i1_ugt(i32 %a, i32 %b) {
+; CHECK-LABEL: icmp_i1_ugt(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<4>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_ugt_param_0];
+; CHECK-NEXT: setp.lt.s32 %p1, %r1, 2;
+; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_ugt_param_1];
+; CHECK-NEXT: setp.gt.s32 %p2, %r2, 1;
+; CHECK-NEXT: or.pred %p3, %p2, %p1;
+; CHECK-NEXT: @%p3 bra $L__BB7_2;
+; CHECK-NEXT: // %bb.1: // %bb1
+; CHECK-NEXT: mov.b32 %r4, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+; CHECK-NEXT: $L__BB7_2: // %bb2
+; CHECK-NEXT: mov.b32 %r3, 127;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %p1 = icmp sgt i32 %a, 1
+ %p2 = icmp sgt i32 %b, 1
+ %c = icmp ugt i1 %p1, %p2
+ br i1 %c, label %bb1, label %bb2
+bb1:
+ ret i32 1
+bb2:
+ ret i32 127
+}
+
+define i32 @icmp_i1_ule(i32 %a, i32 %b) {
+; CHECK-LABEL: icmp_i1_ule(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<4>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_ule_param_0];
+; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;
+; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_ule_param_1];
+; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2;
+; CHECK-NEXT: and.pred %p3, %p1, %p2;
+; CHECK-NEXT: @%p3 bra $L__BB8_2;
+; CHECK-NEXT: // %bb.1: // %bb1
+; CHECK-NEXT: mov.b32 %r4, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+; CHECK-NEXT: $L__BB8_2: // %bb2
+; CHECK-NEXT: mov.b32 %r3, 127;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %p1 = icmp sgt i32 %a, 1
+ %p2 = icmp sgt i32 %b, 1
+ %c = icmp ule i1 %p1, %p2
+ br i1 %c, label %bb1, label %bb2
+bb1:
+ ret i32 1
+bb2:
+ ret i32 127
+}
+
+define i32 @icmp_i1_ult(i32 %a, i32 %b) {
+; CHECK-LABEL: icmp_i1_ult(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<4>;
+; CHECK-NEXT: .reg .b32 %r<5>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.u32 %r1, [icmp_i1_ult_param_0];
+; CHECK-NEXT: setp.gt.s32 %p1, %r1, 1;
+; CHECK-NEXT: ld.param.u32 %r2, [icmp_i1_ult_param_1];
+; CHECK-NEXT: setp.lt.s32 %p2, %r2, 2;
+; CHECK-NEXT: or.pred %p3, %p1, %p2;
+; CHECK-NEXT: @%p3 bra $L__BB9_2;
+; CHECK-NEXT: // %bb.1: // %bb1
+; CHECK-NEXT: mov.b32 %r4, 1;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r4;
+; CHECK-NEXT: ret;
+; CHECK-NEXT: $L__BB9_2: // %bb2
+; CHECK-NEXT: mov.b32 %r3, 127;
+; CHECK-NEXT: st.param.b32 [func_retval0], %r3;
+; CHECK-NEXT: ret;
+ %p1 = icmp sgt i32 %a, 1
+ %p2 = icmp sgt i32 %b, 1
+ %c = icmp ult i1 %p1, %p2
+ br i1 %c, label %bb1, label %bb2
+bb1:
+ ret i32 1
+bb2:
+ ret i32 127
+}
+
More information about the llvm-commits
mailing list