From 97ed3863be4c44912b826d57b46482e2399682a8 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 16 Dec 2025 07:40:05 -0800 Subject: [PATCH 1/2] [SYCL] Implement sycl_ext_oneapi_usm_shortcuts Signed-off-by: Hu, Peisen --- sycl/include/sycl/usm.hpp | 47 +++++++++++++++ sycl/source/detail/usm/usm_impl.cpp | 88 +++++++++++++++++++++++++++++ 2 files changed, 135 insertions(+) diff --git a/sycl/include/sycl/usm.hpp b/sycl/include/sycl/usm.hpp index 43713f84f7092..f89ed157da2d8 100644 --- a/sycl/include/sycl/usm.hpp +++ b/sycl/include/sycl/usm.hpp @@ -349,6 +349,53 @@ __SYCL_EXPORT void release_from_device_copy(const void *Ptr, __SYCL_EXPORT void release_from_device_copy(const void *Ptr, const queue &Queue); +__SYCL_EXPORT void *malloc_device(size_t numBytes, const device &syclDevice, + const property_list &propList = {}); + +template +__SYCL_EXPORT T *malloc_device(size_t count, const device &syclDevice, + const property_list &propList = {}); + +__SYCL_EXPORT void *aligned_alloc_device(size_t alignment, size_t numBytes, + const device &syclDevice, + const property_list &propList = {}); + +template +__SYCL_EXPORT T *aligned_alloc_device(size_t alignment, size_t count, + const device &syclDevice, + const property_list &propList = {}); + +__SYCL_EXPORT void *malloc_shared(size_t numBytes, const device &syclDevice, + const property_list &propList = {}); + +template +__SYCL_EXPORT T *malloc_shared(size_t count, const device &syclDevice, + const property_list &propList = {}); + +__SYCL_EXPORT void *aligned_alloc_shared(size_t alignment, size_t numBytes, + const device &syclDevice, + const property_list &propList = {}); + +template +__SYCL_EXPORT T *aligned_alloc_shared(size_t alignment, size_t count, + const device &syclDevice, + const property_list &propList = {}); + +__SYCL_EXPORT void *malloc(size_t numBytes, const device &syclDevice, + usm::alloc kind, const property_list &propList = {}); + +template +__SYCL_EXPORT T *malloc(size_t count, const device &syclDevice, usm::alloc kind, + const property_list &propList = {}); + +__SYCL_EXPORT void *aligned_alloc(size_t alignment, size_t numBytes, + const device &syclDevice, usm::alloc kind, + const property_list &propList = {}); + +template +__SYCL_EXPORT T *aligned_alloc(size_t alignment, size_t count, + const device &syclDevice, usm::alloc kind, + const property_list &propList = {}); } // namespace ext::oneapi::experimental } // namespace _V1 diff --git a/sycl/source/detail/usm/usm_impl.cpp b/sycl/source/detail/usm/usm_impl.cpp index 486ee62bd9bc4..1a353ef3cf5e2 100644 --- a/sycl/source/detail/usm/usm_impl.cpp +++ b/sycl/source/detail/usm/usm_impl.cpp @@ -651,6 +651,94 @@ void release_from_device_copy(const void *Ptr, const context &Ctxt) { void release_from_device_copy(const void *Ptr, const queue &Queue) { release_from_usm_device_copy(Ptr, Queue.get_context()); } + +void *malloc_device(size_t numBytes, const device &syclDevice, + const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::malloc_device(numBytes, syclDevice, ctxt, propList); +} + +template +T *malloc_device(size_t count, const device &syclDevice, + const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::malloc_device(count, syclDevice, ctxt, propList); +} + +void *aligned_alloc_device(size_t alignment, size_t numBytes, + const device &syclDevice, + const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::aligned_alloc_device(alignment, numBytes, syclDevice, ctxt, + propList); +} + +template +T *aligned_alloc_device(size_t alignment, size_t count, + const device &syclDevice, + const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::aligned_alloc_device(alignment, count, syclDevice, ctxt, + propList); +} + +void *malloc_shared(size_t numBytes, const device &syclDevice, + const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::malloc_shared(numBytes, syclDevice, ctxt, propList); +} + +template +T *malloc_shared(size_t count, const device &syclDevice, + const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::malloc_shared(count, syclDevice, ctxt, propList); +} + +void *aligned_alloc_shared(size_t alignment, size_t numBytes, + const device &syclDevice, + const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::aligned_alloc_shared(alignment, numBytes, syclDevice, ctxt, + propList); +} + +template +T *aligned_alloc_shared(size_t alignment, size_t count, + const device &syclDevice, + const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::aligned_alloc_shared(alignment, count, syclDevice, ctxt, + propList); +} + +void *malloc(size_t numBytes, const device &syclDevice, usm::alloc kind, + const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::malloc(numBytes, syclDevice, ctxt, kind, propList); +} + +template +T *malloc(size_t count, const device &syclDevice, usm::alloc kind, + const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::malloc_shared(count, syclDevice, ctxt, kind, propList); +} + +void *aligned_alloc(size_t alignment, size_t numBytes, const device &syclDevice, + usm::alloc kind, const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::aligned_alloc(alignment, numBytes, syclDevice, ctxt, kind, + propList); +} + +template +T *aligned_alloc(size_t alignment, size_t count, const device &syclDevice, + usm::alloc kind, const property_list &propList) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::aligned_alloc(alignment, count, syclDevice, ctxt, kind, + propList); +} } // namespace ext::oneapi::experimental __SYCL_EXPORT void verifyUSMAllocatorProperties(const property_list &PropList) { From 7d5f369bb1e4cbb0b84e82414e5800be72f1c024 Mon Sep 17 00:00:00 2001 From: "Hu, Peisen" Date: Tue, 16 Dec 2025 07:42:06 -0800 Subject: [PATCH 2/2] [SYCL] Add respective test Signed-off-by: Hu, Peisen --- sycl/test-e2e/USM/usm_shortcuts_utility.cpp | 99 +++++++++++++++++++++ 1 file changed, 99 insertions(+) create mode 100644 sycl/test-e2e/USM/usm_shortcuts_utility.cpp diff --git a/sycl/test-e2e/USM/usm_shortcuts_utility.cpp b/sycl/test-e2e/USM/usm_shortcuts_utility.cpp new file mode 100644 index 0000000000000..c3d6418f47ec7 --- /dev/null +++ b/sycl/test-e2e/USM/usm_shortcuts_utility.cpp @@ -0,0 +1,99 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +//==------ usm_shortcuts_utility.cpp - USM malloc and aligned_alloc test +//-------==// +// +// 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 namespace sycl; +using namespace sycl::ext::oneapi::experimental; + +constexpr int N = 8; + +static void check_and_free(int *array, const device &dev, const context &ctxt, + usm::alloc expected_type) { + // host device treats all allocations as host allocations + assert((get_pointer_type(array, ctxt) == expected_type) && + "Allocation pointer has unexpected type."); + assert((get_pointer_device(array, ctxt) == dev) && + "Allocation pointer has unexpected device associated with it."); + free(array, ctxt); +} + +int main() { + queue q; + auto dev = q.get_device(); + auto ctxt = q.get_context(); + int *array; + + if (dev.get_info()) { + array = (int *)malloc(N * sizeof(int), dev, usm::alloc::host); + check_and_free(array, dev, ctxt, usm::alloc::host); + + array = + (int *)malloc(N * sizeof(int), dev, usm::alloc::host, property_list{}); + check_and_free(array, dev, ctxt, usm::alloc::host); + + array = (int *)aligned_alloc(alignof(long long), N * sizeof(int), dev, + usm::alloc::host); + check_and_free(array, dev, ctxt, usm::alloc::host); + + array = (int *)aligned_alloc(alignof(long long), N * sizeof(int), dev, + usm::alloc::host, property_list{}); + check_and_free(array, dev, ctxt, usm::alloc::host); + } + + if (dev.get_info()) { + array = (int *)malloc_shared(N * sizeof(int), dev); + check_and_free(array, dev, ctxt, usm::alloc::shared); + + array = (int *)malloc_shared( + N * sizeof(int), dev, + property_list{ + ext::intel::experimental::property::usm::buffer_location{2}}); + check_and_free(array, dev, ctxt, usm::alloc::shared); + + array = + (int *)aligned_alloc_shared(alignof(long long), N * sizeof(int), dev); + check_and_free(array, dev, ctxt, usm::alloc::shared); + + array = (int *)aligned_alloc_shared( + alignof(long long), N * sizeof(int), dev, + property_list{ + ext::intel::experimental::property::usm::buffer_location{2}}); + check_and_free(array, dev, ctxt, usm::alloc::shared); + } + + if (dev.get_info()) { + array = (int *)malloc_device(N * sizeof(int), dev); + check_and_free(array, dev, ctxt, usm::alloc::device); + + array = (int *)malloc_device( + N, dev, + property_list{ + ext::intel::experimental::property::usm::buffer_location(2)}); + check_and_free(array, dev, ctxt, usm::alloc::device); + + array = + (int *)aligned_alloc_device(alignof(long long), N * sizeof(int), dev); + check_and_free(array, dev, ctxt, usm::alloc::device); + + array = (int *)aligned_alloc_device(alignof(long long), N * sizeof(int), + dev, property_list{}); + check_and_free(array, dev, ctxt, usm::alloc::device); + } + + return 0; +}