This revision was landed with ongoing or failed builds.
This revision was automatically updated to reflect the committed changes.
Closed by commit rG6781fee08505: Don't permit array bound constant folding
in OpenCL. (authored by rsmith).
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D89520/new/
https://reviews.llvm.org/D89520
Files:
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
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
@@ -3284,10 +3284,10 @@
// 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 @@
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 @@
} 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 @@
return IsGlobalLValue(Val.getLValueBase());
}
-
/// isIntegerConstantExpr - this recursive routine will test if an expression is
/// an integer constant expression.
@@ -15037,15 +15031,20 @@
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: {
Index: clang/lib/AST/Decl.cpp
===================================================================
--- clang/lib/AST/Decl.cpp
+++ clang/lib/AST/Decl.cpp
@@ -2280,7 +2280,9 @@
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 @@
// 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 @@
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;
}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
