Skip to content

Change built-in codegen to remove contract from transcendetals #610

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.

Sign up for GitHub

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

Closed
adelejjeh wants to merge 5 commits into from
Closed

Change built-in codegen to remove contract from transcendetals #610

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
bccec0e
Change built-in codegen to remove contract from transcendetals
adelejjeh Nov 17, 2025
8f5642c
Update to only remove contract for log with AMDGCN
adelejjeh Nov 18, 2025
bc4abcf
Add log10 and fix clang unit test
adelejjeh Nov 18, 2025
1c52b4b
fix typo
adelejjeh Nov 18, 2025
abc8175
Fix small issues
adelejjeh Nov 19, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
26 changes: 25 additions & 1 deletion clang/lib/CodeGen/CGBuiltin.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -582,6 +582,23 @@ 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;
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.
Expand All @@ -596,7 +613,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;
}
}

Expand Down
38 changes: 19 additions & 19 deletions clang/test/Headers/__clang_hip_math.hip
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3720,31 +3720,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) {
Expand Down Expand Up @@ -3992,25 +3992,25 @@ 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(
Expand Down Expand Up @@ -8713,31 +8713,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) {
Expand Down Expand Up @@ -8781,31 +8781,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) {
Expand Down