[SYCL][UR] Avoid redundant copy of misaligned pointers

sycl/source/detail/sycl_mem_obj_t.hpp

@@ -165,16 +165,6 @@ class SYCLMemObjT : public SYCLMemObjI {
            has_property<property::image::use_host_ptr>();
   }
 
-  bool canReadHostPtr(void *HostPtr, const size_t RequiredAlign) {
-    bool Aligned =
-        (reinterpret_cast<std::uintptr_t>(HostPtr) % RequiredAlign) == 0;
-    return Aligned || useHostPtr();
-  }
-
-  bool canReuseHostPtr(void *HostPtr, const size_t RequiredAlign) {
-    return !MHostPtrReadOnly && canReadHostPtr(HostPtr, RequiredAlign);
-  }
-
   void handleHostData(void *HostPtr, const size_t RequiredAlign) {
     MHostPtrProvided = true;
     if (!MHostPtrReadOnly && HostPtr) {
@@ -183,24 +173,18 @@ class SYCLMemObjT : public SYCLMemObjI {
       });
     }
 
-    if (HostPtr) {
-      if (canReuseHostPtr(HostPtr, RequiredAlign)) {
-        MUserPtr = HostPtr;
-      } else if (canReadHostPtr(HostPtr, RequiredAlign)) {
-        MUserPtr = HostPtr;
-        std::lock_guard<std::mutex> Lock(MCreateShadowCopyMtx);
-        MCreateShadowCopy = [this, RequiredAlign, HostPtr]() -> void {
-          setAlign(RequiredAlign);
-          MShadowCopy = allocateHostMem();
-          MUserPtr = MShadowCopy;
-          std::memcpy(MUserPtr, HostPtr, MSizeInBytes);
-        };
-      } else {
+    if (!HostPtr)
+      return;
+
+    MUserPtr = HostPtr;
+    if (MHostPtrReadOnly) {
+      std::lock_guard<std::mutex> Lock(MCreateShadowCopyMtx);
+      MCreateShadowCopy = [this, RequiredAlign, HostPtr]() -> void {
         setAlign(RequiredAlign);
         MShadowCopy = allocateHostMem();
         MUserPtr = MShadowCopy;
         std::memcpy(MUserPtr, HostPtr, MSizeInBytes);
-      }
+      };
     }
   }
 
@@ -214,27 +198,22 @@ class SYCLMemObjT : public SYCLMemObjI {
     MHostPtrProvided = true;
     MSharedPtrStorage = HostPtr;
     MHostPtrReadOnly = IsConstPtr;
-    if (HostPtr) {
-      if (!MHostPtrReadOnly)
-        set_final_data_from_storage();
-
-      if (canReuseHostPtr(HostPtr.get(), RequiredAlign)) {
-        MUserPtr = HostPtr.get();
-      } else if (canReadHostPtr(HostPtr.get(), RequiredAlign)) {
-        MUserPtr = HostPtr.get();
-        std::lock_guard<std::mutex> Lock(MCreateShadowCopyMtx);
-        MCreateShadowCopy = [this, RequiredAlign, HostPtr]() -> void {
-          setAlign(RequiredAlign);
-          MShadowCopy = allocateHostMem();
-          MUserPtr = MShadowCopy;
-          std::memcpy(MUserPtr, HostPtr.get(), MSizeInBytes);
-        };
-      } else {
+
+    if (!HostPtr)
+      return;
+
+    if (!MHostPtrReadOnly)
+      set_final_data_from_storage();
+
+    MUserPtr = HostPtr.get();
+    if (MHostPtrReadOnly) {
+      std::lock_guard<std::mutex> Lock(MCreateShadowCopyMtx);
+      MCreateShadowCopy = [this, RequiredAlign, HostPtr]() -> void {
         setAlign(RequiredAlign);
         MShadowCopy = allocateHostMem();
         MUserPtr = MShadowCopy;
         std::memcpy(MUserPtr, HostPtr.get(), MSizeInBytes);
-      }
+      };
     }
   }
 

sycl/test-e2e/Regression/misaligned_pointer_handling.cpp

@@ -0,0 +1,40 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+//===----------------- misaligned_pointer_handling.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 <cassert>
+#include <cstring>
+#include <stdint.h>
+#include <sycl/detail/core.hpp>
+
+using data_type_t = uint32_t;
+
+void overflow(data_type_t *data) {
+  sycl::buffer<data_type_t, 1> b{data, 1};
+  constexpr data_type_t value = 0xff'ff'ff'ff;
+  sycl::queue q;
+
+  q.submit([&b](sycl::handler &h) {
+    sycl::accessor a{b, h, sycl::read_write};
+    h.parallel_for(sycl::range<1>{1}, [=](auto i) { a[i] += value; });
+  }).wait();
+}
+
+int main() {
+  data_type_t anyData[] = {1, 2};
+  data_type_t before{}, after{};
+  auto unaligned = (data_type_t *)(((uint8_t *)anyData) + 1);
+
+  std::memcpy(&before, unaligned, sizeof(before));
+  overflow(unaligned);
+  std::memcpy(&after, unaligned, sizeof(after));
+
+  assert(after == before - 1);
+}

unified-runtime/source/adapters/level_zero/common/device.cpp

@@ -2292,10 +2292,11 @@ void ZeUSMImportExtension::setZeUSMImport(ur_platform_handle_t_ *Platform) {
     setEnvVar("SYCL_HOST_UNIFIED_MEMORY", "1");
   }
 }
-void ZeUSMImportExtension::doZeUSMImport(ze_driver_handle_t DriverHandle,
-                                         void *HostPtr, size_t Size) {
-  ZE_CALL_NOCHECK(zexDriverImportExternalPointer,
-                  (DriverHandle, HostPtr, Size));
+
+ze_result_t ZeUSMImportExtension::doZeUSMImport(ze_driver_handle_t DriverHandle,
+                                                void *HostPtr, size_t Size) {
+  return ZE_CALL_NOCHECK(zexDriverImportExternalPointer,
+                         (DriverHandle, HostPtr, Size));
 }
 void ZeUSMImportExtension::doZeUSMRelease(ze_driver_handle_t DriverHandle,
                                           void *HostPtr) {

unified-runtime/source/adapters/level_zero/common/helpers/memory_helpers.cpp

@@ -33,8 +33,8 @@ bool maybeImportUSM(ze_driver_handle_t hTranslatedDriver,
 
   if (ret == UR_RESULT_SUCCESS && properties.type == ZE_MEMORY_TYPE_UNKNOWN) {
     // Promote the host ptr to USM host memory
-    ZeUSMImport.doZeUSMImport(hTranslatedDriver, ptr, size);
-    return true;
+    return ZeUSMImport.doZeUSMImport(hTranslatedDriver, ptr, size) ==
+           ZE_RESULT_SUCCESS;
   }
   return false;
 }

unified-runtime/source/adapters/level_zero/common/helpers/shared_helpers.hpp

@@ -283,8 +283,8 @@ class ZeUSMImportExtension {
   ZeUSMImportExtension() : Supported{false}, Enabled{false} {}
 
   void setZeUSMImport(ur_platform_handle_t_ *Platform);
-  void doZeUSMImport(ze_driver_handle_t DriverHandle, void *HostPtr,
-                     size_t Size);
+  ze_result_t doZeUSMImport(ze_driver_handle_t DriverHandle, void *HostPtr,
+                            size_t Size);
   void doZeUSMRelease(ze_driver_handle_t DriverHandle, void *HostPtr);
 };