[OpenCL] Disallow negative attribute arguments

[asavonic]

Summary

Negative arguments in kernel attributes are silently bitcast'ed to
unsigned, for example:

__attribute__((reqd_work_group_size(1, -1, 1)))
__kernel void k() {}

is a complete equivalent of:

__attribute__((reqd_work_group_size(1, 4294967294, 1)))
__kernel void k() {}

This is likely an error, so the patch forbids negative arguments in
several OpenCL attributes. Users who really want 4294967294 can still
use it as an unsigned representation.

[Anastasia]

LGTM!

[bader]

LGTM.
Please, remove "Change-Id: I910b5c077f5f29e02a1572d9202f0fdbea5280fd" from the log message - it not relevant to the project.

[yaxunl]

LGTM. Thanks.

include/clang/Basic/DiagnosticSemaKinds.td

 def err_attribute_argument_n_type : Error<
   "%0 attribute requires parameter %1 to be %select{int or bool|an integer "
   "constant|a string|an identifier}2">;
 def err_attribute_argument_type : Error<
   "%0 attribute requires %select{int or bool|an integer "
   "constant|a string|an identifier}1">;
 def err_attribute_argument_outof_range : Error<
   "%0 attribute requires integer constant between %1 and %2 inclusive">;
+def err_attribute_argument_negative : Error<
+  "negative argument is not allowed for %0 attribute">;
 def err_init_priority_object_attr : Error<
   "can only use 'init_priority' attribute on file-scope definitions "
   "of objects of class type">;
 def err_attribute_argument_vec_type_hint : Error<
   "invalid attribute argument %0 - expecting a vector or vectorizable scalar type">;
 def err_attribute_argument_out_of_bounds : Error<
   "%0 attribute parameter %1 is out of bounds">;
 def err_attribute_only_once_per_parameter : Error<

lib/Sema/SemaDeclAttr.cpp

getAttrName(const AttrInfo &AL) {
   return &AL;
 }
 static const IdentifierInfo *getAttrName(const ParsedAttr &AL) {
   return AL.getName();
 }
 
 /// If Expr is a valid integer constant, get the value of the integer
 /// expression and return success or failure. May output an error.
+///
+/// Negative argument is implicitly converted to unsigned, unless
+/// \p StrictlyUnsigned is true.
 template <typename AttrInfo>
 static bool checkUInt32Argument(Sema &S, const AttrInfo &AI, const Expr *Expr,
-                                uint32_t &Val, unsigned Idx = UINT_MAX) {
+                                uint32_t &Val, unsigned Idx = UINT_MAX,
+                                bool StrictlyUnsigned = false) {
   llvm::APSInt I(32);
   if (Expr->isTypeDependent() || Expr->isValueDependent() ||
       !Expr->isIntegerConstantExpr(I, S.Context)) {
     if (Idx != UINT_MAX)
       S.Diag(getAttrLoc(AI), diag::err_attribute_argument_n_type)
         << getAttrName(AI) << Idx << AANT_ArgumentIntegerConstant
         << Expr->getSourceRange();
     else
       S.Diag(getAttrLoc(AI), diag::err_attribute_argument_type)
         << getAttrName(AI) << AANT_ArgumentIntegerConstant
         << Expr->getSourceRange();
     return false;
   }
 
   if (!I.isIntN(32)) {
     S.Diag(Expr->getExprLoc(), diag::err_ice_too_large)
         << I.toString(10, false) << 32 << /* Unsigned */ 1;
     return false;
   }
 
+  if (StrictlyUnsigned && I.isSigned() && I.isNegative()) {
+    S.Diag(getAttrLoc(AI), diag::err_attribute_argument_negative)
+        << getAttrName(AI);
+    return false;
+  }
+
   Val = (uint32_t)I.getZExtValue();
   return true;
 }
 
 /// Wrapper around checkUInt32Argument, with an extra check to be sure
 /// that the result will fit into a regular (signed) int. All args have the same
 /// purpose as they do in checkUInt32Argument.
 template <typename AttrInfo>

 }
 
 // Handles reqd_work_group_size and work_group_size_hint.
 template <typename WorkGroupAttr>
 static void handleWorkGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) {
   uint32_t WGSize[3];
   for (unsigned i = 0; i < 3; ++i) {
     const Expr *E = AL.getArgAsExpr(i);
-    if (!checkUInt32Argument(S, AL, E, WGSize[i], i))
+    if (!checkUInt32Argument(S, AL, E, WGSize[i], i,
+                             /*StrictlyUnsigned=*/true))
       return;
     if (WGSize[i] == 0) {
       S.Diag(AL.getLoc(), diag::err_attribute_argument_is_zero)
         << AL.getName() << E->getSourceRange();
       return;
     }
   }
 

test/SemaOpenCL/invalid-kernel-attrs.cl

 
 kernel __attribute__((reqd_work_group_size(1,2,0))) void kernel11(){} // expected-error {{'reqd_work_group_size' attribute must be greater than 0}}
 kernel __attribute__((reqd_work_group_size(1,0,2))) void kernel12(){} // expected-error {{'reqd_work_group_size' attribute must be greater than 0}}
 kernel __attribute__((reqd_work_group_size(0,1,2))) void kernel13(){} // expected-error {{'reqd_work_group_size' attribute must be greater than 0}}
 
 __attribute__((intel_reqd_sub_group_size(8))) void kernel14(){} // expected-error {{attribute 'intel_reqd_sub_group_size' can only be applied to an OpenCL kernel}}
 kernel __attribute__((intel_reqd_sub_group_size(0))) void kernel15(){} // expected-error {{'intel_reqd_sub_group_size' attribute must be greater than 0}}
 kernel __attribute__((intel_reqd_sub_group_size(8))) __attribute__((intel_reqd_sub_group_size(16))) void kernel16() {}  //expected-warning{{attribute 'intel_reqd_sub_group_size' is already applied with different parameters}}
+
+__kernel __attribute__((work_group_size_hint(8,-16,32))) void neg1() {} //expected-error{{negative argument is not allowed for 'work_group_size_hint' attribute}}
+__kernel __attribute__((reqd_work_group_size(8,16,-32))) void neg2(){} // expected-error{{negative argument is not allowed for 'reqd_work_group_size' attribute}}
+
+// 4294967294 is a negative integer if treated as signed.
+// Should compile successfully, since we expect an unsigned.
+__kernel __attribute__((reqd_work_group_size(8,16,4294967294))) void ok1(){}