Skip to content

Commit b671979

Browse files
paperchalicepaperchalice
authored
[NVPTX] Remove UnsafeFPMath uses (llvm#151479)
Remove `UnsafeFPMath` in NVPTX part, it blocks some bugfixes related to clang and the ultimate goal is to remove `resetTargetOptions` method in `TargetMachine`, see FIXME in `resetTargetOptions`. See also https://discourse.llvm.org/t/rfc-honor-pragmas-with-ffp-contract-fast https://discourse.llvm.org/t/allowfpopfusion-vs-sdnodeflags-hasallowcontract
1 parent 8432f24 commit b671979

13 files changed

+416
-491
lines changed

‎llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp‎

Lines changed: 1 addition & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -70,7 +70,7 @@ NVPTXDAGToDAGISel::getDivF32Level(const SDNode *N) const {
7070
}
7171

7272
bool NVPTXDAGToDAGISel::usePrecSqrtF32(const SDNode *N) const {
73-
return Subtarget->getTargetLowering()->usePrecSqrtF32(*MF, N);
73+
return Subtarget->getTargetLowering()->usePrecSqrtF32(N);
7474
}
7575

7676
bool NVPTXDAGToDAGISel::useF32FTZ() const {
@@ -82,11 +82,6 @@ bool NVPTXDAGToDAGISel::allowFMA() const {
8282
return TL->allowFMA(*MF, OptLevel);
8383
}
8484

85-
bool NVPTXDAGToDAGISel::allowUnsafeFPMath() const {
86-
const NVPTXTargetLowering *TL = Subtarget->getTargetLowering();
87-
return TL->allowUnsafeFPMath(*MF);
88-
}
89-
9085
bool NVPTXDAGToDAGISel::doRsqrtOpt() const { return EnableRsqrtOpt; }
9186

9287
/// Select - Select instructions not customized! Used for

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

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -44,7 +44,6 @@ class LLVM_LIBRARY_VISIBILITY NVPTXDAGToDAGISel : public SelectionDAGISel {
4444
bool usePrecSqrtF32(const SDNode *N) const;
4545
bool useF32FTZ() const;
4646
bool allowFMA() const;
47-
bool allowUnsafeFPMath() const;
4847
bool doRsqrtOpt() const;
4948

5049
NVPTXScopes Scopes{};

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

Lines changed: 6 additions & 27 deletions
Original file line numberDiff line numberDiff line change
@@ -125,27 +125,18 @@ NVPTXTargetLowering::getDivF32Level(const MachineFunction &MF,
125125
if (UsePrecDivF32.getNumOccurrences() > 0)
126126
return UsePrecDivF32;
127127

128-
// Otherwise, use div.approx if fast math is enabled
129-
if (allowUnsafeFPMath(MF))
130-
return NVPTX::DivPrecisionLevel::Approx;
131-
132128
const SDNodeFlags Flags = N.getFlags();
133129
if (Flags.hasApproximateFuncs())
134130
return NVPTX::DivPrecisionLevel::Approx;
135131

136132
return NVPTX::DivPrecisionLevel::IEEE754;
137133
}
138134

139-
bool NVPTXTargetLowering::usePrecSqrtF32(const MachineFunction &MF,
140-
const SDNode *N) const {
135+
bool NVPTXTargetLowering::usePrecSqrtF32(const SDNode *N) const {
141136
// If nvptx-prec-sqrtf32 is used on the command-line, always honor it
142137
if (UsePrecSqrtF32.getNumOccurrences() > 0)
143138
return UsePrecSqrtF32;
144139

145-
// Otherwise, use sqrt.approx if fast math is enabled
146-
if (allowUnsafeFPMath(MF))
147-
return false;
148-
149140
if (N) {
150141
const SDNodeFlags Flags = N->getFlags();
151142
if (Flags.hasApproximateFuncs())
@@ -1193,8 +1184,7 @@ SDValue NVPTXTargetLowering::getSqrtEstimate(SDValue Operand, SelectionDAG &DAG,
11931184
bool &UseOneConst,
11941185
bool Reciprocal) const {
11951186
if (!(Enabled == ReciprocalEstimate::Enabled ||
1196-
(Enabled == ReciprocalEstimate::Unspecified &&
1197-
!usePrecSqrtF32(DAG.getMachineFunction()))))
1187+
(Enabled == ReciprocalEstimate::Unspecified && !usePrecSqrtF32())))
11981188
return SDValue();
11991189

12001190
if (ExtraSteps == ReciprocalEstimate::Unspecified)
@@ -2851,8 +2841,7 @@ static SDValue lowerROT(SDValue Op, SelectionDAG &DAG) {
28512841
SDLoc(Op), Opcode, DAG);
28522842
}
28532843

2854-
static SDValue lowerFREM(SDValue Op, SelectionDAG &DAG,
2855-
bool AllowUnsafeFPMath) {
2844+
static SDValue lowerFREM(SDValue Op, SelectionDAG &DAG) {
28562845
// Lower (frem x, y) into (sub x, (mul (ftrunc (div x, y)) y)),
28572846
// i.e. "poor man's fmod()". When y is infinite, x is returned. This matches
28582847
// the semantics of LLVM's frem.
@@ -2869,7 +2858,7 @@ static SDValue lowerFREM(SDValue Op, SelectionDAG &DAG,
28692858
SDValue Sub = DAG.getNode(ISD::FSUB, DL, Ty, X, Mul,
28702859
Flags | SDNodeFlags::AllowContract);
28712860

2872-
if (AllowUnsafeFPMath || Flags.hasNoInfs())
2861+
if (Flags.hasNoInfs())
28732862
return Sub;
28742863

28752864
// If Y is infinite, return X
@@ -3014,7 +3003,7 @@ NVPTXTargetLowering::LowerOperation(SDValue Op, SelectionDAG &DAG) const {
30143003
case ISD::CTLZ:
30153004
return lowerCTLZCTPOP(Op, DAG);
30163005
case ISD::FREM:
3017-
return lowerFREM(Op, DAG, allowUnsafeFPMath(DAG.getMachineFunction()));
3006+
return lowerFREM(Op, DAG);
30183007

30193008
default:
30203009
llvm_unreachable("Custom lowering not defined for operation");
@@ -4868,17 +4857,7 @@ bool NVPTXTargetLowering::allowFMA(MachineFunction &MF,
48684857
if (MF.getTarget().Options.AllowFPOpFusion == FPOpFusion::Fast)
48694858
return true;
48704859

4871-
return allowUnsafeFPMath(MF);
4872-
}
4873-
4874-
bool NVPTXTargetLowering::allowUnsafeFPMath(const MachineFunction &MF) const {
4875-
// Honor TargetOptions flags that explicitly say unsafe math is okay.
4876-
if (MF.getTarget().Options.UnsafeFPMath)
4877-
return true;
4878-
4879-
// Allow unsafe math if unsafe-fp-math attribute explicitly says so.
4880-
const Function &F = MF.getFunction();
4881-
return F.getFnAttribute("unsafe-fp-math").getValueAsBool();
4860+
return false;
48824861
}
48834862

48844863
static bool isConstZero(const SDValue &Operand) {

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

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -206,8 +206,7 @@ class NVPTXTargetLowering : public TargetLowering {
206206

207207
// Get whether we should use a precise or approximate 32-bit floating point
208208
// sqrt instruction.
209-
bool usePrecSqrtF32(const MachineFunction &MF,
210-
const SDNode *N = nullptr) const;
209+
bool usePrecSqrtF32(const SDNode *N = nullptr) const;
211210

212211
// Get whether we should use instructions that flush floating-point denormals
213212
// to sign-preserving zero.
@@ -220,7 +219,6 @@ class NVPTXTargetLowering : public TargetLowering {
220219
unsigned combineRepeatedFPDivisors() const override { return 2; }
221220

222221
bool allowFMA(MachineFunction &MF, CodeGenOptLevel OptLevel) const;
223-
bool allowUnsafeFPMath(const MachineFunction &MF) const;
224222

225223
bool isFMAFasterThanFMulAndFAdd(const MachineFunction &MF,
226224
EVT) const override {

‎llvm/lib/Target/NVPTX/NVPTXInstrInfo.td‎

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1133,9 +1133,8 @@ defm FMA_F64 : FMA<F64RT, allow_ftz = false>;
11331133
// sin/cos/tanh
11341134

11351135
class UnaryOpAllowsApproxFn<SDPatternOperator operator>
1136-
: PatFrag<(ops node:$A),
1137-
(operator node:$A), [{
1138-
return allowUnsafeFPMath() || N->getFlags().hasApproximateFuncs();
1136+
: PatFrag<(ops node:$A), (operator node:$A), [{
1137+
return N->getFlags().hasApproximateFuncs();
11391138
}]>;
11401139

11411140
def SIN_APPROX_f32 :

‎llvm/test/CodeGen/NVPTX/bf16x2-instructions-approx.ll‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
11
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
2-
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 --enable-unsafe-fp-math | FileCheck --check-prefixes=CHECK %s
3-
; RUN: %if ptxas-11.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 --enable-unsafe-fp-math | %ptxas-verify -arch=sm_80 %}
2+
; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 | FileCheck --check-prefixes=CHECK %s
3+
; RUN: %if ptxas-11.8 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_80 -mattr=+ptx71 | %ptxas-verify -arch=sm_80 %}
44

55
target datalayout = "e-m:o-i64:64-i128:128-n32:64-S128"
66

@@ -22,7 +22,7 @@ define <2 x bfloat> @test_sin(<2 x bfloat> %a) #0 #1 {
2222
; CHECK-NEXT: cvt.rn.bf16x2.f32 %r5, %r4, %r2;
2323
; CHECK-NEXT: st.param.b32 [func_retval0], %r5;
2424
; CHECK-NEXT: ret;
25-
%r = call <2 x bfloat> @llvm.sin.f16(<2 x bfloat> %a)
25+
%r = call afn <2 x bfloat> @llvm.sin.f16(<2 x bfloat> %a)
2626
ret <2 x bfloat> %r
2727
}
2828

@@ -41,7 +41,7 @@ define <2 x bfloat> @test_cos(<2 x bfloat> %a) #0 #1 {
4141
; CHECK-NEXT: cvt.rn.bf16x2.f32 %r5, %r4, %r2;
4242
; CHECK-NEXT: st.param.b32 [func_retval0], %r5;
4343
; CHECK-NEXT: ret;
44-
%r = call <2 x bfloat> @llvm.cos.f16(<2 x bfloat> %a)
44+
%r = call afn <2 x bfloat> @llvm.cos.f16(<2 x bfloat> %a)
4545
ret <2 x bfloat> %r
4646
}
4747

‎llvm/test/CodeGen/NVPTX/f16-instructions.ll‎

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -886,8 +886,8 @@ define half @test_sqrt(half %a) #0 {
886886
; CHECK: cvt.rn.f16.f32 [[R:%rs[0-9]+]], [[RF]];
887887
; CHECK: st.param.b16 [func_retval0], [[R]];
888888
; CHECK: ret;
889-
define half @test_sin(half %a) #0 #1 {
890-
%r = call half @llvm.sin.f16(half %a)
889+
define half @test_sin(half %a) #0 {
890+
%r = call afn half @llvm.sin.f16(half %a)
891891
ret half %r
892892
}
893893

@@ -900,8 +900,8 @@ define half @test_sin(half %a) #0 #1 {
900900
; CHECK: cvt.rn.f16.f32 [[R:%rs[0-9]+]], [[RF]];
901901
; CHECK: st.param.b16 [func_retval0], [[R]];
902902
; CHECK: ret;
903-
define half @test_cos(half %a) #0 #1 {
904-
%r = call half @llvm.cos.f16(half %a)
903+
define half @test_cos(half %a) #0 {
904+
%r = call afn half @llvm.cos.f16(half %a)
905905
ret half %r
906906
}
907907

@@ -1183,4 +1183,3 @@ define <2 x half> @test_neg_f16x2(<2 x half> noundef %arg) #0 {
11831183
}
11841184

11851185
attributes #0 = { nounwind }
1186-
attributes #1 = { "unsafe-fp-math" = "true" }

‎llvm/test/CodeGen/NVPTX/f16x2-instructions.ll‎

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1674,7 +1674,7 @@ define <2 x half> @test_sqrt(<2 x half> %a) #0 {
16741674
; ret <2 x half> %r
16751675
;}
16761676

1677-
define <2 x half> @test_sin(<2 x half> %a) #0 #1 {
1677+
define <2 x half> @test_sin(<2 x half> %a) #0 {
16781678
; CHECK-LABEL: test_sin(
16791679
; CHECK: {
16801680
; CHECK-NEXT: .reg .b16 %rs<5>;
@@ -1692,11 +1692,11 @@ define <2 x half> @test_sin(<2 x half> %a) #0 #1 {
16921692
; CHECK-NEXT: mov.b32 %r6, {%rs4, %rs3};
16931693
; CHECK-NEXT: st.param.b32 [func_retval0], %r6;
16941694
; CHECK-NEXT: ret;
1695-
%r = call <2 x half> @llvm.sin.f16(<2 x half> %a)
1695+
%r = call afn <2 x half> @llvm.sin.f16(<2 x half> %a)
16961696
ret <2 x half> %r
16971697
}
16981698

1699-
define <2 x half> @test_cos(<2 x half> %a) #0 #1 {
1699+
define <2 x half> @test_cos(<2 x half> %a) #0 {
17001700
; CHECK-LABEL: test_cos(
17011701
; CHECK: {
17021702
; CHECK-NEXT: .reg .b16 %rs<5>;
@@ -1714,7 +1714,7 @@ define <2 x half> @test_cos(<2 x half> %a) #0 #1 {
17141714
; CHECK-NEXT: mov.b32 %r6, {%rs4, %rs3};
17151715
; CHECK-NEXT: st.param.b32 [func_retval0], %r6;
17161716
; CHECK-NEXT: ret;
1717-
%r = call <2 x half> @llvm.cos.f16(<2 x half> %a)
1717+
%r = call afn <2 x half> @llvm.cos.f16(<2 x half> %a)
17181718
ret <2 x half> %r
17191719
}
17201720

@@ -2330,4 +2330,3 @@ define void @test_store_2xhalf(ptr %p1, ptr %p2, <2 x half> %v) {
23302330

23312331

23322332
attributes #0 = { nounwind }
2333-
attributes #1 = { "unsafe-fp-math" = "true" }

‎llvm/test/CodeGen/NVPTX/f32x2-instructions.ll‎

Lines changed: 4 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -1638,7 +1638,7 @@ define <2 x float> @test_sqrt(<2 x float> %a) #0 {
16381638
; ret <2 x float> %r
16391639
;}
16401640

1641-
define <2 x float> @test_sin(<2 x float> %a) #0 #1 {
1641+
define <2 x float> @test_sin(<2 x float> %a) #0 {
16421642
; CHECK-LABEL: test_sin(
16431643
; CHECK: {
16441644
; CHECK-NEXT: .reg .b32 %r<5>;
@@ -1651,11 +1651,11 @@ define <2 x float> @test_sin(<2 x float> %a) #0 #1 {
16511651
; CHECK-NEXT: sin.approx.f32 %r4, %r1;
16521652
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
16531653
; CHECK-NEXT: ret;
1654-
%r = call <2 x float> @llvm.sin(<2 x float> %a)
1654+
%r = call afn <2 x float> @llvm.sin(<2 x float> %a)
16551655
ret <2 x float> %r
16561656
}
16571657

1658-
define <2 x float> @test_cos(<2 x float> %a) #0 #1 {
1658+
define <2 x float> @test_cos(<2 x float> %a) #0 {
16591659
; CHECK-LABEL: test_cos(
16601660
; CHECK: {
16611661
; CHECK-NEXT: .reg .b32 %r<5>;
@@ -1668,7 +1668,7 @@ define <2 x float> @test_cos(<2 x float> %a) #0 #1 {
16681668
; CHECK-NEXT: cos.approx.f32 %r4, %r1;
16691669
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
16701670
; CHECK-NEXT: ret;
1671-
%r = call <2 x float> @llvm.cos(<2 x float> %a)
1671+
%r = call afn <2 x float> @llvm.cos(<2 x float> %a)
16721672
ret <2 x float> %r
16731673
}
16741674

@@ -2157,5 +2157,4 @@ define void @test_trunc_to_v2f16(<2 x float> %a, ptr %p) {
21572157

21582158

21592159
attributes #0 = { nounwind }
2160-
attributes #1 = { "unsafe-fp-math" = "true" }
21612160
attributes #2 = { "denormal-fp-math"="preserve-sign" }

0 commit comments

Comments
 (0)