diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp --- a/clang/lib/AST/ExprConstant.cpp +++ b/clang/lib/AST/ExprConstant.cpp @@ -3293,11 +3293,6 @@ return true; } -static bool IsConstNonVolatile(QualType T) { - Qualifiers Quals = T.getQualifiers(); - return Quals.hasConst() && !Quals.hasVolatile(); -} - /// Get the base index of the given base class within an APValue representing /// the given derived class. static unsigned getBaseIndex(const CXXRecordDecl *Derived, @@ -4015,8 +4010,6 @@ } 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); @@ -14822,7 +14815,6 @@ return IsGlobalLValue(Val.getLValueBase()); } - /// isIntegerConstantExpr - this recursive routine will test if an expression is /// an integer constant expression. @@ -14877,6 +14869,37 @@ return NoDiag(); } +/// Determine if the value of the given declaration is readable in an ICE. +static bool isReadableInICE(const ASTContext &Ctx, const NamedDecl *D) { + if (isa(D)) + return true; + + // 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. + if (!Ctx.getLangOpts().CPlusPlus && !Ctx.getLangOpts().OpenCL) + return false; + + const VarDecl *VD = dyn_cast(D); + if (!VD || isa(VD)) + return false; + + QualType T = VD->getType(); + if (T.isVolatileQualified() || !T->isIntegralOrEnumerationType()) + return false; + if (!T.isConstQualified() && + !(Ctx.getLangOpts().OpenCL && + T.getAddressSpace() == LangAS::opencl_constant)) + return false; + + // Look for a declaration of this variable that has an initializer, and + // check whether it is an ICE. + return VD->getAnyInitializer(VD) && !VD->isWeak() && VD->checkInitIsICE(); +} + static ICEDiag CheckICE(const Expr* E, const ASTContext &Ctx) { assert(!E->isValueDependent() && "Should not see value dependent exprs!"); if (!E->getType()->isIntegralOrEnumerationType()) @@ -15025,33 +15048,8 @@ return CheckICE(cast(E)->getSemanticForm(), Ctx); case Expr::DeclRefExprClass: { - if (isa(cast(E)->getDecl())) + if (isReadableInICE(Ctx, cast(E)->getDecl())) return NoDiag(); - const ValueDecl *D = cast(E)->getDecl(); - if (Ctx.getLangOpts().CPlusPlus && - D && IsConstNonVolatile(D->getType())) { - // Parameter variables are never constants. Without this check, - // getAnyInitializer() can find a default argument, which leads - // to chaos. - if (isa(D)) - return ICEDiag(IK_NotICE, cast(E)->getLocation()); - - // 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. - if (const VarDecl *Dcl = dyn_cast(D)) { - if (!Dcl->getType()->isIntegralOrEnumerationType()) - return ICEDiag(IK_NotICE, cast(E)->getLocation()); - - const VarDecl *VD; - // Look for a declaration of this variable that has an initializer, and - // check whether it is an ICE. - if (Dcl->getAnyInitializer(VD) && !VD->isWeak() && VD->checkInitIsICE()) - return NoDiag(); - else - return ICEDiag(IK_NotICE, cast(E)->getLocation()); - } - } return ICEDiag(IK_NotICE, E->getBeginLoc()); } case Expr::UnaryOperatorClass: { diff --git a/clang/lib/Sema/SemaType.cpp b/clang/lib/Sema/SemaType.cpp --- a/clang/lib/Sema/SemaType.cpp +++ b/clang/lib/Sema/SemaType.cpp @@ -2278,8 +2278,7 @@ // TryFixVariablyModifiedType to convert VLAs to constant array types. ExprResult R = S.VerifyIntegerConstantExpression( ArraySize, &SizeVal, Diagnoser, - (S.LangOpts.GNUMode || S.LangOpts.OpenCL) ? Sema::AllowFold - : Sema::NoFold); + S.LangOpts.GNUMode ? Sema::AllowFold : Sema::NoFold); if (Diagnoser.IsVLA) return ExprResult(); return R; diff --git a/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl b/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl --- a/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl +++ b/clang/test/CodeGenOpenCL/amdgpu-nullptr.cl @@ -104,7 +104,7 @@ // 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 @@ // 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 @@ // 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 @@ // 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 --- 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}}