diff --git a/catch/unit/virtualMemoryManagement/CMakeLists.txt b/catch/unit/virtualMemoryManagement/CMakeLists.txt index f540fdd25b..975308b5c4 100644 --- a/catch/unit/virtualMemoryManagement/CMakeLists.txt +++ b/catch/unit/virtualMemoryManagement/CMakeLists.txt @@ -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") diff --git a/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc b/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc index e05d82f0cf..dd6c19e11b 100644 --- a/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc +++ b/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc @@ -32,7 +32,6 @@ THE SOFTWARE. */ #include - #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 } /** diff --git a/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc b/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc index 4e5a6ff040..be379b7153 100644 --- a/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc +++ b/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc @@ -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 -#include -#include - #include - #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 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(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 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(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 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(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. * @} diff --git a/catch/unit/virtualMemoryManagement/hip_vmm_common.hh b/catch/unit/virtualMemoryManagement/hip_vmm_common.hh index a43af62758..33f38ecfd2 100644 --- a/catch/unit/virtualMemoryManagement/hip_vmm_common.hh +++ b/catch/unit/virtualMemoryManagement/hip_vmm_common.hh @@ -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 +#include +#include +#include +#include +#include +#include +#include +#include +#include +#endif #include "hip_test_context.hh" #define checkVMMSupported(device) \ @@ -46,4 +57,223 @@ THE SOFTWARE. } \ } -constexpr int threadsPerBlk = 64; \ No newline at end of file +#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;