[OpenCL] Refine OpenCLImageAccessAttr to OpenCLAccessAttr

[pxli168]

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.

[Anastasia]

lib/Sema/SemaDeclAttr.cpp

4934

In the case of failure, where would destructor of AccessAttr be called?

Could we invoke it before exiting? Or alternatively it's possible to check an access kind based on string matching from Attr getName() and create AccessAttr later on success.

5694

So where would the attribute check happen? I don't think the comment is very clear.

lib/Sema/SemaType.cpp

6223

remove one -

test/SemaOpenCL/invalid-kernel-attrs.cl

31

Strangely, I don't see this diagnostic being added in this patch.

[pxli168]

lib/Sema/SemaDeclAttr.cpp

4934

Nice catch, I think I will try if getName() can handle this, or I think maybe I need a destructor.

5694

Maybe I don't make my self claer.

The pipe decl like

SomeFn(read_only pipe int p)

It will process read_only twice which is not expected, this will make attr add twice.

I will make the comment clearer to understand.

test/SemaOpenCL/invalid-kernel-attrs.cl

31

As you can see, the read_only used to use with int. Now we check the type, so int could not used for this test more. Just change the in to image1d_t to keep the old test case work.

[pxli168]

Fix destruct problem and fix style in inline comment.

[Anastasia]

I think some tests for new diagnostics are still missing?

include/clang/Basic/DiagnosticSemaKinds.td

7676

before CL2.0 -> earlier than OpenCL2.0 versions

lib/Sema/SemaDeclAttr.cpp

4922

could we write it shorter somehow:

read_write can be used for image types to specify that an image object can be read and written.

5692

Ah, you mean the attr was already process with element type?

I don't think this comment is clear still.

Could you write something like:
Skip if pipe type as already being processes with an element type.

5696

Could you also avoid using break. I think it should be quite easy here.

test/SemaOpenCL/invalid-kernel-attrs.cl

31–32

LG!

[pxli168]

Rewrite some comments and add missing invalid test case.

[Anastasia]

lib/Sema/SemaDeclAttr.cpp

5705

Do you mean to remove "twice" here? As I think that's what you are trying to avoid.

test/Parser/opencl-image-access.cl

2

please, add -verify flag

10

Btw, I can see that read_write is now accepted even if -cl-std=CL1.1. So essentially you don't need to add CL20 checks here.

But I might prepare a quick fix of it afterwards.

test/SemaOpenCL/invalid-access-qualifier.cl

12

this is being tested in test/Parser/opencl-image-access.cl already

[pxli168]

test/Parser/opencl-image-access.cl

10

Yes, this is what this patch added in. It will check if read_write used earlier than CL2.0 as the spec said.

test/SemaOpenCL/invalid-access-qualifier.cl

12

I will remove it.

[pxli168]

Remove repeat test case.

[pxli168]

lib/Sema/SemaDeclAttr.cpp

5068

I used arc to update the patch and it case some line error.
I have check about OpenCL version here and will handle the read_write with image before OpenCL 2.0. As spec mentioned that it is reversed before OpenCL 2.0.

[Anastasia]

lib/Sema/SemaDeclAttr.cpp

5068

Do you think we could just check:

if (S.getLangOpts().OpenCLVersion < 200 || DeclTy->isPipeType())

[pxli168]

Make some optimization

lib/Sema/SemaDeclAttr.cpp

5068

It seems can work, but maybe mislead to others.
But with the comment above it seems ok.

[Anastasia]

Adding a few final comments, otherwise, looks good!

lib/Sema/SemaDeclAttr.cpp

5051

there is only one

5059

OpenCL v2.0 s6.6: -> OpenCL v2.0 s6.6 -

5061

OpenCL v2.0 s6.13.6: -> OpenCL v2.0 s6.13.6 -

5816

Did you mean to say:

"Skip pipe type, otherwise it will be processed twice with its element type"

lib/Sema/SemaType.cpp

6243

OpenCL v2.0 s6.6: -> OpenCL v2.0 s6.6 -

Qualifier -> qualifier

test/Parser/opencl-image-access.cl

2

ping

[Anastasia]

LGTM!

@pekka, any comments here? We will finalize it if not.

lib/Sema/SemaDeclAttr.cpp

5816

. missing

[pekka.jaaskelainen]

Otherwise LGTM.

lib/Sema/SemaChecking.cpp

267–268

This TODO can be removed?

[aaron.ballman]

include/clang/Basic/Attr.td

678

Please, if you are updating this attribute, document it.

lib/Sema/SemaDeclAttr.cpp

5051

Missing punctuation. ;-)

5053

This should be pointing to the attribute, not the declaration, shouldn't it?

5064

Drop the llvm namespace specifier.

5068

This seems like it should also be pointing to the attribute instead of the declaration (I assume the declaration would be fine were it not for the presence of the attribute).

5816

Missing punctuation.

Of bigger concern -- this code is highly suspect -- this is called for every declaration in the source code. Why is this complexity needed here instead of elsewhere?

[Anastasia]

lib/Sema/SemaDeclAttr.cpp

5816

I assume we could do this check in ProcessDeclAttributeList or just after:

if (const AttributeList *Attrs = PD.getTypeObject(i).getAttrs())

[aaron.ballman]

lib/Sema/SemaDeclAttr.cpp

5816

Is this an attribute that's written on the type declarator, and needs to be "slid" onto the declaration? If so, we do that in ProcessDeclAttributes, around line 5794. But we should not be adding ad hoc special cases to filter these out, they should probably be handled in distributeTypeAttrsFromDeclarator() in SemaType.cpp.

[pxli168]

Update doc for OpenCL access qualifier, still find where to handle access qualifier for pipe type.

[pxli168]

Need to refine access qualifier with pipe type.

lib/Sema/SemaDeclAttr.cpp

5053

In this case there will be two attributes, so I think it should just point to the declaration.

5816–5817

I think this attribute one that need to be slid on the declaration. Now I just handle this in ProcessDeclAttributes, with

const ParmVarDecl *PDecl = llvm::dyn_cast<ParmVarDecl>(D);
if (!PDecl ||
    !PDecl->getType().getCanonicalType().getTypePtr()->isPipeType()) {

And I tried to handle it in distributeTypeAttrsFromDeclarator but it seems this attribute does not go through that path and I am still working on how to place this check.

[aaron.ballman]

lib/Sema/SemaDeclAttr.cpp

5053

That's why I think it should be pointing at the attribute. The presence of the second attribute is what triggers the diagnostic, not anything about the declaration. The location should always point to "this thing is what caused the diagnostic."

5069

Spurious TODO?

5816–5817

You should look into why it does not make its way into distributeTypeAttrsFromDeclarator(); as it stands, I am uncomfortable with the current approach.

[pxli168]

Hi Anastasia,
It seems the access qualifier attribute may need to be handled like the you have mentioned in clang-commit that to be in type.
What is you opinion? I can write a patch for that too.

Thanks

lib/Sema/SemaDeclAttr.cpp

5816–5817

Hi Aaron,
It seems that OpenCL type that use this attribute did not go through this path, and I will find a way to handle this.

Thanks

[pxli168]

It seems it is related with the pipe type. I am still working on how to fix this problem.

[pxli168]

Refine the pipe parse to solve the problem the attribute for pipe will be handled twice within Declarator

[Anastasia]

lib/Sema/SemaDeclAttr.cpp

5053

Yes, I think attribute would make more sense.

5068

Yes, w/o this attribute it would assume default access qualifier i.e. read_only. So I agree to point to the attribute here too.

[pxli168]

lib/Sema/SemaDeclAttr.cpp

5053

Already changed to attribute.

5068

Already done.

[Anastasia]

A few small comments!

include/clang/Basic/AttrDocs.td

1578

what about __read_write?

include/clang/Basic/DiagnosticSemaKinds.td

7742

could you write __read_write/read_write please!

test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl

24

This test doesn't seem to work correctly because it didn't detect that we were checking for the wrong error message!

I am thinking it might be the issue of adding "-ferror-limit 100" to the RUN line as it might have stopped to diagnose after certain number of errors given. Could you please double check this?

[Anastasia]

test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl

24

I just checked it and I guess it's passing because it's doing substring check, so no need to add anything to the RUN line. Your change is fine here!

[aaron.ballman]

Mostly minor comments, but I like this approach!

include/clang/Basic/DiagnosticSemaKinds.td

7742

Better yet, use the spelling the user wrote, since you have access to that information already.

lib/Sema/SemaType.cpp

6240

Missing a period at the end of the sentence.

6243

Missing a period at the end of the sentence.

[Anastasia]

include/clang/Basic/AttrDocs.td

1578

Oh, I see it now. May be we could group

__read_only/read_only, __write_only/write_only ...

[pxli168]

Remove test case for access quilifier in test/SemaOpenCL/invalid-kernel-attrs.cl.
Due to the patch http://reviews.llvm.org/D17437.
read_only can only be used in parameters with pipe and image type.

test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl

24

I found these bug by testing with the access qualifier, and find some word missing and just add them in this patch.

[Anastasia]

LGTM!

[aaron.ballman]

LGTM, thanks!

include/clang/Basic/Attr.td

 def OpenCLUnrollHint : InheritableAttr {
   let Spellings = [GNU<"opencl_unroll_hint">];
   let Args = [UnsignedArgument<"UnrollHint">];
   let Documentation = [OpenCLUnrollHintDocs];
 }
 
 // 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">];
   let Subjects = SubjectList<[ParmVar], ErrorDiag>;
   let Accessors = [Accessor<"isReadOnly", [Keyword<"__read_only">,
                                            Keyword<"read_only">]>,
                    Accessor<"isReadWrite", [Keyword<"__read_write">,
                                             Keyword<"read_write">]>,
                    Accessor<"isWriteOnly", [Keyword<"__write_only">,
                                             Keyword<"write_only">]>];
-  let Documentation = [Undocumented];
+  let Documentation = [OpenCLAccessDocs];
 }
 
 def OpenCLPrivateAddressSpace : TypeAttr {
   let Spellings = [Keyword<"__private">, Keyword<"private">];
   let Documentation = [OpenCLAddressSpacePrivateDocs];
 }
 
 def OpenCLGlobalAddressSpace : TypeAttr {

include/clang/Basic/AttrDocs.td

 
 def OpenCLUnrollHintDocs : Documentation {
   let Category = DocCatStmt;
   let Heading = "__attribute__((opencl_unroll_hint))";
   let Content = [{
 The opencl_unroll_hint attribute qualifier can be used to specify that a loop
 (for, while and do loops) can be unrolled. This attribute qualifier can be
 used to specify full unrolling or partial unrolling by a specified amount.
 This is a compiler hint and the compiler may ignore this directive. See
 `OpenCL v2.0 <https://www.khronos.org/registry/cl/specs/opencl-2.0.pdf>`_
 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
 used to allocate the object. OpenCL supports the following address spaces:
 __generic(generic), __global(global), __local(local), __private(private),
 __constant(constant).
 
   .. code-block:: c

include/clang/Basic/DiagnosticSemaKinds.td

def err_opencl_private_ptr_kernel_param : Error<
   "kernel parameter cannot be declared as a pointer to the __private address space">;
 def err_opencl_non_kernel_variable : Error<
   "non-kernel function variable cannot be declared in %0 address space">;
 def err_static_function_scope : Error<
   "variables in function scope cannot be declared static">;
 def err_opencl_bitfields : Error<
   "bitfields are not supported in OpenCL">;
 def err_opencl_vla : Error<
   "variable length arrays are not supported in OpenCL">;
 def err_bad_kernel_param_type : Error<
   "%0 cannot be used as the type of a kernel parameter">;
 def err_record_with_pointers_kernel_param : Error<
   "%select{struct|union}0 kernel parameters may not contain pointers">;
 def note_within_field_of_type : Note<
   "within field of type %0 declared here">;
 def note_illegal_field_declared_here : Note<
   "field of illegal %select{type|pointer type}0 %1 declared here">;
def err_opencl_builtin_pipe_first_arg : Error<
   "first argument to %0 must be a pipe type">;
 def err_opencl_builtin_pipe_arg_num : Error<
     "invalid number of arguments to function: %0">;
 def err_opencl_builtin_pipe_invalid_arg : Error<
   "invalid argument type to function %0 (expecting %1)">;
 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">;
 
 } // end of sema category
 
 let CategoryName = "OpenMP Issue" in {
 // OpenMP support.

lib/CodeGen/CodeGenFunction.cpp

if (ty->isPointerType()) {
         typeQuals += typeQuals.empty() ? "volatile" : " volatile";
       if (isPipe)
         typeQuals = "pipe";
     }
 
     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<OpenCLImageAccessAttr>();
+      const OpenCLAccessAttr *A = parm->getAttr<OpenCLAccessAttr>();
       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"));
 
     // Get argument name.
     argNames.push_back(llvm::MDString::get(Context, parm->getName()));
   }
 
   kernelMDArgs.push_back(llvm::MDNode::get(Context, addressQuals));

lib/Parse/ParseDecl.cpp

if (SS.isNotEmpty()) {
                     /* Don't replace range end. */SourceLocation());
       return;
     }
   }
 
   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()),
         DS.getAttributes(), SourceLocation());
   }
 
   // Not a pointer, C++ reference, or block.
   if (!isPtrOperatorToken(Kind, getLangOpts(), D.getContext())) {

lib/Sema/SemaChecking.cpp

static bool SemaBuiltinSEHScopeCheck(Sema &SemaRef, CallExpr *TheCall,
   return false;
 }
 
 /// Returns readable name for a call.
 static StringRef getFunctionName(CallExpr *Call) {
   return cast<FunctionDecl>(Call->getCalleeDecl())->getName();
 }
 
 /// Returns OpenCL access qual.
-// TODO: Refine OpenCLImageAccessAttr to OpenCLAccessAttr since pipe can use
-// it too
-static OpenCLImageAccessAttr *getOpenCLArgAccess(const Decl *D) {
-  if (D->hasAttr<OpenCLImageAccessAttr>())
-    return D->getAttr<OpenCLImageAccessAttr>();
+static OpenCLAccessAttr *getOpenCLArgAccess(const Decl *D) {
+  if (D->hasAttr<OpenCLAccessAttr>())
+    return D->getAttr<OpenCLAccessAttr>();
   return nullptr;
 }
 
 /// Returns true if pipe element type is different from the pointer.
 static bool checkOpenCLPipeArg(Sema &S, CallExpr *Call) {
   const Expr *Arg0 = Call->getArg(0);
   // First argument type should always be pipe.
   if (!Arg0->getType()->isPipeType()) {
     S.Diag(Call->getLocStart(), diag::err_opencl_builtin_pipe_first_arg)
         << getFunctionName(Call) << Arg0->getSourceRange();
     return true;
   }
-  OpenCLImageAccessAttr *AccessQual =
+  OpenCLAccessAttr *AccessQual =
       getOpenCLArgAccess(cast<DeclRefExpr>(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
   // read_only and write_only, and assumed to be read_only if no qualifier is
   // specified.
   bool isValid = true;
   bool ReadOnly = getFunctionName(Call).find("read") != StringRef::npos;
   if (ReadOnly)

lib/Sema/SemaDeclAttr.cpp

static void handleReleaseCapabilityAttr(Sema &S, Decl *D,
   checkAttrArgsAreCapabilityObjs(S, D, Attr, Args, 0, true);
 
   D->addAttr(::new (S.Context) ReleaseCapabilityAttr(
       Attr.getRange(), S.Context, Args.data(), Args.size(),
       Attr.getAttributeSpellingListIndex()));
 }
 
 static void handleRequiresCapabilityAttr(Sema &S, Decl *D,
                                          const AttributeList &Attr) {
   if (!checkAttributeAtLeastNumArgs(S, Attr, 1))
     return;
 
   // check that all arguments are lockable objects
   SmallVector<Expr*, 1> Args;
   checkAttrArgsAreCapabilityObjs(S, D, Attr, Args);
   if (Args.empty())
     return;
 
   RequiresCapabilityAttr *RCA = ::new (S.Context)
     RequiresCapabilityAttr(Attr.getRange(), S.Context, Args.data(),
                            Args.size(), Attr.getAttributeSpellingListIndex());
 
   D->addAttr(RCA);
 }
 
 static void handleDeprecatedAttr(Sema &S, Decl *D, const AttributeList &Attr) {
   if (auto *NSD = dyn_cast<NamespaceDecl>(D)) {
     if (NSD->isAnonymousNamespace()) {
       S.Diag(Attr.getLoc(), diag::warn_deprecated_anonymous_namespace);
static bool handleCommonAttributeFeatures(Sema &S, Scope *scope, Decl *D,
 
   // Check whether the attribute appertains to the given subject.
   if (!Attr.diagnoseAppertainsTo(S, D))
     return true;
 
   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<OpenCLAccessAttr>()) {
+    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<ParmVarDecl>(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
 //===----------------------------------------------------------------------===//
 
 /// ProcessDeclAttribute - Apply the specific attribute to the specified decl if
 /// the attribute applies to decls.  If the attribute is a type attribute, just
 /// silently ignore it if a GNU attribute.
 static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
   case AttributeList::AT_SysVABI:
   case AttributeList::AT_Pcs:
   case AttributeList::AT_IntelOclBicc:
     handleCallConvAttr(S, D, Attr);
     break;
   case AttributeList::AT_OpenCLKernel:
     handleSimpleAttribute<OpenCLKernelAttr>(S, D, Attr);
     break;
-  case AttributeList::AT_OpenCLImageAccess:
-    handleSimpleAttribute<OpenCLImageAccessAttr>(S, D, Attr);
+  case AttributeList::AT_OpenCLAccess:
+    handleOpenCLAccessAttr(S, D, Attr);
     break;
   case AttributeList::AT_InternalLinkage:
     handleInternalLinkageAttr(S, D, Attr);
     break;
 
   // Microsoft attributes:
   case AttributeList::AT_MSNoVTable:
     handleSimpleAttribute<MSNoVTableAttr>(S, D, Attr);
for ( ; A; A = A->getNext()) {
     if (A->getKind() == AttributeList::UnknownAttribute) {
       S.Diag(A->getLoc(), diag::warn_unknown_attribute_ignored)
         << A->getName() << A->getRange();
     } else {
       S.Diag(A->getLoc(), diag::warn_attribute_not_on_decl)
         << A->getName() << A->getRange();
     }
   }
 }
 
 /// checkUnusedDeclAttributes - Given a declarator which is not being
 /// used to build a declaration, complain about any decl attributes
 /// which might be lying around on it.
 void Sema::checkUnusedDeclAttributes(Declarator &D) {
   ::checkUnusedDeclAttributes(*this, D.getDeclSpec().getAttributes().getList());
   ::checkUnusedDeclAttributes(*this, D.getAttributes());
   for (unsigned i = 0, e = D.getNumTypeObjects(); i != e; ++i)
     ::checkUnusedDeclAttributes(*this, D.getTypeObject(i).getAttrs());
 }
 
 /// DeclClonePragmaWeak - clone existing decl (maybe definition),
 /// \#pragma weak needs a non-definition decl and source may not have one.
 NamedDecl * Sema::DeclClonePragmaWeak(NamedDecl *ND, IdentifierInfo *II,
                                       SourceLocation Loc) {
   assert(isa<FunctionDecl>(ND) || isa<VarDecl>(ND));
   NamedDecl *NewD = nullptr;
   if (FunctionDecl *FD = dyn_cast<FunctionDecl>(ND)) {
     FunctionDecl *NewFD;
     // FIXME: Missing call to CheckFunctionDeclaration().
     // FIXME: Mangling?
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);
 
   // Finally, apply any attributes on the decl itself.
   if (const AttributeList *Attrs = PD.getAttributes())
     ProcessDeclAttributeList(S, D, Attrs);
 }
 
 /// Is the given declaration allowed to use a forbidden type?
 /// If so, it'll still be annotated with an attribute that makes it
 /// illegal to actually use.

lib/Sema/SemaType.cpp

S.Diag(Attr.getLoc(), diag::err_attribute_argument_type)
       << Attr.getName() << AANT_ArgumentIntegerConstant
       << numEltsExpr->getSourceRange();
     Attr.setInvalid();
     return;
   }
   // Only certain element types are supported for Neon vectors.
   if (!isPermittedNeonBaseType(CurType, VecKind, S)) {
     S.Diag(Attr.getLoc(), diag::err_attribute_invalid_vector_type) << CurType;
     Attr.setInvalid();
     return;
   }
 
   // The total size of the vector must be 64 or 128 bits.
   unsigned typeSize = static_cast<unsigned>(S.Context.getTypeSize(CurType));
   unsigned numElts = static_cast<unsigned>(numEltsInt.getZExtValue());
   unsigned vecSize = typeSize * numElts;
   if (vecSize != 64 && vecSize != 128) {
     S.Diag(Attr.getLoc(), diag::err_attribute_bad_neon_vector_size) << CurType;
     Attr.setInvalid();
     return;
   }
 
   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
   // attributes (such as __address_space__, __vector_size__, etc) apply to the
   // type, but others can be present in the type specifiers even though they
   // apply to the decl.  Here we apply type attributes and ignore the rest.
 
   bool hasOpenCLAddressSpace = false;
case AttributeList::AT_NeonVectorType:
                                VectorType::NeonVector);
       attr.setUsedAsTypeAttr();
       break;
     case AttributeList::AT_NeonPolyVectorType:
       HandleNeonVectorTypeAttr(type, attr, state.getSema(),
                                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;
 
     MS_TYPE_ATTRS_CASELIST:
       if (!handleMSPointerTypeQualifierAttr(state, attr, type))
         attr.setUsedAsTypeAttr();
       break;
 

test/Parser/opencl-image-access.cl

-// 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

test/SemaOpenCL/invalid-access-qualifier.cl

+// 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

test/SemaOpenCL/invalid-kernel-attrs.cl

 constant int foo1 __attribute__((reqd_work_group_size(8,16,32))) = 0; // expected-error {{'reqd_work_group_size' attribute only applies to functions}}
 
 constant int foo2 __attribute__((work_group_size_hint(8,16,32))) = 0; // expected-error {{'work_group_size_hint' attribute only applies to functions}}
 
 constant int foo3 __attribute__((vec_type_hint(char))) = 0; // expected-error {{'vec_type_hint' attribute only applies to functions}}
 
 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}}
 kernel __attribute__((reqd_work_group_size(1,0,2))) void kernel12(){} // expected-error {{'reqd_work_group_size' attribute must be greater than 0}}
 kernel __attribute__((reqd_work_group_size(0,1,2))) void kernel13(){} // expected-error {{'reqd_work_group_size' attribute must be greater than 0}}

test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl

void test1(read_only pipe int p, global int* ptr){
   // reserve_read/write_pipe
   reserve_read_pipe(p, ptr);    // expected-error{{invalid argument type to function reserve_read_pipe (expecting 'unsigned int')}}
   work_group_reserve_read_pipe(tmp, tmp);    // expected-error{{first argument to work_group_reserve_read_pipe must be a pipe type}}
   sub_group_reserve_write_pipe(p, tmp);    // expected-error{{invalid pipe access modifier (expecting write_only)}}
 
   // 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){
   int tmp;
   reserve_id_t rid;
 
   // read/write_pipe
   write_pipe(tmp, p);    // expected-error {{first argument to write_pipe must be a pipe type}}
   write_pipe(p);   // expected-error {{invalid number of arguments to function: write_pipe}}
   write_pipe(p, tmp, tmp, ptr);   // expected-error {{invalid argument type to function write_pipe (expecting 'reserve_id_t')}}
   write_pipe(p, rid, rid, ptr);   // expected-error {{invalid argument type to function write_pipe (expecting 'unsigned int')}}
   write_pipe(p, tmp);   // expected-error {{invalid argument type to function write_pipe (expecting 'int *')}}
   read_pipe(p, ptr);    // expected-error {{invalid pipe access modifier (expecting read_only)}}
   read_pipe(p, rid, tmp, ptr);    // expected-error {{invalid pipe access modifier (expecting read_only)}}
 
   // reserve_read/write_pipe
   reserve_write_pipe(p, ptr);    // expected-error{{invalid argument type to function reserve_write_pipe (expecting 'unsigned int')}}
   work_group_reserve_write_pipe(tmp, tmp);    // expected-error{{first argument to work_group_reserve_write_pipe must be a pipe type}}
   sub_group_reserve_read_pipe(p, tmp);    // expected-error{{invalid pipe access modifier (expecting read_only)}}
 
   // 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(){
   int tmp;
   get_pipe_num_packets(tmp);    // expected-error {{first argument to get_pipe_num_packets must be a pipe type}}
   get_pipe_max_packets(tmp);    // expected-error {{first argument to get_pipe_max_packets must be a pipe type}}
 }