@@ -15540,38 +15540,6 @@ static bool isOverflowingIntegerType(ASTContext &Ctx, QualType T) {
1554015540 return Ctx.getIntWidth(T) >= Ctx.getIntWidth(Ctx.IntTy);
1554115541}
1554215542
15543- static Expr *ExpandAMDGPUPredicateBI(ASTContext &Ctx, CallExpr *CE) {
15544- if (!CE->getBuiltinCallee())
15545- return CXXBoolLiteralExpr::Create(Ctx, false, Ctx.BoolTy, CE->getExprLoc());
15546-
15547- if (Ctx.getTargetInfo().getTriple().isSPIRV()) {
15548- CE->setType(Ctx.getLogicalOperationType());
15549- return CE;
15550- }
15551-
15552- bool P = false;
15553- auto &TI = Ctx.getTargetInfo();
15554-
15555- if (CE->getDirectCallee()->getName() == "__builtin_amdgcn_processor_is") {
15556- auto *GFX = dyn_cast<StringLiteral>(CE->getArg(0)->IgnoreParenCasts());
15557- auto TID = TI.getTargetID();
15558- if (GFX && TID) {
15559- auto N = GFX->getString();
15560- P = TI.isValidCPUName(GFX->getString()) && TID->find(N) == 0;
15561- }
15562- } else {
15563- auto *FD = cast<FunctionDecl>(CE->getArg(0)->getReferencedDeclOfCallee());
15564-
15565- StringRef RF = Ctx.BuiltinInfo.getRequiredFeatures(FD->getBuiltinID());
15566- llvm::StringMap<bool> CF;
15567- Ctx.getFunctionFeatureMap(CF, FD);
15568-
15569- P = Builtin::evaluateRequiredTargetFeatures(RF, CF);
15570- }
15571-
15572- return CXXBoolLiteralExpr::Create(Ctx, P, Ctx.BoolTy, CE->getExprLoc());
15573- }
15574-
1557515543ExprResult Sema::CreateBuiltinUnaryOp(SourceLocation OpLoc,
1557615544 UnaryOperatorKind Opc, Expr *InputExpr,
1557715545 bool IsAfterAmp) {
@@ -20465,88 +20433,6 @@ void Sema::DiagnoseEqualityWithExtraParens(ParenExpr *ParenE) {
2046520433 }
2046620434}
2046720435
20468- static bool ValidateAMDGPUPredicateBI(Sema &Sema, CallExpr *CE) {
20469- if (CE->getDirectCallee()->getName() == "__builtin_amdgcn_processor_is") {
20470- auto *GFX = dyn_cast<StringLiteral>(CE->getArg(0)->IgnoreParenCasts());
20471- if (!GFX) {
20472- Sema.Diag(CE->getExprLoc(),
20473- diag::err_amdgcn_processor_is_arg_not_literal);
20474- return false;
20475- }
20476- auto N = GFX->getString();
20477- if (!Sema.getASTContext().getTargetInfo().isValidCPUName(N) &&
20478- (!Sema.getASTContext().getAuxTargetInfo() ||
20479- !Sema.getASTContext().getAuxTargetInfo()->isValidCPUName(N))) {
20480- Sema.Diag(CE->getExprLoc(),
20481- diag::err_amdgcn_processor_is_arg_invalid_value)
20482- << N;
20483- return false;
20484- }
20485- } else {
20486- auto *Arg = CE->getArg(0);
20487- if (!Arg || Arg->getType() != Sema.getASTContext().BuiltinFnTy) {
20488- Sema.Diag(CE->getExprLoc(),
20489- diag::err_amdgcn_is_invocable_arg_invalid_value)
20490- << Arg;
20491- return false;
20492- }
20493- }
20494-
20495- return true;
20496- }
20497-
20498- static Expr *MaybeHandleAMDGPUPredicateBI(Sema &Sema, Expr *E, bool &Invalid) {
20499- if (auto *UO = dyn_cast<UnaryOperator>(E)) {
20500- auto *SE = dyn_cast<CallExpr>(UO->getSubExpr());
20501- if (IsAMDGPUPredicateBI(SE)) {
20502- assert(UO->getOpcode() == UnaryOperator::Opcode::UO_LNot &&
20503- "__builtin_amdgcn_processor_is and __builtin_amdgcn_is_invocable "
20504- "can only be used as operands of logical ops!");
20505-
20506- if (!ValidateAMDGPUPredicateBI(Sema, SE)) {
20507- Invalid = true;
20508- return nullptr;
20509- }
20510-
20511- UO->setSubExpr(ExpandAMDGPUPredicateBI(Sema.getASTContext(), SE));
20512- UO->setType(Sema.getASTContext().getLogicalOperationType());
20513-
20514- return UO;
20515- }
20516- }
20517- if (auto *BO = dyn_cast<BinaryOperator>(E)) {
20518- auto *LHS = dyn_cast<CallExpr>(BO->getLHS());
20519- auto *RHS = dyn_cast<CallExpr>(BO->getRHS());
20520- if (IsAMDGPUPredicateBI(LHS) && IsAMDGPUPredicateBI(RHS)) {
20521- assert(BO->isLogicalOp() &&
20522- "__builtin_amdgcn_processor_is and __builtin_amdgcn_is_invocable "
20523- "can only be used as operands of logical ops!");
20524-
20525- if (!ValidateAMDGPUPredicateBI(Sema, LHS) ||
20526- !ValidateAMDGPUPredicateBI(Sema, RHS)) {
20527- Invalid = true;
20528- return nullptr;
20529- }
20530-
20531- BO->setLHS(ExpandAMDGPUPredicateBI(Sema.getASTContext(), LHS));
20532- BO->setRHS(ExpandAMDGPUPredicateBI(Sema.getASTContext(), RHS));
20533- BO->setType(Sema.getASTContext().getLogicalOperationType());
20534-
20535- return BO;
20536- }
20537- }
20538- if (auto *CE = dyn_cast<CallExpr>(E))
20539- if (IsAMDGPUPredicateBI(CE)) {
20540- if (!ValidateAMDGPUPredicateBI(Sema, CE)) {
20541- Invalid = true;
20542- return nullptr;
20543- }
20544- return ExpandAMDGPUPredicateBI(Sema.getASTContext(), CE);
20545- }
20546-
20547- return nullptr;
20548- }
20549-
2055020436ExprResult Sema::CheckBooleanCondition(SourceLocation Loc, Expr *E,
2055120437 bool IsConstexpr) {
2055220438 DiagnoseAssignmentAsCondition(E);
0 commit comments