From 2d490cfd48c1e20b1737aafca454ef08136f8dc5 Mon Sep 17 00:00:00 2001 From: "Mateusz P. Nowak" Date: Tue, 26 May 2026 16:15:51 +0000 Subject: [PATCH 1/6] Stop pre-aligning host pointers in SYCL; verify L0 USM import --- sycl/source/detail/sycl_mem_obj_t.hpp | 41 ++++++------------- .../level_zero/helpers/memory_helpers.cpp | 8 +++- 2 files changed, 20 insertions(+), 29 deletions(-) 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/unified-runtime/source/adapters/level_zero/helpers/memory_helpers.cpp b/unified-runtime/source/adapters/level_zero/helpers/memory_helpers.cpp index e9b037fec6037..523aa4652c90d 100644 --- a/unified-runtime/source/adapters/level_zero/helpers/memory_helpers.cpp +++ b/unified-runtime/source/adapters/level_zero/helpers/memory_helpers.cpp @@ -34,7 +34,13 @@ 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; + + // doZeUSMImport silently ignores driver-level failures (e.g., misaligned + // ptr), so re-query to confirm the import actually succeeded before + // reporting it to callers. + ret = getMemoryAttrs(hContext, ptr, nullptr, &properties); + return ret == UR_RESULT_SUCCESS && + properties.type != ZE_MEMORY_TYPE_UNKNOWN; } return false; } From a1ac8a736028a54f271f7f03edc2b425bde211a5 Mon Sep 17 00:00:00 2001 From: "Mateusz P. Nowak" Date: Fri, 29 May 2026 11:55:27 +0000 Subject: [PATCH 2/6] Handle misaligned host pointers in opencl's urMemBufferCreate --- .../source/adapters/opencl/memory.cpp | 57 +++++++++++++++++-- 1 file changed, 51 insertions(+), 6 deletions(-) diff --git a/unified-runtime/source/adapters/opencl/memory.cpp b/unified-runtime/source/adapters/opencl/memory.cpp index 3507c003db0b0..9d6318756d695 100644 --- a/unified-runtime/source/adapters/opencl/memory.cpp +++ b/unified-runtime/source/adapters/opencl/memory.cpp @@ -336,10 +336,56 @@ ur_result_t ur_mem_handle_t_::makeWithNative(native_type NativeMem, return UR_RESULT_SUCCESS; } +static bool canUseHostPtrDirectly(ur_context_handle_t hContext, void *HostPtr) { + if (!HostPtr) + return false; + + cl_uint MaxAlignBits = 0; + bool AllUnifiedMem = true; + + for (uint32_t I = 0; I < hContext->DeviceCount; ++I) { + cl_uint AlignBits = 0; + clGetDeviceInfo(hContext->Devices[I]->CLDevice, + CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof(AlignBits), + &AlignBits, nullptr); + if (AlignBits > MaxAlignBits) + MaxAlignBits = AlignBits; + + cl_bool Unified = CL_FALSE; + clGetDeviceInfo(hContext->Devices[I]->CLDevice, + CL_DEVICE_HOST_UNIFIED_MEMORY, sizeof(Unified), &Unified, + nullptr); + if (!Unified) + AllUnifiedMem = false; + } + + size_t RequiredAlign = MaxAlignBits / 8; + bool IsAligned = + RequiredAlign == 0 || + (reinterpret_cast(HostPtr) % RequiredAlign) == 0; + + return IsAligned && AllUnifiedMem; +} + UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreate( ur_context_handle_t hContext, ur_mem_flags_t flags, size_t size, const ur_buffer_properties_t *pProperties, ur_mem_handle_t *phBuffer) { cl_int RetErr = CL_INVALID_OPERATION; + + void *HostPtr = pProperties ? pProperties->pHost : nullptr; + cl_mem_flags CLFlags = convertURMemFlagsToCL(flags); + + // CL_MEM_USE_HOST_PTR requires alignment to CL_DEVICE_MEM_BASE_ADDR_ALIGN + // and unified host memory. When either is missing, fall back to + // CL_MEM_COPY_HOST_PTR so the driver allocates its own aligned storage. + // Write-back to the user pointer is the responsibility of the caller + // (e.g. SYCL's scheduler via urEnqueueMemBufferRead). + if (HostPtr && (CLFlags & CL_MEM_USE_HOST_PTR)) { + if (!canUseHostPtrDirectly(hContext, HostPtr)) { + CLFlags = (CLFlags & ~CL_MEM_USE_HOST_PTR) | CL_MEM_COPY_HOST_PTR; + } + } + if (pProperties) { // TODO: need to check if all properties are supported by OpenCL RT and // ignore unsupported @@ -377,9 +423,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreate( PropertiesIntel.push_back(0); try { - cl_mem Buffer = FuncPtr( - CLContext, PropertiesIntel.data(), static_cast(flags), - size, pProperties->pHost, static_cast(&RetErr)); + cl_mem Buffer = + FuncPtr(CLContext, PropertiesIntel.data(), CLFlags, size, HostPtr, + static_cast(&RetErr)); CL_RETURN_ON_FAILURE(RetErr); auto URMem = std::make_unique(Buffer, hContext); *phBuffer = URMem.release(); @@ -392,11 +438,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreate( } } - void *HostPtr = pProperties ? pProperties->pHost : nullptr; try { cl_mem Buffer = - clCreateBuffer(hContext->CLContext, static_cast(flags), - size, HostPtr, static_cast(&RetErr)); + clCreateBuffer(hContext->CLContext, CLFlags, size, HostPtr, + static_cast(&RetErr)); CL_RETURN_ON_FAILURE(RetErr); auto URMem = std::make_unique(Buffer, hContext); *phBuffer = URMem.release(); From 4a38280dbc4ee8a4abd33fc17c58d10a7e7120cc Mon Sep 17 00:00:00 2001 From: "Mateusz P. Nowak" Date: Thu, 18 Jun 2026 10:59:53 +0000 Subject: [PATCH 3/6] add test --- .../buffer_shadow_copy_platform_policy.cpp | 222 ++++++++++++++++++ .../source/adapters/opencl/memory.cpp | 15 +- 2 files changed, 228 insertions(+), 9 deletions(-) create mode 100644 sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp 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..d1595c3d5eaa2 --- /dev/null +++ b/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp @@ -0,0 +1,222 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#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::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)); + + 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)); + + 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::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); + + // Build a deliberately misaligned copy: offset by 1 byte so that the int* + // is not naturally aligned. + std::vector Storage(sizeof(int) * N + 1); + int *UnalignedPtr = reinterpret_cast(Storage.data() + 1); + std::memcpy(UnalignedPtr, AlignedInput.data(), sizeof(int) * N); + const int *ReadOnlyUnalignedPtr = UnalignedPtr; + + const int ExpectedSum = static_cast((N - 1) * N / 2); + + // --- Read path correctness --- + // Both aligned and misaligned host pointers must produce the correct sum. + const int AlignedSum = runReadOnlySumKernel(Q, AlignedInput.data(), N); + if (AlignedSum != ExpectedSum) { + std::cerr << "Unexpected aligned sum: " << AlignedSum << "\n"; + return 1; + } + + const int MisalignedSum = runReadOnlySumKernel(Q, ReadOnlyUnalignedPtr, N); + if (MisalignedSum != ExpectedSum) { + std::cerr << "Unexpected misaligned sum: " << MisalignedSum << "\n"; + return 1; + } + + // --- Write-back correctness --- + // After the buffer goes out of scope the kernel-written pattern must be + // visible at the original host pointer, even when that pointer is misaligned. + std::vector AlignedWritable(N, 0); + std::vector WritableStorage(sizeof(int) * N + 1, 0); + int *UnalignedWritablePtr = + reinterpret_cast(WritableStorage.data() + 1); + + runWriteKernel(Q, AlignedWritable.data(), N); + runWriteKernel(Q, UnalignedWritablePtr, N); + + if (!checkExpectedPattern(AlignedWritable.data(), N)) { + std::cerr << "Unexpected data in aligned writable buffer\n"; + return 1; + } + if (!checkExpectedPattern(UnalignedWritablePtr, N)) { + std::cerr << "Unexpected data in misaligned writable buffer\n"; + return 1; + } + + // --- Read-only immutability --- + // A read-only buffer must leave the user's source untouched, aligned or not. + if (!runReadOnlyImmutability(Q, AlignedInput.data(), N)) { + std::cerr << "Read-only buffer modified aligned source data\n"; + return 1; + } + if (!runReadOnlyImmutability(Q, ReadOnlyUnalignedPtr, N)) { + std::cerr << "Read-only buffer modified misaligned source data\n"; + return 1; + } + + // --- Mid-lifetime host_accessor (map/unmap) + final write-back --- + // Exercises the bidirectional map/unmap sync path and the final copy-back. + // Expected final pattern at the host pointer: I + 101 (kernel wrote I, host + // added 100 via map/unmap, kernel added 1). + 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; + }; + + 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; + } + + std::vector MidStorage(sizeof(int) * N + 1, 0); + int *UnalignedMidPtr = reinterpret_cast(MidStorage.data() + 1); + if (!runMidLifeHostAccessor(Q, UnalignedMidPtr, N) || + !checkMidLife(UnalignedMidPtr)) { + std::cerr << "Mid-life host_accessor failed on misaligned buffer\n"; + return 1; + } + + return 0; +} diff --git a/unified-runtime/source/adapters/opencl/memory.cpp b/unified-runtime/source/adapters/opencl/memory.cpp index 9d6318756d695..8532805fa9ba4 100644 --- a/unified-runtime/source/adapters/opencl/memory.cpp +++ b/unified-runtime/source/adapters/opencl/memory.cpp @@ -360,9 +360,8 @@ static bool canUseHostPtrDirectly(ur_context_handle_t hContext, void *HostPtr) { } size_t RequiredAlign = MaxAlignBits / 8; - bool IsAligned = - RequiredAlign == 0 || - (reinterpret_cast(HostPtr) % RequiredAlign) == 0; + bool IsAligned = RequiredAlign == 0 || + (reinterpret_cast(HostPtr) % RequiredAlign) == 0; return IsAligned && AllUnifiedMem; } @@ -423,9 +422,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreate( PropertiesIntel.push_back(0); try { - cl_mem Buffer = - FuncPtr(CLContext, PropertiesIntel.data(), CLFlags, size, HostPtr, - static_cast(&RetErr)); + cl_mem Buffer = FuncPtr(CLContext, PropertiesIntel.data(), CLFlags, + size, HostPtr, static_cast(&RetErr)); CL_RETURN_ON_FAILURE(RetErr); auto URMem = std::make_unique(Buffer, hContext); *phBuffer = URMem.release(); @@ -439,9 +437,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreate( } try { - cl_mem Buffer = - clCreateBuffer(hContext->CLContext, CLFlags, size, HostPtr, - static_cast(&RetErr)); + cl_mem Buffer = clCreateBuffer(hContext->CLContext, CLFlags, size, HostPtr, + static_cast(&RetErr)); CL_RETURN_ON_FAILURE(RetErr); auto URMem = std::make_unique(Buffer, hContext); *phBuffer = URMem.release(); From ae3466c7e66113167a6b8872cf7ebad210a64cb8 Mon Sep 17 00:00:00 2001 From: "Mateusz P. Nowak" Date: Mon, 22 Jun 2026 11:57:26 +0000 Subject: [PATCH 4/6] fix opencl path --- .../source/adapters/opencl/memory.cpp | 56 +++---------------- 1 file changed, 7 insertions(+), 49 deletions(-) diff --git a/unified-runtime/source/adapters/opencl/memory.cpp b/unified-runtime/source/adapters/opencl/memory.cpp index 8532805fa9ba4..3507c003db0b0 100644 --- a/unified-runtime/source/adapters/opencl/memory.cpp +++ b/unified-runtime/source/adapters/opencl/memory.cpp @@ -336,55 +336,10 @@ ur_result_t ur_mem_handle_t_::makeWithNative(native_type NativeMem, return UR_RESULT_SUCCESS; } -static bool canUseHostPtrDirectly(ur_context_handle_t hContext, void *HostPtr) { - if (!HostPtr) - return false; - - cl_uint MaxAlignBits = 0; - bool AllUnifiedMem = true; - - for (uint32_t I = 0; I < hContext->DeviceCount; ++I) { - cl_uint AlignBits = 0; - clGetDeviceInfo(hContext->Devices[I]->CLDevice, - CL_DEVICE_MEM_BASE_ADDR_ALIGN, sizeof(AlignBits), - &AlignBits, nullptr); - if (AlignBits > MaxAlignBits) - MaxAlignBits = AlignBits; - - cl_bool Unified = CL_FALSE; - clGetDeviceInfo(hContext->Devices[I]->CLDevice, - CL_DEVICE_HOST_UNIFIED_MEMORY, sizeof(Unified), &Unified, - nullptr); - if (!Unified) - AllUnifiedMem = false; - } - - size_t RequiredAlign = MaxAlignBits / 8; - bool IsAligned = RequiredAlign == 0 || - (reinterpret_cast(HostPtr) % RequiredAlign) == 0; - - return IsAligned && AllUnifiedMem; -} - UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreate( ur_context_handle_t hContext, ur_mem_flags_t flags, size_t size, const ur_buffer_properties_t *pProperties, ur_mem_handle_t *phBuffer) { cl_int RetErr = CL_INVALID_OPERATION; - - void *HostPtr = pProperties ? pProperties->pHost : nullptr; - cl_mem_flags CLFlags = convertURMemFlagsToCL(flags); - - // CL_MEM_USE_HOST_PTR requires alignment to CL_DEVICE_MEM_BASE_ADDR_ALIGN - // and unified host memory. When either is missing, fall back to - // CL_MEM_COPY_HOST_PTR so the driver allocates its own aligned storage. - // Write-back to the user pointer is the responsibility of the caller - // (e.g. SYCL's scheduler via urEnqueueMemBufferRead). - if (HostPtr && (CLFlags & CL_MEM_USE_HOST_PTR)) { - if (!canUseHostPtrDirectly(hContext, HostPtr)) { - CLFlags = (CLFlags & ~CL_MEM_USE_HOST_PTR) | CL_MEM_COPY_HOST_PTR; - } - } - if (pProperties) { // TODO: need to check if all properties are supported by OpenCL RT and // ignore unsupported @@ -422,8 +377,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreate( PropertiesIntel.push_back(0); try { - cl_mem Buffer = FuncPtr(CLContext, PropertiesIntel.data(), CLFlags, - size, HostPtr, static_cast(&RetErr)); + cl_mem Buffer = FuncPtr( + CLContext, PropertiesIntel.data(), static_cast(flags), + size, pProperties->pHost, static_cast(&RetErr)); CL_RETURN_ON_FAILURE(RetErr); auto URMem = std::make_unique(Buffer, hContext); *phBuffer = URMem.release(); @@ -436,9 +392,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urMemBufferCreate( } } + void *HostPtr = pProperties ? pProperties->pHost : nullptr; try { - cl_mem Buffer = clCreateBuffer(hContext->CLContext, CLFlags, size, HostPtr, - static_cast(&RetErr)); + cl_mem Buffer = + clCreateBuffer(hContext->CLContext, static_cast(flags), + size, HostPtr, static_cast(&RetErr)); CL_RETURN_ON_FAILURE(RetErr); auto URMem = std::make_unique(Buffer, hContext); *phBuffer = URMem.release(); From c9e980e65aedbb67682f885ee45fd8dec1886643 Mon Sep 17 00:00:00 2001 From: "Mateusz P. Nowak" Date: Tue, 23 Jun 2026 13:30:31 +0000 Subject: [PATCH 5/6] update the test --- .../buffer_shadow_copy_platform_policy.cpp | 145 +++++++++++------- 1 file changed, 90 insertions(+), 55 deletions(-) diff --git a/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp b/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp index d1595c3d5eaa2..b02d126b710b6 100644 --- a/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp +++ b/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp @@ -6,6 +6,7 @@ #include #include #include +#include #include // Read-only kernel: sum all elements and return the result. @@ -136,72 +137,34 @@ int main() { for (size_t I = 0; I < N; ++I) AlignedInput[I] = static_cast(I); - // Build a deliberately misaligned copy: offset by 1 byte so that the int* - // is not naturally aligned. - std::vector Storage(sizeof(int) * N + 1); - int *UnalignedPtr = reinterpret_cast(Storage.data() + 1); - std::memcpy(UnalignedPtr, AlignedInput.data(), sizeof(int) * N); - const int *ReadOnlyUnalignedPtr = UnalignedPtr; - const int ExpectedSum = static_cast((N - 1) * N / 2); - // --- Read path correctness --- - // Both aligned and misaligned host pointers must produce the correct sum. - const int AlignedSum = runReadOnlySumKernel(Q, AlignedInput.data(), N); - if (AlignedSum != ExpectedSum) { - std::cerr << "Unexpected aligned sum: " << AlignedSum << "\n"; - return 1; - } + 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; + }; - const int MisalignedSum = runReadOnlySumKernel(Q, ReadOnlyUnalignedPtr, N); - if (MisalignedSum != ExpectedSum) { - std::cerr << "Unexpected misaligned sum: " << MisalignedSum << "\n"; + // --- Aligned baseline --- + if (runReadOnlySumKernel(Q, AlignedInput.data(), N) != ExpectedSum) { + std::cerr << "Unexpected aligned sum\n"; return 1; } - // --- Write-back correctness --- - // After the buffer goes out of scope the kernel-written pattern must be - // visible at the original host pointer, even when that pointer is misaligned. std::vector AlignedWritable(N, 0); - std::vector WritableStorage(sizeof(int) * N + 1, 0); - int *UnalignedWritablePtr = - reinterpret_cast(WritableStorage.data() + 1); - runWriteKernel(Q, AlignedWritable.data(), N); - runWriteKernel(Q, UnalignedWritablePtr, N); - if (!checkExpectedPattern(AlignedWritable.data(), N)) { std::cerr << "Unexpected data in aligned writable buffer\n"; return 1; } - if (!checkExpectedPattern(UnalignedWritablePtr, N)) { - std::cerr << "Unexpected data in misaligned writable buffer\n"; - return 1; - } - // --- Read-only immutability --- - // A read-only buffer must leave the user's source untouched, aligned or not. if (!runReadOnlyImmutability(Q, AlignedInput.data(), N)) { std::cerr << "Read-only buffer modified aligned source data\n"; return 1; } - if (!runReadOnlyImmutability(Q, ReadOnlyUnalignedPtr, N)) { - std::cerr << "Read-only buffer modified misaligned source data\n"; - return 1; - } - - // --- Mid-lifetime host_accessor (map/unmap) + final write-back --- - // Exercises the bidirectional map/unmap sync path and the final copy-back. - // Expected final pattern at the host pointer: I + 101 (kernel wrote I, host - // added 100 via map/unmap, kernel added 1). - 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; - }; std::vector AlignedMid(N, 0); if (!runMidLifeHostAccessor(Q, AlignedMid.data(), N) || @@ -210,12 +173,84 @@ int main() { return 1; } - std::vector MidStorage(sizeof(int) * N + 1, 0); - int *UnalignedMidPtr = reinterpret_cast(MidStorage.data() + 1); - if (!runMidLifeHostAccessor(Q, UnalignedMidPtr, N) || - !checkMidLife(UnalignedMidPtr)) { - std::cerr << "Mid-life host_accessor failed on misaligned 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; From 518244369a76f7d16be0beff73fa69a6f9809a24 Mon Sep 17 00:00:00 2001 From: "Mateusz P. Nowak" Date: Mon, 22 Jun 2026 11:57:26 +0000 Subject: [PATCH 6/6] fix opencl path --- .../buffer_shadow_copy_platform_policy.cpp | 12 ++++++++---- .../source/adapters/level_zero/common.hpp | 4 ++-- .../source/adapters/level_zero/device.cpp | 8 ++++---- .../adapters/level_zero/helpers/memory_helpers.cpp | 14 +++++--------- 4 files changed, 19 insertions(+), 19 deletions(-) diff --git a/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp b/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp index b02d126b710b6..9177be4aab34d 100644 --- a/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp +++ b/sycl/test-e2e/Regression/buffer_shadow_copy_platform_policy.cpp @@ -11,7 +11,8 @@ // 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::buffer Buf(HostPtr, sycl::range<1>(N), + {sycl::property::buffer::use_host_ptr{}}); sycl::buffer SumBuf(1); Q.submit([&](sycl::handler &CGH) { @@ -33,7 +34,8 @@ static int runReadOnlySumKernel(sycl::queue &Q, const int *HostPtr, size_t N) { // 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::buffer Buf(HostPtr, sycl::range<1>(N), + {sycl::property::buffer::use_host_ptr{}}); Q.submit([&](sycl::handler &CGH) { auto OutAcc = Buf.get_access(CGH); @@ -67,7 +69,8 @@ static bool checkExpectedPattern(const int *Ptr, size_t N) { // 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::buffer Buf(HostPtr, sycl::range<1>(N), + {sycl::property::buffer::use_host_ptr{}}); Q.submit([&](sycl::handler &CGH) { auto Acc = Buf.get_access(CGH); @@ -110,7 +113,8 @@ static bool runReadOnlyImmutability(sycl::queue &Q, const int *HostPtr, std::memcpy(Orig.data(), HostPtr, sizeof(int) * N); { - sycl::buffer Buf(HostPtr, sycl::range<1>(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); 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 0bed7e269142c..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,15 +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); - - // doZeUSMImport silently ignores driver-level failures (e.g., misaligned - // ptr), so re-query to confirm the import actually succeeded before - // reporting it to callers. - ret = getMemoryAttrs(hContext, ptr, nullptr, &properties); - return ret == UR_RESULT_SUCCESS && - properties.type != ZE_MEMORY_TYPE_UNKNOWN; + // 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; }