Index: include/clang/Basic/Attr.td =================================================================== --- include/clang/Basic/Attr.td +++ include/clang/Basic/Attr.td @@ -657,7 +657,7 @@ // 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">]; Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -7703,6 +7703,14 @@ 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 read_write can not be used for %0 %select{|earlier than OpenCL2.0 version}1">; +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">; Index: lib/CodeGen/CodeGenFunction.cpp =================================================================== --- lib/CodeGen/CodeGenFunction.cpp +++ lib/CodeGen/CodeGenFunction.cpp @@ -561,15 +561,14 @@ 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")); Index: lib/Sema/SemaChecking.cpp =================================================================== --- lib/Sema/SemaChecking.cpp +++ lib/Sema/SemaChecking.cpp @@ -266,9 +266,9 @@ /// 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; } @@ -281,7 +281,7 @@ << 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 Index: lib/Sema/SemaDeclAttr.cpp =================================================================== --- lib/Sema/SemaDeclAttr.cpp +++ lib/Sema/SemaDeclAttr.cpp @@ -5042,6 +5042,40 @@ return false; } +static void handleOpenCLAccessAttr(Sema &S, Decl *D, + const AttributeList &Attr) { + if (D->isInvalidDecl()) + return; + + // Check if there only one access qualifier + if (D->hasAttr()) { + S.Diag(D->getLocation(), 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 = llvm::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(D->getLocation(), diag::err_opencl_invalid_read_write) + << PDecl->getType() << DeclTy->isImageType() << D->getSourceRange(); + D->setInvalidDecl(true); + return; + } + } + } + + D->addAttr(::new (S.Context) OpenCLAccessAttr( + Attr.getRange(), S.Context, Attr.getAttributeSpellingListIndex())); +} + //===----------------------------------------------------------------------===// // Top Level Sema Entry Points //===----------------------------------------------------------------------===// @@ -5453,8 +5487,8 @@ 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); @@ -5786,13 +5820,19 @@ if (const AttributeList *Attrs = PD.getDeclSpec().getAttributes().getList()) ProcessDeclAttributeList(S, D, Attrs); - // Walk the declarator structure, applying decl attributes that were in a type - // position to the decl itself. This handles cases like: - // int *__attr__(x)** D; - // when X is a decl attribute. - for (unsigned i = 0, e = PD.getNumTypeObjects(); i != e; ++i) - if (const AttributeList *Attrs = PD.getTypeObject(i).getAttrs()) - ProcessDeclAttributeList(S, D, Attrs, /*IncludeCXX11Attributes=*/false); + // Skip pipe type, it will be processed twice with its element type + const ParmVarDecl *PDecl = llvm::dyn_cast(D); + if (!PDecl || + !PDecl->getType().getCanonicalType().getTypePtr()->isPipeType()) { + // Walk the declarator structure, applying decl attributes that were in a + // type position to the decl itself. This handles cases like: + // int *__attr__(x)** D; + // when X is a decl attribute. + for (unsigned i = 0, e = PD.getNumTypeObjects(); i != e; ++i) { + if (const AttributeList *Attrs = PD.getTypeObject(i).getAttrs()) + ProcessDeclAttributeList(S, D, Attrs, /*IncludeCXX11Attributes=*/false); + } + } // Finally, apply any attributes on the decl itself. if (const AttributeList *Attrs = PD.getAttributes()) Index: lib/Sema/SemaType.cpp =================================================================== --- lib/Sema/SemaType.cpp +++ lib/Sema/SemaType.cpp @@ -6218,6 +6218,17 @@ 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 @@ -6313,9 +6324,8 @@ 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; Index: test/Parser/opencl-image-access.cl =================================================================== --- test/Parser/opencl-image-access.cl +++ test/Parser/opencl-image-access.cl @@ -1,14 +1,18 @@ // RUN: %clang_cc1 %s -fsyntax-only +// RUN: %clang_cc1 %s -fsyntax-only -cl-std=CL2.0 -DCL20 __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 Index: test/SemaOpenCL/invalid-access-qualifier.cl =================================================================== --- /dev/null +++ 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 Index: test/SemaOpenCL/invalid-kernel-attrs.cl =================================================================== --- test/SemaOpenCL/invalid-kernel-attrs.cl +++ test/SemaOpenCL/invalid-kernel-attrs.cl @@ -28,8 +28,8 @@ 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}} + read_only image1d_t i; // expected-error {{'read_only' attribute only applies to parameters}} + __write_only image2d_t 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}}