Correct P2P memory access section

Signed-off-by: Jan Stephan <[email protected]>

Author: j-stephan

docs/how-to/hip_runtime_api/multi_device.rst

@@ -73,7 +73,10 @@ applications that require frequent data exchange between GPUs, as it eliminates
 the need to transfer data through the host memory.
 
 By adding peer-to-peer access to the example referenced in
-:ref:`multi_device_selection`, data can be copied between devices:
+:ref:`multi_device_selection`, data can be efficiently copied between devices.
+If peer-to-peer access is not activated, the call to :cpp:func:`hipMemcpy`
+still works but internally uses a staging buffer in host memory, which incurs a
+performance penalty.
 
 .. tab-set::
 
@@ -82,13 +85,13 @@ By adding peer-to-peer access to the example referenced in
         .. literalinclude:: ../../tools/example_codes/p2p_memory_access.hip
             :start-after: // [sphinx-start]
             :end-before: // [sphinx-end]
-            :emphasize-lines: 31-37, 51-55
+            :emphasize-lines: 43-49, 63-67
             :language: cpp
 
     .. tab-item:: without peer-to-peer
 
-        .. literalinclude:: ../../tools/example_codes/p2p_memory_access.hip
+        .. literalinclude:: ../../tools/example_codes/p2p_memory_access_host_staging.hip
             :start-after: // [sphinx-start]
             :end-before: // [sphinx-end]
-            :emphasize-lines: 43-49, 53, 58
+            :emphasize-lines: 55-57
             :language: cpp

docs/tools/example_codes/p2p_memory_access_failed.hip renamed to docs/tools/example_codes/p2p_memory_access_host_staging.hip

@@ -53,7 +53,7 @@ int main()
     if(deviceCount < 2)
     {
         std::cout << "This example requires at least two HIP devices." << std::endl;
-        return EXIT_FAILURE;
+        return EXIT_SUCCESS;
     }
 
     double* deviceData0;
@@ -75,13 +75,9 @@ int main()
     simpleKernel<<<1000, 128>>>(deviceData1); // Launch kernel on device 1
     HIP_CHECK(hipDeviceSynchronize());
 
-    // Attempt to use deviceData0 on device 1 (This will not work as deviceData0 is allocated on device 0)
+    // Use deviceData0 on device 1. This works but incurs a performance penalty.
     HIP_CHECK(hipSetDevice(deviceId1));
-    hipError_t err = hipMemcpy(deviceData1, deviceData0, size, hipMemcpyDeviceToDevice); // This should fail
-    if (err != hipSuccess)
-    {
-        std::cout << "Error: Cannot access deviceData0 from device 1, deviceData0 is on device 0" << std::endl;
-    }
+    HIP_CHECK(hipMemcpy(deviceData1, deviceData0, size, hipMemcpyDeviceToDevice));
 
     // Copy result from device 0
     double hostData0[1024];

docs/tools/update_example_codes.py

@@ -264,8 +264,8 @@
     "docs/tools/example_codes/p2p_memory_access.hip"
 )
 urllib.request.urlretrieve(
-    "https://raw.githubusercontent.com/ROCm/rocm-examples/amd-staging/HIP-Doc/Programming-Guide/Using-HIP-Runtime-API/Multi-Device-Management/p2p_memory_access_failed/main.hip",
-    "docs/tools/example_codes/p2p_memory_access_failed.hip"
+    "https://raw.githubusercontent.com/ROCm/rocm-examples/amd-staging/HIP-Doc/Programming-Guide/Using-HIP-Runtime-API/Multi-Device-Management/p2p_memory_access_host_staging/main.hip",
+    "docs/tools/example_codes/p2p_memory_access_host_staging.hip"
 )
 
 # Reference examples from HIP-Doc / Reference