diff --git a/clang/lib/Sema/SemaAttr.cpp b/clang/lib/Sema/SemaAttr.cpp index 7729c113e422e..8a9a1e62770e1 100644 --- a/clang/lib/Sema/SemaAttr.cpp +++ b/clang/lib/Sema/SemaAttr.cpp @@ -1381,9 +1381,15 @@ void Sema::ActOnPragmaVisibility(const IdentifierInfo* VisType, void Sema::ActOnPragmaFPContract(SourceLocation Loc, LangOptions::FPModeKind FPC) { FPOptionsOverride NewFPFeatures = CurFPFeatureOverrides(); + bool HasContractOverride = NewFPFeatures.hasFPContractModeOverride(); + auto DefContractMode = getLangOpts().getDefaultFPContractMode(); + switch (FPC) { case LangOptions::FPM_On: - NewFPFeatures.setAllowFPContractWithinStatement(); + if ((HasContractOverride && + NewFPFeatures.getFPContractModeOverride() == LangOptions::FPM_Off) || + (!HasContractOverride && DefContractMode == LangOptions::FPM_Off)) + NewFPFeatures.setAllowFPContractWithinStatement(); break; case LangOptions::FPM_Fast: case LangOptions::FPM_FastHonorPragmas: diff --git a/clang/test/CodeGen/ffp-contract-fhp-pragma-override.cpp b/clang/test/CodeGen/ffp-contract-fhp-pragma-override.cpp index ff35c9204c79c..fe2d52fae5cf0 100644 --- a/clang/test/CodeGen/ffp-contract-fhp-pragma-override.cpp +++ b/clang/test/CodeGen/ffp-contract-fhp-pragma-override.cpp @@ -2,15 +2,16 @@ float fp_contract_on_1(float a, float b, float c) { // CHECK-LABEL: fp_contract_on_1fff( - // CHECK: call float @llvm.fmuladd.f32(float {{.*}}, float {{.*}}, float {{.*}}) + // CHECK: fmul contract float + // CHECK: fadd contract float #pragma STDC FP_CONTRACT ON return a * b + c; } float fp_contract_on_2(float a, float b, float c) { // CHECK-LABEL: fp_contract_on_2fff( - // CHECK: fmul float - // CHECK: fadd float + // CHECK: fmul contract float + // CHECK: fadd contract float #pragma STDC FP_CONTRACT ON float t = a * b; return t + c; @@ -52,15 +53,16 @@ float fp_contract_default_2(float a, float b, float c) { float fp_contract_clang_on_1(float a, float b, float c) { // CHECK-LABEL: fp_contract_clang_on_1fff( - // CHECK: call float @llvm.fmuladd.f32(float {{.*}}, float {{.*}}, float {{.*}}) + // CHECK: fmul contract float + // CHECK: fadd contract float #pragma clang fp contract(on) return a * b + c; } float fp_contract_clang_on_2(float a, float b, float c) { // CHECK-LABEL: fp_contract_clang_on_2fff( - // CHECK: fmul float - // CHECK: fadd float + // CHECK: fmul contract float + // CHECK: fadd contract float #pragma clang fp contract(on) float t = a * b; return t + c; @@ -104,14 +106,15 @@ float fp_contract_clang_fast_2(float a, float b, float c) { float fp_contract_global_on_1(float a, float b, float c) { // CHECK-LABEL: fp_contract_global_on_1fff( - // CHECK: call float @llvm.fmuladd.f32(float {{.*}}, float {{.*}}, float {{.*}}) + // CHECK: fmul contract float + // CHECK: fadd contract float return a * b + c; } float fp_contract_global_on_2(float a, float b, float c) { // CHECK-LABEL: fp_contract_global_on_2fff( - // CHECK: fmul float - // CHECK: fadd float + // CHECK: fmul contract float + // CHECK: fadd contract float float t = a * b; return t + c; } diff --git a/clang/test/CodeGen/fp-contract-fast-pragma.cpp b/clang/test/CodeGen/fp-contract-fast-pragma.cpp index 0bb01d6e17a1d..f0399dbdc9039 100644 --- a/clang/test/CodeGen/fp-contract-fast-pragma.cpp +++ b/clang/test/CodeGen/fp-contract-fast-pragma.cpp @@ -87,17 +87,28 @@ float fp_contract_6(float a, float b, float c) { return a * b + c; } +// 'on' after 'fast' shouldn't disable 'fast' +#pragma clang fp contract(fast) +#pragma clang fp contract(on) +float fp_contract_7(float a, float b, float c) { + // COMMON: _Z13fp_contract_7fff + // CHECK: %[[M:.+]] = fmul contract float %a, %b + // CHECK-NEXT: fadd contract float %[[M]], %c + // STRICT: %[[M:.+]] = tail call contract float @llvm.experimental.constrained.fmul.f32(float %a, float %b, metadata !"round.tonearest", metadata !"fpexcept.strict") + // STRICT-NEXT: tail call contract float @llvm.experimental.constrained.fadd.f32(float %[[M]], float %c, metadata !"round.tonearest", metadata !"fpexcept.strict") + return a * b + c; +} #pragma clang fp contract(fast) -float fp_contract_7(float a) { -// COMMON: _Z13fp_contract_7f +float fp_contract_8(float a) { +// COMMON: _Z13fp_contract_8f // CHECK: tail call contract float @llvm.sqrt.f32(float %a) // STRICT: tail call contract float @llvm.experimental.constrained.sqrt.f32(float %a, metadata !"round.tonearest", metadata !"fpexcept.strict") return __builtin_sqrtf(a); } -float fp_contract_8(float a) { -// COMMON: _Z13fp_contract_8f +float fp_contract_9(float a) { +// COMMON: _Z13fp_contract_9f // CHECK: tail call float @llvm.sqrt.f32(float %a) // STRICT: tail call float @llvm.experimental.constrained.sqrt.f32(float %a, metadata !"round.tonearest", metadata !"fpexcept.strict") #pragma clang fp contract(off) diff --git a/clang/test/CodeGen/fp-function-attrs.cpp b/clang/test/CodeGen/fp-function-attrs.cpp index 3775bd5452d78..3b3122f492fc1 100644 --- a/clang/test/CodeGen/fp-function-attrs.cpp +++ b/clang/test/CodeGen/fp-function-attrs.cpp @@ -51,7 +51,7 @@ float test_contract_on_pragma(float a, float b, float c) { // CHECK: define{{.*}} float @_Z23test_contract_on_pragmafff(float noundef nofpclass(nan inf) %a, float noundef nofpclass(nan inf) %b, float noundef nofpclass(nan inf) %c) // CHECK: fmul fast float {{%.+}}, {{%.+}} -// CHECK: fadd reassoc nnan ninf nsz arcp afn float {{%.+}}, {{%.+}} +// CHECK: fadd fast float {{%.+}}, {{%.+}} // CHECK: attributes [[FAST_ATTRS]] = { {{.*}}"no-infs-fp-math"="true" {{.*}}"no-nans-fp-math"="true" "no-signed-zeros-fp-math"="true"{{.*}} } // CHECK: attributes [[PRECISE_ATTRS]] = { {{.*}}"no-infs-fp-math"="false" {{.*}}"no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false"{{.*}} } diff --git a/clang/test/CodeGenCUDA/fp-contract.cu b/clang/test/CodeGenCUDA/fp-contract.cu index d6c796a817cbf..b3fc8e30e8990 100644 --- a/clang/test/CodeGenCUDA/fp-contract.cu +++ b/clang/test/CodeGenCUDA/fp-contract.cu @@ -268,8 +268,7 @@ __host__ __device__ float func2(float a, float b, float c) { // COMMON-LABEL: _Z5func3fff // NV-OPT-FAST: fma.rn.f32 // NV-OPT-FAST-NEXT: st.param.b32 -// NV-OPT-FASTSTD: mul.rn.f32 -// NV-OPT-FASTSTD: add.rn.f32 +// NV-OPT-FASTSTD: fma.rn.f32 // NV-OPT-FASTSTD-NEXT: st.param.b32 // NV-OPT-ON: mul.rn.f32 // NV-OPT-ON: add.rn.f32 @@ -278,8 +277,8 @@ __host__ __device__ float func2(float a, float b, float c) { // NV-OPT-OFF: add.rn.f32 // NV-OPT-OFF-NEXT: st.param.b32 -// AMD-OPT-FAST-IR: fmul float -// AMD-OPT-FAST-IR: fadd float +// AMD-OPT-FAST-IR: fmul contract float +// AMD-OPT-FAST-IR: fadd contract float // AMD-OPT-ON-IR: fmul float // AMD-OPT-ON-IR: fadd float // AMD-OPT-OFF-IR: fmul float @@ -287,9 +286,8 @@ __host__ __device__ float func2(float a, float b, float c) { // AMD-OPT-FAST: v_fmac_f32_e32 // AMD-OPT-FAST-NEXT: s_setpc_b64 -// AMD-OPT-FASTSTD: v_mul_f32_e32 -// AMD-OPT-FASTSTD-NEXT: v_add_f32_e32 -// AMD-OPT-FASTSTD-NEXT: s_setpc_b64 +// AMD-OPT-FASTSTD: v_fmac_f32_e32 +// AMD-OPT-FASTSTD: s_setpc_b64 // AMD-OPT-ON: v_mul_f32_e32 // AMD-OPT-ON-NEXT: v_add_f32_e32 // AMD-OPT-ON-NEXT: s_setpc_b64