diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt b/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt index f0ad472af7..3bf34bb852 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt @@ -32,13 +32,12 @@ if(UNIX) if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} - hipMemImportFromShareableHandle.cc hipMemGetHandleForAddressRange.cc) endif() - -set(TEST_SRC - ${TEST_SRC} - hipMemExportToShareableHandle.cc) + set(TEST_SRC + ${TEST_SRC} + hipMemExportToShareableHandle.cc + hipMemImportFromShareableHandle.cc) endif() if(HIP_PLATFORM MATCHES "amd") diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetHandleForAddressRange.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetHandleForAddressRange.cc index 027f8f6e61..475e5c2250 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetHandleForAddressRange.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetHandleForAddressRange.cc @@ -94,27 +94,36 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_Negative") { constexpr int sizeBytes = size * sizeof(int); HIP_CHECK(hipMalloc(&dptr, sizeBytes)); + #if HT_AMD + hipDeviceptr_t nptr = nullptr; + #else + hipDeviceptr_t nptr = 0; + #endif + SECTION("nullptr") { - HIP_CHECK_ERROR(hipMemGetHandleForAddressRange(&handle, nullptr, sizeBytes, + HIP_CHECK_ERROR(hipMemGetHandleForAddressRange(&handle, nptr, sizeBytes, hipMemRangeHandleTypeDmaBufFd, 0), hipErrorInvalidValue); } SECTION("size 0") { HIP_CHECK_ERROR( - hipMemGetHandleForAddressRange(&handle, dptr, 0, hipMemRangeHandleTypeDmaBufFd, 0), + hipMemGetHandleForAddressRange(&handle, reinterpret_cast(dptr), + 0, hipMemRangeHandleTypeDmaBufFd, 0), hipErrorInvalidValue); } SECTION("Invalid Handle type") { - HIP_CHECK_ERROR(hipMemGetHandleForAddressRange(&handle, dptr, sizeBytes, - static_cast(-1), 0), + HIP_CHECK_ERROR(hipMemGetHandleForAddressRange(&handle, + reinterpret_cast(dptr), sizeBytes, + static_cast(-1), 0), hipErrorInvalidValue); } SECTION("Invalid Flags") { - HIP_CHECK_ERROR(hipMemGetHandleForAddressRange(&handle, dptr, sizeBytes, - hipMemRangeHandleTypeDmaBufFd, 0xFF), + HIP_CHECK_ERROR(hipMemGetHandleForAddressRange(&handle, + reinterpret_cast(dptr), sizeBytes, + hipMemRangeHandleTypeDmaBufFd, 0xFF), hipErrorInvalidValue); } @@ -123,7 +132,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_Negative") { HIP_CHECK(hipMalloc(&devMem, sizeBytes)); HIP_CHECK(hipFree(devMem)); - HIP_CHECK_ERROR(hipMemGetHandleForAddressRange(&handle, devMem, sizeBytes, + HIP_CHECK_ERROR(hipMemGetHandleForAddressRange(&handle, reinterpret_cast(devMem), sizeBytes, hipMemRangeHandleTypeDmaBufFd, 0), hipErrorInvalidValue); } @@ -131,7 +140,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_Negative") { SECTION("With Host memory") { int* hptr = new int[size]; HIP_CHECK_ERROR( - hipMemGetHandleForAddressRange(&handle, hptr, sizeBytes, hipMemRangeHandleTypeDmaBufFd, 0), + hipMemGetHandleForAddressRange(&handle, reinterpret_cast(hptr), sizeBytes, hipMemRangeHandleTypeDmaBufFd, 0), hipErrorInvalidValue); delete[] hptr; } @@ -146,9 +155,10 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_Negative") { assert(granularity > 0); size_t size_mem = ((granularity + sizeBytes - 1) / granularity) * granularity; - hipDeviceptr_t ptrA = nullptr; - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, granularity, 0, 0)); - REQUIRE(ptrA != nullptr); + hipDeviceptr_t ptrA; + HIP_CHECK(hipMemAddressReserve(reinterpret_cast(&ptrA), size_mem, granularity, 0, 0)); + + REQUIRE(reinterpret_cast(ptrA) != nullptr); HIP_CHECK_ERROR( hipMemGetHandleForAddressRange(&handle, ptrA, size_mem, hipMemRangeHandleTypeDmaBufFd, 0), @@ -189,7 +199,7 @@ hipDeviceptr_t createVirtualMemoryAndFillData(int size, int* reservedAddrSize, i size_t granularity = GetGranularity(device); if (granularity <= 0) { std::cout << "Invalid Granularity" << std::endl; - return nullptr; + return 0; } int* srcHostMem = reinterpret_cast(malloc(size * sizeof(int))); @@ -198,21 +208,21 @@ hipDeviceptr_t createVirtualMemoryAndFillData(int size, int* reservedAddrSize, i } size_t size_mem = ((granularity + (size * sizeof(int)) - 1) / granularity) * granularity; - hipDeviceptr_t ptrA = nullptr; - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, granularity, 0, 0)); - REQUIRE(ptrA != nullptr); + hipDeviceptr_t ptrA; + HIP_CHECK(hipMemAddressReserve(reinterpret_cast(&ptrA), size_mem, granularity, 0, 0)); + REQUIRE(reinterpret_cast(ptrA) != nullptr); hipMemGenericAllocationHandle_t handle = GetPhysicalMemory(device, size_mem); - HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemMap(reinterpret_cast(ptrA), size_mem, 0, handle, 0)); hipMemAccessDesc accessDesc = {}; accessDesc.location.type = hipMemLocationTypeDevice; accessDesc.location.id = device; accessDesc.flags = hipMemAccessFlagsProtReadWrite; - HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemSetAccess(reinterpret_cast(ptrA), size_mem, &accessDesc, 1)); - HIP_CHECK(hipMemcpy(ptrA, srcHostMem, size * sizeof(int), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(reinterpret_cast(ptrA), srcHostMem, size * sizeof(int), hipMemcpyHostToDevice)); *reservedAddrSize = size_mem; return ptrA; @@ -291,12 +301,13 @@ bool validateHandle(int handle, int size, int device = 0) { TEST_CASE("Unit_hipMemGetHandleForAddressRange_DeviceMemory") { constexpr int size = 1024; constexpr int sizeBytes = size * sizeof(int); + CTX_CREATE(); void* srcDevMem = createDeviceMemoryAndFillData(size); REQUIRE(srcDevMem != nullptr); int handle = -1; - HIP_CHECK(hipMemGetHandleForAddressRange(&handle, srcDevMem, sizeBytes, + HIP_CHECK(hipMemGetHandleForAddressRange(&handle, reinterpret_cast(srcDevMem), sizeBytes, hipMemRangeHandleTypeDmaBufFd, 0)); REQUIRE(handle > 0); @@ -307,6 +318,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_DeviceMemory") { REQUIRE(validateHandle(handle, size) == true); HIP_CHECK(hipFree(srcDevMem)); + CTX_DESTROY(); } /** @@ -324,6 +336,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_DeviceMemory") { * - HIP_VERSION >= 7.0 */ TEST_CASE("Unit_hipMemGetHandleForAddressRange_VM") { + CTX_CREATE(); hipDevice_t device; constexpr int kDeviceId = 0; HIP_CHECK(hipDeviceGet(&device, kDeviceId)); @@ -332,10 +345,10 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_VM") { constexpr int size = 1024; constexpr int sizeBytes = size * sizeof(int); - hipDeviceptr_t ptrA = nullptr; + hipDeviceptr_t ptrA; int reservedAddrSize; ptrA = createVirtualMemoryAndFillData(size, &reservedAddrSize); - REQUIRE(ptrA != nullptr); + REQUIRE(reinterpret_cast(ptrA) != nullptr); int handle = -1; HIP_CHECK( @@ -344,8 +357,9 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_VM") { REQUIRE(validateHandle(handle, size) == true); - HIP_CHECK(hipMemUnmap(ptrA, reservedAddrSize)); - HIP_CHECK(hipMemAddressFree(ptrA, reservedAddrSize)); + HIP_CHECK(hipMemUnmap(reinterpret_cast(ptrA), reservedAddrSize)); + HIP_CHECK(hipMemAddressFree(reinterpret_cast(ptrA), reservedAddrSize)); + CTX_DESTROY(); } /** @@ -364,6 +378,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_VM") { */ TEST_CASE("Unit_hipMemGetHandleForAddressRange_DeviceMemory_InAnotherDevice", "[multigpu]") { + CTX_CREATE(); int deviceCount = 0; HIP_CHECK(hipGetDeviceCount(&deviceCount)); if (deviceCount < 2) { @@ -387,7 +402,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_DeviceMemory_InAnotherDevice", REQUIRE(srcDevMem != nullptr); int handle = -1; - HIP_CHECK(hipMemGetHandleForAddressRange(&handle, srcDevMem, sizeBytes, + HIP_CHECK(hipMemGetHandleForAddressRange(&handle, reinterpret_cast(srcDevMem), sizeBytes, hipMemRangeHandleTypeDmaBufFd, 0)); REQUIRE(handle > 0); @@ -398,7 +413,8 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_DeviceMemory_InAnotherDevice", REQUIRE(validateHandle(handle, size, dstDeviceId) == true); HIP_CHECK(hipFree(srcDevMem)); - HIP_CHECK(hipDeviceReset()) + HIP_CHECK(hipDeviceReset()); + CTX_DESTROY(); } /** @@ -417,6 +433,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_DeviceMemory_InAnotherDevice", */ TEST_CASE("Unit_hipMemGetHandleForAddressRange_VM_InAnotherDevice", "[multigpu]") { + CTX_CREATE(); int deviceCount = 0; HIP_CHECK(hipGetDeviceCount(&deviceCount)); if (deviceCount < 2) { @@ -436,10 +453,10 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_VM_InAnotherDevice", constexpr int size = 1024; constexpr int sizeBytes = kNumElemsSize * sizeof(int); - hipDeviceptr_t ptrA = nullptr; + hipDeviceptr_t ptrA; int reservedAddrSize; ptrA = createVirtualMemoryAndFillData(size, &reservedAddrSize); - REQUIRE(ptrA != nullptr); + REQUIRE(reinterpret_cast(ptrA) != nullptr); int handle = 0; HIP_CHECK( @@ -453,10 +470,11 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_VM_InAnotherDevice", REQUIRE(validateHandle(handle, size, dstDeviceId) == true); - HIP_CHECK(hipMemUnmap(ptrA, reservedAddrSize)); - HIP_CHECK(hipMemAddressFree(ptrA, reservedAddrSize)); + HIP_CHECK(hipMemUnmap(reinterpret_cast(ptrA), reservedAddrSize)); + HIP_CHECK(hipMemAddressFree(reinterpret_cast(ptrA), reservedAddrSize)); - HIP_CHECK(hipDeviceReset()) + HIP_CHECK(hipDeviceReset()); + CTX_DESTROY(); } #if __linux__ @@ -486,6 +504,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_MulProc_Socket_DeviceMem") { if (pid == 0) { // child REQUIRE(close(fd[1]) == 0); REQUIRE(close(fdSig[0]) == 0); + CTX_CREATE(); // Wait for parent process to create the socket. int size_mem = 0; @@ -511,6 +530,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_MulProc_Socket_DeviceMem") { // Validate the handle REQUIRE(validateHandle(shHandle, size_mem / sizeof(int))); + CTX_DESTROY(); checkSysCallErrors(sockObj.closeThisSock()); REQUIRE(close(fd[0]) == 0); @@ -520,6 +540,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_MulProc_Socket_DeviceMem") { REQUIRE(close(fd[0]) == 0); REQUIRE(close(fdSig[1]) == 0); + CTX_CREATE(); constexpr int size = 1024; constexpr int sizeBytes = size * sizeof(int); @@ -532,7 +553,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_MulProc_Socket_DeviceMem") { REQUIRE(srcDevMem != nullptr); int handle = -1; - HIP_CHECK(hipMemGetHandleForAddressRange(&handle, srcDevMem, sizeBytes, + HIP_CHECK(hipMemGetHandleForAddressRange(&handle, reinterpret_cast(srcDevMem), sizeBytes, hipMemRangeHandleTypeDmaBufFd, 0)); int size_mem = sizeBytes; @@ -548,6 +569,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_MulProc_Socket_DeviceMem") { int status; REQUIRE(wait(&status) >= 0); REQUIRE(status == 0); + CTX_DESTROY(); // Free all resources checkSysCallErrors(sockObj.closeThisSock()); // HIP_CHECK(hipMemRelease(handle)); @@ -582,6 +604,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_MulProc_Socket_VM") { if (pid == 0) { // child REQUIRE(close(fd[1]) == 0); REQUIRE(close(fdSig[0]) == 0); + CTX_CREATE(); // Wait for parent process to create the socket. int size_mem = 0; REQUIRE(read(fd[0], &size_mem, sizeof(int)) >= 0); @@ -604,6 +627,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_MulProc_Socket_VM") { // Validate handle REQUIRE(validateHandle(shHandle, size_mem / sizeof(int))); + CTX_DESTROY(); checkSysCallErrors(sockObj.closeThisSock()); REQUIRE(close(fd[0]) == 0); @@ -612,6 +636,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_MulProc_Socket_VM") { } else { // parent REQUIRE(close(fd[0]) == 0); REQUIRE(close(fdSig[1]) == 0); + CTX_CREATE(); constexpr int N = 1024; int buffer_size = N * sizeof(int); @@ -620,10 +645,10 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_MulProc_Socket_VM") { HIP_CHECK(hipDeviceGet(&device, 0)); checkVMMSupported(device); - hipDeviceptr_t ptrA = nullptr; + hipDeviceptr_t ptrA; int reservedAddrSize; ptrA = createVirtualMemoryAndFillData(N, &reservedAddrSize); - REQUIRE(ptrA != nullptr); + REQUIRE(reinterpret_cast(ptrA) != nullptr); int handle = -1; HIP_CHECK(hipMemGetHandleForAddressRange(&handle, ptrA, buffer_size, @@ -649,8 +674,9 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_MulProc_Socket_VM") { REQUIRE(close(fd[1]) == 0); REQUIRE(close(fdSig[0]) == 0); - HIP_CHECK(hipMemUnmap(ptrA, reservedAddrSize)); - HIP_CHECK(hipMemAddressFree(ptrA, reservedAddrSize)); + HIP_CHECK(hipMemUnmap(reinterpret_cast(ptrA), reservedAddrSize)); + HIP_CHECK(hipMemAddressFree(reinterpret_cast(ptrA), reservedAddrSize)); + CTX_DESTROY(); } } @@ -665,7 +691,7 @@ void launchForDevMem() { void* srcDevMem = createDeviceMemoryAndFillData(size); int handle = -1; - HIP_CHECK(hipMemGetHandleForAddressRange(&handle, srcDevMem, sizeBytes, + HIP_CHECK(hipMemGetHandleForAddressRange(&handle, reinterpret_cast(srcDevMem), sizeBytes, hipMemRangeHandleTypeDmaBufFd, 0)); REQUIRE(handle > 0); HIP_CHECK(hipFree(srcDevMem)); @@ -678,18 +704,18 @@ void launchForVM() { constexpr int size = 1024; constexpr int sizeBytes = size * sizeof(int); - hipDeviceptr_t ptrA = nullptr; + hipDeviceptr_t ptrA; int reservedAddrSize; ptrA = createVirtualMemoryAndFillData(size, &reservedAddrSize); - REQUIRE(ptrA != nullptr); + REQUIRE(reinterpret_cast(ptrA) != nullptr); int handle = -1; HIP_CHECK( hipMemGetHandleForAddressRange(&handle, ptrA, sizeBytes, hipMemRangeHandleTypeDmaBufFd, 0)); REQUIRE(handle > 0); - HIP_CHECK(hipMemUnmap(ptrA, reservedAddrSize)); - HIP_CHECK(hipMemAddressFree(ptrA, reservedAddrSize)); + HIP_CHECK(hipMemUnmap(reinterpret_cast(ptrA), reservedAddrSize)); + HIP_CHECK(hipMemAddressFree(reinterpret_cast(ptrA), reservedAddrSize)); } /** @@ -757,7 +783,7 @@ TEST_CASE("Unit_hipMemGetHandleForAddressRange_DifferentOffsets") { for (int i = 0; i < size; i++) { handle = -1; - HIP_CHECK(hipMemGetHandleForAddressRange(&handle, dptr + i, sizeBytes - (i * sizeof(int)), + HIP_CHECK(hipMemGetHandleForAddressRange(&handle, reinterpret_cast(dptr + i), sizeBytes - (i * sizeof(int)), hipMemRangeHandleTypeDmaBufFd, 0)); REQUIRE(handle > 0); } diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc index a69fdfdc3c..80b892ca9a 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc @@ -58,7 +58,7 @@ static __global__ void square_kernel(int* Buff) { * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemImportFromShareableHandle_Positive_Basic") { - CTX_DESTROY(); + CTX_CREATE(); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); @@ -85,6 +85,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_Positive_Basic") { hipMemHandleTypePosixFileDescriptor)); HIP_CHECK(hipMemRelease(handle)); HIP_CHECK(hipMemRelease(imported_handle)); + CTX_DESTROY(); } @@ -168,6 +169,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ChldUseHdl") { if (pid == 0) { // child REQUIRE(close(fd[1]) == 0); REQUIRE(close(fdSig[0]) == 0); + CTX_CREATE(); // Wait for parent process to create the socket. size_t size_mem = 0; REQUIRE(read(fd[0], &size_mem, sizeof(size_t)) >= 0); @@ -185,7 +187,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ChldUseHdl") { reinterpret_cast(static_cast(shHandle)), hipMemHandleTypePosixFileDescriptor)); // Allocate virtual address range - hipDeviceptr_t ptrA; + void* ptrA; HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, imported_handle, 0)); // Set access @@ -201,11 +203,11 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ChldUseHdl") { A_h[idx] = idx; C_h[idx] = idx * idx; } - HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyHtoD(reinterpret_cast(ptrA), A_h.data(), buffer_size)); // Invoke kernel hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0, reinterpret_cast(ptrA)); - HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast(ptrA), buffer_size)); HIP_CHECK(hipDeviceSynchronize()); // validate REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); @@ -213,6 +215,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ChldUseHdl") { // free resources HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); checkSysCallErrors(sockObj.closeThisSock()); REQUIRE(close(fd[0]) == 0); REQUIRE(close(fdSig[1]) == 0); @@ -220,6 +223,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ChldUseHdl") { } else { // parent REQUIRE(close(fd[0]) == 0); REQUIRE(close(fdSig[1]) == 0); + CTX_CREATE(); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); @@ -257,6 +261,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ChldUseHdl") { // Free all resources checkSysCallErrors(sockObj.closeThisSock()); HIP_CHECK(hipMemRelease(handle)); + CTX_DESTROY(); REQUIRE(close(fd[1]) == 0); REQUIRE(close(fdSig[0]) == 0); } @@ -297,6 +302,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ParntChldUseHdl") { if (pid == 0) { // child REQUIRE(close(fd[1]) == 0); REQUIRE(close(fdSig[0]) == 0); + CTX_CREATE(); // Wait for parent process to create the socket. size_t size_mem = 0; REQUIRE(read(fd[0], &size_mem, sizeof(size_t)) >= 0); @@ -318,7 +324,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ParntChldUseHdl") { reinterpret_cast(static_cast(shHandle)), hipMemHandleTypePosixFileDescriptor)); // Allocate virtual address range - hipDeviceptr_t ptrA; + void* ptrA; HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, imported_handle, 0)); // Set access @@ -328,7 +334,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ParntChldUseHdl") { 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)); + HIP_CHECK(hipMemcpyHtoD(reinterpret_cast(ptrA), A_h.data(), buffer_size)); // Invoke kernel hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0, reinterpret_cast(ptrA)); @@ -340,10 +346,12 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ParntChldUseHdl") { checkSysCallErrors(sockObj.closeThisSock()); REQUIRE(close(fd[0]) == 0); REQUIRE(close(fdSig[1]) == 0); + CTX_DESTROY(); exit(0); } else { // parent REQUIRE(close(fd[0]) == 0); REQUIRE(close(fdSig[1]) == 0); + CTX_CREATE(); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); checkVMMSupported(device); @@ -367,7 +375,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ParntChldUseHdl") { hipMemHandleTypePosixFileDescriptor, 0)); // Allocate virtual address range - hipDeviceptr_t ptrA; + void* ptrA; HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); // Set access @@ -392,7 +400,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ParntChldUseHdl") { REQUIRE(status == 0); // Check results of Vmm data processing in child - HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast(ptrA), buffer_size)); // validate REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); @@ -401,6 +409,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ParntChldUseHdl") { HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); HIP_CHECK(hipMemRelease(handle)); checkSysCallErrors(sockObj.closeThisSock()); + CTX_DESTROY(); REQUIRE(close(fd[1]) == 0); REQUIRE(close(fdSig[0]) == 0); } @@ -438,6 +447,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_GrndChldUseHdl") { // Wait for parent process to create the socket. size_t size_mem = 0; REQUIRE(read(fd[0], &size_mem, sizeof(size_t)) >= 0); + CTX_CREATE(); // Open Socket as client ipcSocketCom sockObj(false); @@ -456,7 +466,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_GrndChldUseHdl") { reinterpret_cast(static_cast(shHandle)), hipMemHandleTypePosixFileDescriptor)); // Allocate virtual address range - hipDeviceptr_t ptrA; + void* ptrA; HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, imported_handle, 0)); // Set access @@ -472,11 +482,11 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_GrndChldUseHdl") { A_h[idx] = idx; C_h[idx] = idx * idx; } - HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyHtoD(reinterpret_cast(ptrA), A_h.data(), buffer_size)); // Invoke kernel hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0, reinterpret_cast(ptrA)); - HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), reinterpret_cast(ptrA), buffer_size)); HIP_CHECK(hipDeviceSynchronize()); // validate REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); @@ -484,6 +494,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_GrndChldUseHdl") { // free resources HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); checkSysCallErrors(sockObj.closeThisSock()); REQUIRE(close(fd[0]) == 0); REQUIRE(close(fdSig[1]) == 0); @@ -503,7 +514,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_GrndChldUseHdl") { REQUIRE(close(fdpid[1]) == 0); int pid_grChld = 0; REQUIRE(read(fdpid[0], &pid_grChld, sizeof(pid_grChld)) >= 0); - + CTX_CREATE(); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); checkVMMSupported(device); @@ -541,6 +552,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_GrndChldUseHdl") { // Free all resources HIP_CHECK(hipMemRelease(handle)); + CTX_DESTROY(); checkSysCallErrors(sockObj.closeThisSock()); REQUIRE(close(fd[1]) == 0); REQUIRE(close(fdSig[0]) == 0);