Added support for hipMemGetAddressRange

Change-Id: I99a796a4eb765152cf15a12d6a86b58684d34f50

hipamd/include/hip/hcc_detail/hip_runtime_api.h

@@ -1281,6 +1281,18 @@ hipError_t  hipDeviceEnablePeerAccess (int  peerDeviceId, unsigned int flags);
  */
 hipError_t  hipDeviceDisablePeerAccess (int peerDeviceId);
 
+/**
+ * @brief Get information on memory allocations.
+ *
+ * @param [out] pbase - BAse pointer address
+ * @param [out] psize - Size of allocation
+ * @param [in]  dptr- Device Pointer
+ *
+ * @returns #hipSuccess, #hipErrorInvalidDevicePointer
+ *
+ * @see hipCtxCreate, hipCtxDestroy, hipCtxGetFlags, hipCtxPopCurrent, hipCtxGetCurrent, hipCtxSetCurrent, hipCtxPushCurrent, hipCtxSetCacheConfig, hipCtxSynchronize, hipCtxGetDevice
+ */
+hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr );
 
 #ifndef USE_PEER_NON_UNIFIED
 #define USE_PEER_NON_UNIFIED 1

hipamd/include/hip/nvcc_detail/hip_runtime_api.h

@@ -639,6 +639,11 @@ inline static hipError_t  hipCtxEnablePeerAccess ( hipCtx_t peerCtx, unsigned in
     return hipCUResultTohipError(cuCtxEnablePeerAccess(peerCtx, flags));
 }
 
+inline static hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr )
+{
+    return hipCUResultTohipError(cuMemGetAddressRange( pbase , psize , dptr));
+}
+
 inline static hipError_t hipMemcpyPeer ( void* dst, int  dstDevice, const void* src, int  srcDevice, size_t count )
 {
     return hipCUDAErrorTohipError(cudaMemcpyPeer(dst, dstDevice, src, srcDevice, count));

hipamd/src/hip_memory.cpp

@@ -28,6 +28,7 @@ THE SOFTWARE.
 #include "hip_hcc.h"
 #include "trace_helper.h"
 #include "hip/hcc_detail/hip_texture.h"
+#include <hc_am.hpp>
 
 //-------------------------------------------------------------------------------------------------
 //-------------------------------------------------------------------------------------------------
@@ -1037,3 +1038,17 @@ hipError_t hipFreeArray(hipArray* array)
     return ihipLogStatus(hipStatus);
 }
 
+hipError_t hipMemGetAddressRange ( hipDeviceptr_t* pbase, size_t* psize, hipDeviceptr_t dptr )
+{
+    HIP_INIT_API ( pbase , psize , dptr );
+    hipError_t hipStatus = hipSuccess;
+    hc::accelerator acc;
+    hc::AmPointerInfo amPointerInfo( NULL , NULL , 0 , acc , 0 , 0 );
+    am_status_t status = hc::am_memtracker_getinfo( &amPointerInfo , dptr );
+    if (status == AM_SUCCESS) {
+        *pbase = amPointerInfo._devicePointer;
+        *psize = amPointerInfo._sizeBytes;
+    }
+    hipStatus = hipErrorInvalidDevicePointer;
+    return hipStatus;
+}