Skip to content

Commit 925df13

Browse files
adelejjehadelejjeh
committed
[AMDGPU] Update log lowering to remove contract for AMDGCN backend
1 parent c6775e2 commit 925df13

File tree

2 files changed

+47
-21
lines changed

2 files changed

+47
-21
lines changed

‎clang/lib/CodeGen/CGBuiltin.cpp‎

Lines changed: 27 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -582,6 +582,25 @@ static Value *EmitISOVolatileStore(CodeGenFunction &CGF, const CallExpr *E) {
582582
return Store;
583583
}
584584

585+
// Check if an intrinsic is a transcendental function that is unsafe to
586+
// contract.
587+
static bool isUnsafeToContract(unsigned IntrinsicID, CodeGenFunction &CGF) {
588+
switch (IntrinsicID) {
589+
// The implementation for log in the AMDGCN backend uses a refinement
590+
// algorithm that requires intermediate rounding. The contract flag would
591+
// allow FMA formation that recomputes products, breaking the refinement
592+
// algorithm.
593+
case Intrinsic::log:
594+
case Intrinsic::log10:
595+
if ((CGF.getTarget().getTriple().isAMDGCN() ||
596+
CGF.getTarget().getTriple().isSPIRV()) &&
597+
CGF.getLangOpts().HIP)
598+
return true;
599+
return false;
600+
default:
601+
return false;
602+
}
603+
}
585604
// Emit a simple mangled intrinsic that has 1 argument and a return type
586605
// matching the argument type. Depending on mode, this may be a constrained
587606
// floating-point intrinsic.
@@ -596,7 +615,14 @@ Value *emitUnaryMaybeConstrainedFPBuiltin(CodeGenFunction &CGF,
596615
return CGF.Builder.CreateConstrainedFPCall(F, { Src0 });
597616
} else {
598617
Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
599-
return CGF.Builder.CreateCall(F, Src0);
618+
CallInst *Call = CGF.Builder.CreateCall(F, Src0);
619+
620+
// Check if the intrinsic is unsafe to contract
621+
if (isUnsafeToContract(IntrinsicID, CGF)) {
622+
Call->setHasAllowContract(false);
623+
}
624+
625+
return Call;
600626
}
601627
}
602628

‎clang/test/Headers/__clang_hip_math.hip‎

