Author: lijinpei-amd Date: 2026-05-27T21:12:44-05:00 New Revision: f9185454a9a70fd351cb8a1ba3b2b3e33e064c06
URL: https://github.com/llvm/llvm-project/commit/f9185454a9a70fd351cb8a1ba3b2b3e33e064c06 DIFF: https://github.com/llvm/llvm-project/commit/f9185454a9a70fd351cb8a1ba3b2b3e33e064c06.diff LOG: [clang][AMDGPU] Fix -ast-print crash on expanded predicate builtins (#199963) ExpandAMDGPUPredicateBuiltIn synthesized an IntegerLiteral typed _Bool/bool — a shape no other producer creates, and one that StmtPrinter::VisitIntegerLiteral has no case for. -ast-print on the resulting if-condition hit llvm_unreachable. Emit the canonical boolean literal instead: - C++, C23, OpenCL, HIP: CXXBoolLiteralExpr 'bool' - pre-C23 C: IntegerLiteral 'int' In the C case this matches what <stdbool.h>'s true/false macros expand to. Fixes #199563 Added: clang/test/AST/ast-print-amdgcn-predicate.c Modified: clang/include/clang/Sema/Sema.h clang/include/clang/Sema/SemaObjC.h clang/lib/Sema/SemaAMDGPU.cpp clang/lib/Sema/SemaExpr.cpp clang/lib/Sema/SemaExprObjC.cpp Removed: ################################################################################ diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index e71794b2d92c9..79ad4e7aa75c0 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -8577,6 +8577,9 @@ class Sema final : public SemaBase { /// ActOnCXXBoolLiteral - Parse {true,false} literals. ExprResult ActOnCXXBoolLiteral(SourceLocation OpLoc, tok::TokenKind Kind); + /// Build a boolean-typed literal expression. + ExprResult BuildBoolLiteral(SourceLocation Loc, bool Value); + /// ActOnCXXNullPtrLiteral - Parse 'nullptr'. ExprResult ActOnCXXNullPtrLiteral(SourceLocation Loc); diff --git a/clang/include/clang/Sema/SemaObjC.h b/clang/include/clang/Sema/SemaObjC.h index ed08ff0acf89d..f5cb8ea4b034a 100644 --- a/clang/include/clang/Sema/SemaObjC.h +++ b/clang/include/clang/Sema/SemaObjC.h @@ -658,7 +658,10 @@ class SemaObjC : public SemaBase { /// or "id" if NSNumber is unavailable. ExprResult BuildObjCNumericLiteral(SourceLocation AtLoc, Expr *Number); ExprResult ActOnObjCBoolLiteral(SourceLocation AtLoc, SourceLocation ValueLoc, - bool Value); + bool Value) { + return BuildObjCNumericLiteral( + AtLoc, SemaRef.BuildBoolLiteral(ValueLoc, Value).get()); + } ExprResult BuildObjCArrayLiteral(SourceRange SR, MultiExprArg Elements); /// BuildObjCBoxedExpr - builds an ObjCBoxedExpr AST node for the diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 1d2b3898c92d6..c7420f0559a51 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -760,13 +760,11 @@ Expr *SemaAMDGPU::ExpandAMDGPUPredicateBuiltIn(Expr *E) { CallExpr *CE = cast<CallExpr>(E->IgnoreParens()); ASTContext &Ctx = getASTContext(); QualType BoolTy = Ctx.getLogicalOperationType(); - llvm::APInt False = llvm::APInt::getZero(Ctx.getIntWidth(BoolTy)); - llvm::APInt True = llvm::APInt::getAllOnes(Ctx.getIntWidth(BoolTy)); SourceLocation Loc = CE->getExprLoc(); if (!CE->getBuiltinCallee()) return *ExpandedPredicates - .insert(IntegerLiteral::Create(Ctx, False, BoolTy, Loc)) + .insert(SemaRef.BuildBoolLiteral(Loc, false).get()) .first; bool P = false; @@ -824,9 +822,7 @@ Expr *SemaAMDGPU::ExpandAMDGPUPredicateBuiltIn(Expr *E) { P = Builtin::evaluateRequiredTargetFeatures(RF, CF); } - return *ExpandedPredicates - .insert( - IntegerLiteral::Create(Ctx, P ? True : False, BoolTy, Loc)) + return *ExpandedPredicates.insert(SemaRef.BuildBoolLiteral(Loc, P).get()) .first; } diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index db0fbdc546fbc..ad6e7183cb3a4 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -3737,6 +3737,20 @@ ExprResult Sema::ActOnIntegerConstant(SourceLocation Loc, int64_t Val) { Context.IntTy, Loc); } +ExprResult Sema::BuildBoolLiteral(SourceLocation Loc, bool Value) { + ExprResult Inner; + if (getLangOpts().CPlusPlus) { + Inner = ActOnCXXBoolLiteral(Loc, Value ? tok::kw_true : tok::kw_false); + } else { + // C doesn't actually have a way to represent literal values of type + // _Bool. So, we'll use 0/1 and implicit cast to _Bool. + Inner = ActOnIntegerConstant(Loc, Value ? 1 : 0); + Inner = + ImpCastExprToType(Inner.get(), Context.BoolTy, CK_IntegralToBoolean); + } + return Inner; +} + static Expr *BuildFloatingLiteral(Sema &S, NumericLiteralParser &Literal, QualType Ty, SourceLocation Loc) { const llvm::fltSemantics &Format = S.Context.getFloatTypeSemantics(Ty); diff --git a/clang/lib/Sema/SemaExprObjC.cpp b/clang/lib/Sema/SemaExprObjC.cpp index bc191359db839..c166e0c9e7b72 100644 --- a/clang/lib/Sema/SemaExprObjC.cpp +++ b/clang/lib/Sema/SemaExprObjC.cpp @@ -434,24 +434,6 @@ ExprResult SemaObjC::BuildObjCNumericLiteral(SourceLocation AtLoc, return SemaRef.MaybeBindToTemporary(NumberLiteral); } -ExprResult SemaObjC::ActOnObjCBoolLiteral(SourceLocation AtLoc, - SourceLocation ValueLoc, bool Value) { - ASTContext &Context = getASTContext(); - ExprResult Inner; - if (getLangOpts().CPlusPlus) { - Inner = SemaRef.ActOnCXXBoolLiteral(ValueLoc, - Value ? tok::kw_true : tok::kw_false); - } else { - // C doesn't actually have a way to represent literal values of type - // _Bool. So, we'll use 0/1 and implicit cast to _Bool. - Inner = SemaRef.ActOnIntegerConstant(ValueLoc, Value ? 1 : 0); - Inner = SemaRef.ImpCastExprToType(Inner.get(), Context.BoolTy, - CK_IntegralToBoolean); - } - - return BuildObjCNumericLiteral(AtLoc, Inner.get()); -} - /// Check that the given expression is a valid element of an Objective-C /// collection literal. static ExprResult CheckObjCCollectionLiteralElement(Sema &S, Expr *Element, diff --git a/clang/test/AST/ast-print-amdgcn-predicate.c b/clang/test/AST/ast-print-amdgcn-predicate.c new file mode 100644 index 0000000000000..2b2c6e21ee2ad --- /dev/null +++ b/clang/test/AST/ast-print-amdgcn-predicate.c @@ -0,0 +1,88 @@ +// REQUIRES: amdgpu-registered-target + +// Regression test for issue #199563: -ast-print on +// __builtin_amdgcn_is_invocable / __builtin_amdgcn_processor_is used to +// crash because the expansion produced an IntegerLiteral typed _Bool/bool. + +// C (C11 and C23): IntegerLiteral 'int' implicit-cast to _Bool. +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -std=c11 \ +// RUN: -ast-print %s | FileCheck --check-prefix=INT-TRUE %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -std=c11 \ +// RUN: -ast-print %s | FileCheck --check-prefix=INT-FALSE %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -std=c11 \ +// RUN: -ast-dump %s | FileCheck --check-prefix=INT-DUMP %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -std=c23 \ +// RUN: -ast-print %s | FileCheck --check-prefix=INT-TRUE %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -std=c23 \ +// RUN: -ast-print %s | FileCheck --check-prefix=INT-FALSE %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -std=c23 \ +// RUN: -ast-dump %s | FileCheck --check-prefix=INT-DUMP %s + +// C++ / HIP device: CXXBoolLiteralExpr 'bool'. +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x c++ \ +// RUN: -std=c++17 -ast-print %s | FileCheck --check-prefix=BOOL-TRUE %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -x c++ \ +// RUN: -std=c++17 -ast-print %s | FileCheck --check-prefix=BOOL-FALSE %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x c++ \ +// RUN: -std=c++17 -ast-dump %s | FileCheck --check-prefix=BOOL-DUMP %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x hip \ +// RUN: -fcuda-is-device -ast-print %s \ +// RUN: | FileCheck --check-prefix=BOOL-TRUE %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -x hip \ +// RUN: -fcuda-is-device -ast-print %s \ +// RUN: | FileCheck --check-prefix=BOOL-FALSE %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x hip \ +// RUN: -fcuda-is-device -ast-dump %s \ +// RUN: | FileCheck --check-prefix=BOOL-DUMP %s + +#if defined(__HIP__) || defined(__CUDA__) +#define __device__ __attribute__((device)) +#else +#define __device__ +#endif + +__device__ void use_is_invocable(void) { + if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_flat_atomic_fadd_f32)) + (void)0; +} + +__device__ void use_processor_is(void) { + if (__builtin_amdgcn_processor_is("gfx942")) + (void)0; +} + +// INT-TRUE-LABEL: use_is_invocable +// INT-TRUE: if (1) +// INT-TRUE-LABEL: use_processor_is +// INT-TRUE: if (1) + +// INT-FALSE-LABEL: use_is_invocable +// INT-FALSE: if (0) +// INT-FALSE-LABEL: use_processor_is +// INT-FALSE: if (0) + +// BOOL-TRUE-LABEL: use_is_invocable +// BOOL-TRUE: if (true) +// BOOL-TRUE-LABEL: use_processor_is +// BOOL-TRUE: if (true) + +// BOOL-FALSE-LABEL: use_is_invocable +// BOOL-FALSE: if (false) +// BOOL-FALSE-LABEL: use_processor_is +// BOOL-FALSE: if (false) + +// INT-DUMP-LABEL: FunctionDecl {{.*}} use_is_invocable +// INT-DUMP: IfStmt +// INT-DUMP-NEXT: ImplicitCastExpr {{.*}} {{'_Bool'|'bool'}} <IntegralToBoolean> +// INT-DUMP-NEXT: IntegerLiteral {{.*}} 'int' 1 +// INT-DUMP-LABEL: FunctionDecl {{.*}} use_processor_is +// INT-DUMP: IfStmt +// INT-DUMP-NEXT: ImplicitCastExpr {{.*}} {{'_Bool'|'bool'}} <IntegralToBoolean> +// INT-DUMP-NEXT: IntegerLiteral {{.*}} 'int' 1 + +// BOOL-DUMP-LABEL: FunctionDecl {{.*}} use_is_invocable +// BOOL-DUMP: IfStmt +// BOOL-DUMP-NEXT: CXXBoolLiteralExpr {{.*}} 'bool' true +// BOOL-DUMP-LABEL: FunctionDecl {{.*}} use_processor_is +// BOOL-DUMP: IfStmt +// BOOL-DUMP-NEXT: CXXBoolLiteralExpr {{.*}} 'bool' true _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
