From d353e6d748e0c151884345f2c08f1b988c5e7b94 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Thu, 6 Sep 2018 11:54:09 +0000 Subject: [PATCH] [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 llvm-svn: 341539 --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 ++ clang/lib/Sema/SemaDeclAttr.cpp | 14 ++++++++++++-- clang/test/SemaOpenCL/invalid-kernel-attrs.cl | 7 +++++++ 3 files changed, 21 insertions(+), 2 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index c1f0774..03d8f70 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -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">; def err_init_priority_object_attr : Error< "can only use 'init_priority' attribute on file-scope definitions " "of objects of class type">; diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index d72b8fb..488299a 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/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 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) diff --git a/clang/test/SemaOpenCL/invalid-kernel-attrs.cl b/clang/test/SemaOpenCL/invalid-kernel-attrs.cl index a19a989..9a0bcd8 100644 --- a/clang/test/SemaOpenCL/invalid-kernel-attrs.cl +++ b/clang/test/SemaOpenCL/invalid-kernel-attrs.cl @@ -37,3 +37,10 @@ kernel __attribute__((reqd_work_group_size(0,1,2))) void kernel13(){} // expecte __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(){} -- 2.7.4