[SYCL] Add diagnostic on command group with multiple actions

sycl/include/CL/sycl/handler.hpp

@@ -190,6 +190,14 @@ class handler {
     return Storage;
   }
 
+  void throwIfActionIsCreated() {
+    if (detail::CG::NONE != MCGType)
+      throw sycl::runtime_error("Attempt to set multiple actions for the "
+                                "command group. Command group must consist of "
+                                "a single kernel or explicit memory operation.",
+                                CL_INVALID_OPERATION);
+  }
+
   // The method extracts and prepares kernel arguments from the lambda using
   // integration header.
   void
@@ -676,6 +684,7 @@ class handler {
   // single_task version with a kernel represented as a lambda.
   template <typename KernelName = detail::auto_name, typename KernelType>
   void single_task(KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -693,6 +702,7 @@ class handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for(range<Dims> NumWorkItems, KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -706,6 +716,7 @@ class handler {
 
   // Similar to single_task, but passed lambda will be executed on host.
   template <typename FuncT> void run_on_host_intel(FuncT Func) {
+    throwIfActionIsCreated();
     MNDRDesc.set(range<1>{1});
 
     MArgs = std::move(MAssociatedAccesors);
@@ -719,6 +730,7 @@ class handler {
             int Dims>
   void parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
                     KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -735,6 +747,7 @@ class handler {
   template <typename KernelName = detail::auto_name, typename KernelType,
             int Dims>
   void parallel_for(nd_range<Dims> ExecutionRange, KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -750,6 +763,7 @@ class handler {
             int Dims>
   void parallel_for_work_group(range<Dims> NumWorkGroups,
                                KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -766,6 +780,7 @@ class handler {
   void parallel_for_work_group(range<Dims> NumWorkGroups,
                                range<Dims> WorkGroupSize,
                                KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -780,6 +795,7 @@ class handler {
   // single_task version with a kernel represented as a sycl::kernel.
   // The kernel invocation method has no functors and cannot be called on host.
   void single_task(kernel SyclKernel) {
+    throwIfActionIsCreated();
     verifySyclKernelInvoc(SyclKernel);
     MNDRDesc.set(range<1>{1});
     MSyclKernel = detail::getSyclObjImpl(std::move(SyclKernel));
@@ -792,6 +808,7 @@ class handler {
   // functors and cannot be called on host.
   template <int Dims>
   void parallel_for(range<Dims> NumWorkItems, kernel SyclKernel) {
+    throwIfActionIsCreated();
     verifySyclKernelInvoc(SyclKernel);
     MSyclKernel = detail::getSyclObjImpl(std::move(SyclKernel));
     MNDRDesc.set(std::move(NumWorkItems));
@@ -805,6 +822,7 @@ class handler {
   template <int Dims>
   void parallel_for(range<Dims> NumWorkItems, id<Dims> workItemOffset,
                     kernel SyclKernel) {
+    throwIfActionIsCreated();
     verifySyclKernelInvoc(SyclKernel);
     MSyclKernel = detail::getSyclObjImpl(std::move(SyclKernel));
     MNDRDesc.set(std::move(NumWorkItems), std::move(workItemOffset));
@@ -817,6 +835,7 @@ class handler {
   // method has no functors and cannot be called on host.
   template <int Dims>
   void parallel_for(nd_range<Dims> NDRange, kernel SyclKernel) {
+    throwIfActionIsCreated();
     verifySyclKernelInvoc(SyclKernel);
     MSyclKernel = detail::getSyclObjImpl(std::move(SyclKernel));
     MNDRDesc.set(std::move(NDRange));
@@ -833,6 +852,7 @@ class handler {
   // which is used otherwise.
   template <typename KernelName = detail::auto_name, typename KernelType>
   void single_task(kernel SyclKernel, KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -855,6 +875,7 @@ class handler {
             int Dims>
   void parallel_for(kernel SyclKernel, range<Dims> NumWorkItems,
                     KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -877,6 +898,7 @@ class handler {
             int Dims>
   void parallel_for(kernel SyclKernel, range<Dims> NumWorkItems,
                     id<Dims> WorkItemOffset, KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -899,6 +921,7 @@ class handler {
             int Dims>
   void parallel_for(kernel SyclKernel, nd_range<Dims> NDRange,
                     KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -924,6 +947,7 @@ class handler {
             int Dims>
   void parallel_for_work_group(kernel SyclKernel, range<Dims> NumWorkGroups,
                                KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -943,6 +967,7 @@ class handler {
   void parallel_for_work_group(kernel SyclKernel, range<Dims> NumWorkGroups,
                                range<Dims> WorkGroupSize,
                                KernelType KernelFunc) {
+    throwIfActionIsCreated();
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
@@ -977,6 +1002,7 @@ class handler {
             access::placeholder IsPlaceholder = access::placeholder::false_t>
   void copy(accessor<T_Src, Dims, AccessMode, AccessTarget, IsPlaceholder> Src,
             shared_ptr_class<T_Dst> Dst) {
+    throwIfActionIsCreated();
     static_assert(isValidTargetForExplicitOp(AccessTarget),
                   "Invalid accessor target for the copy method.");
     // Make sure data shared_ptr points to is not released until we finish
@@ -993,6 +1019,7 @@ class handler {
   void
   copy(shared_ptr_class<T_Src> Src,
        accessor<T_Dst, Dims, AccessMode, AccessTarget, IsPlaceholder> Dst) {
+    throwIfActionIsCreated();
     static_assert(isValidTargetForExplicitOp(AccessTarget),
                   "Invalid accessor target for the copy method.");
     // Make sure data shared_ptr points to is not released until we finish
@@ -1008,6 +1035,7 @@ class handler {
             access::placeholder IsPlaceholder = access::placeholder::false_t>
   void copy(accessor<T_Src, Dims, AccessMode, AccessTarget, IsPlaceholder> Src,
             T_Dst *Dst) {
+    throwIfActionIsCreated();
     static_assert(isValidTargetForExplicitOp(AccessTarget),
                   "Invalid accessor target for the copy method.");
 #ifndef __SYCL_DEVICE_ONLY__
@@ -1047,6 +1075,7 @@ class handler {
   void
   copy(const T_Src *Src,
        accessor<T_Dst, Dims, AccessMode, AccessTarget, IsPlaceholder> Dst) {
+    throwIfActionIsCreated();
     static_assert(isValidTargetForExplicitOp(AccessTarget),
                   "Invalid accessor target for the copy method.");
 #ifndef __SYCL_DEVICE_ONLY__
@@ -1124,6 +1153,7 @@ class handler {
             accessor<T_Dst, Dims_Dst, AccessMode_Dst, AccessTarget_Dst,
                      IsPlaceholder_Dst>
                 Dst) {
+    throwIfActionIsCreated();
     static_assert(isValidTargetForExplicitOp(AccessTarget_Src),
                   "Invalid source accessor target for the copy method.");
     static_assert(isValidTargetForExplicitOp(AccessTarget_Dst),
@@ -1177,6 +1207,7 @@ class handler {
             access::placeholder IsPlaceholder = access::placeholder::false_t>
   void
   update_host(accessor<T, Dims, AccessMode, AccessTarget, IsPlaceholder> Acc) {
+    throwIfActionIsCreated();
     static_assert(isValidTargetForExplicitOp(AccessTarget),
                   "Invalid accessor target for the update_host method.");
     MCGType = detail::CG::UPDATE_HOST;
@@ -1198,6 +1229,7 @@ class handler {
             access::placeholder IsPlaceholder = access::placeholder::false_t>
   void fill(accessor<T, Dims, AccessMode, AccessTarget, IsPlaceholder> Dst,
             const T &Pattern) {
+    throwIfActionIsCreated();
     // TODO add check:T must be an integral scalar value or a SYCL vector type
     static_assert(isValidTargetForExplicitOp(AccessTarget),
                   "Invalid accessor target for the fill method.");
@@ -1229,6 +1261,7 @@ class handler {
 
   // Copy memory from the source to the destination.
   void memcpy(void *Dest, const void *Src, size_t Count) {
+    throwIfActionIsCreated();
     MSrcPtr = const_cast<void *>(Src);
     MDstPtr = Dest;
     MLength = Count;
@@ -1237,6 +1270,7 @@ class handler {
 
   // Fill the memory pointed to by the destination with the given bytes.
   void memset(void *Dest, int Value, size_t Count) {
+    throwIfActionIsCreated();
     MDstPtr = Dest;
     MPattern.push_back((char)Value);
     MLength = Count;
@@ -1245,6 +1279,7 @@ class handler {
 
   // Prefetch the memory pointed to by the pointer.
   void prefetch(const void *Ptr, size_t Count) {
+    throwIfActionIsCreated();
     MDstPtr = const_cast<void *>(Ptr);
     MLength = Count;
     MCGType = detail::CG::PREFETCH_USM;

sycl/test/basic_tests/diagnostics/handler.cpp

@@ -0,0 +1,37 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out | FileCheck %s
+//==------------------- handler.cpp ----------------------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl.hpp>
+
+using namespace cl;
+
+int main() {
+
+  bool Failed = false;
+
+  sycl::queue Queue([](sycl::exception_list ExceptionList) {
+    if (ExceptionList.size() != 1) {
+      std::cerr << "Should be one exception in exception list" << std::endl;
+      std::abort();
+    }
+    std::rethrow_exception(*ExceptionList.begin());
+  });
+
+  try {
+    Queue.submit([&](sycl::handler &CGH) {
+      CGH.single_task<class Dummy1>([]() {});
+      CGH.single_task<class Dummy2>([]() {});
+    });
+    Queue.throw_asynchronous();
+  } catch (sycl::exception &E) {
+    // CHECK: Attempt to set multiple actions for the command group
+    std::cout << E.what() << std::endl;
+  }
+}