llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Mikołaj Piróg (mikolaj-pirog)
<details>
<summary>Changes</summary>
As in title. These are 2 changes to how fp contract pragmas interact with each
other and the CLI options:
- Code like this, compiled under -ffp-contract=fast
```
//
float foo(float a, float b, float c, float d, float e) {
#pragma STDC FP_CONTRACT ON
float mul = c * d;
float add = mul + e;
return add;
}
```cpp
will now have 'contract' set on fp ops, i.e. `on` pragma won't overwrite `fast`
CLI option.
- Code like this:
```
float bar(float a, float b, float c, float d, float e) {
#pragma clang fp contract(fast)
#pragma clang fp contract(on)
float mul = c * d;
float add = mul + e;
return add;
}
```cpp
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
``````````
</details>
https://github.com/llvm/llvm-project/pull/169859
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
