[clang] 6781fee - Don't permit array bound constant folding in OpenCL.
Richard Smith via cfe-commits
cfe-commits at lists.llvm.org
Tue Oct 20 16:52:37 PDT 2020
Author: Richard Smith
Date: 2020-10-20T16:52:28-07:00
New Revision: 6781fee085058913444e0c5937da9c0e7e928db5
URL: https://github.com/llvm/llvm-project/commit/6781fee085058913444e0c5937da9c0e7e928db5
DIFF: https://github.com/llvm/llvm-project/commit/6781fee085058913444e0c5937da9c0e7e928db5.diff
LOG: Don't permit array bound constant folding in OpenCL.
Permitting non-standards-driven "do the best you can" constant-folding
of array bounds is permitted solely as a GNU compatibility feature. We
should not be doing it in any language mode that is attempting to be
conforming.
>From https://reviews.llvm.org/D20090 it appears the intent here was to
permit `__constant int` globals to be used in array bounds, but the
change in that patch only added half of the functionality necessary to
support that in the constant evaluator. This patch adds the other half
of the functionality and turns off constant folding for array bounds in
OpenCL.
I couldn't find any spec justification for accepting the kinds of cases
that D20090 accepts, so a reference to where in the OpenCL specification
this is permitted would be useful.
Note that this change also affects the code generation in one test:
because after 'const int n = 0' we now treat 'n' as a constant
expression with value 0, it's now a null pointer, so '(local int *)n'
forms a null pointer rather than a zero pointer.
Reviewed By: Anastasia
Differential Revision: https://reviews.llvm.org/D89520
Added:
Modified:
clang/lib/AST/Decl.cpp
clang/lib/AST/ExprConstant.cpp
clang/lib/Sema/SemaType.cpp
clang/test/CodeGenOpenCL/amdgpu-nullptr.cl
clang/test/SemaOpenCL/address-spaces.cl
Removed:
################################################################################
diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp
index b3a67715d604..80bc22b61b70 100644
--- a/clang/lib/AST/Decl.cpp
+++ b/clang/lib/AST/Decl.cpp
@@ -2280,7 +2280,9 @@ void VarDecl::setInit(Expr *I) {
bool VarDecl::mightBeUsableInConstantExpressions(const ASTContext &C) const {
const LangOptions &Lang = C.getLangOpts();
- if (!Lang.CPlusPlus)
+ // OpenCL permits const integral variables to be used in constant
+ // expressions, like in C++98.
+ if (!Lang.CPlusPlus && !Lang.OpenCL)
return false;
// Function parameters are never usable in constant expressions.
@@ -2299,7 +2301,7 @@ bool VarDecl::mightBeUsableInConstantExpressions(const ASTContext &C) const {
// Only const objects can be used in constant expressions in C++. C++98 does
// not require the variable to be non-volatile, but we consider this to be a
// defect.
- if (!getType().isConstQualified() || getType().isVolatileQualified())
+ if (!getType().isConstant(C) || getType().isVolatileQualified())
return false;
// In C++, const, non-volatile variables of integral or enumeration types
@@ -2325,14 +2327,14 @@ bool VarDecl::isUsableInConstantExpressions(const ASTContext &Context) const {
if (!DefVD->mightBeUsableInConstantExpressions(Context))
return false;
// ... and its initializer is a constant initializer.
- if (!DefVD->hasConstantInitialization())
+ if (Context.getLangOpts().CPlusPlus && !DefVD->hasConstantInitialization())
return false;
// C++98 [expr.const]p1:
// An integral constant-expression can involve only [...] const variables
// or static data members of integral or enumeration types initialized with
// [integer] constant expressions (dcl.init)
- if (Context.getLangOpts().CPlusPlus && !Context.getLangOpts().CPlusPlus11 &&
- !DefVD->hasICEInitializer(Context))
+ if ((Context.getLangOpts().CPlusPlus || Context.getLangOpts().OpenCL) &&
+ !Context.getLangOpts().CPlusPlus11 && !DefVD->hasICEInitializer(Context))
return false;
return true;
}
diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp
index f4664bb62414..b2ddca7f7ce0 100644
--- a/clang/lib/AST/ExprConstant.cpp
+++ b/clang/lib/AST/ExprConstant.cpp
@@ -3284,10 +3284,10 @@ static bool evaluateVarDeclInit(EvalInfo &Info, const Expr *E,
// FIXME: We don't diagnose cases that aren't potentially usable in constant
// expressions here; doing so would regress diagnostics for things like
// reading from a volatile constexpr variable.
- if ((!VD->hasConstantInitialization() &&
+ if ((Info.getLangOpts().CPlusPlus && !VD->hasConstantInitialization() &&
VD->mightBeUsableInConstantExpressions(Info.Ctx)) ||
- (Info.getLangOpts().CPlusPlus && !Info.getLangOpts().CPlusPlus11 &&
- !VD->hasICEInitializer(Info.Ctx))) {
+ ((Info.getLangOpts().CPlusPlus || Info.getLangOpts().OpenCL) &&
+ !Info.getLangOpts().CPlusPlus11 && !VD->hasICEInitializer(Info.Ctx))) {
Info.CCEDiag(E, diag::note_constexpr_var_init_non_constant, 1) << VD;
NoteLValueLocation(Info, Base);
}
@@ -3997,10 +3997,7 @@ static CompleteObject findCompleteObject(EvalInfo &Info, const Expr *E,
return CompleteObject();
}
- // In OpenCL if a variable is in constant address space it is a const value.
- bool IsConstant = BaseType.isConstQualified() ||
- (Info.getLangOpts().OpenCL &&
- BaseType.getAddressSpace() == LangAS::opencl_constant);
+ bool IsConstant = BaseType.isConstant(Info.Ctx);
// Unless we're looking at a local variable or argument in a constexpr call,
// the variable we're reading must be const.
@@ -4021,8 +4018,6 @@ static CompleteObject findCompleteObject(EvalInfo &Info, const Expr *E,
} else if (VD->isConstexpr()) {
// OK, we can read this variable.
} else if (BaseType->isIntegralOrEnumerationType()) {
- // In OpenCL if a variable is in constant address space it is a const
- // value.
if (!IsConstant) {
if (!IsAccess)
return CompleteObject(LVal.getLValueBase(), nullptr, BaseType);
@@ -14834,7 +14829,6 @@ bool Expr::EvalResult::isGlobalLValue() const {
return IsGlobalLValue(Val.getLValueBase());
}
-
/// isIntegerConstantExpr - this recursive routine will test if an expression is
/// an integer constant expression.
@@ -15037,15 +15031,20 @@ static ICEDiag CheckICE(const Expr* E, const ASTContext &Ctx) {
return CheckICE(cast<CXXRewrittenBinaryOperator>(E)->getSemanticForm(),
Ctx);
case Expr::DeclRefExprClass: {
- if (isa<EnumConstantDecl>(cast<DeclRefExpr>(E)->getDecl()))
+ const NamedDecl *D = cast<DeclRefExpr>(E)->getDecl();
+ if (isa<EnumConstantDecl>(D))
return NoDiag();
- const VarDecl *VD = dyn_cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
- if (VD && VD->isUsableInConstantExpressions(Ctx)) {
- // C++ 7.1.5.1p2
- // A variable of non-volatile const-qualified integral or enumeration
- // type initialized by an ICE can be used in ICEs.
+
+ // C++ and OpenCL (FIXME: spec reference?) allow reading const-qualified
+ // integer variables in constant expressions:
+ //
+ // C++ 7.1.5.1p2
+ // A variable of non-volatile const-qualified integral or enumeration
+ // type initialized by an ICE can be used in ICEs.
+ const VarDecl *VD = dyn_cast<VarDecl>(D);
+ if (VD && VD->isUsableInConstantExpressions(Ctx))
return NoDiag();
- }
+
return ICEDiag(IK_NotICE, E->getBeginLoc());
}
case Expr::UnaryOperatorClass: {
diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp
index a862b4e955d6..136f03a69b25 100644
--- a/clang/lib/Sema/SemaType.cpp
+++ b/clang/lib/Sema/SemaType.cpp
@@ -2273,9 +2273,8 @@ static ExprResult checkArraySize(Sema &S, Expr *&ArraySize,
}
} Diagnoser(VLADiag, VLAIsError);
- ExprResult R = S.VerifyIntegerConstantExpression(
- ArraySize, &SizeVal, Diagnoser,
- S.LangOpts.OpenCL ? Sema::AllowFold : Sema::NoFold);
+ ExprResult R =
+ S.VerifyIntegerConstantExpression(ArraySize, &SizeVal, Diagnoser);
if (Diagnoser.IsVLA)
return ExprResult();
return R;
diff --git a/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl b/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl
index ddf358c1188e..5f0a1d063fab 100644
--- a/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl
+++ b/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl
@@ -104,7 +104,7 @@ int fold_int5_local = (int) &((local StructTy1*)0)->p2;
// NOOPT: @test_static_var_private.sp2 = internal addrspace(1) global i8 addrspace(5)* addrspacecast (i8* null to i8 addrspace(5)*), align 4
// NOOPT: @test_static_var_private.sp3 = internal addrspace(1) global i8 addrspace(5)* addrspacecast (i8* null to i8 addrspace(5)*), align 4
// NOOPT: @test_static_var_private.sp4 = internal addrspace(1) global i8 addrspace(5)* null, align 4
-// NOOPT: @test_static_var_private.sp5 = internal addrspace(1) global i8 addrspace(5)* null, align 4
+// NOOPT: @test_static_var_private.sp5 = internal addrspace(1) global i8 addrspace(5)* addrspacecast (i8* null to i8 addrspace(5)*), align 4
// NOOPT: @test_static_var_private.SS1 = internal addrspace(1) global %struct.StructTy1 { i8 addrspace(5)* addrspacecast (i8* null to i8 addrspace(5)*), i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(4)* null, i8 addrspace(1)* null, i8* null }, align 8
// NOOPT: @test_static_var_private.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 8
@@ -123,7 +123,7 @@ void test_static_var_private(void) {
// NOOPT: @test_static_var_local.sp2 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), align 4
// NOOPT: @test_static_var_local.sp3 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), align 4
// NOOPT: @test_static_var_local.sp4 = internal addrspace(1) global i8 addrspace(3)* null, align 4
-// NOOPT: @test_static_var_local.sp5 = internal addrspace(1) global i8 addrspace(3)* null, align 4
+// NOOPT: @test_static_var_local.sp5 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), align 4
// NOOPT: @test_static_var_local.SS1 = internal addrspace(1) global %struct.StructTy1 { i8 addrspace(5)* addrspacecast (i8* null to i8 addrspace(5)*), i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(4)* null, i8 addrspace(1)* null, i8* null }, align 8
// NOOPT: @test_static_var_local.SS2 = internal addrspace(1) global %struct.StructTy2 zeroinitializer, align 8
void test_static_var_local(void) {
@@ -142,7 +142,7 @@ void test_static_var_local(void) {
// NOOPT: store i8 addrspace(5)* addrspacecast (i8* null to i8 addrspace(5)*), i8 addrspace(5)* addrspace(5)* %sp1, align 4
// NOOPT: store i8 addrspace(5)* addrspacecast (i8* null to i8 addrspace(5)*), i8 addrspace(5)* addrspace(5)* %sp2, align 4
// NOOPT: store i8 addrspace(5)* null, i8 addrspace(5)* addrspace(5)* %sp3, align 4
-// NOOPT: store i8 addrspace(5)* null, i8 addrspace(5)* addrspace(5)* %sp4, align 4
+// NOOPT: store i8 addrspace(5)* addrspacecast (i8* null to i8 addrspace(5)*), i8 addrspace(5)* addrspace(5)* %sp4, align 4
// NOOPT: %[[SS1:.*]] = bitcast %struct.StructTy1 addrspace(5)* %SS1 to i8 addrspace(5)*
// NOOPT: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 8 %[[SS1]], i8 addrspace(4)* align 8 bitcast (%struct.StructTy1 addrspace(4)* @__const.test_func_scope_var_private.SS1 to i8 addrspace(4)*), i64 32, i1 false)
// NOOPT: %[[SS2:.*]] = bitcast %struct.StructTy2 addrspace(5)* %SS2 to i8 addrspace(5)*
@@ -162,7 +162,7 @@ void test_func_scope_var_private(void) {
// NOOPT: store i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(3)* addrspace(5)* %sp1, align 4
// NOOPT: store i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(3)* addrspace(5)* %sp2, align 4
// NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)* addrspace(5)* %sp3, align 4
-// NOOPT: store i8 addrspace(3)* null, i8 addrspace(3)* addrspace(5)* %sp4, align 4
+// NOOPT: store i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), i8 addrspace(3)* addrspace(5)* %sp4, align 4
// NOOPT: %[[SS1:.*]] = bitcast %struct.StructTy1 addrspace(5)* %SS1 to i8 addrspace(5)*
// NOOPT: call void @llvm.memcpy.p5i8.p4i8.i64(i8 addrspace(5)* align 8 %[[SS1]], i8 addrspace(4)* align 8 bitcast (%struct.StructTy1 addrspace(4)* @__const.test_func_scope_var_local.SS1 to i8 addrspace(4)*), i64 32, i1 false)
// NOOPT: %[[SS2:.*]] = bitcast %struct.StructTy2 addrspace(5)* %SS2 to i8 addrspace(5)*
diff --git a/clang/test/SemaOpenCL/address-spaces.cl b/clang/test/SemaOpenCL/address-spaces.cl
index e9825fd1682a..2c4f0fdec524 100644
--- a/clang/test/SemaOpenCL/address-spaces.cl
+++ b/clang/test/SemaOpenCL/address-spaces.cl
@@ -4,6 +4,13 @@
__constant int ci = 1;
+// __constant ints are allowed in constant expressions.
+enum use_ci_in_enum { enumerator = ci };
+typedef int use_ci_in_array_bound[ci];
+
+// general constant folding of array bounds is not permitted
+typedef int folding_in_array_bounds[&ci + 3 - &ci]; // expected-error-re {{{{variable length arrays are not supported in OpenCL|array size is not a constant expression}}}} expected-note {{cannot refer to element 3}}
+
__kernel void foo(__global int *gip) {
__local int li;
__local int lj = 2; // expected-error {{'__local' variable cannot have an initializer}}
More information about the cfe-commits
mailing list