Skip to content

Commit 002f4eb

Browse files
lijinpei-amdlijinpei-amd
committed
[clang][AMDGPU] Fix -ast-print crash on expanded predicate builtins
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. Factor the canonical-boolean-literal logic that already lived inline in SemaObjC::ActOnObjCBoolLiteral into a shared Sema::BuildBoolLiteral helper, and use it from both ObjC and the AMDGPU predicate expansion: - C++: CXXBoolLiteralExpr 'bool' - C: IntegerLiteral 'int' implicit-cast to _Bool Fixes #199563
1 parent 3a8b5e3 commit 002f4eb

6 files changed

Lines changed: 111 additions & 25 deletions

File tree

‎clang/include/clang/Sema/Sema.h‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -8577,6 +8577,9 @@ class Sema final : public SemaBase {
85778577
/// ActOnCXXBoolLiteral - Parse {true,false} literals.
85788578
ExprResult ActOnCXXBoolLiteral(SourceLocation OpLoc, tok::TokenKind Kind);
85798579

8580+
/// Build a boolean-typed literal expression.
8581+
ExprResult BuildBoolLiteral(SourceLocation Loc, bool Value);
8582+
85808583
/// ActOnCXXNullPtrLiteral - Parse 'nullptr'.
85818584
ExprResult ActOnCXXNullPtrLiteral(SourceLocation Loc);
85828585

‎clang/include/clang/Sema/SemaObjC.h‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -658,7 +658,10 @@ class SemaObjC : public SemaBase {
658658
/// or "id" if NSNumber is unavailable.
659659
ExprResult BuildObjCNumericLiteral(SourceLocation AtLoc, Expr *Number);
660660
ExprResult ActOnObjCBoolLiteral(SourceLocation AtLoc, SourceLocation ValueLoc,
661-
bool Value);
661+
bool Value) {
662+
return BuildObjCNumericLiteral(
663+
AtLoc, SemaRef.BuildBoolLiteral(ValueLoc, Value).get());
664+
}
662665
ExprResult BuildObjCArrayLiteral(SourceRange SR, MultiExprArg Elements);
663666

664667
/// BuildObjCBoxedExpr - builds an ObjCBoxedExpr AST node for the

‎clang/lib/Sema/SemaAMDGPU.cpp‎

Lines changed: 2 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -760,13 +760,11 @@ Expr *SemaAMDGPU::ExpandAMDGPUPredicateBuiltIn(Expr *E) {
760760
CallExpr *CE = cast<CallExpr>(E->IgnoreParens());
761761
ASTContext &Ctx = getASTContext();
762762
QualType BoolTy = Ctx.getLogicalOperationType();
763-
llvm::APInt False = llvm::APInt::getZero(Ctx.getIntWidth(BoolTy));
764-
llvm::APInt True = llvm::APInt::getAllOnes(Ctx.getIntWidth(BoolTy));
765763
SourceLocation Loc = CE->getExprLoc();
766764

767765
if (!CE->getBuiltinCallee())
768766
return *ExpandedPredicates
769-
.insert(IntegerLiteral::Create(Ctx, False, BoolTy, Loc))
767+
.insert(SemaRef.BuildBoolLiteral(Loc, false).get())
770768
.first;
771769

772770
bool P = false;
@@ -824,9 +822,7 @@ Expr *SemaAMDGPU::ExpandAMDGPUPredicateBuiltIn(Expr *E) {
824822
P = Builtin::evaluateRequiredTargetFeatures(RF, CF);
825823
}
826824

827-
return *ExpandedPredicates
828-
.insert(
829-
IntegerLiteral::Create(Ctx, P ? True : False, BoolTy, Loc))
825+
return *ExpandedPredicates.insert(SemaRef.BuildBoolLiteral(Loc, P).get())
830826
.first;
831827
}
832828

‎clang/lib/Sema/SemaExpr.cpp‎

Lines changed: 14 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3737,6 +3737,20 @@ ExprResult Sema::ActOnIntegerConstant(SourceLocation Loc, int64_t Val) {
37373737
Context.IntTy, Loc);
37383738
}
37393739

3740+
ExprResult Sema::BuildBoolLiteral(SourceLocation Loc, bool Value) {
3741+
ExprResult Inner;
3742+
if (getLangOpts().CPlusPlus) {
3743+
Inner = ActOnCXXBoolLiteral(Loc, Value ? tok::kw_true : tok::kw_false);
3744+
} else {
3745+
// C doesn't actually have a way to represent literal values of type
3746+
// _Bool. So, we'll use 0/1 and implicit cast to _Bool.
3747+
Inner = ActOnIntegerConstant(Loc, Value ? 1 : 0);
3748+
Inner =
3749+
ImpCastExprToType(Inner.get(), Context.BoolTy, CK_IntegralToBoolean);
3750+
}
3751+
return Inner;
3752+
}
3753+
37403754
static Expr *BuildFloatingLiteral(Sema &S, NumericLiteralParser &Literal,
37413755
QualType Ty, SourceLocation Loc) {
37423756
const llvm::fltSemantics &Format = S.Context.getFloatTypeSemantics(Ty);

‎clang/lib/Sema/SemaExprObjC.cpp‎

Lines changed: 0 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -434,24 +434,6 @@ ExprResult SemaObjC::BuildObjCNumericLiteral(SourceLocation AtLoc,
434434
return SemaRef.MaybeBindToTemporary(NumberLiteral);
435435
}
436436

437-
ExprResult SemaObjC::ActOnObjCBoolLiteral(SourceLocation AtLoc,
438-
SourceLocation ValueLoc, bool Value) {
439-
ASTContext &Context = getASTContext();
440-
ExprResult Inner;
441-
if (getLangOpts().CPlusPlus) {
442-
Inner = SemaRef.ActOnCXXBoolLiteral(ValueLoc,
443-
Value ? tok::kw_true : tok::kw_false);
444-
} else {
445-
// C doesn't actually have a way to represent literal values of type
446-
// _Bool. So, we'll use 0/1 and implicit cast to _Bool.
447-
Inner = SemaRef.ActOnIntegerConstant(ValueLoc, Value ? 1 : 0);
448-
Inner = SemaRef.ImpCastExprToType(Inner.get(), Context.BoolTy,
449-
CK_IntegralToBoolean);
450-
}
451-
452-
return BuildObjCNumericLiteral(AtLoc, Inner.get());
453-
}
454-
455437
/// Check that the given expression is a valid element of an Objective-C
456438
/// collection literal.
457439
static ExprResult CheckObjCCollectionLiteralElement(Sema &S, Expr *Element,

‎clang/test/AST/ast-print-amdgcn-predicate.c‎

Lines changed: 88 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,88 @@
1+
// REQUIRES: amdgpu-registered-target
2+
3+
// Regression test for issue #199563: -ast-print on
4+
// __builtin_amdgcn_is_invocable / __builtin_amdgcn_processor_is used to
5+
// crash because the expansion produced an IntegerLiteral typed _Bool/bool.
6+
7+
// C (C11 and C23): IntegerLiteral 'int' implicit-cast to _Bool.
8+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -std=c11 \
9+
// RUN: -ast-print %s | FileCheck --check-prefix=INT-TRUE %s
10+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -std=c11 \
11+
// RUN: -ast-print %s | FileCheck --check-prefix=INT-FALSE %s
12+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -std=c11 \
13+
// RUN: -ast-dump %s | FileCheck --check-prefix=INT-DUMP %s
14+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -std=c23 \
15+
// RUN: -ast-print %s | FileCheck --check-prefix=INT-TRUE %s
16+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -std=c23 \
17+
// RUN: -ast-print %s | FileCheck --check-prefix=INT-FALSE %s
18+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -std=c23 \
19+
// RUN: -ast-dump %s | FileCheck --check-prefix=INT-DUMP %s
20+
21+
// C++ / HIP device: CXXBoolLiteralExpr 'bool'.
22+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x c++ \
23+
// RUN: -std=c++17 -ast-print %s | FileCheck --check-prefix=BOOL-TRUE %s
24+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -x c++ \
25+
// RUN: -std=c++17 -ast-print %s | FileCheck --check-prefix=BOOL-FALSE %s
26+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x c++ \
27+
// RUN: -std=c++17 -ast-dump %s | FileCheck --check-prefix=BOOL-DUMP %s
28+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x hip \
29+
// RUN: -fcuda-is-device -ast-print %s \
30+
// RUN: | FileCheck --check-prefix=BOOL-TRUE %s
31+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -x hip \
32+
// RUN: -fcuda-is-device -ast-print %s \
33+
// RUN: | FileCheck --check-prefix=BOOL-FALSE %s
34+
// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx942 -x hip \
35+
// RUN: -fcuda-is-device -ast-dump %s \
36+
// RUN: | FileCheck --check-prefix=BOOL-DUMP %s
37+
38+
#if defined(__HIP__) || defined(__CUDA__)
39+
#define __device__ __attribute__((device))
40+
#else
41+
#define __device__
42+
#endif
43+
44+
__device__ void use_is_invocable(void) {
45+
if (__builtin_amdgcn_is_invocable(__builtin_amdgcn_flat_atomic_fadd_f32))
46+
(void)0;
47+
}
48+
49+
__device__ void use_processor_is(void) {
50+
if (__builtin_amdgcn_processor_is("gfx942"))
51+
(void)0;
52+
}
53+
54+
// INT-TRUE-LABEL: use_is_invocable
55+
// INT-TRUE: if (1)
56+
// INT-TRUE-LABEL: use_processor_is
57+
// INT-TRUE: if (1)
58+
59+
// INT-FALSE-LABEL: use_is_invocable
60+
// INT-FALSE: if (0)
61+
// INT-FALSE-LABEL: use_processor_is
62+
// INT-FALSE: if (0)
63+
64+
// BOOL-TRUE-LABEL: use_is_invocable
65+
// BOOL-TRUE: if (true)
66+
// BOOL-TRUE-LABEL: use_processor_is
67+
// BOOL-TRUE: if (true)
68+
69+
// BOOL-FALSE-LABEL: use_is_invocable
70+
// BOOL-FALSE: if (false)
71+
// BOOL-FALSE-LABEL: use_processor_is
72+
// BOOL-FALSE: if (false)
73+
74+
// INT-DUMP-LABEL: FunctionDecl {{.*}} use_is_invocable
75+
// INT-DUMP: IfStmt
76+
// INT-DUMP-NEXT: ImplicitCastExpr {{.*}} {{'_Bool'|'bool'}} <IntegralToBoolean>
77+
// INT-DUMP-NEXT: IntegerLiteral {{.*}} 'int' 1
78+
// INT-DUMP-LABEL: FunctionDecl {{.*}} use_processor_is
79+
// INT-DUMP: IfStmt
80+
// INT-DUMP-NEXT: ImplicitCastExpr {{.*}} {{'_Bool'|'bool'}} <IntegralToBoolean>
81+
// INT-DUMP-NEXT: IntegerLiteral {{.*}} 'int' 1
82+
83+
// BOOL-DUMP-LABEL: FunctionDecl {{.*}} use_is_invocable
84+
// BOOL-DUMP: IfStmt
85+
// BOOL-DUMP-NEXT: CXXBoolLiteralExpr {{.*}} 'bool' true
86+
// BOOL-DUMP-LABEL: FunctionDecl {{.*}} use_processor_is
87+
// BOOL-DUMP: IfStmt
88+
// BOOL-DUMP-NEXT: CXXBoolLiteralExpr {{.*}} 'bool' true

0 commit comments

Comments
 (0)