diff --git a/catch/unit/device/CMakeLists.txt b/catch/unit/device/CMakeLists.txt index 59f64028f..7ca9e1f95 100644 --- a/catch/unit/device/CMakeLists.txt +++ b/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) diff --git a/catch/unit/device/hipIpcMemCopy.cc b/catch/unit/device/hipIpcMemCopy.cc new file mode 100644 index 000000000..2a172c49a --- /dev/null +++ b/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 +#include + +#include +#include + +#include +#include +#include + +std::string hipHandleToHex(const hipIpcMemHandle_t &h) { + std::ostringstream oss; + auto bytes = reinterpret_cast(&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(&A_h), Nbytes, hipHostMallocDefault)); + int *A_d{nullptr}; + HIP_CHECK(hipMalloc(reinterpret_cast(&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); +} diff --git a/catch/unit/device/hipIpcMemCopyChildProc_Exe.cc b/catch/unit/device/hipIpcMemCopyChildProc_Exe.cc new file mode 100644 index 000000000..581785b00 --- /dev/null +++ b/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 +#include +#include +#include +#include + +#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(&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(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(&C_h), Nbytes, hipHostMallocDefault)); + memset(reinterpret_cast(C_h), 0, Nbytes); + + HIP_CHECK(hipMalloc(&C_d, Nbytes)); + HIP_CHECK(hipIpcOpenMemHandle(reinterpret_cast(&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(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(B_d))); + HIP_CHECK(hipFree(C_d)); + + return (IfTestPassed == true); +} \ No newline at end of file