Stream Ordered Memory Allocator#

HIP Runtime API Reference: Stream Ordered Memory Allocator
Stream Ordered Memory Allocator
Collaboration diagram for Stream Ordered Memory Allocator:

Functions

hipError_t hipMallocAsync (void **dev_ptr, size_t size, hipStream_t stream)
 Allocates memory with stream ordered semantics. More...
 
hipError_t hipFreeAsync (void *dev_ptr, hipStream_t stream)
 Frees memory with stream ordered semantics. More...
 
hipError_t hipMemPoolTrimTo (hipMemPool_t mem_pool, size_t min_bytes_to_hold)
 Releases freed memory back to the OS. More...
 
hipError_t hipMemPoolSetAttribute (hipMemPool_t mem_pool, hipMemPoolAttr attr, void *value)
 Sets attributes of a memory pool. More...
 
hipError_t hipMemPoolGetAttribute (hipMemPool_t mem_pool, hipMemPoolAttr attr, void *value)
 Gets attributes of a memory pool. More...
 
hipError_t hipMemPoolSetAccess (hipMemPool_t mem_pool, const hipMemAccessDesc *desc_list, size_t count)
 Controls visibility of the specified pool between devices. More...
 
hipError_t hipMemPoolGetAccess (hipMemAccessFlags *flags, hipMemPool_t mem_pool, hipMemLocation *location)
 Returns the accessibility of a pool from a device. More...
 
hipError_t hipMemPoolCreate (hipMemPool_t *mem_pool, const hipMemPoolProps *pool_props)
 Creates a memory pool. More...
 
hipError_t hipMemPoolDestroy (hipMemPool_t mem_pool)
 Destroys the specified memory pool. More...
 
hipError_t hipMallocFromPoolAsync (void **dev_ptr, size_t size, hipMemPool_t mem_pool, hipStream_t stream)
 Allocates memory from a specified pool with stream ordered semantics. More...
 
hipError_t hipMemPoolExportToShareableHandle (void *shared_handle, hipMemPool_t mem_pool, hipMemAllocationHandleType handle_type, unsigned int flags)
 Exports a memory pool to the requested handle type. More...
 
hipError_t hipMemPoolImportFromShareableHandle (hipMemPool_t *mem_pool, void *shared_handle, hipMemAllocationHandleType handle_type, unsigned int flags)
 Imports a memory pool from a shared handle. More...
 
hipError_t hipMemPoolExportPointer (hipMemPoolPtrExportData *export_data, void *dev_ptr)
 Export data to share a memory pool allocation between processes. More...
 
hipError_t hipMemPoolImportPointer (void **dev_ptr, hipMemPool_t mem_pool, hipMemPoolPtrExportData *export_data)
 Import a memory pool allocation from another process. More...
 

Detailed Description



This section describes Stream Ordered Memory Allocator functions of HIP runtime API.

The asynchronous allocator allows the user to allocate and free in stream order. All asynchronous accesses of the allocation must happen between the stream executions of the allocation and the free. If the memory is accessed outside of the promised stream order, a use before allocation / use after free error will cause undefined behavior.

The allocator is free to reallocate the memory as long as it can guarantee that compliant memory accesses will not overlap temporally. The allocator may refer to internal stream ordering as well as inter-stream dependencies (such as HIP events and null stream dependencies) when establishing the temporal guarantee. The allocator may also insert inter-stream dependencies to establish the temporal guarantee. Whether or not a device supports the integrated stream ordered memory allocator may be queried by calling hipDeviceGetAttribute with the device attribute hipDeviceAttributeMemoryPoolsSupported

Note
APIs in this section are implemented on Linux, under development on Windows.

Function Documentation

◆ hipFreeAsync()

hipError_t hipFreeAsync ( void *  dev_ptr,
hipStream_t  stream 
)

Frees memory with stream ordered semantics.

Inserts a free operation into stream. The allocation must not be used after stream execution reaches the free. After this API returns, accessing the memory from any subsequent work launched on the GPU or querying its pointer attributes results in undefined behavior.

Note
During stream capture, this function results in the creation of a free node and must therefore be passed the address of a graph allocation.
Parameters
[in]dev_ptrPointer to device memory to free
[in]streamThe stream, where the destruciton will occur according to the execution order
Returns
hipSuccess, hipErrorInvalidValue, hipErrorNotSupported
See also
hipMallocFromPoolAsync, hipMallocAsync, hipMemPoolTrimTo, hipMemPoolGetAttribute, hipDeviceSetMemPool, hipMemPoolSetAttribute, hipMemPoolSetAccess, hipMemPoolGetAccess
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMallocAsync()

hipError_t hipMallocAsync ( void **  dev_ptr,
size_t  size,
hipStream_t  stream 
)

Allocates memory with stream ordered semantics.

Inserts a memory allocation operation into stream. A pointer to the allocated memory is returned immediately in *dptr. The allocation must not be accessed until the allocation operation completes. The allocation comes from the memory pool associated with the stream's device.

Note
The default memory pool of a device contains device memory from that device.
Basic stream ordering allows future work submitted into the same stream to use the allocation. Stream query, stream synchronize, and HIP events can be used to guarantee that the allocation operation completes before work submitted in a separate stream runs.
During stream capture, this function results in the creation of an allocation node. In this case, the allocation is owned by the graph instead of the memory pool. The memory pool's properties are used to set the node's creation parameters.
Parameters
[out]dev_ptrReturned device pointer of memory allocation
[in]sizeNumber of bytes to allocate
[in]streamThe stream establishing the stream ordering contract and the memory pool to allocate from
Returns
hipSuccess, hipErrorInvalidValue, hipErrorNotSupported, hipErrorOutOfMemory
See also
hipMallocFromPoolAsync, hipFreeAsync, hipMemPoolTrimTo, hipMemPoolGetAttribute, hipDeviceSetMemPool, hipMemPoolSetAttribute, hipMemPoolSetAccess, hipMemPoolGetAccess
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMallocFromPoolAsync()

hipError_t hipMallocFromPoolAsync ( void **  dev_ptr,
size_t  size,
hipMemPool_t  mem_pool,
hipStream_t  stream 
)

Allocates memory from a specified pool with stream ordered semantics.

Inserts an allocation operation into stream. A pointer to the allocated memory is returned immediately in dev_ptr. The allocation must not be accessed until the allocation operation completes. The allocation comes from the specified memory pool.

Note
The specified memory pool may be from a device different than that of the specified stream.

Basic stream ordering allows future work submitted into the same stream to use the allocation. Stream query, stream synchronize, and HIP events can be used to guarantee that the allocation operation completes before work submitted in a separate stream runs.

Note
During stream capture, this function results in the creation of an allocation node. In this case, the allocation is owned by the graph instead of the memory pool. The memory pool's properties are used to set the node's creation parameters.
Parameters
[out]dev_ptrReturned device pointer
[in]sizeNumber of bytes to allocate
[in]mem_poolThe pool to allocate from
[in]streamThe stream establishing the stream ordering semantic
Returns
hipSuccess, hipErrorInvalidValue, hipErrorNotSupported, hipErrorOutOfMemory
See also
hipMallocAsync, hipFreeAsync, hipMemPoolGetAttribute, hipMemPoolCreate hipMemPoolTrimTo, hipDeviceSetMemPool, hipMemPoolSetAttribute, hipMemPoolSetAccess, hipMemPoolGetAccess,
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMemPoolCreate()

hipError_t hipMemPoolCreate ( hipMemPool_t mem_pool,
const hipMemPoolProps pool_props 
)

Creates a memory pool.

Creates a HIP memory pool and returns the handle in mem_pool. The pool_props determines the properties of the pool such as the backing device and IPC capabilities.

By default, the memory pool will be accessible from the device it is allocated on.

Parameters
[out]mem_poolContains createed memory pool
[in]pool_propsMemory pool properties
Note
Specifying hipMemHandleTypeNone creates a memory pool that will not support IPC.
Returns
hipSuccess, hipErrorInvalidValue, hipErrorNotSupported
See also
hipMallocFromPoolAsync, hipMallocAsync, hipFreeAsync, hipMemPoolGetAttribute, hipMemPoolDestroy, hipMemPoolTrimTo, hipDeviceSetMemPool, hipMemPoolSetAttribute, hipMemPoolSetAccess, hipMemPoolGetAccess
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMemPoolDestroy()

hipError_t hipMemPoolDestroy ( hipMemPool_t  mem_pool)

Destroys the specified memory pool.

If any pointers obtained from this pool haven't been freed or the pool has free operations that haven't completed when hipMemPoolDestroy is invoked, the function will return immediately and the resources associated with the pool will be released automatically once there are no more outstanding allocations.

Destroying the current mempool of a device sets the default mempool of that device as the current mempool for that device.

Parameters
[in]mem_poolMemory pool for destruction
Note
A device's default memory pool cannot be destroyed.
Returns
hipSuccess, hipErrorInvalidValue
See also
hipMallocFromPoolAsync, hipMallocAsync, hipFreeAsync, hipMemPoolGetAttribute, hipMemPoolCreate hipMemPoolTrimTo, hipDeviceSetMemPool, hipMemPoolSetAttribute, hipMemPoolSetAccess, hipMemPoolGetAccess
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMemPoolExportPointer()

hipError_t hipMemPoolExportPointer ( hipMemPoolPtrExportData export_data,
void *  dev_ptr 
)

Export data to share a memory pool allocation between processes.

