Skip to content

Commit 0c0f08a

Browse files
AlexMacleanAlexMaclean
committed
Add TLI::supportMismatchCopysign to prevent mismatch float types
1 parent f13f2f6 commit 0c0f08a

File tree

4 files changed

+97
-14
lines changed

4 files changed

+97
-14
lines changed

‎llvm/include/llvm/CodeGen/TargetLowering.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -4401,6 +4401,10 @@ class TargetLowering : public TargetLoweringBase {
44014401
/// Return true if the target supports ptrauth operand bundles.
44024402
virtual bool supportPtrAuthBundles() const { return false; }
44034403

4404+
/// Return true if the target supports FCOPYSIGN nodes with operands of
4405+
/// mismatched types.
4406+
virtual bool supportMismatchCopysign() const { return true; }
4407+
44044408
/// Perform necessary initialization to handle a subset of CSRs explicitly
44054409
/// via copies. This function is called at the beginning of instruction
44064410
/// selection.

‎llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp

Lines changed: 10 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -17686,7 +17686,11 @@ SDValue DAGCombiner::visitFSQRT(SDNode *N) {
1768617686
/// copysign(x, fp_extend(y)) -> copysign(x, y)
1768717687
/// copysign(x, fp_round(y)) -> copysign(x, y)
1768817688
/// Operands to the functions are the type of X and Y respectively.
17689-
static inline bool CanCombineFCOPYSIGN_EXTEND_ROUND(EVT XTy, EVT YTy) {
17689+
static inline bool CanCombineFCOPYSIGN_EXTEND_ROUND(EVT XTy, EVT YTy,
17690+
const TargetLowering &TLI) {
17691+
if (!TLI.supportMismatchCopysign())
17692+
return false;
17693+
1769017694
// Always fold no-op FP casts.
1769117695
if (XTy == YTy)
1769217696
return true;
@@ -17701,14 +17705,15 @@ static inline bool CanCombineFCOPYSIGN_EXTEND_ROUND(EVT XTy, EVT YTy) {
1770117705
return !YTy.isVector() || EnableVectorFCopySignExtendRound;
1770217706
}
1770317707

17704-
static inline bool CanCombineFCOPYSIGN_EXTEND_ROUND(SDNode *N) {
17708+
static inline bool CanCombineFCOPYSIGN_EXTEND_ROUND(SDNode *N,
17709+
const TargetLowering &TLI) {
1770517710
SDValue N1 = N->getOperand(1);
1770617711
if (N1.getOpcode() != ISD::FP_EXTEND &&
1770717712
N1.getOpcode() != ISD::FP_ROUND)
1770817713
return false;
1770917714
EVT N1VT = N1->getValueType(0);
1771017715
EVT N1Op0VT = N1->getOperand(0).getValueType();
17711-
return CanCombineFCOPYSIGN_EXTEND_ROUND(N1VT, N1Op0VT);
17716+
return CanCombineFCOPYSIGN_EXTEND_ROUND(N1VT, N1Op0VT, TLI);
1771217717
}
1771317718

1771417719
SDValue DAGCombiner::visitFCOPYSIGN(SDNode *N) {
@@ -17752,7 +17757,7 @@ SDValue DAGCombiner::visitFCOPYSIGN(SDNode *N) {
1775217757

1775317758
// copysign(x, fp_extend(y)) -> copysign(x, y)
1775417759
// copysign(x, fp_round(y)) -> copysign(x, y)
17755-
if (CanCombineFCOPYSIGN_EXTEND_ROUND(N))
17760+
if (CanCombineFCOPYSIGN_EXTEND_ROUND(N, TLI))
1775617761
return DAG.getNode(ISD::FCOPYSIGN, DL, VT, N0, N1.getOperand(0));
1775717762

1775817763
// We only take the sign bit from the sign operand.
@@ -18105,8 +18110,7 @@ SDValue DAGCombiner::visitFP_ROUND(SDNode *N) {
1810518110
// eliminate the fp_round on Y. The second step requires an additional
1810618111
// predicate to match the implementation above.
1810718112
if (N0.getOpcode() == ISD::FCOPYSIGN && N0->hasOneUse() &&
18108-
CanCombineFCOPYSIGN_EXTEND_ROUND(VT,
18109-
N0.getValueType())) {
18113+
CanCombineFCOPYSIGN_EXTEND_ROUND(VT, N0.getValueType(), TLI)) {
1811018114
SDValue Tmp = DAG.getNode(ISD::FP_ROUND, SDLoc(N0), VT,
1811118115
N0.getOperand(0), N1);
1811218116
AddToWorklist(Tmp.getNode());

‎llvm/lib/Target/NVPTX/NVPTXISelLowering.h

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -596,6 +596,8 @@ class NVPTXTargetLowering : public TargetLowering {
596596
// instruction, so we say that ctlz is cheap to speculate.
597597
bool isCheapToSpeculateCtlz(Type *Ty) const override { return true; }
598598

599+
bool supportMismatchCopysign() const override { return false; }
600+
599601
AtomicExpansionKind shouldCastAtomicLoadInIR(LoadInst *LI) const override {
600602
return AtomicExpansionKind::None;
601603
}

‎llvm/test/CodeGen/NVPTX/copysign.ll

Lines changed: 81 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -5,35 +5,108 @@
55
target triple = "nvptx64-nvidia-cuda"
66
target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
77

8-
define float @fcopysign_f(float %a, float %b) {
9-
; CHECK-LABEL: fcopysign_f(
8+
define float @fcopysign_f_f(float %a, float %b) {
9+
; CHECK-LABEL: fcopysign_f_f(
1010
; CHECK: {
1111
; CHECK-NEXT: .reg .f32 %f<4>;
1212
; CHECK-EMPTY:
1313
; CHECK-NEXT: // %bb.0:
14-
; CHECK-NEXT: ld.param.f32 %f1, [fcopysign_f_param_0];
15-
; CHECK-NEXT: ld.param.f32 %f2, [fcopysign_f_param_1];
14+
; CHECK-NEXT: ld.param.f32 %f1, [fcopysign_f_f_param_0];
15+
; CHECK-NEXT: ld.param.f32 %f2, [fcopysign_f_f_param_1];
1616
; CHECK-NEXT: copysign.f32 %f3, %f2, %f1;
1717
; CHECK-NEXT: st.param.f32 [func_retval0+0], %f3;
1818
; CHECK-NEXT: ret;
1919
%val = call float @llvm.copysign.f32(float %a, float %b)
2020
ret float %val
2121
}
2222

23-
define double @fcopysign_d(double %a, double %b) {
24-
; CHECK-LABEL: fcopysign_d(
23+
define double @fcopysign_d_d(double %a, double %b) {
24+
; CHECK-LABEL: fcopysign_d_d(
2525
; CHECK: {
2626
; CHECK-NEXT: .reg .f64 %fd<4>;
2727
; CHECK-EMPTY:
2828
; CHECK-NEXT: // %bb.0:
29-
; CHECK-NEXT: ld.param.f64 %fd1, [fcopysign_d_param_0];
30-
; CHECK-NEXT: ld.param.f64 %fd2, [fcopysign_d_param_1];
29+
; CHECK-NEXT: ld.param.f64 %fd1, [fcopysign_d_d_param_0];
30+
; CHECK-NEXT: ld.param.f64 %fd2, [fcopysign_d_d_param_1];
3131
; CHECK-NEXT: copysign.f64 %fd3, %fd2, %fd1;
3232
; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd3;
3333
; CHECK-NEXT: ret;
3434
%val = call double @llvm.copysign.f64(double %a, double %b)
3535
ret double %val
3636
}
3737

38+
define float @fcopysign_f_d(float %a, double %b) {
39+
; CHECK-LABEL: fcopysign_f_d(
40+
; CHECK: {
41+
; CHECK-NEXT: .reg .f32 %f<4>;
42+
; CHECK-NEXT: .reg .f64 %fd<2>;
43+
; CHECK-EMPTY:
44+
; CHECK-NEXT: // %bb.0:
45+
; CHECK-NEXT: ld.param.f32 %f1, [fcopysign_f_d_param_0];
46+
; CHECK-NEXT: ld.param.f64 %fd1, [fcopysign_f_d_param_1];
47+
; CHECK-NEXT: cvt.rn.f32.f64 %f2, %fd1;
48+
; CHECK-NEXT: copysign.f32 %f3, %f2, %f1;
49+
; CHECK-NEXT: st.param.f32 [func_retval0+0], %f3;
50+
; CHECK-NEXT: ret;
51+
%c = fptrunc double %b to float
52+
%val = call float @llvm.copysign.f32(float %a, float %c)
53+
ret float %val
54+
}
55+
56+
define float @fcopysign_f_h(float %a, half %b) {
57+
; CHECK-LABEL: fcopysign_f_h(
58+
; CHECK: {
59+
; CHECK-NEXT: .reg .b16 %rs<2>;
60+
; CHECK-NEXT: .reg .f32 %f<4>;
61+
; CHECK-EMPTY:
62+
; CHECK-NEXT: // %bb.0:
63+
; CHECK-NEXT: ld.param.f32 %f1, [fcopysign_f_h_param_0];
64+
; CHECK-NEXT: ld.param.b16 %rs1, [fcopysign_f_h_param_1];
65+
; CHECK-NEXT: cvt.f32.f16 %f2, %rs1;
66+
; CHECK-NEXT: copysign.f32 %f3, %f2, %f1;
67+
; CHECK-NEXT: st.param.f32 [func_retval0+0], %f3;
68+
; CHECK-NEXT: ret;
69+
%c = fpext half %b to float
70+
%val = call float @llvm.copysign.f32(float %a, float %c)
71+
ret float %val
72+
}
73+
74+
define double @fcopysign_d_f(double %a, float %b) {
75+
; CHECK-LABEL: fcopysign_d_f(
76+
; CHECK: {
77+
; CHECK-NEXT: .reg .f32 %f<2>;
78+
; CHECK-NEXT: .reg .f64 %fd<4>;
79+
; CHECK-EMPTY:
80+
; CHECK-NEXT: // %bb.0:
81+
; CHECK-NEXT: ld.param.f64 %fd1, [fcopysign_d_f_param_0];
82+
; CHECK-NEXT: ld.param.f32 %f1, [fcopysign_d_f_param_1];
83+
; CHECK-NEXT: cvt.f64.f32 %fd2, %f1;
84+
; CHECK-NEXT: copysign.f64 %fd3, %fd2, %fd1;
85+
; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd3;
86+
; CHECK-NEXT: ret;
87+
%c = fpext float %b to double
88+
%val = call double @llvm.copysign.f64(double %a, double %c)
89+
ret double %val
90+
}
91+
92+
define double @fcopysign_d_h(double %a, half %b) {
93+
; CHECK-LABEL: fcopysign_d_h(
94+
; CHECK: {
95+
; CHECK-NEXT: .reg .b16 %rs<2>;
96+
; CHECK-NEXT: .reg .f64 %fd<4>;
97+
; CHECK-EMPTY:
98+
; CHECK-NEXT: // %bb.0:
99+
; CHECK-NEXT: ld.param.f64 %fd1, [fcopysign_d_h_param_0];
100+
; CHECK-NEXT: ld.param.b16 %rs1, [fcopysign_d_h_param_1];
101+
; CHECK-NEXT: cvt.f64.f16 %fd2, %rs1;
102+
; CHECK-NEXT: copysign.f64 %fd3, %fd2, %fd1;
103+
; CHECK-NEXT: st.param.f64 [func_retval0+0], %fd3;
104+
; CHECK-NEXT: ret;
105+
%c = fpext half %b to double
106+
%val = call double @llvm.copysign.f64(double %a, double %c)
107+
ret double %val
108+
}
109+
110+
38111
declare float @llvm.copysign.f32(float, float)
39112
declare double @llvm.copysign.f64(double, double)

0 commit comments

Comments
 (0)