[SYCL] Defer shadow copy creation in SYCLMemObjT (#11348)

#10334 causes a performance regression
since `HostPtr` can't be reused when it's read-only. This PR fixes the
regression by deferring the copy operation to the creation of a writable
accessor. It includes following the changes:

- A new `SYCLMemObjT::MCreateShadowCopy` to defer allocation. When the
`HostPtr` cannot be reused since it's read-only,
`SYCLMemObjT::handleHostData` sets this member to a function that will
allocate the shadow copy.
- A new `SYCLMemObjT::handleWriteAccessorCreation` member function. This
function calls `SYCLMemObjT::MCreateShadowCopy` and updates any existing
`MAllocaCommands` if `MUserPtr` changed.
- Whenever a writable host or device accessor is created,
`handleWriteAccessorCreation` gets called to ensure that any required
memory allocation occurs.

With this change, the allocation and copying overhead occurs during the
creation of the first writable accessor. There's no overhead if all of
the relevant accessors use `sycl::access_mode::read`.

---------

Signed-off-by: Michael Aziz <[email protected]>

Author: 0x12CC

sycl/source/detail/scheduler/graph_builder.cpp

@@ -521,6 +521,11 @@ Command *
 Scheduler::GraphBuilder::addHostAccessor(Requirement *Req,
                                          std::vector<Command *> &ToEnqueue) {
 
+  if (Req->MAccessMode != sycl::access_mode::read) {
+    auto SYCLMemObj = static_cast<detail::SYCLMemObjT *>(Req->MSYCLMemObj);
+    SYCLMemObj->handleWriteAccessorCreation();
+  }
+
   const QueueImplPtr &HostQueue = getInstance().getDefaultHostQueue();
 
   MemObjRecord *Record = getOrInsertMemObjRecord(HostQueue, Req, ToEnqueue);

sycl/source/detail/sycl_mem_obj_t.cpp

@@ -233,6 +233,19 @@ void SYCLMemObjT::detachMemoryObject(
     Scheduler::getInstance().deferMemObjRelease(Self);
 }
 
+void SYCLMemObjT::handleWriteAccessorCreation() {
+  const auto InitialUserPtr = MUserPtr;
+  MCreateShadowCopy();
+  MCreateShadowCopy = []() -> void {};
+  if (MRecord != nullptr && MUserPtr != InitialUserPtr) {
+    for (auto &it : MRecord->MAllocaCommands) {
+      if (it->MMemAllocation == InitialUserPtr) {
+        it->MMemAllocation = MUserPtr;
+      }
+    }
+  }
+}
+
 } // namespace detail
 } // namespace _V1
 } // namespace sycl

sycl/source/detail/sycl_mem_obj_t.hpp

@@ -173,10 +173,14 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
            has_property<property::image::use_host_ptr>();
   }
 
-  bool canReuseHostPtr(void *HostPtr, const size_t RequiredAlign) {
+  bool canReadHostPtr(void *HostPtr, const size_t RequiredAlign) {
     bool Aligned =
         (reinterpret_cast<std::uintptr_t>(HostPtr) % RequiredAlign) == 0;
-    return !MHostPtrReadOnly && (Aligned || useHostPtr());
+    return Aligned || useHostPtr();
+  }
+
+  bool canReuseHostPtr(void *HostPtr, const size_t RequiredAlign) {
+    return !MHostPtrReadOnly && canReadHostPtr(HostPtr, RequiredAlign);
   }
 
   void handleHostData(void *HostPtr, const size_t RequiredAlign) {
@@ -190,6 +194,14 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
     if (HostPtr) {
       if (canReuseHostPtr(HostPtr, RequiredAlign)) {
         MUserPtr = HostPtr;
+      } else if (canReadHostPtr(HostPtr, RequiredAlign)) {
+        MUserPtr = HostPtr;
+        MCreateShadowCopy = [this, RequiredAlign, HostPtr]() -> void {
+          setAlign(RequiredAlign);
+          MShadowCopy = allocateHostMem();
+          MUserPtr = MShadowCopy;
+          std::memcpy(MUserPtr, HostPtr, MSizeInBytes);
+        };
       } else {
         setAlign(RequiredAlign);
         MShadowCopy = allocateHostMem();
@@ -213,9 +225,17 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
       if (!MHostPtrReadOnly)
         set_final_data_from_storage();
 
-      if (canReuseHostPtr(HostPtr.get(), RequiredAlign))
+      if (canReuseHostPtr(HostPtr.get(), RequiredAlign)) {
+        MUserPtr = HostPtr.get();
+      } else if (canReadHostPtr(HostPtr.get(), RequiredAlign)) {
         MUserPtr = HostPtr.get();
-      else {
+        MCreateShadowCopy = [this, RequiredAlign, HostPtr]() -> void {
+          setAlign(RequiredAlign);
+          MShadowCopy = allocateHostMem();
+          MUserPtr = MShadowCopy;
+          std::memcpy(MUserPtr, HostPtr.get(), MSizeInBytes);
+        };
+      } else {
         setAlign(RequiredAlign);
         MShadowCopy = allocateHostMem();
         MUserPtr = MShadowCopy;
@@ -248,6 +268,8 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
   static size_t getBufSizeForContext(const ContextImplPtr &Context,
                                      pi_native_handle MemObject);
 
+  void handleWriteAccessorCreation();
+
   void *allocateMem(ContextImplPtr Context, bool InitFromUserData,
                     void *HostPtr,
                     sycl::detail::pi::PiEvent &InteropEvent) override {
@@ -349,6 +371,10 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI {
   bool MIsInternal = false;
   // The number of graphs which are currently using this memory object.
   std::atomic<size_t> MGraphUseCount = 0;
+  // Function which creates a shadow copy of the host pointer. This is used to
+  // defer the memory allocation and copying to the point where a writable
+  // accessor is created.
+  std::function<void(void)> MCreateShadowCopy = []() -> void {};
   bool MOwnNativeHandle = true;
 };
 } // namespace detail

sycl/source/handler.cpp

@@ -527,6 +527,10 @@ void handler::associateWithHandlerCommon(detail::AccessorImplPtr AccImpl,
                           "are not allowed to be used in command graphs.");
   }
   detail::Requirement *Req = AccImpl.get();
+  if (Req->MAccessMode != sycl::access_mode::read) {
+    auto SYCLMemObj = static_cast<detail::SYCLMemObjT *>(Req->MSYCLMemObj);
+    SYCLMemObj->handleWriteAccessorCreation();
+  }
   // Add accessor to the list of requirements.
   if (Req->MAccessRange.size() != 0)
     CGData.MRequirements.push_back(Req);

sycl/test-e2e/Basic/host_defer_copy.cpp

@@ -0,0 +1,41 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <cstdlib>
+#include <iostream>
+#include <sycl/sycl.hpp>
+
+constexpr int N = 10 * 1024 * 1024;
+
+int main() {
+  std::vector<int> vec(N, 1);
+  const int *const host_address = &vec[0];
+  {
+    // Create a buffer with a read-only hostData pointer.
+    sycl::buffer<int, 1> buf(static_cast<const int *>(vec.data()),
+                             sycl::range<1>{N});
+
+    // Assert that the hostData pointer is being reused.
+    {
+      sycl::host_accessor<int, 1, sycl::access_mode::read> r_acc{buf};
+      assert(&r_acc[0] == host_address && "hostData was copied");
+    }
+
+    // Assert that creating a writeable accessor copies the data and the
+    // hostData pointer is not being reused.
+    {
+      sycl::host_accessor<int, 1, sycl::access_mode::write> rw_acc{buf};
+      assert(&rw_acc[0] != host_address &&
+             "writable accessor references read-only hostData");
+
+      rw_acc[0] = 0;
+      assert(rw_acc[0] == 0 && "failed to write to accessor");
+    }
+  }
+
+  // Assert that the vector was never modified (since hostData is read-only).
+  assert(vec[0] == 1 && "read-only hostData was modified");
+
+  std::cout << "Test passed!" << std::endl;
+  return EXIT_SUCCESS;
+}

sycl/test/abi/sycl_symbols_linux.dump

@@ -3810,6 +3810,7 @@ _ZN4sycl3_V16detail11SYCLMemObjT16determineHostPtrERKSt10shared_ptrINS1_12contex
 _ZN4sycl3_V16detail11SYCLMemObjT16updateHostMemoryEPv
 _ZN4sycl3_V16detail11SYCLMemObjT16updateHostMemoryEv
 _ZN4sycl3_V16detail11SYCLMemObjT20getBufSizeForContextERKSt10shared_ptrINS1_12context_implEEm
+_ZN4sycl3_V16detail11SYCLMemObjT27handleWriteAccessorCreationEv
 _ZN4sycl3_V16detail11SYCLMemObjTC1EmRKNS0_7contextEbNS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE
 _ZN4sycl3_V16detail11SYCLMemObjTC1EmRKNS0_7contextEbNS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE23_pi_image_channel_order22_pi_image_channel_typeNS0_5rangeILi3EEEjm
 _ZN4sycl3_V16detail11SYCLMemObjTC1EmRKNS0_7contextEmNS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE

sycl/test/abi/sycl_symbols_windows.dump

@@ -933,6 +933,7 @@
 ?begin_recording@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA_NAEAVqueue@67@@Z
 ?begin_recording@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA_NAEBV?$vector@Vqueue@_V1@sycl@@V?$allocator@Vqueue@_V1@sycl@@@std@@@std@@@Z
 ?build_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBV?$kernel_bundle@$0A@@23@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@5@AEBVproperty_list@23@@Z
+?canReadHostPtr@SYCLMemObjT@detail@_V1@sycl@@QEAA_NPEAX_K@Z
 ?canReuseHostPtr@SYCLMemObjT@detail@_V1@sycl@@QEAA_NPEAX_K@Z
 ?cancel_fusion@fusion_wrapper@experimental@codeplay@ext@_V1@sycl@@QEAAXXZ
 ?category@exception@_V1@sycl@@QEBAAEBVerror_category@std@@XZ
@@ -1289,6 +1290,7 @@
 ?handleHostData@SYCLMemObjT@detail@_V1@sycl@@QEAAXPEAX_K@Z
 ?handleHostData@SYCLMemObjT@detail@_V1@sycl@@QEAAXPEBX_K@Z
 ?handleRelease@buffer_plain@detail@_V1@sycl@@IEBAXXZ
+?handleWriteAccessorCreation@SYCLMemObjT@detail@_V1@sycl@@QEAAXXZ
 ?has@device@_V1@sycl@@QEBA_NW4aspect@23@@Z
 ?has@platform@_V1@sycl@@QEBA_NW4aspect@23@@Z
 ?hasUserDataPtr@SYCLMemObjT@detail@_V1@sycl@@UEBA_NXZ