[SYCL] Add diagnostic for annotated_pointers and annotated_args (#7053)

Types cannot be nested.

Signed-off-by: Elizabeth Andrews <elizabeth.andrews@intel.com>

clang/include/clang/Basic/Attr.td

@@ -1302,11 +1302,11 @@ def SYCLType: InheritableAttr {
                            ["accessor", "local_accessor", "spec_constant",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
-                            "private_memory", "aspect"],
+                            "private_memory", "aspect", "annotated_ptr", "annotated_arg"],
                            ["accessor", "local_accessor", "spec_constant",
                             "specialization_id", "kernel_handler", "buffer_location",
                             "no_alias", "accessor_property_list", "group",
-                            "private_memory", "aspect"]>];
+                            "private_memory", "aspect", "annotated_ptr", "annotated_arg"]>];
   // Only used internally by SYCL implementation
   let Documentation = [InternalOnly];
 }

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10390,8 +10390,8 @@ def warn_opencl_generic_address_space_arg : Warning<
   "passing non-generic address space pointer to %0"
   " may cause dynamic conversion affecting performance">,
   InGroup<Conversion>, DefaultIgnore;
-def err_bad_union_kernel_param_members : Error<
-  "%0 cannot be used inside a union kernel parameter">;
+def err_bad_kernel_param_data_members : Error<
+  "%0 cannot be a data member of a %select{union|struct}1 kernel parameter">;
 
 // OpenCL v2.0 s6.13.6 -- Builtin Pipe Functions
 def err_opencl_builtin_pipe_first_arg : Error<

clang/lib/Sema/SemaSYCL.cpp

@@ -1549,6 +1549,14 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
   bool IsInvalid = false;
   DiagnosticsEngine &Diag;
   bool IsSIMD = false;
+  // Keeps track of whether we are currently handling fields inside a struct.
+  // Fields of kernel functor or direct kernel captures will have a depth 0.
+  int StructFieldDepth = 0;
+  // Initialize with -1 so that fields of a base class of the kernel functor
+  // has depth 0. Visitor method enterStruct increments this to 0 when the base
+  // class is entered.
+  int StructBaseDepth = -1;
+
   // Check whether the object should be disallowed from being copied to kernel.
   // Return true if not copyable, false if copyable.
   bool checkNotCopyableToKernel(const FieldDecl *FD, QualType FieldTy) {
@@ -1633,6 +1641,16 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
   bool checkSyclSpecialType(QualType Ty, SourceRange Loc) {
     assert(isSyclSpecialType(Ty, SemaRef) &&
            "Should only be called on sycl special class types.");
+
+    // Annotated pointers and annotated arguments must be captured
+    // directly by the SYCL kernel.
+    if ((isSyclType(Ty, SYCLTypeAttr::annotated_ptr) ||
+         isSyclType(Ty, SYCLTypeAttr::annotated_arg)) &&
+         (StructFieldDepth > 0 || StructBaseDepth > 0))
+      return SemaRef.Diag(Loc.getBegin(),
+                          diag::err_bad_kernel_param_data_members)
+             << Ty << /*Struct*/ 1;
+
     const RecordDecl *RecD = Ty->getAsRecordDecl();
     if (IsSIMD && !isSyclAccessorType(Ty))
       return SemaRef.Diag(Loc.getBegin(),
@@ -1715,6 +1733,28 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
     IsInvalid = true;
     return isValid();
   }
+
+  bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final {
+    ++StructFieldDepth;
+    return true;
+  }
+
+  bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType) final {
+    --StructFieldDepth;
+    return true;
+  }
+
+  bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                   QualType FieldTy) final {
+    ++StructBaseDepth;
+    return true;
+  }
+
+  bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
+                   QualType FieldTy) final {
+    --StructBaseDepth;
+    return true;
+  }
 };
 
 // A type to check the validity of accessing accessor/sampler/stream
