diff --git a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h index 0d08563fde..1c50c33359 100644 --- a/projects/clr/hipamd/include/hcc_detail/hip_runtime_api.h +++ b/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 /** diff --git a/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h index a51bca2d02..67091b8bfd 100644 --- a/projects/clr/hipamd/include/nvcc_detail/hip_runtime_api.h +++ b/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 )); diff --git a/projects/clr/hipamd/src/hip_peer.cpp b/projects/clr/hipamd/src/hip_peer.cpp index fd51599815..8f75d99918 100644 --- a/projects/clr/hipamd/src/hip_peer.cpp +++ b/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); +}