Don't permit array bound constant folding in OpenCL.

[rsmith]

Summary

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.

[Anastasia]

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.

Thanks for fixing this. I agree that the original change was not compliant with the spec. OpenCL indeed doesn't allow constant folding for array bounds. The idea of the change was to allow using expressions that are compile time constant in the array bound because this doesn't result in VLA.

Regarding the spec reference, I think we can refer to the section 6.5.3 describing variables in the __constant address space:

These variables  are  required  to  be  initialized  and  the  values  used  to  initialize  these  variables  must  be  a compile time constant. Writing to such a variable results in a compile-time error.

I.e. the __constant address space variables are semantically similar to constexpr in C++.

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.

I am slightly confused about this case because technically the value of 'n' can be overwritten by casting away const. However, I suspect this is compliant C99 behavior?

This example:

void test(void) {
  const int x = 0;
  const int* ptrx = &x;
  int* ptry = (int*)ptrx;
  *ptry = 100; 
  int *sp5 = ( int*)x;
}

shows that the IR produced is indeed supposed to have NULL despite of the fact that the value of initializing variable is not 0 at the time of the last initialization.

store i32* null, i32** %4, align 8

[rsmith]

In D89520#2334477, @Anastasia wrote:

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.

Thanks for fixing this. I agree that the original change was not compliant with the spec. OpenCL indeed doesn't allow constant folding for array bounds. The idea of the change was to allow using expressions that are compile time constant in the array bound because this doesn't result in VLA.

Regarding the spec reference, I think we can refer to the section 6.5.3 describing variables in the __constant address space:

These variables  are  required  to  be  initialized  and  the  values  used  to  initialize  these  variables  must  be  a compile time constant. Writing to such a variable results in a compile-time error.

I.e. the __constant address space variables are semantically similar to constexpr in C++.

I don't see any way to read this wording that way. That talks about how the variables are initialized and says they're immutable, but I can't find anything in the OpenCL specification that says they can be used in integer constant expressions. The C definition of "integral constant expression" does not allow the use of variables:

C11 6.6/6: "An integer constant expression shall have integer type and shall only have operands that are integer constants, enumeration constants, character constants, sizeof expressions whose results are integer constants, _Alignof expressions, and floating constants that are the immediate operands of casts."

... and absent a modification to that, nor would OpenCL. That said, I'm happy to take your word for it that the OpenCL specification intends for such variable uses to be permitted in constant expressions.

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.

I am slightly confused about this case because technically the value of 'n' can be overwritten by casting away const.

That would result in undefined behavior. And if not, the same consideration would apply to array bounds :)

However, I suspect this is compliant C99 behavior?

Not exactly, because C99 doesn't permit usage of variables in integer constant expressions; this is a consequence of the presumptive OpenCL rule that does permit such things. In C, the closest analogue would be something like:

enum { e = 0 };
int *p = (int*)0; // null pointer constant, not a bit-cast of zero to a pointer

This example:

void test(void) {
  const int x = 0;
  const int* ptrx = &x;
  int* ptry = (int*)ptrx;
  *ptry = 100; 
  int *sp5 = ( int*)x;
}

shows that the IR produced is indeed supposed to have NULL despite of the fact that the value of initializing variable is not 0 at the time of the last initialization.

This example has undefined behavior due to writing to a const variable. Also, in C it initializes sp5 to a zero pointer, not necessarily a null pointer, because x is not an integer constant expression -- though you can't see the difference here because we're in an address space where they're the same. Under the presumptive OpenCL rule, x would be an integer constant expression, so the above example would initialize sp5 to a null pointer instead of a zero pointer, just like in C++98.

[yaxunl]

I am OK with the changes regarding null pointer. I guess people seldom set pointer to zero address in OpenCL.

[rsmith]

Rebase on D89523.

[Anastasia]

In D89520#2335628, @rsmith wrote:

In D89520#2334477, @Anastasia wrote:

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.