@@ -1734,7 +1774,8 @@ class SyclKernelUnionChecker : public SyclKernelFieldHandler {
   bool checkType(SourceLocation Loc, QualType Ty) {
     if (UnionCount) {
       IsInvalid = true;
-      Diag.Report(Loc, diag::err_bad_union_kernel_param_members) << Ty;
+      Diag.Report(Loc, diag::err_bad_kernel_param_data_members)
+          << Ty << /*Union*/ 0;
     }
     return isValid();
   }

clang/test/SemaSYCL/Inputs/sycl.hpp

@@ -403,6 +403,23 @@ class __SYCL_TYPE(spec_constant) spec_constant {
 private:
   T DefaultValue;
 };
+
+template <typename T, typename... Props>
+class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_arg) annotated_arg {
+  T obj;
+  #ifdef __SYCL_DEVICE_ONLY__
+    void __init(T _obj) {}
+  #endif
+};
+
+template <typename T, typename... Props>
+class __attribute__((sycl_special_class)) __SYCL_TYPE(annotated_ptr) annotated_ptr {
+  T* obj;
+  #ifdef __SYCL_DEVICE_ONLY__
+    void __init(T* _obj) {}
+  #endif
+};
+
 } // namespace experimental
 } // namespace oneapi
 } // namespace ext

clang/test/SemaSYCL/annotated_arg_or_ptr.cpp

@@ -0,0 +1,62 @@
+// RUN: %clang_cc1 -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -verify %s
+
+// Test diagnostic for nested annotated_arg and annotated_ptr type.
+
+#include "sycl.hpp"
+
+sycl::queue myQueue;
+
+struct MockProperty {};
+
+struct WrappedAnnotatedTypes {
+  // expected-error@+1 3{{'sycl::ext::oneapi::experimental::annotated_arg<int, MockProperty>' cannot be a data member of a struct kernel parameter}}
+  sycl::ext::oneapi::experimental::annotated_arg<int, MockProperty> AA;
+  // expected-error@+1 3{{'sycl::ext::oneapi::experimental::annotated_ptr<int, MockProperty>' cannot be a data member of a struct kernel parameter}}
+  sycl::ext::oneapi::experimental::annotated_ptr<int, MockProperty> AP;
+  sycl::accessor<int, 1, sycl::access::mode::read_write> Acc;
+};
+
+struct KernelBase {
+  sycl::ext::oneapi::experimental::annotated_arg<int, MockProperty> BaseAA; // OK
+  sycl::ext::oneapi::experimental::annotated_ptr<int, MockProperty> BaseAP; // OK
+  WrappedAnnotatedTypes NestedInBase; // Error
+};
+
+struct KernelFunctor : KernelBase  {
+  sycl::ext::oneapi::experimental::annotated_arg<int, MockProperty> AA; // OK
+  sycl::ext::oneapi::experimental::annotated_ptr<int, MockProperty> AP; // OK
+  void operator()() const {}
+};
+
+struct KernelFunctor2  {
+  WrappedAnnotatedTypes NestedInField; // Error
+  void operator()() const {}
+};
+
+int main() {
+  sycl::ext::oneapi::experimental::annotated_arg<int, MockProperty> AA;
+  sycl::ext::oneapi::experimental::annotated_ptr<int, MockProperty> AP;
+  WrappedAnnotatedTypes Obj;
+  myQueue.submit([&](sycl::handler &h) {
+    // expected-note@+1 {{in instantiation of}}
+    h.single_task<class kernel_half>(
+        [=]() {
+          (void)AA;  // OK
+	  (void)AP;  // OK
+	  (void)Obj; // Error
+        });
+  });
+
+  myQueue.submit([&](sycl::handler &h) {
+    KernelFunctor f;
+    // expected-note@+1 {{in instantiation of}}
+    h.single_task(f);
+  });
+
+  myQueue.submit([&](sycl::handler &h) {
+    KernelFunctor2 f2;
+    // expected-note@+1 {{in instantiation of}}
+    h.single_task(f2);
+  });
+}
+

clang/test/SemaSYCL/union-kernel-param-neg.cpp

@@ -8,7 +8,7 @@ using namespace sycl;
 
 union union_with_sampler {
   sycl::sampler smpl;
-  // expected-error@-1 {{'sycl::sampler' cannot be used inside a union kernel parameter}}
+  // expected-error@-1 {{'sycl::sampler' cannot be a data member of a union kernel parameter}}
 };
 
 template <typename name, typename Func>
@@ -23,7 +23,7 @@ int main() {
 
   union union_with_accessor {
     Accessor member_acc[1];
-    // expected-error@-1 {{'Accessor' (aka 'accessor<int, 1, access::mode::read_write, access::target::global_buffer>') cannot be used inside a union kernel parameter}}
+    // expected-error@-1 {{'Accessor' (aka 'accessor<int, 1, access::mode::read_write, access::target::global_buffer>') cannot be a data member of a union kernel parameter}}
   } union_acc;
 
   union_with_sampler Sampler;