AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-468182 - [catch2][dtest] Adding scenarios for hipMemImportFromShareableHandle and hipMemExportToShareableHandle

Browse Source 
Change-Id: I46e2bb60acc354c2db02a150848d9be4386bff30
This commit is contained in:
Rupam Chetia
2024-07-16 16:10:23 +05:30
committed by Rakesh Roy
parent 0a22d14775
commit 2eb6754de5
4 changed files with 665 additions and 105 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/unit/virtualMemoryManagement/CMakeLists.txt
+3 -2
View File
@@ -26,14 +26,15 @@ if(HIP_PLATFORM MATCHES "nvidia")
set(TEST_SRC
${TEST_SRC}
hipMemMapArrayAsync.cc)
endif()
if(UNIX) # Disabled on AMD due to defect EXSWHTEC-375
if(UNIX)
if(HIP_PLATFORM MATCHES "amd")
set(TEST_SRC
${TEST_SRC}
hipMemExportToShareableHandle.cc
hipMemImportFromShareableHandle.cc)
endif()
endif()
if(HIP_PLATFORM MATCHES "amd")
catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc
+4 -12
View File
@@ -32,7 +32,6 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
/**
@@ -55,7 +54,7 @@ TEST_CASE("Unit_hipMemExportToShareableHandle_Positive_Basic") {
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
prop.requestedHandleType = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
@@ -94,7 +93,7 @@ TEST_CASE("Unit_hipMemExportToShareableHandle_Negative_Parameters") {
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
prop.requestedHandleType = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
@@ -106,12 +105,13 @@ TEST_CASE("Unit_hipMemExportToShareableHandle_Negative_Parameters") {
HIP_CHECK(hipMemCreate(&handle, granularity * 2, &prop, 0));
void* shareable_handle = nullptr;
#if HT_NVIDIA
SECTION("shareableHandle == nullptr") {
HIP_CHECK_ERROR(
hipMemExportToShareableHandle(nullptr, handle, hipMemHandleTypePosixFileDescriptor, 0),
hipErrorInvalidValue);
}
#endif
#if HT_AMD
SECTION("handle == nullptr") {
@@ -134,14 +134,6 @@ TEST_CASE("Unit_hipMemExportToShareableHandle_Negative_Parameters") {
}
HIP_CHECK(hipMemRelease(handle));
#if HT_AMD // segfaults on NVIDIA
SECTION("released handle") {
HIP_CHECK_ERROR(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 0),
hipErrorInvalidValue);
}
#endif
}
/**
catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc
+425 -88
View File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2023-2024 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
@@ -30,14 +30,22 @@ THE SOFTWARE.
* Imports an allocation from a requested shareable handle type.
*/
#include <unistd.h>
#include <sys/types.h>
#include <sys/wait.h>
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
#define DATA_SIZE (1 << 13)
#define THREADS_PER_BLOCK 512
typedef int ShareableHandle;
/**
Kernel to perform Square of input data.
*/
static __global__ void square_kernel(int* Buff) {
int i = threadIdx.x + blockDim.x * blockIdx.x;
int temp = Buff[i] * Buff[i];
Buff[i] = temp;
}
/**
* Test Description
* ------------------------
@@ -58,7 +66,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_Positive_Basic") {
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
prop.requestedHandleType = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
@@ -68,87 +76,15 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_Positive_Basic") {
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, granularity * 2, &prop, 0));
void* shareable_handle = nullptr;
ShareableHandle shareable_handle;
HIP_CHECK(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 0));
hipMemGenericAllocationHandle_t imported_handle;
HIP_CHECK(hipMemImportFromShareableHandle(&imported_handle, shareable_handle,
HIP_CHECK(hipMemImportFromShareableHandle(&imported_handle, &shareable_handle,
hipMemHandleTypePosixFileDescriptor));
HIP_CHECK(hipMemRelease(handle));
}
/**
* Test Description
* ------------------------
* - Basic multiprocess sanity test.
* ------------------------
* - unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc
* Test requirements
* ------------------------
* - Host specific (LINUX)
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemImportFromShareableHandle_Positive_MultiProc") {
int fd[2];
REQUIRE(pipe(fd) == 0);
auto pid = fork();
REQUIRE(pid >= 0);
if (pid == 0) { // child
REQUIRE(close(fd[1]) == 0);
void* shareable_handle = nullptr;
REQUIRE(read(fd[0], &shareable_handle, sizeof(shareable_handle)) >= 0);
REQUIRE(close(fd[0]) == 0);
REQUIRE(shareable_handle != nullptr);
HIP_CHECK(hipFree(0));
hipMemGenericAllocationHandle_t imported_handle;
HIP_CHECK(hipMemImportFromShareableHandle(&imported_handle, shareable_handle,
hipMemHandleTypePosixFileDescriptor));
exit(0);
} else { // parent
REQUIRE(close(fd[0]) == 0);
HIP_CHECK(hipFree(0));
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device);
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
size_t granularity;
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, granularity * 2, &prop, 0));
void* shareable_handle = nullptr;
HIP_CHECK(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 0));
REQUIRE(write(fd[1], &shareable_handle, sizeof(shareable_handle)) >= 0);
REQUIRE(close(fd[1]) == 0);
REQUIRE(wait(NULL) >= 0);
HIP_CHECK(hipMemRelease(handle));
}
}
/**
* Test Description
* ------------------------
@@ -169,7 +105,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_Negative_Parameters") {
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
prop.requestedHandleType = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
@@ -200,15 +136,416 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_Negative_Parameters") {
hipErrorInvalidValue);
}
SECTION("invalid handleType") {
HIP_CHECK_ERROR(
hipMemImportFromShareableHandle(&imported_handle, shareable_handle, hipMemHandleTypeWin32),
hipErrorNotSupported);
}
HIP_CHECK(hipMemRelease(handle));
}
/**
* Test Description
* ------------------------
* - Multiprocess functionality test. Create Vmm handle in Parent
* Process and export it to Child Process using Sockets. The Child
* Process imports this handle via sockets and uses this handle
* to perform VMM operations.
* ------------------------
* - unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc
* Test requirements
* ------------------------
* - Host specific (LINUX)
* - HIP_VERSION >= 6.2
*/
TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ChldUseHdl") {
constexpr int N = DATA_SIZE;
size_t buffer_size = N * sizeof(int);
int fd[2], fdSig[2];
REQUIRE(pipe(fd) == 0);
REQUIRE(pipe(fdSig) == 0);
auto pid = fork();
REQUIRE(pid >= 0);
if (pid == 0) { // child
REQUIRE(close(fd[1]) == 0);
REQUIRE(close(fdSig[0]) == 0);
// Wait for parent process to create the socket.
size_t size_mem = 0;
REQUIRE(read(fd[0], &size_mem, sizeof(size_t)) >= 0);
// Open Socket as client
ipcSocketCom sockObj(false);
hipShareableHdl shHandle;
// Signal Parent process that Child is ready to receive msg
int sig = 0;
REQUIRE(write(fdSig[1], &sig, sizeof(int)) >= 0);
// receive message from parent provess
checkSysCallErrors(sockObj.recvShareableHdl(&shHandle));
hipMemGenericAllocationHandle_t imported_handle;
// import the sareable handle
HIP_CHECK(hipMemImportFromShareableHandle(&imported_handle, &shHandle,
hipMemHandleTypePosixFileDescriptor));
// Allocate virtual address range
hipDeviceptr_t ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, imported_handle, 0));
// Set access
hipMemAccessDesc accessDesc = {};
accessDesc.location.type = hipMemLocationTypeDevice;
accessDesc.location.id = 0;
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
// Make the address accessible to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
std::vector<int> A_h(N), B_h(N), C_h(N);
// Initialize with data
for (size_t idx = 0; idx < N; idx++) {
A_h[idx] = idx;
C_h[idx] = idx * idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
// Invoke kernel
hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0,
static_cast<int*>(ptrA));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipDeviceSynchronize());
// validate
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
// free resources
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
checkSysCallErrors(sockObj.closeThisSock());
REQUIRE(close(fd[0]) == 0);
REQUIRE(close(fdSig[1]) == 0);
exit(0);
} else { // parent
REQUIRE(close(fd[0]) == 0);
REQUIRE(close(fdSig[1]) == 0);
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device);
// Set property
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleType = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
// Set Granularity of the VMM memory
size_t granularity;
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
hipShareableHdl shareable_handle;
HIP_CHECK(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 0));
// Create the socket for communication as Server
ipcSocketCom sockObj(true);
// Signal child process that socket is ready
REQUIRE(write(fd[1], &size_mem, sizeof(size_t)) >= 0);
// Wait for the child process to receive msg
int sig = 0;
REQUIRE(read(fdSig[0], &sig, sizeof(int)) >= 0);
checkSysCallErrors(
sockObj.sendShareableHdl(shareable_handle, pid));
// Wait for child process to exit.
int status;
REQUIRE(wait(&status) >= 0);
REQUIRE(status == 0);
// Free all resources
checkSysCallErrors(sockObj.closeThisSock());
HIP_CHECK(hipMemRelease(handle));
REQUIRE(close(fd[1]) == 0);
REQUIRE(close(fdSig[0]) == 0);
}
}
/**
* Test Description
* ------------------------
* - Multiprocess functionality test. Create Vmm handle in Parent
* Process and export it to Child Process using Sockets. The Child
* Process imports this handle via sockets. Both Parent and Child Process
* uses this handle to perform VMM operations.
* ------------------------
* - unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc
* Test requirements
* ------------------------
* - Host specific (LINUX)
* - HIP_VERSION >= 6.2
*/
TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ParntChldUseHdl") {
constexpr int N = DATA_SIZE;
size_t buffer_size = N * sizeof(int);
int fd[2], fdSig[2];
REQUIRE(pipe(fd) == 0);
REQUIRE(pipe(fdSig) == 0);
// Create data buffer
std::vector<int> A_h(N), B_h(N), C_h(N);
// Initialize with data
for (size_t idx = 0; idx < N; idx++) {
A_h[idx] = idx;
C_h[idx] = idx * idx;
}
auto pid = fork();
REQUIRE(pid >= 0);
if (pid == 0) { // child
REQUIRE(close(fd[1]) == 0);
REQUIRE(close(fdSig[0]) == 0);
// Wait for parent process to create the socket.
size_t size_mem = 0;
REQUIRE(read(fd[0], &size_mem, sizeof(size_t)) >= 0);
// Open Socket as client
ipcSocketCom sockObj(false);
hipShareableHdl shHandle;
// Signal Parent process that Child is ready to receive msg
int sig = 0;
REQUIRE(write(fdSig[1], &sig, sizeof(int)) >= 0);
// receive message from parent provess
checkSysCallErrors(sockObj.recvShareableHdl(&shHandle));
hipMemGenericAllocationHandle_t imported_handle;
// import the sareable handle
HIP_CHECK(hipMemImportFromShareableHandle(&imported_handle, &shHandle,
hipMemHandleTypePosixFileDescriptor));
// Allocate virtual address range
hipDeviceptr_t ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, imported_handle, 0));
// Set access
hipMemAccessDesc accessDesc = {};
accessDesc.location.type = hipMemLocationTypeDevice;
accessDesc.location.id = 0;
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
// Make the address accessible to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
// Invoke kernel
hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0,
static_cast<int*>(ptrA));
HIP_CHECK(hipDeviceSynchronize());
// free resources
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
checkSysCallErrors(sockObj.closeThisSock());
REQUIRE(close(fd[0]) == 0);
REQUIRE(close(fdSig[1]) == 0);
exit(0);
} else { // parent
REQUIRE(close(fd[0]) == 0);
REQUIRE(close(fdSig[1]) == 0);
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device);
// Set property
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleType = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
// Set Granularity of the VMM memory
size_t granularity;
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
hipShareableHdl shareable_handle;
HIP_CHECK(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 0));
// Allocate virtual address range
hipDeviceptr_t ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
// Set access
hipMemAccessDesc accessDesc = {};
accessDesc.location.type = hipMemLocationTypeDevice;
accessDesc.location.id = device;
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
// Make the address accessible to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
// Create the socket for communication as Server
ipcSocketCom sockObj(true);
// Signal child process that socket is ready
REQUIRE(write(fd[1], &size_mem, sizeof(size_t)) >= 0);
// Wait for the child process to receive msg
int sig = 0;
REQUIRE(read(fdSig[0], &sig, sizeof(int)) >= 0);
checkSysCallErrors(sockObj.sendShareableHdl(shareable_handle, pid));
// Wait for child process to exit.
int status;
REQUIRE(wait(&status) >= 0);
REQUIRE(status == 0);
// Check results of Vmm data processing in child
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
// validate
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
// Free all resources
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
HIP_CHECK(hipMemRelease(handle));
checkSysCallErrors(sockObj.closeThisSock());
REQUIRE(close(fd[1]) == 0);
REQUIRE(close(fdSig[0]) == 0);
}
}
/**
* Test Description
* ------------------------
* - Multiprocess functionality test. Create Vmm handle in Parent
* Process and export it to Grand Child Process using Sockets. The Grand
* Child Process imports this handle via sockets. The Grand Child Process
* uses this handle to perform VMM operations.
* ------------------------
* - unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc
* Test requirements
* ------------------------
* - Host specific (LINUX)
* - HIP_VERSION >= 6.2
*/
TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_GrndChldUseHdl") {
constexpr int N = DATA_SIZE;
size_t buffer_size = N * sizeof(int);
int fd[2], fdSig[2], fdpid[2];
REQUIRE(pipe(fd) == 0);
REQUIRE(pipe(fdSig) == 0);
REQUIRE(pipe(fdpid) == 0);
auto pid = fork();
REQUIRE(pid >= 0);
if (pid == 0) { // child
auto pid2 = fork();
if (pid2 == 0) { // grandchild
REQUIRE(close(fd[1]) == 0);
REQUIRE(close(fdSig[0]) == 0);
// Wait for parent process to create the socket.
size_t size_mem = 0;
REQUIRE(read(fd[0], &size_mem, sizeof(size_t)) >= 0);
// Open Socket as client
ipcSocketCom sockObj(false);
hipShareableHdl shHandle;
// Signal Parent process that Child is ready to receive msg
int sig = 0;
REQUIRE(write(fdSig[1], &sig, sizeof(int)) >= 0);
// receive message from parent provess
checkSysCallErrors(sockObj.recvShareableHdl(&shHandle));
hipMemGenericAllocationHandle_t imported_handle;
// import the sareable handle
HIP_CHECK(hipMemImportFromShareableHandle(&imported_handle, &shHandle,
hipMemHandleTypePosixFileDescriptor));
// Allocate virtual address range
hipDeviceptr_t ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, imported_handle, 0));
// Set access
hipMemAccessDesc accessDesc = {};
accessDesc.location.type = hipMemLocationTypeDevice;
accessDesc.location.id = 0;
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
// Make the address accessible to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
std::vector<int> A_h(N), B_h(N), C_h(N);
// Initialize with data
for (size_t idx = 0; idx < N; idx++) {
A_h[idx] = idx;
C_h[idx] = idx * idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
// Invoke kernel
hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0,
static_cast<int*>(ptrA));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipDeviceSynchronize());
// validate
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
// free resources
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
checkSysCallErrors(sockObj.closeThisSock());
REQUIRE(close(fd[0]) == 0);
REQUIRE(close(fdSig[1]) == 0);
exit(0);
} else {
int status;
REQUIRE(close(fdpid[0]) == 0);
REQUIRE(write(fdpid[1], &pid2, sizeof(pid2)) >= 0);
REQUIRE(wait(&status) >= 0);
REQUIRE(status == 0);
REQUIRE(close(fdpid[1]) == 0);
exit(0);
}
} else { // parent
REQUIRE(close(fd[0]) == 0);
REQUIRE(close(fdSig[1]) == 0);
REQUIRE(close(fdpid[1]) == 0);
int pid_grChld = 0;
REQUIRE(read(fdpid[0], &pid_grChld, sizeof(pid_grChld)) >= 0);
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device);
// Set property
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleType = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
// Set Granularity of the VMM memory
size_t granularity;
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
hipShareableHdl shareable_handle;
HIP_CHECK(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 0));
// Create the socket for communication as Server
ipcSocketCom sockObj(true);
// Signal child process that socket is ready
REQUIRE(write(fd[1], &size_mem, sizeof(size_t)) >= 0);
// Wait for the child process to receive msg
int sig = 0;
REQUIRE(read(fdSig[0], &sig, sizeof(int)) >= 0);
checkSysCallErrors(sockObj.sendShareableHdl(shareable_handle, pid_grChld));
// Wait for child process to exit.
int status;
REQUIRE(wait(&status) >= 0);
REQUIRE(status == 0);
// Free all resources
HIP_CHECK(hipMemRelease(handle));
checkSysCallErrors(sockObj.closeThisSock());
REQUIRE(close(fd[1]) == 0);
REQUIRE(close(fdSig[0]) == 0);
REQUIRE(close(fdpid[0]) == 0);
}
}
/**
* End doxygen group VirtualMemoryManagementTest.
* @}
catch/unit/virtualMemoryManagement/hip_vmm_common.hh
+233 -3
View File
@@ -1,5 +1,5 @@
/*
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
Copyright (c) 2023 - 2024 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
@@ -21,7 +21,18 @@ THE SOFTWARE.
*/
#pragma once
#ifdef __linux__
#include <unistd.h>
#include <sys/types.h>
#include <sys/wait.h>
#include <fcntl.h>
#include <sys/mman.h>
#include <unistd.h>
#include <errno.h>
#include <sys/socket.h>
#include <memory.h>
#include <sys/un.h>
#endif
#include "hip_test_context.hh"
#define checkVMMSupported(device) \
@@ -46,4 +57,223 @@ THE SOFTWARE.
} \
}
constexpr int threadsPerBlk = 64;
#ifdef __linux__
#define checkSysCallErrors(result) \
if (result == -1) { \
fprintf(stderr, "Failure at %u %s\n", __LINE__, __FILE__); exit(EXIT_FAILURE); \
}
typedef pid_t Process;
typedef int hipShareableHdl;
struct ipcHdl {
int socket;
char *name;
};
class ipcSocketCom {
ipcHdl *handle;
// method to create socket from server
int createSocket() {
int server_fd;
struct sockaddr_un servaddr;
char name[16];
// Create a unique socket name based on current pid
sprintf(name, "%u", getpid());
// Create the socket handle
handle = new ipcHdl;
if (nullptr == handle) {
perror("Socket failure: Handle memory allocation failed");
return -1;
}
memset(handle, 0, sizeof(*handle));
handle->socket = -1;
handle->name = NULL;
// Creating socket
if ((server_fd = socket(AF_UNIX, SOCK_DGRAM, 0)) == 0) {
perror("Socket failure: Socket creation failed");
return -1;
}
unlink(name);
bzero(&servaddr, sizeof(servaddr));
servaddr.sun_family = AF_UNIX;
size_t len = strlen(name);
if (len > (sizeof(servaddr.sun_path) - 1)) {
perror("Socket failure: Cannot bind provided name to socket. Name too large");
return -1;
}
strncpy(servaddr.sun_path, name, len);
if (bind(server_fd, (struct sockaddr *)&servaddr, SUN_LEN(&servaddr)) < 0) {
perror("Socket failure: Binding socket failed");
return -1;
}
handle->name = new char[strlen(name) + 1];
strcpy(handle->name, name);
handle->socket = server_fd;
return 0;
}
// method to create socket from client
int openSocket() {
int sock = 0;
struct sockaddr_un cliaddr;
handle = new ipcHdl;
if (nullptr == handle) {
perror("Socket failure: Handle memory allocation failed");
return -1;
}
memset(handle, 0, sizeof(*handle));
if ((sock = socket(AF_UNIX, SOCK_DGRAM, 0)) < 0) {
perror("IPC failure:Socket creation error");
return -1;
}
bzero(&cliaddr, sizeof(cliaddr));
cliaddr.sun_family = AF_UNIX;
char name[16];
// Create a unique socket name based on current process id.
sprintf(name, "%u", getpid());
strcpy(cliaddr.sun_path, name);
if (bind(sock, (struct sockaddr *)&cliaddr, sizeof(cliaddr)) < 0) {
perror("Socket failure: Binding socket failed");
return -1;
}
handle->socket = sock;
handle->name = new char[strlen(name) + 1];
strcpy(handle->name, name);
return 0;
}
// method to close socket
int closeSocket() {
if (!handle) {
return -1;
}
if (handle->name) {
unlink(handle->name);
delete[] handle->name;
}
close(handle->socket);
delete handle;
return 0;
}
public:
ipcSocketCom() = default;
ipcSocketCom(bool isServer) {
if (isServer) {
checkSysCallErrors(createSocket());
} else {
checkSysCallErrors(openSocket());
}
}
~ipcSocketCom() {
}
int closeThisSock() {
return closeSocket();
}
// method to receive shareable handle via socket
int recvShareableHdl(hipShareableHdl *shHandle) {
int dummy_data;
struct msghdr msg;
struct iovec iov[1];
// Union to guarantee alignment requirements for control array
union {
struct cmsghdr cm;
char control[CMSG_SPACE(sizeof(int))];
} control_un;
struct cmsghdr *cmptr;
ssize_t n;
int receivedfd;
msg.msg_name = NULL;
msg.msg_namelen = 0;
msg.msg_control = control_un.control;
msg.msg_controllen = sizeof(control_un.control);
iov[0].iov_base = &dummy_data;
iov[0].iov_len = sizeof(dummy_data);
msg.msg_iov = iov;
msg.msg_iovlen = 1;
if ((n = recvmsg(handle->socket, &msg, 0)) <= 0) {
perror("Socket failure: Receiving data over socket failed");
return -1;
}
if (((cmptr = CMSG_FIRSTHDR(&msg)) != NULL) &&
(cmptr->cmsg_len == CMSG_LEN(sizeof(int)))) {
if ((cmptr->cmsg_level != SOL_SOCKET) || (cmptr->cmsg_type != SCM_RIGHTS)) {
return -1;
}
memmove(&receivedfd, CMSG_DATA(cmptr), sizeof(receivedfd));
*(int *)shHandle = receivedfd;
} else {
return -1;
}
return 0;
}
// method to send shareable handle via sockets
int sendShareableHdl(hipShareableHdl shareableHdl, Process process) {
struct msghdr msg;
struct iovec iov[1];
int dummy_data = 0;
union {
struct cmsghdr cm;
char control[CMSG_SPACE(sizeof(int))];
} control_un;
struct cmsghdr *cmptr;
struct sockaddr_un cliaddr;
// Construct client address to send this SHareable handle to
bzero(&cliaddr, sizeof(cliaddr));
cliaddr.sun_family = AF_UNIX;
char temp[10];
sprintf(temp, "%u", process);
strcpy(cliaddr.sun_path, temp);
// Send corresponding shareable handle to the client
int sendfd = (int)shareableHdl;
msg.msg_control = control_un.control;
msg.msg_controllen = sizeof(control_un.control);
cmptr = CMSG_FIRSTHDR(&msg);
cmptr->cmsg_len = CMSG_LEN(sizeof(int));
cmptr->cmsg_level = SOL_SOCKET;
cmptr->cmsg_type = SCM_RIGHTS;
memmove(CMSG_DATA(cmptr), &sendfd, sizeof(sendfd));
msg.msg_name = (void *)&cliaddr;
msg.msg_namelen = sizeof(struct sockaddr_un);
iov[0].iov_base = &dummy_data;
iov[0].iov_len = sizeof(dummy_data);
msg.msg_iov = iov;
msg.msg_iovlen = 1;
ssize_t sendResult = sendmsg(handle->socket, &msg, 0);
if (sendResult <= 0) {
perror("Socket failure: Sending data over socket failed");
return -1;
}
return 0;
}
};
#endif
constexpr int threadsPerBlk = 64;