Allow passing default constructed accessors to kernels (#10780)

Prior to this PR, we would throw or crash when a default constructed
accessor was passed to a kernel if it was not filtered out, even if it
wasn't really accessed. Passing a default constructed accessor to a
kernel is allowed by SYCL2020, and this PR ensures our implementation
allows it.

---------

Signed-off-by: Maronas, Marcos <[email protected]>

Author: maarquitos14

sycl/source/detail/scheduler/commands.cpp

@@ -1537,9 +1537,10 @@ AllocaCommandBase *ExecCGCommand::getAllocaForReq(Requirement *Req) {
     if (Dep.MDepRequirement == Req)
       return Dep.MAllocaCmd;
   }
-  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
-                        "Alloca for command not found " +
-                            codeToString(PI_ERROR_INVALID_OPERATION));
+  // Default constructed accessors do not add dependencies, but they can be
+  // passed to commands. Simply return nullptr, since they are empty and don't
+  // really require any memory.
+  return nullptr;
 }
 
 std::vector<std::shared_ptr<const void>>
@@ -2213,11 +2214,15 @@ void SetArgBasedOnType(
     break;
   case kernel_param_kind_t::kind_accessor: {
     Requirement *Req = (Requirement *)(Arg.MPtr);
-    assert(getMemAllocationFunc != nullptr &&
-           "We should have caught this earlier.");
 
+    // getMemAllocationFunc is nullptr when there are no requirements. However,
+    // we may pass default constructed accessors to a command, which don't add
+    // requirements. In such case, getMemAllocationFunc is nullptr, but it's a
+    // valid case, so we need to properly handle it.
     sycl::detail::pi::PiMem MemArg =
-        (sycl::detail::pi::PiMem)getMemAllocationFunc(Req);
+        getMemAllocationFunc
+            ? (sycl::detail::pi::PiMem)getMemAllocationFunc(Req)
+            : nullptr;
     if (Context.get_backend() == backend::opencl) {
       // clSetKernelArg (corresponding to piKernelSetArg) returns an error
       // when MemArg is null, which is the case when zero-sized buffers are
@@ -2831,7 +2836,9 @@ pi_int32 ExecCGCommand::enqueueImpQueue() {
 
     auto getMemAllocationFunc = [this](Requirement *Req) {
       AllocaCommandBase *AllocaCmd = getAllocaForReq(Req);
-      return AllocaCmd->getMemAllocation();
+      // getAllocaForReq may return nullptr if Req is a default constructed
+      // accessor. Simply return nullptr in such a case.
+      return AllocaCmd ? AllocaCmd->getMemAllocation() : nullptr;
     };
 
     const std::shared_ptr<detail::kernel_impl> &SyclKernel =

sycl/test-e2e/Basic/accessor/accessor.cpp

@@ -1491,5 +1491,72 @@ int main() {
     assert(!result);
   }
 
+  // default constructed accessor can be passed to a kernel.
+  {
+    AccT acc;
+    sycl::queue q;
+    bool result;
+    {
+      sycl::buffer<bool, 1> Buf{&result, sycl::range<1>{1}};
+      // We are passing a default constructed accessor and a non default
+      // constructed accessor with storage. Default constructed accessors can be
+      // passed to commands, but trying to access the (non-existing) underlying
+      // storage is UB. This test should work, since the access to the default
+      // constructed accessor must never be reached.
+      try {
+        q.submit([&](sycl::handler &cgh) {
+          sycl::accessor res_acc{Buf, cgh};
+          cgh.single_task<class def_ctor_kernel>([=] {
+            if (false)
+              res_acc[0] = acc[0];
+          });
+        });
+        q.wait_and_throw();
+      } catch (sycl::exception &e) {
+        assert("Unexpected exception");
+      } catch (...) {
+        std::cout << "Some other unexpected exception (line " << __LINE__ << ")"
+                  << std::endl;
+        return 1;
+      }
+    }
+    assert(!result);
+  }
+
+  // default constructed accessor can be passed to a kernel (2).
+  {
+    using AccT = sycl::accessor<int, 1, sycl::access::mode::read_write>;
+    AccT acc;
+    assert(acc.empty());
+    sycl::queue q;
+    bool result;
+    {
+      // We are passing only a default constructed accessor. Default constructed
+      // accessors can be passed to commands, but trying to access the
+      // (non-existing) underlying storage is UB. This test should work, since
+      // the access to the default constructed accessor must never be reached.
+      // The difference with the previous test case is that in this case the
+      // task will not have any requirements, while the previous one does have
+      // one requirement for the non default constructed accessor, testing
+      // different code paths.
+      try {
+        q.submit([&](sycl::handler &cgh) {
+          cgh.single_task<class def_ctor_kernel2>([=] {
+            if (!acc.empty())
+              acc[0] = 1;
+          });
+        });
+        q.wait_and_throw();
+      } catch (sycl::exception &e) {
+        assert("Unexpected exception");
+      } catch (...) {
+        std::cout << "Some other unexpected exception (line " << __LINE__ << ")"
+                  << std::endl;
+        return 1;
+      }
+    }
+    assert(!result);
+  }
+
   std::cout << "Test passed" << std::endl;
 }