rsmith created this revision.
rsmith added reviewers: Anastasia, bader, yaxunl.
Herald added subscribers: kerbowa, nhaehnle, jvesely.
Herald added a project: clang.
rsmith requested review of this revision.
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 <https://reviews.llvm.org/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.
Repository:
rG LLVM Github Monorepo
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
@@ -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;
Index: clang/lib/AST/ExprConstant.cpp
===================================================================
--- clang/lib/AST/ExprConstant.cpp
+++ 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<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())
@@ -15025,33 +15048,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
