This is an archive of the discontinued LLVM Phabricator instance.

Differential D74673

CGBuiltin: Remove uses of deprecated CreateCall overloads
ClosedPublic

Authored by nhaehnle on Feb 15 2020, 12:50 PM.

Details

Reviewers
t.p.northover
arsenm
Commits
rGbf197304a66a: CGBuiltin: Remove uses of deprecated CreateCall overloads

Diff Detail

Event Timeline

nhaehnle created this revision.Feb 15 2020, 12:50 PM
Herald added a project: Restricted Project. · View Herald TranscriptFeb 15 2020, 12:50 PM
Herald added a subscriber: cfe-commits. · View Herald Transcript
nhaehnle added a child revision: D74676: OpaquePtr: IRBuilder: Add deprecated attribute on CreateCall methods.Feb 15 2020, 12:51 PM
arsenm accepted this revision.Feb 15 2020, 1:05 PM
This revision is now accepted and ready to land.Feb 15 2020, 1:05 PM
Herald added a subscriber: wdng. · View Herald TranscriptFeb 15 2020, 1:05 PM
Harbormaster failed remote builds in B46578: Diff 244835!Feb 15 2020, 1:05 PM
Closed by commit rGbf197304a66a: CGBuiltin: Remove uses of deprecated CreateCall overloads (authored by nhaehnle). · Explain WhyFeb 17 2020, 3:30 PM
This revision was automatically updated to reflect the committed changes.

Revision Contents

PathSize
clang/
lib/
CodeGen/
4 lines

Diff 245047

clang/lib/CodeGen/CGBuiltin.cpp

  • This file is larger than 256 KB, so syntax highlighting is disabled by default.
Show First 20 Lines • Show All 13,297 Lines • ▼ Show 20 LinesValue *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_uicmpl: case AMDGPU::BI__builtin_amdgcn_uicmpl:
case AMDGPU::BI__builtin_amdgcn_sicmp: case AMDGPU::BI__builtin_amdgcn_sicmp:
case AMDGPU::BI__builtin_amdgcn_sicmpl: { case AMDGPU::BI__builtin_amdgcn_sicmpl: {
llvm::Value *Src0 = EmitScalarExpr(E->getArg(0)); llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));
llvm::Value *Src1 = EmitScalarExpr(E->getArg(1)); llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));
llvm::Value *Src2 = EmitScalarExpr(E->getArg(2)); llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));
// FIXME-GFX10: How should 32 bit mask be handled? // FIXME-GFX10: How should 32 bit mask be handled?
Value *F = CGM.getIntrinsic(Intrinsic::amdgcn_icmp, Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_icmp,
{ Builder.getInt64Ty(), Src0->getType() }); { Builder.getInt64Ty(), Src0->getType() });
return Builder.CreateCall(F, { Src0, Src1, Src2 }); return Builder.CreateCall(F, { Src0, Src1, Src2 });
} }
case AMDGPU::BI__builtin_amdgcn_fcmp: case AMDGPU::BI__builtin_amdgcn_fcmp:
case AMDGPU::BI__builtin_amdgcn_fcmpf: { case AMDGPU::BI__builtin_amdgcn_fcmpf: {
llvm::Value *Src0 = EmitScalarExpr(E->getArg(0)); llvm::Value *Src0 = EmitScalarExpr(E->getArg(0));
llvm::Value *Src1 = EmitScalarExpr(E->getArg(1)); llvm::Value *Src1 = EmitScalarExpr(E->getArg(1));
llvm::Value *Src2 = EmitScalarExpr(E->getArg(2)); llvm::Value *Src2 = EmitScalarExpr(E->getArg(2));
// FIXME-GFX10: How should 32 bit mask be handled? // FIXME-GFX10: How should 32 bit mask be handled?
Value *F = CGM.getIntrinsic(Intrinsic::amdgcn_fcmp, Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_fcmp,
{ Builder.getInt64Ty(), Src0->getType() }); { Builder.getInt64Ty(), Src0->getType() });
return Builder.CreateCall(F, { Src0, Src1, Src2 }); return Builder.CreateCall(F, { Src0, Src1, Src2 });
} }
case AMDGPU::BI__builtin_amdgcn_class: case AMDGPU::BI__builtin_amdgcn_class:
case AMDGPU::BI__builtin_amdgcn_classf: case AMDGPU::BI__builtin_amdgcn_classf:
case AMDGPU::BI__builtin_amdgcn_classh: case AMDGPU::BI__builtin_amdgcn_classh:
return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class); return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_class);
case AMDGPU::BI__builtin_amdgcn_fmed3f: case AMDGPU::BI__builtin_amdgcn_fmed3f:
▲ Show 20 Lines • Show All 1,831 Lines • Show Last 20 Lines