@@ -774,18 +774,6 @@ static Value *emitBuiltinWithOneOverloadedType(CodeGenFunction &CGF,
774
774
return CGF.Builder.CreateCall(F, Args, Name);
775
775
}
776
776
777
- // Emit an intrinsic that has 4 operands of the same type as its result.
778
- static Value *emitQuaternaryBuiltin(CodeGenFunction &CGF, const CallExpr *E,
779
- unsigned IntrinsicID) {
780
- llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
781
- llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));
782
- llvm::Value *Src2 = CGF.EmitScalarExpr(E->getArg(2));
783
- llvm::Value *Src3 = CGF.EmitScalarExpr(E->getArg(3));
784
-
785
- Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
786
- return CGF.Builder.CreateCall(F, {Src0, Src1, Src2, Src3});
787
- }
788
-
789
777
// Emit an intrinsic that has 1 float or double operand, and 1 integer.
790
778
static Value *emitFPIntBuiltin(CodeGenFunction &CGF,
791
779
const CallExpr *E,
@@ -20443,7 +20431,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
20443
20431
}
20444
20432
case AMDGPU::BI__builtin_amdgcn_bitop3_b32:
20445
20433
case AMDGPU::BI__builtin_amdgcn_bitop3_b16:
20446
- return emitQuaternaryBuiltin(*this, E, Intrinsic::amdgcn_bitop3);
20434
+ return emitBuiltinWithOneOverloadedType<4>(*this, E,
20435
+ Intrinsic::amdgcn_bitop3);
20447
20436
case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc:
20448
20437
return emitBuiltinWithOneOverloadedType<4>(
20449
20438
*this, E, Intrinsic::amdgcn_make_buffer_rsrc);
0 commit comments