@@ -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,
@@ -20429,7 +20417,8 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
2042920417 }
2043020418 case AMDGPU::BI__builtin_amdgcn_bitop3_b32:
2043120419 case AMDGPU::BI__builtin_amdgcn_bitop3_b16:
20432- return emitQuaternaryBuiltin(*this, E, Intrinsic::amdgcn_bitop3);
20420+ return emitBuiltinWithOneOverloadedType<4>(*this, E,
20421+ Intrinsic::amdgcn_bitop3);
2043320422 case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc:
2043420423 return emitBuiltinWithOneOverloadedType<4>(
2043520424 *this, E, Intrinsic::amdgcn_make_buffer_rsrc);
0 commit comments