[SYCL] Implement sycl_ext_oneapi_usm_shortcuts

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<T>(count, syclDevice, ctxt, kind, propList);
+return sycl::malloc<T>(count, syclDevice, ctxt, kind, propList);
 ----
 
 '''

sycl/include/sycl/usm.hpp

@@ -11,6 +11,7 @@
 #include <sycl/detail/common.hpp>
 #include <sycl/detail/export.hpp>
 #include <sycl/device.hpp>
+#include <sycl/platform.hpp>
 #include <sycl/property_list.hpp>
 #include <sycl/queue.hpp>
 #include <sycl/usm/usm_pointer_info.hpp>
@@ -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 <typename T>
+__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<T>(count, syclDevice, ctxt, propList);
+}
+
+__SYCL_EXPORT void *aligned_alloc_device(size_t alignment, size_t numBytes,
+                                         const device &syclDevice,
+                                         const property_list &propList = {});
+
+template <typename T>
+__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<T>(alignment, count, syclDevice, ctxt,
+                                       propList);
+}
+
+__SYCL_EXPORT void *malloc_shared(size_t numBytes, const device &syclDevice,
+                                  const property_list &propList = {});
+
+template <typename T>
+__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<T>(count, syclDevice, ctxt, propList);
+}
+
+__SYCL_EXPORT void *aligned_alloc_shared(size_t alignment, size_t numBytes,
+                                         const device &syclDevice,
+                                         const property_list &propList = {});
+
+template <typename T>
+__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<T>(alignment, count, syclDevice, ctxt,
+                                       propList);
+}
+
+__SYCL_EXPORT void *malloc(size_t numBytes, const device &syclDevice,
+                           usm::alloc kind, const property_list &propList = {});
+
+template <typename T>
+__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<T>(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 <typename T>
+__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<T>(alignment, count, syclDevice, ctxt, kind,
+                                propList);
+}
 } // namespace ext::oneapi::experimental
 
 } // namespace _V1

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) {

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.

sycl/test-e2e/USM/usm_shortcuts_utility.cpp

@@ -0,0 +1,148 @@
+// 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 <sycl/detail/core.hpp>
+
+#include <sycl/ext/intel/experimental/usm_properties.hpp>
+#include <sycl/usm.hpp>
+
+#include <cassert>
+
+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<sycl::info::device::usm_host_allocations>()) {
+    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<int>(N * sizeof(int), dev, usm::alloc::host);
+    check_and_free(array, dev, ctxt, usm::alloc::host);
+
+    array =
+        malloc<int>(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<int>(alignof(long long), N * sizeof(int), dev,
+                               usm::alloc::host);
+    check_and_free(array, dev, ctxt, usm::alloc::host);
+
+    array = aligned_alloc<int>(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<sycl::info::device::usm_shared_allocations>()) {
+    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 = malloc_shared<int>(N * sizeof(int), dev);
+    check_and_free(array, dev, ctxt, usm::alloc::shared);
+
+    array = malloc_shared<int>(
+        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);
+
+    array = aligned_alloc_shared<int>(alignof(long long), N * sizeof(int), dev);
+    check_and_free(array, dev, ctxt, usm::alloc::shared);
+
+    array = aligned_alloc_shared<int>(
+        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<sycl::info::device::usm_device_allocations>()) {
+    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 = malloc_device<int>(N * sizeof(int), dev);
+    check_and_free(array, dev, ctxt, usm::alloc::device);
+
+    array = malloc_device<int>(
+        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);
+
+    array = aligned_alloc_device<int>(alignof(long long), N * sizeof(int), dev);
+    check_and_free(array, dev, ctxt, usm::alloc::device);
+
+    array = aligned_alloc_device<int>(alignof(long long), N * sizeof(int), dev,
+                                      property_list{});
+    check_and_free(array, dev, ctxt, usm::alloc::device);
+  }
+
+  return 0;
+}

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

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:

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: