intel · Seanst98 · Apr 7, 2025
@@ -75,9 +75,13 @@ class __SYCL_EXPORT memory_pool {
   sycl::usm::alloc get_alloc_kind() const;
   size_t get_threshold() const;
   size_t get_reserved_size_current() const;
+  size_t get_reserved_size_high() const;
   size_t get_used_size_current() const;
+  size_t get_used_size_high() const;
 
   void increase_threshold_to(size_t newThreshold);
+  void reset_reserved_size_high();
+  void reset_used_size_high();
 
   // Property getters.
   template <typename PropertyT> bool has_property() const noexcept {

@@ -35,16 +35,32 @@ __SYCL_EXPORT size_t memory_pool::get_reserved_size_current() const {
   return impl->get_reserved_size_current();
 }
 
+__SYCL_EXPORT size_t memory_pool::get_reserved_size_high() const {
+  return impl->get_reserved_size_high();
+}
+
 __SYCL_EXPORT size_t memory_pool::get_used_size_current() const {
   return impl->get_used_size_current();
 }
 
+__SYCL_EXPORT size_t memory_pool::get_used_size_high() const {
+  return impl->get_used_size_high();
+}
+
 __SYCL_EXPORT void memory_pool::increase_threshold_to(size_t newThreshold) {
   // Only increase.
   if (newThreshold > get_threshold())
     impl->set_new_threshold(newThreshold);
 }
 
+__SYCL_EXPORT void memory_pool::reset_reserved_size_high() {
+  impl->reset_reserved_size_high();
+}
+
+__SYCL_EXPORT void memory_pool::reset_used_size_high() {
+  impl->reset_used_size_high();
+}
+
 __SYCL_EXPORT memory_pool::memory_pool(const sycl::context &ctx,
                                        const sycl::device &dev,
                                        sycl::usm::alloc kind,

@@ -97,14 +97,19 @@ int main() {
 
     size_t ReleaseThresholdGet = MemPool1.get_threshold();
     size_t ReservedSizeCurrent = MemPool1.get_reserved_size_current();
+    size_t ReservedSizeHigh = MemPool1.get_reserved_size_high();
     size_t UsedSizeCurrent = MemPool1.get_used_size_current();
+    size_t UsedSizeHigh = MemPool1.get_used_size_high();
 #ifdef VERBOSE_PRINT
     std::cout << "Memory pool release threshold: " << ReleaseThresholdGet
               << std::endl;
     std::cout << "Memory pool current reserved size: " << ReservedSizeCurrent
               << std::endl;
+    std::cout << "Memory pool high reserved size: " << ReservedSizeHigh
+              << std::endl;
     std::cout << "Memory pool current used size: " << UsedSizeCurrent
               << std::endl;
+    std::cout << "Memory pool high used size: " << UsedSizeHigh << std::endl;
 #endif
 
     // Set new threshold -- then check getter
@@ -120,25 +125,50 @@ int main() {
     void *dummyPtr = syclexp::async_malloc_from_pool(Q, 2048, MemPool1);
 
     ReservedSizeCurrent = MemPool1.get_reserved_size_current();
+    ReservedSizeHigh = MemPool1.get_reserved_size_high();
     UsedSizeCurrent = MemPool1.get_used_size_current();
+    UsedSizeHigh = MemPool1.get_used_size_high();
 #ifdef VERBOSE_PRINT
     std::cout << "Memory pool current reserved size: " << ReservedSizeCurrent
               << std::endl;
+    std::cout << "Memory pool high reserved size: " << ReservedSizeHigh
+              << std::endl;
     std::cout << "Memory pool current used size: " << UsedSizeCurrent
               << std::endl;
+    std::cout << "Memory pool high used size: " << UsedSizeHigh << std::endl;
 #endif
 
     // We don't know what the exact sizes of each could be - but they must each
     // be greater than 0
     assert(ReservedSizeCurrent > 0 &&
            "Pool reserved size has not increased despite allocating memory!");
+    assert(ReservedSizeHigh > 0 && "Pool reserved size high has not increased "
+                                   "despite allocating memory!");
     assert(UsedSizeCurrent > 0 &&
            "Pool used size has not increased despite allocating memory!");
+    assert(UsedSizeHigh > 0 &&
+           "Pool used size high has not increased despite allocating memory!");
 
     // Free that allocation and wait to release back to OS
     syclexp::async_free(Q, dummyPtr);
     Q.wait_and_throw();
 
+    // Reset high watermarks and check
+    MemPool1.reset_reserved_size_high();
+    MemPool1.reset_used_size_high();
+
+    ReservedSizeHigh = MemPool1.get_reserved_size_high();
+    UsedSizeHigh = MemPool1.get_used_size_high();
+#ifdef VERBOSE_PRINT
+    std::cout << "Memory pool high reserved size: " << ReservedSizeHigh
+              << std::endl;
+    std::cout << "Memory pool high used size: " << UsedSizeHigh << std::endl;
+#endif
+    assert(ReservedSizeHigh == 0 &&
+           "Pool reserved size high not equal to zero despite resetting!");
+    assert(UsedSizeHigh == 0 &&
+           "Pool used size high not equal to zero despite resetting!");
+
     // Default memory pool
     syclexp::memory_pool DefaultPool =
         Ctx.ext_oneapi_get_default_memory_pool(Dev, sycl::usm::alloc::device);

@@ -2992,7 +2992,9 @@ _ZN4sycl3_V13ext6oneapi10level_zero6detail11make_deviceERKNS0_8platformEm
 _ZN4sycl3_V13ext6oneapi12experimental10async_freeERKNS0_5queueEPvRKNS0_6detail13code_locationE
 _ZN4sycl3_V13ext6oneapi12experimental10async_freeERNS0_7handlerEPv
 _ZN4sycl3_V13ext6oneapi12experimental10mem_adviseENS0_5queueEPvmiRKNS0_6detail13code_locationE
+_ZN4sycl3_V13ext6oneapi12experimental11memory_pool20reset_used_size_highEv
 _ZN4sycl3_V13ext6oneapi12experimental11memory_pool21increase_threshold_toEm
+_ZN4sycl3_V13ext6oneapi12experimental11memory_pool24reset_reserved_size_highEv
 _ZN4sycl3_V13ext6oneapi12experimental11memory_poolC1ERKNS0_7contextERKNS0_6deviceENS0_3usm5allocERKNS0_13property_listE
 _ZN4sycl3_V13ext6oneapi12experimental11memory_poolC2ERKNS0_7contextERKNS0_6deviceENS0_3usm5allocERKNS0_13property_listE
 _ZN4sycl3_V13ext6oneapi12experimental12async_mallocERKNS0_5queueENS0_3usm5allocEmRKNS0_6detail13code_locationE
@@ -3640,7 +3642,9 @@ _ZNK4sycl3_V13ext6oneapi12experimental11memory_pool11getPropListEv
 _ZNK4sycl3_V13ext6oneapi12experimental11memory_pool11get_contextEv
 _ZNK4sycl3_V13ext6oneapi12experimental11memory_pool13get_thresholdEv
 _ZNK4sycl3_V13ext6oneapi12experimental11memory_pool14get_alloc_kindEv
+_ZNK4sycl3_V13ext6oneapi12experimental11memory_pool18get_used_size_highEv
 _ZNK4sycl3_V13ext6oneapi12experimental11memory_pool21get_used_size_currentEv
+_ZNK4sycl3_V13ext6oneapi12experimental11memory_pool22get_reserved_size_highEv
 _ZNK4sycl3_V13ext6oneapi12experimental11memory_pool25get_reserved_size_currentEv
 _ZNK4sycl3_V13ext6oneapi12experimental12physical_mem10get_deviceEv
 _ZNK4sycl3_V13ext6oneapi12experimental12physical_mem11get_contextEv

@@ -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
@@ -4173,6 +4171,7 @@
 ?get_range@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AV?$range@$02@56@XZ
 ?get_range@image_plain@detail@_V1@sycl@@IEBA?AV?$range@$02@34@XZ
 ?get_reserved_size_current@memory_pool@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ
+?get_reserved_size_high@memory_pool@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ
 ?get_root_nodes@modifiable_command_graph@detail@experimental@oneapi@ext@_V1@sycl@@QEBA?AV?$vector@Vnode@experimental@oneapi@ext@_V1@sycl@@V?$allocator@Vnode@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@XZ
 ?get_size@image_plain@detail@_V1@sycl@@IEBA_KXZ
 ?get_size@stream@_V1@sycl@@QEBA_KXZ
@@ -4183,6 +4182,7 @@
 ?get_type@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AW4image_type@23456@XZ
 ?get_type@node@experimental@oneapi@ext@_V1@sycl@@QEBA?AW4node_type@23456@XZ
 ?get_used_size_current@memory_pool@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ
+?get_used_size_high@memory_pool@experimental@oneapi@ext@_V1@sycl@@QEBA_KXZ
 ?get_wait_list@event@_V1@sycl@@QEAA?AV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@XZ
 ?get_width@stream@_V1@sycl@@QEBA_KXZ
 ?get_work_item_buffer_size@stream@_V1@sycl@@QEBA_KXZ
@@ -4335,6 +4335,8 @@
 ?reserve_virtual_mem@experimental@oneapi@ext@_V1@sycl@@YA_K_K0AEBVcontext@45@@Z
 ?reset@filter_selector@ONEAPI@_V1@sycl@@QEBAXXZ
 ?reset@filter_selector@oneapi@ext@_V1@sycl@@QEBAXXZ
+?reset_reserved_size_high@memory_pool@experimental@oneapi@ext@_V1@sycl@@QEAAXXZ
+?reset_used_size_high@memory_pool@experimental@oneapi@ext@_V1@sycl@@QEAAXXZ
 ?sampledImageConstructorNotification@detail@_V1@sycl@@YAXPEAX0AEBV?$optional@W4image_target@_V1@sycl@@@std@@PEBXIAEBUcode_location@123@@Z
 ?sampledImageConstructorNotification@image_plain@detail@_V1@sycl@@IEAAXAEBUcode_location@234@PEAXPEBXIQEA_KW4image_format@34@AEBUimage_sampler@34@@Z
 ?sampledImageDestructorNotification@image_plain@detail@_V1@sycl@@IEAAXPEAX@Z
@@ -4347,6 +4349,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
@@ -4422,6 +4425,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