Initial support for HIP managed memory

- Call the new ROCclr interfaces for HMM

Change-Id: I2cd1bf438f712a9e9e328340e7d0c025257ca6c1


[ROCm/clr commit: b182ad9574]

projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h

@@ -178,8 +178,10 @@ enum hipLimit_t {
     0x80000000  ///< Allocate non-coherent memory. Overrides HIP_COHERENT_HOST_ALLOC for specific
                 ///< allocation.
 
-#define hipMemAttachGlobal 0x0
-#define hipMemAttachHost 0x1
+#define hipMemAttachGlobal  0x01    ///< Memory can be accessed by any stream on any device
+#define hipMemAttachHost    0x02    ///< Memory cannot be accessed by any stream on any device
+#define hipMemAttachSingle  0x04    ///< Memory can only be accessed by a single stream on
+                                    ///< the associated device
 
 #define hipDeviceMallocDefault 0x0
 #define hipDeviceMallocFinegrained 0x1  ///< Memory is allocated in fine grained region of device.
@@ -217,6 +219,41 @@ enum hipLimit_t {
 #define hipCooperativeLaunchMultiDeviceNoPreSync 0x01
 #define hipCooperativeLaunchMultiDeviceNoPostSync 0x02
 
+#define hipCpuDeviceId ((int)-1)
+#define hipInvalidDeviceId ((int)-2)
+
+/*
+ * @brief HIP Memory Advise values
+ * @enum
+ * @ingroup Enumerations
+ */
+typedef enum hipMemoryAdvise {
+    hipMemAdviseSetReadMostly = 1,          ///< Data will mostly be read and only occassionally
+                                            ///< be written to
+    hipMemAdviseUnsetReadMostly = 2,        ///< Undo the effect of hipMemAdviseSetReadMostly
+    hipMemAdviseSetPreferredLocation = 3,   ///< Set the preferred location for the data as
+                                            ///< the specified device
+    hipMemAdviseUnsetPreferredLocation = 4, ///< Clear the preferred location for the data
+    hipMemAdviseSetAccessedBy = 5,          ///< Data will be accessed by the specified device,
+                                            ///< so prevent page faults as much as possible
+    hipMemAdviseUnsetAccessedBy = 6         ///< Let the Unified Memory subsystem decide on
+                                            ///< the page faulting policy for the specified device
+} hipMemoryAdvise;
+
+/*
+ * @brief HIP range attributes
+ * @enum
+ * @ingroup Enumerations
+ */
+typedef enum hipMemRangeAttribute {
+    hipMemRangeAttributeReadMostly = 1,         ///< Whether the range will mostly be read and
+                                                ///< only occassionally be written to
+    hipMemRangeAttributePreferredLocation = 2,  ///< The preferred location of the range
+    hipMemRangeAttributeAccessedBy = 3,         ///< Memory range has cudaMemAdviseSetAccessedBy
+                                                ///< set for specified device
+    hipMemRangeAttributeLastPrefetchLocation = 4,///< The last location to which the range was prefetched
+} hipMemRangeAttribute;
+
 /*
  * @brief hipJitOption
  * @enum
@@ -1180,15 +1217,18 @@ hipError_t hipMemAllocHost(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.
+ * @brief Allocates memory that will be automatically managed by AMD HMM.
  *
- *  @param[out] ptr Pointer to the allocated managed memory
- *  @param[in]  size Requested memory size
- *  @param[in]  flags must be either hipMemAttachGlobal/hipMemAttachHost
+ * @param [out] dev_ptr - pointer to allocated device memory
+ * @param [in]  size    - requested allocation size in bytes
+ * @param [in]  flags   - must be either hipMemAttachGlobal or hipMemAttachHost
+ *                        (defaults to hipMemAttachGlobal)
  *
- *  @return #hipSuccess, #hipErrorOutOfMemory
+ * @returns #hipSuccess, #hipErrorMemoryAllocation, #hipErrorNotSupported, #hipErrorInvalidValue
  */
-hipError_t hipMallocManaged(void** devPtr, size_t size, unsigned int flags __dparm(0));
+hipError_t hipMallocManaged(void** dev_ptr,
+                            size_t size,
+                            unsigned int flags __dparm(hipMemAttachGlobal));
 
 /**
  *  @brief Allocate device accessible page locked host memory [Deprecated]
@@ -3369,7 +3409,6 @@ hipError_t __hipPopCallConfiguration(dim3 *gridDim,
  * @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice
  *
  */
-
 hipError_t hipLaunchKernel(const void* function_address,
                            dim3 numBlocks,
                            dim3 dimBlocks,
@@ -3377,6 +3416,92 @@ hipError_t hipLaunchKernel(const void* function_address,
                            size_t sharedMemBytes __dparm(0),
                            hipStream_t stream __dparm(0));
 
+/**
+ * @brief Prefetches memory to the specified destination device using AMD HMM.
+ *
+ * @param [in] dev_ptr  pointer to be prefetched
+ * @param [in] count    size in bytes for prefetching
+ * @param [in] device   destination device to prefetch to
+ * @param [in] stream   stream to enqueue prefetch operation
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue
+ */
+hipError_t hipMemPrefetchAsync(const void* dev_ptr,
+                               size_t count,
+                               int device,
+                               hipStream_t stream __dparm(0));
+
+/**
+ * @brief Advise about the usage of a given memory range to AMD HMM.
+ *
+ * @param [in] dev_ptr  pointer to memory to set the advice for
+ * @param [in] count    size in bytes of the memory range
+ * @param [in] advice   advice to be applied for the specified memory range
+ * @param [in] device   device to apply the advice for
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue
+ */
+hipError_t hipMemAdvise(const void* dev_ptr,
+                        size_t count,
+                        hipMemoryAdvise advice,
+                        int device);
+
+/**
+ * @brief Query an attribute of a given memory range in AMD HMM.
+ *
+ * @param [in/out] data   a pointer to a memory location where the result of each
+ *                        attribute query will be written to
+ * @param [in] data_size  the size of data
+ * @param [in] attribute  the attribute to query
+ * @param [in] dev_ptr    start of the range to query
+ * @param [in] count      size of the range to query
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue
+ */
+hipError_t hipMemRangeGetAttribute(void* data,
+                                   size_t data_size,
+                                   hipMemRangeAttribute attribute,
+                                   const void* dev_ptr,
+                                   size_t count);
+
+/**
+ * @brief Query attributes of a given memory range in AMD HMM.
+ *
+ * @param [in/out] data     a two-dimensional array containing pointers to memory locations
+ *                          where the result of each attribute query will be written to
+ * @param [in] data_sizes   an array, containing the sizes of each result
+ * @param [in] attributes   the attribute to query
+ * @param [in] num_attributes  an array of attributes to query (numAttributes and the number
+ *                          of attributes in this array should match)
+ * @param [in] dev_ptr      start of the range to query
+ * @param [in] count        size of the range to query
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue
+ */
+hipError_t hipMemRangeGetAttributes(void** data,
+                                    size_t* data_sizes,
+                                    hipMemRangeAttribute* attributes,
+                                    size_t num_attributes,
+                                    const void* dev_ptr,
+                                    size_t count);
+
+/**
+ * @brief Attach memory to a stream asynchronously in AMD HMM.
+ *
+ * @param [in] stream     - stream in which to enqueue the attach operation
+ * @param [in] dev_ptr    - pointer to memory (must be a pointer to managed memory or
+ *                          to a valid host-accessible region of system-allocated memory)
+ * @param [in] length     - length of memory (defaults to zero)
+ * @param [in] flags      - must be one of cudaMemAttachGlobal, cudaMemAttachHost or
+ *                          cudaMemAttachSingle (defaults to cudaMemAttachSingle)
+ *
+ * @returns #hipSuccess, #hipErrorInvalidValue
+ */
+hipError_t hipStreamAttachMemAsync(hipStream_t stream,
+                                   hipDeviceptr_t* dev_ptr,
+                                   size_t length __dparm(0),
+                                   unsigned int flags __dparm(hipMemAttachSingle));
+
 #if __HIP_ROCclr__ || !defined(__HCC__)
 //TODO: Move this to hip_ext.h
 hipError_t hipExtLaunchKernel(const void* function_address, dim3 numBlocks, dim3 dimBlocks,

projects/clr/hipamd/rocclr/CMakeLists.txt

@@ -81,6 +81,7 @@ add_library(hip64 OBJECT
  hip_error.cpp
  hip_event.cpp
  hip_global.cpp
+ hip_hmm.cpp
  hip_memory.cpp
  hip_module.cpp
  hip_peer.cpp

projects/clr/hipamd/rocclr/hip_hcc.def.in

@@ -80,6 +80,7 @@ hipMallocManaged
 hipArrayCreate
 hipArray3DCreate
 hipMallocArray
+hipMemAdvise
 hipMemAllocPitch
 hipMallocPitch
 hipMemcpy
@@ -111,7 +112,10 @@ hipMemGetAddressRange
 hipGetSymbolAddress
 hipGetSymbolSize
 hipMemGetInfo
+hipMemPrefetchAsync
 hipMemPtrGetInfo
+hipMemRangeGetAttribute
+hipMemRangeGetAttributes
 hipMemset
 hipMemsetAsync
 hipMemsetD8
@@ -154,6 +158,7 @@ hipGetDeviceFlags
 hipSetDevice
 hipSetDeviceFlags
 hipStreamAddCallback
+hipStreamAttachMemAsync
 hipStreamCreate
 hipStreamCreateWithFlags
 hipStreamCreateWithPriority

projects/clr/hipamd/rocclr/hip_hcc.map.in

@@ -62,6 +62,7 @@ global:
     hipGetErrorName;
     hipGetErrorString;
     hipGetLastError;
+    hipMemAdvise;
     hipMemAllocHost;
     hipHostAlloc;
     hipHostFree;
@@ -112,7 +113,10 @@ global:
     hipGetSymbolAddress;
     hipGetSymbolSize;
     hipMemGetInfo;
+    hipMemPrefetchAsync;
     hipMemPtrGetInfo;
+    hipMemRangeGetAttribute;
+    hipMemRangeGetAttributes;
     hipMemset;
     hipMemsetAsync;
     hipMemsetD8;
@@ -154,6 +158,7 @@ global:
     hipSetDevice;
     hipSetDeviceFlags;
     hipStreamAddCallback;
+    hipStreamAttachMemAsync;
     hipStreamCreate;
     hipStreamCreateWithFlags;
     hipStreamCreateWithPriority;

projects/clr/hipamd/rocclr/hip_hmm.cpp

@@ -0,0 +1,215 @@
+/* Copyright (c) 2020-present Advanced Micro Devices, Inc.
+
+ Permission is hereby granted, free of charge, to any person obtaining a copy
+ of this software and associated documentation files (the "Software"), to deal
+ in the Software without restriction, including without limitation the rights
+ to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+ copies of the Software, and to permit persons to whom the Software is
+ furnished to do so, subject to the following conditions:
+
+ The above copyright notice and this permission notice shall be included in
+ all copies or substantial portions of the Software.
+
+ THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+ IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+ FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+ AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+ THE SOFTWARE. */
+
+#include <hip/hip_runtime.h>
+#include "hip_internal.hpp"
+#include "hip_conversions.hpp"
+#include "platform/context.hpp"
+#include "platform/command.hpp"
+#include "platform/memory.hpp"
+
+// Forward declaraiton of a static function
+static hipError_t ihipMallocManaged(void** ptr, size_t size);
+
+// Make sure HIP defines match ROCclr to avoid double conversion
+static_assert(hipCpuDeviceId == amd::CpuDeviceId, "CPU device ID mismatch with ROCclr!");
+static_assert(hipInvalidDeviceId == amd::InvalidDeviceId,
+              "Invalid device ID mismatch with ROCclr!");
+
+static_assert(static_cast<uint32_t>(hipMemAdviseSetReadMostly) ==
+              amd::MemoryAdvice::SetReadMostly, "Enum mismatch with ROCclr!");
+static_assert(static_cast<uint32_t>(hipMemAdviseUnsetReadMostly) ==
+              amd::MemoryAdvice::UnsetReadMostly, "Enum mismatch with ROCclr!");
+static_assert(static_cast<uint32_t>(hipMemAdviseSetPreferredLocation) ==
+              amd::MemoryAdvice::SetPreferredLocation, "Enum mismatch with ROCclr!");
+static_assert(static_cast<uint32_t>(hipMemAdviseUnsetPreferredLocation) ==
+              amd::MemoryAdvice::UnsetPreferredLocation, "Enum mismatch with ROCclr!");
+static_assert(static_cast<uint32_t>(hipMemAdviseSetAccessedBy) ==
+              amd::MemoryAdvice::SetAccessedBy, "Enum mismatch with ROCclr!");
+static_assert(static_cast<uint32_t>(hipMemAdviseUnsetAccessedBy) ==
+              amd::MemoryAdvice::UnsetAccessedBy, "Enum mismatch with ROCclr!");
+
+static_assert(static_cast<uint32_t>(hipMemRangeAttributeReadMostly) ==
+              amd::MemRangeAttribute::ReadMostly, "Enum mismatch with ROCclr!");
+static_assert(static_cast<uint32_t>(hipMemRangeAttributePreferredLocation) ==
+              amd::MemRangeAttribute::PreferredLocation, "Enum mismatch with ROCclr!");
+static_assert(static_cast<uint32_t>(hipMemRangeAttributeAccessedBy) ==
+              amd::MemRangeAttribute::AccessedBy, "Enum mismatch with ROCclr!");
+static_assert(static_cast<uint32_t>(hipMemRangeAttributeLastPrefetchLocation) ==
+              amd::MemRangeAttribute::LastPrefetchLocation, "Enum mismatch with ROCclr!");
+
+// ================================================================================================
+hipError_t hipMallocManaged(void** dev_ptr, size_t size, unsigned int flags) {
+  HIP_INIT_API(hipMallocManaged, dev_ptr, size, flags);
+
+  if ((dev_ptr == nullptr) || (flags != hipMemAttachGlobal)) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  HIP_RETURN(ihipMallocManaged(dev_ptr, size), *dev_ptr);
+}
+
+// ================================================================================================
+hipError_t hipMemPrefetchAsync(const void* dev_ptr, size_t count, int device,
+                               hipStream_t stream) {
+  HIP_INIT_API(hipMemPrefetchAsync, dev_ptr, count, device, stream);
+
+  if ((dev_ptr == nullptr) || (count == 0) || (stream == nullptr)) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  amd::HostQueue* queue = nullptr;
+  bool cpu_access = (device == hipCpuDeviceId) ? true : false;
+
+  // Pick the specified stream or Null one from the provided device
+  if (stream != nullptr) {
+    queue = hip::getQueue(stream);
+  } else {
+    if (!cpu_access) {
+      queue = g_devices[device]->NullStream();
+    } else {
+      queue = hip::getCurrentDevice()->NullStream();
+    }
+  }
+  if (queue == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  amd::Command::EventWaitList waitList;
+  amd::SvmPrefetchAsyncCommand* command =
+      new amd::SvmPrefetchAsyncCommand(*queue, waitList, dev_ptr, count, cpu_access);
+  if (command == nullptr) {
+    return hipErrorOutOfMemory;
+  }
+
+  if (!command->validateMemory()) {
+    delete command;
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  command->enqueue();
+  command->release();
+
+  HIP_RETURN(hipErrorInvalidValue);
+}
+
+// ================================================================================================
+hipError_t hipMemAdvise(const void* dev_ptr, size_t count, hipMemoryAdvise advice, int device) {
+  HIP_INIT_API(hipMemAdvise, dev_ptr, count, advice, device);
+
+  if ((dev_ptr == nullptr) || (count == 0) || (device >= g_devices.size())) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  amd::Device* dev = g_devices[device]->devices()[0];
+
+  // Set the allocation attributes in AMD HMM
+  if (!dev->SetSvmAttributes(dev_ptr, count, static_cast<amd::MemoryAdvice>(advice))) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  HIP_RETURN(hipSuccess);
+}
+
+// ================================================================================================
+hipError_t hipMemRangeGetAttribute(void* data, size_t data_size, hipMemRangeAttribute attribute,
+                                   const void* dev_ptr, size_t count) {
+  HIP_INIT_API(hipMemRangeGetAttribute, data, data_size, attribute, dev_ptr, count);
+
+  if ((data == nullptr) || (data_size == 0) || (dev_ptr == nullptr) || (count == 0)) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  // Shouldn't matter for which device the interface is called
+  amd::Device* dev = g_devices[0]->devices()[0];
+
+  // Get the allocation attribute from AMD HMM
+  if (!dev->GetSvmAttributes(&data, &data_size, reinterpret_cast<int*>(&attribute), 1,
+                             dev_ptr, count)) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  HIP_RETURN(hipErrorInvalidValue);
+}
+
+// ================================================================================================
+hipError_t hipMemRangeGetAttributes(void** data, size_t* data_sizes,
+                                    hipMemRangeAttribute* attributes, size_t num_attributes,
+                                    const void* dev_ptr, size_t count) {
+  HIP_INIT_API(hipMemRangeGetAttributes, data, data_sizes,
+               attributes, num_attributes, dev_ptr, count);
+
+  if ((data == nullptr) || (data_sizes == nullptr) || (attributes == nullptr) ||
+      (num_attributes == 0) || (dev_ptr == nullptr) || (count == 0)) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  // Shouldn't matter for which device the interface is called
+  amd::Device* dev = g_devices[0]->devices()[0];
+  // Get the allocation attributes from AMD HMM
+  if (!dev->GetSvmAttributes(data, data_sizes, reinterpret_cast<int*>(attributes),
+      num_attributes, dev_ptr, count)) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  HIP_RETURN(hipErrorInvalidValue);
+}
+
+// ================================================================================================
+hipError_t hipStreamAttachMemAsync(hipStream_t stream, hipDeviceptr_t* dev_ptr,
+                                   size_t length, unsigned int flags) {
+  HIP_INIT_API(hipStreamAttachMemAsync, stream, dev_ptr, length, flags);
+
+  if ((stream == nullptr) || (dev_ptr == nullptr) || (length == 0)) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  // Unclear what should be done for this interface in AMD HMM, since it's generic SVM alloc
+  HIP_RETURN(hipErrorInvalidValue);
+}
+
+// ================================================================================================
+static hipError_t ihipMallocManaged(void** ptr, size_t size) {
+  if (size == 0) {
+    *ptr = nullptr;
+    return hipSuccess;
+  } else if (ptr == nullptr) {
+    return hipErrorInvalidValue;
+  }
+
+  assert((hip::host_device->asContext()!= nullptr) && "Current host context must be valid");
+  amd::Context& ctx = *hip::host_device->asContext();
+
+  const amd::Device& dev = *ctx.devices()[0];
+  // For now limit to the max allocation size on the device.
+  // The apps should be able to go over theCould you  limit allocation in the future
+  if (dev.info().maxMemAllocSize_ < size) {
+    return hipErrorMemoryAllocation;
+  }
+
+  // Allocate SVM fine grain buffer with the forced host pointer, avoiding explicit memory
+  // allocation in the device driver
+  *ptr = amd::SvmBuffer::malloc(ctx, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_ALLOC_HOST_PTR,
+                                size, dev.info().memBaseAddrAlign_);
+  if (*ptr == nullptr) {
+    return hipErrorMemoryAllocation;
+  }
+
+  ClPrint(amd::LOG_INFO, amd::LOG_API, "%-5d: [%zx] ihipMallocManaged ptr=0x%zx",  getpid(),
+    std::this_thread::get_id(), *ptr);
+  return hipSuccess;
+}

projects/clr/hipamd/rocclr/hip_memory.cpp

@@ -105,6 +105,7 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
   return hipSuccess;
 }
 
+// ================================================================================================
 hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind,
                       amd::HostQueue& queue, bool isAsync = false) {
   if (sizeBytes == 0) {
@@ -154,36 +155,26 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
               *srcMemory->asBuffer(), sOffset, sizeBytes, dst);
     isAsync = false;
   } else if ((srcMemory != nullptr) && (dstMemory != nullptr)) {
-    if (queueDevice != srcMemory->getContext().devices()[0]) {
-      amd::Coord3D srcOffset(sOffset, 0, 0);
-      amd::Coord3D dstOffset(dOffset, 0, 0);
-      amd::Coord3D copySize(sizeBytes, 1, 1);
+    if ((kind == hipMemcpyDeviceToDevice) &&
+        // Check if the queue device doesn't match the device on any memory object. Hence
+        // it's a P2P transfer, because the app has requested access to another GPU
+        (srcMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0])) {
       command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
-                *srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
-      command->enqueue();
-      if (!isAsync) {
-        command->awaitCompletion();
+          *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
+      if (command == nullptr) {
+        return hipErrorOutOfMemory;
       }
-      command->release();
-      return hipSuccess;
-    }
-    if (queueDevice != dstMemory->getContext().devices()[0]) {
-      amd::Coord3D srcOffset(sOffset, 0, 0);
-      amd::Coord3D dstOffset(dOffset, 0, 0);
-      amd::Coord3D copySize(sizeBytes, 1, 1);
-      command = new amd::CopyMemoryP2PCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
-                *srcMemory->asBuffer(),*dstMemory->asBuffer(), srcOffset, dstOffset, copySize);
-      command->enqueue();
-      if (!isAsync) {
-        command->awaitCompletion();
+      // Make sure runtime has valid memory for the command execution. P2P access
+      // requires page table mapping on the current device to another GPU memory
+      if (!static_cast<amd::CopyMemoryP2PCommand*>(command)->validateMemory()) {
+        delete command;
+        return hipErrorInvalidValue;
       }
-      command->release();
-      return hipSuccess;
+    } else {
+      command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
+          *srcMemory->asBuffer(), *dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
     }
-    command = new amd::CopyMemoryCommand(queue, CL_COMMAND_COPY_BUFFER, waitList,
-              *srcMemory->asBuffer(),*dstMemory->asBuffer(), sOffset, dOffset, sizeBytes);
   }
-
   if (command == nullptr) {
     return hipErrorOutOfMemory;
   }
@@ -201,6 +192,7 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
   return hipSuccess;
 }
 
+// ================================================================================================
 hipError_t hipExtMallocWithFlags(void** ptr, size_t sizeBytes, unsigned int flags) {
   HIP_INIT_API(hipExtMallocWithFlags, ptr, sizeBytes, flags);
 
@@ -245,17 +237,6 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
   HIP_RETURN(ihipMalloc(ptr, sizeBytes, ihipFlags), *ptr);
 }
 
-hipError_t hipMallocManaged(void** devPtr, size_t size,
-                            unsigned int flags) {
-  HIP_INIT_API(hipMallocManaged, devPtr, size, flags);
-
-  if (flags != hipMemAttachGlobal) {
-    HIP_RETURN(hipErrorInvalidValue);
-  }
-
-  HIP_RETURN(ihipMalloc(devPtr, size, CL_MEM_SVM_FINE_GRAIN_BUFFER), *devPtr);
-}
-
 hipError_t hipFree(void* ptr) {
   HIP_INIT_API(hipFree, ptr);