From eff4f59a7d3241c3f5cd370640587a632aa7e216 Mon Sep 17 00:00:00 2001 From: Rakesh Roy Date: Fri, 8 Mar 2024 12:01:07 +0530 Subject: [PATCH] SWDEV-1 - Revert "EXSWHTEC-57 - Implement tests for Import/Export hipMemPool APIs #445" This reverts commit d1500f261269e40ae1eceec1b5884558aa570efc. Change-Id: I4955ae6db9385969b736f4ef85dff8d291c06c13 --- catch/unit/CMakeLists.txt | 1 - catch/unit/stream_ordered/CMakeLists.txt | 30 - .../stream_ordered/helper_multiprocess.cc | 518 ------------------ .../stream_ordered/helper_multiprocess.hh | 119 ---- .../stream_ordered/hipMemPoolExportImport.cc | 495 ----------------- .../hipMemPoolExportImportIPC.cc | 422 -------------- 6 files changed, 1585 deletions(-) delete mode 100644 catch/unit/stream_ordered/CMakeLists.txt delete mode 100644 catch/unit/stream_ordered/helper_multiprocess.cc delete mode 100644 catch/unit/stream_ordered/helper_multiprocess.hh delete mode 100644 catch/unit/stream_ordered/hipMemPoolExportImport.cc delete mode 100644 catch/unit/stream_ordered/hipMemPoolExportImportIPC.cc diff --git a/catch/unit/CMakeLists.txt b/catch/unit/CMakeLists.txt index 6b63292c91..d43e957c42 100644 --- a/catch/unit/CMakeLists.txt +++ b/catch/unit/CMakeLists.txt @@ -22,7 +22,6 @@ add_subdirectory(rtc) add_subdirectory(deviceLib) add_subdirectory(graph) add_subdirectory(memory) -add_subdirectory(stream_ordered) add_subdirectory(stream) add_subdirectory(event) add_subdirectory(occupancy) diff --git a/catch/unit/stream_ordered/CMakeLists.txt b/catch/unit/stream_ordered/CMakeLists.txt deleted file mode 100644 index e9a5e56a8d..0000000000 --- a/catch/unit/stream_ordered/CMakeLists.txt +++ /dev/null @@ -1,30 +0,0 @@ -# Copyright (c) 2023 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. - -# Common Tests - Test independent of all platforms -set(COMMON_SHARED_SRC helper_multiprocess.cc) - -set(TEST_SRC - hipMemPoolExportImport.cc - hipMemPoolExportImportIPC.cc) - -hip_add_exe_to_target(NAME StreamOrderedTest - TEST_SRC ${TEST_SRC} - TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) diff --git a/catch/unit/stream_ordered/helper_multiprocess.cc b/catch/unit/stream_ordered/helper_multiprocess.cc deleted file mode 100644 index 77797fa71b..0000000000 --- a/catch/unit/stream_ordered/helper_multiprocess.cc +++ /dev/null @@ -1,518 +0,0 @@ -/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved. - * - * Redistribution and use in source and binary forms, with or without - * modification, are permitted provided that the following conditions - * are met: - * * Redistributions of source code must retain the above copyright - * notice, this list of conditions and the following disclaimer. - * * Redistributions in binary form must reproduce the above copyright - * notice, this list of conditions and the following disclaimer in the - * documentation and/or other materials provided with the distribution. - * * Neither the name of NVIDIA CORPORATION nor the names of its - * contributors may be used to endorse or promote products derived - * from this software without specific prior written permission. - * - * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY - * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE - * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR - * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR - * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, - * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, - * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR - * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY - * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT - * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE - * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - */ - -#include "helper_multiprocess.hh" -#include -#include - -int sharedMemoryCreate(const char* name, size_t sz, sharedMemoryInfo* info) { -#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) - info->size = sz; - info->shmHandle = - CreateFileMapping(INVALID_HANDLE_VALUE, NULL, PAGE_READWRITE, 0, (DWORD)sz, name); - if (info->shmHandle == 0) { - return GetLastError(); - } - - info->addr = MapViewOfFile(info->shmHandle, FILE_MAP_ALL_ACCESS, 0, 0, sz); - if (info->addr == NULL) { - return GetLastError(); - } - - return 0; -#else - int status = 0; - - info->size = sz; - - info->shmFd = shm_open(name, O_RDWR | O_CREAT, 0777); - if (info->shmFd < 0) { - return errno; - } - - status = ftruncate(info->shmFd, sz); - if (status != 0) { - return status; - } - - info->addr = mmap(0, sz, PROT_READ | PROT_WRITE, MAP_SHARED, info->shmFd, 0); - if (info->addr == NULL) { - return errno; - } - - return 0; -#endif -} - -int sharedMemoryOpen(const char* name, size_t sz, sharedMemoryInfo* info) { -#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) - info->size = sz; - - info->shmHandle = OpenFileMapping(FILE_MAP_ALL_ACCESS, FALSE, name); - if (info->shmHandle == 0) { - return GetLastError(); - } - - info->addr = MapViewOfFile(info->shmHandle, FILE_MAP_ALL_ACCESS, 0, 0, sz); - if (info->addr == NULL) { - return GetLastError(); - } - - return 0; -#else - info->size = sz; - - info->shmFd = shm_open(name, O_RDWR, 0777); - if (info->shmFd < 0) { - return errno; - } - - info->addr = mmap(0, sz, PROT_READ | PROT_WRITE, MAP_SHARED, info->shmFd, 0); - if (info->addr == NULL) { - return errno; - } - - return 0; -#endif -} - -void sharedMemoryClose(sharedMemoryInfo* info) { -#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) - if (info->addr) { - UnmapViewOfFile(info->addr); - } - if (info->shmHandle) { - CloseHandle(info->shmHandle); - } -#else - if (info->addr) { - munmap(info->addr, info->size); - } - if (info->shmFd) { - close(info->shmFd); - } -#endif -} - -int spawnProcess(Process* process, const char* app, char* const* args) { -#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) - STARTUPINFO si = {}; - memset(&si, 0, sizeof(STARTUPINFO)); - BOOL status; - std::string arg_string; - memset(process, 0, sizeof(*process)); - - while (*args) { - arg_string.append(*args).append(1, ' '); - args++; - } - - status = - CreateProcess(app, LPSTR(arg_string.c_str()), NULL, NULL, FALSE, 0, NULL, NULL, &si, process); - - return status ? 0 : GetLastError(); -#else - *process = fork(); - if (*process == 0) { - if (0 > execvp(app, args)) { - return errno; - } - } else if (*process < 0) { - return errno; - } - return 0; -#endif -} - -int waitProcess(Process* process) { -#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) - DWORD exitCode; - WaitForSingleObject(process->hProcess, INFINITE); - GetExitCodeProcess(process->hProcess, &exitCode); - CloseHandle(process->hProcess); - CloseHandle(process->hThread); - return (int)exitCode; -#else - int status = 0; - do { - if (0 > waitpid(*process, &status, 0)) { - return errno; - } - } while (!WIFEXITED(status)); - return WEXITSTATUS(status); -#endif -} - -#if defined(__linux__) -int ipcCreateSocket(ipcHandle*& handle, const char* name, const std::vector& /*processes*/) { - int server_fd; - struct sockaddr_un servaddr; - - handle = new ipcHandle; - memset(handle, 0, sizeof(*handle)); - handle->socket = -1; - handle->socketName = NULL; - - // Creating socket file descriptor - if ((server_fd = socket(AF_UNIX, SOCK_DGRAM, 0)) == 0) { - perror("IPC 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("IPC 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("IPC failure: Binding socket failed"); - return -1; - } - - handle->socketName = new char[strlen(name) + 1]; - strcpy(handle->socketName, name); - handle->socket = server_fd; - return 0; -} - -int ipcOpenSocket(ipcHandle*& handle) { - int sock = 0; - struct sockaddr_un cliaddr; - - handle = new ipcHandle; - 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 temp[10]; - - // Create unique name for the socket. - sprintf(temp, "%u", getpid()); - - strcpy(cliaddr.sun_path, temp); - if (bind(sock, (struct sockaddr*)&cliaddr, sizeof(cliaddr)) < 0) { - perror("IPC failure: Binding socket failed"); - return -1; - } - - handle->socket = sock; - handle->socketName = new char[strlen(temp) + 1]; - strcpy(handle->socketName, temp); - - return 0; -} - -int ipcCloseSocket(ipcHandle* handle) { - if (!handle) { - return -1; - } - - if (handle->socketName) { - unlink(handle->socketName); - delete[] handle->socketName; - } - close(handle->socket); - delete handle; - return 0; -} - -int ipcRecvShareableHandle(ipcHandle* handle, ShareableHandle* shHandle) { - struct msghdr msg = {}; - memset(&msg, 0, sizeof(struct msghdr)); - 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; - int receivedfd; - char dummy_buffer[1]; - - msg.msg_control = control_un.control; - msg.msg_controllen = sizeof(control_un.control); - - iov[0].iov_base = (void*)dummy_buffer; - iov[0].iov_len = sizeof(dummy_buffer); - - msg.msg_iov = iov; - msg.msg_iovlen = 1; - - if (recvmsg(handle->socket, &msg, 0) <= 0) { - perror("IPC 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; -} - -int ipcRecvDataFromClient(ipcHandle* serverHandle, void* data, size_t size) { - ssize_t readResult; - struct sockaddr_un cliaddr; - socklen_t len = sizeof(cliaddr); - - readResult = recvfrom(serverHandle->socket, data, size, 0, (struct sockaddr*)&cliaddr, &len); - if (readResult == -1) { - perror("IPC failure: Receiving data over socket failed"); - return -1; - } - return 0; -} - -int ipcSendDataToServer(ipcHandle* handle, const char* serverName, const void* data, size_t size) { - ssize_t sendResult; - struct sockaddr_un serveraddr; - - bzero(&serveraddr, sizeof(serveraddr)); - serveraddr.sun_family = AF_UNIX; - strncpy(serveraddr.sun_path, serverName, sizeof(serveraddr.sun_path) - 1); - - sendResult = - sendto(handle->socket, data, size, 0, (struct sockaddr*)&serveraddr, sizeof(serveraddr)); - if (sendResult <= 0) { - perror("IPC failure: Sending data over socket failed"); - } - - return 0; -} - -int ipcSendShareableHandle(ipcHandle* handle, const std::vector& shareableHandles, - Process process, int data) { - struct msghdr msg; - struct iovec iov[1]; - - 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)shareableHandles[data]; - - 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 = (void*)""; - iov[0].iov_len = 1; - msg.msg_iov = iov; - msg.msg_iovlen = 1; - - ssize_t sendResult = sendmsg(handle->socket, &msg, 0); - if (sendResult <= 0) { - perror("IPC failure: Sending data over socket failed"); - return -1; - } - - return 0; -} - -int ipcSendShareableHandles(ipcHandle* handle, const std::vector& shareableHandles, - const std::vector& processes) { - // Send all shareable handles to every single process. - for (unsigned int i = 0; i < shareableHandles.size(); i++) { - for (unsigned int j = 0; j < processes.size(); j++) { - checkIpcErrors(ipcSendShareableHandle(handle, shareableHandles, processes[j], i)); - } - } - return 0; -} - -int ipcRecvShareableHandles(ipcHandle* handle, std::vector& shareableHandles) { - for (unsigned int i = 0; i < shareableHandles.size(); i++) { - checkIpcErrors(ipcRecvShareableHandle(handle, &shareableHandles[i])); - } - return 0; -} - -int ipcCloseShareableHandle(ShareableHandle shHandle) { return close(shHandle); } - -#elif defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) -// Generic name to build individual Mailslot names by appending process ids. -LPTSTR SlotName = (LPTSTR)TEXT("\\\\.\\mailslot\\sample_mailslot_"); - -int ipcCreateSocket(ipcHandle*& handle, const char*, const std::vector& processes) { - handle = new ipcHandle; - handle->hMailslot.resize(processes.size()); - - // Open Mailslots of all clients and store respective handles. - for (unsigned int i = 0; i < handle->hMailslot.size(); ++i) { - std::basic_string childSlotName(SlotName); - char tempBuf[20]; - _itoa_s(processes[i].dwProcessId, tempBuf, 10); - childSlotName += TEXT(tempBuf); - - HANDLE hFile = - CreateFile(TEXT(childSlotName.c_str()), GENERIC_WRITE, FILE_SHARE_READ, - (LPSECURITY_ATTRIBUTES)NULL, OPEN_EXISTING, FILE_ATTRIBUTE_NORMAL, (HANDLE)NULL); - if (hFile == INVALID_HANDLE_VALUE) { - printf("IPC failure: Opening Mailslot by CreateFile failed with %lu\n", GetLastError()); - return -1; - } - handle->hMailslot[i] = hFile; - } - return 0; -} - -int ipcOpenSocket(ipcHandle*& handle) { - handle = new ipcHandle; - HANDLE hSlot; - - std::basic_string clientSlotName(SlotName); - char tempBuf[20]; - _itoa_s(GetCurrentProcessId(), tempBuf, 10); - clientSlotName += TEXT(tempBuf); - - hSlot = CreateMailslot((LPSTR)clientSlotName.c_str(), 0, MAILSLOT_WAIT_FOREVER, - (LPSECURITY_ATTRIBUTES)NULL); - if (hSlot == INVALID_HANDLE_VALUE) { - printf("IPC failure: CreateMailslot failed for client with %lu\n", GetLastError()); - return -1; - } - - handle->hMailslot.push_back(hSlot); - return 0; -} - -int ipcSendData(HANDLE mailslot, const void* data, size_t sz) { - BOOL result; - DWORD cbWritten; - - result = WriteFile(mailslot, data, (DWORD)sz, &cbWritten, (LPOVERLAPPED)NULL); - if (!result) { - printf("IPC failure: WriteFile failed with %lu.\n", GetLastError()); - return -1; - } - return 0; -} - -int ipcRecvData(ipcHandle* handle, void* data, size_t sz) { - DWORD cbRead = 0; - - if (!ReadFile(handle->hMailslot[0], data, (DWORD)sz, &cbRead, NULL)) { - printf("IPC failure: ReadFile failed with %lu.\n", GetLastError()); - return -1; - } - - if (sz != (size_t)cbRead) { - printf("IPC failure: ReadFile didn't receive the expected number of bytes\n"); - return -1; - } - - return 0; -} - -int ipcSendShareableHandles(ipcHandle* handle, const std::vector& shareableHandles, - const std::vector& processes) { - // Send all shareable handles to every single process. - for (unsigned int i = 0; i < processes.size(); i++) { - HANDLE hProcess = OpenProcess(PROCESS_DUP_HANDLE, FALSE, processes[i].dwProcessId); - if (hProcess == INVALID_HANDLE_VALUE) { - printf("IPC failure: OpenProcess failed (%lu)\n", GetLastError()); - return -1; - } - - for (unsigned int j = 0; j < shareableHandles.size(); j++) { - HANDLE hDup = INVALID_HANDLE_VALUE; - // Duplicate the handle into the target process's space - if (!DuplicateHandle(GetCurrentProcess(), shareableHandles[j], hProcess, &hDup, 0, FALSE, - DUPLICATE_SAME_ACCESS)) { - printf("IPC failure: DuplicateHandle failed (%lu)\n", GetLastError()); - return -1; - } - checkIpcErrors(ipcSendData(handle->hMailslot[i], &hDup, sizeof(hDup))); - } - CloseHandle(hProcess); - } - return 0; -} - -int ipcRecvShareableHandles(ipcHandle* handle, std::vector& shareableHandles) { - for (unsigned int i = 0; i < shareableHandles.size(); i++) { - checkIpcErrors(ipcRecvData(handle, &shareableHandles[i], sizeof(shareableHandles[i]))); - } - return 0; -} - -int ipcCloseSocket(ipcHandle* handle) { - for (unsigned int i = 0; i < handle->hMailslot.size(); i++) { - CloseHandle(handle->hMailslot[i]); - } - delete handle; - return 0; -} - -int ipcCloseShareableHandle(ShareableHandle shHandle) { - CloseHandle(shHandle); - return 0; -} - -#endif diff --git a/catch/unit/stream_ordered/helper_multiprocess.hh b/catch/unit/stream_ordered/helper_multiprocess.hh deleted file mode 100644 index b9bb4e16ea..0000000000 --- a/catch/unit/stream_ordered/helper_multiprocess.hh +++ /dev/null @@ -1,119 +0,0 @@ -/* Copyright (c) 2022, NVIDIA CORPORATION. All rights reserved. - * - * Redistribution and use in source and binary forms, with or without - * modification, are permitted provided that the following conditions - * are met: - * * Redistributions of source code must retain the above copyright - * notice, this list of conditions and the following disclaimer. - * * Redistributions in binary form must reproduce the above copyright - * notice, this list of conditions and the following disclaimer in the - * documentation and/or other materials provided with the distribution. - * * Neither the name of NVIDIA CORPORATION nor the names of its - * contributors may be used to endorse or promote products derived - * from this software without specific prior written permission. - * - * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS ``AS IS'' AND ANY - * EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE - * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR - * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR - * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, - * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, - * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR - * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY - * OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT - * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE - * OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - */ - -#ifndef HELPER_MULTIPROCESS_H -#define HELPER_MULTIPROCESS_H - -#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) -#ifndef WIN32_LEAN_AND_MEAN -#define WIN32_LEAN_AND_MEAN -#endif -#include -#include -#include -#include -#include -#include -#include -#include -#else -#include -#include -#include -#include -#include -#include -#include -#include -#include -#include -#endif -#include - -typedef struct sharedMemoryInfo_st { - void* addr; - size_t size; -#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) - HANDLE shmHandle; -#else - int shmFd; -#endif -} sharedMemoryInfo; - -int sharedMemoryCreate(const char* name, size_t sz, sharedMemoryInfo* info); - -int sharedMemoryOpen(const char* name, size_t sz, sharedMemoryInfo* info); - -void sharedMemoryClose(sharedMemoryInfo* info); - - -#if defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) -typedef PROCESS_INFORMATION Process; -#else -typedef pid_t Process; -#endif - -int spawnProcess(Process* process, const char* app, char* const* args); - -int waitProcess(Process* process); - -#define checkIpcErrors(ipcFuncResult) \ - if (ipcFuncResult == -1) { \ - fprintf(stderr, "Failure at %u %s\n", __LINE__, __FILE__); \ - exit(EXIT_FAILURE); \ - } - -#if defined(__linux__) -struct ipcHandle_st { - int socket; - char* socketName; -}; -typedef int ShareableHandle; -#elif defined(WIN32) || defined(_WIN32) || defined(WIN64) || defined(_WIN64) -struct ipcHandle_st { - std::vector - hMailslot; // 1 Handle in case of child and `num children` Handles for parent. -}; -typedef HANDLE ShareableHandle; -#endif - -typedef struct ipcHandle_st ipcHandle; - -int ipcCreateSocket(ipcHandle*& handle, const char* name, const std::vector& processes); - -int ipcOpenSocket(ipcHandle*& handle); - -int ipcCloseSocket(ipcHandle* handle); - -int ipcRecvShareableHandles(ipcHandle* handle, std::vector& shareableHandles); - -int ipcSendShareableHandles(ipcHandle* handle, const std::vector& shareableHandles, - const std::vector& processes); - -int ipcCloseShareableHandle(ShareableHandle shHandle); - -#endif // HELPER_MULTIPROCESS_H diff --git a/catch/unit/stream_ordered/hipMemPoolExportImport.cc b/catch/unit/stream_ordered/hipMemPoolExportImport.cc deleted file mode 100644 index 94e2954692..0000000000 --- a/catch/unit/stream_ordered/hipMemPoolExportImport.cc +++ /dev/null @@ -1,495 +0,0 @@ -/* - Copyright (c) 2023 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, INNCLUDING 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 ANNY 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. - */ - -#ifdef _WIN64 -#define NOMINMAX -#endif /* _WIN64 */ - -#include -#include -#include - -/** - * @addtogroup hipMemPoolExportToShareableHandle hipMemPoolExportToShareableHandle - * @{ - * @ingroup StreamOTest - * `hipMemPoolExportToShareableHandle(void* shared_handle, hipMemPool_t mem_pool, - * hipMemAllocationHandleType handle_type, unsigned int flags)` - Exports a memory pool to the - * requested handle type. - */ - -/** - * Test Description - * ------------------------ - * - Basic test to verify exporting/importing a shareable handle on a single device in a single - * process. - * Test source - * ------------------------ - * - /unit/memory/hipMemPoolExportImport.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.0 - */ -TEST_CASE("Unit_hipMemPoolExportImport_Functional") { - int mem_pool_support = 0; - HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); - if (!mem_pool_support) { - SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); - return; - } - - int shareable_handle; - hipMemPoolPtrExportData export_ptr; - void* ptr; - - hipMemAllocationHandleType handle_type = hipMemHandleTypePosixFileDescriptor; - HIP_CHECK(hipSetDevice(0)); - StreamGuard stream(Streams::withFlags, hipStreamNonBlocking); - - hipMemPool_t mempool; - hipMemPoolProps pool_props; - memset(&pool_props, 0, sizeof(hipMemPoolProps)); - pool_props.allocType = hipMemAllocationTypePinned; - pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; - pool_props.location.type = hipMemLocationTypeDevice; - pool_props.location.id = 0; - - HIP_CHECK(hipMemPoolCreate(&mempool, &pool_props)); - - // Allocate memory in a stream from the pool just created - HIP_CHECK(hipMallocFromPoolAsync(&ptr, kPageSize, mempool, stream.stream())); - - HIP_CHECK(hipMemPoolExportToShareableHandle(&shareable_handle, mempool, handle_type, 0)); - - memset((void*)&export_ptr, 0, sizeof(hipMemPoolPtrExportData)); - HIP_CHECK(hipMemPoolExportPointer(&export_ptr, reinterpret_cast(ptr))); - - LinearAllocGuard host_ptr(LinearAllocs::hipHostMalloc, kPageSize); - - hipMemPool_t shared_mempool; - int* shared_ptr; - - HIP_CHECK(hipMemPoolImportFromShareableHandle( - &shared_mempool, reinterpret_cast(shareable_handle), handle_type, 0)); - - hipMemAccessFlags access_flags; - hipMemLocation location; - location.type = hipMemLocationTypeDevice; - location.id = 0; - HIP_CHECK(hipMemPoolGetAccess(&access_flags, shared_mempool, &location)); - if (access_flags != hipMemAccessFlagsProtReadWrite) { - hipMemAccessDesc desc; - memset(&desc, 0, sizeof(hipMemAccessDesc)); - desc.location.type = hipMemLocationTypeDevice; - desc.location.id = 0; - desc.flags = hipMemAccessFlagsProtReadWrite; - HIP_CHECK(hipMemPoolSetAccess(shared_mempool, &desc, 1)); - } - - HIP_CHECK( - hipMemPoolImportPointer(reinterpret_cast(&shared_ptr), shared_mempool, &export_ptr)); - - const auto element_count = kPageSize / sizeof(int); - constexpr auto thread_count = 1024; - const auto block_count = element_count / thread_count + 1; - int expected_value = 12; - VectorSet<<>>(shared_ptr, expected_value, - element_count); - HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipStreamSynchronize(stream.stream())); - - // Copy the buffer locally - HIP_CHECK(hipMemcpyAsync(host_ptr.host_ptr(), shared_ptr, kPageSize, hipMemcpyDeviceToHost, - stream.stream())); - HIP_CHECK(hipStreamSynchronize(stream.stream())); - - // Check if the content is as expected - ArrayFindIfNot(host_ptr.host_ptr(), expected_value, element_count); - - // Free the memory before the exporter process frees it - HIP_CHECK(hipFreeAsync(shared_ptr, stream.stream())); - - // And wait for all the queued up work to complete - HIP_CHECK(hipStreamSynchronize(stream.stream())); - - HIP_CHECK(hipFreeAsync(ptr, stream.stream())); - HIP_CHECK(hipStreamSynchronize(stream.stream())); - HIP_CHECK(hipMemPoolDestroy(mempool)); -} - -/** - * Test Description - * ------------------------ - * - Test to verify hipMemPoolExportToShareableHandle behavior with invalid arguments: - * -# Invalid shareable handle - * -# Invalid Memory Pool - * -# Invalid flag - * -# Invalid Memory Pool properties - * -# Invalid Memory Handle type - * - * Test source - * ------------------------ - * - /unit/memory/hipMemPoolExportImport.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.0 - */ -TEST_CASE("Unit_hipMemPoolExportToShareableHandle_Negative_Parameters") { - int mem_pool_support = 0; - HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); - if (!mem_pool_support) { - SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); - return; - } - - int device_id = 0; - HIP_CHECK(hipSetDevice(device_id)); - - hipMemPool_t mempool; - hipMemPoolProps pool_props; - memset(&pool_props, 0, sizeof(hipMemPoolProps)); - pool_props.allocType = hipMemAllocationTypePinned; - pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; - pool_props.location.type = hipMemLocationTypeDevice; - pool_props.location.id = 0; - HIP_CHECK(hipMemPoolCreate(&mempool, &pool_props)); - - SECTION("Invalid shareable handle") { - HIP_CHECK_ERROR( - hipMemPoolExportToShareableHandle(nullptr, mempool, hipMemHandleTypePosixFileDescriptor, 0), - hipErrorInvalidValue); - } - - SECTION("Invalid Memory Pool") { - int share_handle; - - HIP_CHECK_ERROR(hipMemPoolExportToShareableHandle(&share_handle, nullptr, - hipMemHandleTypePosixFileDescriptor, 0), - hipErrorInvalidValue); - } - - SECTION("Invalid flag") { - int share_handle; - - HIP_CHECK_ERROR(hipMemPoolExportToShareableHandle(&share_handle, mempool, - hipMemHandleTypePosixFileDescriptor, 1), - hipErrorInvalidValue); - } - - SECTION("Invalid Memory Pool properties") { - int share_handle; - pool_props.handleTypes = hipMemHandleTypeNone; - hipMemPool_t mempool_none; - HIP_CHECK(hipMemPoolCreate(&mempool_none, &pool_props)); - - HIP_CHECK_ERROR(hipMemPoolExportToShareableHandle(&share_handle, mempool_none, - hipMemHandleTypePosixFileDescriptor, 0), - hipErrorInvalidValue); - pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; - HIP_CHECK(hipMemPoolDestroy(mempool_none)); - } - - SECTION("Invalid Memory Handle type") { - int share_handle; - - HIP_CHECK_ERROR( - hipMemPoolExportToShareableHandle(&share_handle, mempool, hipMemHandleTypeNone, 0), - hipErrorInvalidValue); - } - - HIP_CHECK(hipMemPoolDestroy(mempool)); -} - -/** - * End doxygen group hipMemPoolExportToShareableHandle. - * @} - */ - - -/** - * @addtogroup hipMemPoolImportFromShareableHandle hipMemPoolImportFromShareableHandle - * @{ - * @ingroup StreamOTest - * `hipMemPoolImportFromShareableHandle(hipMemPool_t* mem_pool, void* shared_handle, - * hipMemAllocationHandleType handle_type,unsigned int flags)` - Imports a memory pool from a shared - * handle. - * ________________________ - * Test cases from other APIs: - * - @ref Unit_hipMemPoolExportImport_Functional - * - @ref Unit_hipMemPoolExportImport_IPC_Functional - * - @ref Unit_hipMemPoolExportImport_MultipleDevices_IPC_Functional - */ - -/** - * Test Description - * ------------------------ - * - Test to verify hipMemPoolImportFromShareableHandle behavior with invalid arguments: - * -# Invalid shareable handle - * -# Invalid Memory Pool - * -# Invalid flag - * -# Invalid Memory Handle type - * - * Test source - * ------------------------ - * - /unit/memory/hipMemPoolExportImport.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.0 - */ -TEST_CASE("Unit_hipMemPoolImportFromShareableHandle_Negative_Parameters") { - int mem_pool_support = 0; - HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); - if (!mem_pool_support) { - SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); - return; - } - - int device_id = 0; - HIP_CHECK(hipSetDevice(device_id)); - - int share_handle; - hipMemPool_t mempool; - hipMemPoolProps pool_props; - memset(&pool_props, 0, sizeof(hipMemPoolProps)); - pool_props.allocType = hipMemAllocationTypePinned; - pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; - pool_props.location.type = hipMemLocationTypeDevice; - pool_props.location.id = 0; - HIP_CHECK(hipMemPoolCreate(&mempool, &pool_props)); - HIP_CHECK(hipMemPoolExportToShareableHandle(&share_handle, mempool, - hipMemHandleTypePosixFileDescriptor, 0)); - - hipMemPool_t shared_mempool = nullptr; - - SECTION("Invalid shareable handle") { - HIP_CHECK_ERROR(hipMemPoolImportFromShareableHandle(&shared_mempool, nullptr, - hipMemHandleTypePosixFileDescriptor, 0), - hipErrorInvalidValue); - } - - SECTION("Invalid Memory Pool") { - HIP_CHECK_ERROR(hipMemPoolImportFromShareableHandle(nullptr, &share_handle, - hipMemHandleTypePosixFileDescriptor, 0), - hipErrorInvalidValue); - } - - SECTION("Invalid flag") { - HIP_CHECK_ERROR(hipMemPoolImportFromShareableHandle(&shared_mempool, &share_handle, - hipMemHandleTypePosixFileDescriptor, 1), - hipErrorInvalidValue); - } - - SECTION("Invalid Memory Handle type") { - HIP_CHECK_ERROR(hipMemPoolImportFromShareableHandle(&shared_mempool, &share_handle, - hipMemHandleTypeNone, 0), - hipErrorInvalidValue); - } -} - -/** - * End doxygen group hipMemPoolImportFromShareableHandle. - * @} - */ - - -/** - * @addtogroup hipMemPoolExportPointer hipMemPoolExportPointer - * @{ - * @ingroup StreamOTest - * `hipMemPoolExportPointer(hipMemPoolPtrExportData* export_data, void* dev_ptr)` - Export data to share a memory pool allocation between processes. - * ________________________ - * Test cases from other APIs: - * - @ref Unit_hipMemPoolExportImport_Functional - * - @ref Unit_hipMemPoolExportImport_IPC_Functional - * - @ref Unit_hipMemPoolExportImport_MultipleDevices_IPC_Functional - */ - -/** - * Test Description - * ------------------------ - * - Test to verify hipMemPoolExportPointer behavior with invalid arguments: - * -# Invalid exported data - * -# Invalid device pointer - * - * Test source - * ------------------------ - * - /unit/memory/hipMemPoolExportImport.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.0 - */ -TEST_CASE("Unit_hipMemPoolExportPointer_Negative_Parameters") { - int mem_pool_support = 0; - HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); - if (!mem_pool_support) { - SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); - return; - } - - int device_id = 0; - HIP_CHECK(hipSetDevice(device_id)); - - void* ptr; - hipMemPoolPtrExportData export_ptr; - hipMemAllocationHandleType handle_type = hipMemHandleTypePosixFileDescriptor; - - StreamGuard stream(Streams::withFlags, hipStreamNonBlocking); - int share_handle; - hipMemPool_t mempool; - hipMemPoolProps pool_props; - memset(&pool_props, 0, sizeof(hipMemPoolProps)); - pool_props.allocType = hipMemAllocationTypePinned; - pool_props.handleTypes = handle_type; - pool_props.location.type = hipMemLocationTypeDevice; - pool_props.location.id = 0; - HIP_CHECK(hipMemPoolCreate(&mempool, &pool_props)); - - HIP_CHECK(hipMallocFromPoolAsync(&ptr, kPageSize, mempool, stream.stream())); - - HIP_CHECK(hipMemPoolExportToShareableHandle(&share_handle, mempool, handle_type, 0)); - - memset(&export_ptr, 0, sizeof(hipMemPoolPtrExportData)); - - SECTION("Invalid exported data") { - HIP_CHECK_ERROR(hipMemPoolExportPointer(nullptr, reinterpret_cast(ptr)), - hipErrorInvalidValue); - } - - SECTION("Invalid device pointer") { - HIP_CHECK_ERROR(hipMemPoolExportPointer(&export_ptr, nullptr), hipErrorInvalidValue); - } - - HIP_CHECK(hipFreeAsync(ptr, stream.stream())); - HIP_CHECK(hipStreamSynchronize(stream.stream())); - HIP_CHECK(hipMemPoolDestroy(mempool)); -} - -/** - * End doxygen group hipMemPoolExportPointer. - * @} - */ - - -/** - * @addtogroup hipMemPoolImportPointer hipMemPoolImportPointer - * @{ - * @ingroup StreamOTest - * `hipMemPoolImportPointer(void** dev_ptr, hipMemPool_t mem_pool, hipMemPoolPtrExportData* export_data)` - Import a memory pool allocation from another process. - * ________________________ - * Test cases from other APIs: - * - @ref Unit_hipMemPoolExportImport_Functional - * - @ref Unit_hipMemPoolExportImport_IPC_Functional - * - @ref Unit_hipMemPoolExportImport_MultipleDevices_IPC_Functional - */ - -/** - * Test Description - * ------------------------ - * - Test to verify hipMemPoolImportPointer behavior with invalid arguments: - * -# Invalid device ptr - * -# Invalid Memory Pool - * -# Invalid exported data - * - * Test source - * ------------------------ - * - /unit/memory/hipMemPoolExportImport.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.0 - */ -TEST_CASE("Unit_hipMemPoolImportPointer_Negative_Parameters") { - int mem_pool_support = 0; - HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); - if (!mem_pool_support) { - SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); - return; - } - - int device_id = 0; - HIP_CHECK(hipSetDevice(device_id)); - - void* ptr; - hipMemPoolPtrExportData export_ptr; - hipMemAllocationHandleType handle_type = hipMemHandleTypePosixFileDescriptor; - - StreamGuard stream(Streams::withFlags, hipStreamNonBlocking); - int share_handle; - hipMemPool_t mempool; - hipMemPoolProps pool_props; - memset(&pool_props, 0, sizeof(hipMemPoolProps)); - pool_props.allocType = hipMemAllocationTypePinned; - pool_props.handleTypes = handle_type; - pool_props.location.type = hipMemLocationTypeDevice; - pool_props.location.id = 0; - HIP_CHECK(hipMemPoolCreate(&mempool, &pool_props)); - - // Allocate memory in a stream from the pool just created - HIP_CHECK(hipMallocFromPoolAsync(&ptr, kPageSize, mempool, stream.stream())); - - HIP_CHECK(hipMemPoolExportToShareableHandle(&share_handle, mempool, handle_type, 0)); - - memset((void*)&export_ptr, 0, sizeof(hipMemPoolPtrExportData)); - HIP_CHECK(hipMemPoolExportPointer(&export_ptr, reinterpret_cast(ptr))); - - hipMemPool_t shared_mempool; - int* shared_ptr = nullptr; - - HIP_CHECK(hipMemPoolImportFromShareableHandle( - &shared_mempool, reinterpret_cast(share_handle), handle_type, 0)); - - hipMemAccessFlags access_flags; - hipMemLocation location; - location.type = hipMemLocationTypeDevice; - location.id = 0; - HIP_CHECK(hipMemPoolGetAccess(&access_flags, shared_mempool, &location)); - if (access_flags != hipMemAccessFlagsProtReadWrite) { - hipMemAccessDesc desc; - memset(&desc, 0, sizeof(hipMemAccessDesc)); - desc.location.type = hipMemLocationTypeDevice; - desc.location.id = 0; - desc.flags = hipMemAccessFlagsProtReadWrite; - HIP_CHECK(hipMemPoolSetAccess(shared_mempool, &desc, 1)); - } - - SECTION("Invalid device ptr") { - HIP_CHECK_ERROR(hipMemPoolImportPointer(nullptr, shared_mempool, &export_ptr), - hipErrorInvalidValue); - } - - SECTION("Invalid Memory Pool") { - HIP_CHECK_ERROR( - hipMemPoolImportPointer(reinterpret_cast(&shared_ptr), nullptr, &export_ptr), - hipErrorInvalidValue); - } - - SECTION("Invalid exported data") { - HIP_CHECK_ERROR( - hipMemPoolImportPointer(reinterpret_cast(&shared_ptr), shared_mempool, nullptr), - hipErrorInvalidValue); - } - - HIP_CHECK(hipFreeAsync(ptr, stream.stream())); - HIP_CHECK(hipStreamSynchronize(stream.stream())); - HIP_CHECK(hipMemPoolDestroy(mempool)); -} - -/** - * End doxygen group hipMemPoolImportPointer. - * @} - */ diff --git a/catch/unit/stream_ordered/hipMemPoolExportImportIPC.cc b/catch/unit/stream_ordered/hipMemPoolExportImportIPC.cc deleted file mode 100644 index 1310b6ecf6..0000000000 --- a/catch/unit/stream_ordered/hipMemPoolExportImportIPC.cc +++ /dev/null @@ -1,422 +0,0 @@ -/* - Copyright (c) 2023 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, INNCLUDING 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 ANNY 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. - */ - -#ifdef _WIN64 -#define NOMINMAX -#endif /* _WIN64 */ - -#include "helper_multiprocess.hh" - -#include -#include -#include - -/** - * @addtogroup hipMemPoolExportToShareableHandle hipMemPoolExportToShareableHandle - * @{ - * @ingroup StreamOTest - * `hipMemPoolExportToShareableHandle(void* shared_handle, hipMemPool_t mem_pool, - * hipMemAllocationHandleType handle_type, unsigned int flags)` - Exports a memory pool to the - * requested handle type. - */ - -#ifdef __linux__ - -static const char shm_name[] = "mempool_test_shm"; -static const char ipc_name[] = "mempool_test_pipe"; - -static constexpr int kMaxDevices = 8; - -typedef struct shmStruct_st { - Process processes[kMaxDevices]; - hipMemPoolPtrExportData exportPtrData[kMaxDevices]; -} shmStruct; - -typedef struct ipcBarrier { - int count; - bool sense; - bool allExit; -} ipcBarrier_t; - -typedef struct ipcDevices { - int count; - int ordinals[kMaxDevices]; -} ipcDevices_t; - -static ipcBarrier_t* g_Barrier{}; -static bool g_procSense; -static int g_processCnt; - -/* - Get device with P2P access to device 0. -*/ -static void get_devices(ipcDevices_t* devices) { - pid_t pid = fork(); - - if (!pid) { - // HIP APIs are called in child process, - // to avoid HIP initialization in main process. - int i, device_count; - HIP_CHECK(hipGetDeviceCount(&device_count)); - - int mem_pool_support = 0; - HIP_CHECK(hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, 0)); - if (!mem_pool_support) { - devices->count = 0; - exit(EXIT_SUCCESS); - } - - // Device 0 - devices->ordinals[0] = 0; - devices->count = 1; - - if (device_count < 2) { - exit(EXIT_SUCCESS); - } - - int can_peer_access_0i, can_peer_access_i0; - for (i = 1; i < device_count; i++) { - HIP_CHECK(hipDeviceCanAccessPeer(&can_peer_access_0i, 0, i)); - HIP_CHECK(hipDeviceCanAccessPeer(&can_peer_access_i0, i, 0)); - HIP_CHECK( - hipDeviceGetAttribute(&mem_pool_support, hipDeviceAttributeMemoryPoolsSupported, i)); - - if (can_peer_access_0i * can_peer_access_i0 * mem_pool_support) { - devices->ordinals[i] = i; - INFO("Two-way peer access is available between GPU" << devices->ordinals[0] << " and GPU" - << devices->ordinals[i]); - devices->count += 1; - if (devices->count >= kMaxDevices) break; - } else { - break; - } - } - - exit(EXIT_SUCCESS); - } else { - int status; - waitpid(pid, &status, 0); - HIP_ASSERT(!status); - } -} - -/* - Calling process waits for other processes to signal/complete. -*/ -static void process_barrier() { - int newCount = __sync_add_and_fetch(&g_Barrier->count, 1); - - if (newCount == g_processCnt) { - g_Barrier->count = 0; - g_Barrier->sense = !g_procSense; - - } else { - while (g_Barrier->sense == g_procSense) { - if (!g_Barrier->allExit) { - sched_yield(); - } else { - exit(EXIT_FAILURE); - } - } - } - - g_procSense = !g_procSense; -} - -/* Child process(es) import shared memory pool and check if allocated memory can be accessed and - * used*/ -static void child_process(int id) { - volatile shmStruct* shm = NULL; - hipStream_t stream; - sharedMemoryInfo info; - void* ptr; - - LinearAllocGuard host_ptr(LinearAllocs::hipHostMalloc, kPageSize); - - ipcHandle* ipc_child_handle = NULL; - checkIpcErrors(ipcOpenSocket(ipc_child_handle)); - - // wait for parent process to create shared memory - process_barrier(); - - if (sharedMemoryOpen(shm_name, sizeof(shmStruct), &info) != 0) { - INFO("Failed to create shared memory slab\n"); - exit(EXIT_FAILURE); - } - shm = reinterpret_cast(info.addr); - shm->processes[id] = getpid(); - - // wait for parent process to send shareable handle - process_barrier(); - - // Receive allocation handle shared by parent. - std::vector sh_handle(1); - checkIpcErrors(ipcRecvShareableHandles(ipc_child_handle, sh_handle)); - - HIP_CHECK(hipSetDevice(0)); - HIP_CHECK(hipStreamCreateWithFlags(&stream, hipStreamNonBlocking)); - - hipMemPool_t pool; - - hipMemAllocationHandleType handle_type = hipMemHandleTypePosixFileDescriptor; - - // Import mem pool from all the devices created in the master - // process using shareable handles received via socket - // and import the pointer to the allocated buffer using - // exportData filled in shared memory by the master process. - HIP_CHECK(hipMemPoolImportFromShareableHandle(&pool, reinterpret_cast(sh_handle[0]), - handle_type, 0)); - - hipMemAccessFlags access_flags; - hipMemLocation location; - location.type = hipMemLocationTypeDevice; - location.id = 0; - HIP_CHECK(hipMemPoolGetAccess(&access_flags, pool, &location)); - if (access_flags != hipMemAccessFlagsProtReadWrite) { - hipMemAccessDesc desc; - memset(&desc, 0, sizeof(hipMemAccessDesc)); - desc.location.type = hipMemLocationTypeDevice; - desc.location.id = 0; - desc.flags = hipMemAccessFlagsProtReadWrite; - HIP_CHECK(hipMemPoolSetAccess(pool, &desc, 1)); - } - - // Import the allocation from memory pool using the opaque export data retrieved through - // the shared memory - HIP_CHECK(hipMemPoolImportPointer(&ptr, pool, - const_cast(&shm->exportPtrData[id]))); - - // Since we have imported allocations shared by the parent with us, we can - // close this ShareableHandle. - checkIpcErrors(ipcCloseShareableHandle(sh_handle[0])); - - // Since we have imported allocations shared by the parent with us, we can - // close the socket - checkIpcErrors(ipcCloseSocket(ipc_child_handle)); - - // Child processed accesses imported buffer - const auto element_count = kPageSize / sizeof(int); - constexpr auto thread_count = 1024; - const auto block_count = element_count / thread_count + 1; - int expected_value = 12 + id; - VectorSet<<>>((int*)ptr, expected_value, element_count); - HIP_CHECK(hipGetLastError()); - HIP_CHECK(hipStreamSynchronize(stream)); - - // Copy the buffer locally - HIP_CHECK(hipMemcpyAsync(host_ptr.host_ptr(), ptr, kPageSize, hipMemcpyDeviceToHost, stream)); - HIP_CHECK(hipStreamSynchronize(stream)); - - INFO("Process " << id << "verifying...\n"); - - // Check if the content is as expected - ArrayFindIfNot(host_ptr.host_ptr(), expected_value, element_count); - - // Free the memory before the exporter process frees it - HIP_CHECK(hipFreeAsync(ptr, stream)); - - // And wait for all the queued up work to complete - HIP_CHECK(hipStreamSynchronize(stream)); - HIP_CHECK(hipStreamDestroy(stream)); -} - -static void parent_process(int dev_count) { - sharedMemoryInfo info; - int i; - volatile shmStruct* shm = NULL; - std::vector ptrs; - std::vector child_processes; - - if (sharedMemoryCreate(shm_name, sizeof(*shm), &info) != 0) { - INFO("Failed to create shared memory slab\n"); - exit(EXIT_FAILURE); - } - shm = (volatile shmStruct*)info.addr; - memset((void*)shm, 0, sizeof(*shm)); - - // wait for child processes to insert their pids into shared memory - process_barrier(); - - std::vector shareable_handles(dev_count); - std::vector streams(dev_count); - std::vector pools(dev_count); - - // Now allocate memory for each process and fill the shared - // memory buffer with the export data and get mempool handles to communicate - for (i = 0; i < dev_count; i++) { - void* ptr; - HIP_CHECK(hipSetDevice(i)); - HIP_CHECK(hipStreamCreateWithFlags(&streams[i], hipStreamNonBlocking)); - // Allocate an explicit pool with IPC capabilities - hipMemPoolProps pool_props; - memset(&pool_props, 0, sizeof(hipMemPoolProps)); - pool_props.allocType = hipMemAllocationTypePinned; - pool_props.handleTypes = hipMemHandleTypePosixFileDescriptor; - - pool_props.location.type = hipMemLocationTypeDevice; - pool_props.location.id = i; - - HIP_CHECK(hipMemPoolCreate(&pools[i], &pool_props)); - - // Query the shareable handle for the pool - hipMemAllocationHandleType handle_type = hipMemHandleTypePosixFileDescriptor; - // Allocate memory in a stream from the pool just created - HIP_CHECK(hipMallocFromPoolAsync(&ptr, kPageSize, pools[i], streams[i])); - - HIP_CHECK(hipMemPoolExportToShareableHandle(&shareable_handles[i], pools[i], handle_type, 0)); - - // Memset handle to 0 to make sure call to hipMemPoolImportPointer in - // child process will fail if the following call to hipMemPoolExportPointer fails. - memset((void*)&shm->exportPtrData[i], 0, sizeof(hipMemPoolPtrExportData)); - HIP_CHECK( - hipMemPoolExportPointer(const_cast(&shm->exportPtrData[i]), ptr)); - ptrs.push_back(ptr); - child_processes.push_back(static_cast(shm->processes[i])); - } - - ipcHandle* ipc_parent_handle; - checkIpcErrors(ipcCreateSocket(ipc_parent_handle, ipc_name, child_processes)); - - for (i = 0; i < dev_count; i++) { - std::vector current_handle(1, shareable_handles[i]); - std::vector current_process(1, child_processes[i]); - checkIpcErrors(ipcSendShareableHandles(ipc_parent_handle, current_handle, current_process)); - } - - // Close the shareable handles as they are not needed anymore. - for (int i = 0; i < dev_count; i++) { - checkIpcErrors(ipcCloseShareableHandle(shareable_handles[i])); - } - - checkIpcErrors(ipcCloseSocket(ipc_parent_handle)); - - process_barrier(); - - // And wait for them to finish - for (i = 0; i < child_processes.size(); i++) { - if (waitProcess(&child_processes[i]) != EXIT_SUCCESS) { - INFO("Process " << i << " failed!\n"); - exit(EXIT_FAILURE); - } - } - - // Clean up! - for (i = 0; i < dev_count; i++) { - HIP_CHECK(hipSetDevice(i)); - HIP_CHECK(hipFreeAsync(ptrs[i], streams[i])); - HIP_CHECK(hipStreamSynchronize(streams[i])); - HIP_CHECK(hipMemPoolDestroy(pools[i])); - HIP_CHECK(hipStreamDestroy(streams[i])); - } - - sharedMemoryClose(&info); -} - -/** - * Test Description - * ------------------------ - * - Test to verify exporting/importing a shareable handle on a single device between parent and - * child process using IPC mechanisms - shared memory and sockets. - * Test source - * ------------------------ - * - /unit/memory/hipMemPoolExportImportIPC.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.0 - */ -TEST_CASE("Unit_hipMemPoolExportImport_IPC_Functional") { - ipcDevices_t* shm_devices; - shm_devices = reinterpret_cast( - mmap(NULL, sizeof(*shm_devices), PROT_READ | PROT_WRITE, MAP_SHARED | MAP_ANONYMOUS, 0, 0)); - REQUIRE(MAP_FAILED != shm_devices); - // Barrier is used to synchronize created processes - g_Barrier = reinterpret_cast( - mmap(NULL, sizeof(*g_Barrier), PROT_READ | PROT_WRITE, MAP_SHARED | MAP_ANONYMOUS, 0, 0)); - memset(g_Barrier, 0, sizeof(*g_Barrier)); - - // set local barrier sense flag - g_procSense = 0; - - get_devices(shm_devices); - if (!shm_devices->count) { - SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); - return; - } - // Set device count to 1 - shm_devices->count = 1; - g_processCnt = shm_devices->count + 1; - int index = 0; - - Process process = fork(); - if (process != 0) { - parent_process(shm_devices->count); - } else { - child_process(index); - } -} - -/** - * Test Description - * ------------------------ - * - Test to verify exporting/importing a shareable handle on multiple devices between parent and - * child processes using IPC mechanisms - shared memory and sockets. - * Test source - * ------------------------ - * - /unit/memory/hipMemPoolExportImportIPC.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.0 - */ -TEST_CASE("Unit_hipMemPoolExportImport_MultipleDevices_IPC_Functional") { - ipcDevices_t* shm_devices; - shm_devices = reinterpret_cast( - mmap(NULL, sizeof(*shm_devices), PROT_READ | PROT_WRITE, MAP_SHARED | MAP_ANONYMOUS, 0, 0)); - REQUIRE(MAP_FAILED != shm_devices); - // Barrier is used to synchronize processes created. - g_Barrier = reinterpret_cast( - mmap(NULL, sizeof(*g_Barrier), PROT_READ | PROT_WRITE, MAP_SHARED | MAP_ANONYMOUS, 0, 0)); - memset(g_Barrier, 0, sizeof(*g_Barrier)); - - // set local barrier sense flag - g_procSense = 0; - - get_devices(shm_devices); - if (!shm_devices->count) { - SUCCEED("Runtime doesn't support Memory Pool. Skip the test case."); - return; - } - g_processCnt = shm_devices->count + 1; - - int index = 0; - - for (int i = 1; i < g_processCnt; i++) { - Process process = fork(); - if (!process) { - index = i; - break; - } - } - - if (index == 0) { - parent_process(shm_devices->count); - } else { - child_process(index - 1); - } -} -#endif