https://github.com/adelejjeh updated https://github.com/llvm/llvm-project/pull/168770
>From 925df13db78370da9785bcb0fa90b89c09ce2470 Mon Sep 17 00:00:00 2001 From: Adel Ejjeh <[email protected]> Date: Wed, 19 Nov 2025 15:19:27 -0600 Subject: [PATCH] [AMDGPU] Update log lowering to remove contract for AMDGCN backend --- clang/lib/CodeGen/CGBuiltin.cpp | 28 ++++++++++++++++- clang/test/Headers/__clang_hip_math.hip | 40 ++++++++++++------------- 2 files changed, 47 insertions(+), 21 deletions(-) diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 3079f8ab7229e..99d1811d4787a 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -582,6 +582,25 @@ static Value *EmitISOVolatileStore(CodeGenFunction &CGF, const CallExpr *E) { return Store; } +// Check if an intrinsic is a transcendental function that is unsafe to +// contract. +static bool isUnsafeToContract(unsigned IntrinsicID, CodeGenFunction &CGF) { + switch (IntrinsicID) { + // The implementation for log in the AMDGCN backend uses a refinement + // algorithm that requires intermediate rounding. The contract flag would + // allow FMA formation that recomputes products, breaking the refinement + // algorithm. + case Intrinsic::log: + case Intrinsic::log10: + if ((CGF.getTarget().getTriple().isAMDGCN() || + CGF.getTarget().getTriple().isSPIRV()) && + CGF.getLangOpts().HIP) + return true; + return false; + default: + return false; + } +} // Emit a simple mangled intrinsic that has 1 argument and a return type // matching the argument type. Depending on mode, this may be a constrained // floating-point intrinsic. @@ -596,7 +615,14 @@ Value *emitUnaryMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, return CGF.Builder.CreateConstrainedFPCall(F, { Src0 }); } else { Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType()); - return CGF.Builder.CreateCall(F, Src0); + CallInst *Call = CGF.Builder.CreateCall(F, Src0); + + // Check if the intrinsic is unsafe to contract + if (isUnsafeToContract(IntrinsicID, CGF)) { + Call->setHasAllowContract(false); + } + + return Call; } } diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index 7e2691633c215..4eda85247d161 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -3673,31 +3673,31 @@ extern "C" __device__ long long int test_llround(double x) { // DEFAULT-LABEL: define dso_local noundef float @test_log10f( // DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // DEFAULT-NEXT: [[ENTRY:.*:]] -// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]]) +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]]) // DEFAULT-NEXT: ret float [[TMP0]] // // FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_log10f( // FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // FINITEONLY-NEXT: [[ENTRY:.*:]] -// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: define dso_local noundef float @test_log10f( // APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // APPROX-NEXT: [[ENTRY:.*:]] -// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]]) +// APPROX-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]]) // APPROX-NEXT: ret float [[TMP0]] // // NCRDIV-LABEL: define dso_local noundef float @test_log10f( // NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // NCRDIV-NEXT: [[ENTRY:.*:]] -// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]]) +// NCRDIV-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]]) // NCRDIV-NEXT: ret float [[TMP0]] // // AMDGCNSPIRV-LABEL: define spir_func noundef float @test_log10f( // AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] { // AMDGCNSPIRV-NEXT: [[ENTRY:.*:]] -// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log10.f32(float [[X]]) +// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call noundef addrspace(4) float @llvm.log10.f32(float [[X]]) // AMDGCNSPIRV-NEXT: ret float [[TMP0]] // extern "C" __device__ float test_log10f(float x) { @@ -3945,31 +3945,31 @@ extern "C" __device__ double test_logb(double x) { // DEFAULT-LABEL: define dso_local noundef float @test_logf( // DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // DEFAULT-NEXT: [[ENTRY:.*:]] -// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]]) +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]]) // DEFAULT-NEXT: ret float [[TMP0]] // // FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test_logf( // FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // FINITEONLY-NEXT: [[ENTRY:.*:]] -// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log.f32(float nofpclass(nan inf) [[X]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf noundef float @llvm.log.f32(float nofpclass(nan inf) [[X]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: define dso_local noundef float @test_logf( // APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // APPROX-NEXT: [[ENTRY:.*:]] -// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]]) +// APPROX-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]]) // APPROX-NEXT: ret float [[TMP0]] // // NCRDIV-LABEL: define dso_local noundef float @test_logf( // NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // NCRDIV-NEXT: [[ENTRY:.*:]] -// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]]) +// NCRDIV-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]]) // NCRDIV-NEXT: ret float [[TMP0]] // // AMDGCNSPIRV-LABEL: define spir_func noundef float @test_logf( // AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] { // AMDGCNSPIRV-NEXT: [[ENTRY:.*:]] -// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log.f32(float [[X]]) +// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call noundef addrspace(4) float @llvm.log.f32(float [[X]]) // AMDGCNSPIRV-NEXT: ret float [[TMP0]] // extern "C" __device__ float test_logf(float x) { @@ -8600,31 +8600,31 @@ extern "C" __device__ float test___fsub_rn(float x, float y) { // DEFAULT-LABEL: define dso_local noundef float @test___log10f( // DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // DEFAULT-NEXT: [[ENTRY:.*:]] -// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]]) +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]]) // DEFAULT-NEXT: ret float [[TMP0]] // // FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test___log10f( // FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // FINITEONLY-NEXT: [[ENTRY:.*:]] -// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf noundef float @llvm.log10.f32(float nofpclass(nan inf) [[X]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: define dso_local noundef float @test___log10f( // APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // APPROX-NEXT: [[ENTRY:.*:]] -// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]]) +// APPROX-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]]) // APPROX-NEXT: ret float [[TMP0]] // // NCRDIV-LABEL: define dso_local noundef float @test___log10f( // NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // NCRDIV-NEXT: [[ENTRY:.*:]] -// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log10.f32(float [[X]]) +// NCRDIV-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log10.f32(float [[X]]) // NCRDIV-NEXT: ret float [[TMP0]] // // AMDGCNSPIRV-LABEL: define spir_func noundef float @test___log10f( // AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] { // AMDGCNSPIRV-NEXT: [[ENTRY:.*:]] -// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log10.f32(float [[X]]) +// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call noundef addrspace(4) float @llvm.log10.f32(float [[X]]) // AMDGCNSPIRV-NEXT: ret float [[TMP0]] // extern "C" __device__ float test___log10f(float x) { @@ -8668,31 +8668,31 @@ extern "C" __device__ float test___log2f(float x) { // DEFAULT-LABEL: define dso_local noundef float @test___logf( // DEFAULT-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // DEFAULT-NEXT: [[ENTRY:.*:]] -// DEFAULT-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]]) +// DEFAULT-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]]) // DEFAULT-NEXT: ret float [[TMP0]] // // FINITEONLY-LABEL: define dso_local noundef nofpclass(nan inf) float @test___logf( // FINITEONLY-SAME: float noundef nofpclass(nan inf) [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // FINITEONLY-NEXT: [[ENTRY:.*:]] -// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf contract noundef float @llvm.log.f32(float nofpclass(nan inf) [[X]]) +// FINITEONLY-NEXT: [[TMP0:%.*]] = tail call nnan ninf noundef float @llvm.log.f32(float nofpclass(nan inf) [[X]]) // FINITEONLY-NEXT: ret float [[TMP0]] // // APPROX-LABEL: define dso_local noundef float @test___logf( // APPROX-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // APPROX-NEXT: [[ENTRY:.*:]] -// APPROX-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]]) +// APPROX-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]]) // APPROX-NEXT: ret float [[TMP0]] // // NCRDIV-LABEL: define dso_local noundef float @test___logf( // NCRDIV-SAME: float noundef [[X:%.*]]) local_unnamed_addr #[[ATTR3]] { // NCRDIV-NEXT: [[ENTRY:.*:]] -// NCRDIV-NEXT: [[TMP0:%.*]] = tail call contract noundef float @llvm.log.f32(float [[X]]) +// NCRDIV-NEXT: [[TMP0:%.*]] = tail call noundef float @llvm.log.f32(float [[X]]) // NCRDIV-NEXT: ret float [[TMP0]] // // AMDGCNSPIRV-LABEL: define spir_func noundef float @test___logf( // AMDGCNSPIRV-SAME: float noundef [[X:%.*]]) local_unnamed_addr addrspace(4) #[[ATTR3]] { // AMDGCNSPIRV-NEXT: [[ENTRY:.*:]] -// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call contract noundef addrspace(4) float @llvm.log.f32(float [[X]]) +// AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = tail call noundef addrspace(4) float @llvm.log.f32(float [[X]]) // AMDGCNSPIRV-NEXT: ret float [[TMP0]] // extern "C" __device__ float test___logf(float x) { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
