From 2672818bd43482ee2235583b1b02788ec77623bb Mon Sep 17 00:00:00 2001 From: kweronsx Date: Wed, 26 Aug 2026 14:37:23 +0200 Subject: [PATCH] [SYCL][UR] Avoid redundant copy of misaligned pointers Add regression test --- sycl/source/detail/sycl_mem_obj_t.hpp | 61 ++++++------------- .../misaligned_pointer_handling.cpp | 40 ++++++++++++ .../adapters/level_zero/common/device.cpp | 9 +-- .../common/helpers/memory_helpers.cpp | 4 +- .../common/helpers/shared_helpers.hpp | 4 +- 5 files changed, 69 insertions(+), 49 deletions(-) create mode 100644 sycl/test-e2e/Regression/misaligned_pointer_handling.cpp diff --git a/sycl/source/detail/sycl_mem_obj_t.hpp b/sycl/source/detail/sycl_mem_obj_t.hpp index 2ebb814ff4a97..a54ffa352bc4c 100644 --- a/sycl/source/detail/sycl_mem_obj_t.hpp +++ b/sycl/source/detail/sycl_mem_obj_t.hpp @@ -165,16 +165,6 @@ class SYCLMemObjT : public SYCLMemObjI { has_property(); } - bool canReadHostPtr(void *HostPtr, const size_t RequiredAlign) { - bool Aligned = - (reinterpret_cast(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 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 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 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 Lock(MCreateShadowCopyMtx); + MCreateShadowCopy = [this, RequiredAlign, HostPtr]() -> void { setAlign(RequiredAlign); MShadowCopy = allocateHostMem(); MUserPtr = MShadowCopy; std::memcpy(MUserPtr, HostPtr.get(), MSizeInBytes); - } + }; } } diff --git a/sycl/test-e2e/Regression/misaligned_pointer_handling.cpp b/sycl/test-e2e/Regression/misaligned_pointer_handling.cpp new file mode 100644 index 0000000000000..7812151f709ac --- /dev/null +++ b/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 +#include +#include +#include + +using data_type_t = uint32_t; + +void overflow(data_type_t *data) { + sycl::buffer 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); +} diff --git a/unified-runtime/source/adapters/level_zero/common/device.cpp b/unified-runtime/source/adapters/level_zero/common/device.cpp index 2a5941534d374..697c836a898af 100644 --- a/unified-runtime/source/adapters/level_zero/common/device.cpp +++ b/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) { diff --git a/unified-runtime/source/adapters/level_zero/common/helpers/memory_helpers.cpp b/unified-runtime/source/adapters/level_zero/common/helpers/memory_helpers.cpp index 2ac58f9591d3f..a8197e6259ae9 100644 --- a/unified-runtime/source/adapters/level_zero/common/helpers/memory_helpers.cpp +++ b/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; } diff --git a/unified-runtime/source/adapters/level_zero/common/helpers/shared_helpers.hpp b/unified-runtime/source/adapters/level_zero/common/helpers/shared_helpers.hpp index 447719f9a08ce..557042d66f7c0 100644 --- a/unified-runtime/source/adapters/level_zero/common/helpers/shared_helpers.hpp +++ b/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); };