@@ -774,18 +774,6 @@ static Value *emitBuiltinWithOneOverloadedType(CodeGenFunction &CGF,
774774 return CGF.Builder.CreateCall(F, Args, Name);
775775}
776776
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-
789777// Emit an intrinsic that has 1 float or double operand, and 1 integer.
790778static Value *emitFPIntBuiltin(CodeGenFunction &CGF,
791779 const CallExpr *E,
@@ -20443,7 +20431,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
2044320431 }
2044420432 case AMDGPU::BI__builtin_amdgcn_bitop3_b32:
2044520433 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);
2044720436 case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc:
2044820437 return emitBuiltinWithOneOverloadedType<4>(
2044920438 *this, E, Intrinsic::amdgcn_make_buffer_rsrc);
0 commit comments