diff --git a/sycl/doc/extensions/proposed/sycl_ext_oneapi_usm_shortcuts.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_oneapi_usm_shortcuts.asciidoc index d76512a1c07fe..80d85e63a513b 100644 --- a/sycl/doc/extensions/proposed/sycl_ext_oneapi_usm_shortcuts.asciidoc +++ b/sycl/doc/extensions/proposed/sycl_ext_oneapi_usm_shortcuts.asciidoc @@ -322,7 +322,7 @@ _Effects_: Equivalent to: [source,c++,indent=2] ---- sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); -return sycl::malloc_shared(count, syclDevice, ctxt, kind, propList); +return sycl::malloc(count, syclDevice, ctxt, kind, propList); ---- ''' diff --git a/sycl/include/sycl/usm.hpp b/sycl/include/sycl/usm.hpp index 43713f84f7092..4b7bb1529d75d 100644 --- a/sycl/include/sycl/usm.hpp +++ b/sycl/include/sycl/usm.hpp @@ -11,6 +11,7 @@ #include #include #include +#include #include #include #include @@ -349,6 +350,74 @@ __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::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::malloc_device(count, syclDevice, ctxt, 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::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::aligned_alloc_device(alignment, count, syclDevice, ctxt, + 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::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::malloc_shared(count, syclDevice, ctxt, 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::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::aligned_alloc_shared(alignment, count, syclDevice, ctxt, + 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::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::malloc(count, syclDevice, ctxt, kind, 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 = {}) { + sycl::context ctxt = syclDevice.get_platform().khr_get_default_context(); + return sycl::aligned_alloc(alignment, count, syclDevice, ctxt, kind, + 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 3c714ee4cd26c..d4ae709903748 100644 --- a/sycl/source/detail/usm/usm_impl.cpp +++ b/sycl/source/detail/usm/usm_impl.cpp @@ -608,6 +608,47 @@ 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); +} + +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); +} + +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); +} + +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); +} + +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); +} + +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); +} } // namespace ext::oneapi::experimental __SYCL_EXPORT void verifyUSMAllocatorProperties(const property_list &PropList) { diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index 2b6a584bd7f5c..443966f689075 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -117,6 +117,7 @@ inline namespace _V1 { #define SYCL_EXT_INTEL_EVENT_MODE 1 #define SYCL_EXT_ONEAPI_TANGLE 1 #define SYCL_EXT_ONEAPI_INTER_PROCESS_COMMUNICATION 1 +#define SYCL_EXT_ONEAPI_USM_SHORTCUTS 1 // Unfinished KHR extensions. These extensions are only available if the // __DPCPP_ENABLE_UNFINISHED_KHR_EXTENSIONS macro is defined. 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..7d9eb73614f90 --- /dev/null +++ b/sycl/test-e2e/USM/usm_shortcuts_utility.cpp @@ -0,0 +1,131 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +//==------ usm_shortcuts_utility.cpp - USM shortcuts 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 = malloc(N * sizeof(int), dev, usm::alloc::host); + check_and_free(array, dev, ctxt, usm::alloc::host); + + array = + 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); + + array = aligned_alloc(alignof(long long), N * sizeof(int), dev, + usm::alloc::host); + check_and_free(array, dev, ctxt, usm::alloc::host); + + array = 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{}); + check_and_free(array, dev, ctxt, usm::alloc::shared); + + array = malloc_shared(N * sizeof(int), dev); + check_and_free(array, dev, ctxt, usm::alloc::shared); + + array = malloc_shared(N * sizeof(int), dev, property_list{}); + 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{}); + check_and_free(array, dev, ctxt, usm::alloc::shared); + + array = aligned_alloc_shared(alignof(long long), N * sizeof(int), dev); + check_and_free(array, dev, ctxt, usm::alloc::shared); + + array = aligned_alloc_shared(alignof(long long), N * sizeof(int), dev, + property_list{}); + 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{}); + check_and_free(array, dev, ctxt, usm::alloc::device); + + array = malloc_device(N * sizeof(int), dev); + check_and_free(array, dev, ctxt, usm::alloc::device); + + array = malloc_device(N, dev, property_list{}); + 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); + + array = aligned_alloc_device(alignof(long long), N * sizeof(int), dev); + check_and_free(array, dev, ctxt, usm::alloc::device); + + array = aligned_alloc_device(alignof(long long), N * sizeof(int), dev, + property_list{}); + check_and_free(array, dev, ctxt, usm::alloc::device); + } + + return 0; +} diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index f8af016c99c09..38d79e646ff1c 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3008,6 +3008,9 @@ _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_22bind _ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental12physical_memC1ERKNS0_6deviceERKNS0_7contextEm _ZN4sycl3_V13ext6oneapi12experimental12physical_memC2ERKNS0_6deviceERKNS0_7contextEm +_ZN4sycl3_V13ext6oneapi12experimental13aligned_allocEmmRKNS0_6deviceENS0_3usm5allocERKNS0_13property_listE +_ZN4sycl3_V13ext6oneapi12experimental13malloc_deviceEmRKNS0_6deviceERKNS0_13property_listE +_ZN4sycl3_V13ext6oneapi12experimental13malloc_sharedEmRKNS0_6deviceERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental14free_image_memENS3_16image_mem_handleENS3_10image_typeERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental14free_image_memENS3_16image_mem_handleENS3_10image_typeERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental15alloc_image_memERKNS3_16image_descriptorERKNS0_5queueE @@ -3020,6 +3023,8 @@ _ZN4sycl3_V13ext6oneapi12experimental16free_virtual_memEmmRKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental19get_mem_granularityERKNS0_6deviceERKNS0_7contextENS3_16granularity_modeE _ZN4sycl3_V13ext6oneapi12experimental19get_mem_granularityERKNS0_7contextENS3_16granularity_modeE _ZN4sycl3_V13ext6oneapi12experimental19reserve_virtual_memEmmRKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental20aligned_alloc_deviceEmmRKNS0_6deviceERKNS0_13property_listE +_ZN4sycl3_V13ext6oneapi12experimental20aligned_alloc_sharedEmmRKNS0_6deviceERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental20destroy_image_handleERNS3_20sampled_image_handleERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental20destroy_image_handleERNS3_20sampled_image_handleERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental20destroy_image_handleERNS3_22unsampled_image_handleERKNS0_5queueE @@ -3127,6 +3132,7 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail30dynamic_work_group_memory_baseC2Em _ZN4sycl3_V13ext6oneapi12experimental6detail30make_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageENS0_6detail11string_viewESt6vectorISt4pairISA_SA_ESaISD_EE _ZN4sycl3_V13ext6oneapi12experimental6detail30make_kernel_bundle_from_sourceERKNS0_7contextENS3_15source_languageERKSt6vectorISt4byteSaISA_EES9_ISt4pairINS0_6detail11string_viewESH_ESaISI_EE _ZN4sycl3_V13ext6oneapi12experimental6detail33export_device_mem_win32_nt_handleEPvRKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental6mallocEmRKNS0_6deviceENS0_3usm5allocERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6memcpyENS0_5queueEPvPKvmRKNS0_6detail13code_locationE _ZN4sycl3_V13ext6oneapi12experimental6memsetENS0_5queueEPvimRKNS0_6detail13code_locationE _ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_5queueE diff --git a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp index 17ae5f95c67f2..024a9c10557d2 100644 --- a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp @@ -185,5 +185,6 @@ // CHECK-NEXT: ext/oneapi/experimental/event_mode_property.hpp // CHECK-NEXT: sycl_span.hpp // CHECK-NEXT: usm.hpp +// CHECK-NEXT: platform.hpp // CHECK-NEXT: usm/usm_pointer_info.hpp // CHECK-EMPTY: diff --git a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp index 0311e59e7abe5..aa6d427576ef3 100644 --- a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp +++ b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp @@ -172,6 +172,7 @@ // CHECK-NEXT: kernel.hpp // CHECK-NEXT: sampler.hpp // CHECK-NEXT: sycl_span.hpp +// CHECK-NEXT: platform.hpp // CHECK-NEXT: usm/usm_pointer_info.hpp // CHECK-NEXT: usm/usm_allocator.hpp // CHECK-EMPTY: