P2p checkpoint.

- set USE_PEER_TO_PEER=3 (requires HCC "am_memtracker_update_peers") - when enabling peer, turn it on for previously allocated memory. - hipDeviceCanAccessPeer is no longer self-ware (self does not qualify as a peer) - device peerlist always includes self, so when we call allow_access we never remove self access. - hipDeviceReset() removes old peer mappings. [ROCm/hip commit: 1f53c55d3e]
2016-04-11 12:52:18 -05:00
@@ -38,7 +38,7 @@ THE SOFTWARE.
 // Compile peer-to-peer support.
 // >= 2 : use HCC hc:accelerator::get_is_peer
 // >= 3 : use hc::am_memtracker_update_peers(...)
-#define USE_PEER_TO_PEER 2
+#define USE_PEER_TO_PEER 3

 //#define INLINE static inline

@@ -526,6 +526,7 @@ public:

    bool addPeer(ihipDevice_t *peer);
    bool removePeer(ihipDevice_t *peer);
+    void resetPeers(ihipDevice_t *thisDevice);

    uint32_t peerCnt() const { return _peerCnt; };
    hsa_agent_t *peerAgents() const { return _peerAgents; };
@@ -907,10 +907,33 @@ hipError_t hipMemGetInfo  (size_t * free, size_t * total)   ;
 *
 * Returns "1" in @p canAccessPeer if the specified @p device is capable
 * of directly accessing memory physically located on peerDevice , or "0" if not.
+ *
+ * Returns "0" in @p canAccessPeer if deviceId == peerDeviceId, and both are valid devices : a device is not a peer of itself.
+ *
+ *
+ *
+ * @returns #hipSuccess, 
+ * @returns #hipErrorInvalidDevice if deviceId or peerDeviceId are not valid devices
 */
 hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDeviceId);


+/**
+ * @brief Enable direct access from current device's virtual address space to memory allocations physically located on a peer device.  
+ *
+ * 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] peerDeviceId
+ * @param [in] flags
+ *
+ * Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, 
+ * @returns #hipErrorPeerAccessAlreadyEnabled if peer access is already enabled for this device.
+ */
+hipError_t  hipDeviceEnablePeerAccess (int  peerDeviceId, unsigned int flags);
+

 /**
 * @brief Disable direct access from current device's virtual address space to memory allocations physically located on a peer device.  
@@ -923,20 +946,6 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int deviceId, int peerDev
 */
 hipError_t  hipDeviceDisablePeerAccess (int peerDeviceId);

