Add hipMallocManaged default functional support (#1036)

* Add hipMallocManaged default functional support

* Fix build error

* Add dtest


[ROCm/clr commit: 94769fc8dd]

projects/clr/hipamd/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]
  *

projects/clr/hipamd/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

projects/clr/hipamd/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)) {

projects/clr/hipamd/src/hip_memory.cpp

@@ -301,9 +301,7 @@ hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flag
     return ihipLogStatus(hip_status);
 }
 
-hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
-    HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags);
-    HIP_SET_DEVICE();
+hipError_t ihipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
     hipError_t hip_status = hipSuccess;
 
     if (HIP_SYNC_HOST_ALLOC) {
@@ -368,6 +366,25 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
     if (HIP_SYNC_HOST_ALLOC) {
         hipDeviceSynchronize();
     }
+    return hip_status;
+}
+
+hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
+    HIP_INIT_SPECIAL_API(hipHostMalloc, (TRACE_MEM), ptr, sizeBytes, flags);
+    HIP_SET_DEVICE();
+    hipError_t hip_status = hipSuccess;
+    hip_status = ihipHostMalloc(ptr, sizeBytes, flags);
+    return ihipLogStatus(hip_status);
+}
+
+hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags) {
+    HIP_INIT_SPECIAL_API(hipMallocManaged, (TRACE_MEM), devPtr, size, flags);
+    HIP_SET_DEVICE();
+    hipError_t hip_status = hipSuccess;
+    if(flags != hipMemAttachGlobal)
+        hip_status = hipErrorInvalidValue;
+    else
+        hip_status = ihipHostMalloc(devPtr, size, hipHostMallocDefault);
     return ihipLogStatus(hip_status);
 }
 

projects/clr/hipamd/tests/src/runtimeApi/memory/hipMallocManaged.cpp

@@ -0,0 +1,52 @@
+#include <hip/hip_runtime.h>
+#include <math.h>
+#include "test_common.h"
+
+/* HIT_START
+ * BUILD: %t %s ../../test_common.cpp
+ * RUN: %t -N 256M
+ * HIT_END
+ */
+
+__global__
+void add(int n, float *x, float *y)
+{
+    int index = blockIdx.x * blockDim.x + threadIdx.x;
+    int stride = blockDim.x * gridDim.x;
+    for (int i = index; i < n; i += stride)
+        y[i] = x[i] + y[i];
+}
+
+int main(int argc, char *argv[])
+{
+    HipTest::parseStandardArguments(argc, argv, true);
+    int numElements = N;
+    bool testResult = true;
+    float *A, *B;
+
+    hipMallocManaged(&A, numElements*sizeof(float));
+    hipMallocManaged(&B, numElements*sizeof(float));
+
+    for (int i = 0; i < numElements; i++) {
+        A[i] = 1.0f;
+        B[i] = 2.0f;
+    }
+
+    int blockSize = 256;
+    int numBlocks = (numElements + blockSize - 1) / blockSize;
+    dim3 dimGrid(numBlocks, 1, 1);
+    dim3 dimBlock(blockSize, 1, 1);
+    hipLaunchKernelGGL(add, dimGrid, dimBlock, 0, 0, numElements, A, B);
+
+    hipDeviceSynchronize();
+
+    float maxError = 0.0f;
+    for (int i = 0; i < numElements; i++)
+        maxError = fmax(maxError, fabs(B[i]-3.0f));
+
+    hipFree(A);
+    hipFree(B);
+    if(maxError == 0.0f)
+        passed();
+    failed("Output Mismatch\n");
+}