AI/rocm-systems
2
0
انشعاب 0
کد مسائل تقاضاهای واکشی Actions Packages پروژه‌ها انتشارها دانشنامه فعالیت

EXSWHTEC-57 - Implement tests for Import/Export hipMemPool APIs #445

فهرست منبع 
Change-Id: Ia40ead5612819d1e87a4c526adf77be20ebc4ff2


[ROCm/hip-tests commit: d1500f2612]
This commit is contained in:
Nives Vukovic
2024-02-02 12:16:12 +05:30
کامیت شده توسط Rakesh Roy
والد e97080ec7b
کامیت c26a76e4e9
6فایلهای تغییر یافته به همراه1578 افزوده شده و 0 حذف شده
دانلود پرونده وصله دانلود فایل تغییرات diff Expand all files Collapse all files
projects/hip-tests/catch/unit/CMakeLists.txt
+1
مشاهده پرونده
@@ -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)
projects/hip-tests/catch/unit/stream_ordered/CMakeLists.txt
+30
مشاهده پرونده
@@ -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})
projects/hip-tests/catch/unit/stream_ordered/helper_multiprocess.cc
+518
مشاهده پرونده
@@ -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 <cstdlib>
#include <string>
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<Process>& 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<ShareableHandle>& 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<ShareableHandle>& shareableHandles,
const std::vector<Process>& 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<ShareableHandle>& 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<Process>& 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<TCHAR> 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<TCHAR> 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<ShareableHandle>& shareableHandles,
const std::vector<Process>& 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<ShareableHandle>& 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
projects/hip-tests/catch/unit/stream_ordered/helper_multiprocess.hh
+119
مشاهده پرونده
@@ -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 <windows.h>
#include <iostream>
#include <stdio.h>
#include <tchar.h>
#include <strsafe.h>
#include <sddl.h>
#include <aclapi.h>
#include <winternl.h>
#else
#include <stdio.h>
#include <fcntl.h>
#include <sys/mman.h>
#include <unistd.h>
#include <errno.h>
#include <sys/wait.h>
#include <sys/types.h>
#include <sys/socket.h>
#include <memory.h>
#include <sys/un.h>
#endif
#include <vector>
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<HANDLE>
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<Process>& processes);
int ipcOpenSocket(ipcHandle*& handle);
int ipcCloseSocket(ipcHandle* handle);
int ipcRecvShareableHandles(ipcHandle* handle, std::vector<ShareableHandle>& shareableHandles);
int ipcSendShareableHandles(ipcHandle* handle, const std::vector<ShareableHandle>& shareableHandles,
const std::vector<Process>& processes);
int ipcCloseShareableHandle(ShareableHandle shHandle);
#endif // HELPER_MULTIPROCESS_H
projects/hip-tests/catch/unit/stream_ordered/hipMemPoolExportImport.cc
+491
مشاهده پرونده
@@ -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 <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <utils.hh>
#include <resource_guards.hh>
/**
* @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<void*>(ptr)));
LinearAllocGuard<int> host_ptr(LinearAllocs::hipHostMalloc, kPageSize);
hipMemPool_t shared_mempool;
int* shared_ptr;
HIP_CHECK(hipMemPoolImportFromShareableHandle(
&shared_mempool, reinterpret_cast<void*>(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<void**>(&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<<<block_count, thread_count, 0, stream.stream()>>>(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<void*>(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<void*>(ptr)));
hipMemPool_t shared_mempool;
int* shared_ptr = nullptr;
HIP_CHECK(hipMemPoolImportFromShareableHandle(
&shared_mempool, reinterpret_cast<void*>(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<void**>(&shared_ptr), nullptr, &export_ptr),
hipErrorInvalidValue);
}
SECTION("Invalid exported data") {
HIP_CHECK_ERROR(
hipMemPoolImportPointer(reinterpret_cast<void**>(&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.
* @}
*/
projects/hip-tests/catch/unit/stream_ordered/hipMemPoolExportImportIPC.cc
+419
مشاهده پرونده
@@ -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 <hip_test_common.hh>
#include <hip_test_checkers.hh>
#include <utils.hh>
#include <resource_guards.hh>
/**
* @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<int> 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<volatile shmStruct*>(info.addr);
shm->processes[id] = getpid();
// wait for parent process to send shareable handle
process_barrier();
// Receive allocation handle shared by parent.
std::vector<ShareableHandle> 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<void*>(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<hipMemPoolPtrExportData*>(&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<<<block_count, thread_count, 0, stream>>>((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<void*> ptrs;
std::vector<Process> 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<ShareableHandle> shareable_handles(dev_count);
std::vector<hipStream_t> streams(dev_count);
std::vector<hipMemPool_t> 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<hipMemPoolPtrExportData*>(&shm->exportPtrData[i]), ptr));
ptrs.push_back(ptr);
child_processes.push_back(static_cast<Process>(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<ShareableHandle> current_handle(1, shareable_handles[i]);
std::vector<ShareableHandle> 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<ipcDevices_t*>(
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<ipcBarrier_t*>(
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<ipcDevices_t*>(
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<ipcBarrier_t*>(
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