On Thu, Sep 6, 2018 at 7:54 AM, Andrew Savonichev via cfe-commits <cfe-commits@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@lists.llvm.org > http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
