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}} _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
