[SYCL] Mark kernel object as invalid if incorrect accessor format (#4274)

If kernel object is marked invalid, we do not continue processing
it, i.e. the compiler does not proceed to generate the corresponding
openCL kernel. Invalid captures/kernel object fields should cause
kernel object to be marked as invalid. This patch fixes a bug where
invalid accessors did not do so.

Part 1 of fix for  - #4115

Signed-off-by: Elizabeth Andrews <[email protected]>

Author: elizabethandrews

clang/lib/Sema/SemaSYCL.cpp

@@ -1605,71 +1605,71 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
     return false;
   }
 
-  void checkPropertyListType(TemplateArgument PropList, SourceLocation Loc) {
-    if (PropList.getKind() != TemplateArgument::ArgKind::Type) {
-      SemaRef.Diag(Loc,
-                   diag::err_sycl_invalid_accessor_property_template_param);
-      return;
-    }
+  bool checkPropertyListType(TemplateArgument PropList, SourceLocation Loc) {
+    if (PropList.getKind() != TemplateArgument::ArgKind::Type)
+      return SemaRef.Diag(
+          Loc, diag::err_sycl_invalid_accessor_property_template_param);
+
     QualType PropListTy = PropList.getAsType();
-    if (!Util::isAccessorPropertyListType(PropListTy)) {
-      SemaRef.Diag(Loc,
-                   diag::err_sycl_invalid_accessor_property_template_param);
-      return;
-    }
+    if (!Util::isAccessorPropertyListType(PropListTy))
+      return SemaRef.Diag(
+          Loc, diag::err_sycl_invalid_accessor_property_template_param);
+
     const auto *AccPropListDecl =
         cast<ClassTemplateSpecializationDecl>(PropListTy->getAsRecordDecl());
-    if (AccPropListDecl->getTemplateArgs().size() != 1) {
-      SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_param_number)
-          << "accessor_property_list";
-      return;
-    }
+    if (AccPropListDecl->getTemplateArgs().size() != 1)
+      return SemaRef.Diag(Loc,
+                          diag::err_sycl_invalid_property_list_param_number)
+             << "accessor_property_list";
+
     const auto TemplArg = AccPropListDecl->getTemplateArgs()[0];
-    if (TemplArg.getKind() != TemplateArgument::ArgKind::Pack) {
-      SemaRef.Diag(Loc,
-                   diag::err_sycl_invalid_accessor_property_list_template_param)
-          << /*accessor_property_list*/ 0 << /*parameter pack*/ 0;
-      return;
-    }
+    if (TemplArg.getKind() != TemplateArgument::ArgKind::Pack)
+      return SemaRef.Diag(
+                 Loc,
+                 diag::err_sycl_invalid_accessor_property_list_template_param)
+             << /*accessor_property_list*/ 0 << /*parameter pack*/ 0;
+
     for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin();
          Prop != TemplArg.pack_end(); ++Prop) {
-      if (Prop->getKind() != TemplateArgument::ArgKind::Type) {
-        SemaRef.Diag(
-            Loc, diag::err_sycl_invalid_accessor_property_list_template_param)
-            << /*accessor_property_list pack argument*/ 1 << /*type*/ 1;
-        return;
-      }
+      if (Prop->getKind() != TemplateArgument::ArgKind::Type)
+        return SemaRef.Diag(
+                   Loc,
+                   diag::err_sycl_invalid_accessor_property_list_template_param)
+               << /*accessor_property_list pack argument*/ 1 << /*type*/ 1;
       QualType PropTy = Prop->getAsType();
-      if (Util::isSyclBufferLocationType(PropTy))
-        checkBufferLocationType(PropTy, Loc);
+      if (Util::isSyclBufferLocationType(PropTy) &&
+          checkBufferLocationType(PropTy, Loc))
+        return true;
     }
+    return false;
   }
 