Thanks for fixing this. I agree that the original change was not compliant with the spec. OpenCL indeed doesn't allow constant folding for array bounds. The idea of the change was to allow using expressions that are compile time constant in the array bound because this doesn't result in VLA.

Regarding the spec reference, I think we can refer to the section 6.5.3 describing variables in the __constant address space:

These variables  are  required  to  be  initialized  and  the  values  used  to  initialize  these  variables  must  be  a compile time constant. Writing to such a variable results in a compile-time error.

I.e. the __constant address space variables are semantically similar to constexpr in C++.

I don't see any way to read this wording that way. That talks about how the variables are initialized and says they're immutable, but I can't find anything in the OpenCL specification that says they can be used in integer constant expressions. The C definition of "integral constant expression" does not allow the use of variables:

C11 6.6/6: "An integer constant expression shall have integer type and shall only have operands that are integer constants, enumeration constants, character constants, sizeof expressions whose results are integer constants, _Alignof expressions, and floating constants that are the immediate operands of casts."

... and absent a modification to that, nor would OpenCL. That said, I'm happy to take your word for it that the OpenCL specification intends for such variable uses to be permitted in constant expressions.

Yes, I agree it is not explicit. I will attempt to improve it. However considering that your change doesn't add this feature but modifies existing implementation, I think it makes sense to go ahead.

[Anastasia]

LGTM! Thanks

[rsmith]

Rebase.

clang/lib/AST/ExprConstant.cpp

if (VD->isWeak()) {
     NoteLValueLocation(Info, Base);
     return false;
   }
 
   Result = VD->getEvaluatedValue();
   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,
                              const CXXRecordDecl *Base) {
   Base = Base->getCanonicalDecl();
   unsigned Index = 0;
   for (CXXRecordDecl::base_class_const_iterator I = Derived->bases_begin(),
          E = Derived->bases_end(); I != E; ++I, ++Index) {
if (!Frame) {
         // evaluation.
       } else if (isModification(AK)) {
         // All the remaining cases do not permit modification of the object.
         Info.FFDiag(E, diag::note_constexpr_modify_global);
         return CompleteObject();
       } 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);
           if (Info.getLangOpts().CPlusPlus) {
             Info.FFDiag(E, diag::note_constexpr_ltor_non_const_int, 1) << VD;
             Info.Note(VD->getLocation(), diag::note_declared_at);
           } else {
             Info.FFDiag(E);
void Expr::EvaluateForOverflow(const ASTContext &Ctx) const {
   }
 }
 
 bool Expr::EvalResult::isGlobalLValue() const {
   assert(Val.isLValue());
   return IsGlobalLValue(Val.getLValueBase());
 }
 
-
 /// isIntegerConstantExpr - this recursive routine will test if an expression is
 /// an integer constant expression.
 
 /// FIXME: Pass up a reason why! Invalid operation in i-c-e, division by zero,
 /// comma, etc
 
 // CheckICE - This function does the fundamental ICE checking: the returned
 // ICEDiag contains an ICEKind indicating whether the expression is an ICE,
static ICEDiag CheckEvalInICE(const Expr* E, const ASTContext &Ctx) {
   Info.InConstantContext = true;
   if (!::EvaluateAsRValue(E, EVResult, Ctx, Info) || EVResult.HasSideEffects ||
       !EVResult.Val.isInt())
     return ICEDiag(IK_NotICE, E->getBeginLoc());
 
   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())
     return ICEDiag(IK_NotICE, E->getBeginLoc());
 
   switch (E->getStmtClass()) {
 #define ABSTRACT_STMT(Node)
 #define STMT(Node, Base) case Expr::Node##Class:
case Expr::CXXOperatorCallExprClass: {
     if (CE->getBuiltinCallee())
       return CheckEvalInICE(E, Ctx);
     return ICEDiag(IK_NotICE, E->getBeginLoc());
   }
   case Expr::CXXRewrittenBinaryOperatorClass:
     return CheckICE(cast<CXXRewrittenBinaryOperator>(E)->getSemanticForm(),
                     Ctx);
   case Expr::DeclRefExprClass: {
-    if (isa<EnumConstantDecl>(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());
-      }
-    }
+    if (isReadableInICE(Ctx, cast<DeclRefExpr>(E)->getDecl()))
     return ICEDiag(IK_NotICE, E->getBeginLoc());
   }
   case Expr::UnaryOperatorClass: {
     const UnaryOperator *Exp = cast<UnaryOperator>(E);
     switch (Exp->getOpcode()) {
     case UO_PostInc:
     case UO_PostDec:
     case UO_PreInc:

clang/lib/Sema/SemaType.cpp

public:
     }
 
     Sema::SemaDiagnosticBuilder diagnoseFold(Sema &S,
                                              SourceLocation Loc) override {
       return S.Diag(Loc, diag::ext_vla_folded_to_constant);
     }
   } 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;
 }
 
 /// Build an array type.
 ///
 /// \param T The type of each element in the array.

