-
Notifications
You must be signed in to change notification settings - Fork 15.4k
[clang] Make fp contract(on) not overwrite earlier contract(fast) and CLI --ffp-contract=fast(-honor-pragmas) #169859
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.
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
base: main
Are you sure you want to change the base?
[clang] Make fp contract(on) not overwrite earlier contract(fast) and CLI --ffp-contract=fast(-honor-pragmas) #169859
Conversation
|
@llvm/pr-subscribers-clang Author: Mikołaj Piróg (mikolaj-pirog) ChangesAs in title. These are 2 changes to how fp contract pragmas interact with each other and the CLI options:
float bar(float a, float b, float c, float d, float e) { return add; will now have 'contract' set as well -- `on` doesn't overwrite previous `fast`.
---
Full diff: https://github.com/llvm/llvm-project/pull/169859.diff
5 Files Affected:
- (modified) clang/lib/Sema/SemaAttr.cpp (+7-1)
- (modified) clang/test/CodeGen/ffp-contract-fhp-pragma-override.cpp (+12-9)
- (modified) clang/test/CodeGen/fp-contract-fast-pragma.cpp (+15-4)
- (modified) clang/test/CodeGen/fp-function-attrs.cpp (+1-1)
- (modified) clang/test/CodeGenCUDA/fp-contract.cu (+5-7)
``````````diff
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
|
|
My understanding of interactions of this flag is based on treating |
As in title. These are 2 changes to how fp contract pragmas interact with each other and the CLI options:
will now have 'contract' set on fp ops, i.e.
onpragma won't overwritefastCLI option.will now have 'contract' set as well --
ondoesn't overwrite previousfast.