Constructs export_data for sharing a specific allocation from an already shared memory pool. The recipient process can import the allocation with the hipMemPoolImportPointer api. The data is not a handle and may be shared through any IPC mechanism.

Parameters
[out]export_dataReturned export data
[in]dev_ptrPointer to memory being exported
Returns
hipSuccess, hipErrorInvalidValue, hipErrorOutOfMemory
See also
hipMemPoolImportPointer
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMemPoolExportToShareableHandle()

hipError_t hipMemPoolExportToShareableHandle ( void *  shared_handle,
hipMemPool_t  mem_pool,
hipMemAllocationHandleType  handle_type,
unsigned int  flags 
)

Exports a memory pool to the requested handle type.

Given an IPC capable mempool, create an OS handle to share the pool with another process. A recipient process can convert the shareable handle into a mempool with hipMemPoolImportFromShareableHandle. Individual pointers can then be shared with the hipMemPoolExportPointer and hipMemPoolImportPointer APIs. The implementation of what the shareable handle is and how it can be transferred is defined by the requested handle type.

Note
: To create an IPC capable mempool, create a mempool with a hipMemAllocationHandleType other than hipMemHandleTypeNone.
Parameters
[out]shared_handlePointer to the location in which to store the requested handle
[in]mem_poolPool to export
[in]handle_typeThe type of handle to create
[in]flagsMust be 0
Returns
hipSuccess, hipErrorInvalidValue, hipErrorOutOfMemory
See also
hipMemPoolImportFromShareableHandle
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMemPoolGetAccess()

hipError_t hipMemPoolGetAccess ( hipMemAccessFlags flags,
hipMemPool_t  mem_pool,
hipMemLocation location 
)

Returns the accessibility of a pool from a device.

Returns the accessibility of the pool's memory from the specified location.

Parameters
[out]flagsAccessibility of the memory pool from the specified location/device
[in]mem_poolMemory pool being queried
[in]locationLocation/device for memory pool access
Returns
hipSuccess, hipErrorInvalidValue
See also
hipMallocFromPoolAsync, hipMallocAsync, hipFreeAsync, hipMemPoolGetAttribute, hipMemPoolTrimTo, hipDeviceSetMemPool, hipMemPoolSetAttribute, hipMemPoolSetAccess
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMemPoolGetAttribute()

hipError_t hipMemPoolGetAttribute ( hipMemPool_t  mem_pool,
hipMemPoolAttr  attr,
void *  value 
)

Gets attributes of a memory pool.

Supported attributes are:

  • hipMemPoolAttrReleaseThreshold: (value type = cuuint64_t) Amount of reserved memory in bytes to hold onto before trying to release memory back to the OS. When more than the release threshold bytes of memory are held by the memory pool, the allocator will try to release memory back to the OS on the next call to stream, event or context synchronize. (default 0)
  • hipMemPoolReuseFollowEventDependencies: (value type = int) Allow hipMallocAsync to use memory asynchronously freed in another stream as long as a stream ordering dependency of the allocating stream on the free action exists. HIP events and null stream interactions can create the required stream ordered dependencies. (default enabled)
  • hipMemPoolReuseAllowOpportunistic: (value type = int) Allow reuse of already completed frees when there is no dependency between the free and allocation. (default enabled)
  • hipMemPoolReuseAllowInternalDependencies: (value type = int) Allow hipMallocAsync to insert new stream dependencies in order to establish the stream ordering required to reuse a piece of memory released by hipFreeAsync (default enabled).
Parameters
[in]mem_poolThe memory pool to get attributes of
[in]attrThe attribute to get
[in]valueRetrieved value
Returns
hipSuccess, hipErrorInvalidValue
See also
hipMallocFromPoolAsync, hipMallocAsync, hipFreeAsync, hipMemPoolTrimTo, hipDeviceSetMemPool, hipMemPoolSetAttribute, hipMemPoolSetAccess, hipMemPoolGetAccess
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMemPoolImportFromShareableHandle()

hipError_t hipMemPoolImportFromShareableHandle ( hipMemPool_t mem_pool,
void *  shared_handle,
hipMemAllocationHandleType  handle_type,
unsigned int  flags 
)

Imports a memory pool from a shared handle.

Specific allocations can be imported from the imported pool with hipMemPoolImportPointer.

Note
Imported memory pools do not support creating new allocations. As such imported memory pools may not be used in hipDeviceSetMemPool or hipMallocFromPoolAsync calls.
Parameters
[out]mem_poolReturned memory pool
[in]shared_handleOS handle of the pool to open
[in]handle_typeThe type of handle being imported
[in]flagsMust be 0
Returns
hipSuccess, hipErrorInvalidValue, hipErrorOutOfMemory
See also
hipMemPoolExportToShareableHandle
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMemPoolImportPointer()

