SWDEV-470372 - Added hipExtHostAlloc API

This change adds a new HIP API `hipExtHostAlloc` which preserves
the functionality of `hipHostMalloc`.

Change-Id: I2198ecd7a10a1ceb28b8fd32b16625b1b6d3abaf

docs/reference/deprecated_api_list.rst

@@ -43,8 +43,8 @@ These APIs are now marked as deprecated because there are better alternate inter
 Memory management
 ============================================================
 
-* ``hipMallocHost`` (replaced with ``hipHostMalloc``)
-* ``hipMemAllocHost`` (replaced with ``hipHostMalloc``)
+* ``hipMallocHost`` (replaced with ``hipHostAlloc``)
+* ``hipMemAllocHost`` (replaced with ``hipHostAlloc``)
 * ``hipMemcpyToArray``
 * ``hipMemcpyFromArray``
 

include/hip/hip_runtime_api.h

@@ -701,7 +701,7 @@ enum hipLimit_t {
  * non-coherent host memory visible to the host. The flag is a no-op on CUDA platforms.*/
 #define hipEventReleaseToSystem  0x80000000
 
-//Flags that can be used with hipHostMalloc/hipHostAlloc.
+//Flags that can be used with hipExtHostAlloc/hipHostAlloc.
 /** Default pinned memory allocation on the host.*/
 #define hipHostMallocDefault 0x0
 #define hipHostAllocDefault 0x0
@@ -726,12 +726,15 @@ enum hipLimit_t {
 * @note  This numa allocation flag is applicable on Linux, under development on Windows.
 */
 #define hipHostMallocNumaUser  0x20000000
+#define hipExtHostAllocNumaUser  0x20000000
 
 /** Allocate coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific allocation.*/
 #define hipHostMallocCoherent  0x40000000
+#define hipExtHostAllocCoherent  0x40000000
 
 /** Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific allocation.*/
 #define hipHostMallocNonCoherent  0x80000000
+#define hipExtHostAllocNonCoherent  0x80000000
 
 /** Memory can be accessed by any stream on any device*/
 #define hipMemAttachGlobal  0x01
@@ -3352,7 +3355,7 @@ hipError_t hipExternalMemoryGetMappedMipmappedArray(hipMipmappedArray_t* mipmap,
  *  @return #hipSuccess, #hipErrorOutOfMemory, #hipErrorInvalidValue (bad context, null *ptr)
  *
  *  @see hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D, hipMalloc3DArray,
- * hipHostFree, hipHostMalloc
+ * hipFreeHost, hipHostAlloc
  */
 hipError_t hipMalloc(void** ptr, size_t size);
 /**
@@ -3372,7 +3375,7 @@ hipError_t hipMalloc(void** ptr, size_t size);
  *  @return #hipSuccess, #hipErrorOutOfMemory, #hipErrorInvalidValue (bad context, null *ptr)
  *
  *  @see hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D, hipMalloc3DArray,
- * hipHostFree, hipHostMalloc
+ * hipFreeHost, hipExtHostMalloc
  */
 hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags);
 
@@ -3397,9 +3400,9 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag
  *
  *  @return #hipSuccess, #hipErrorOutOfMemory
  *
- *  @warning  This API is deprecated, use hipHostMalloc() instead
+ *  @warning  This API is deprecated, use hipHostAlloc() instead
  */
-DEPRECATED("use hipHostMalloc instead")
+DEPRECATED("use hipHostAlloc instead")
 hipError_t hipMallocHost(void** ptr, size_t size);
 /**
  *  @brief Allocate pinned host memory [Deprecated]
@@ -3411,9 +3414,9 @@ hipError_t hipMallocHost(void** ptr, size_t size);
  *
  *  @return #hipSuccess, #hipErrorOutOfMemory
  *
- *  @warning  This API is deprecated, use hipHostMalloc() instead
+ *  @warning  This API is deprecated, use hipHostAlloc() instead
  */
-DEPRECATED("use hipHostMalloc instead")
+DEPRECATED("use hipHostAlloc instead")
 hipError_t hipMemAllocHost(void** ptr, size_t size);
 // end doxygen deprecated management memory
 /**
@@ -3449,9 +3452,42 @@ hipError_t hipMemAllocHost(void** ptr, size_t size);
  *
  *  @return #hipSuccess, #hipErrorOutOfMemory
  *
- *  @see hipSetDeviceFlags, hipHostFree
+ *
+ *  @see hipSetDeviceFlags, hiptFreeHost
  */
 hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags);
+/**
+ *  @brief Allocates device accessible page locked (pinned) host memory
+ *
+ *  This API allocates pinned host memory which is mapped into the address space of all GPUs
+ *  in the system, the memory can be accessed directly by the GPU device, and can be read or
+ *  written with much higher bandwidth than pageable memory obtained with functions such as
+ *  malloc().
+ *
+ *  Using the pinned host memory, applications can implement faster data transfers for HostToDevice
+ *  and DeviceToHost. The runtime tracks the hipExtHostAlloc allocations and can avoid some of the
+ *  setup required for regular unpinned memory.
+ *
+ *  When the memory accesses are infrequent, zero-copy memory can be a good choice, for coherent
+ *  allocation. GPU can directly access the host memory over the CPU/GPU interconnect, without need
+ *  to copy the data.
+ *
+ *  Currently the allocation granularity is 4KB for the API.
+ *
+ *  Developers need to choose proper allocation flag with consideration of synchronization.
+ *
+ *  @param[out] ptr Pointer to the allocated host pinned memory
+ *  @param[in]  size Requested memory size in bytes
+ *  If size is 0, no memory is allocated, *ptr returns nullptr, and hipSuccess is returned.
+ *  @param[in]  flags Type of host memory allocation
+ *
+ *  If no input for flags, it will be the default pinned memory allocation on the host.
+ *
+ *  @return #hipSuccess, #hipErrorOutOfMemory
+ *
+ *  @see hipSetDeviceFlags, hipFreeHost
+ */
+hipError_t hipExtHostAlloc(void** ptr, size_t size, unsigned int flags);
 /**
  *-------------------------------------------------------------------------------------------------
  *-------------------------------------------------------------------------------------------------
@@ -4035,25 +4071,25 @@ hipError_t hipMemPoolImportPointer(
  */
 hipError_t hipHostAlloc(void** ptr, size_t size, unsigned int flags);
 /**
- *  @brief Get Device pointer from Host Pointer allocated through hipHostMalloc
+ *  @brief Get Device pointer from Host Pointer allocated through hipHostAlloc
  *
  *  @param[out] devPtr Device Pointer mapped to passed host pointer
- *  @param[in]  hstPtr Host Pointer allocated through hipHostMalloc
+ *  @param[in]  hstPtr Host Pointer allocated through hipHostAlloc
  *  @param[in]  flags Flags to be passed for extension
  *
  *  @return #hipSuccess, #hipErrorInvalidValue, #hipErrorOutOfMemory
  *
- *  @see hipSetDeviceFlags, hipHostMalloc
+ *  @see hipSetDeviceFlags, hipHostAlloc
  */
 hipError_t hipHostGetDevicePointer(void** devPtr, void* hstPtr, unsigned int flags);
 /**
  *  @brief Return flags associated with host pointer
  *
  *  @param[out] flagsPtr Memory location to store flags
- *  @param[in]  hostPtr Host Pointer allocated through hipHostMalloc
+ *  @param[in]  hostPtr Host Pointer allocated through hipHostAlloc
  *  @return #hipSuccess, #hipErrorInvalidValue
  *
- *  @see hipHostMalloc
+ *  @see hipHostAlloc
  */
 hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr);
 /**
@@ -4117,8 +4153,8 @@ hipError_t hipHostUnregister(void* hostPtr);
  *
  *  @return Error code
  *
- *  @see hipMalloc, hipFree, hipMallocArray, hipFreeArray, hipHostFree, hipMalloc3D,
- * hipMalloc3DArray, hipHostMalloc
+ *  @see hipMalloc, hipFree, hipMallocArray, hipFreeArray, hipFreeHost, hipMalloc3D,
+ * hipMalloc3DArray, hipHostAlloc
  */
 hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height);
 /**
@@ -4140,8 +4176,8 @@ hipError_t hipMallocPitch(void** ptr, size_t* pitch, size_t width, size_t height
  *
  *  @return Error code
  *
- *  @see hipMalloc, hipFree, hipMallocArray, hipFreeArray, hipHostFree, hipMalloc3D,
- * hipMalloc3DArray, hipHostMalloc
+ *  @see hipMalloc, hipFree, hipMallocArray, hipFreeArray, hiptFreeHost, hipMalloc3D,
+ * hipMalloc3DArray, hipHostAlloc
  */
 hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInBytes, size_t height,
    unsigned int elementSizeBytes);
