P2P checkpoint.

Maintain enabled peer tables for each device.

include/hcc_detail/hip_hcc.h

@@ -519,9 +519,11 @@ public:
     // "Allocate" a stream ID:
     ihipStream_t::SeqNum_t  incStreamId() { return _stream_id++; };
 
-    void recomputePeerAgents();
-    void addPeer(ihipDevice_t *peer);
-    void removePeer(ihipDevice_t *peer);
+    bool addPeer(ihipDevice_t *peer);
+    bool removePeer(ihipDevice_t *peer);
+
+    uint32_t peerCnt() const { return _peerCnt; };
+    uint32_t peerAgents() const { return _peerAgents; };
 
 
 private:
@@ -532,6 +534,8 @@ private:
     std::list<ihipDevice_t*>  _peers;     // list of enabled peer devices.
     uint32_t                  _peerCnt;     // number of enabled peers
     hsa_agent_t              *_peerAgents;  // efficient packed array of enabled agents (to use for allocations.)
+private:
+    void recomputePeerAgents();
 };
 
 // Note Mutex selected based on DeviceMutex
@@ -560,6 +564,8 @@ public: // Functions:
     void locked_waitAllStreams();
     void locked_syncDefaultStream(bool waitOnSelf);
 
+    ihipDeviceCritical_t  &criticalData() { return _criticalData; }; // TODO, move private.  Fix P2P.
+
 public: // Data, set at initialization:
     unsigned                _device_index; // index into g_devices.
 

include/hcc_detail/hip_runtime_api.h

@@ -913,21 +913,21 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDev
 
 
 /**
- * @brief Disable registering memory on peerDevice for direct access from the current device.
+ * @brief Disable direct access from current device's virtual address space to memory allocations physically located on a peer device.  
  *
- * If there are any allocations on peerDevice which were registered in the current device using hipPeerRegister() then these allocations will be automatically unregistered.
  * Returns hipErrorPeerAccessNotEnabled if direct access to memory on peerDevice has not yet been enabled from the current device.
  *
- * @param [in] peerDevice
+ * @param [in] peerDeviceId
+ *
  * TODO:cudaErrorPeerAccessNotEnabled and cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown
  * Returns #hipSuccess, #hipErrorUnknown
  */
 hipError_t  hipDeviceDisablePeerAccess (int peerDeviceId);
 
 /**
- * @brief Enable registering memory on peerDevice for direct access from the current device.
+ * @brief Enable direct access from current device's virtual address space to memory allocations physically located on a peer device.  
  *
- * @param [in] peerDevice
+ * @param [in] peerDeviceId
  * @param [in] flags
  *
  * TODO:cudaErrorInvalidDevice error not supported in HIP, return hipErrorUnknown
@@ -939,14 +939,14 @@ hipError_t  hipDeviceEnablePeerAccess (int  peerDeviceId, unsigned int flags);
  * @brief Copies memory from one device to memory on another device.
  *
  * @param [out] dst - Destination device pointer.
- * @param [in] dstDevice - Destination device
+ * @param [in] dstDeviceId - Destination device
  * @param [in] src - Source device pointer
- * @param [in] srcDevice - Source device
+ * @param [in] srcDeviceId - Source device
  * @param [in] sizeBytes - Size of memory copy in bytes
  *
  * Returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice
  */
-hipError_t hipMemcpyPeer ( void* dst, int  dstDevice, const void* src, int  srcDevice, size_t sizeBytes );
+hipError_t hipMemcpyPeer (void* dst, int dstDeviceId, const void* src, int srcDeviceId, size_t sizeBytes);
 
 /**
  * @brief Copies memory from one device to memory on another device.
@@ -961,7 +961,7 @@ hipError_t hipMemcpyPeer ( void* dst, int  dstDevice, const void* src, int  srcD
  * Returns #hipSuccess, #hipErrorInvalidValue, #hipErrorInvalidDevice
  */
 #if __cplusplus
-hipError_t hipMemcpyPeerAsync ( void* dst, int  dstDevice, const void* src, int  srcDevice, size_t sizeBytes, hipStream_t stream=0 );
+hipError_t hipMemcpyPeerAsync ( void* dst, int  dstDeviceId, const void* src, int  srcDevice, size_t sizeBytes, hipStream_t stream=0 );
 #else
 hipError_t hipMemcpyPeerAsync(void* dst, int dstDevice, const void* src, int srcDevice, size_t sizeBytes, hipStream_t stream);
 #endif

include/hip_runtime_api.h

@@ -138,7 +138,7 @@ typedef struct hipPointerAttribute_t {
  * @enum
  * @ingroup Enumerations
  */
-// Developer note - when updating these, update the hipErrorName and hipErrorString functions
+// Developer note - when updating these, update the hipErrorName and hipErrorString functions in NVCC and HCC paths
 typedef enum hipError_t {
      hipSuccess = 0                   ///< Successful completion.
     ,hipErrorMemoryAllocation         ///< Memory allocation error.
@@ -155,6 +155,8 @@ typedef enum hipError_t {
     ,hipErrorNoDevice                 ///< Call to hipGetDeviceCount returned 0 devices
     ,hipErrorNotReady                 ///< Indicates that asynchronous operations enqueued earlier are not ready.  This is not actually an error, but is used to distinguish from hipSuccess (which indicates completion).  APIs that return this error include hipEventQuery and hipStreamQuery.
     ,hipErrorUnknown                  ///< Unknown error.
+    ,hipErrorPeerAccessNotEnabled     ///< Peer access was never enabled from the current device.
+    ,hipErrorPeerAccessAlreadyEnabled ///< Peer access was already enabled from the current device.
     ,hipErrorRuntimeMemory            ///< HSA runtime memory call returned error.  Typically not seen in production systems.
     ,hipErrorRuntimeOther             ///< HSA runtime call other than memory returned error.  Typically not seen in production systems.
     ,hipErrorTbd                      ///< Marker that more error codes are needed.