[llvm] [libsycl] Initial prefetch implementation for queue (PR #212104)
via llvm-commits
llvm-commits at lists.llvm.org
Wed Aug 12 06:08:30 PDT 2026
https://github.com/Robertkq updated https://github.com/llvm/llvm-project/pull/212104
>From 8713a2fd0292b964b7a7b430ffbd7d0319a00e6e Mon Sep 17 00:00:00 2001
From: Robertkq <robertvuia06 at gmail.com>
Date: Sat, 25 Jul 2026 14:35:09 +0300
Subject: [PATCH 1/7] Initial implementation of prefetch for queue
---
libsycl/include/sycl/__impl/queue.hpp | 35 ++++++++++++++++++++++++++
libsycl/src/detail/queue_impl.cpp | 27 ++++++++++++++++++++
libsycl/src/detail/queue_impl.hpp | 9 +++++++
libsycl/src/queue.cpp | 8 ++++++
libsycl/unittests/queue/CMakeLists.txt | 1 +
libsycl/unittests/queue/prefetch.cpp | 15 +++++++++++
6 files changed, 95 insertions(+)
create mode 100644 libsycl/unittests/queue/prefetch.cpp
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index 32b057104a165..30cc89720a8e3 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -413,6 +413,41 @@ class _LIBSYCL_EXPORT queue {
event memcpy(void *dest, const void *src, std::size_t numBytes,
const std::vector<event> &depEvents);
+ /// Provides hints to the runtime library that data should be made available
+ /// on a device earlier than Unified Shared Memory would normally require it
+ /// to be available.
+ ///
+ /// \param ptr is a USM pointer to the memory to be prefetched to the device.
+ /// \param numBytes is a number of bytes to be prefetched.
+ /// \return an event representing prefetch operation.
+ event prefetch(void *ptr, std::size_t numBytes) {
+ return prefetch(ptr, numBytes, std::vector<event>{});
+ }
+
+ /// Provides hints to the runtime library that data should be made available
+ /// on a device earlier than Unified Shared Memory would normally require it
+ /// to be available.
+ ///
+ /// \param ptr is a USM pointer to the memory to be prefetched to the device.
+ /// \param numBytes is a number of bytes to be prefetched.
+ /// \param depEvent is an event that specifies the kernel dependencies.
+ /// \return an event representing prefetch operation.
+ event prefetch(void *ptr, std::size_t numBytes, event depEvent) {
+ return prefetch(ptr, numBytes, std::vector<event>{depEvent});
+ }
+
+ /// Provides hints to the runtime library that data should be made available
+ /// on a device earlier than Unified Shared Memory would normally require it
+ /// to be available.
+ ///
+ /// \param ptr is a USM pointer to the memory to be prefetched to the device.
+ /// \param numBytes is a number of bytes to be prefetched.
+ /// \param depEvents is a vector of events that specify the kernel
+ /// dependencies.
+ /// \return an event representing prefetch operation.
+ event prefetch(void *ptr, std::size_t numBytes,
+ const std::vector<event> &depEvents);
+
private:
template <typename KernelName, int Dims, template <int> class Range,
typename... Rest>
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index 430c0eb302f72..cb6e0f0777ebc 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -202,6 +202,33 @@ QueueImpl::memcpy(void *Dest, const void *Src, std::size_t NumBytes,
return createEvent();
}
+EventImplPtr QueueImpl::prefetch(void *Ptr, std::size_t NumBytes,
+ const std::vector<EventImplPtr> &DepEvents) {
+ checkEventsPlatformMatch(DepEvents, MDevice.getPlatformImpl());
+
+ if (NumBytes == 0) {
+ handleEventDependencies(DepEvents);
+ return createEvent();
+ }
+ if (!Ptr) {
+ throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
+ "Nullptr argument in prefetch operation");
+ }
+
+ auto [TargetDevice, IsHost] = getAllocDevice(Ptr);
+
+ std::size_t Count = 1;
+ const void *Mems[] = {Ptr};
+ const std::size_t Sizes[] = {NumBytes};
+
+ ol_mem_migration_flags_t Flag = OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE;
+
+ handleEventDependencies(DepEvents);
+ callAndThrow(olMemPrefetch, MOffloadQueue, Count, Mems, Sizes, Flag);
+
+ return nullptr;
+}
+
void QueueImpl::handleEventDependencies(const std::vector<EventImplPtr> &Deps) {
// TODO: liboffload supports only in-order queues and no cross context waiting
// is available now that means that this code is excessive but correct. I
diff --git a/libsycl/src/detail/queue_impl.hpp b/libsycl/src/detail/queue_impl.hpp
index e062546de5c5f..c889a189a7ff4 100644
--- a/libsycl/src/detail/queue_impl.hpp
+++ b/libsycl/src/detail/queue_impl.hpp
@@ -117,6 +117,15 @@ class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
EventImplPtr memcpy(void *Dest, const void *Src, std::size_t NumBytes,
const std::vector<EventImplPtr> &DepEvents);
+ /// Submits a prefetch operation for a USM pointer.
+ ///
+ /// \param Ptr is a USM pointer to the memory to be prefetched to the device.
+ /// \param NumBytes is a number of bytes to be prefetched.
+ /// \param DepEvents is a vector of dependencies for the operation.
+ /// \return an event impl object that represents the status of the operation.
+ EventImplPtr prefetch(void *Ptr, std::size_t NumBytes,
+ const std::vector<EventImplPtr> &DepEvents);
+
private:
void handleEventDependencies(const std::vector<EventImplPtr> &Dep);
EventImplPtr createEvent(std::vector<EventImplPtr> &&Deps = {});
diff --git a/libsycl/src/queue.cpp b/libsycl/src/queue.cpp
index 6d15537636936..c257a09367f3f 100644
--- a/libsycl/src/queue.cpp
+++ b/libsycl/src/queue.cpp
@@ -47,6 +47,14 @@ event queue::memcpy(void *dest, const void *src, std::size_t numBytes,
return detail::createSyclObjFromImpl<event>(EventImplPtr);
}
+event queue::prefetch(void *ptr, std::size_t numBytes,
+ const std::vector<event> &depEvents) {
+ std::shared_ptr<detail::EventImpl> EventImplPtr =
+ impl->prefetch(ptr, numBytes, detail::getSyclObjImpls(depEvents));
+ assert(EventImplPtr);
+ return detail::createSyclObjFromImpl<event>(EventImplPtr);
+}
+
event queue::getLastEvent() {
return detail::createSyclObjFromImpl<event>(impl->getLastEvent());
}
diff --git a/libsycl/unittests/queue/CMakeLists.txt b/libsycl/unittests/queue/CMakeLists.txt
index ea98c6d0d48fe..418fd1789ca76 100644
--- a/libsycl/unittests/queue/CMakeLists.txt
+++ b/libsycl/unittests/queue/CMakeLists.txt
@@ -1,4 +1,5 @@
add_sycl_unittest(QueueTests
+ prefetch.cpp
memcpy.cpp
queue.cpp
sycl_kernel_launch.cpp
diff --git a/libsycl/unittests/queue/prefetch.cpp b/libsycl/unittests/queue/prefetch.cpp
new file mode 100644
index 0000000000000..4068c165a3734
--- /dev/null
+++ b/libsycl/unittests/queue/prefetch.cpp
@@ -0,0 +1,15 @@
+#include <mock/helpers.hpp>
+
+#include <sycl/__impl/device.hpp>
+#include <sycl/__impl/queue.hpp>
+
+#include <detail/device_impl.hpp>
+#include <detail/queue_impl.hpp>
+
+#include <gmock/gmock.h>
+#include <gtest/gtest.h>
+
+using namespace sycl;
+using namespace ::testing;
+
+TEST(Queue, Prefetch) {}
\ No newline at end of file
>From d8697b9e47c010b032f8af50c81524d64105d5cb Mon Sep 17 00:00:00 2001
From: Robertkq <robertvuia06 at gmail.com>
Date: Sat, 25 Jul 2026 23:32:40 +0300
Subject: [PATCH 2/7] Switch between appropriate flag & introduce unit tests
for normal use case and no-op case
---
libsycl/src/detail/queue_impl.cpp | 6 ++--
libsycl/unittests/mock/helpers.hpp | 3 ++
libsycl/unittests/mock/mock.cpp | 7 +++++
libsycl/unittests/queue/prefetch.cpp | 47 +++++++++++++++++++++++++++-
4 files changed, 60 insertions(+), 3 deletions(-)
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index cb6e0f0777ebc..8b5cb7230878d 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -221,12 +221,14 @@ EventImplPtr QueueImpl::prefetch(void *Ptr, std::size_t NumBytes,
const void *Mems[] = {Ptr};
const std::size_t Sizes[] = {NumBytes};
- ol_mem_migration_flags_t Flag = OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE;
+ ol_mem_migration_flags_t Flag = (MDevice.isCPU())
+ ? OL_MEM_MIGRATION_FLAG_DEVICE_TO_HOST
+ : OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE;
handleEventDependencies(DepEvents);
callAndThrow(olMemPrefetch, MOffloadQueue, Count, Mems, Sizes, Flag);
- return nullptr;
+ return createEvent();
}
void QueueImpl::handleEventDependencies(const std::vector<EventImplPtr> &Deps) {
diff --git a/libsycl/unittests/mock/helpers.hpp b/libsycl/unittests/mock/helpers.hpp
index 2435a49b22f29..d989d3a5bc0e9 100644
--- a/libsycl/unittests/mock/helpers.hpp
+++ b/libsycl/unittests/mock/helpers.hpp
@@ -125,6 +125,9 @@ class MockLiboffload {
(ol_queue_handle_t Queue, void *DstPtr,
ol_device_handle_t DstDevice, const void *SrcPtr,
ol_device_handle_t SrcDevice, size_t Size));
+ MOCK_METHOD(ol_result_t, olMemPrefetch,
+ (ol_queue_handle_t Queue, size_t Count, const void **Mems,
+ const size_t *Sizes, ol_mem_migration_flags_t Flags));
MOCK_METHOD(ol_result_t, olGetMemInfo,
(const void *Ptr, ol_mem_info_t PropName, size_t PropSize,
void *PropValue));
diff --git a/libsycl/unittests/mock/mock.cpp b/libsycl/unittests/mock/mock.cpp
index 0a2eefb887c4b..01430524b1b5e 100644
--- a/libsycl/unittests/mock/mock.cpp
+++ b/libsycl/unittests/mock/mock.cpp
@@ -116,6 +116,13 @@ ol_result_t olMemcpy(ol_queue_handle_t Queue, void *DstPtr,
SrcDevice, Size);
}
+ol_result_t olMemPrefetch(ol_queue_handle_t Queue, size_t Count,
+ const void **Mems, const size_t *Sizes,
+ ol_mem_migration_flags_t Flags) {
+ return mock::getMockLiboffload().olMemPrefetch(Queue, Count, Mems, Sizes,
+ Flags);
+}
+
ol_result_t olGetMemInfo(const void *Ptr, ol_mem_info_t PropName,
size_t PropSize, void *PropValue) {
return mock::getMockLiboffload().olGetMemInfo(Ptr, PropName, PropSize,
diff --git a/libsycl/unittests/queue/prefetch.cpp b/libsycl/unittests/queue/prefetch.cpp
index 4068c165a3734..025271ac6980d 100644
--- a/libsycl/unittests/queue/prefetch.cpp
+++ b/libsycl/unittests/queue/prefetch.cpp
@@ -12,4 +12,49 @@
using namespace sycl;
using namespace ::testing;
-TEST(Queue, Prefetch) {}
\ No newline at end of file
+TEST(Queue, TwoPrefetches) {
+ constexpr std::size_t NumBytes = 1024;
+ constexpr int NPrefetches = 2;
+
+ mock::MockWrapper Mock;
+ queue Q;
+
+ void *Ptr = reinterpret_cast<void *>(1);
+
+ bool IsCpu = Q.get_device().is_cpu();
+ ol_mem_migration_flags_t ExpectedFlag =
+ IsCpu ? OL_MEM_MIGRATION_FLAG_DEVICE_TO_HOST
+ : OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE;
+
+ EXPECT_CALL(Mock.get(), olMemPrefetch(_, 1, _, _, ExpectedFlag))
+ .Times(NPrefetches)
+ .WillRepeatedly([&](ol_queue_handle_t Queue, size_t Count,
+ const void **Mems, const size_t *Sizes,
+ ol_mem_migration_flags_t Flags) -> ol_result_t {
+ EXPECT_NE(Queue, nullptr);
+ EXPECT_EQ(Count, 1u);
+ EXPECT_EQ(Mems[0], Ptr);
+ EXPECT_EQ(Sizes[0], NumBytes);
+ EXPECT_EQ(Flags, ExpectedFlag);
+ return OL_SUCCESS;
+ });
+
+ EXPECT_CALL(Mock.get(), olCreateEvent(_, _, _)).Times(NPrefetches);
+
+ event Event = Q.prefetch(Ptr, NumBytes);
+
+ // second prefetch depeends on the first one
+ EXPECT_CALL(Mock.get(), olWaitEvents(_, _, 1));
+ Q.prefetch(Ptr, NumBytes, Event);
+}
+
+TEST(Queue, PrefetchZeroBytes) {
+ mock::MockWrapper Mock;
+ queue Q;
+
+ EXPECT_CALL(Mock.get(), olMemPrefetch(_, _, _, _, _)).Times(0);
+ EXPECT_CALL(Mock.get(), olWaitEvents(_, _, 1)).Times(1);
+
+ event Event = Q.prefetch(nullptr, 0);
+ Q.prefetch(nullptr, 0, Event);
+}
>From b957cfbf192b34b0c7da4db4d03afb38b821cb18 Mon Sep 17 00:00:00 2001
From: Robertkq <robertvuia06 at gmail.com>
Date: Sun, 26 Jul 2026 12:29:25 +0300
Subject: [PATCH 3/7] remove unecessary includes for tests
---
libsycl/unittests/queue/prefetch.cpp | 4 ----
1 file changed, 4 deletions(-)
diff --git a/libsycl/unittests/queue/prefetch.cpp b/libsycl/unittests/queue/prefetch.cpp
index 025271ac6980d..6f14dcb4e449b 100644
--- a/libsycl/unittests/queue/prefetch.cpp
+++ b/libsycl/unittests/queue/prefetch.cpp
@@ -1,11 +1,7 @@
#include <mock/helpers.hpp>
-#include <sycl/__impl/device.hpp>
#include <sycl/__impl/queue.hpp>
-#include <detail/device_impl.hpp>
-#include <detail/queue_impl.hpp>
-
#include <gmock/gmock.h>
#include <gtest/gtest.h>
>From ab204f81b17ac1087437ce22450b5c8e71933e60 Mon Sep 17 00:00:00 2001
From: Robertkq <robertvuia06 at gmail.com>
Date: Tue, 28 Jul 2026 17:22:18 +0300
Subject: [PATCH 4/7] address reviews
---
libsycl/include/sycl/__impl/queue.hpp | 6 +-
libsycl/src/detail/queue_impl.cpp | 4 +-
libsycl/test/usm/aligned_alloc_functions.cpp | 83 ++++++++++++++++++++
libsycl/unittests/mock/helpers.cpp | 14 +++-
libsycl/unittests/queue/prefetch.cpp | 4 +-
5 files changed, 101 insertions(+), 10 deletions(-)
create mode 100644 libsycl/test/usm/aligned_alloc_functions.cpp
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index 30cc89720a8e3..bd6613e95324b 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -413,7 +413,7 @@ class _LIBSYCL_EXPORT queue {
event memcpy(void *dest, const void *src, std::size_t numBytes,
const std::vector<event> &depEvents);
- /// Provides hints to the runtime library that data should be made available
+ /// Provides hints to the runtime library that data can be made available
/// on a device earlier than Unified Shared Memory would normally require it
/// to be available.
///
@@ -424,7 +424,7 @@ class _LIBSYCL_EXPORT queue {
return prefetch(ptr, numBytes, std::vector<event>{});
}
- /// Provides hints to the runtime library that data should be made available
+ /// Provides hints to the runtime library that data can be made available
/// on a device earlier than Unified Shared Memory would normally require it
/// to be available.
///
@@ -436,7 +436,7 @@ class _LIBSYCL_EXPORT queue {
return prefetch(ptr, numBytes, std::vector<event>{depEvent});
}
- /// Provides hints to the runtime library that data should be made available
+ /// Provides hints to the runtime library that data can be made available
/// on a device earlier than Unified Shared Memory would normally require it
/// to be available.
///
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index 8b5cb7230878d..e44a35ecf4f4b 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -215,9 +215,7 @@ EventImplPtr QueueImpl::prefetch(void *Ptr, std::size_t NumBytes,
"Nullptr argument in prefetch operation");
}
- auto [TargetDevice, IsHost] = getAllocDevice(Ptr);
-
- std::size_t Count = 1;
+ constexpr std::size_t Count = 1;
const void *Mems[] = {Ptr};
const std::size_t Sizes[] = {NumBytes};
diff --git a/libsycl/test/usm/aligned_alloc_functions.cpp b/libsycl/test/usm/aligned_alloc_functions.cpp
new file mode 100644
index 0000000000000..a422e7e28bf6e
--- /dev/null
+++ b/libsycl/test/usm/aligned_alloc_functions.cpp
@@ -0,0 +1,83 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <cassert>
+#include <cstddef>
+#include <cstdint>
+#include <iostream>
+#include <sycl/sycl.hpp>
+
+using namespace sycl;
+
+void check_aligned_allocation(size_t Alignment, auto AllocFn, queue &q) {
+ constexpr size_t N = 10;
+ void *ptrs[N];
+
+ // 1. Allocate multiple blocks
+ for (size_t i = 0; i < N; ++i) {
+ ptrs[i] = AllocFn();
+ assert(ptrs[i] != nullptr && "Allocation returned nullptr!");
+
+ // 2. Bitwise alignment check
+ auto addr = reinterpret_cast<uintptr_t>(ptrs[i]);
+ if ((addr & (Alignment - 1)) != 0) {
+ std::cerr << "Address " << ptrs[i] << " not aligned to " << Alignment
+ << " bytes!\n";
+ assert(false && "Alignment check failed");
+ }
+ }
+
+ // 3. Cleanup
+ for (size_t i = 0; i < N; ++i) {
+ free(ptrs[i], q);
+ }
+}
+
+int main() {
+ queue q;
+ context ctx = q.get_context();
+ device d = q.get_device();
+
+ size_t Alignments[] = {16, 32, 64, 128, 256, 512};
+ size_t Size = 1024;
+
+ for (size_t Align : Alignments) {
+ // Test aligned_alloc_device
+ check_aligned_allocation(
+ Align, [&]() { return aligned_alloc_device(Align, Size, q); }, q);
+ check_aligned_allocation(
+ Align, [&]() { return aligned_alloc_device(Align, Size, d, ctx); }, q);
+
+ // Test aligned_alloc_host
+ if (d.has(aspect::usm_host_allocations)) {
+ check_aligned_allocation(
+ Align, [&]() { return aligned_alloc_host(Align, Size, q); }, q);
+ check_aligned_allocation(
+ Align, [&]() { return aligned_alloc_host(Align, Size, ctx); }, q);
+ }
+
+ // Test aligned_alloc_shared
+ if (d.has(aspect::usm_shared_allocations)) {
+ check_aligned_allocation(
+ Align, [&]() { return aligned_alloc_shared(Align, Size, q); }, q);
+ check_aligned_allocation(
+ Align, [&]() { return aligned_alloc_shared(Align, Size, d, ctx); },
+ q);
+ }
+
+ // Test generic aligned_alloc
+ check_aligned_allocation(
+ Align,
+ [&]() { return aligned_alloc(Align, Size, q, usm::alloc::device); }, q);
+ check_aligned_allocation(
+ Align,
+ [&]() {
+ return aligned_alloc(Align, Size, d, ctx, usm::alloc::device);
+ },
+ q);
+ }
+
+ std::cout << "All aligned USM E2E tests passed!\n";
+ return 0;
+}
\ No newline at end of file
diff --git a/libsycl/unittests/mock/helpers.cpp b/libsycl/unittests/mock/helpers.cpp
index 65c3ebb078d57..780a302c81849 100644
--- a/libsycl/unittests/mock/helpers.cpp
+++ b/libsycl/unittests/mock/helpers.cpp
@@ -307,7 +307,19 @@ void mock::MockLiboffload::initDefault() {
EXPECT_NE(SrcDevice, nullptr);
return OL_SUCCESS;
});
-
+ ON_CALL(*this, olMemPrefetch)
+ .WillByDefault([this](ol_queue_handle_t Queue, size_t Count,
+ const void **Mems, const size_t *Sizes,
+ ol_mem_migration_flags_t Flags) -> ol_result_t {
+ if (!Queue)
+ return makeEmptyStrError(OL_ERRC_INVALID_NULL_HANDLE);
+ if (Count > 0 && (Mems == nullptr || Sizes == nullptr))
+ return makeEmptyStrError(OL_ERRC_INVALID_NULL_POINTER);
+ if ((Flags & ~(OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE |
+ OL_MEM_MIGRATION_FLAG_DEVICE_TO_HOST)) != 0)
+ return makeEmptyStrError(OL_ERRC_INVALID_ENUMERATION);
+ return OL_SUCCESS;
+ });
ON_CALL(*this, olGetMemInfo)
.WillByDefault([this](const void *Ptr, ol_mem_info_t PropName,
size_t PropSize, void *PropValue) -> ol_result_t {
diff --git a/libsycl/unittests/queue/prefetch.cpp b/libsycl/unittests/queue/prefetch.cpp
index 6f14dcb4e449b..869aa992fb045 100644
--- a/libsycl/unittests/queue/prefetch.cpp
+++ b/libsycl/unittests/queue/prefetch.cpp
@@ -28,10 +28,8 @@ TEST(Queue, TwoPrefetches) {
const void **Mems, const size_t *Sizes,
ol_mem_migration_flags_t Flags) -> ol_result_t {
EXPECT_NE(Queue, nullptr);
- EXPECT_EQ(Count, 1u);
EXPECT_EQ(Mems[0], Ptr);
EXPECT_EQ(Sizes[0], NumBytes);
- EXPECT_EQ(Flags, ExpectedFlag);
return OL_SUCCESS;
});
@@ -39,7 +37,7 @@ TEST(Queue, TwoPrefetches) {
event Event = Q.prefetch(Ptr, NumBytes);
- // second prefetch depeends on the first one
+ // second prefetch depends on the first one
EXPECT_CALL(Mock.get(), olWaitEvents(_, _, 1));
Q.prefetch(Ptr, NumBytes, Event);
}
>From d800151f1e8106e8f43e6ece7a052ba066ef35ed Mon Sep 17 00:00:00 2001
From: Robertkq <robertvuia06 at gmail.com>
Date: Fri, 31 Jul 2026 19:04:09 +0300
Subject: [PATCH 5/7] unconditionally prefetch for device
---
libsycl/src/detail/queue_impl.cpp | 5 ++---
libsycl/unittests/queue/prefetch.cpp | 6 ++----
2 files changed, 4 insertions(+), 7 deletions(-)
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index e44a35ecf4f4b..a2df1bd815993 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -219,9 +219,8 @@ EventImplPtr QueueImpl::prefetch(void *Ptr, std::size_t NumBytes,
const void *Mems[] = {Ptr};
const std::size_t Sizes[] = {NumBytes};
- ol_mem_migration_flags_t Flag = (MDevice.isCPU())
- ? OL_MEM_MIGRATION_FLAG_DEVICE_TO_HOST
- : OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE;
+ constexpr ol_mem_migration_flags_t Flag =
+ OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE;
handleEventDependencies(DepEvents);
callAndThrow(olMemPrefetch, MOffloadQueue, Count, Mems, Sizes, Flag);
diff --git a/libsycl/unittests/queue/prefetch.cpp b/libsycl/unittests/queue/prefetch.cpp
index 869aa992fb045..0abed1dd3f8f5 100644
--- a/libsycl/unittests/queue/prefetch.cpp
+++ b/libsycl/unittests/queue/prefetch.cpp
@@ -17,10 +17,8 @@ TEST(Queue, TwoPrefetches) {
void *Ptr = reinterpret_cast<void *>(1);
- bool IsCpu = Q.get_device().is_cpu();
- ol_mem_migration_flags_t ExpectedFlag =
- IsCpu ? OL_MEM_MIGRATION_FLAG_DEVICE_TO_HOST
- : OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE;
+ constexpr ol_mem_migration_flags_t ExpectedFlag =
+ OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE;
EXPECT_CALL(Mock.get(), olMemPrefetch(_, 1, _, _, ExpectedFlag))
.Times(NPrefetches)
>From a12bddbc6914291795c2a7f4ed18d38b1de14845 Mon Sep 17 00:00:00 2001
From: Robertkq <robertvuia06 at gmail.com>
Date: Mon, 10 Aug 2026 17:48:54 +0300
Subject: [PATCH 6/7] remove miscommited file & fix expected usage of
olMemPrefetch flags
---
libsycl/test/usm/aligned_alloc_functions.cpp | 83 --------------------
libsycl/unittests/mock/helpers.cpp | 3 +-
2 files changed, 1 insertion(+), 85 deletions(-)
delete mode 100644 libsycl/test/usm/aligned_alloc_functions.cpp
diff --git a/libsycl/test/usm/aligned_alloc_functions.cpp b/libsycl/test/usm/aligned_alloc_functions.cpp
deleted file mode 100644
index a422e7e28bf6e..0000000000000
--- a/libsycl/test/usm/aligned_alloc_functions.cpp
+++ /dev/null
@@ -1,83 +0,0 @@
-// REQUIRES: any-device
-// RUN: %clangxx -fsycl %s -o %t.out
-// RUN: %t.out
-
-#include <cassert>
-#include <cstddef>
-#include <cstdint>
-#include <iostream>
-#include <sycl/sycl.hpp>
-
-using namespace sycl;
-
-void check_aligned_allocation(size_t Alignment, auto AllocFn, queue &q) {
- constexpr size_t N = 10;
- void *ptrs[N];
-
- // 1. Allocate multiple blocks
- for (size_t i = 0; i < N; ++i) {
- ptrs[i] = AllocFn();
- assert(ptrs[i] != nullptr && "Allocation returned nullptr!");
-
- // 2. Bitwise alignment check
- auto addr = reinterpret_cast<uintptr_t>(ptrs[i]);
- if ((addr & (Alignment - 1)) != 0) {
- std::cerr << "Address " << ptrs[i] << " not aligned to " << Alignment
- << " bytes!\n";
- assert(false && "Alignment check failed");
- }
- }
-
- // 3. Cleanup
- for (size_t i = 0; i < N; ++i) {
- free(ptrs[i], q);
- }
-}
-
-int main() {
- queue q;
- context ctx = q.get_context();
- device d = q.get_device();
-
- size_t Alignments[] = {16, 32, 64, 128, 256, 512};
- size_t Size = 1024;
-
- for (size_t Align : Alignments) {
- // Test aligned_alloc_device
- check_aligned_allocation(
- Align, [&]() { return aligned_alloc_device(Align, Size, q); }, q);
- check_aligned_allocation(
- Align, [&]() { return aligned_alloc_device(Align, Size, d, ctx); }, q);
-
- // Test aligned_alloc_host
- if (d.has(aspect::usm_host_allocations)) {
- check_aligned_allocation(
- Align, [&]() { return aligned_alloc_host(Align, Size, q); }, q);
- check_aligned_allocation(
- Align, [&]() { return aligned_alloc_host(Align, Size, ctx); }, q);
- }
-
- // Test aligned_alloc_shared
- if (d.has(aspect::usm_shared_allocations)) {
- check_aligned_allocation(
- Align, [&]() { return aligned_alloc_shared(Align, Size, q); }, q);
- check_aligned_allocation(
- Align, [&]() { return aligned_alloc_shared(Align, Size, d, ctx); },
- q);
- }
-
- // Test generic aligned_alloc
- check_aligned_allocation(
- Align,
- [&]() { return aligned_alloc(Align, Size, q, usm::alloc::device); }, q);
- check_aligned_allocation(
- Align,
- [&]() {
- return aligned_alloc(Align, Size, d, ctx, usm::alloc::device);
- },
- q);
- }
-
- std::cout << "All aligned USM E2E tests passed!\n";
- return 0;
-}
\ No newline at end of file
diff --git a/libsycl/unittests/mock/helpers.cpp b/libsycl/unittests/mock/helpers.cpp
index 780a302c81849..c6c82b922faff 100644
--- a/libsycl/unittests/mock/helpers.cpp
+++ b/libsycl/unittests/mock/helpers.cpp
@@ -315,8 +315,7 @@ void mock::MockLiboffload::initDefault() {
return makeEmptyStrError(OL_ERRC_INVALID_NULL_HANDLE);
if (Count > 0 && (Mems == nullptr || Sizes == nullptr))
return makeEmptyStrError(OL_ERRC_INVALID_NULL_POINTER);
- if ((Flags & ~(OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE |
- OL_MEM_MIGRATION_FLAG_DEVICE_TO_HOST)) != 0)
+ if (Flags != OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE)
return makeEmptyStrError(OL_ERRC_INVALID_ENUMERATION);
return OL_SUCCESS;
});
>From 0900e9dd7f68e3a4439d96c7f4f1bbfd9d239b08 Mon Sep 17 00:00:00 2001
From: Robertkq <robertvuia06 at gmail.com>
Date: Wed, 12 Aug 2026 16:08:11 +0300
Subject: [PATCH 7/7] Add E2E test for queue prefetch & update default mock
impl
---
libsycl/test/usm/prefetch.cpp | 41 ++++++++++++++++++++++++++++++
libsycl/unittests/mock/helpers.cpp | 12 ++++-----
2 files changed, 47 insertions(+), 6 deletions(-)
create mode 100644 libsycl/test/usm/prefetch.cpp
diff --git a/libsycl/test/usm/prefetch.cpp b/libsycl/test/usm/prefetch.cpp
new file mode 100644
index 0000000000000..8f41bbccecbf8
--- /dev/null
+++ b/libsycl/test/usm/prefetch.cpp
@@ -0,0 +1,41 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <cstddef>
+
+using namespace sycl;
+
+constexpr std::size_t Count = 1024;
+constexpr std::size_t NumBytes = Count * sizeof(int);
+
+int main() {
+ queue Q;
+
+ int *SharedData = malloc_shared<int>(Count, Q);
+ assert(SharedData != nullptr);
+
+ for (std::size_t I = 0; I < Count; ++I)
+ SharedData[I] = static_cast<int>(I);
+
+ event E1 = Q.prefetch(SharedData, NumBytes);
+
+ event E2 = Q.prefetch(SharedData, NumBytes / 2, E1);
+
+ event E3 = Q.prefetch(nullptr, 0, E2);
+ E3.wait();
+
+ // Verify memory access on device after prefetch
+ Q.parallel_for(range<1>{Count}, [=](id<1> Idx) {
+ SharedData[Idx] += 1;
+ }).wait();
+
+ for (std::size_t I = 0; I < Count; ++I)
+ assert(SharedData[I] == static_cast<int>(I + 1));
+
+ free(SharedData, Q);
+
+ return 0;
+}
\ No newline at end of file
diff --git a/libsycl/unittests/mock/helpers.cpp b/libsycl/unittests/mock/helpers.cpp
index c6c82b922faff..d9dddc22f0f87 100644
--- a/libsycl/unittests/mock/helpers.cpp
+++ b/libsycl/unittests/mock/helpers.cpp
@@ -311,12 +311,12 @@ void mock::MockLiboffload::initDefault() {
.WillByDefault([this](ol_queue_handle_t Queue, size_t Count,
const void **Mems, const size_t *Sizes,
ol_mem_migration_flags_t Flags) -> ol_result_t {
- if (!Queue)
- return makeEmptyStrError(OL_ERRC_INVALID_NULL_HANDLE);
- if (Count > 0 && (Mems == nullptr || Sizes == nullptr))
- return makeEmptyStrError(OL_ERRC_INVALID_NULL_POINTER);
- if (Flags != OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE)
- return makeEmptyStrError(OL_ERRC_INVALID_ENUMERATION);
+ EXPECT_NE(Queue, nullptr);
+ if (Count > 0) {
+ EXPECT_NE(Mems, nullptr);
+ EXPECT_NE(Sizes, nullptr);
+ }
+ EXPECT_EQ(Flags, OL_MEM_MIGRATION_FLAG_HOST_TO_DEVICE);
return OL_SUCCESS;
});
ON_CALL(*this, olGetMemInfo)
More information about the llvm-commits
mailing list