clang/test/CodeGenOpenCL/amdgpu-nullptr.cl

 
 
 // Test static variable initialization.
 
 // NOOPT: @test_static_var_private.sp1 = internal addrspace(1) global i8 addrspace(5)* addrspacecast (i8* null to i8 addrspace(5)*), align 4
 // 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
 
 void test_static_var_private(void) {
   static private char *sp1 = 0;
   static private char *sp2 = NULL;
   static private char *sp3;
   static private char *sp4 = (private char*)((void)0, 0);
   const int x = 0;
   static private char *sp5 = (private char*)x;
   static StructTy1 SS1;
   static StructTy2 SS2;
 }
 
 // NOOPT: @test_static_var_local.sp1 = internal addrspace(1) global i8 addrspace(3)* addrspacecast (i8* null to i8 addrspace(3)*), align 4
 // 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) {
   static local char *sp1 = 0;
   static local char *sp2 = NULL;
   static local char *sp3;
   static local char *sp4 = (local char*)((void)0, 0);
   const int x = 0;
   static local char *sp5 = (local char*)x;
   static StructTy1 SS1;
   static StructTy2 SS2;
 }
 
 // Test function-scope variable initialization.
 // NOOPT-LABEL: @test_func_scope_var_private(
 // 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)*
 // NOOPT: call void @llvm.memset.p5i8.i64(i8 addrspace(5)* align 8 %[[SS2]], i8 0, i64 24, i1 false)
 void test_func_scope_var_private(void) {
   private char *sp1 = 0;
   private char *sp2 = NULL;
   private char *sp3 = (private char*)((void)0, 0);
   const int x = 0;
   private char *sp4 = (private char*)x;
   StructTy1 SS1 = {0, 0, 0, 0, 0};
   StructTy2 SS2 = {0, 0, 0};
 }
 
 // Test function-scope variable initialization.
 // NOOPT-LABEL: @test_func_scope_var_local(
 // 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)*
 // NOOPT: call void @llvm.memset.p5i8.i64(i8 addrspace(5)* align 8 %[[SS2]], i8 0, i64 24, i1 false)
 void test_func_scope_var_local(void) {
   local char *sp1 = 0;
   local char *sp2 = NULL;
   local char *sp3 = (local char*)((void)0, 0);

clang/test/SemaOpenCL/address-spaces.cl

 // RUN: %clang_cc1 %s -verify -pedantic -fsyntax-only
 // RUN: %clang_cc1 %s -cl-std=CL2.0 -verify -pedantic -fsyntax-only
 // RUN: %clang_cc1 %s -cl-std=clc++ -verify -pedantic -fsyntax-only
 
 __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}}
 
   int *ip;
 #if ((!__OPENCL_CPP_VERSION__) && (__OPENCL_C_VERSION__ < 200))
   ip = gip; // expected-error {{assigning '__global int *__private' to '__private int *__private' changes address space of pointer}}
   ip = &li; // expected-error {{assigning '__local int *' to '__private int *__private' changes address space of pointer}}