@@ -4153,14 +4189,16 @@ hipError_t hipMemAllocPitch(hipDeviceptr_t* dptr, size_t* pitch, size_t widthInB
  *  @param[in] ptr Pointer to memory to be freed
  *  @return #hipSuccess
  *  @return #hipErrorInvalidDevicePointer (if pointer is invalid, including host pointers allocated
- * with hipHostMalloc)
+ * with hipHostAlloc)
  *
- *  @see hipMalloc, hipMallocPitch, hipMallocArray, hipFreeArray, hipHostFree, hipMalloc3D,
- * hipMalloc3DArray, hipHostMalloc
+ *  @see hipMalloc, hipMallocPitch, hipMallocArray, hipFreeArray, hipFreeHost, hipMalloc3D,
+ * hipMalloc3DArray, hipHostAlloc
  */
 hipError_t hipFree(void* ptr);
 /**
  *  @brief Frees page-locked memory
+ *  This API performs an implicit hipDeviceSynchronize() call.
+ *  If pointer is NULL, the hip runtime is initialized and hipSuccess is returned.
  *
  *  @ingroup MemoryD
  *
@@ -4182,7 +4220,8 @@ hipError_t hipFreeHost(void* ptr);
  * hipMalloc)
  *
  *  @see hipMalloc, hipMallocPitch, hipFree, hipMallocArray, hipFreeArray, hipMalloc3D,
- * hipMalloc3DArray, hipHostMalloc
+ * hipMalloc3DArray, hipHostAlloc
+ * 
  */
 hipError_t hipHostFree(void* ptr);
 /**
@@ -4588,7 +4627,7 @@ hipError_t hipMemcpyFromSymbolAsync(void* dst, const void* symbol,
  *  @brief Copy data from src to dst asynchronously.
  *
  *  @warning If host or dest are not pinned, the memory copy will be performed synchronously.  For
- * best performance, use hipHostMalloc to allocate host memory that is transferred asynchronously.
+ * best performance, use hipHostAlloc to allocate host memory that is transferred asynchronously.
  *
  *  @warning on HCC hipMemcpyAsync does not support overlapped H2D and D2H copies.
  *  For hipMemcpy, the copy is always performed by the device associated with the specified stream.
@@ -4802,7 +4841,7 @@ hipError_t hipMemPtrGetInfo(void* ptr, size_t* size);
  *  @param[in]   flags  Requested properties of allocated array
  *  @return      #hipSuccess, #hipErrorOutOfMemory
  *
- *  @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostMalloc, hipHostFree
+ *  @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostAlloc, hipFreeHost
  */
 hipError_t hipMallocArray(hipArray_t* array, const hipChannelFormatDesc* desc, size_t width,
                           size_t height __dparm(0), unsigned int flags __dparm(hipArrayDefault));
