[OpenCL] Disallow negative attribute arguments

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.

Reviewers: Anastasia, yaxunl, bader

Reviewed By: Anastasia, yaxunl, bader

Subscribers: bader, cfe-commits

Differential Revision: https://reviews.llvm.org/D50259

git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@341539 91177308-0d34-0410-b5e6-96231b3b80d8

1 files changed, 12 insertions, 2 deletions

diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp
index d72b8fbb1a..488299a970 100644
--- a/lib/Sema/SemaDeclAttr.cpp
+++ b/lib/Sema/SemaDeclAttr.cpp
@@ -227,9 +227,13 @@ static SourceLocation getAttrLoc(const ParsedAttr &AL) { return AL.getLoc(); }
 
 /// 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)) {
@@ -249,6 +253,11 @@ static bool checkUInt32Argument(Sema &S, const AttrInfo &AI, const Expr *Expr,
     return false;
   }
 
+  if (StrictlyUnsigned && I.isSigned() && I.isNegative()) {
+    S.Diag(getAttrLoc(AI), diag::err_attribute_argument_negative) << AI;
+    return false;
+  }
+
   Val = (uint32_t)I.getZExtValue();
   return true;
 }
@@ -2766,7 +2775,8 @@ 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)