From c26a76e4e9e1e48c286aa0580745da7760b14fc4 Mon Sep 17 00:00:00 2001 From: Nives Vukovic Date: Fri, 2 Feb 2024 12:16:12 +0530 Subject: [PATCH] EXSWHTEC-57 - Implement tests for Import/Export hipMemPool APIs #445 Change-Id: Ia40ead5612819d1e87a4c526adf77be20ebc4ff2 [ROCm/hip-tests commit: d1500f261269e40ae1eceec1b5884558aa570efc] --- projects/hip-tests/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 | 491 +++++++++++++++++ .../hipMemPoolExportImportIPC.cc | 419 ++++++++++++++ 6 files changed, 1578 insertions(+) create mode 100644 projects/hip-tests/catch/unit/stream_ordered/CMakeLists.txt create mode 100644 projects/hip-tests/catch/unit/stream_ordered/helper_multiprocess.cc create mode 100644 projects/hip-tests/catch/unit/stream_ordered/helper_multiprocess.hh create mode 100644 projects/hip-tests/catch/unit/stream_ordered/hipMemPoolExportImport.cc create mode 100644 projects/hip-tests/catch/unit/stream_ordered/hipMemPoolExportImportIPC.cc diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index 37f8b73cc5..59f8f6ad47 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -22,6 +22,7 @@ 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/projects/hip-tests/catch/unit/stream_ordered/CMakeLists.txt b/projects/hip-tests/catch/unit/stream_ordered/CMakeLists.txt new file mode 100644 index 0000000000..e9a5e56a8d --- /dev/null +++ b/projects/hip-tests/catch/unit/stream_ordered/CMakeLists.txt @@ -0,0 +1,30 @@ +# 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/projects/hip-tests/catch/unit/stream_ordered/helper_multiprocess.cc b/projects/hip-tests/catch/unit/stream_ordered/helper_multiprocess.cc new file mode 100644 index 0000000000..4ab32e64e7 --- /dev/null +++ b/projects/hip-tests/catch/unit/stream_ordered/helper_multiprocess.cc @@ -0,0 +1,518 @@ +/* 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 = {0}; + BOOL status; + size_t arglen = 0; + size_t argIdx = 0; + 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 = {0}; + 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 (int i = 0; i < shareableHandles.size(); i++) { + for (int j = 0; j < processes.size(); j++) { + checkIpcErrors(ipcSendShareableHandle(handle, shareableHandles, processes[j], i)); + } + } + return 0; +} + +int ipcRecvShareableHandles(ipcHandle* handle, std::vector& shareableHandles) { + for (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* name, const std::vector& processes) { + handle = new ipcHandle; + handle->hMailslot.resize(processes.size()); + + // Open Mailslots of all clients and store respective handles. + for (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 %d\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 %d\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 %d.\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 %d.\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 (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 (%d)\n", GetLastError()); + return -1; + } + + for (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 (%d)\n", GetLastError()); + return -1; + } + checkIpcErrors(ipcSendData(handle->hMailslot[i], &hDup, sizeof(hDup))); + } + CloseHandle(hProcess); + } + return 0; +} + +int ipcRecvShareableHandles(ipcHandle* handle, std::vector& shareableHandles) { + for (int i = 0; i < shareableHandles.size(); i++) { + checkIpcErrors(ipcRecvData(handle, &shareableHandles[i], sizeof(shareableHandles[i]))); + } + return 0; +} + +int ipcCloseSocket(ipcHandle* handle) { + for (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/projects/hip-tests/catch/unit/stream_ordered/helper_multiprocess.hh b/projects/hip-tests/catch/unit/stream_ordered/helper_multiprocess.hh new file mode 100644 index 0000000000..b9bb4e16ea --- /dev/null +++ b/projects/hip-tests/catch/unit/stream_ordered/helper_multiprocess.hh @@ -0,0 +1,119 @@ +/* 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/projects/hip-tests/catch/unit/stream_ordered/hipMemPoolExportImport.cc b/projects/hip-tests/catch/unit/stream_ordered/hipMemPoolExportImport.cc new file mode 100644 index 0000000000..58ada79b7c --- /dev/null +++ b/projects/hip-tests/catch/unit/stream_ordered/hipMemPoolExportImport.cc @@ -0,0 +1,491 @@ +/* + 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. + */ +#include +#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/projects/hip-tests/catch/unit/stream_ordered/hipMemPoolExportImportIPC.cc b/projects/hip-tests/catch/unit/stream_ordered/hipMemPoolExportImportIPC.cc new file mode 100644 index 0000000000..9635f93b63 --- /dev/null +++ b/projects/hip-tests/catch/unit/stream_ordered/hipMemPoolExportImportIPC.cc @@ -0,0 +1,419 @@ +/* + 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. + */ + +#include "helper_multiprocess.hh" + +#include +#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, cudaStreamNonBlocking)); + + 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