pxli168 updated this revision to Diff 49007.
http://reviews.llvm.org/D16040
Files:
include/clang/Basic/Attr.td
include/clang/Basic/AttrDocs.td
include/clang/Basic/DiagnosticSemaKinds.td
lib/CodeGen/CodeGenFunction.cpp
lib/Parse/ParseDecl.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
test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl
Index: test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl
===================================================================
--- test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl
+++ test/SemaOpenCL/invalid-pipe-builtin-cl2.0.cl
@@ -21,7 +21,7 @@
// 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 @@
// 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(){
Index: test/SemaOpenCL/invalid-kernel-attrs.cl
===================================================================
--- test/SemaOpenCL/invalid-kernel-attrs.cl
+++ test/SemaOpenCL/invalid-kernel-attrs.cl
@@ -28,8 +28,6 @@
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}}
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
@@ -6237,6 +6237,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
@@ -6332,9 +6343,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
@@ -5043,6 +5043,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(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
//===----------------------------------------------------------------------===//
@@ -5440,8 +5474,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);
Index: lib/Sema/SemaChecking.cpp
===================================================================
--- lib/Sema/SemaChecking.cpp
+++ lib/Sema/SemaChecking.cpp
@@ -265,11 +265,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;
}
@@ -282,7 +280,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/Parse/ParseDecl.cpp
===================================================================
--- lib/Parse/ParseDecl.cpp
+++ lib/Parse/ParseDecl.cpp
@@ -4989,7 +4989,8 @@
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()),
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
@@ -7735,6 +7735,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 %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">;
Index: include/clang/Basic/AttrDocs.td
===================================================================
--- include/clang/Basic/AttrDocs.td
+++ include/clang/Basic/AttrDocs.td
@@ -1581,6 +1581,32 @@
}];
}
+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
Index: include/clang/Basic/Attr.td
===================================================================
--- include/clang/Basic/Attr.td
+++ include/clang/Basic/Attr.td
@@ -664,7 +664,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">];
@@ -675,7 +675,7 @@
Keyword<"read_write">]>,
Accessor<"isWriteOnly", [Keyword<"__write_only">,
Keyword<"write_only">]>];
- let Documentation = [Undocumented];
+ let Documentation = [OpenCLAccessDocs];
}
def OpenCLPrivateAddressSpace : TypeAttr {
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
