SWDEV-333557 - add hipIpcMemCopyTest

catch/unit/device/CMakeLists.txt

@@ -28,6 +28,7 @@ set(TEST_SRC
     hipDeviceSetGetMemPool.cc
     hipInit.cc
     hipDriverGetVersion.cc
+    hipIpcMemCopy.cc
 )
 
 if(HIP_PLATFORM MATCHES "amd" AND BUILD_SHARED_LIBS)
@@ -56,10 +57,12 @@ set_property(GLOBAL APPEND PROPERTY G_INSTALL_EXE_TARGETS getDeviceCount)
 add_executable(chkUUIDFrmChildProc_Exe EXCLUDE_FROM_ALL chkUUIDFrmChildProc_Exe.cc)
 add_executable(chkUUIDInGrandChild_Exe EXCLUDE_FROM_ALL chkUUIDInGrandChild_Exe.cc)
 add_executable(setuuidGetDevCount EXCLUDE_FROM_ALL setuuidGetDevCount_Exe.cc)
+add_executable(hipIpcMemCopyChildProc_Exe EXCLUDE_FROM_ALL hipIpcMemCopyChildProc_Exe.cc)
 set_property(GLOBAL APPEND PROPERTY G_INSTALL_EXE_TARGETS
              chkUUIDFrmChildProc_Exe
              chkUUIDInGrandChild_Exe
-             setuuidGetDevCount)
+             setuuidGetDevCount
+             hipIpcMemCopyChildProc_Exe)
 
 if(UNIX)
 add_executable(getUUIDfrmRocinfo EXCLUDE_FROM_ALL getUUIDfrmRocinfo_Exe.cc)
@@ -85,7 +88,8 @@ hip_add_exe_to_target(NAME DeviceTest
                       TEST_SRC ${TEST_SRC}
                       TEST_TARGET_NAME build_tests
                       COMPILE_OPTIONS -std=c++17)
-add_dependencies(build_tests getDeviceCount chkUUIDFrmChildProc_Exe chkUUIDInGrandChild_Exe setuuidGetDevCount multipleUUID setEnvInChildProc uuidList)
+add_dependencies(build_tests getDeviceCount chkUUIDFrmChildProc_Exe chkUUIDInGrandChild_Exe 
+                 setuuidGetDevCount multipleUUID setEnvInChildProc uuidList hipIpcMemCopyChildProc_Exe)
 #Disabled below two executable due to the defect ticket SWDEV-467665
 if(0)
 add_dependencies(build_tests passUUIDToGrandChild_Exe ResetUUIDInChild_Exe)

catch/unit/device/hipIpcMemCopy.cc

@@ -0,0 +1,76 @@
+/*
+Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+#include <hip_test_common.hh>
+#include <hip_test_process.hh>
+
+#include <string>
+#include <cstdint>
+
+#include <sstream>
+#include <iomanip>
+#include <iostream>
+
+std::string hipHandleToHex(const hipIpcMemHandle_t &h) {
+    std::ostringstream oss;
+    auto bytes = reinterpret_cast<const unsigned char*>(&h);
+    for (size_t i = 0; i < sizeof(h); ++i) {
+        oss << std::hex << std::setw(2) << std::setfill('0') << int(bytes[i]);
+    }
+    return oss.str();
+}
+
+/**
+ * Test Description
+ * ------------------------
+ *  - Verifies IPC copy with hipIpcGetMemHandle() and hipIpcOpenMemHandle()
+ *    by copying data between two processes and verifying result.
+ *  - Spawns child process and waits for it to finish.
+ *  - Child process reads the data and check copy result.
+ * Test source
+ * ------------------------
+ *  - unit/multiproc/hipIpcMemCopyTest.cc
+ *  - unit/multiproc/hipIpcMemCopyTest_child.cc
+ * Test requirements
+ * ------------------------
+ *  - HIP_VERSION >= 6.4
+ */
+TEST_CASE("Unit_hipIpcMemCopyTest_validation") {
+  size_t N = 1024;
+  size_t Nbytes = N * sizeof(int);
+  int *A_h{nullptr};
+  HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&A_h), Nbytes, hipHostMallocDefault));
+  int *A_d{nullptr};
+  HIP_CHECK(hipMalloc(reinterpret_cast<void**>(&A_d), Nbytes));
+
+  for (int i = 0; i < N; i ++){
+    A_h[i] = 6;
+  }
+
+  hipIpcMemHandle_t memHandle;
+
+  HIP_CHECK(hipMalloc(&A_d, Nbytes));
+  HIP_CHECK(hipIpcGetMemHandle(&memHandle,
+                               A_d));
+  HIP_CHECK(hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice));
+
+  std::string hex = hipHandleToHex(memHandle);
+
+  hip::SpawnProc proc("hipIpcMemCopyChildProc_Exe", false);
+  REQUIRE(proc.run(hex) == 1);
+}

catch/unit/device/hipIpcMemCopyChildProc_Exe.cc

@@ -0,0 +1,98 @@
+/*
+Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include <hip/hip_runtime.h>
+#include <iostream>
+#include <cstring>
+#include <cstdint>
+#include <sstream>
+
+#define HIP_CHECK(error)\
+{\
+  hipError_t localError = error;\
+  if (localError != hipSuccess) {\
+    printf("error: '%s'(%d) from %s at %s:%d\n", \
+           hipGetErrorString(localError), \
+            localError, #error, __FUNCTION__, __LINE__);\
+    exit(0);\
+  }\
+}
+
+hipIpcMemHandle_t hexToHipHandle(const std::string &hex) {
+  if (hex.size() != sizeof(hipIpcMemHandle_t)*2) {
+    printf("Invalid hex string length\n");
+  }
+
+  hipIpcMemHandle_t h{};
+  auto bytes = reinterpret_cast<unsigned char*>(&h);
+  for (size_t i = 0; i < sizeof(h); i++) {
+      unsigned int byte;
+      std::stringstream ss(hex.substr(i*2, 2));
+      ss >> std::hex >> byte;
+      bytes[i] = static_cast<unsigned char>(byte);
+  }
+  return h;
+}
+
+int main(int argc, char** argv) {
+  if (argc != 2) {
+    return -1;
+  }
+
+  hipIpcMemHandle_t memHandle = hexToHipHandle(argv[1]);
+  bool IfTestPassed = true;
+
+  size_t N = 1024;
+  size_t Nbytes = N * sizeof(int);
+  int *B_d{nullptr}, *C_d{nullptr};
+  int *C_h{nullptr};
+
+  HIP_CHECK(hipHostMalloc(reinterpret_cast<void**>(&C_h), Nbytes, hipHostMallocDefault));
+  memset(reinterpret_cast<void*>(C_h), 0, Nbytes);
+
+  HIP_CHECK(hipMalloc(&C_d, Nbytes));
+  HIP_CHECK(hipIpcOpenMemHandle(reinterpret_cast<void **>(&B_d),
+                                memHandle,
+                                hipIpcMemLazyEnablePeerAccess));
+
+  HIP_CHECK(hipMemcpy(C_d, B_d, Nbytes, hipMemcpyDeviceToDevice));
+  HIP_CHECK(hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));
+  for (size_t i = 0; i < N; i ++) {
+    if (C_h[i] != 6) {
+      printf("mismatch at index: %zu with %d", i, C_h[i]);
+      IfTestPassed = false;
+      break;
+    }
+  }
+
+  // Checking if the data obtained from IPC shared memory is consistent
+  memset(reinterpret_cast<void*>(C_h), 0, Nbytes);
+  HIP_CHECK(hipMemcpy(C_h, B_d, Nbytes, hipMemcpyDeviceToHost));
+  for (size_t i = 0; i < N; i ++) {
+    if (C_h[i] != 6) {
+      printf("mismatch at index: %zu with %d", i, C_h[i]);
+      IfTestPassed = false;
+    }
+  }
+
+  HIP_CHECK(hipIpcCloseMemHandle(reinterpret_cast<void*>(B_d)));
+  HIP_CHECK(hipFree(C_d));
+
+  return (IfTestPassed == true);
+}