Lines changed: 20 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -3673,31 +3673,31 @@ extern "C" __device__ long long int test_llround(double x) {
36733673
// DEFAULT-LABEL: define dso_local noundef float @test_log10f(
36743674
// DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
36753675
// DEFAULT-NEXT: [[ENTRY:.*:]]
3676-
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]])
3676+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]])
36773677
// DEFAULT-NEXT: ret float [[TMP0]]
36783678
//
36793679
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_log10f(
36803680
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
36813681
// FINITEONLY-NEXT: [[ENTRY:.*:]]
3682-
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X]])
3682+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X]])
36833683
// FINITEONLY-NEXT: ret float [[TMP0]]
36843684
//
36853685
// APPROX-LABEL: define dso_local noundef float @test_log10f(
36863686
// APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
36873687
// APPROX-NEXT: [[ENTRY:.*:]]
3688-
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]])
3688+
// APPROX-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]])
36893689
// APPROX-NEXT: ret float [[TMP0]]
36903690
//
36913691
// NCRDIV-LABEL: define dso_local noundef float @test_log10f(
36923692
// NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
36933693
// NCRDIV-NEXT: [[ENTRY:.*:]]
3694-
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]])
3694+
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]])
36953695
// NCRDIV-NEXT: ret float [[TMP0]]
36963696
//
36973697
// AMDGCNSPIRV-LABEL: define spir_func noundef float @test_log10f(
36983698
// AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
36993699
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
3700-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log10.f32(float [[X]])
3700+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call noundef addrspace(4) float @llvm.log10.f32(float [[X]])
37013701
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
37023702
//
37033703
extern "C" __device__ float test_log10f(float x) {
@@ -3945,31 +3945,31 @@ extern "C" __device__ double test_logb(double x) {
39453945
// DEFAULT-LABEL: define dso_local noundef float @test_logf(
39463946
// DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
39473947
// DEFAULT-NEXT: [[ENTRY:.*:]]
3948-
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]])
3948+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]])
39493949
// DEFAULT-NEXT: ret float [[TMP0]]
39503950
//
39513951
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_logf(
39523952
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
39533953
// FINITEONLY-NEXT: [[ENTRY:.*:]]
3954-
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log.f32(float nofpclass(nan inf) [[X]])
3954+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf noundef float @llvm.log.f32(float nofpclass(nan inf) [[X]])
39553955
// FINITEONLY-NEXT: ret float [[TMP0]]
39563956
//
39573957
// APPROX-LABEL: define dso_local noundef float @test_logf(
39583958
// APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
39593959
// APPROX-NEXT: [[ENTRY:.*:]]
3960-
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]])
3960+
// APPROX-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]])
39613961
// APPROX-NEXT: ret float [[TMP0]]
39623962
//
39633963
// NCRDIV-LABEL: define dso_local noundef float @test_logf(
39643964
// NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
39653965
// NCRDIV-NEXT: [[ENTRY:.*:]]
3966-
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]])
3966+
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]])
39673967
// NCRDIV-NEXT: ret float [[TMP0]]
39683968
//
39693969
// AMDGCNSPIRV-LABEL: define spir_func noundef float @test_logf(
39703970
// AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
39713971
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
3972-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log.f32(float [[X]])
3972+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call noundef addrspace(4) float @llvm.log.f32(float [[X]])
39733973
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
39743974
//
39753975
extern "C" __device__ float test_logf(float x) {
@@ -8600,31 +8600,31 @@ extern "C" __device__ float test___fsub_rn(float x, float y) {
86008600
// DEFAULT-LABEL: define dso_local noundef float @test___log10f(
86018601
// DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
86028602
// DEFAULT-NEXT: [[ENTRY:.*:]]
8603-
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]])
8603+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]])
86048604
// DEFAULT-NEXT: ret float [[TMP0]]
86058605
//
86068606
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test___log10f(
86078607
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
86088608
// FINITEONLY-NEXT: [[ENTRY:.*:]]
8609-
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X]])
8609+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X]])
86108610
// FINITEONLY-NEXT: ret float [[TMP0]]
86118611
//
86128612
// APPROX-LABEL: define dso_local noundef float @test___log10f(
86138613
// APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
86148614
// APPROX-NEXT: [[ENTRY:.*:]]
8615-
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]])
8615+
// APPROX-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]])
86168616
// APPROX-NEXT: ret float [[TMP0]]
86178617
//
86188618
// NCRDIV-LABEL: define dso_local noundef float @test___log10f(
86198619
// NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
86208620
// NCRDIV-NEXT: [[ENTRY:.*:]]
8621-
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]])
8621+
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]])
86228622
// NCRDIV-NEXT: ret float [[TMP0]]
86238623
//
86248624
// AMDGCNSPIRV-LABEL: define spir_func noundef float @test___log10f(
86258625
// AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
86268626
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
8627-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log10.f32(float [[X]])
8627+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call noundef addrspace(4) float @llvm.log10.f32(float [[X]])
86288628
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
86298629
//
86308630
extern "C" __device__ float test___log10f(float x) {
@@ -8668,31 +8668,31 @@ extern "C" __device__ float test___log2f(float x) {
86688668
// DEFAULT-LABEL: define dso_local noundef float @test___logf(
86698669
// DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
86708670
// DEFAULT-NEXT: [[ENTRY:.*:]]
8671-
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]])
8671+
// DEFAULT-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]])
86728672
// DEFAULT-NEXT: ret float [[TMP0]]
86738673
//
86748674
// FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test___logf(
86758675
// FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
86768676
// FINITEONLY-NEXT: [[ENTRY:.*:]]
8677-
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log.f32(float nofpclass(nan inf) [[X]])
8677+
// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf noundef float @llvm.log.f32(float nofpclass(nan inf) [[X]])
86788678
// FINITEONLY-NEXT: ret float [[TMP0]]
86798679
//
86808680
// APPROX-LABEL: define dso_local noundef float @test___logf(
86818681
// APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
86828682
// APPROX-NEXT: [[ENTRY:.*:]]
8683-
// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]])
8683+
// APPROX-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]])
86848684
// APPROX-NEXT: ret float [[TMP0]]
86858685
//
86868686
// NCRDIV-LABEL: define dso_local noundef float @test___logf(
86878687
// NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] {
86888688
// NCRDIV-NEXT: [[ENTRY:.*:]]
8689-
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]])
8689+
// NCRDIV-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]])
86908690
// NCRDIV-NEXT: ret float [[TMP0]]
86918691
//
86928692
// AMDGCNSPIRV-LABEL: define spir_func noundef float @test___logf(
86938693
// AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] {
86948694
// AMDGCNSPIRV-NEXT: [[ENTRY:.*:]]
8695-
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log.f32(float [[X]])
8695+
// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call noundef addrspace(4) float @llvm.log.f32(float [[X]])
86968696
// AMDGCNSPIRV-NEXT: ret float [[TMP0]]
86978697
//
86988698
extern "C" __device__ float test___logf(float x) {

0 commit comments

Comments
 (0)