hipError_t hipMemPoolImportPointer ( void **  dev_ptr,
hipMemPool_t  mem_pool,
hipMemPoolPtrExportData export_data 
)

Import a memory pool allocation from another process.

Returns in dev_ptr a pointer to the imported memory. The imported memory must not be accessed before the allocation operation completes in the exporting process. The imported memory must be freed from all importing processes before being freed in the exporting process. The pointer may be freed with hipFree or hipFreeAsync. If hipFreeAsync is used, the free must be completed on the importing process before the free operation on the exporting process.

Note
The hipFreeAsync api may be used in the exporting process before the hipFreeAsync operation completes in its stream as long as the hipFreeAsync in the exporting process specifies a stream with a stream dependency on the importing process's hipFreeAsync.
Parameters
[out]dev_ptrPointer to imported memory
[in]mem_poolMemory pool from which to import a pointer
[in]export_dataData specifying the memory to import
Returns
hipSuccess, hipErrorInvalidValue, hipErrorNotInitialized, hipErrorOutOfMemory
See also
hipMemPoolExportPointer
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMemPoolSetAccess()

hipError_t hipMemPoolSetAccess ( hipMemPool_t  mem_pool,
const hipMemAccessDesc desc_list,
size_t  count 
)

Controls visibility of the specified pool between devices.

Parameters
[in]mem_poolMemory pool for acccess change
[in]desc_listArray of access descriptors. Each descriptor instructs the access to enable for a single gpu
[in]countNumber of descriptors in the map array.
Returns
hipSuccess, hipErrorInvalidValue
See also
hipMallocFromPoolAsync, hipMallocAsync, hipFreeAsync, hipMemPoolGetAttribute, hipMemPoolTrimTo, hipDeviceSetMemPool, hipMemPoolSetAttribute, hipMemPoolGetAccess
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMemPoolSetAttribute()

hipError_t hipMemPoolSetAttribute ( hipMemPool_t  mem_pool,
hipMemPoolAttr  attr,
void *  value 
)

Sets attributes of a memory pool.

Supported attributes are:

  • hipMemPoolAttrReleaseThreshold: (value type = cuuint64_t) Amount of reserved memory in bytes to hold onto before trying to release memory back to the OS. When more than the release threshold bytes of memory are held by the memory pool, the allocator will try to release memory back to the OS on the next call to stream, event or context synchronize. (default 0)
  • hipMemPoolReuseFollowEventDependencies: (value type = int) Allow hipMallocAsync to use memory asynchronously freed in another stream as long as a stream ordering dependency of the allocating stream on the free action exists. HIP events and null stream interactions can create the required stream ordered dependencies. (default enabled)
  • hipMemPoolReuseAllowOpportunistic: (value type = int) Allow reuse of already completed frees when there is no dependency between the free and allocation. (default enabled)
  • hipMemPoolReuseAllowInternalDependencies: (value type = int) Allow hipMallocAsync to insert new stream dependencies in order to establish the stream ordering required to reuse a piece of memory released by hipFreeAsync (default enabled).
Parameters
[in]mem_poolThe memory pool to modify
[in]attrThe attribute to modify
[in]valuePointer to the value to assign
Returns
hipSuccess, hipErrorInvalidValue
See also
hipMallocFromPoolAsync, hipMallocAsync, hipFreeAsync, hipMemPoolGetAttribute, hipMemPoolTrimTo, hipDeviceSetMemPool, hipMemPoolSetAccess, hipMemPoolGetAccess
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.

◆ hipMemPoolTrimTo()

hipError_t hipMemPoolTrimTo ( hipMemPool_t  mem_pool,
size_t  min_bytes_to_hold 
)

Releases freed memory back to the OS.

Releases memory back to the OS until the pool contains fewer than min_bytes_to_keep reserved bytes, or there is no more memory that the allocator can safely release. The allocator cannot release OS allocations that back outstanding asynchronous allocations. The OS allocations may happen at different granularity from the user allocations.

Note
: Allocations that have not been freed count as outstanding.
: Allocations that have been asynchronously freed but whose completion has not been observed on the host (eg. by a synchronize) can count as outstanding.
Parameters
[in]mem_poolThe memory pool to trim allocations
[in]min_bytes_to_holdIf the pool has less than min_bytes_to_hold reserved, then the TrimTo operation is a no-op. Otherwise the memory pool will contain at least min_bytes_to_hold bytes reserved after the operation.
Returns
hipSuccess, hipErrorInvalidValue
See also
hipMallocFromPoolAsync, hipMallocAsync, hipFreeAsync, hipMemPoolGetAttribute, hipDeviceSetMemPool, hipMemPoolSetAttribute, hipMemPoolSetAccess, hipMemPoolGetAccess
Warning
: This API is marked as beta, meaning, while this is feature complete, it is still open to changes and may have outstanding issues.
Note
This API is implemented on Linux, under development on Windows.