diff --git a/sycl/source/detail/sycl_mem_obj_t.hpp b/sycl/source/detail/sycl_mem_obj_t.hpp index 4893d154a7566..60593231848a7 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) { @@ -184,10 +174,16 @@ class SYCLMemObjT : public SYCLMemObjI { } if (HostPtr) { - if (canReuseHostPtr(HostPtr, RequiredAlign)) { - MUserPtr = HostPtr; - } else if (canReadHostPtr(HostPtr, RequiredAlign)) { - MUserPtr = HostPtr; + // Pass the user pointer to UR unchanged. Each adapter is responsible + // for handling pointers it cannot use directly (misaligned, not + // importable, etc.) by allocating its own backing storage and copying. + MUserPtr = HostPtr; + + // For a read-only host pointer we still need a writable backing store + // if the user later creates a write accessor. Defer the allocation + // until that happens. This is adapter-independent: the language rule + // is that we may not write through a const user pointer. + if (MHostPtrReadOnly) { std::lock_guard Lock(MCreateShadowCopyMtx); MCreateShadowCopy = [this, RequiredAlign, HostPtr]() -> void { setAlign(RequiredAlign); @@ -195,11 +191,6 @@ class SYCLMemObjT : public SYCLMemObjI { MUserPtr = MShadowCopy; std::memcpy(MUserPtr, HostPtr, MSizeInBytes); }; - } else { - setAlign(RequiredAlign); - MShadowCopy = allocateHostMem(); - MUserPtr = MShadowCopy; - std::memcpy(MUserPtr, HostPtr, MSizeInBytes); } } } @@ -218,10 +209,9 @@ class SYCLMemObjT : public SYCLMemObjI { if (!MHostPtrReadOnly) set_final_data_from_storage(); - if (canReuseHostPtr(HostPtr.get(), RequiredAlign)) { - MUserPtr = HostPtr.get(); - } else if (canReadHostPtr(HostPtr.get(), RequiredAlign)) { - MUserPtr = HostPtr.get(); + MUserPtr = HostPtr.get(); + + if (MHostPtrReadOnly) { std::lock_guard Lock(MCreateShadowCopyMtx); MCreateShadowCopy = [this, RequiredAlign, HostPtr]() -> void { setAlign(RequiredAlign); @@ -229,11 +219,6 @@ class SYCLMemObjT : public SYCLMemObjI { MUserPtr = MShadowCopy; std::memcpy(MUserPtr, HostPtr.get(), MSizeInBytes); }; - } else { - setAlign(RequiredAlign); - MShadowCopy = allocateHostMem(); - MUserPtr = MShadowCopy; - std::memcpy(MUserPtr, HostPtr.get(), MSizeInBytes); } } } diff --git a/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp b/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp new file mode 100644 index 0000000000000..9177be4aab34d --- /dev/null +++ b/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp @@ -0,0 +1,261 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#include + +#include +#include +#include +#include +#include + +// Read-only kernel: sum all elements and return the result. +static int runReadOnlySumKernel(sycl::queue &Q, const int *HostPtr, size_t N) { + sycl::buffer Buf(HostPtr, sycl::range<1>(N), + {sycl::property::buffer::use_host_ptr{}}); + sycl::buffer SumBuf(1); + + Q.submit([&](sycl::handler &CGH) { + auto InAcc = Buf.get_access(CGH); + auto SumAcc = SumBuf.get_access(CGH); + CGH.single_task([=]() { + int Sum = 0; + for (size_t I = 0; I < N; ++I) + Sum += InAcc[I]; + SumAcc[0] = Sum; + }); + }); + Q.wait_and_throw(); + + auto SumHostAcc = SumBuf.get_host_access(); + return SumHostAcc[0]; +} + +// Writable kernel path; buffer destruction happens at scope exit. +static void runWriteKernel(sycl::queue &Q, int *HostPtr, size_t N) { + { + sycl::buffer Buf(HostPtr, sycl::range<1>(N), + {sycl::property::buffer::use_host_ptr{}}); + + Q.submit([&](sycl::handler &CGH) { + auto OutAcc = Buf.get_access(CGH); + CGH.single_task([=]() { + for (size_t I = 0; I < N; ++I) + OutAcc[I] = static_cast(I * 3 + 7); + }); + }); + Q.wait_and_throw(); + } +} + +// Verifies host-side result after writable-buffer destruction. +static bool checkExpectedPattern(const int *Ptr, size_t N) { + std::vector Tmp(N); + std::memcpy(Tmp.data(), Ptr, sizeof(int) * N); + for (size_t I = 0; I < N; ++I) { + if (Tmp[I] != static_cast(I * 3 + 7)) + return false; + } + return true; +} + +// Exercises the map/unmap path: a mid-lifetime host_accessor on a live buffer. +// This is the path that, on adapters keeping a separate working buffer (e.g. +// L0v2 integrated with a non-importable pointer), must keep the user pointer +// and the device-visible buffer in sync in BOTH directions. The buffer is +// local to this function, so it is destroyed on return; the caller then checks +// the final write-back at HostPtr. +// +// Sequence: kernel writes I -> host reads (map READ) and verifies -> host adds +// 100 (unmap WRITE-back) -> kernel adds 1 -> scope exit writes I+101 to host. +static bool runMidLifeHostAccessor(sycl::queue &Q, int *HostPtr, size_t N) { + sycl::buffer Buf(HostPtr, sycl::range<1>(N), + {sycl::property::buffer::use_host_ptr{}}); + + Q.submit([&](sycl::handler &CGH) { + auto Acc = Buf.get_access(CGH); + CGH.single_task([=]() { + for (size_t I = 0; I < N; ++I) + Acc[I] = static_cast(I); + }); + }); + Q.wait_and_throw(); + + // Mid-lifetime host access: map READ must observe the kernel's writes, and + // the modification must survive back to the device on unmap WRITE. + { + sycl::host_accessor HAcc(Buf, sycl::read_write); + for (size_t I = 0; I < N; ++I) { + if (HAcc[I] != static_cast(I)) + return false; + HAcc[I] += 100; + } + } + + Q.submit([&](sycl::handler &CGH) { + auto Acc = Buf.get_access(CGH); + CGH.single_task([=]() { + for (size_t I = 0; I < N; ++I) + Acc[I] += 1; + }); + }); + Q.wait_and_throw(); + + return true; +} + +// A read-only buffer must never modify the user's source data, regardless of +// whether the backend aliases the pointer (zero-copy) or copies it. Portable +// across all backends because a read-only access performs no writes. +static bool runReadOnlyImmutability(sycl::queue &Q, const int *HostPtr, + size_t N) { + std::vector Orig(N); + std::memcpy(Orig.data(), HostPtr, sizeof(int) * N); + + { + sycl::buffer Buf(HostPtr, sycl::range<1>(N), + {sycl::property::buffer::use_host_ptr{}}); + sycl::buffer SumBuf(1); + Q.submit([&](sycl::handler &CGH) { + auto In = Buf.get_access(CGH); + auto S = SumBuf.get_access(CGH); + CGH.single_task([=]() { + int Sum = 0; + for (size_t I = 0; I < N; ++I) + Sum += In[I]; + S[0] = Sum; + }); + }); + Q.wait_and_throw(); + } + + return std::memcmp(Orig.data(), HostPtr, sizeof(int) * N) == 0; +} + +int main() { + constexpr size_t N = 32; + sycl::queue Q; + + // Build aligned reference data. + std::vector AlignedInput(N); + for (size_t I = 0; I < N; ++I) + AlignedInput[I] = static_cast(I); + + const int ExpectedSum = static_cast((N - 1) * N / 2); + + auto checkMidLife = [](const int *Ptr) { + std::vector Tmp(N); + std::memcpy(Tmp.data(), Ptr, sizeof(int) * N); + for (size_t I = 0; I < N; ++I) + if (Tmp[I] != static_cast(I + 101)) + return false; + return true; + }; + + // --- Aligned baseline --- + if (runReadOnlySumKernel(Q, AlignedInput.data(), N) != ExpectedSum) { + std::cerr << "Unexpected aligned sum\n"; + return 1; + } + + std::vector AlignedWritable(N, 0); + runWriteKernel(Q, AlignedWritable.data(), N); + if (!checkExpectedPattern(AlignedWritable.data(), N)) { + std::cerr << "Unexpected data in aligned writable buffer\n"; + return 1; + } + + if (!runReadOnlyImmutability(Q, AlignedInput.data(), N)) { + std::cerr << "Read-only buffer modified aligned source data\n"; + return 1; + } + + std::vector AlignedMid(N, 0); + if (!runMidLifeHostAccessor(Q, AlignedMid.data(), N) || + !checkMidLife(AlignedMid.data())) { + std::cerr << "Mid-life host_accessor failed on aligned buffer\n"; + return 1; + } + + // --- Misaligned variants --- + // Cover several offsets. 1 byte is not even int-aligned; 4 bytes is + // int-aligned but typically below a backend's required base-address + // alignment (e.g. CL_DEVICE_MEM_BASE_ADDR_ALIGN is normally 64-128 B); + // 64 bytes exercises the boundary case where some backends still require + // a stricter (e.g. 128 B / page) alignment for zero-copy import. + constexpr size_t Offsets[] = {1, 4, 64}; + for (size_t Offset : Offsets) { + std::vector ROStorage(sizeof(int) * N + Offset); + int *UnalignedPtr = reinterpret_cast(ROStorage.data() + Offset); + std::memcpy(UnalignedPtr, AlignedInput.data(), sizeof(int) * N); + const int *ReadOnlyUnalignedPtr = UnalignedPtr; + + if (runReadOnlySumKernel(Q, ReadOnlyUnalignedPtr, N) != ExpectedSum) { + std::cerr << "Unexpected misaligned sum (offset=" << Offset << ")\n"; + return 1; + } + + std::vector WritableStorage(sizeof(int) * N + Offset, 0); + int *UnalignedWritablePtr = + reinterpret_cast(WritableStorage.data() + Offset); + runWriteKernel(Q, UnalignedWritablePtr, N); + if (!checkExpectedPattern(UnalignedWritablePtr, N)) { + std::cerr << "Unexpected data in misaligned writable buffer (offset=" + << Offset << ")\n"; + return 1; + } + + if (!runReadOnlyImmutability(Q, ReadOnlyUnalignedPtr, N)) { + std::cerr << "Read-only buffer modified misaligned source data (offset=" + << Offset << ")\n"; + return 1; + } + + std::vector MidStorage(sizeof(int) * N + Offset, 0); + int *UnalignedMidPtr = reinterpret_cast(MidStorage.data() + Offset); + if (!runMidLifeHostAccessor(Q, UnalignedMidPtr, N) || + !checkMidLife(UnalignedMidPtr)) { + std::cerr << "Mid-life host_accessor failed on misaligned buffer (offset=" + << Offset << ")\n"; + return 1; + } + } + + // --- Aligned-but-non-importable: raw heap allocation --- + // A plain new[] returns a pointer that is naturally aligned for int but + // is not USM-imported, not pinned, and not part of any device-visible + // allocation. On L0 this exercises the path where maybeImportUSM either + // succeeds (and the pointer is promoted) or fails and the adapter must + // fall back to its own backing storage with explicit copies. + { + std::unique_ptr HeapInput(new int[N]); + for (size_t I = 0; I < N; ++I) + HeapInput[I] = static_cast(I); + + if (runReadOnlySumKernel(Q, HeapInput.get(), N) != ExpectedSum) { + std::cerr << "Unexpected sum on aligned heap pointer\n"; + return 1; + } + + std::unique_ptr HeapWritable(new int[N]()); + runWriteKernel(Q, HeapWritable.get(), N); + if (!checkExpectedPattern(HeapWritable.get(), N)) { + std::cerr << "Unexpected data in aligned heap writable buffer\n"; + return 1; + } + + if (!runReadOnlyImmutability(Q, HeapInput.get(), N)) { + std::cerr << "Read-only buffer modified aligned heap source data\n"; + return 1; + } + + std::unique_ptr HeapMid(new int[N]()); + if (!runMidLifeHostAccessor(Q, HeapMid.get(), N) || + !checkMidLife(HeapMid.get())) { + std::cerr << "Mid-life host_accessor failed on aligned heap buffer\n"; + return 1; + } + } + + return 0; +} diff --git a/unified-runtime/source/adapters/level_zero/common.hpp b/unified-runtime/source/adapters/level_zero/common.hpp index af5976fa8aa53..b70c4fb2bf00a 100644 --- a/unified-runtime/source/adapters/level_zero/common.hpp +++ b/unified-runtime/source/adapters/level_zero/common.hpp @@ -328,8 +328,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); }; diff --git a/unified-runtime/source/adapters/level_zero/device.cpp b/unified-runtime/source/adapters/level_zero/device.cpp index 638df58eea1c3..49e690af3b5f5 100644 --- a/unified-runtime/source/adapters/level_zero/device.cpp +++ b/unified-runtime/source/adapters/level_zero/device.cpp @@ -2388,10 +2388,10 @@ 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/helpers/memory_helpers.cpp b/unified-runtime/source/adapters/level_zero/helpers/memory_helpers.cpp index 51a75bbd5fc75..44f5e792b638a 100644 --- a/unified-runtime/source/adapters/level_zero/helpers/memory_helpers.cpp +++ b/unified-runtime/source/adapters/level_zero/helpers/memory_helpers.cpp @@ -31,9 +31,11 @@ bool maybeImportUSM(ze_driver_handle_t hTranslatedDriver, auto ret = getMemoryAttrs(hContext, ptr, nullptr, &properties); 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; + // Promote the host ptr to USM host memory. The driver-level import can + // fail (e.g., for a misaligned ptr), so report success to callers only + // when the underlying L0 call actually succeeded. + return ZeUSMImport.doZeUSMImport(hTranslatedDriver, ptr, size) == + ZE_RESULT_SUCCESS; } return false; }