@@ -4855,7 +4894,7 @@ hipError_t hipMalloc3D(hipPitchedPtr* pitchedDevPtr, hipExtent extent);
  *  @param[in]  array  Pointer to array to free
  *  @return     #hipSuccess, #hipErrorInvalidValue, #hipErrorNotInitialized
  *
- *  @see hipMalloc, hipMallocPitch, hipFree, hipMallocArray, hipHostMalloc, hipHostFree
+ *  @see hipMalloc, hipMallocPitch, hipFree, hipMallocArray, hipHostAlloc, hipFreeHost
  */
 hipError_t hipFreeArray(hipArray_t array);
 /**
@@ -4867,7 +4906,7 @@ hipError_t hipFreeArray(hipArray_t array);
  *  @param[in]   flags  Requested properties of allocated array
  *  @return      #hipSuccess, #hipErrorOutOfMemory
  *
- *  @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostMalloc, hipHostFree
+ *  @see hipMalloc, hipMallocPitch, hipFree, hipFreeArray, hipHostAlloc, hipFreeHost
  */
 hipError_t hipMalloc3DArray(hipArray_t* array, const struct hipChannelFormatDesc* desc,
                             struct hipExtent extent, unsigned int flags);
@@ -9585,6 +9624,23 @@ template <class T>
 static inline hipError_t hipMallocPitch(T** devPtr, size_t* pitch, size_t width, size_t height) {
     return hipMallocPitch((void**)devPtr, pitch, width, height);
 }