-  void checkBufferLocationType(QualType PropTy, SourceLocation Loc) {
+  bool checkBufferLocationType(QualType PropTy, SourceLocation Loc) {
     const auto *PropDecl =
         cast<ClassTemplateSpecializationDecl>(PropTy->getAsRecordDecl());
-    if (PropDecl->getTemplateArgs().size() != 1) {
-      SemaRef.Diag(Loc, diag::err_sycl_invalid_property_list_param_number)
-          << "buffer_location";
-      return;
-    }
+    if (PropDecl->getTemplateArgs().size() != 1)
+      return SemaRef.Diag(Loc,
+                          diag::err_sycl_invalid_property_list_param_number)
+             << "buffer_location";
+
     const auto BufferLoc = PropDecl->getTemplateArgs()[0];
-    if (BufferLoc.getKind() != TemplateArgument::ArgKind::Integral) {
-      SemaRef.Diag(Loc,
-                   diag::err_sycl_invalid_accessor_property_list_template_param)
-          << /*buffer_location*/ 2 << /*non-negative integer*/ 2;
-      return;
-    }
+    if (BufferLoc.getKind() != TemplateArgument::ArgKind::Integral)
+      return SemaRef.Diag(
+                 Loc,
+                 diag::err_sycl_invalid_accessor_property_list_template_param)
+             << /*buffer_location*/ 2 << /*non-negative integer*/ 2;
+
     int LocationID = static_cast<int>(BufferLoc.getAsIntegral().getExtValue());
-    if (LocationID < 0) {
-      SemaRef.Diag(Loc,
-                   diag::err_sycl_invalid_accessor_property_list_template_param)
-          << /*buffer_location*/ 2 << /*non-negative integer*/ 2;
-      return;
-    }
+    if (LocationID < 0)
+      return SemaRef.Diag(
+                 Loc,
+                 diag::err_sycl_invalid_accessor_property_list_template_param)
+             << /*buffer_location*/ 2 << /*non-negative integer*/ 2;
+
+    return false;
   }
 
-  void checkAccessorType(QualType Ty, SourceRange Loc) {
+  bool checkAccessorType(QualType Ty, SourceRange Loc) {
     assert(Util::isSyclAccessorType(Ty) &&
            "Should only be called on SYCL accessor types.");
 
@@ -1678,13 +1678,13 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
             dyn_cast<ClassTemplateSpecializationDecl>(RecD)) {
       const TemplateArgumentList &TAL = CTSD->getTemplateArgs();
       TemplateArgument TA = TAL.get(0);
-      const QualType TemplateArgTy = TA.getAsType();
+      llvm::DenseSet<QualType> Visited;
+      checkSYCLType(SemaRef, TA.getAsType(), Loc, Visited);
 
       if (TAL.size() > 5)
-        checkPropertyListType(TAL.get(5), Loc.getBegin());
-      llvm::DenseSet<QualType> Visited;
-      checkSYCLType(SemaRef, TemplateArgTy, Loc, Visited);
+        return checkPropertyListType(TAL.get(5), Loc.getBegin());
     }
+    return false;
   }
 
 public:
@@ -1706,12 +1706,12 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler {
 
   bool handleSyclAccessorType(const CXXRecordDecl *, const CXXBaseSpecifier &BS,
                               QualType FieldTy) final {
-    checkAccessorType(FieldTy, BS.getBeginLoc());
+    IsInvalid |= checkAccessorType(FieldTy, BS.getBeginLoc());
     return isValid();
   }
 
   bool handleSyclAccessorType(FieldDecl *FD, QualType FieldTy) final {
-    checkAccessorType(FieldTy, FD->getLocation());
+    IsInvalid |= checkAccessorType(FieldTy, FD->getLocation());
     return isValid();
   }
 

clang/test/SemaSYCL/buffer_location.cpp

@@ -54,6 +54,15 @@ int main() {
                      cl::sycl::ext::oneapi::accessor_property_list<
                          another_property>>
       accessorC;
+
+  cl::sycl::kernel_single_task<class kernel_function>(
+      [=]() {
+        // expected-no-diagnostics
+        Obj.use();
+        accessorA.use();
+        accessorB.use();
+        accessorC.use();
+      });
 #else
   cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
                      cl::sycl::access::target::global_buffer,
@@ -72,23 +81,18 @@ int main() {
                          buffer_location<1>,
                          buffer_location<2>>>
       accessorF;
-#endif
   cl::sycl::kernel_single_task<class kernel_function>(
       [=]() {
-#ifndef TRIGGER_ERROR
-        // expected-no-diagnostics
-        Obj.use();
-        accessorA.use();
-        accessorB.use();
-        accessorC.use();
-#else
         //expected-error@+1{{buffer_location template parameter must be a non-negative integer}}
         accessorD.use();
         //expected-error@+1{{sixth template parameter of the accessor must be of accessor_property_list type}}
         accessorE.use();
+      });
+  cl::sycl::kernel_single_task<class kernel_function>(
+      [=]() {
         //expected-error@+1{{can't apply buffer_location property twice to the same accessor}}
         accessorF.use();
-#endif
       });
+#endif
   return 0;
 }