Add hipMallocManaged default functional support (#1036)

* Add hipMallocManaged default functional support

* Fix build error

* Add dtest

include/hip/hcc_detail/hip_runtime_api.h

@@ -177,6 +177,9 @@ enum hipLimit_t {
     0x80000000  ///< Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific
                 ///< allocation.
 
+#define hipMemAttachGlobal 0x0
+#define hipMemAttachHost 0x1
+
 #define hipDeviceMallocDefault 0x0
 #define hipDeviceMallocFinegrained 0x1  ///< Memory is allocated in fine grained region of device.
 
@@ -1103,6 +1106,17 @@ hipError_t hipMallocHost(void** ptr, size_t size);
  */
 hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags);
 
+/**
+ *  @brief Allocates memory that will be automatically managed by the Unified Memory system.
+ *
+ *  @param[out] ptr Pointer to the allocated managed memory
+ *  @param[in]  size Requested memory size
+ *  @param[in]  flags must be either hipMemAttachGlobal/hipMemAttachHost
+ *
+ *  @return #hipSuccess, #hipErrorMemoryAllocation
+ */
+hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags __dparm(0));
+
 /**
  *  @brief Allocate device accessible page locked host memory [Deprecated]
  *

include/hip/hip_runtime_api.h

@@ -337,6 +337,12 @@ static inline hipError_t hipHostMalloc(T** ptr, size_t size,
                                        unsigned int flags = hipHostMallocDefault) {
     return hipHostMalloc((void**)ptr, size, flags);
 }
+
+template <class T>
+static inline hipError_t hipMallocManaged(T** devPtr, size_t size,
+                                       unsigned int flags = hipMemAttachGlobal) {
+    return hipMallocManaged((void**)devPtr, size, flags);
+}
 #endif
 
 #endif

include/hip/nvcc_detail/hip_runtime_api.h

@@ -111,6 +111,9 @@ typedef enum hipChannelFormatKind {
 #define hipHostMallocCoherent 0x0
 #define hipHostMallocNonCoherent 0x0
 
+#define hipMemAttachGlobal cudaMemAttachGlobal
+#define hipMemAttachHost cudaMemAttachHost
+
 #define hipHostRegisterDefault cudaHostRegisterDefault
 #define hipHostRegisterPortable cudaHostRegisterPortable
 #define hipHostRegisterMapped cudaHostRegisterMapped
@@ -420,6 +423,10 @@ inline static hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int fla
     return hipCUDAErrorTohipError(cudaHostAlloc(ptr, size, flags));
 }
 
+inline static hipError_t hipMallocManaged(void** ptr, size_t size, unsigned int flags) {
+    return hipCUDAErrorTohipError(cudaMallocManaged(ptr, size, flags));
+}
+
 inline static hipError_t hipMallocArray(hipArray** array, const struct hipChannelFormatDesc* desc,
                                         size_t width, size_t height,
                                         unsigned int flags __dparm(hipArrayDefault)) {