+/**
+ * @brief: C++ wrapper for hipExtHostAlloc
+ * @ingroup Memory
+ * Provide an override to automatically typecast the pointer type from void**, and also provide a
+ * default for the flags.
+ *
+ * __HIP_DISABLE_CPP_FUNCTIONS__ macro can be defined to suppress these
+ * wrappers. It is useful for applications which need to obtain decltypes of
+ * HIP runtime APIs.
+ *
+ * @see hipExtHostAlloc
+ */
+template <class T>
+static inline hipError_t hipExtHostAlloc(T** ptr, size_t size,
+                                       unsigned int flags = hipHostAllocDefault) {
+    return hipExtHostAlloc((void**)ptr, size, flags);
+}
 /**
  * @brief: C++ wrapper for hipHostMalloc
  * @ingroup Memory
@@ -9599,9 +9655,10 @@ static inline hipError_t hipMallocPitch(T** devPtr, size_t* pitch, size_t width,
  */
 template <class T>
 static inline hipError_t hipHostMalloc(T** ptr, size_t size,
-                                       unsigned int flags = hipHostMallocDefault) {
+                                       unsigned int flags = hipHostAllocDefault) {
     return hipHostMalloc((void**)ptr, size, flags);
 }
+
 /**
  * @brief: C++ wrapper for hipHostAlloc
  * @ingroup Memory

util/vim/hip.vim

@@ -105,6 +105,7 @@ syn keyword hipFunctionName hipEventSynchronize
 syn keyword hipFunctionName hipFree 
 syn keyword hipFunctionName hipFreeArray 
 syn keyword hipFunctionName hipHostMalloc
+syn keyword hipFunctionName hipExtHostAlloc
 syn keyword hipFunctionName hipHostAlloc
 syn keyword hipFunctionName hipHostFree
 syn keyword hipFunctionName hipFreeHost
@@ -189,6 +190,15 @@ syn keyword hipFlags hipHostMallocMapped
 syn keyword hipFlags hipHostMallocWriteCombined
 syn keyword hipFlags hipHostMallocCoherent
 syn keyword hipFlags hipHostMallocNonCoherent
+syn keyword hipFlags hipHostMallocNumaUser
+
+syn keyword hipFlags hipHostAllocDefault
+syn keyword hipFlags hipHostAllocPortable
+syn keyword hipFlags hipHostAllocMapped
+syn keyword hipFlags hipHostAllocWriteCombined
+syn keyword hipFlags hipExtHostAllocCoherent
+syn keyword hipFlags hipExtHostAllocNonCoherent
+syn keyword hipFlags hipExtHostAllocNumaUser
 
 syn keyword hipFlags hipHostRegisterDefault
 syn keyword hipFlags hipHostRegisterPortable