r341539 - [OpenCL] Disallow negative attribute arguments

[Aaron Ballman via cfe-commits]

On Thu, Sep 6, 2018 at 7:54 AM, Andrew Savonichev via cfe-commits
<cfe-commits at lists.llvm.org> wrote:
> Author: asavonic
> Date: Thu Sep  6 04:54:09 2018
> New Revision: 341539
>
> URL: http://llvm.org/viewvc/llvm-project?rev=341539&view=rev
> Log:
> [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
>
> Modified:
>     cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
>     cfe/trunk/lib/Sema/SemaDeclAttr.cpp
>     cfe/trunk/test/SemaOpenCL/invalid-kernel-attrs.cl
>
> Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td
> URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=341539&r1=341538&r2=341539&view=diff
> ==============================================================================
> --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original)
> +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Thu Sep  6 04:54:09 2018
> @@ -2529,6 +2529,8 @@ def err_attribute_argument_type : Error<
>    "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">;

I don't think we need a new diagnostic here as we already have
err_attribute_requires_positive_integer.

~Aaron

>  def err_init_priority_object_attr : Error<
>    "can only use 'init_priority' attribute on file-scope definitions "
>    "of objects of class type">;
>
> Modified: cfe/trunk/lib/Sema/SemaDeclAttr.cpp
> URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaDeclAttr.cpp?rev=341539&r1=341538&r2=341539&view=diff
> ==============================================================================
> --- cfe/trunk/lib/Sema/SemaDeclAttr.cpp (original)
> +++ cfe/trunk/lib/Sema/SemaDeclAttr.cpp Thu Sep  6 04:54:09 2018
> @@ -227,9 +227,13 @@ static SourceLocation getAttrLoc(const P
>
>  /// 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,
>      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,
>    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)
>
> Modified: cfe/trunk/test/SemaOpenCL/invalid-kernel-attrs.cl
> URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/SemaOpenCL/invalid-kernel-attrs.cl?rev=341539&r1=341538&r2=341539&view=diff
> ==============================================================================
> --- cfe/trunk/test/SemaOpenCL/invalid-kernel-attrs.cl (original)
> +++ cfe/trunk/test/SemaOpenCL/invalid-kernel-attrs.cl Thu Sep  6 04:54:09 2018
> @@ -37,3 +37,10 @@ kernel __attribute__((reqd_work_group_si
>  __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(){}
>
>
> _______________________________________________
> cfe-commits mailing list
> cfe-commits at lists.llvm.org
> http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits