[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.

5692

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–32

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.

5692

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–32

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

9

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

11 ↗

(On Diff #45681)

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

[pxli168]

test/Parser/opencl-image-access.cl

9

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

11 ↗

(On Diff #45681)

I will remove it.

[pxli168]

Remove repeat test case.

[pxli168]

lib/Sema/SemaDeclAttr.cpp

5067

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

5067

Do you think we could just check:

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

[pxli168]

Make some optimization

lib/Sema/SemaDeclAttr.cpp

5067

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

5050

there is only one

5058

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

5060

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

5823

Did you mean to say:

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

lib/Sema/SemaType.cpp

6224

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

5823

. missing

[pekka.jaaskelainen]

Otherwise LGTM.

lib/Sema/SemaChecking.cpp

267

This TODO can be removed?

[aaron.ballman]

include/clang/Basic/Attr.td

671

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

lib/Sema/SemaDeclAttr.cpp

5050

Missing punctuation. ;-)

5052

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

5063

Drop the llvm namespace specifier.

5067

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).

5823

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

5823

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

5823

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

5052

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

5804–5816

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

5046

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."

5062

Spurious TODO?

5804–5816

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

5804–5816

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

5046

Yes, I think attribute would make more sense.

5061

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

5046

Already changed to attribute.

5061

Already done.

[Anastasia]

A few small comments!

include/clang/Basic/AttrDocs.td

1572 ↗

(On Diff #48452)

what about __read_write?

include/clang/Basic/DiagnosticSemaKinds.td

7713

could you write __read_write/read_write please!

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

24 ↗

(On Diff #48452)

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 ↗

(On Diff #48452)

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

7713

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

lib/Sema/SemaType.cpp

6228

Missing a period at the end of the sentence.

6231

Missing a period at the end of the sentence.

[Anastasia]

include/clang/Basic/AttrDocs.td

1572 ↗

(On Diff #48452)

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 ↗

(On Diff #48452)

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

 
 // 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">];

include/clang/Basic/DiagnosticSemaKinds.td

 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{|before CL2.0}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">;

lib/CodeGen/CodeGenFunction.cpp

     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"));
 

lib/Sema/SemaChecking.cpp

 }
 
 /// Returns OpenCL access qual.
-// TODO:Refine OpenCLImageAccessAttr to OpenCLArgAccessAttr 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;
 }
 

         << getFunctionName(Call) << Arg0->getSourceRange();
     return true;
   }
-  OpenCLImageAccessAttr *AccessQual =
+  OpenCLAccessAttr *AccessQual =
       getOpenCLArgAccess(cast<DeclRefExpr>(Arg0)->getDecl());
   // Validates the access modifier is compatible with the call.
   // From OpenCL C Specification 6.13.16 the access qualifiers for

lib/Sema/SemaDeclAttr.cpp

   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<OpenCLAccessAttr>()) {
+    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 type The __read_write
+  // (or read_write)qualifier must be used with image object arguments of
+  // kernels and of user-defined functions to declare if the image object is
+  // being both read and written by the kernel
+  // 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<ParmVarDecl>(D)) {
+    const Type *DeclTy = PDecl->getType().getCanonicalType().getTypePtr();
+    if (Attr.getName()->getName().find("read_write") != StringRef::npos) {
+      if (DeclTy->isPipeType() ||
+          (S.getLangOpts().OpenCLVersion < 200 && DeclTy->isImageType())) {
+        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
 //===----------------------------------------------------------------------===//

   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);

   // 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)
+  for (unsigned i = 0, e = PD.getNumTypeObjects(); i != e; ++i) {
+    // The pipe type has the same form:
+    // attr pipe int p;
+    // which will process attr for twice, but we do not need to process that for
+    // twice, so skip this process if it is pipe type
+    if (const ParmVarDecl *PDecl = llvm::dyn_cast<ParmVarDecl>(D)) {
+      if (PDecl->getType().getCanonicalType().getTypePtr()->isPipeType())
+        break;
+    }
     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())

lib/Sema/SemaType.cpp

   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

                                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;
 

test/Parser/opencl-image-access.cl

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

test/SemaOpenCL/invalid-kernel-attrs.cl

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