llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: lijinpei-amd <details> <summary>Changes</summary> 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 --- Full diff: https://github.com/llvm/llvm-project/pull/199963.diff 2 Files Affected: - (modified) clang/lib/Sema/SemaAMDGPU.cpp (+13-9) - (added) clang/test/AST/ast-print-amdgcn-predicate.c (+104) ``````````diff diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 1d2b3898c92d6..3f0ce5b3a080b 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -760,14 +760,21 @@ 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(); + // A plain IntegerLiteral whose type is _Bool/bool is not a shape any other + // clang producer creates, and consumers (e.g. StmtPrinter) assume it cannot + // occur — see issue #199563. + auto MakePredicateLiteral = [&](bool Val) -> Expr * { + if (Ctx.getLangOpts().Bool) + return SemaRef + .ActOnCXXBoolLiteral(Loc, Val ? tok::kw_true : tok::kw_false) + .get(); + return SemaRef.ActOnIntegerConstant(Loc, Val ? 1 : 0).get(); + }; + if (!CE->getBuiltinCallee()) - return *ExpandedPredicates - .insert(IntegerLiteral::Create(Ctx, False, BoolTy, Loc)) - .first; + return *ExpandedPredicates.insert(MakePredicateLiteral(false)).first; bool P = false; unsigned BI = CE->getBuiltinCallee(); @@ -824,10 +831,7 @@ Expr *SemaAMDGPU::ExpandAMDGPUPredicateBuiltIn(Expr *E) { P = Builtin::evaluateRequiredTargetFeatures(RF, CF); } - return *ExpandedPredicates - .insert( - IntegerLiteral::Create(Ctx, P ? True : False, BoolTy, Loc)) - .first; + return *ExpandedPredicates.insert(MakePredicateLiteral(P)).first; } bool SemaAMDGPU::IsPredicate(Expr *E) const { 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..1da7f78f8ac2f --- /dev/null +++ b/clang/test/AST/ast-print-amdgcn-predicate.c @@ -0,0 +1,104 @@ +// REQUIRES: amdgpu-registered-target + +// Regression test for https://github.com/llvm/llvm-project/issues/199563: +// SemaAMDGPU::ExpandAMDGPUPredicateBuiltIn used to synthesize an +// IntegerLiteral of type _Bool/bool to replace +// __builtin_amdgcn_is_invocable / __builtin_amdgcn_processor_is calls. No +// other clang producer creates that shape, and StmtPrinter has no case for +// it, so -ast-print on the resulting if-condition crashed via +// llvm_unreachable. The fix synthesises a CXXBoolLiteralExpr when +// LangOpts.Bool is set (C++, C23, OpenCL, HIP) and an int IntegerLiteral +// otherwise (pre-C23 C). The matrix below covers every dispatch arm. + +// C11 (pre-C23, LangOpts.Bool=false): IntegerLiteral 'int'. +// 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 + +// C23 (LangOpts.Bool=true): CXXBoolLiteralExpr 'bool'. +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -std=c23 \ +// RUN: -ast-print %s | FileCheck --check-prefix=BOOL-TRUE %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -std=c23 \ +// RUN: -ast-print %s | FileCheck --check-prefix=BOOL-FALSE %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -std=c23 \ +// RUN: -ast-dump %s | FileCheck --check-prefix=BOOL-DUMP %s + +// C++17 (LangOpts.Bool=true): same as C23. +// 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 + +// HIP device compilation (the originally-reported scenario): C++ under HIP +// mode with -fcuda-is-device. LangOpts.Bool=true, so same as C23/C++17. +// 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 + +// `__builtin_amdgcn_is_invocable` evaluates its argument unevaluated, so the +// builtin reference itself does not need to be callable in this context. +__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; +} + +// -ast-print: only the if-condition shape matters; anchor each check to its +// function so the two if-stmts can't cross-match. + +// 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) + +// -ast-dump: confirm the AST node class under each function's IfStmt. + +// INT-DUMP-LABEL: FunctionDecl {{.*}} use_is_invocable +// INT-DUMP: IfStmt +// INT-DUMP-NEXT: IntegerLiteral {{.*}} 'int' 1 +// INT-DUMP-LABEL: FunctionDecl {{.*}} use_processor_is +// INT-DUMP: IfStmt +// 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 `````````` </details> https://github.com/llvm/llvm-project/pull/199963 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
