From dec3869d6d4c1eb2eb685b74ccb22b9dd41fce00 Mon Sep 17 00:00:00 2001 From: "Arandjelovic, Marko" Date: Tue, 13 May 2025 00:21:38 +0200 Subject: [PATCH] SWDEV-499927 - Enable Virtual Memory tests on NV platform (#79) --- catch/unit/memory/hip_vmm_common.hh | 51 ---- .../virtualMemoryManagement/CMakeLists.txt | 26 +- .../hipMemAddressFree.cc | 4 +- .../hipMemAddressReserve.cc | 17 +- .../virtualMemoryManagement/hipMemCreate.cc | 40 ++- .../hipMemExportToShareableHandle.cc | 10 +- .../hipMemGetAllocationGranularity.cc | 6 +- ...hipMemGetAllocationPropertiesFromHandle.cc | 7 +- .../hipMemImportFromShareableHandle.cc | 18 +- .../unit/virtualMemoryManagement/hipMemMap.cc | 97 ++----- .../virtualMemoryManagement/hipMemRelease.cc | 6 +- .../hipMemSetGetAccess.cc | 265 ++++++++---------- .../virtualMemoryManagement/hipMemUnmap.cc | 6 +- .../virtualMemoryManagement/hipMemVmm_old.cc | 25 +- .../virtualMemoryManagement/hip_vmm_common.hh | 11 - 15 files changed, 256 insertions(+), 333 deletions(-) delete mode 100644 catch/unit/memory/hip_vmm_common.hh diff --git a/catch/unit/memory/hip_vmm_common.hh b/catch/unit/memory/hip_vmm_common.hh deleted file mode 100644 index 24ecb6408a..0000000000 --- a/catch/unit/memory/hip_vmm_common.hh +++ /dev/null @@ -1,51 +0,0 @@ -/* -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. -*/ - -#pragma once - -#include "hip_test_context.hh" - -#define checkVMMSupported(device) {\ - int value = 0;\ - hipDeviceAttribute_t\ - attr = hipDeviceAttributeVirtualMemoryManagementSupported;\ - HIP_CHECK(\ - hipDeviceGetAttribute(&value, attr, device));\ - if (value == 0) {\ - HipTest::HIP_SKIP_TEST("Machine does not support VMM. Skipping Test..");\ - return;\ - }\ -} - -#define checkVMMSupportedRetVal(device) {\ - int value = 0;\ - hipDeviceAttribute_t\ - attr = hipDeviceAttributeVirtualMemoryManagementSupported;\ - HIP_CHECK(\ - hipDeviceGetAttribute(&value, attr, device));\ - if (value == 0) {\ - HipTest::HIP_SKIP_TEST("Machine does not support VMM. Skipping Test..");\ - return true;\ - }\ -} - -constexpr int threadsPerBlk = 64; \ No newline at end of file diff --git a/catch/unit/virtualMemoryManagement/CMakeLists.txt b/catch/unit/virtualMemoryManagement/CMakeLists.txt index 3fd7375cfd..e7bd4ff2b6 100644 --- a/catch/unit/virtualMemoryManagement/CMakeLists.txt +++ b/catch/unit/virtualMemoryManagement/CMakeLists.txt @@ -32,26 +32,32 @@ if(UNIX) if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} - hipMemExportToShareableHandle.cc hipMemImportFromShareableHandle.cc) endif() + +set(TEST_SRC + ${TEST_SRC} + hipMemExportToShareableHandle.cc) endif() if(HIP_PLATFORM MATCHES "amd") set(TEST_SRC ${TEST_SRC} - hipMemAddressFree.cc - hipMemAddressReserve.cc - hipMemCreate.cc - hipMemSetGetAccess.cc - hipMemGetAllocationPropertiesFromHandle.cc - hipMemMap.cc - hipMemRelease.cc - hipMemUnmap.cc - hipMemVmm_old.cc hipGetProcAddressVmmApis.cc) endif() +set(TEST_SRC + ${TEST_SRC} + hipMemAddressFree.cc + hipMemAddressReserve.cc + hipMemCreate.cc + hipMemSetGetAccess.cc + hipMemGetAllocationPropertiesFromHandle.cc + hipMemMap.cc + hipMemRelease.cc + hipMemUnmap.cc + hipMemVmm_old.cc) + hip_add_exe_to_target(NAME VirtualMemoryManagementTest TEST_SRC ${TEST_SRC} TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) diff --git a/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc b/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc index 470d290a5b..284be767f5 100644 --- a/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc +++ b/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc @@ -50,6 +50,7 @@ TEST_CASE("Unit_hipMemAddressFree_negative") { size_t buffer_size = N * sizeof(int); int deviceId = 0; hipDevice_t device; + CTX_CREATE(); HIP_CHECK(hipDeviceGet(&device, deviceId)); checkVMMSupported(device); hipMemAllocationProp prop{}; @@ -65,12 +66,13 @@ TEST_CASE("Unit_hipMemAddressFree_negative") { HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); SECTION("nullptr to devptr") { - REQUIRE(hipMemAddressFree(nullptr, size_mem) == hipErrorInvalidValue); + REQUIRE(hipMemAddressFree((hipDeviceptr_t)nullptr, size_mem) == hipErrorInvalidValue); } SECTION("pass zero to size") { REQUIRE(hipMemAddressFree(ptrA, 0) == hipErrorInvalidValue); } HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** diff --git a/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc b/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc index a28e12b74c..6005668b49 100644 --- a/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc +++ b/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc @@ -56,6 +56,7 @@ TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") { constexpr int initializer = 0; int deviceId = 0; hipDevice_t device; + CTX_CREATE(); HIP_CHECK(hipDeviceGet(&device, deviceId)); checkVMMSupported(device); hipMemAllocationProp prop{}; @@ -82,7 +83,7 @@ TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") { for (int iter = 0; iter < 12; iter++) { alignmnt = alignmnt * 2; HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, alignmnt, 0, 0)); - REQUIRE((reinterpret_cast(ptrA) % alignmnt) == 0); + REQUIRE((reinterpret_cast(ptrA) % alignmnt) == 0); std::fill(B_h.begin(), B_h.end(), initializer); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); // Set access @@ -99,6 +100,7 @@ TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") { HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); } HIP_CHECK(hipMemRelease(handle)); + CTX_DESTROY(); } /** @@ -117,6 +119,7 @@ TEST_CASE("Unit_hipMemAddressReserve_Negative") { size_t buffer_size = N * sizeof(int); int deviceId = 0; hipDevice_t device; + CTX_CREATE(); HIP_CHECK(hipDeviceGet(&device, deviceId)); checkVMMSupported(device); hipMemAllocationProp prop{}; @@ -135,18 +138,28 @@ TEST_CASE("Unit_hipMemAddressReserve_Negative") { } SECTION("pass size as 0") { +#if HT_AMD REQUIRE(hipMemAddressReserve(&ptrA, 0, 0, 0, 0) == hipErrorMemoryAllocation); +#else + REQUIRE(hipMemAddressReserve(&ptrA, 0, 0, 0, 0) == hipErrorInvalidValue); +#endif } #if HT_NVIDIA SECTION("pass non power of two for alignment") { - REQUIRE(hipMemAddressReserve(&ptrA, size_mem, 3, 0, 0) == hipErrorMemoryAllocation); + REQUIRE(hipMemAddressReserve(&ptrA, size_mem, 3, 0, 0) == hipErrorInvalidValue); } #endif SECTION("pass size as non multiple of host page size") { +#if HT_AMD REQUIRE(hipMemAddressReserve(&ptrA, (size_mem - 1), 0, 0, 0) == hipErrorMemoryAllocation); +#else + REQUIRE(hipMemAddressReserve(&ptrA, (size_mem - 1), 0, 0, 0) == hipErrorInvalidValue); +#endif } + + CTX_DESTROY(); } /** diff --git a/catch/unit/virtualMemoryManagement/hipMemCreate.cc b/catch/unit/virtualMemoryManagement/hipMemCreate.cc index 5501d4e90e..42de5181d0 100644 --- a/catch/unit/virtualMemoryManagement/hipMemCreate.cc +++ b/catch/unit/virtualMemoryManagement/hipMemCreate.cc @@ -63,6 +63,7 @@ static __global__ void square_kernel(int* Buff) { TEST_CASE("Unit_hipMemCreate_BasicAllocateDeAlloc_MultGranularity") { size_t granularity = 0; int deviceId = 0; + CTX_CREATE(); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); checkVMMSupported(device); @@ -79,6 +80,8 @@ TEST_CASE("Unit_hipMemCreate_BasicAllocateDeAlloc_MultGranularity") { HIP_CHECK(hipMemCreate(&handle, granularity * mul, &prop, 0)); HIP_CHECK(hipMemRelease(handle)); } + + CTX_DESTROY(); } /** @@ -98,6 +101,7 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap") { size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -136,6 +140,7 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap") { HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); HIP_CHECK(hipMemRelease(handle)); + CTX_DESTROY(); } /** @@ -155,6 +160,7 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") { size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -193,6 +199,7 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") { REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** @@ -212,6 +219,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") { size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -248,12 +256,13 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") { 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)); + reinterpret_cast(ptrA)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); HIP_CHECK(hipDeviceSynchronize()); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** @@ -274,6 +283,7 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") { constexpr int numOfBuffers = NUM_OF_BUFFERS; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -295,9 +305,9 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") { hipDeviceptr_t ptrA; HIP_CHECK(hipMemAddressReserve(&ptrA, (numOfBuffers * size_mem), 0, 0, 0)); for (int idx = 0; idx < numOfBuffers; idx++) { - uint64_t uiptr = reinterpret_cast(ptrA); + unsigned long long uiptr = reinterpret_cast(ptrA); uiptr = uiptr + idx * size_mem; - HIP_CHECK(hipMemMap(reinterpret_cast(uiptr), size_mem, 0, handle[idx], 0)); + HIP_CHECK(hipMemMap(reinterpret_cast(uiptr), size_mem, 0, handle[idx], 0)); HIP_CHECK(hipMemRelease(handle[idx])); } hipMemAccessDesc accessDesc = {}; @@ -315,18 +325,19 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") { } HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), numOfBuffers * buffer_size)); // Launch square kernel - hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0, - static_cast(ptrA)); + hipLaunchKernelGGL(square_kernel, dim3((N * numOfBuffers) / THREADS_PER_BLOCK), + dim3(THREADS_PER_BLOCK), 0, 0, reinterpret_cast(ptrA)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, numOfBuffers * buffer_size)); HIP_CHECK(hipDeviceSynchronize()); // Validate Results REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); for (int idx = 0; idx < numOfBuffers; idx++) { - uint64_t uiptr = reinterpret_cast(ptrA); + unsigned long long uiptr = reinterpret_cast(ptrA); uiptr = uiptr + idx * size_mem; - HIP_CHECK(hipMemUnmap(reinterpret_cast(uiptr), size_mem)); + HIP_CHECK(hipMemUnmap(reinterpret_cast(uiptr), size_mem)); } HIP_CHECK(hipMemAddressFree(ptrA, (numOfBuffers * size_mem))); + CTX_DESTROY(); } /** @@ -346,6 +357,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithMemset") { constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); constexpr int init_val = 0; + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -373,7 +385,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithMemset") { // Make the address accessible to GPU 0 HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); std::vector A_h(N); - HIP_CHECK(hipMemset(ptrA, init_val, buffer_size)); + HIP_CHECK(hipMemset(reinterpret_cast(ptrA), init_val, buffer_size)); HIP_CHECK(hipMemcpyDtoH(A_h.data(), ptrA, buffer_size)); for (int idx = 0; idx < N; idx++) { REQUIRE(A_h[idx] == init_val); @@ -381,6 +393,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithMemset") { HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); HIP_CHECK(hipMemRelease(handle)); + CTX_DESTROY(); } /** @@ -397,6 +410,7 @@ TEST_CASE("Unit_hipMemCreate_Negative") { size_t granularity = 0; int deviceId = 0; hipDevice_t device; + CTX_CREATE(); HIP_CHECK(hipDeviceGet(&device, deviceId)); checkVMMSupported(device); hipMemGenericAllocationHandle_t handle; @@ -433,15 +447,25 @@ TEST_CASE("Unit_hipMemCreate_Negative") { SECTION("pass location id as -1") { prop.location.id = -1; // set to non existing device +#if HT_AMD REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidValue); +#else + REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidDevice); +#endif } SECTION("pass location id as > highest device number") { int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); prop.location.id = numDevices; // set to non existing device +#if HT_AMD REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidValue); +#else + REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidDevice); +#endif } + + CTX_DESTROY(); } /** diff --git a/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc b/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc index e7a859ec7d..3c1dd98833 100644 --- a/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc +++ b/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc @@ -113,13 +113,13 @@ TEST_CASE("Unit_hipMemExportToShareableHandle_Negative_Parameters") { } #endif -#if HT_AMD SECTION("handle == nullptr") { - HIP_CHECK_ERROR(hipMemExportToShareableHandle(&shareable_handle, nullptr, - hipMemHandleTypePosixFileDescriptor, 0), - hipErrorInvalidValue); + HIP_CHECK_ERROR( + hipMemExportToShareableHandle(&shareable_handle, + (hipMemGenericAllocationHandle_t)nullptr, + hipMemHandleTypePosixFileDescriptor, 0), + hipErrorInvalidValue); } -#endif SECTION("invalid handleType") { HIP_CHECK_ERROR( diff --git a/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc b/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc index 0cc76ed026..724645887b 100644 --- a/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc +++ b/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc @@ -139,20 +139,18 @@ TEST_CASE("Unit_hipMemGetAllocationGranularity_NegativeTests") { REQUIRE(hipErrorInvalidValue == hipMemGetAllocationGranularity(nullptr, &prop, hipMemAllocationGranularityMinimum)); } -#if HT_AMD // segfaults on NVIDIA SECTION("Prop is nullptr") { REQUIRE( hipErrorInvalidValue == hipMemGetAllocationGranularity(&granularity, nullptr, hipMemAllocationGranularityMinimum)); } -#endif -#if HT_NVIDIA + SECTION("flag is invalid") { REQUIRE(hipErrorInvalidValue == hipMemGetAllocationGranularity(&granularity, &prop, (hipMemAllocationGranularity_flags)0xff)); } -#endif + #if HT_AMD // succeeds on NVIDIA SECTION("device id > highest device id") { int numDevices = 0; diff --git a/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc b/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc index ddc6f7c0fe..91dd0a2e8c 100644 --- a/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc +++ b/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc @@ -47,6 +47,7 @@ THE SOFTWARE. */ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") { hipDevice_t device; + CTX_CREATE(); HIP_CHECK(hipDeviceGet(&device, 0)); checkVMMSupported(device); hipMemGenericAllocationHandle_t handle; @@ -71,6 +72,7 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") { REQUIRE(prop_temp.location.type == prop.location.type); REQUIRE(prop_temp.location.id == prop.location.id); HIP_CHECK(hipMemRelease(handle)); + CTX_DESTROY(); } /** @@ -84,6 +86,7 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") { * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_Negative") { + CTX_CREATE(); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); checkVMMSupported(device); @@ -110,10 +113,12 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_Negative") { SECTION("null handle") { prop.location.type = hipMemLocationTypeInvalid; - REQUIRE(hipMemGetAllocationPropertiesFromHandle(&prop_temp, nullptr) == hipErrorInvalidValue); + REQUIRE(hipMemGetAllocationPropertiesFromHandle( + &prop_temp, (hipMemGenericAllocationHandle_t) nullptr) == hipErrorInvalidValue); } HIP_CHECK(hipMemRelease(handle)); + CTX_DESTROY(); } /** diff --git a/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc b/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc index 72670d7755..30100ca19a 100644 --- a/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc +++ b/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") { - HIP_CHECK(hipFree(0)); + CTX_DESTROY(); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); @@ -83,6 +83,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_Positive_Basic") { HIP_CHECK(hipMemImportFromShareableHandle(&imported_handle, &shareable_handle, hipMemHandleTypePosixFileDescriptor)); HIP_CHECK(hipMemRelease(handle)); + CTX_DESTROY(); } /** @@ -97,7 +98,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_Positive_Basic") { * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemImportFromShareableHandle_Negative_Parameters") { - HIP_CHECK(hipFree(0)); + CTX_CREATE(); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); @@ -122,13 +123,11 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_Negative_Parameters") { hipMemGenericAllocationHandle_t imported_handle; -#if HT_AMD SECTION("handle == nullptr") { HIP_CHECK_ERROR(hipMemImportFromShareableHandle(nullptr, shareable_handle, hipMemHandleTypePosixFileDescriptor), hipErrorInvalidValue); } -#endif SECTION("shareableHandle == nullptr") { HIP_CHECK_ERROR(hipMemImportFromShareableHandle(&imported_handle, nullptr, @@ -137,6 +136,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_Negative_Parameters") { } HIP_CHECK(hipMemRelease(handle)); + CTX_DESTROY(); } /** @@ -201,7 +201,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ChldUseHdl") { 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)); + reinterpret_cast(ptrA)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); HIP_CHECK(hipDeviceSynchronize()); // validate @@ -329,7 +329,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ParntChldUseHdl") { 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)); + reinterpret_cast(ptrA)); HIP_CHECK(hipDeviceSynchronize()); // free resources @@ -342,7 +342,6 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_ParntChldUseHdl") { } else { // parent REQUIRE(close(fd[0]) == 0); REQUIRE(close(fdSig[1]) == 0); - hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); checkVMMSupported(device); @@ -453,7 +452,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_GrndChldUseHdl") { // import the sareable handle HIP_CHECK(hipMemImportFromShareableHandle(&imported_handle, &shHandle, - hipMemHandleTypePosixFileDescriptor)); + hipMemHandleTypePosixFileDescriptor)); // Allocate virtual address range hipDeviceptr_t ptrA; HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); @@ -474,7 +473,7 @@ TEST_CASE("Unit_hipMemImportFromShareableHandle_MulProc_GrndChldUseHdl") { 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)); + reinterpret_cast(ptrA)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); HIP_CHECK(hipDeviceSynchronize()); // validate @@ -502,6 +501,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); + hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); checkVMMSupported(device); diff --git a/catch/unit/virtualMemoryManagement/hipMemMap.cc b/catch/unit/virtualMemoryManagement/hipMemMap.cc index 310976028d..1bbca28ecb 100644 --- a/catch/unit/virtualMemoryManagement/hipMemMap.cc +++ b/catch/unit/virtualMemoryManagement/hipMemMap.cc @@ -40,6 +40,15 @@ constexpr int N = (1 << 13); constexpr int num_buf = 3; constexpr int initializer = 0; +/** + 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 * ------------------------ @@ -56,6 +65,7 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") { constexpr int iterations = 20; size_t granularity = 0; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -93,17 +103,17 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") { HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); -#if HT_NVIDIA - square_kernel<<>>(static_cast(ptrA)); + square_kernel<<>>( + reinterpret_cast(ptrA)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); HIP_CHECK(hipStreamSynchronize(0)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); -#endif HIP_CHECK(hipMemUnmap(ptrA, size_mem)); } // Release resources HIP_CHECK(hipMemRelease(handle)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** @@ -121,6 +131,7 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") { TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") { size_t granularity = 0; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -160,13 +171,11 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") { HIP_CHECK(hipMemcpyHtoD(ptrA[buf], A_h.data(), buffer_size)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); -#if HT_NVIDIA square_kernel<<>>( - static_cast(ptrA[buf])); + reinterpret_cast(ptrA[buf])); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size)); HIP_CHECK(hipStreamSynchronize(0)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); -#endif HIP_CHECK(hipMemUnmap(ptrA[buf], size_mem)); } // Release resources @@ -174,6 +183,8 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") { for (int buf = 0; buf < num_buf; buf++) { HIP_CHECK(hipMemAddressFree(ptrA[buf], size_mem)); } + + CTX_DESTROY(); } /** @@ -191,6 +202,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") { TEST_CASE("Unit_hipMemMap_PhysicalMemory_Map2MultVMMs") { size_t granularity = 0; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -247,6 +259,8 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemory_Map2MultVMMs") { for (int buf = 0; buf < num_buf; buf++) { HIP_CHECK(hipMemAddressFree(ptrA[buf], size_mem)); } + + CTX_DESTROY(); } /** @@ -337,6 +351,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_MultiDev") { TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") { size_t granularity = 0; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -378,7 +393,7 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") { HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); #if HT_NVIDIA - square_kernel<<>>(static_cast(ptrA)); + square_kernel<<>>(reinterpret_cast(ptrA)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); HIP_CHECK(hipStreamSynchronize(0)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); @@ -390,6 +405,8 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") { HIP_CHECK(hipMemRelease(handle[buf])); } HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + + CTX_DESTROY(); } /** @@ -483,60 +500,6 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_MultiGPU") { HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); } -/** - * Test Description - * ------------------------ - * - Check if a partial part of a physical chunk can be mapped/unmapped - * to a smaller vmm address. - * ------------------------ - * - unit/virtualMemoryManagement/hipMemMap.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.1 - */ -TEST_CASE("Unit_hipMemMap_MapPartialPhysicalMem") { - int deviceId = 0; - size_t granularity = 0; - size_t buffer_size = N * sizeof(int); - hipDevice_t device; - HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device); - hipMemAllocationProp prop{}; - prop.type = hipMemAllocationTypePinned; - prop.location.type = hipMemLocationTypeDevice; - prop.location.id = device; // Current Devices - HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); - REQUIRE(granularity > 0); - size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; - hipMemGenericAllocationHandle_t handle; - // Allocate host memory and intialize data - std::vector A_h(N), B_h(N); - // Initialize with data - for (size_t idx = 0; idx < N; idx++) { - A_h[idx] = idx; - } - // Allocate a bigger physical memory chunk of twice size_mem - HIP_CHECK(hipMemCreate(&handle, 2 * size_mem, &prop, 0)); - // Allocate virtual address range of size size_mem - hipDeviceptr_t ptrA; - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); - hipMemAccessDesc accessDesc = {}; - accessDesc.location.type = hipMemLocationTypeDevice; - accessDesc.location.id = device; - accessDesc.flags = hipMemAccessFlagsProtReadWrite; - std::fill(B_h.begin(), B_h.end(), initializer); - HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); - HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); - HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); - HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); - REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); - HIP_CHECK(hipMemUnmap(ptrA, size_mem)); - // Release resources - HIP_CHECK(hipMemRelease(handle)); - HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); -} - /** * Test Description * ------------------------ @@ -553,6 +516,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") { size_t granularity = 0; size_t buffer_size = N * sizeof(int); hipDevice_t device; + CTX_CREATE(); HIP_CHECK(hipDeviceGet(&device, deviceId)); checkVMMSupported(device); hipMemAllocationProp prop{}; @@ -589,6 +553,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") { // Release resources HIP_CHECK(hipMemRelease(handle)); HIP_CHECK(hipMemAddressFree(ptrA, 2 * size_mem)); + CTX_DESTROY(); } /** @@ -604,6 +569,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") { TEST_CASE("Unit_hipMemMap_negative") { size_t granularity = 0; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -624,19 +590,16 @@ TEST_CASE("Unit_hipMemMap_negative") { HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); SECTION("nullptr to ptrA") { - REQUIRE(hipMemMap(nullptr, size_mem, 0, handle, 0) == hipErrorInvalidValue); + REQUIRE(hipMemMap((hipDeviceptr_t)nullptr, size_mem, 0, handle, 0) == hipErrorInvalidValue); } SECTION("pass zero to size") { - REQUIRE(hipMemMap(&ptrA, 0, 0, handle, 0) == hipErrorInvalidValue); - } - - SECTION("pass negative to offset") { - REQUIRE(hipMemMap(&ptrA, size_mem, -1, handle, 0) == hipErrorInvalidValue); + REQUIRE(hipMemMap(ptrA, 0, 0, handle, 0) == hipErrorInvalidValue); } HIP_CHECK(hipMemRelease(handle)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** diff --git a/catch/unit/virtualMemoryManagement/hipMemRelease.cc b/catch/unit/virtualMemoryManagement/hipMemRelease.cc index b41276e058..20358fb807 100644 --- a/catch/unit/virtualMemoryManagement/hipMemRelease.cc +++ b/catch/unit/virtualMemoryManagement/hipMemRelease.cc @@ -42,7 +42,11 @@ THE SOFTWARE. * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemRelease_negative") { - SECTION("Nullptr to handle") { REQUIRE(hipMemRelease(nullptr) == hipErrorInvalidValue); } + CTX_CREATE(); + SECTION("Nullptr to handle") { + REQUIRE(hipMemRelease((hipMemGenericAllocationHandle_t) nullptr) == hipErrorInvalidValue); + } + CTX_DESTROY(); } /** diff --git a/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc b/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc index 250ee8911f..1fe32c4e87 100644 --- a/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc +++ b/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc @@ -79,6 +79,7 @@ TEST_CASE("Unit_hipMemSetAccess_SetGet") { size_t buffer_size = N * sizeof(int); int deviceId = 0; hipDevice_t device; + CTX_CREATE(); HIP_CHECK(hipDeviceGet(&device, deviceId)); checkVMMSupported(device); hipMemAllocationProp prop{}; @@ -123,6 +124,7 @@ TEST_CASE("Unit_hipMemSetAccess_SetGet") { } HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** @@ -181,7 +183,7 @@ TEST_CASE("Unit_hipMemSetAccess_MultDevSetGet") { accessDesc[1].location.id = device1; accessDesc[1].flags = hipMemAccessFlagsProtReadWrite; // Make the address accessible to GPU 0 and 1 - HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc[0], 2)); + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, accessDesc, 2)); // Validate using hipMemGetAccess() hipMemLocation location; location.type = hipMemLocationTypeDevice; @@ -214,6 +216,7 @@ TEST_CASE("Unit_hipMemSetAccess_EntireVMMRangeSetGet") { size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -248,12 +251,13 @@ TEST_CASE("Unit_hipMemSetAccess_EntireVMMRangeSetGet") { unsigned long long flags = 0; // NOLINT HIP_CHECK(hipMemGetAccess(&flags, &location, ptrA)); REQUIRE(flags == hipMemAccessFlagsProtReadWrite); - uint64_t uiptr = reinterpret_cast(ptrA); + unsigned long long uiptr = reinterpret_cast(ptrA); uiptr += (size_mem - 1); - HIP_CHECK(hipMemGetAccess(&flags, &location, reinterpret_cast(uiptr))); + HIP_CHECK(hipMemGetAccess(&flags, &location, reinterpret_cast(uiptr))); REQUIRE(flags == hipMemAccessFlagsProtReadWrite); HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** @@ -270,6 +274,7 @@ TEST_CASE("Unit_hipMemGetAccess_NegTst") { size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -307,12 +312,13 @@ TEST_CASE("Unit_hipMemGetAccess_NegTst") { REQUIRE(status == hipErrorInvalidValue); status = hipMemGetAccess(&flags, nullptr, ptrA); REQUIRE(status == hipErrorInvalidValue); - uint64_t uiptr = reinterpret_cast(ptrA); + unsigned long long uiptr = reinterpret_cast(ptrA); uiptr += size_mem; - status = hipMemGetAccess(&flags, &location, reinterpret_cast(uiptr)); + status = hipMemGetAccess(&flags, &location, reinterpret_cast(uiptr)); REQUIRE(status == hipErrorInvalidValue); HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** @@ -332,16 +338,20 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") { size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0, devicecount = 0; - hipDevice_t device; HIP_CHECK(hipGetDeviceCount(&devicecount)); + if (devicecount < 2) { + HipTest::HIP_SKIP_TEST("Machine is Single GPU. Skipping Test.."); + return; + } for (deviceId = 0; deviceId < devicecount; deviceId++) { - HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device); + HIP_CHECK(hipSetDevice(deviceId)); + checkVMMSupported(deviceId); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; - prop.location.id = device; // Current Devices + prop.location.id = deviceId; // Current Devices HIP_CHECK( hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); @@ -357,7 +367,7 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") { // Set access hipMemAccessDesc accessDesc = {}; accessDesc.location.type = hipMemLocationTypeDevice; - accessDesc.location.id = device; + accessDesc.location.id = deviceId; accessDesc.flags = hipMemAccessFlagsProtReadWrite; // Make the address accessible to GPU deviceId std::vector A_h(N), B_h(N); @@ -371,16 +381,16 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") { for (int idx = 0; idx < N; idx++) { A_h[idx] = idx * idx; } - HIP_CHECK(hipSetDevice(deviceId)); // Launch square kernel hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0, - static_cast(ptrA)); + reinterpret_cast(ptrA)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); HIP_CHECK(hipDeviceSynchronize()); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); } + CTX_DESTROY(); } /** @@ -402,6 +412,7 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") { size_t buffer_size = N * sizeof(int); int dev = 0; hipDevice_t device; + CTX_CREATE(); HIP_CHECK(hipDeviceGet(&device, dev)); checkVMMSupported(device); hipMemAllocationProp prop{}; @@ -427,17 +438,7 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") { hipMemAccessDesc accessDesc = {}; accessDesc.location.type = hipMemLocationTypeDevice; accessDesc.location.id = device; - SECTION("Change ReadWrite to Read") { - accessDesc.flags = hipMemAccessFlagsProtReadWrite; - HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); - HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); - // Change property of virtual memory range to read only - accessDesc.flags = hipMemAccessFlagsProtRead; - HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); - // validate - HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); - REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); - } + SECTION("Change Read to ReadWrite") { accessDesc.flags = hipMemAccessFlagsProtRead; HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); @@ -448,6 +449,7 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") { HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); } + SECTION("Change Inaccessible to ReadWrite") { accessDesc.flags = hipMemAccessFlagsProtNone; HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); @@ -458,22 +460,26 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") { HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); } + #if HT_NVIDIA SECTION("Check error while writing on Read-Only memory") { accessDesc.flags = hipMemAccessFlagsProtRead; HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); REQUIRE(hipErrorInvalidValue == hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); } + SECTION("Check error while writing on inaccessible memory") { accessDesc.flags = hipMemAccessFlagsProtNone; HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); REQUIRE(hipErrorInvalidValue == hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); } #endif + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); // Release resources HIP_CHECK(hipMemRelease(handle)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** @@ -489,6 +495,7 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") { * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemSetAccess_Vmm2UnifiedMemCpy") { + CTX_CREATE(); auto managed = HmmAttrPrint(); if (managed != 1) { HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory.Skipping Test.."); @@ -531,7 +538,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2UnifiedMemCpy") { ptrA_h[idx] = idx; } HIP_CHECK(hipMemcpyHtoD(ptrA, ptrA_h, buffer_size)); - HIP_CHECK(hipMalloc(&ptrB, buffer_size)); + HIP_CHECK(hipMalloc(reinterpret_cast(&ptrB), buffer_size)); HIP_CHECK(hipMemcpyDtoD(ptrB, ptrA, buffer_size)); HIP_CHECK(hipMemcpyDtoH(ptrB_h, ptrB, buffer_size)); bool bPassed = true; @@ -542,11 +549,12 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2UnifiedMemCpy") { } } REQUIRE(bPassed == true); - HIP_CHECK(hipFree(ptrB)); - HIP_CHECK(hipFree(ptrA_h)); - HIP_CHECK(hipFree(ptrB_h)); + HIP_CHECK(hipFree(reinterpret_cast(ptrB))); + HIP_CHECK(hipFree(reinterpret_cast(ptrA_h))); + HIP_CHECK(hipFree(reinterpret_cast(ptrB_h))); HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** @@ -565,6 +573,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2DevMemCpy") { size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -597,13 +606,14 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2DevMemCpy") { A_h[idx] = idx; } HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); - HIP_CHECK(hipMalloc(&ptrB, buffer_size)); + HIP_CHECK(hipMalloc(reinterpret_cast(&ptrB), buffer_size)); HIP_CHECK(hipMemcpyDtoD(ptrB, ptrA, buffer_size)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrB, buffer_size)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); - HIP_CHECK(hipFree(ptrB)); + HIP_CHECK(hipFree(reinterpret_cast(ptrB))); HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** @@ -622,6 +632,13 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") { size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); + int devicecount = 0; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + if (devicecount < 2) { + HipTest::HIP_SKIP_TEST("Machine is Single GPU. Skipping Test.."); + return; + } int deviceId = 0, value = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -654,8 +671,6 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") { A_h[idx] = idx; } HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); - int devicecount = 0; - HIP_CHECK(hipGetDeviceCount(&devicecount)); // Check Peer Access for (deviceId = 1; deviceId < devicecount; deviceId++) { int canAccessPeer = 0; @@ -674,15 +689,22 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") { break; } HIP_CHECK(hipSetDevice(deviceId)); + hipMemAccessDesc access = {}; + access.location.type = hipMemLocationTypeDevice; + access.location.id = deviceId; + access.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &access, 1)); hipDeviceptr_t dptr_peer; - HIP_CHECK(hipMalloc(&dptr_peer, buffer_size)); + HIP_CHECK(hipMalloc(reinterpret_cast(&dptr_peer), buffer_size)); HIP_CHECK(hipMemcpyDtoD(dptr_peer, ptrA, buffer_size)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), dptr_peer, buffer_size)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); - HIP_CHECK(hipFree(dptr_peer)); + HIP_CHECK(hipFree(reinterpret_cast(dptr_peer))); } HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** @@ -701,6 +723,13 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy") { size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); + int devicecount = 0; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + if (devicecount < 2) { + HipTest::HIP_SKIP_TEST("Machine is Single GPU. Skipping Test.."); + return; + } int deviceId = 0, value = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -733,8 +762,6 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy") { A_h[idx] = idx; } HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); - int devicecount = 0; - HIP_CHECK(hipGetDeviceCount(&devicecount)); // Check Peer Access for (deviceId = 1; deviceId < devicecount; deviceId++) { std::fill(B_h.begin(), B_h.end(), initializer); @@ -763,14 +790,16 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy") { } HIP_CHECK(hipSetDevice(deviceId)); hipDeviceptr_t dptr_peer; - HIP_CHECK(hipMalloc(&dptr_peer, buffer_size)); - HIP_CHECK(hipMemcpyPeer(dptr_peer, deviceId, ptrA, 0, buffer_size)); + HIP_CHECK(hipMalloc(reinterpret_cast(&dptr_peer), buffer_size)); + HIP_CHECK(hipMemcpyPeer(reinterpret_cast(dptr_peer), deviceId, + reinterpret_cast(ptrA), 0, buffer_size)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), dptr_peer, buffer_size)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); - HIP_CHECK(hipFree(dptr_peer)); + HIP_CHECK(hipFree(reinterpret_cast(dptr_peer))); } HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** @@ -790,6 +819,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMMemCpy") { size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -834,6 +864,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMMemCpy") { HIP_CHECK(hipMemUnmap(ptrB, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrB, size_mem)); + CTX_DESTROY(); } /** @@ -853,6 +884,13 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") { size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); + CTX_CREATE(); + int devicecount = 0; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + if (devicecount < 2) { + HipTest::HIP_SKIP_TEST("Machine is Single GPU. Skipping Test.."); + return; + } int deviceId = 0, value = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -885,8 +923,6 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") { A_h[idx] = idx; } HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); - int devicecount = 0; - HIP_CHECK(hipGetDeviceCount(&devicecount)); for (deviceId = 1; deviceId < devicecount; deviceId++) { int canAccessPeer = 0; hipDevice_t device_other; @@ -918,7 +954,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") { // Allocate virtual address range hipDeviceptr_t ptrB; HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem_loc, 0, 0, 0)); - HIP_CHECK(hipMemMap(ptrB, size_mem_loc, 0, handle, 0)); + HIP_CHECK(hipMemMap(ptrB, size_mem_loc, 0, handle_loc, 0)); HIP_CHECK(hipMemRelease(handle_loc)); // Set access hipMemAccessDesc accessDesc_loc = {}; @@ -927,7 +963,8 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") { accessDesc_loc.flags = hipMemAccessFlagsProtReadWrite; // Make the address accessible to GPU 0 HIP_CHECK(hipMemSetAccess(ptrB, size_mem_loc, &accessDesc_loc, 1)); - HIP_CHECK(hipMemcpyPeer(ptrB, deviceId, ptrA, 0, buffer_size)); + HIP_CHECK(hipMemcpyPeer(reinterpret_cast(ptrB), deviceId, reinterpret_cast(ptrA), + 0, buffer_size)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrB, buffer_size)); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); HIP_CHECK(hipMemUnmap(ptrB, size_mem_loc)); @@ -935,6 +972,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") { } HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } class vmm_resize_class { @@ -1021,9 +1059,9 @@ class vmm_resize_class { if (idx == 0) { HIP_CHECK(hipMemMap(ptrVmm, vsize[idx], 0, myhandle, 0)); } else { - uint64_t uiptr = reinterpret_cast(ptrVmm); + unsigned long long uiptr = reinterpret_cast(ptrVmm); uiptr = uiptr + vsize[idx - 1]; - HIP_CHECK(hipMemMap(reinterpret_cast(uiptr), vsize[idx], 0, myhandle, 0)); + HIP_CHECK(hipMemMap(reinterpret_cast(uiptr), vsize[idx], 0, myhandle, 0)); } idx++; } @@ -1063,6 +1101,7 @@ TEST_CASE("Unit_hipMemSetAccess_GrowVMM") { size_t buffer_size = N * sizeof(int); int deviceId = 0; hipDevice_t device; + CTX_CREATE(); HIP_CHECK(hipDeviceGet(&device, deviceId)); checkVMMSupported(device); // Create VMM Object of size buffer_size @@ -1090,9 +1129,9 @@ TEST_CASE("Unit_hipMemSetAccess_GrowVMM") { } int* ptrB_h = static_cast(malloc(buffer_size_new)); REQUIRE(ptrB_h != nullptr); - uint64_t uiptr = reinterpret_cast(ptr); + unsigned long long uiptr = reinterpret_cast(ptr); uiptr = uiptr + buffer_size; - HIP_CHECK(hipMemcpyHtoD(reinterpret_cast(uiptr), ptrA_h, (buffer_size_new - buffer_size))); + HIP_CHECK(hipMemcpyHtoD(reinterpret_cast(uiptr), ptrA_h, (buffer_size_new - buffer_size))); HIP_CHECK(hipMemcpyDtoH(ptrB_h, ptr, buffer_size_new)); bool bPassed = true; for (int idx = 0; idx < Nnew; idx++) { @@ -1105,6 +1144,7 @@ TEST_CASE("Unit_hipMemSetAccess_GrowVMM") { free(ptrB_h); free(ptrA_h); resizeobj.free_vmm(); + CTX_DESTROY(); } std::atomic bTestPassed{1}; @@ -1122,6 +1162,7 @@ void test_thread(hipDevice_t device) { ptrA_h[idx] = idx; } // Copy to VMM + CTX_CREATE(); HIP_CHECK(hipMemcpyHtoD(ptr, ptrA_h, buffer_size)); int* ptrB_h = static_cast(malloc(buffer_size)); REQUIRE(ptrB_h != nullptr); @@ -1141,6 +1182,7 @@ void test_thread(hipDevice_t device) { free(ptrB_h); free(ptrA_h); vmmobj.free_vmm(); + CTX_DESTROY(); } /** @@ -1156,6 +1198,7 @@ void test_thread(hipDevice_t device) { * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemSetAccess_Multithreaded") { + CTX_CREATE(); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); @@ -1169,98 +1212,9 @@ TEST_CASE("Unit_hipMemSetAccess_Multithreaded") { T[i].join(); } REQUIRE(1 == bTestPassed.load()); + CTX_DESTROY(); } -#ifdef __linux__ - -bool test_mprocess() { - int fd[2]; - bool testResult = false; - pid_t childpid; - int testResultChild = 0; - int deviceId = 0; - constexpr int N = DATA_SIZE; - size_t buffer_size = N * sizeof(int); - // create pipe descriptors - pipe(fd); - // fork process - childpid = fork(); - if (childpid > 0) { // Parent - close(fd[1]); - hipDeviceptr_t ptr; - hipDevice_t device; - HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupportedRetVal(device); - // Create VMM Object of size buffer_size - vmm_resize_class vmmobj(&ptr, device, buffer_size); - // Inititalize Host Buffer - std::vector A_h(N), B_h(N); - for (int idx = 0; idx < N; idx++) { - A_h[idx] = idx; - } - // Copy to VMM - HIP_CHECK(hipMemcpyHtoD(ptr, A_h.data(), buffer_size)); - HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptr, buffer_size)); - bool bPassed = std::equal(B_h.begin(), B_h.end(), A_h.data()); - vmmobj.free_vmm(); - // parent will wait to read the device cnt - read(fd[0], &testResultChild, sizeof(int)); - if (testResultChild == 0) { - testResult = bPassed & false; - } else { - testResult = bPassed & true; - } - // close the read-descriptor - close(fd[0]); - // wait for child exit - wait(NULL); - } else if (!childpid) { // Child - close(fd[0]); - hipDeviceptr_t ptr; - hipDevice_t device; - - HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupportedRetVal(device); - // Create VMM Object of size buffer_size - vmm_resize_class vmmobj(&ptr, device, buffer_size); - // Inititalize Host Buffer - std::vector A_h(N), B_h(N); - for (int idx = 0; idx < N; idx++) { - A_h[idx] = idx; - } - // Copy to VMM - HIP_CHECK(hipMemcpyHtoD(ptr, A_h.data(), buffer_size)); - HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptr, buffer_size)); - int result = 0; - if (true == std::equal(B_h.begin(), B_h.end(), A_h.data())) { - result = 1; - } - vmmobj.free_vmm(); - // send the value on the write-descriptor: - write(fd[1], &result, sizeof(int)); - // close the write descriptor: - close(fd[1]); - exit(0); - } - return testResult; -} - -/** - * Test Description - * ------------------------ - * - Multiprocess test: Allocate unique virtual memory chunks from - * multiple processes. Transfer data to these chunks from host and - * execute kernel function on these data. Validate the results. - * ------------------------ - * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.1 - */ -TEST_CASE("Unit_hipMemSetAccess_MultiProc") { REQUIRE(true == test_mprocess()); } - -#endif - /** * Test Description * ------------------------ @@ -1277,6 +1231,7 @@ TEST_CASE("Unit_hipMemSetAccess_negative") { size_t buffer_size = N * sizeof(int); int deviceId = 0; hipDevice_t device; + CTX_CREATE(); HIP_CHECK(hipDeviceGet(&device, deviceId)); checkVMMSupported(device); hipMemAllocationProp prop{}; @@ -1301,65 +1256,79 @@ TEST_CASE("Unit_hipMemSetAccess_negative") { accessDesc.flags = hipMemAccessFlagsProtReadWrite; SECTION("nullptr to ptrA") { - REQUIRE(hipMemSetAccess(nullptr, size_mem, &accessDesc, 1) == hipErrorInvalidValue); + REQUIRE(hipMemSetAccess((hipDeviceptr_t) nullptr, size_mem, &accessDesc, 1) == + hipErrorInvalidValue); } SECTION("pass zero to size") { - REQUIRE(hipMemSetAccess(&ptrA, 0, &accessDesc, 1) == hipErrorInvalidValue); + REQUIRE(hipMemSetAccess(ptrA, 0, &accessDesc, 1) == hipErrorInvalidValue); } SECTION("pass a size greater than reserved size") { - REQUIRE(hipMemSetAccess(&ptrA, size_mem + 1, &accessDesc, 1) == hipErrorInvalidValue); + REQUIRE(hipMemSetAccess(ptrA, size_mem + 1, &accessDesc, 1) == hipErrorInvalidValue); } SECTION("pass a size less than reserved size") { - REQUIRE(hipMemSetAccess(&ptrA, size_mem - 1, &accessDesc, 1) == hipErrorInvalidValue); +#if HT_AMD + REQUIRE(hipMemSetAccess(ptrA, size_mem - 1, &accessDesc, 1) == hipSuccess); +#else + REQUIRE(hipMemSetAccess(ptrA, size_mem - 1, &accessDesc, 1) == hipErrorInvalidValue); +#endif } SECTION("invalid location type") { accessDesc.location.type = hipMemLocationTypeInvalid; - REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); +#if HT_AMD + REQUIRE(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1) == hipSuccess); +#else + REQUIRE(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); +#endif } SECTION("invalid id") { accessDesc.location.id = -1; - REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); + REQUIRE(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); } SECTION("pass location id as > highest device number") { int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); accessDesc.location.id = numDevices; // set to non existing device - REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); + REQUIRE(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); } SECTION("invalid flag") { accessDesc.flags = static_cast(-1); - REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); +#if HT_AMD + REQUIRE(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1) == hipSuccess); +#else + REQUIRE(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); +#endif } SECTION(" pass zero to count") { - REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 0) == hipErrorInvalidValue); + REQUIRE(hipMemSetAccess(ptrA, size_mem, &accessDesc, 0) == hipErrorInvalidValue); } SECTION("pass desc as nullptr") { - REQUIRE(hipMemSetAccess(&ptrA, size_mem, nullptr, 1) == hipErrorInvalidValue); + REQUIRE(hipMemSetAccess(ptrA, size_mem, nullptr, 1) == hipErrorInvalidValue); } SECTION("uninitialized virtual memory") { hipDeviceptr_t ptrB; HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem, 0, 0, 0)); - REQUIRE(hipMemSetAccess(&ptrB, size_mem, &accessDesc, 1) == hipErrorInvalidValue); + REQUIRE(hipMemSetAccess(ptrB, size_mem, &accessDesc, 1) == hipErrorInvalidValue); HIP_CHECK(hipMemAddressFree(ptrB, size_mem)); } HIP_CHECK(hipMemUnmap(ptrA, size_mem)); SECTION("unmapped virtual memory") { - REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); + REQUIRE(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); } HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); HIP_CHECK(hipMemRelease(handle)); + CTX_DESTROY(); } /** diff --git a/catch/unit/virtualMemoryManagement/hipMemUnmap.cc b/catch/unit/virtualMemoryManagement/hipMemUnmap.cc index 67561fc83c..7acdc6c4a5 100644 --- a/catch/unit/virtualMemoryManagement/hipMemUnmap.cc +++ b/catch/unit/virtualMemoryManagement/hipMemUnmap.cc @@ -51,6 +51,7 @@ TEST_CASE("Unit_hipMemUnmap_negative") { int deviceId = 0; hipDevice_t device; + CTX_CREATE(); HIP_CHECK(hipDeviceGet(&device, deviceId)); checkVMMSupported(device); @@ -72,7 +73,9 @@ TEST_CASE("Unit_hipMemUnmap_negative") { HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); - SECTION("nullptr to ptrA") { REQUIRE(hipMemUnmap(nullptr, size_mem) == hipErrorInvalidValue); } + SECTION("nullptr to ptrA") { + REQUIRE(hipMemUnmap((hipDeviceptr_t) nullptr, size_mem) == hipErrorInvalidValue); + } SECTION("pass zero to size") { REQUIRE(hipMemUnmap(ptrA, 0) == hipErrorInvalidValue); } @@ -85,6 +88,7 @@ TEST_CASE("Unit_hipMemUnmap_negative") { HIP_CHECK(hipMemRelease(handle)); HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + CTX_DESTROY(); } /** diff --git a/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc b/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc index c2258f057d..17303eb035 100644 --- a/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc +++ b/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc @@ -26,11 +26,6 @@ THE SOFTWARE. */ #include -#include -#include -#include -#include -#include #include #include @@ -39,8 +34,8 @@ THE SOFTWARE. /* This testcase verifies HIP Mem VMM API basic scenario - supported on all devices */ - TEST_CASE("Unit_hipMemVmm_Basic") { + CTX_CREATE(); int vmm = 0; HIP_CHECK(hipDeviceGetAttribute(&vmm, hipDeviceAttributeVirtualMemoryManagementSupported, 0)); INFO("hipDeviceAttributeVirtualMemoryManagementSupported: " << vmm); @@ -54,7 +49,7 @@ TEST_CASE("Unit_hipMemVmm_Basic") { size_t granularity = 0; - hipMemAllocationProp memAllocationProp; + hipMemAllocationProp memAllocationProp{}; memAllocationProp.type = hipMemAllocationTypePinned; memAllocationProp.location.id = 0; memAllocationProp.location.type = hipMemLocationTypeDevice; @@ -62,14 +57,15 @@ TEST_CASE("Unit_hipMemVmm_Basic") { HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &memAllocationProp, hipMemAllocationGranularityRecommended)); - size_t size = 4 * 1024; + size_t size = granularity; void* reservedAddress{nullptr}; - HIP_CHECK(hipMemAddressReserve(&reservedAddress, size, granularity, nullptr, 0)); + HIP_CHECK(hipMemAddressReserve(reinterpret_cast(&reservedAddress), size, 0, + (hipDeviceptr_t) nullptr, 0)); - hipMemGenericAllocationHandle_t gaHandle{nullptr}; + hipMemGenericAllocationHandle_t gaHandle; HIP_CHECK(hipMemCreate(&gaHandle, size, &memAllocationProp, 0)); - HIP_CHECK(hipMemMap(reservedAddress, size, 0, gaHandle, 0)); + HIP_CHECK(hipMemMap(reinterpret_cast(reservedAddress), size, 0, gaHandle, 0)); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); @@ -80,7 +76,7 @@ TEST_CASE("Unit_hipMemVmm_Basic") { std::vector values(size); const char value = 1; - HIP_CHECK(hipMemSetAccess(reservedAddress, size, &desc, 1)); + HIP_CHECK(hipMemSetAccess(reinterpret_cast(reservedAddress), size, &desc, 1)); HIP_CHECK(hipMemset(reservedAddress, value, size)); HIP_CHECK(hipMemcpy(&values[0], reservedAddress, size, hipMemcpyDeviceToHost)); @@ -88,8 +84,9 @@ TEST_CASE("Unit_hipMemVmm_Basic") { REQUIRE(values[i] == value); } - HIP_CHECK(hipMemUnmap(reservedAddress, size)); + HIP_CHECK(hipMemUnmap(reinterpret_cast(reservedAddress), size)); HIP_CHECK(hipMemRelease(gaHandle)); - HIP_CHECK(hipMemAddressFree(reservedAddress, size)); + HIP_CHECK(hipMemAddressFree(reinterpret_cast(reservedAddress), size)); + CTX_DESTROY(); } diff --git a/catch/unit/virtualMemoryManagement/hip_vmm_common.hh b/catch/unit/virtualMemoryManagement/hip_vmm_common.hh index 33f38ecfd2..ff754ca0d3 100644 --- a/catch/unit/virtualMemoryManagement/hip_vmm_common.hh +++ b/catch/unit/virtualMemoryManagement/hip_vmm_common.hh @@ -46,17 +46,6 @@ THE SOFTWARE. } \ } -#define checkVMMSupportedRetVal(device) \ - { \ - int value = 0; \ - hipDeviceAttribute_t attr = hipDeviceAttributeVirtualMemoryManagementSupported; \ - HIP_CHECK(hipDeviceGetAttribute(&value, attr, device)); \ - if (value == 0) { \ - HipTest::HIP_SKIP_TEST("Machine does not support VMM. Skipping Test.."); \ - return true; \ - } \ - } - #ifdef __linux__ #define checkSysCallErrors(result) \ if (result == -1) { \