From bccec0e45047ef1d58a6e73013e0c4a8bb8e47a8 Mon Sep 17 00:00:00 2001 From: Adel Ejjeh Date: Mon, 17 Nov 2025 14:31:30 -0600 Subject: [PATCH 1/5] Change built-in codegen to remove contract from transcendetals --- clang/lib/CodeGen/CGBuiltin.cpp | 53 +++++++++++++++++++++++++++++++-- 1 file changed, 50 insertions(+), 3 deletions(-) diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 16074215e8275..12dba9d471213 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -582,6 +582,28 @@ static Value *EmitISOVolatileStore(CodeGenFunction &CGF, const CallExpr *E) { return Store; } +// Check if an intrinsic is a transcendental function that uses polynomial +// refinement and should not have the contract flag (which would allow +// incorrect FMA contraction of error compensation terms). +static bool isTranscendentalIntrinsic(unsigned IntrinsicID) { + switch (IntrinsicID) { + case Intrinsic::log: + case Intrinsic::log2: + case Intrinsic::log10: + case Intrinsic::exp: + case Intrinsic::exp2: + case Intrinsic::sin: + case Intrinsic::cos: + case Intrinsic::sinh: + case Intrinsic::cosh: + case Intrinsic::tan: + case Intrinsic::tanh: + case Intrinsic::pow: + 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. @@ -596,7 +618,18 @@ 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); + + llvm::errs() << "Contract: " << Call->hasAllowContract() << "\n"; + // Transcendental intrinsics often expand to polynomial approximations with + // error compensation that require intermediate rounding. The contract flag + // would allow FMA formation that recomputes products, breaking the + // refinement algorithm. See SWDEV-561934 for details. + if (isTranscendentalIntrinsic(IntrinsicID)) { + Call->setHasAllowContract(false); + } + + return Call; } } @@ -614,7 +647,14 @@ static Value *emitBinaryMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, return CGF.Builder.CreateConstrainedFPCall(F, { Src0, Src1 }); } else { Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType()); - return CGF.Builder.CreateCall(F, { Src0, Src1 }); + CallInst *Call = CGF.Builder.CreateCall(F, { Src0, Src1 }); + + // Transcendental intrinsics should not have contract flag + if (isTranscendentalIntrinsic(IntrinsicID)) { + Call->setHasAllowContract(false); + } + + return Call; } } @@ -635,7 +675,14 @@ emitBinaryExpMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, const CallExpr *E, Function *F = CGF.CGM.getIntrinsic(IntrinsicID, {Src0->getType(), Src1->getType()}); - return CGF.Builder.CreateCall(F, {Src0, Src1}); + CallInst *Call = CGF.Builder.CreateCall(F, {Src0, Src1}); + + // Transcendental intrinsics should not have contract flag + if (isTranscendentalIntrinsic(IntrinsicID)) { + Call->setHasAllowContract(false); + } + + return Call; } // Emit an intrinsic that has 3 operands of the same type as its result. From 8f5642c6eb2e07a3042cd6e9fe8710da8d4fa90e Mon Sep 17 00:00:00 2001 From: Adel Ejjeh Date: Tue, 18 Nov 2025 15:47:20 -0600 Subject: [PATCH 2/5] Update to only remove contract for log with AMDGCN --- clang/lib/CodeGen/CGBuiltin.cpp | 46 +++++++++------------------------ 1 file changed, 12 insertions(+), 34 deletions(-) diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 12dba9d471213..c2c7171c65397 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -582,24 +582,18 @@ static Value *EmitISOVolatileStore(CodeGenFunction &CGF, const CallExpr *E) { return Store; } -// Check if an intrinsic is a transcendental function that uses polynomial -// refinement and should not have the contract flag (which would allow -// incorrect FMA contraction of error compensation terms). -static bool isTranscendentalIntrinsic(unsigned IntrinsicID) { +// 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::log2: - case Intrinsic::log10: - case Intrinsic::exp: - case Intrinsic::exp2: - case Intrinsic::sin: - case Intrinsic::cos: - case Intrinsic::sinh: - case Intrinsic::cosh: - case Intrinsic::tan: - case Intrinsic::tanh: - case Intrinsic::pow: - return true; + if ((CGF.getTarget().getTriple().isAMDGCN() || + CGF.getTarget().getTriple().isSPIRV()) && + CGF.getLangOpts().HIP) + return true; default: return false; } @@ -620,12 +614,8 @@ Value *emitUnaryMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType()); CallInst *Call = CGF.Builder.CreateCall(F, Src0); - llvm::errs() << "Contract: " << Call->hasAllowContract() << "\n"; - // Transcendental intrinsics often expand to polynomial approximations with - // error compensation that require intermediate rounding. The contract flag - // would allow FMA formation that recomputes products, breaking the - // refinement algorithm. See SWDEV-561934 for details. - if (isTranscendentalIntrinsic(IntrinsicID)) { + // Check if the intrinsic is unsafe to contract + if (isUnsafeToContract(IntrinsicID, CGF)) { Call->setHasAllowContract(false); } @@ -648,12 +638,6 @@ static Value *emitBinaryMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, } else { Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType()); CallInst *Call = CGF.Builder.CreateCall(F, { Src0, Src1 }); - - // Transcendental intrinsics should not have contract flag - if (isTranscendentalIntrinsic(IntrinsicID)) { - Call->setHasAllowContract(false); - } - return Call; } } @@ -676,12 +660,6 @@ emitBinaryExpMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, const CallExpr *E, Function *F = CGF.CGM.getIntrinsic(IntrinsicID, {Src0->getType(), Src1->getType()}); CallInst *Call = CGF.Builder.CreateCall(F, {Src0, Src1}); - - // Transcendental intrinsics should not have contract flag - if (isTranscendentalIntrinsic(IntrinsicID)) { - Call->setHasAllowContract(false); - } - return Call; } From bc4abcf2895bc560754646535b1da99866f5fa22 Mon Sep 17 00:00:00 2001 From: Adel Ejjeh Date: Tue, 18 Nov 2025 16:54:57 -0600 Subject: [PATCH 3/5] Add log10 and fix clang unit test --- clang/lib/CodeGen/CGBuiltin.cpp | 2 ++ clang/test/Headers/__clang_hip_math.hip | 40 ++++++++++++------------- 2 files changed, 22 insertions(+), 20 deletions(-) diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index c2c7171c65397..7e88ae264bb4a 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -35,6 +35,7 @@ #include "llvm/IR/MatrixBuilder.h" #include "llvm/Support/ConvertUTF.h" #include "llvm/Support/ScopedPrinter.h" +#include #include #include @@ -590,6 +591,7 @@ static bool isUnsafeToContract(unsigned IntrinsicID, CodeGenFunction &CGF) { // 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) diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index 22c0689a4552e..f9817ce038575 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -45,7 +45,7 @@ // RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=AMDGCNSPIRV %s #define BOOL_TYPE int -typedef unsigned long long uint64_t; +// typedef unsigned long long uint64_t; // CHECK-LABEL: define dso_local i64 @test___make_mantissa_base8( // CHECK-SAME: ptr noundef readonly captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] { @@ -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) { @@ -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( @@ -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) { @@ -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) { From 1c52b4b23d24e79e2ded0f1139b83eb9e37afa06 Mon Sep 17 00:00:00 2001 From: Adel Ejjeh Date: Tue, 18 Nov 2025 17:07:01 -0600 Subject: [PATCH 4/5] fix typo --- clang/test/Headers/__clang_hip_math.hip | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/Headers/__clang_hip_math.hip b/clang/test/Headers/__clang_hip_math.hip index f9817ce038575..e54b75810a6f3 100644 --- a/clang/test/Headers/__clang_hip_math.hip +++ b/clang/test/Headers/__clang_hip_math.hip @@ -45,7 +45,7 @@ // RUN: -D__HIPCC_RTC__ | FileCheck -check-prefixes=AMDGCNSPIRV %s #define BOOL_TYPE int -// typedef unsigned long long uint64_t; +typedef unsigned long long uint64_t; // CHECK-LABEL: define dso_local i64 @test___make_mantissa_base8( // CHECK-SAME: ptr noundef readonly captures(none) [[P:%.*]]) local_unnamed_addr #[[ATTR2:[0-9]+]] { From abc8175b896d8d7b623cb1c8ba9cbdc84b680feb Mon Sep 17 00:00:00 2001 From: Adel Ejjeh Date: Wed, 19 Nov 2025 11:03:36 -0600 Subject: [PATCH 5/5] Fix small issues --- clang/lib/CodeGen/CGBuiltin.cpp | 7 ++----- 1 file changed, 2 insertions(+), 5 deletions(-) diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 7e88ae264bb4a..b3aca39e4b18a 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -35,7 +35,6 @@ #include "llvm/IR/MatrixBuilder.h" #include "llvm/Support/ConvertUTF.h" #include "llvm/Support/ScopedPrinter.h" -#include #include #include @@ -639,8 +638,7 @@ static Value *emitBinaryMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, return CGF.Builder.CreateConstrainedFPCall(F, { Src0, Src1 }); } else { Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType()); - CallInst *Call = CGF.Builder.CreateCall(F, { Src0, Src1 }); - return Call; + return CGF.Builder.CreateCall(F, { Src0, Src1 }); } } @@ -661,8 +659,7 @@ emitBinaryExpMaybeConstrainedFPBuiltin(CodeGenFunction &CGF, const CallExpr *E, Function *F = CGF.CGM.getIntrinsic(IntrinsicID, {Src0->getType(), Src1->getType()}); - CallInst *Call = CGF.Builder.CreateCall(F, {Src0, Src1}); - return Call; + return CGF.Builder.CreateCall(F, {Src0, Src1}); } // Emit an intrinsic that has 3 operands of the same type as its result.