rsmith updated this revision to Diff 298764.
rsmith added a comment.
Rebase on D89523 <https://reviews.llvm.org/D89523>.
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D89520/new/
https://reviews.llvm.org/D89520
Files:
clang/lib/AST/ExprConstant.cpp
clang/lib/Sema/SemaType.cpp
clang/test/CodeGenOpenCL/amdgpu-nullptr.cl
clang/test/SemaOpenCL/address-spaces.cl
Index: clang/test/SemaOpenCL/address-spaces.cl
===================================================================
--- clang/test/SemaOpenCL/address-spaces.cl
+++ 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}}
Index: clang/test/CodeGenOpenCL/amdgpu-nullptr.cl
===================================================================
--- clang/test/CodeGenOpenCL/amdgpu-nullptr.cl
+++ 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)*
Index: clang/lib/Sema/SemaType.cpp
===================================================================
--- clang/lib/Sema/SemaType.cpp
+++ clang/lib/Sema/SemaType.cpp
@@ -2273,9 +2273,8 @@
}
} 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;
Index: clang/lib/AST/ExprConstant.cpp
===================================================================
--- clang/lib/AST/ExprConstant.cpp
+++ clang/lib/AST/ExprConstant.cpp
@@ -3298,11 +3298,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,
@@ -4020,8 +4015,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);
@@ -14825,7 +14818,6 @@
return IsGlobalLValue(Val.getLValueBase());
}
-
/// isIntegerConstantExpr - this recursive routine will test if an expression is
/// an integer constant expression.
@@ -14880,6 +14872,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<EnumConstantDecl>(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<VarDecl>(D);
+ if (!VD || isa<ParmVarDecl>(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())
@@ -15028,33 +15051,8 @@
return CheckICE(cast<CXXRewrittenBinaryOperator>(E)->getSemanticForm(),
Ctx);
case Expr::DeclRefExprClass: {
- if (isa<EnumConstantDecl>(cast<DeclRefExpr>(E)->getDecl()))
+ if (isReadableInICE(Ctx, cast<DeclRefExpr>(E)->getDecl()))
return NoDiag();
- const ValueDecl *D = cast<DeclRefExpr>(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<ParmVarDecl>(D))
- return ICEDiag(IK_NotICE, cast<DeclRefExpr>(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<VarDecl>(D)) {
- if (!Dcl->getType()->isIntegralOrEnumerationType())
- return ICEDiag(IK_NotICE, cast<DeclRefExpr>(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<DeclRefExpr>(E)->getLocation());
- }
- }
return ICEDiag(IK_NotICE, E->getBeginLoc());
}
case Expr::UnaryOperatorClass: {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