-/**
- * @brief Enable direct access from current device's virtual address space to memory allocations physically located on a peer device.  
- *
- * 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] peerDeviceId
- * @param [in] flags
- *
- * Returns #hipSuccess, #hipErrorInvalidDevice, #hipErrorInvalidValue, #hipErrorPeerAccessAlreadyEnabled
- */
-hipError_t  hipDeviceEnablePeerAccess (int  peerDeviceId, unsigned int flags);

 /**
 * @brief Copies memory from one device to memory on another device.
@@ -127,7 +127,7 @@ void printDeviceProp (int deviceId)
    for (int i=0; i<deviceCnt; i++) {
        int isPeer;
        hipDeviceCanAccessPeer(&isPeer, deviceId, i);
-        if ((i != deviceId) && isPeer) {
+        if (isPeer) {
            cout << "device#" << i << " ";
        }
    }
@@ -174,6 +174,7 @@ hipError_t hipDeviceReset(void)
    if (device) {
        //---
        //Wait for pending activity to complete?  TODO - check if this is required behavior:
+        //TODO, also we have small window between wait and reset.
        
        device->locked_waitAllStreams();

@@ -240,6 +240,15 @@ bool ihipDeviceCriticalBase_t<DeviceMutex>::removePeer(ihipDevice_t *peer)
    }
 }

+
+template<>
+void ihipDeviceCriticalBase_t<DeviceMutex>::resetPeers(ihipDevice_t *thisDevice)
+{
+    _peers.clear();
+    _peerCnt = 0;
+    addPeer(thisDevice); // peer-list always contains self agent.
+}
+
 //-------------------------------------------------------------------------------------------------

 //---
@@ -444,12 +453,18 @@ void ihipDevice_t::locked_reset()
    // Reset and remove streams:
    crit->streams().clear();

-#if USE_PEER_TO_PEER==2
-    // remove peer mappings to this device?  Call removePeer on all other devices?
+
+
+#if USE_PEER_TO_PEER>=2
+    // This resest peer list to just me:
+    crit->resetPeers(this);
+
 #endif

    // Reset and release all memory stored in the tracker:
+    // Reset will remove peer mapping so don't need to do this explicitly.
    am_memtracker_reset(_acc);
+
 };


@@ -474,10 +489,13 @@ void ihipDevice_t::init(unsigned device_index, unsigned deviceCnt, hc::accelerat

    getProperties(&_props);

+    _criticalData.init(deviceCnt);
+
+    locked_reset();
+
    _default_stream = new ihipStream_t(device_index, acc.get_default_view(), hipStreamDefault);
    locked_addStream(_default_stream);
-
-    _criticalData.init(deviceCnt);
+    

    tprintf(DB_SYNC, "created device with default_stream=%p\n", _default_stream);

@@ -132,7 +132,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
            hc::am_memtracker_update(*ptr, device->_device_index, 0);
            {
                LockedAccessor_DeviceCrit_t crit(device->criticalData());
-                if (crit->peerCnt()) {
+                if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved:
                    hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
                    if (hsa_status != HSA_STATUS_SUCCESS) {
                        hip_status = hipErrorMemoryAllocation; 
@@ -173,8 +173,9 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
            }else{
                hc::am_memtracker_update(*ptr, device->_device_index, flags);
                {
+                    // TODO - allow_access only works for device memory, need to change am_alloc to allocate host directly.
                    LockedAccessor_DeviceCrit_t crit(device->criticalData());
-                    if (crit->peerCnt()) {
+                    if (crit->peerCnt() > 1) { // peerCnt includes self so only call allow_access if other peers involved:
                        hsa_status_t hsa_status = hsa_amd_agents_allow_access(crit->peerCnt(), crit->peerAgents(), NULL, *ptr);
                        if (hsa_status != HSA_STATUS_SUCCESS) {
                            hip_status = hipErrorMemoryAllocation; 
@@ -17,6 +17,8 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 THE SOFTWARE.
 */

+#include <hc_am.hpp>
+
 #include "hip_runtime.h"
 #include "hcc_detail/hip_hcc.h"
 #include "hcc_detail/trace_helper.h"
@@ -35,11 +37,15 @@ hipError_t hipDeviceCanAccessPeer (int* canAccessPeer, int  deviceId, int peerDe
    auto peerDevice = ihipGetDevice(peerDeviceId);

    if ((thisDevice != NULL) && (peerDevice != NULL)) {
+        if (deviceId == peerDeviceId) {
+            *canAccessPeer = 0;
+        } else {
 #if USE_PEER_TO_PEER>=2
-        *canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc);
+            *canAccessPeer = peerDevice->_acc.get_is_peer(thisDevice->_acc);
 #else
-        *canAccessPeer = 0;
+            *canAccessPeer = 0;
 #endif
+        }

    } else {
        *canAccessPeer = 0;
@@ -69,15 +75,15 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId)
 #endif
        if (! canAccessPeer) {
            err = hipErrorInvalidDevice;  // P2P not allowed between these devices.
+        } else if (thisDevice == peerDevice)  {
+            err = hipErrorInvalidDevice;  // Can't disable peer access to self.
        } else {
-
-
            LockedAccessor_DeviceCrit_t crit(thisDevice->criticalData());
            bool changed = crit->removePeer(peerDevice);
            if (changed) {
 #if USE_PEER_TO_PEER>=3
                // Update the peers for all memory already saved in the tracker:
-                am_memtracker_update_peers(device->_acc, crit->peerCnt(), crit->peerAgents());
+                am_memtracker_update_peers(thisDevice->_acc, crit->peerCnt(), crit->peerAgents());
 #endif
            } else {
                err = hipErrorPeerAccessNotEnabled; // never enabled P2P access.
@@ -92,9 +98,6 @@ hipError_t hipDeviceDisablePeerAccess (int peerDeviceId)
 };


-/**
- * @warning Need to update this function when RT supports P2P
- */
 //---
 // Enable registering memory on peerDevice for direct access from the current device.
 hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
@@ -113,7 +116,7 @@ hipError_t hipDeviceEnablePeerAccess (int peerDeviceId, unsigned int flags)
            bool isNewPeer = crit->addPeer(peerDevice);
            if (isNewPeer) {
 #if USE_PEER_TO_PEER>=3
-                am_memtracker_update_peers(device->_acc, crit->peerCnt(), crit->peerAgents());
+                am_memtracker_update_peers(thisDevice->_acc, crit->peerCnt(), crit->peerAgents());
 #endif
            } else {
                err = hipErrorPeerAccessAlreadyEnabled;
@@ -8,6 +8,10 @@ include_directories( ${PROJECT_SOURCE_DIR}/include )
 set (HIP_Unit_Test_VERSION_MAJOR 1)
 set (HIP_Unit_Test_VERSION_MINOR 0)

+if(NOT DEFINED HIP_MULTI_GPU)
+    set(HIP_MULTI_GPU 0 CACHE BOOL "Run tests requiring more than one GPU")
+endif()
+
 if(NOT DEFINED HIP_BUILD_LOCAL)
    if(NOT DEFINED ENV{HIP_BUILD_LOCAL})
        set(HIP_BUILD_LOCAL 0 CACHE BOOL "Build HIP in local folder")
@@ -218,6 +222,12 @@ make_test(hipFuncDeviceSynchronize " ")
 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-serial" --tests 0x1) 
 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-pyramid" --tests 0x4) 
 make_named_test (hipMultiThreadDevice "hipMultiThreadDevice-nearzero" --tests 0x10) 
-make_test(hipPeerToPeer_simple " " )
+
+if (${HIP_MULTI_GPU})
+    make_test(hipPeerToPeer_simple )                  # use current device for copy, this fails. 
+    make_test(hipPeerToPeer_simple --memcpyWithPeer)  
+    make_test(hipPeerToPeer_simple --mirrorPeers)    # mirror mapping: test to ensure mirror doesn't destroy orig mapping.
+
+endif()

 make_hipify_test(specialFunc.cu )
@@ -1,4 +1,3 @@
-
 /*
 Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.

@@ -52,9 +51,11 @@ void parseMyArguments(int argc, char *argv[])
 };


-int main(int argc, char *argv[])
+//---
+// Test which enables peer2peer first, then allocates the memory.
+void enablePeerFirst()
 {
-    parseMyArguments(argc, argv);
+    printf ("\n==testing: %s\n", __func__);

    int deviceCnt;

@@ -74,6 +75,110 @@ int main(int argc, char *argv[])

    assert(canAccessPeer);

+    HIPCHECK (hipSetDevice(currentDevice));
+    HIPCHECK(hipDeviceReset());
+    HIPCHECK (hipSetDevice(peerDevice));
+    HIPCHECK(hipDeviceReset());
+
+    HIPCHECK(hipSetDevice(currentDevice));
+    HIPCHECK(hipDeviceEnablePeerAccess(peerDevice, 0));
+
+    if (p_mirrorPeers) {
+        int canAccessPeer;
+        HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, peerDevice, currentDevice));
+        assert(canAccessPeer);
+
+        HIPCHECK(hipSetDevice(peerDevice));
+            HIPCHECK(hipDeviceEnablePeerAccess(currentDevice, 0));
+        }
+
+        size_t Nbytes = N*sizeof(char);
+
+        char *A_d0, *A_d1;
+        char *A_h;
+
+        A_h = (char*)malloc(Nbytes);
+
+        // allocate and initialize memory on device0
+        HIPCHECK (hipSetDevice(currentDevice));
+        HIPCHECK (hipMalloc(&A_d0, Nbytes) );
+        HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); 
+
+        // allocate and initialize memory on peer device
+        HIPCHECK (hipSetDevice(peerDevice));
+        HIPCHECK (hipMalloc(&A_d1, Nbytes) );
+        HIPCHECK ( hipMemset(A_d1, 0x13, Nbytes) ); 
+
+
+
+        // Device0 push to device1, using P2P:
+        HIPCHECK (hipSetDevice(p_memcpyWithPeer ? peerDevice : currentDevice));
+        HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault));
+
+        // Copy data back to host:
+        HIPCHECK (hipSetDevice(peerDevice));
+        HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost));
+
+        // Check host data:
+        for (int i=0; i<N; i++) {
+            if (A_h[i] != memsetval) {
+                failed("mismatch at index:%d computed:0x%02x, golden memsetval:0x%02x\n", i, (int)A_h[i], (int)memsetval);
+            }
+        }
+    }
+
+
+    //---
+    // Test which allocated memory first, then enables peer2peer.
+    // Enabling peer needs to scan all allocated memory and enable peer access.
+    void allocMemoryFirst()
+    {
+        printf ("\n==testing: %s\n", __func__);
+        int deviceCnt;
+
+        HIPCHECK(hipGetDeviceCount(&deviceCnt));
+
+        int currentDevice = p_gpuDevice;
+        int peerDevice    = (p_peerDevice == -1) ? ((currentDevice + 1) % deviceCnt) : p_peerDevice;
+
+        printf ("N=%zu  device=%d peerDevice=%d (%d devices total)\n", N, currentDevice, peerDevice, deviceCnt);
+
+        // Must be on a multi-gpu system:
+        assert (currentDevice != peerDevice);
+
+        int canAccessPeer;
+        HIPCHECK(hipDeviceCanAccessPeer(&canAccessPeer, currentDevice, peerDevice));
+        printf ("dev#%d canAccessPeer:#%d=%d\n", currentDevice, peerDevice, canAccessPeer);
+
+        assert(canAccessPeer);
+
+        HIPCHECK (hipSetDevice(currentDevice));
+    HIPCHECK(hipDeviceReset());
+    HIPCHECK (hipSetDevice(peerDevice));
+    HIPCHECK(hipDeviceReset());
+
+
+    size_t Nbytes = N*sizeof(char);
+
+    char *A_d0, *A_d1;
+    char *A_h;
+
+    A_h = (char*)malloc(Nbytes);
+
+    //---
+    // allocate and initialize memory on device0
+    HIPCHECK (hipSetDevice(currentDevice));
+    HIPCHECK (hipMalloc(&A_d0, Nbytes) );
+    HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); 
+
+    // allocate and initialize memory on peer device
+    HIPCHECK (hipSetDevice(peerDevice));
+    HIPCHECK (hipMalloc(&A_d1, Nbytes) );
+    HIPCHECK ( hipMemset(A_d1, 0x13, Nbytes) ); 
+
+
+    //---
+    //Enable peer access, for memory already allocated:
    HIPCHECK(hipSetDevice(currentDevice));
    HIPCHECK(hipDeviceEnablePeerAccess(peerDevice, 0));

@@ -86,25 +191,9 @@ int main(int argc, char *argv[])
        HIPCHECK(hipDeviceEnablePeerAccess(currentDevice, 0));
    }

-    size_t Nbytes = N*sizeof(char);
-
-    char *A_d0, *A_d1;
-    char *A_h;
-
-    A_h = (char*)malloc(Nbytes);
-
-    // allocate and initialize memory on device0
-    HIPCHECK (hipSetDevice(currentDevice));
-    HIPCHECK (hipMalloc(&A_d0, Nbytes) );
-    HIPCHECK ( hipMemset(A_d0, memsetval, Nbytes) ); 
-
-    // allocate and initialize memory on peer device
-    HIPCHECK (hipSetDevice(peerDevice));
-    HIPCHECK (hipMalloc(&A_d1, Nbytes) );
-    HIPCHECK ( hipMemset(A_d1, 0x13, Nbytes) ); 
-
-

+    //---
+    // Copies to test functionality:
    // Device0 push to device1, using P2P:
    HIPCHECK (hipSetDevice(p_memcpyWithPeer ? peerDevice : currentDevice));
    HIPCHECK (hipMemcpy(A_d1, A_d0, Nbytes, hipMemcpyDefault));
@@ -113,13 +202,29 @@ int main(int argc, char *argv[])
    HIPCHECK (hipSetDevice(peerDevice));
    HIPCHECK (hipMemcpy(A_h, A_d1, Nbytes, hipMemcpyDeviceToHost));

+
+    //---
    // Check host data:
    for (int i=0; i<N; i++) {
        if (A_h[i] != memsetval) {
            failed("mismatch at index:%d computed:0x%02x, golden memsetval:0x%02x\n", i, (int)A_h[i], (int)memsetval);
        }
    }
+}
+
+
+
+int main(int argc, char *argv[])
+{
+    parseMyArguments(argc, argv);
+
+    if (p_tests & 0x1) {
+        enablePeerFirst();
+    }
+
+    if (p_tests & 0x2) {
+        allocMemoryFirst();
+    }

    passed();
-
 }