intel · Seanst98 · Apr 7, 2025
@@ -78,6 +78,7 @@ class __SYCL_EXPORT memory_pool {
   size_t get_used_size_current() const;
 
   void increase_threshold_to(size_t newThreshold);
+  void trim_to(size_t minBytesToKeep);
 
   // Property getters.
   template <typename PropertyT> bool has_property() const noexcept {

@@ -45,6 +45,10 @@ __SYCL_EXPORT void memory_pool::increase_threshold_to(size_t newThreshold) {
     impl->set_new_threshold(newThreshold);
 }
 
+__SYCL_EXPORT void memory_pool::trim_to(size_t minBytesToKeep) {
+  impl->trim_to(minBytesToKeep);
+}
+
 __SYCL_EXPORT memory_pool::memory_pool(const sycl::context &ctx,
                                        const sycl::device &dev,
                                        sycl::usm::alloc kind,

@@ -221,6 +221,18 @@ void memory_pool_impl::reset_used_size_high() {
           static_cast<void *>(&resetVal), 8 /*uint64_t*/);
 }
 
+void memory_pool_impl::trim_to(size_t minBytesToKeep) {
+  ur_context_handle_t C = MContextImplPtr->getHandleRef();
+  std::shared_ptr<sycl::detail::device_impl> DevImpl =
+      sycl::detail::getSyclObjImpl(MDevice);
+  ur_device_handle_t Device = DevImpl->getHandleRef();
+  const sycl::detail::AdapterPtr &Adapter = MContextImplPtr->getAdapter();
+
+  Adapter
+      ->call<sycl::errc::runtime, sycl::detail::UrApiKind::urUSMPoolTrimToExp>(
+          C, Device, MPoolHandle, minBytesToKeep);
+}
+
 } // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1

@@ -49,6 +49,7 @@ class memory_pool_impl {
   void set_new_threshold(size_t newThreshold);
   void reset_reserved_size_high();
   void reset_used_size_high();
+  void trim_to(size_t minBytesToKeep);
 
 private:
   std::shared_ptr<sycl::detail::context_impl> MContextImplPtr;

@@ -117,7 +117,7 @@ int main() {
 #endif
 
     // Allocate memory to check queries
-    void *dummyPtr = syclexp::async_malloc_from_pool(Q, 2048, MemPool1);
+    void *dummyPtr = syclexp::async_malloc_from_pool(Q, 4096, MemPool1);
 
     ReservedSizeCurrent = MemPool1.get_reserved_size_current();
     UsedSizeCurrent = MemPool1.get_used_size_current();
@@ -135,6 +135,20 @@ int main() {
     assert(UsedSizeCurrent > 0 &&
            "Pool used size has not increased despite allocating memory!");
 
+    // Pool trimming
+    // Implementation may or may not trim to exact amount. We know it will not
+    // trim to below that amount - amount previously reserved must be greater
+    // than minBytesToKeep
+    size_t minBytesToKeep = 3072;
+    MemPool1.trim_to(minBytesToKeep);
+    ReservedSizeCurrent = MemPool1.get_reserved_size_current();
+#ifdef VERBOSE_PRINT
+    std::cout << "Memory pool current reserved size after trimming: "
+              << ReservedSizeCurrent << std::endl;
+#endif
+    assert(ReservedSizeCurrent >= minBytesToKeep &&
+           "Pool reserved size current has decreased too far after trimming!");
+
     // Free that allocation and wait to release back to OS
     syclexp::async_free(Q, dummyPtr);
     Q.wait_and_throw();

@@ -2993,6 +2993,7 @@ _ZN4sycl3_V13ext6oneapi12experimental10async_freeERKNS0_5queueEPvRKNS0_6detail13
 _ZN4sycl3_V13ext6oneapi12experimental10async_freeERNS0_7handlerEPv
 _ZN4sycl3_V13ext6oneapi12experimental10mem_adviseENS0_5queueEPvmiRKNS0_6detail13code_locationE
 _ZN4sycl3_V13ext6oneapi12experimental11memory_pool21increase_threshold_toEm
+_ZN4sycl3_V13ext6oneapi12experimental11memory_pool7trim_toEm
 _ZN4sycl3_V13ext6oneapi12experimental11memory_poolC1ERKNS0_7contextERKNS0_6deviceENS0_3usm5allocERKNS0_13property_listE
 _ZN4sycl3_V13ext6oneapi12experimental11memory_poolC2ERKNS0_7contextERKNS0_6deviceENS0_3usm5allocERKNS0_13property_listE
 _ZN4sycl3_V13ext6oneapi12experimental12async_mallocERKNS0_5queueENS0_3usm5allocEmRKNS0_6detail13code_locationE

@@ -354,7 +354,9 @@
 ??0dynamic_command_group@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV?$command_graph@$0A@@12345@AEBV?$vector@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@V?$allocator@V?$function@$$A6AXAEAVhandler@_V1@sycl@@@Z@std@@@2@@std@@@Z
 ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV0123456@@Z
 ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV0123456@@Z
+??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@@Z
 ??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@_KPEBX@Z
+??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ
 ??0event@_V1@sycl@@AEAA@V?$shared_ptr@Vevent_impl@detail@_V1@sycl@@@std@@@Z
 ??0event@_V1@sycl@@QEAA@$$QEAV012@@Z
 ??0event@_V1@sycl@@QEAA@AEBV012@@Z
@@ -505,14 +507,10 @@
 ??1exception@_V1@sycl@@UEAA@XZ
 ??1exception_list@_V1@sycl@@QEAA@XZ
 ??1executable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ
-??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@V?$command_graph@$0A@@23456@@Z
-?updateWorkGroupMem@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAX_K@Z
 ??1filter_selector@ONEAPI@_V1@sycl@@UEAA@XZ
 ??1filter_selector@oneapi@ext@_V1@sycl@@UEAA@XZ
 ??1fusion_wrapper@experimental@codeplay@ext@_V1@sycl@@QEAA@XZ
 ??1gpu_selector@_V1@sycl@@UEAA@XZ
-??0dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ
-?setArgHelper@handler@_V1@sycl@@AEAAXHAEAVdynamic_work_group_memory_base@detail@experimental@oneapi@ext@23@@Z
 ??1handler@_V1@sycl@@AEAA@XZ
 ??1image_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ
 ??1image_mem_impl@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ
@@ -4347,6 +4345,7 @@
 ?select_device@filter_selector@oneapi@ext@_V1@sycl@@UEBA?AVdevice@45@XZ
 ?setArgHelper@handler@_V1@sycl@@AEAAXH$$QEAVraw_kernel_arg@experimental@oneapi@ext@23@@Z
 ?setArgHelper@handler@_V1@sycl@@AEAAXH$$QEAVsampler@23@@Z
+?setArgHelper@handler@_V1@sycl@@AEAAXHAEAVdynamic_work_group_memory_base@detail@experimental@oneapi@ext@23@@Z
 ?setArgHelper@handler@_V1@sycl@@AEAAXHAEAVwork_group_memory_impl@detail@23@@Z
 ?setArgsHelper@handler@_V1@sycl@@AEAAXH@Z
 ?setArgsToAssociatedAccessors@handler@_V1@sycl@@AEAAXXZ
@@ -4411,6 +4410,7 @@
 ?throwIfActionIsCreated@handler@_V1@sycl@@AEAAXXZ
 ?throwOnKernelParameterMisuseHelper@handler@_V1@sycl@@AEBAXHP6A?AUkernel_param_desc_t@detail@23@H@Z@Z
 ?throw_asynchronous@queue@_V1@sycl@@QEAAXXZ
+?trim_to@memory_pool@experimental@oneapi@ext@_V1@sycl@@QEAAX_K@Z
 ?unmap@experimental@oneapi@ext@_V1@sycl@@YAXPEBX_KAEBVcontext@45@@Z
 ?unsampledImageConstructorNotification@detail@_V1@sycl@@YAXPEAX0AEBV?$optional@W4image_target@_V1@sycl@@@std@@W4mode@access@23@PEBXIAEBUcode_location@123@@Z
 ?unsampledImageConstructorNotification@image_plain@detail@_V1@sycl@@IEAAXAEBUcode_location@234@PEAXPEBXIQEA_KW4image_format@34@@Z
@@ -4422,6 +4422,7 @@
 ?updateAccessor@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXPEBVAccessorBaseHost@267@@Z
 ?updateValue@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXPEBVraw_kernel_arg@34567@_K@Z
 ?updateValue@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAXPEBX_K@Z
+?updateWorkGroupMem@dynamic_parameter_base@detail@experimental@oneapi@ext@_V1@sycl@@IEAAX_K@Z
 ?use_kernel_bundle@handler@_V1@sycl@@QEAAXAEBV?$kernel_bundle@$01@23@@Z
 ?verifyDeviceHasProgressGuarantee@handler@_V1@sycl@@AEAAXW4forward_progress_guarantee@experimental@oneapi@ext@23@W4execution_scope@56723@1@Z
 ?verifyReductionProps@detail@_V1@sycl@@YAXAEBVproperty_list@23@@Z