From 779beb09f2e401130b3d090e307c65bef1615405 Mon Sep 17 00:00:00 2001 From: Rupam Chetia Date: Tue, 16 Jul 2024 16:29:26 +0530 Subject: [PATCH] SWDEV-471105 - [catch2][dtest] Adding functional tests for hipMemPoolExportToShareableHandle, hipMemPoolExportPointer, hipMemPoolImportFromShareableHandle and hipMemPoolImportPointer Change-Id: Ief0f0d7f4c05d741e49eaffa54f11338c4bdc30f [ROCm/hip-tests commit: 9fd38dd3e1df4adb67214cd73fe924bcdc2c1e5f] --- .../catch/unit/memory/CMakeLists.txt | 7 + .../unit/memory/hipMemPoolExportPointer.cc | 81 +++ .../hipMemPoolExportToShareableHandle.cc | 517 ++++++++++++++++++ .../hipMemPoolImportFromShareableHandle.cc | 86 +++ .../unit/memory/hipMemPoolImportPointer.cc | 94 ++++ .../catch/unit/memory/mempool_common.hh | 263 ++++++++- 6 files changed, 1030 insertions(+), 18 deletions(-) create mode 100644 projects/hip-tests/catch/unit/memory/hipMemPoolExportPointer.cc create mode 100644 projects/hip-tests/catch/unit/memory/hipMemPoolExportToShareableHandle.cc create mode 100644 projects/hip-tests/catch/unit/memory/hipMemPoolImportFromShareableHandle.cc create mode 100644 projects/hip-tests/catch/unit/memory/hipMemPoolImportPointer.cc diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index 336a71b816..a125dbd4fe 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -83,6 +83,13 @@ set(TEST_SRC hipMemsetD32Async.cc hipMemsetD8Async.cc) +if(UNIX) + set(TEST_SRC ${TEST_SRC} hipMemPoolExportPointer.cc + hipMemPoolExportToShareableHandle.cc + hipMemPoolImportFromShareableHandle.cc + hipMemPoolImportPointer.cc) +endif() + if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} diff --git a/projects/hip-tests/catch/unit/memory/hipMemPoolExportPointer.cc b/projects/hip-tests/catch/unit/memory/hipMemPoolExportPointer.cc new file mode 100644 index 0000000000..4400ed3a84 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemPoolExportPointer.cc @@ -0,0 +1,81 @@ +/** +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** + * @addtogroup hipMemPoolExportPointer hipMemPoolExportPointer + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemPoolExportPointer(hipMemPoolPtrExportData* export_data, void* dev_ptr) ` - + * Export a memory pool allocation to another process. + */ + +#include "mempool_common.hh" + +constexpr int DATA_SIZE = 1024*1024; +constexpr size_t byte_size = DATA_SIZE * sizeof(int); + +/** + * Test Description + * ------------------------ + * - Negative Tests for hipMemPoolExportPointer. + * ------------------------ + * - unit/memory/hipMemPoolExportPointer.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemPoolExportPointer_Negative") { + hipMemPoolPtrExportData ptrExp; + hipShareableHdl sharedHandle; + hipMemPoolProps pool_props{}; + hipMemPool_t mempoolPfd; + checkMempoolSupported(0) + + // Create mempool with Posix File Descriptor + pool_props.allocType = hipMemAllocationTypePinned; + pool_props.location.id = 0; + pool_props.location.type = hipMemLocationTypeDevice; + pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; + HIP_CHECK(hipMemPoolCreate(&mempoolPfd, &pool_props)); + int *A_d; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A_d), + byte_size, mempoolPfd, 0)); + HIP_CHECK(hipStreamSynchronize(0)); + HIP_CHECK(hipMemPoolExportToShareableHandle(&sharedHandle, + mempoolPfd, hipMemHandleTypePosixFileDescriptor, 0)); + SECTION("Passing nullptr as export data") { + HIP_CHECK_ERROR(hipMemPoolExportPointer(nullptr, A_d), + hipErrorInvalidValue); + } + SECTION("Passing nullptr as device memory ptr") { + HIP_CHECK_ERROR(hipMemPoolExportPointer(&ptrExp, nullptr), + hipErrorInvalidValue); + } + HIP_CHECK(hipFree(reinterpret_cast(A_d))); + HIP_CHECK(hipMemPoolDestroy(mempoolPfd)); +} + +/** +* End doxygen group MemoryTest. +* @} +*/ diff --git a/projects/hip-tests/catch/unit/memory/hipMemPoolExportToShareableHandle.cc b/projects/hip-tests/catch/unit/memory/hipMemPoolExportToShareableHandle.cc new file mode 100644 index 0000000000..b9ec3ffca4 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemPoolExportToShareableHandle.cc @@ -0,0 +1,517 @@ +/* +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** + * @addtogroup hipMemPoolExportToShareableHandle hipMemPoolExportToShareableHandle + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemPoolExportToShareableHandle(void* shared_handle, + hipMemPool_t mem_pool, + hipMemAllocationHandleType handle_type, + unsigned int flags) ` - + * Exports a memory pool to the requested handle type. + */ + +#include "mempool_common.hh" + +constexpr int DATA_SIZE = 1024*1024; +constexpr size_t byte_size = DATA_SIZE * sizeof(int); + +/** + Kernel to perform Square of input data. + */ +static __global__ void square_kernel(int* Buff) { + int i = threadIdx.x + blockDim.x * blockIdx.x; + int temp = Buff[i] * Buff[i]; + Buff[i] = temp; +} + +/** + Fill with input and expected output data. + */ +static void fill_data(std::vector &A_h, std::vector &B_h, + std::vector &C_h) { + for (int i = 0; i < DATA_SIZE; i++) { + A_h[i] = i % 1024; + B_h[i] = 0; + C_h[i] = A_h[i]*A_h[i]; + } +} + +/** + * Test Description + * ------------------------ + * - Create mempool handle and allocate a memory chunk. Export + * the mempool and the pointer to the chunk. In the same process, + * Import the handle and the pointer in the same process. Use the + * pointer in kernel launch. + * ------------------------ + * - unit/memory/hipMemPoolExportImportToShareableHandle.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemPoolExportToShareableHandle_SameProc") { + hipMemPoolPtrExportData ptrExp; + hipShareableHdl sharedHandle; + std::vector A_h(DATA_SIZE), B_h(DATA_SIZE), C_h(DATA_SIZE); + fill_data(A_h, B_h, C_h); + hipMemPoolProps pool_props{}; + hipMemPool_t mempool, mempoolImp; + checkMempoolSupported(0) + HIP_CHECK(hipSetDevice(0)); + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + // Create mempool + pool_props.allocType = hipMemAllocationTypePinned; + pool_props.location.id = 0; + pool_props.location.type = hipMemLocationTypeDevice; + pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; + HIP_CHECK(hipMemPoolCreate(&mempool, &pool_props)); + // Allocate device memory from mempool + int *A_d; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A_d), + byte_size, mempool, stream)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h.data(), byte_size, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + // Export mempool + HIP_CHECK(hipMemPoolExportToShareableHandle(&sharedHandle, + mempool, hipMemHandleTypePosixFileDescriptor, 0)); + // Export A_d + HIP_CHECK(hipMemPoolExportPointer(&ptrExp, A_d)); + // Import mempool + HIP_CHECK(hipMemPoolImportFromShareableHandle(&mempoolImp, + (void*)sharedHandle, + hipMemHandleTypePosixFileDescriptor, 0)); + // Import and use pointer + void *ptrImp; + HIP_CHECK(hipMemPoolImportPointer(&ptrImp, mempoolImp, &ptrExp)); + square_kernel<<>>((int*)ptrImp); + HIP_CHECK(hipMemcpyAsync(B_h.data(), ptrImp, byte_size, + hipMemcpyDeviceToHost, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); + HIP_CHECK(hipFreeAsync(reinterpret_cast(A_d), stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipMemPoolDestroy(mempool)); +} + +/** + * Test Description + * ------------------------ + * - Multiprocess functionality test. Create mempool handle and + * allocate a memory chunk. Export the mempool and the pointer to + * the chunk. Import the mempool and the pointer in child process. + * Copy data to the memory chunk and launch kernel to perform + * operations on the data. + * ------------------------ + * - unit/memory/hipMemPoolExportImportToShareableHandle.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemPoolExportToShareableHandle_ChldUseHdl") { + std::vector A_h(DATA_SIZE), B_h(DATA_SIZE), C_h(DATA_SIZE); + fill_data(A_h, B_h, C_h); + int fd[2], fdSig[2]; + REQUIRE(pipe(fd) == 0); + REQUIRE(pipe(fdSig) == 0); + + auto pid = fork(); + REQUIRE(pid >= 0); + + if (pid == 0) { // child + REQUIRE(close(fd[1]) == 0); + REQUIRE(close(fdSig[0]) == 0); + // Wait for parent process to create the socket. + hipMemPoolPtrExportData ptrExp; + REQUIRE(read(fd[0], &ptrExp, sizeof(hipMemPoolPtrExportData)) >= 0); + // Open Socket as client + ipcSocketCom sockObj(false); + // Signal Parent process that Child is ready to receive msg + int sig = 0; + REQUIRE(write(fdSig[1], &sig, sizeof(int)) >= 0); + hipShareableHdl shdl; + // receive message from parent provess + checkSysCallErrors(sockObj.recvShareableHdl(&shdl)); + // Import mempool + hipMemPool_t mempoolImp; + HIP_CHECK(hipMemPoolImportFromShareableHandle(&mempoolImp, (void*)shdl, + hipMemHandleTypePosixFileDescriptor, 0)); + // Import and use pointer + void *ptrImp; + HIP_CHECK(hipMemPoolImportPointer(&ptrImp, mempoolImp, &ptrExp)); + square_kernel<<>>((int*)ptrImp); + HIP_CHECK(hipStreamSynchronize(0)); + // Import and use pointer + REQUIRE(close(fd[0]) == 0); + REQUIRE(close(fdSig[1]) == 0); + checkSysCallErrors(sockObj.closeThisSock()); + exit(0); + } else { // parent + REQUIRE(close(fd[0]) == 0); + REQUIRE(close(fdSig[1]) == 0); + + hipMemPoolProps pool_props{}; + checkMempoolSupported(0) + // Set property + hipMemPool_t mempool; + pool_props.allocType = hipMemAllocationTypePinned; + pool_props.location.id = 0; + pool_props.location.type = hipMemLocationTypeDevice; + pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; + HIP_CHECK(hipMemPoolCreate(&mempool, &pool_props)); + // Export mempool + hipShareableHdl shdl; + HIP_CHECK(hipMemPoolExportToShareableHandle(&shdl, mempool, + hipMemHandleTypePosixFileDescriptor, 0)); + // Allocate device memory from mempool + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + int *A_d; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A_d), + byte_size, mempool, stream)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h.data(), byte_size, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + hipMemPoolPtrExportData ptrExp; + // Export A_d + HIP_CHECK(hipMemPoolExportPointer(&ptrExp, A_d)); + // Create the socket for communication as Server + ipcSocketCom sockObj(true); + // Signal child process that socket is ready and share ptr to child + REQUIRE(write(fd[1], &ptrExp, sizeof(hipMemPoolPtrExportData)) >= 0); + // Wait for the child process to receive msg + int sig = 0; + REQUIRE(read(fdSig[0], &sig, sizeof(int)) >= 0); + checkSysCallErrors(sockObj.sendShareableHdl(shdl, pid)); + // Wait for child process to exit. + int status; + REQUIRE(wait(&status) >= 0); + REQUIRE(status == 0); + HIP_CHECK(hipMemcpyAsync(B_h.data(), A_d, byte_size, + hipMemcpyDeviceToHost, stream)); + // Free all resources + HIP_CHECK(hipFreeAsync(reinterpret_cast(A_d), stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipMemPoolDestroy(mempool)); + REQUIRE(close(fd[1]) == 0); + REQUIRE(close(fdSig[0]) == 0); + checkSysCallErrors(sockObj.closeThisSock()); + } +} + +/** + * Test Description + * ------------------------ + * - Multiprocess functionality test. Create mempool handle and + * allocate a memory chunk. Export the mempool and the pointer to + * the chunk. Import the mempool and the pointer in child process. + * In parent process change mempool property. Verify the change in + * child process. + * ------------------------ + * - unit/memory/hipMemPoolExportImportToShareableHandle.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.2 + */ +#if HT_AMD +TEST_CASE("Unit_hipMemPoolExportToShareableHandle_ChldCheckAccess") { + int fd[2], fdSig[2]; + REQUIRE(pipe(fd) == 0); + REQUIRE(pipe(fdSig) == 0); + + auto pid = fork(); + REQUIRE(pid >= 0); + + if (pid == 0) { // child + REQUIRE(close(fd[1]) == 0); + REQUIRE(close(fdSig[0]) == 0); + // Wait for parent process to create the socket. + int sig = 0; + REQUIRE(read(fd[0], &sig, sizeof(int)) >= 0); + // Open Socket as client + ipcSocketCom sockObj(false); + // Signal Parent process that Child is ready to receive msg + REQUIRE(write(fdSig[1], &sig, sizeof(int)) >= 0); + hipShareableHdl shdl; + // receive message from parent provess + checkSysCallErrors(sockObj.recvShareableHdl(&shdl)); + // Import mempool + hipMemPool_t mempoolImp; + HIP_CHECK(hipMemPoolImportFromShareableHandle(&mempoolImp, (void*)shdl, + hipMemHandleTypePosixFileDescriptor, 0)); + // Get and validate access for all devices + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + for (int dev = 0; dev < numDevices; dev++) { + hipMemAccessFlags flags; + hipMemLocation location; + location.type = hipMemLocationTypeDevice; + location.id = dev; + HIP_CHECK(hipMemPoolGetAccess(&flags, mempoolImp, &location)); + REQUIRE(flags == hipMemAccessFlagsProtReadWrite); + } + // Import and use pointer + REQUIRE(close(fd[0]) == 0); + REQUIRE(close(fdSig[1]) == 0); + checkSysCallErrors(sockObj.closeThisSock()); + exit(0); + } else { // parent + REQUIRE(close(fd[0]) == 0); + REQUIRE(close(fdSig[1]) == 0); + + hipMemPoolProps pool_props{}; + checkMempoolSupported(0) + // Set property + hipMemPool_t mempool; + pool_props.allocType = hipMemAllocationTypePinned; + pool_props.location.id = 0; + pool_props.location.type = hipMemLocationTypeDevice; + pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; + HIP_CHECK(hipMemPoolCreate(&mempool, &pool_props)); + // Set access to all devices + int numDevices = 0; + HIP_CHECK(hipGetDeviceCount(&numDevices)); + for (int dev = 0; dev < numDevices; dev++) { + checkMempoolSupported(dev) + hipMemAccessDesc accessDesc; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = dev; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + HIP_CHECK(hipMemPoolSetAccess(mempool, &accessDesc, 1)); + } + // Export mempool + hipShareableHdl shdl; + HIP_CHECK(hipMemPoolExportToShareableHandle(&shdl, mempool, + hipMemHandleTypePosixFileDescriptor, 0)); + // Create the socket for communication as Server + ipcSocketCom sockObj(true); + // Signal child process that socket is ready + int sig = 0; + REQUIRE(write(fd[1], &sig, sizeof(int)) >= 0); + // Wait for the child process to receive msg + REQUIRE(read(fdSig[0], &sig, sizeof(int)) >= 0); + checkSysCallErrors(sockObj.sendShareableHdl(shdl, pid)); + // Wait for child process to exit. + int status; + REQUIRE(wait(&status) >= 0); + REQUIRE(status == 0); + HIP_CHECK(hipMemPoolDestroy(mempool)); + REQUIRE(close(fd[1]) == 0); + REQUIRE(close(fdSig[0]) == 0); + checkSysCallErrors(sockObj.closeThisSock()); + } +} +#endif +/** + * Test Description + * ------------------------ + * - Multiprocess functionality test. Create mempool handle and + * allocate a memory chunk. Export the mempool and the pointer to + * the chunk. Import the mempool and the pointer in grandchild process. + * Copy data to the memory chunk and launch kernel to perform + * operations on the data. + + * ------------------------ + * - unit/memory/hipMemPoolExportImportToShareableHandle.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemPoolExportToShareableHandle_GrndChldUseHdl") { + std::vector A_h(DATA_SIZE), B_h(DATA_SIZE), C_h(DATA_SIZE); + fill_data(A_h, B_h, C_h); + int fd[2], fdSig[2], fdpid[2]; + REQUIRE(pipe(fd) == 0); + REQUIRE(pipe(fdSig) == 0); + REQUIRE(pipe(fdpid) == 0); + auto pid = fork(); + REQUIRE(pid >= 0); + + if (pid == 0) { // child + auto pid2 = fork(); + if (pid2 == 0) { // grandchild + REQUIRE(close(fd[1]) == 0); + REQUIRE(close(fdSig[0]) == 0); + // Wait for parent process to create the socket. + hipMemPoolPtrExportData ptrExp; + REQUIRE(read(fd[0], &ptrExp, sizeof(hipMemPoolPtrExportData)) >= 0); + // Open Socket as client + ipcSocketCom sockObj(false); + hipShareableHdl shdl; + // Signal Parent process that Child is ready to receive msg + int sig = 0; + REQUIRE(write(fdSig[1], &sig, sizeof(int)) >= 0); + // receive message from parent provess + checkSysCallErrors(sockObj.recvShareableHdl(&shdl)); + // Import mempool + hipMemPool_t mempoolImp; + HIP_CHECK(hipMemPoolImportFromShareableHandle(&mempoolImp, (void*)shdl, + hipMemHandleTypePosixFileDescriptor, 0)); + // Import and use pointer + void *ptrImp; + HIP_CHECK(hipMemPoolImportPointer(&ptrImp, mempoolImp, &ptrExp)); + square_kernel<<>>((int*)ptrImp); + HIP_CHECK(hipStreamSynchronize(0)); + REQUIRE(close(fd[0]) == 0); + REQUIRE(close(fdSig[1]) == 0); + checkSysCallErrors(sockObj.closeThisSock()); + exit(0); + } else { + int status; + REQUIRE(close(fdpid[0]) == 0); + REQUIRE(write(fdpid[1], &pid2, sizeof(pid2)) >= 0); + REQUIRE(wait(&status) >= 0); + REQUIRE(status == 0); + REQUIRE(close(fdpid[1]) == 0); + exit(0); + } + } else { // parent + REQUIRE(close(fd[0]) == 0); + REQUIRE(close(fdSig[1]) == 0); + REQUIRE(close(fdpid[1]) == 0); + int pid_grChld = 0; + REQUIRE(read(fdpid[0], &pid_grChld, sizeof(pid_grChld)) >= 0); + + hipMemPoolProps pool_props{}; + checkMempoolSupported(0) + // Set property + hipMemPool_t mempool; + pool_props.allocType = hipMemAllocationTypePinned; + pool_props.location.id = 0; + pool_props.location.type = hipMemLocationTypeDevice; + pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; + HIP_CHECK(hipMemPoolCreate(&mempool, &pool_props)); + // Export mempool + hipShareableHdl shdl; + HIP_CHECK(hipMemPoolExportToShareableHandle(&shdl, mempool, + hipMemHandleTypePosixFileDescriptor, 0)); + // Allocate device memory from mempool + hipStream_t stream; + HIP_CHECK(hipStreamCreate(&stream)); + int *A_d; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A_d), + byte_size, mempool, stream)); + HIP_CHECK(hipMemcpyAsync(A_d, A_h.data(), byte_size, + hipMemcpyHostToDevice, stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + hipMemPoolPtrExportData ptrExp; + // Export A_d + HIP_CHECK(hipMemPoolExportPointer(&ptrExp, A_d)); + + // Create the socket for communication as Server + ipcSocketCom sockObj(true); + // Signal child process that socket is ready and share ptr to child + REQUIRE(write(fd[1], &ptrExp, sizeof(hipMemPoolPtrExportData)) >= 0); + // Wait for the child process to receive msg + int sig = 0; + REQUIRE(read(fdSig[0], &sig, sizeof(int)) >= 0); + checkSysCallErrors(sockObj.sendShareableHdl(shdl, pid_grChld)); + // Wait for child process to exit. + int status; + REQUIRE(wait(&status) >= 0); + REQUIRE(status == 0); + HIP_CHECK(hipMemcpyAsync(B_h.data(), A_d, byte_size, + hipMemcpyDeviceToHost, stream)); + // Free all resources + HIP_CHECK(hipFreeAsync(reinterpret_cast(A_d), stream)); + HIP_CHECK(hipStreamSynchronize(stream)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); + // Free all resources + HIP_CHECK(hipStreamDestroy(stream)); + HIP_CHECK(hipMemPoolDestroy(mempool)); + REQUIRE(close(fd[1]) == 0); + REQUIRE(close(fdSig[0]) == 0); + REQUIRE(close(fdpid[0]) == 0); + checkSysCallErrors(sockObj.closeThisSock()); + } +} +/** + * Test Description + * ------------------------ + * - Negative Tests for hipMemPoolExportToShareableHandle. + * ------------------------ + * - unit/memory/hipMemPoolExportImportToShareableHandle.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemPoolExportToShareableHandle_Negative") { + hipShareableHdl sharedHandle; + hipMemPoolProps pool_props{}; + hipMemPool_t mempoolPfd, mempoolwoPfd; + checkMempoolSupported(0) + + // Create mempool with Posix File Descriptor + pool_props.allocType = hipMemAllocationTypePinned; + pool_props.location.id = 0; + pool_props.location.type = hipMemLocationTypeDevice; + pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; + HIP_CHECK(hipMemPoolCreate(&mempoolPfd, &pool_props)); + + // Create mempool without File Descriptor + pool_props.allocType = hipMemAllocationTypePinned; + pool_props.location.id = 0; + pool_props.location.type = hipMemLocationTypeDevice; + pool_props.handleTypes = hipMemHandleTypeNone; + HIP_CHECK(hipMemPoolCreate(&mempoolwoPfd, &pool_props)); + SECTION("Passing nullptr as handle") { + HIP_CHECK_ERROR(hipMemPoolExportToShareableHandle(nullptr, + mempoolPfd, hipMemHandleTypePosixFileDescriptor, 0), + hipErrorInvalidValue); + } + SECTION("Passing nullptr as mempool") { + HIP_CHECK_ERROR(hipMemPoolExportToShareableHandle(&sharedHandle, + nullptr, hipMemHandleTypePosixFileDescriptor, 0), + hipErrorInvalidValue); + } + SECTION("Passing invalid handle type") { + HIP_CHECK_ERROR(hipMemPoolExportToShareableHandle(&sharedHandle, + mempoolPfd, hipMemHandleTypeNone, 0), + hipErrorInvalidValue); + } + SECTION("Passing mempool without file descriptor") { + HIP_CHECK_ERROR(hipMemPoolExportToShareableHandle(&sharedHandle, + mempoolwoPfd, hipMemHandleTypePosixFileDescriptor, 0), + hipErrorInvalidValue); + } + HIP_CHECK(hipMemPoolDestroy(mempoolPfd)); + HIP_CHECK(hipMemPoolDestroy(mempoolwoPfd)); +} + +/** +* End doxygen group MemoryTest. +* @} +*/ diff --git a/projects/hip-tests/catch/unit/memory/hipMemPoolImportFromShareableHandle.cc b/projects/hip-tests/catch/unit/memory/hipMemPoolImportFromShareableHandle.cc new file mode 100644 index 0000000000..e5c5d05c27 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemPoolImportFromShareableHandle.cc @@ -0,0 +1,86 @@ +/** +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** + * @addtogroup hipMemPoolImportFromShareableHandle hipMemPoolImportFromShareableHandle + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemPoolImportFromShareableHandle( + hipMemPool_t* mem_pool, + void* shared_handle, + hipMemAllocationHandleType handle_type, + unsigned int flags) ` - + * Imports a memory pool from a shared handle. + */ + +#include "mempool_common.hh" + +/** + * Test Description + * ------------------------ + * - Negative Tests for hipMemPoolImportFromShareableHandle. + * ------------------------ + * - unit/memory/hipMemPoolImportFromShareableHandle.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemPoolImportFromShareableHandle_Negative") { + hipShareableHdl sharedHandle; + hipMemPoolProps pool_props{}; + hipMemPool_t mempoolPfd; + checkMempoolSupported(0) + + // Create mempool with Posix File Descriptor + pool_props.allocType = hipMemAllocationTypePinned; + pool_props.location.id = 0; + pool_props.location.type = hipMemLocationTypeDevice; + pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; + HIP_CHECK(hipMemPoolCreate(&mempoolPfd, &pool_props)); + + HIP_CHECK(hipMemPoolExportToShareableHandle(&sharedHandle, + mempoolPfd, hipMemHandleTypePosixFileDescriptor, 0)); + hipMemPool_t mempoolImp; + SECTION("Passing nullptr as imported mempool") { + HIP_CHECK_ERROR(hipMemPoolImportFromShareableHandle(nullptr, + (void*)sharedHandle, + hipMemHandleTypePosixFileDescriptor, 0), + hipErrorInvalidValue); + } + SECTION("Passing nullptr as handle") { + HIP_CHECK_ERROR(hipMemPoolImportFromShareableHandle(&mempoolImp, + nullptr, hipMemHandleTypePosixFileDescriptor, 0), + hipErrorInvalidValue); + } + SECTION("Passing invalid handle type") { + HIP_CHECK_ERROR(hipMemPoolImportFromShareableHandle(&mempoolImp, + (void*)sharedHandle, hipMemHandleTypeNone, 0), + hipErrorInvalidValue); + } + HIP_CHECK(hipMemPoolDestroy(mempoolPfd)); +} + +/** +* End doxygen group MemoryTest. +* @} +*/ diff --git a/projects/hip-tests/catch/unit/memory/hipMemPoolImportPointer.cc b/projects/hip-tests/catch/unit/memory/hipMemPoolImportPointer.cc new file mode 100644 index 0000000000..8ae1e21908 --- /dev/null +++ b/projects/hip-tests/catch/unit/memory/hipMemPoolImportPointer.cc @@ -0,0 +1,94 @@ +/** +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 WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS 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 IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** + * @addtogroup hipMemPoolImportPointer hipMemPoolImportPointer + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemPoolImportPointer( + void** dev_ptr, + hipMemPool_t mem_pool, + hipMemPoolPtrExportData* export_data) ` - + * Import a memory pool allocation from another process. + */ + +#include "mempool_common.hh" + +constexpr int DATA_SIZE = 1024*1024; +constexpr size_t byte_size = DATA_SIZE * sizeof(int); + +/** + * Test Description + * ------------------------ + * - Negative Tests for hipMemPoolImportPointer. + * ------------------------ + * - unit/memory/hipMemPoolImportPointer.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.2 + */ +TEST_CASE("Unit_hipMemPoolImportPointer_Negative") { + hipMemPoolPtrExportData ptrExp; + hipShareableHdl sharedHandle; + hipMemPoolProps pool_props{}; + hipMemPool_t mempoolPfd; + checkMempoolSupported(0) + + // Create mempool with Posix File Descriptor + pool_props.allocType = hipMemAllocationTypePinned; + pool_props.location.id = 0; + pool_props.location.type = hipMemLocationTypeDevice; + pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; + HIP_CHECK(hipMemPoolCreate(&mempoolPfd, &pool_props)); + int *A_d; + HIP_CHECK(hipMallocFromPoolAsync(reinterpret_cast(&A_d), + byte_size, mempoolPfd, 0)); + HIP_CHECK(hipStreamSynchronize(0)); + HIP_CHECK(hipMemPoolExportToShareableHandle(&sharedHandle, + mempoolPfd, hipMemHandleTypePosixFileDescriptor, 0)); + HIP_CHECK(hipMemPoolExportPointer(&ptrExp, A_d)); + hipMemPool_t mempoolImp; + HIP_CHECK(hipMemPoolImportFromShareableHandle(&mempoolImp, + (void*)sharedHandle, + hipMemHandleTypePosixFileDescriptor, 0)); + void *ptrImp; + SECTION("Passing nullptr as import data") { + HIP_CHECK_ERROR(hipMemPoolImportPointer(nullptr, mempoolImp, &ptrExp), + hipErrorInvalidValue); + } + SECTION("Passing nullptr as imported mempool") { + HIP_CHECK_ERROR(hipMemPoolImportPointer(&ptrImp, nullptr, &ptrExp), + hipErrorInvalidValue); + } + SECTION("Passing nullptr as exported pointer") { + HIP_CHECK_ERROR(hipMemPoolImportPointer(&ptrImp, mempoolImp, nullptr), + hipErrorInvalidValue); + } + HIP_CHECK(hipFree(reinterpret_cast(A_d))); + HIP_CHECK(hipMemPoolDestroy(mempoolPfd)); +} + +/** +* End doxygen group MemoryTest. +* @} +*/ diff --git a/projects/hip-tests/catch/unit/memory/mempool_common.hh b/projects/hip-tests/catch/unit/memory/mempool_common.hh index 21479005a4..05876fe721 100644 --- a/projects/hip-tests/catch/unit/memory/mempool_common.hh +++ b/projects/hip-tests/catch/unit/memory/mempool_common.hh @@ -23,29 +23,27 @@ #include #include +#ifdef __linux__ +#include +#include +#include +#include +#include +#include +#include +#include +#include +#endif + namespace { constexpr auto wait_ms = 500; } // anonymous namespace -/** - * Local Function to test if Hip Stream Ordered Memory allocator - * functionality is supoorted. - */ -static bool isStrmOrdMemAllocSupported(int dev) { - int deviceSupportsMemoryPools = 0; - bool supported = false; - HIP_CHECK(hipDeviceGetAttribute(&deviceSupportsMemoryPools, - hipDeviceAttributeMemoryPoolsSupported, dev)); - if (deviceSupportsMemoryPools != 0) { - supported = true; - } else { - supported = false; - } - return supported; -} - #define checkMempoolSupported(device) {\ - if (false == isStrmOrdMemAllocSupported(device)) {\ + int deviceSupportsMemoryPools = 0;\ + HIP_CHECK(hipDeviceGetAttribute(&deviceSupportsMemoryPools,\ + hipDeviceAttributeMemoryPoolsSupported, device));\ + if (0 == deviceSupportsMemoryPools) {\ HipTest::HIP_SKIP_TEST("Memory Pool not supported. Skipping Test..");\ return;\ }\ @@ -450,3 +448,232 @@ class streamMemAllocTest { free(C_h); } }; + +#ifdef __linux__ + +#define checkSysCallErrors(result) \ + if (result == -1) { \ + fprintf(stderr, "Failure at %u %s\n", __LINE__, __FILE__); exit(EXIT_FAILURE); \ + } + +#ifdef HT_AMD +typedef int64_t hipShareableHdl; +#else +typedef int hipShareableHdl; +#endif + +typedef pid_t Process; + +struct ipcHdl { + int socket; + char *name; +}; + +class ipcSocketCom { + ipcHdl *handle; + // method to create socket from server + int createSocket() { + int server_fd; + struct sockaddr_un servaddr; + + char name[16]; + // Create a unique socket name based on current pid + sprintf(name, "%u", getpid()); + + // Create the socket handle + handle = new ipcHdl; + if (nullptr == handle) { + perror("Socket failure: Handle memory allocation failed"); + return -1; + } + + memset(handle, 0, sizeof(*handle)); + handle->socket = -1; + handle->name = NULL; + + // Creating socket + if ((server_fd = socket(AF_UNIX, SOCK_DGRAM, 0)) == 0) { + perror("Socket failure: Socket creation failed"); + return -1; + } + + unlink(name); + bzero(&servaddr, sizeof(servaddr)); + servaddr.sun_family = AF_UNIX; + + size_t len = strlen(name); + if (len > (sizeof(servaddr.sun_path) - 1)) { + perror("Socket failure: Cannot bind provided name to socket. Name too large"); + return -1; + } + + strncpy(servaddr.sun_path, name, len); + + if (bind(server_fd, (struct sockaddr *)&servaddr, SUN_LEN(&servaddr)) < 0) { + perror("Socket failure: Binding socket failed"); + return -1; + } + + handle->name = new char[strlen(name) + 1]; + strcpy(handle->name, name); + handle->socket = server_fd; + return 0; + } + // method to create socket from client + int openSocket() { + int sock = 0; + struct sockaddr_un cliaddr; + + handle = new ipcHdl; + if (nullptr == handle) { + perror("Socket failure: Handle memory allocation failed"); + return -1; + } + memset(handle, 0, sizeof(*handle)); + + if ((sock = socket(AF_UNIX, SOCK_DGRAM, 0)) < 0) { + perror("IPC failure:Socket creation error"); + return -1; + } + + bzero(&cliaddr, sizeof(cliaddr)); + cliaddr.sun_family = AF_UNIX; + char name[16]; + + // Create a unique socket name based on current process id. + sprintf(name, "%u", getpid()); + + strcpy(cliaddr.sun_path, name); + if (bind(sock, (struct sockaddr *)&cliaddr, sizeof(cliaddr)) < 0) { + perror("Socket failure: Binding socket failed"); + return -1; + } + + handle->socket = sock; + handle->name = new char[strlen(name) + 1]; + strcpy(handle->name, name); + + return 0; + } + // method to close socket + int closeSocket() { + if (!handle) { + return -1; + } + + if (handle->name) { + unlink(handle->name); + delete[] handle->name; + } + close(handle->socket); + delete handle; + return 0; + } +public: + ipcSocketCom() = default; + ipcSocketCom(bool isServer) { + if (isServer) { + checkSysCallErrors(createSocket()); + } else { + checkSysCallErrors(openSocket()); + } + } + ~ipcSocketCom() { + } + int closeThisSock() { + return closeSocket(); + } + // method to receive shareable handle via socket + int recvShareableHdl(hipShareableHdl *shHandle) { + struct msghdr msg; + struct iovec iov[1]; + + // Union to guarantee alignment requirements for control array + union { + struct cmsghdr cm; + char control[CMSG_SPACE(sizeof(int))]; + } control_un; + + struct cmsghdr *cmptr; + ssize_t n; + int receivedfd; + int dummy_data; + + msg.msg_name = NULL; + msg.msg_namelen = 0; + msg.msg_control = control_un.control; + msg.msg_controllen = sizeof(control_un.control); + iov[0].iov_base = &dummy_data; + iov[0].iov_len = sizeof(dummy_data); + msg.msg_iov = iov; + msg.msg_iovlen = 1; + + if ((n = recvmsg(handle->socket, &msg, 0)) <= 0) { + perror("Socket failure: Receiving data over socket failed"); + return -1; + } + + if (((cmptr = CMSG_FIRSTHDR(&msg)) != NULL) && + (cmptr->cmsg_len == CMSG_LEN(sizeof(int)))) { + if ((cmptr->cmsg_level != SOL_SOCKET) || (cmptr->cmsg_type != SCM_RIGHTS)) { + return -1; + } + + memmove(&receivedfd, CMSG_DATA(cmptr), sizeof(receivedfd)); + *(int *)shHandle = receivedfd; + } else { + return -1; + } + + return 0; + } + // method to send shareable handle via sockets + int sendShareableHdl(hipShareableHdl shareableHdl, Process process) { + struct msghdr msg; + struct iovec iov[1]; + int dummy_data = 0; + + union { + struct cmsghdr cm; + char control[CMSG_SPACE(sizeof(int))]; + } control_un; + + struct cmsghdr *cmptr; + struct sockaddr_un cliaddr; + + // Construct client address to send this SHareable handle to + bzero(&cliaddr, sizeof(cliaddr)); + cliaddr.sun_family = AF_UNIX; + char temp[10]; + sprintf(temp, "%u", process); + strcpy(cliaddr.sun_path, temp); + + // Send corresponding shareable handle to the client + int sendfd = (int)shareableHdl; + + msg.msg_control = control_un.control; + msg.msg_controllen = sizeof(control_un.control); + + cmptr = CMSG_FIRSTHDR(&msg); + cmptr->cmsg_len = CMSG_LEN(sizeof(int)); + cmptr->cmsg_level = SOL_SOCKET; + cmptr->cmsg_type = SCM_RIGHTS; + + memmove(CMSG_DATA(cmptr), &sendfd, sizeof(sendfd)); + + msg.msg_name = (void *)&cliaddr; + msg.msg_namelen = sizeof(struct sockaddr_un); + iov[0].iov_base = &dummy_data; + iov[0].iov_len = sizeof(dummy_data); + msg.msg_iov = iov; + msg.msg_iovlen = 1; + + ssize_t sendResult = sendmsg(handle->socket, &msg, 0); + if (sendResult <= 0) { + perror("Socket failure: Sending data over socket failed"); + return -1; + } + return 0; + } +}; +#endif