From 2d12e2d6c969a4a6249b87cff4c337ed141bd65f Mon Sep 17 00:00:00 2001 From: Xiuli Pan Date: Fri, 26 Feb 2016 03:13:03 +0000 Subject: [PATCH] [OpenCL] Refine OpenCLImageAccessAttr to OpenCLAccessAttr Summary: OpenCL access qualifiers are now not only used for image types, refine it to avoid misleading, Add semacheck for OpenCL access qualifier as well as test caees. Reviewers: pekka.jaaskelainen, Anastasia, aaron.ballman Subscribers: aaron.ballman, cfe-commits Differential Revision: http://reviews.llvm.org/D16040 git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@261961 91177308-0d34-0410-b5e6-96231b3b80d8 --- include/clang/Basic/Attr.td | 4 +- include/clang/Basic/AttrDocs.td | 26 +++++++++++++ include/clang/Basic/DiagnosticSemaKinds.td | 8 ++++ lib/CodeGen/CodeGenFunction.cpp | 7 ++-- lib/Parse/ParseDecl.cpp | 3 +- lib/Sema/SemaChecking.cpp | 10 ++--- lib/Sema/SemaDeclAttr.cpp | 38 ++++++++++++++++++- lib/Sema/SemaType.cpp | 16 ++++++-- test/Parser/opencl-image-access.cl | 9 ++++- test/SemaOpenCL/invalid-access-qualifier.cl | 14 +++++++ test/SemaOpenCL/invalid-kernel-attrs.cl | 2 - test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl | 4 +- 12 files changed, 117 insertions(+), 24 deletions(-) create mode 100644 test/SemaOpenCL/invalid-access-qualifier.cl diff --git a/include/clang/Basic/Attr.td b/include/clang/Basic/Attr.td index c7f559d9b3..b77a533719 100644 --- a/include/clang/Basic/Attr.td +++ b/include/clang/Basic/Attr.td @@ -664,7 +664,7 @@ def OpenCLUnrollHint : InheritableAttr { // This attribute is both a type attribute, and a declaration attribute (for // parameter variables). -def OpenCLImageAccess : Attr { +def OpenCLAccess : Attr { let Spellings = [Keyword<"__read_only">, Keyword<"read_only">, Keyword<"__write_only">, Keyword<"write_only">, Keyword<"__read_write">, Keyword<"read_write">]; @@ -675,7 +675,7 @@ def OpenCLImageAccess : Attr { Keyword<"read_write">]>, Accessor<"isWriteOnly", [Keyword<"__write_only">, Keyword<"write_only">]>]; - let Documentation = [Undocumented]; + let Documentation = [OpenCLAccessDocs]; } def OpenCLPrivateAddressSpace : TypeAttr { diff --git a/include/clang/Basic/AttrDocs.td b/include/clang/Basic/AttrDocs.td index 5ed44de420..f1a6ae2b1c 100644 --- a/include/clang/Basic/AttrDocs.td +++ b/include/clang/Basic/AttrDocs.td @@ -1581,6 +1581,32 @@ s6.11.5 for details. }]; } +def OpenCLAccessDocs : Documentation { + let Category = DocCatStmt; + let Content = [{ +The access qualifiers must be used with image object arguments or pipe arguments +to declare if they are being read or written by a kernel or function. + +The read_only/__read_only, write_only/__write_only and read_write/__read_write +names are reserved for use as access qualifiers and shall not be used otherwise. + + .. code-block:: c + kernel void + foo (read_only image2d_t imageA, + write_only image2d_t imageB) + { + ... + } + +In the above example imageA is a read-only 2D image object, and imageB is a +write-only 2D image object. + +The read_write (or __read_write) qualifier can not be used with pipe. + +More details can be found in the OpenCL C language Spec v2.0, Section 6.6. + }]; +} + def DocOpenCLAddressSpaces : DocumentationCategory<"OpenCL Address Spaces"> { let Content = [{ The address space qualifier may be used to specify the region of memory that is diff --git a/include/clang/Basic/DiagnosticSemaKinds.td b/include/clang/Basic/DiagnosticSemaKinds.td index 70a841711a..7f5ecca9b6 100644 --- a/include/clang/Basic/DiagnosticSemaKinds.td +++ b/include/clang/Basic/DiagnosticSemaKinds.td @@ -7735,6 +7735,14 @@ def err_opencl_builtin_pipe_invalid_arg : Error< def err_opencl_builtin_pipe_invalid_access_modifier : Error< "invalid pipe access modifier (expecting %0)">; +// OpenCL access qualifier +def err_opencl_invalid_access_qualifier : Error< + "access qualifier can only be used for pipe and image type">; +def err_opencl_invalid_read_write : Error< + "access qualifier %0 can not be used for %1 %select{|earlier than OpenCL2.0 version}2">; +def err_opencl_multiple_access_qualifiers : Error< + "multiple access qualifiers">; + // OpenCL Section 6.8.g def err_opencl_unknown_type_specifier : Error< "OpenCL does not support the '%0' %select{type qualifier|storage class specifier}1">; diff --git a/lib/CodeGen/CodeGenFunction.cpp b/lib/CodeGen/CodeGenFunction.cpp index 3d0b17b924..8e327c2bc4 100644 --- a/lib/CodeGen/CodeGenFunction.cpp +++ b/lib/CodeGen/CodeGenFunction.cpp @@ -561,15 +561,14 @@ static void GenOpenCLArgMetadata(const FunctionDecl *FD, llvm::Function *Fn, argTypeQuals.push_back(llvm::MDString::get(Context, typeQuals)); // Get image and pipe access qualifier: - // FIXME: now image and pipe share the same access qualifier maybe we can - // refine it to OpenCL access qualifier and also handle write_read if (ty->isImageType()|| ty->isPipeType()) { - const OpenCLImageAccessAttr *A = parm->getAttr(); + const OpenCLAccessAttr *A = parm->getAttr(); if (A && A->isWriteOnly()) accessQuals.push_back(llvm::MDString::get(Context, "write_only")); + else if (A && A->isReadWrite()) + accessQuals.push_back(llvm::MDString::get(Context, "read_write")); else accessQuals.push_back(llvm::MDString::get(Context, "read_only")); - // FIXME: what about read_write? } else accessQuals.push_back(llvm::MDString::get(Context, "none")); diff --git a/lib/Parse/ParseDecl.cpp b/lib/Parse/ParseDecl.cpp index 14b11285ea..ab6f3ccc2d 100644 --- a/lib/Parse/ParseDecl.cpp +++ b/lib/Parse/ParseDecl.cpp @@ -4989,7 +4989,8 @@ void Parser::ParseDeclaratorInternal(Declarator &D, tok::TokenKind Kind = Tok.getKind(); if (D.getDeclSpec().isTypeSpecPipe() && !isPipeDeclerator(D)) { - DeclSpec &DS = D.getMutableDeclSpec(); + DeclSpec DS(AttrFactory); + ParseTypeQualifierListOpt(DS); D.AddTypeInfo( DeclaratorChunk::getPipe(DS.getTypeQualifiers(), DS.getPipeLoc()), diff --git a/lib/Sema/SemaChecking.cpp b/lib/Sema/SemaChecking.cpp index a1f975a826..134248dbe2 100644 --- a/lib/Sema/SemaChecking.cpp +++ b/lib/Sema/SemaChecking.cpp @@ -265,11 +265,9 @@ static StringRef getFunctionName(CallExpr *Call) { } /// Returns OpenCL access qual. -// TODO: Refine OpenCLImageAccessAttr to OpenCLAccessAttr since pipe can use -// it too -static OpenCLImageAccessAttr *getOpenCLArgAccess(const Decl *D) { - if (D->hasAttr()) - return D->getAttr(); +static OpenCLAccessAttr *getOpenCLArgAccess(const Decl *D) { + if (D->hasAttr()) + return D->getAttr(); return nullptr; } @@ -282,7 +280,7 @@ static bool checkOpenCLPipeArg(Sema &S, CallExpr *Call) { << getFunctionName(Call) << Arg0->getSourceRange(); return true; } - OpenCLImageAccessAttr *AccessQual = + OpenCLAccessAttr *AccessQual = getOpenCLArgAccess(cast(Arg0)->getDecl()); // Validates the access qualifier is compatible with the call. // OpenCL v2.0 s6.13.16 - The access qualifiers for pipe should only be diff --git a/lib/Sema/SemaDeclAttr.cpp b/lib/Sema/SemaDeclAttr.cpp index 3f6af3d765..b031e38c75 100644 --- a/lib/Sema/SemaDeclAttr.cpp +++ b/lib/Sema/SemaDeclAttr.cpp @@ -5043,6 +5043,40 @@ static bool handleCommonAttributeFeatures(Sema &S, Scope *scope, Decl *D, return false; } +static void handleOpenCLAccessAttr(Sema &S, Decl *D, + const AttributeList &Attr) { + if (D->isInvalidDecl()) + return; + + // Check if there is only one access qualifier. + if (D->hasAttr()) { + S.Diag(Attr.getLoc(), diag::err_opencl_multiple_access_qualifiers) + << D->getSourceRange(); + D->setInvalidDecl(true); + return; + } + + // OpenCL v2.0 s6.6 - read_write can be used for image types to specify that an + // image object can be read and written. + // OpenCL v2.0 s6.13.6 - A kernel cannot read from and write to the same pipe + // object. Using the read_write (or __read_write) qualifier with the pipe + // qualifier is a compilation error. + if (const ParmVarDecl *PDecl = dyn_cast(D)) { + const Type *DeclTy = PDecl->getType().getCanonicalType().getTypePtr(); + if (Attr.getName()->getName().find("read_write") != StringRef::npos) { + if (S.getLangOpts().OpenCLVersion < 200 || DeclTy->isPipeType()) { + S.Diag(Attr.getLoc(), diag::err_opencl_invalid_read_write) + << Attr.getName() << PDecl->getType() << DeclTy->isImageType(); + D->setInvalidDecl(true); + return; + } + } + } + + D->addAttr(::new (S.Context) OpenCLAccessAttr( + Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex())); +} + //===----------------------------------------------------------------------===// // Top Level Sema Entry Points //===----------------------------------------------------------------------===// @@ -5440,8 +5474,8 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D, case AttributeList::AT_OpenCLKernel: handleSimpleAttribute(S, D, Attr); break; - case AttributeList::AT_OpenCLImageAccess: - handleSimpleAttribute(S, D, Attr); + case AttributeList::AT_OpenCLAccess: + handleOpenCLAccessAttr(S, D, Attr); break; case AttributeList::AT_InternalLinkage: handleInternalLinkageAttr(S, D, Attr); diff --git a/lib/Sema/SemaType.cpp b/lib/Sema/SemaType.cpp index fb37743a71..dbc325c149 100644 --- a/lib/Sema/SemaType.cpp +++ b/lib/Sema/SemaType.cpp @@ -6237,6 +6237,17 @@ static void HandleNeonVectorTypeAttr(QualType& CurType, CurType = S.Context.getVectorType(CurType, numElts, VecKind); } +/// Handle OpenCL Access Qualifier Attribute. +static void HandleOpenCLAccessAttr(QualType &CurType, const AttributeList &Attr, + Sema &S) { + // OpenCL v2.0 s6.6 - Access qualifier can used only for image and pipe type. + if (!(CurType->isImageType() || CurType->isPipeType())) { + S.Diag(Attr.getLoc(), diag::err_opencl_invalid_access_qualifier); + Attr.setInvalid(); + return; + } +} + static void processTypeAttrs(TypeProcessingState &state, QualType &type, TypeAttrLocation TAL, AttributeList *attrs) { // Scan through and apply attributes to this type where it makes sense. Some @@ -6332,9 +6343,8 @@ static void processTypeAttrs(TypeProcessingState &state, QualType &type, VectorType::NeonPolyVector); attr.setUsedAsTypeAttr(); break; - case AttributeList::AT_OpenCLImageAccess: - // FIXME: there should be some type checking happening here, I would - // imagine, but the original handler's checking was entirely superfluous. + case AttributeList::AT_OpenCLAccess: + HandleOpenCLAccessAttr(type, attr, state.getSema()); attr.setUsedAsTypeAttr(); break; diff --git a/test/Parser/opencl-image-access.cl b/test/Parser/opencl-image-access.cl index e08d129214..99ced8e32b 100644 --- a/test/Parser/opencl-image-access.cl +++ b/test/Parser/opencl-image-access.cl @@ -1,14 +1,19 @@ -// RUN: %clang_cc1 %s -fsyntax-only +// RUN: %clang_cc1 %s -fsyntax-only -verify +// RUN: %clang_cc1 %s -fsyntax-only -verify -cl-std=CL2.0 -DCL20 +// expected-no-diagnostics __kernel void f__ro(__read_only image2d_t a) { } __kernel void f__wo(__write_only image2d_t a) { } +#if CL20 __kernel void f__rw(__read_write image2d_t a) { } - +#endif __kernel void fro(read_only image2d_t a) { } __kernel void fwo(write_only image2d_t a) { } +#if CL20 __kernel void frw(read_write image2d_t a) { } +#endif diff --git a/test/SemaOpenCL/invalid-access-qualifier.cl b/test/SemaOpenCL/invalid-access-qualifier.cl new file mode 100644 index 0000000000..b090bb5733 --- /dev/null +++ b/test/SemaOpenCL/invalid-access-qualifier.cl @@ -0,0 +1,14 @@ +// RUN: %clang_cc1 -verify %s +// RUN: %clang_cc1 -verify -cl-std=CL2.0 -DCL20 %s + +void test1(read_only int i){} // expected-error{{access qualifier can only be used for pipe and image type}} + +void test2(read_only write_only image1d_t i){} // expected-error{{multiple access qualifiers}} + +void test3(read_only read_only image1d_t i){} // expected-error{{multiple access qualifiers}} + +#ifdef CL20 +void test4(read_write pipe int i){} // expected-error{{access qualifier 'read_write' can not be used for 'pipe'}} +#else +void test4(__read_write image1d_t i){} // expected-error{{access qualifier '__read_write' can not be used for 'image1d_t' earlier than OpenCL2.0 version}} +#endif diff --git a/test/SemaOpenCL/invalid-kernel-attrs.cl b/test/SemaOpenCL/invalid-kernel-attrs.cl index 4b4fdf79e3..cedbb06646 100644 --- a/test/SemaOpenCL/invalid-kernel-attrs.cl +++ b/test/SemaOpenCL/invalid-kernel-attrs.cl @@ -28,8 +28,6 @@ constant int foo3 __attribute__((vec_type_hint(char))) = 0; // expected-error {{ void f_kernel_image2d_t( kernel image2d_t image ) { // expected-error {{'kernel' attribute only applies to functions}} int __kernel x; // expected-error {{'__kernel' attribute only applies to functions}} - read_only int i; // expected-error {{'read_only' attribute only applies to parameters}} - __write_only int j; // expected-error {{'__write_only' attribute only applies to parameters}} } kernel __attribute__((reqd_work_group_size(1,2,0))) void kernel11(){} // expected-error {{'reqd_work_group_size' attribute must be greater than 0}} diff --git a/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl b/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl index 3b613b57cf..7836a25e7f 100644 --- a/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl +++ b/test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl @@ -21,7 +21,7 @@ void test1(read_only pipe int p, global int* ptr){ // commit_read/write_pipe commit_read_pipe(tmp, rid); // expected-error{{first argument to commit_read_pipe must be a pipe type}} work_group_commit_read_pipe(p, tmp); // expected-error{{invalid argument type to function work_group_commit_read_pipe (expecting 'reserve_id_t')}} - sub_group_commit_write_pipe(p, tmp); // expected-error{{nvalid pipe access modifier (expecting write_only)}} + sub_group_commit_write_pipe(p, tmp); // expected-error{{invalid pipe access modifier (expecting write_only)}} } void test2(write_only pipe int p, global int* ptr){ @@ -45,7 +45,7 @@ void test2(write_only pipe int p, global int* ptr){ // commit_read/write_pipe commit_write_pipe(tmp, rid); // expected-error{{first argument to commit_write_pipe must be a pipe type}} work_group_commit_write_pipe(p, tmp); // expected-error{{invalid argument type to function work_group_commit_write_pipe (expecting 'reserve_id_t')}} - sub_group_commit_read_pipe(p, tmp); // expected-error{{nvalid pipe access modifier (expecting read_only)}} + sub_group_commit_read_pipe(p, tmp); // expected-error{{invalid pipe access modifier (expecting read_only)}} } void test3(){ -- 2.40.0