pxli168 updated this revision to Diff 47455.
pxli168 marked 5 inline comments as done.

http://reviews.llvm.org/D16040

Files:
  include/clang/Basic/Attr.td
  include/clang/Basic/DiagnosticSemaKinds.td
  lib/CodeGen/CodeGenFunction.cpp
  lib/Sema/SemaChecking.cpp
  lib/Sema/SemaDeclAttr.cpp
  lib/Sema/SemaType.cpp
  test/Parser/opencl-image-access.cl
  test/SemaOpenCL/invalid-access-qualifier.cl
  test/SemaOpenCL/invalid-kernel-attrs.cl

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}}
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/Parser/opencl-image-access.cl
===================================================================
--- test/Parser/opencl-image-access.cl
+++ 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
Index: lib/Sema/SemaType.cpp
===================================================================
--- lib/Sema/SemaType.cpp
+++ lib/Sema/SemaType.cpp
@@ -6225,6 +6225,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
@@ -6320,9 +6331,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: 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 is 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 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<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(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<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);
@@ -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, otherwise it will be processed twice with its element type
+  const ParmVarDecl *PDecl = llvm::dyn_cast<ParmVarDecl>(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/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<OpenCLImageAccessAttr>())
-    return D->getAttr<OpenCLImageAccessAttr>();
+static OpenCLAccessAttr *getOpenCLArgAccess(const Decl *D) {
+  if (D->hasAttr<OpenCLAccessAttr>())
+    return D->getAttr<OpenCLAccessAttr>();
   return nullptr;
 }
 
@@ -281,7 +281,7 @@
         << 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
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<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"));
 
Index: include/clang/Basic/DiagnosticSemaKinds.td
===================================================================
--- include/clang/Basic/DiagnosticSemaKinds.td
+++ include/clang/Basic/DiagnosticSemaKinds.td
@@ -7706,6 +7706,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: 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">];
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to