Addition of hipCtxEnablePeerAccess and hipCtxDisablePeerAccess functions

Change-Id: I381c8cbbde17eae7d9bb5d4cb1596cebf4bda039


[ROCm/clr commit: a1d401336b]

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

@@ -1084,6 +1084,35 @@ hipError_t hipCtxSynchronize ( void );
 
 hipError_t hipCtxGetFlags ( unsigned int* flags );
 
+/**
+ * @brief Enables direct access to memory allocations in a peer context.
+ *
+ * Memory which already allocated on peer device will be mapped into the address space of the current device.  In addition, all
+ * future memory allocations on peerDeviceId will be mapped into the address space of the current device when the memory is allocated.
+ * The peer memory remains accessible from the current device until a call to hipDeviceDisablePeerAccess or hipDeviceReset.
+ *
+ *
+ * @param [in] peerCtx
+ * @param [in] flags
+ *
+ * Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue,
+ * @returns #hipErrorPeerAccessAlreadyEnabled if peer access is already enabled for this device.
+ * @warning PeerToPeer support is experimental.
+ */
+hipError_t  hipCtxEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags);
+
+/**
+ * @brief Disable direct access from current device's virtual address space to memory allocations physically located on a peer device through contex.Disables direct access to memory allocations in a peer context and unregisters any registered allocations.
+ *
+ * Returns hipErrorPeerAccessNotEnabled if direct access to memory on peerDevice has not yet been enabled from the current device.
+ *
+ * @param [in] peerCtx
+ *
+ * @returns #hipSuccess, #hipErrorPeerAccessNotEnabled
+ * @warning PeerToPeer support is experimental.
+ */
+hipError_t  hipCtxDisablePeerAccess (hipCtx_t peerCtx);
+
 
 // TODO-ctx
 /**

projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h

@@ -474,6 +474,16 @@ inline static hipError_t  hipDeviceEnablePeerAccess ( int  peerDevice, unsigned
     return hipCUDAErrorTohipError(cudaDeviceEnablePeerAccess ( peerDevice, flags ));
 }
 
+inline static hipError_t  hipCtxDisablePeerAccess ( hipCtx_t peerCtx )
+{
+    return hipCUDAErrorTohipError(cudaCtxDisablePeerAccess ( peerCtx ));
+}
+
+inline static hipError_t  hipCtxEnablePeerAccess ( hipCtx_t peerCtx, unsigned int  flags )
+{
+    return hipCUDAErrorTohipError(cudaCtxEnablePeerAccess ( peerCtx, flags ));
+}
+
 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 ));

projects/clr/hipamd/src/hip_peer.cpp

@@ -67,7 +67,7 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, hipCtx_t thisCtx, hipCtx_
 //---
 // Disable visibility of this device into memory allocated on peer device.
 // Remove this device from peer device peerlist.
-hipError_t hipDeviceDisablePeerAccess (hipCtx_t peerCtx)
+hipError_t ihipDisablePeerAccess (hipCtx_t peerCtx)
 {
     HIP_INIT_API(peerCtx);
 
@@ -109,7 +109,7 @@ hipError_t hipDeviceDisablePeerAccess (hipCtx_t peerCtx)
 //---
 // Allow the current device to see all memory allocated on peerDevice.
 // This should add this device to the peer-device peer list.
-hipError_t hipDeviceEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags)
+hipError_t ihipEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags)
 {
     HIP_INIT_API(peerCtx, flags);
 
@@ -175,7 +175,7 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId)
 {
     HIP_INIT_API(peerDeviceId);
 
-    return hipDeviceDisablePeerAccess(ihipGetPrimaryCtx(peerDeviceId));
+    return ihipDisablePeerAccess(ihipGetPrimaryCtx(peerDeviceId));
 }
 
 
@@ -183,7 +183,7 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
 {
     HIP_INIT_API(peerDeviceId, flags);
 
-    return hipDeviceEnablePeerAccess(ihipGetPrimaryCtx(peerDeviceId), flags);
+    return ihipEnablePeerAccess(ihipGetPrimaryCtx(peerDeviceId), flags);
 }
 
 
@@ -200,6 +200,16 @@ hipError_t hipMemcpyPeerAsync (void* dst, int  dstDevice, const void* src, int
     return hipMemcpyPeerAsync(dst, ihipGetPrimaryCtx(dstDevice), src, ihipGetPrimaryCtx(srcDevice), sizeBytes, stream);
 }
 
+hipError_t hipCtxEnablePeerAccess (hipCtx_t peerCtx, unsigned int flags)
+{
+    HIP_INIT_API(peerCtx, flags);
 
+    return ihipEnablePeerAccess(peerCtx, flags);
+}
 
+hipError_t hipCtxDisablePeerAccess (hipCtx_t peerCtx)
+{
+    HIP_INIT_API(peerCtx);
 
+    return ihipDisablePeerAccess(peerCtx);
+}