diff --git a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json index ada918a267..453c984199 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json @@ -92,6 +92,8 @@ "Unit_hipMemcpy_Positive_Synchronization_Behavior", "=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/210 ===", "Unit_StaticAssert_Positive_Basic", - "Unit_StaticAssert_Negative_Basic" + "Unit_StaticAssert_Negative_Basic", + "Unit_hipMemImportFromShareableHandle_Positive_MultiProc", + "Unit_hipMemMapArrayAsync_Positive_Basic" ] } diff --git a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json index 3e7785a3e8..5d118b16e6 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json +++ b/projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json @@ -44,6 +44,7 @@ "Performance_hipMemsetD32", "Performance_hipMemsetD32Async", "Unit_hipMemcpyParam2D_Positive_Synchronization_Behavior", - "Unit_hipMemcpy_Positive_Synchronization_Behavior" + "Unit_hipMemcpy_Positive_Synchronization_Behavior", + "Unit_hipMemMapArrayAsync_Positive_Basic" ] } diff --git a/projects/hip-tests/catch/unit/memory/hipMallocManagedCommon.hh b/projects/hip-tests/catch/include/hipMallocManagedCommon.hh similarity index 100% rename from projects/hip-tests/catch/unit/memory/hipMallocManagedCommon.hh rename to projects/hip-tests/catch/include/hipMallocManagedCommon.hh diff --git a/projects/hip-tests/catch/include/hip_test_defgroups.hh b/projects/hip-tests/catch/include/hip_test_defgroups.hh index 8191c8b96f..2a8413d79f 100644 --- a/projects/hip-tests/catch/include/hip_test_defgroups.hh +++ b/projects/hip-tests/catch/include/hip_test_defgroups.hh @@ -103,16 +103,9 @@ THE SOFTWARE. */ /** -* @defgroup KernelTest Kernel Functions Management -* @{ -* This section describes the various kernel functions invocation. -* @} -*/ - -/** - * @defgroup AtomicsTest Device Atomics + * @defgroup KernelTest Kernel Functions Management * @{ - * This section describes tests for the Device Atomic APIs. + * This section describes the various kernel functions invocation. * @} */ @@ -140,7 +133,8 @@ THE SOFTWARE. /** * @defgroup PeerToPeerTest PeerToPeer Device Memory Access * @{ - * This section describes tests for the PeerToPeer device memory access functions of HIP runtime API. + * This section describes tests for the PeerToPeer device memory access functions of HIP runtime + * API. * @warning PeerToPeer support is experimental. * @} */ @@ -200,124 +194,6 @@ THE SOFTWARE. * @} */ -/** - * @defgroup AtomicsTest Device Atomics - * @{ - * This section describes tests for the Device Atomic APIs. - */ - -/** - * @addtogroup atomicAdd atomicAdd - * @{ - * @ingroup AtomicsTest - */ - -/** - * Test Description - * ------------------------ - * - Compiles atomicAdd with invalid parameters. - * - Compiles the source with specialized Python tool. - * -# Utilizes sub-process to invoke compilation of faulty source. - * -# Performs post-processing of compiler output and counts errors. - * Test source - * ------------------------ - * - unit/atomics/CMakeLists.txt - * Test requirements - * ------------------------ - * - HIP_VERSION >= 5.2 - */ -TEST_CASE("Unit_atomicAdd_Negative_Parameters") {} -/** - * End doxygen group atomicAdd. - * @} - */ - -/** - * @addtogroup atomicSub atomicSub - * @{ - * @ingroup AtomicsTest - */ - -/** - * Test Description - * ------------------------ - * - Compiles atomicSub with invalid parameters. - * - Compiles the source with specialized Python tool. - * -# Utilizes sub-process to invoke compilation of faulty source. - * -# Performs post-processing of compiler output and counts errors. - * Test source - * ------------------------ - * - unit/atomics/CMakeLists.txt - * Test requirements - * ------------------------ - * - HIP_VERSION >= 5.2 - */ -TEST_CASE("Unit_atomicSub_Negative_Parameters") {} -/** - * End doxygen group atomicSub. - * @} - */ - -/** - * @addtogroup atomicInc atomicInc - * @{ - * @ingroup AtomicsTest - */ - -/** - * Test Description - * ------------------------ - * - Compiles atomicInc with invalid parameters. - * - Compiles the source with specialized Python tool. - * -# Utilizes sub-process to invoke compilation of faulty source. - * -# Performs post-processing of compiler output and counts errors. - * Test source - * ------------------------ - * - unit/atomics/CMakeLists.txt - * Test requirements - * ------------------------ - * - HIP_VERSION >= 5.2 - */ -TEST_CASE("Unit_atomicInc_Negative_Parameters") {} -/** - * End doxygen group atomicInc. - * @} - */ - -/** - * @addtogroup atomicDec atomicDec - * @{ - * @ingroup AtomicsTest - */ - -/** - * Test Description - * ------------------------ - * - Compiles atomicDec with invalid parameters. - * - Compiles the source with specialized Python tool. - * -# Utilizes sub-process to invoke compilation of faulty source. - * -# Performs post-processing of compiler output and counts errors. - * Test source - * ------------------------ - * - unit/atomics/CMakeLists.txt - * Test requirements - * ------------------------ - * - HIP_VERSION >= 5.2 - */ -TEST_CASE("Unit_atomicDec_Negative_Parameters") {} -/** - * End doxygen group atomicDec. - * @} - */ - -/** - * End doxygen group AtomicsTest. - * @defgroup MathTest Math Device Functions - * @{ - * This section describes tests for device math functions of HIP runtime API. - * @} - */ - /** * @defgroup MathTest Math Device Functions * @{ @@ -347,62 +223,8 @@ TEST_CASE("Unit_atomicDec_Negative_Parameters") {} */ /** - * @defgroup DeviceLanguageTest Device Language + * @defgroup VirtualMemoryManagementTest Virtual Memory Management APIs * @{ - * This section describes tests for the Device Language API. - */ - -/** - * @addtogroup launch_bounds launch_bounds - * @{ - * @ingroup DeviceLanguageTest - */ - -/** - * Test Description - * ------------------------ - * - Validates handling of invalid arguments: - * -# Compiles kernels that are not created appropriately: - * - Maximum number of threads is 0 - * - Maximum number of threads is not integer value - * - Mimimum number of warps is not integer value - * -# Expected output: compiler error - * Test source - * ------------------------ - * - unit/launch_bounds/CMakeLists.txt - * Test requirements - * ------------------------ - * - HIP_VERSION >= 5.2 - */ -TEST_CASE("Unit_Kernel_Launch_bounds_Negative_Parameters_CompilerError") {} - -/** - * Test Description - * ------------------------ - * - Validates handling of invalid arguments: - * -# Compiles kernels that are not created appropriately: - * - Maximum number of threads is negative - * - Mimimum number of warps is negative - * - Validates handling of invalid arguments: - * -# Expected output: parse error - * Test source - * ------------------------ - * - unit/launch_bounds/CMakeLists.txt - * Test requirements - * ------------------------ - * - HIP_VERSION >= 5.2 - */ -TEST_CASE("Unit_Kernel_Launch_bounds_Negative_Parameters_ParseError") {} - -/** - * End doxygen group launch_bounds. + * This section describes the virtual memory management types & functions of HIP runtime API. * @} */ - -/** - * End doxygen group DeviceLanguageTest. - * @} - * @defgroup VectorTypeTest Vector types - * @{ - * This section describes tests for the Vector type functions and operators. - */ diff --git a/projects/hip-tests/catch/unit/CMakeLists.txt b/projects/hip-tests/catch/unit/CMakeLists.txt index 59f8f6ad47..6b63292c91 100644 --- a/projects/hip-tests/catch/unit/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/CMakeLists.txt @@ -52,6 +52,7 @@ add_subdirectory(p2p) add_subdirectory(gcc) add_subdirectory(syncthreads) add_subdirectory(threadfence) +add_subdirectory(virtualMemoryManagement) if(HIP_PLATFORM STREQUAL "amd") add_subdirectory(callback) diff --git a/projects/hip-tests/catch/unit/memory/CMakeLists.txt b/projects/hip-tests/catch/unit/memory/CMakeLists.txt index 4ef22c8013..47ae8aa9ed 100644 --- a/projects/hip-tests/catch/unit/memory/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/memory/CMakeLists.txt @@ -94,7 +94,8 @@ if(HIP_PLATFORM MATCHES "amd") hipMemAddressFree.cc hipMemAddressReserve.cc hipMemRelease.cc - hipMemGetAllocationPropertiesFromHandle.cc) + hipMemGetAllocationPropertiesFromHandle.cc + hipArray.cc) else() set(TEST_SRC ${TEST_SRC} hipGetSymbolSizeAddress.cc) endif() diff --git a/projects/hip-tests/catch/unit/memory/hip_vmm_common.hh b/projects/hip-tests/catch/unit/memory/hip_vmm_common.hh deleted file mode 100644 index 24ecb6408a..0000000000 --- a/projects/hip-tests/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/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt b/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt new file mode 100644 index 0000000000..f540fdd25b --- /dev/null +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt @@ -0,0 +1,55 @@ +# 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. + +set(TEST_SRC + hipMemGetAllocationGranularity.cc + hipMemRetainAllocationHandle.cc) + +if(HIP_PLATFORM MATCHES "nvidia") +set(TEST_SRC + ${TEST_SRC} + hipMemMapArrayAsync.cc) + +if(UNIX) # Disabled on AMD due to defect EXSWHTEC-375 + set(TEST_SRC + ${TEST_SRC} + hipMemExportToShareableHandle.cc + hipMemImportFromShareableHandle.cc) +endif() + +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) +endif() + +hip_add_exe_to_target(NAME VirtualMemoryManagementTest + TEST_SRC ${TEST_SRC} + TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC}) \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/memory/hipMemAddressFree.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc similarity index 84% rename from projects/hip-tests/catch/unit/memory/hipMemAddressFree.cc rename to projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc index 6890a26d8c..6171f830ac 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemAddressFree.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc @@ -23,12 +23,13 @@ THE SOFTWARE. /** * @addtogroup hipMemAddressFree hipMemAddressFree * @{ - * @ingroup MemoryTest + * @ingroup VirtualMemoryManagementTest * `hipError_t hipMemAddressFree (void* devPtr, size_t size)` - * Frees an address range reservation made via hipMemAddressReserve. */ #include + #include "hip_vmm_common.hh" #define DATA_SIZE (1 << 13) @@ -38,7 +39,7 @@ THE SOFTWARE. * ------------------------ * - Negative Tests * ------------------------ - * - catch\unit\memory\hipMemAddressFree.cc + * - unit/virtualMemoryManagement/hipMemAddressFree.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -50,16 +51,15 @@ TEST_CASE("Unit_hipMemAddressFree_negative") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate virtual address range hipDeviceptr_t ptrA; HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); @@ -68,9 +68,7 @@ TEST_CASE("Unit_hipMemAddressFree_negative") { REQUIRE(hipMemAddressFree(nullptr, size_mem) == hipErrorInvalidValue); } - SECTION("pass zero to size") { - REQUIRE(hipMemAddressFree(ptrA, 0) == hipErrorInvalidValue); - } + SECTION("pass zero to size") { REQUIRE(hipMemAddressFree(ptrA, 0) == hipErrorInvalidValue); } HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); } diff --git a/projects/hip-tests/catch/unit/memory/hipMemAddressReserve.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc similarity index 83% rename from projects/hip-tests/catch/unit/memory/hipMemAddressReserve.cc rename to projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc index c6a76fc56c..c19f8ada43 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemAddressReserve.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc @@ -23,7 +23,7 @@ THE SOFTWARE. /** * @addtogroup hipMemAddressReserve hipMemAddressReserve * @{ - * @ingroup MemoryTest + * @ingroup VirtualMemoryManagementTest * `hipError_t hipMemAddressReserve (void** ptr, * size_t size, * size_t alignment, @@ -33,6 +33,7 @@ THE SOFTWARE. */ #include + #include "hip_vmm_common.hh" #define DATA_SIZE (1 << 13) @@ -43,7 +44,7 @@ THE SOFTWARE. * - Verify if reserved address returned by hipMemAddressReserve * for different alignment values are correctly aligned. * ------------------------ - * - catch\unit\memory\hipMemAddressReserve.cc + * - unit/virtualMemoryManagement/hipMemAddressReserve.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -56,16 +57,15 @@ TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate virtual address range hipDeviceptr_t ptrA; size_t alignmnt = 1; @@ -106,7 +106,7 @@ TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") { * ------------------------ * - Negative Tests * ------------------------ - * - catch\unit\memory\hipMemAddressReserve.cc + * - unit/virtualMemoryManagement/hipMemAddressReserve.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -118,35 +118,33 @@ TEST_CASE("Unit_hipMemAddressReserve_Negative") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate virtual address range hipDeviceptr_t ptrA; + SECTION("Nullptr to ptr") { - REQUIRE(hipMemAddressReserve(nullptr, size_mem, 0, 0, 0) == - hipErrorInvalidValue); + REQUIRE(hipMemAddressReserve(nullptr, size_mem, 0, 0, 0) == hipErrorInvalidValue); } SECTION("pass size as 0") { - REQUIRE(hipMemAddressReserve(&ptrA, 0, 0, 0, 0) == - hipErrorMemoryAllocation); + REQUIRE(hipMemAddressReserve(&ptrA, 0, 0, 0, 0) == hipErrorMemoryAllocation); } + #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) == hipErrorMemoryAllocation); } #endif + SECTION("pass size as non multiple of host page size") { - REQUIRE(hipMemAddressReserve(&ptrA, (size_mem - 1), 0, 0, 0) == - hipErrorMemoryAllocation); + REQUIRE(hipMemAddressReserve(&ptrA, (size_mem - 1), 0, 0, 0) == hipErrorMemoryAllocation); } } diff --git a/projects/hip-tests/catch/unit/memory/hipMemCreate.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc similarity index 84% rename from projects/hip-tests/catch/unit/memory/hipMemCreate.cc rename to projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc index 83f6ff7c01..4cc52786ad 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemCreate.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc @@ -19,20 +19,23 @@ 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. */ + /** * @addtogroup hipMemCreate hipMemCreate * @{ - * @ingroup MemoryTest + * @ingroup VirtualMemoryManagementTest * `hipError_t hipMemCreate (hipMemGenericAllocationHandle_t* handle, * size_t size, * const hipMemAllocationProp* prop, * unsigned long long flags)` - * Creates a memory allocation described by the properties and size. */ -#include "hip_vmm_common.hh" + #include #include +#include "hip_vmm_common.hh" + #define THREADS_PER_BLOCK 512 #define NUM_OF_BUFFERS 3 #define DATA_SIZE (1 << 13) @@ -52,7 +55,7 @@ static __global__ void square_kernel(int* Buff) { * - Allocate physical memories for different multiples of * granularity and deallocate them. * ------------------------ - * - catch\unit\memory\hipMemCreate.cc + * - unit/virtualMemoryManagement/hipMemCreate.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -62,19 +65,18 @@ TEST_CASE("Unit_hipMemCreate_BasicAllocateDeAlloc_MultGranularity") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); hipMemGenericAllocationHandle_t handle; // Allocate physical memory for (int mul = 1; mul < 64; mul++) { - HIP_CHECK(hipMemCreate(&handle, granularity*mul, &prop, 0)); + HIP_CHECK(hipMemCreate(&handle, granularity * mul, &prop, 0)); HIP_CHECK(hipMemRelease(handle)); } } @@ -87,7 +89,7 @@ TEST_CASE("Unit_hipMemCreate_BasicAllocateDeAlloc_MultGranularity") { * and back to host. Verify the result. Release handle at end after * unmapping VMM range. * ------------------------ - * - catch\unit\memory\hipMemCreate.cc + * - unit/virtualMemoryManagement/hipMemCreate.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -99,17 +101,15 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); @@ -146,7 +146,7 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap") { * to VMM memory and back to host. Verify the result. Release * handle before the VMM range is used. * ------------------------ - * - catch\unit\memory\hipMemCreate.cc + * - unit/virtualMemoryManagement/hipMemCreate.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -158,17 +158,15 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); @@ -205,12 +203,11 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") { * to device, launch kernel to square the data, copy data back * to host. Verify the result. * ------------------------ - * - catch\unit\memory\hipMemCreate.cc + * - unit/virtualMemoryManagement/hipMemCreate.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 */ - TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") { size_t granularity = 0; constexpr int N = DATA_SIZE; @@ -218,17 +215,15 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); @@ -248,12 +243,12 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") { // Initialize with data for (size_t idx = 0; idx < N; idx++) { A_h[idx] = idx; - C_h[idx] = 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)); + 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()); REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); @@ -269,7 +264,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") { * device permission, copy data from host to device, launch kernel * to square the data, copy data back to host. Verify the result. * ------------------------ - * - catch\unit\memory\hipMemCreate.cc + * - unit/virtualMemoryManagement/hipMemCreate.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -282,17 +277,15 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle[NUM_OF_BUFFERS]; // Allocate 3 physical memory chunks for (int count = 0; count < numOfBuffers; count++) { @@ -304,8 +297,7 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") { for (int idx = 0; idx < numOfBuffers; idx++) { uint64_t 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,16 +307,16 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") { // Make the address accessible to GPU 0 HIP_CHECK(hipMemSetAccess(ptrA, (numOfBuffers * size_mem), &accessDesc, 1)); std::vector A_h(numOfBuffers * size_mem), B_h(numOfBuffers * size_mem), - C_h(numOfBuffers * size_mem); + C_h(numOfBuffers * size_mem); // Fill Data for (size_t idx = 0; idx < (numOfBuffers * N); idx++) { A_h[idx] = idx; - C_h[idx] = idx*idx; + C_h[idx] = idx * idx; } 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 / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0, + static_cast(ptrA)); HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, numOfBuffers * buffer_size)); HIP_CHECK(hipDeviceSynchronize()); // Validate Results @@ -344,7 +336,7 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") { * to the VMM address range. Memset the VMM address range with initial * value. Validate. * ------------------------ - * - catch\unit\memory\hipMemCreate.cc + * - unit/virtualMemoryManagement/hipMemCreate.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -357,17 +349,15 @@ TEST_CASE("Unit_hipMemCreate_ChkWithMemset") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); @@ -398,7 +388,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithMemset") { * ------------------------ * - Negative Tests * ------------------------ - * - catch\unit\memory\hipMemCreate.cc + * - unit/virtualMemoryManagement/hipMemCreate.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -408,58 +398,48 @@ TEST_CASE("Unit_hipMemCreate_Negative") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemGenericAllocationHandle_t handle; hipMemAllocationProp prop = {}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Device HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); SECTION("Nullptr to handle") { - REQUIRE(hipMemCreate(nullptr, granularity, &prop, 0) == - hipErrorInvalidValue); + REQUIRE(hipMemCreate(nullptr, granularity, &prop, 0) == hipErrorInvalidValue); } SECTION("Nullptr to prop") { - REQUIRE(hipMemCreate(&handle, granularity, nullptr, 0) == - hipErrorInvalidValue); + REQUIRE(hipMemCreate(&handle, granularity, nullptr, 0) == hipErrorInvalidValue); } - SECTION("pass size as 0") { - REQUIRE(hipMemCreate(&handle, 0, &prop, 0) == hipErrorInvalidValue); - } + SECTION("pass size as 0") { REQUIRE(hipMemCreate(&handle, 0, &prop, 0) == hipErrorInvalidValue); } SECTION("Pass prop type as invalid") { prop.type = hipMemAllocationTypeInvalid; - REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == - hipErrorInvalidValue); + REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidValue); } SECTION("pass location as invalid") { prop.location.type = hipMemLocationTypeInvalid; - REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == - hipErrorInvalidValue); + REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidValue); } SECTION("non multiple of granularity") { - REQUIRE(hipMemCreate(&handle, (granularity - 1), &prop, 0) == - hipErrorInvalidValue); + REQUIRE(hipMemCreate(&handle, (granularity - 1), &prop, 0) == hipErrorInvalidValue); } SECTION("pass location id as -1") { prop.location.id = -1; // set to non existing device - REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == - hipErrorInvalidValue); + REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidValue); } SECTION("pass location id as > highest device number") { int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); prop.location.id = numDevices; // set to non existing device - REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == - hipErrorInvalidValue); + REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidValue); } } diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc new file mode 100644 index 0000000000..d5c4b5394e --- /dev/null +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc @@ -0,0 +1,145 @@ +/* +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. +*/ + +/** + * @addtogroup hipMemExportToShareableHandle hipMemExportToShareableHandle + * @{ + * @ingroup VirtualMemoryManagementTest + * `hipError_t hipMemExportToShareableHandle(void *shareableHandle, + * hipMemGenericAllocationHandle_t handle, + * hipMemAllocationHandleType handleType, + * unsigned long long flags)` - + * Exports an allocation to a requested shareable handle type. + */ + +#include + +#include "hip_vmm_common.hh" + +/** + * Test Description + * ------------------------ + * - Basic sanity test. + * ------------------------ + * - unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemExportToShareableHandle_Positive_Basic") { + 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(shareable_handle != nullptr); + + HIP_CHECK(hipMemRelease(handle)); +} + +/** + * Test Description + * ------------------------ + * - Negative parameters test. + * ------------------------ + * - unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemExportToShareableHandle_Negative_Parameters") { + 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; + + SECTION("shareableHandle == nullptr") { + HIP_CHECK_ERROR( + hipMemExportToShareableHandle(nullptr, handle, hipMemHandleTypePosixFileDescriptor, 0), + hipErrorInvalidValue); + } + +#if HT_AMD + SECTION("handle == nullptr") { + HIP_CHECK_ERROR(hipMemExportToShareableHandle(&shareable_handle, nullptr, + hipMemHandleTypePosixFileDescriptor, 0), + hipErrorInvalidValue); + } +#endif + + SECTION("invalid handleType") { + HIP_CHECK_ERROR( + hipMemExportToShareableHandle(&shareable_handle, handle, hipMemHandleTypeWin32, 0), + hipErrorInvalidValue); + } + + SECTION("non-zero flags") { + HIP_CHECK_ERROR(hipMemExportToShareableHandle(&shareable_handle, handle, + hipMemHandleTypePosixFileDescriptor, 1), + hipErrorInvalidValue); + } + + HIP_CHECK(hipMemRelease(handle)); + +#if HT_AMD // segfaults on NVIDIA + SECTION("released handle") { + HIP_CHECK_ERROR(hipMemExportToShareableHandle(&shareable_handle, handle, + hipMemHandleTypePosixFileDescriptor, 0), + hipErrorInvalidValue); + } +#endif +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/memory/hipMemGetAllocationGranularity.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc similarity index 73% rename from projects/hip-tests/catch/unit/memory/hipMemGetAllocationGranularity.cc rename to projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc index 5f5821de48..0f45a53810 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemGetAllocationGranularity.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc @@ -23,7 +23,7 @@ THE SOFTWARE. /** * @addtogroup hipMemGetAllocationGranularity hipMemGetAllocationGranularity * @{ - * @ingroup MemoryTest + * @ingroup VirtualMemoryManagementTest * `hipError_t hipMemGetAllocationGranularity (size_t* granularity, * const hipMemAllocationProp* prop, * hipMemAllocationGranularity_flags option)` - @@ -33,14 +33,13 @@ THE SOFTWARE. #include #include #include + #include "hip_vmm_common.hh" /** local function to invoke hipMemGetAllocationGranularity. */ -void getGranularity(size_t *granularity, - hipMemAllocationGranularity_flags option, - int device) { +void getGranularity(size_t* granularity, hipMemAllocationGranularity_flags option, int device) { hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; @@ -54,16 +53,17 @@ void getGranularity(size_t *granularity, * - Functional Test to get granularity size for * hipMemAllocationGranularityMinimum option. * ------------------------ - * - catch\unit\memory\hipMemGetAllocationGranularity.cc + * - unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemGetAllocationGranularity_MinGranularity") { + HIP_CHECK(hipFree(0)); size_t granularity = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); - checkVMMSupported(device) + checkVMMSupported(device); getGranularity(&granularity, hipMemAllocationGranularityMinimum, 0); REQUIRE(granularity > 0); } @@ -74,16 +74,17 @@ TEST_CASE("Unit_hipMemGetAllocationGranularity_MinGranularity") { * - Functional Test to get granularity size for * hipMemAllocationGranularityRecommended option. * ------------------------ - * - catch\unit\memory\hipMemGetAllocationGranularity.cc + * - unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemGetAllocationGranularity_RecommendedGranularity") { + HIP_CHECK(hipFree(0)); size_t granularity = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); - checkVMMSupported(device) + checkVMMSupported(device); getGranularity(&granularity, hipMemAllocationGranularityRecommended, 0); REQUIRE(granularity > 0); } @@ -94,21 +95,21 @@ TEST_CASE("Unit_hipMemGetAllocationGranularity_RecommendedGranularity") { * - Functional Test to get granularity size for * hipMemAllocationGranularityMinimum option for all GPUs. * ------------------------ - * - catch\unit\memory\hipMemGetAllocationGranularity.cc + * - unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemGetAllocationGranularity_AllGPUs") { + HIP_CHECK(hipFree(0)); int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); for (int dev = 0; dev < numDevices; dev++) { size_t granularity = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, dev)); - checkVMMSupported(device) - getGranularity(&granularity, hipMemAllocationGranularityRecommended, - dev); + checkVMMSupported(device); + getGranularity(&granularity, hipMemAllocationGranularityRecommended, dev); REQUIRE(granularity > 0); } } @@ -118,16 +119,17 @@ TEST_CASE("Unit_hipMemGetAllocationGranularity_AllGPUs") { * ------------------------ * - Negative Tests * ------------------------ - * - catch\unit\memory\hipMemGetAllocationGranularity.cc + * - unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemGetAllocationGranularity_NegativeTests") { + HIP_CHECK(hipFree(0)); size_t granularity = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; @@ -135,45 +137,48 @@ TEST_CASE("Unit_hipMemGetAllocationGranularity_NegativeTests") { SECTION("Granularity is nullptr") { REQUIRE(hipErrorInvalidValue == - hipMemGetAllocationGranularity(nullptr, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(nullptr, &prop, hipMemAllocationGranularityMinimum)); } +#if HT_AMD // segfaults on NVIDIA SECTION("Prop is nullptr") { - REQUIRE(hipErrorInvalidValue == - hipMemGetAllocationGranularity(&granularity, nullptr, - hipMemAllocationGranularityMinimum)); + REQUIRE( + hipErrorInvalidValue == + hipMemGetAllocationGranularity(&granularity, nullptr, hipMemAllocationGranularityMinimum)); } +#endif #if HT_NVIDIA SECTION("flag is invalid") { REQUIRE(hipErrorInvalidValue == - hipMemGetAllocationGranularity(&granularity, &prop, - (hipMemAllocationGranularity_flags)0xff)); + hipMemGetAllocationGranularity(&granularity, &prop, + (hipMemAllocationGranularity_flags)0xff)); } #endif +#if HT_AMD // succeeds on NVIDIA SECTION("device id > highest device id") { int numDevices = 0; HIP_CHECK(hipGetDeviceCount(&numDevices)); prop.location.id = numDevices; // set to non existing device - REQUIRE(hipErrorInvalidValue == - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + REQUIRE( + hipErrorInvalidValue == + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); } SECTION("device id < lowest device id") { prop.location.id = -1; // set to non existing device - REQUIRE(hipErrorInvalidValue == - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + REQUIRE( + hipErrorInvalidValue == + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); } SECTION("allocation type as invalid") { prop.type = hipMemAllocationTypeInvalid; - REQUIRE(hipErrorInvalidValue == - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + REQUIRE( + hipErrorInvalidValue == + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); } SECTION("location type as invalid") { prop.location.type = hipMemLocationTypeInvalid; - REQUIRE(hipErrorInvalidValue == - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + REQUIRE( + hipErrorInvalidValue == + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); } +#endif } diff --git a/projects/hip-tests/catch/unit/memory/hipMemGetAllocationPropertiesFromHandle.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc similarity index 84% rename from projects/hip-tests/catch/unit/memory/hipMemGetAllocationPropertiesFromHandle.cc rename to projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc index 5c348a85f8..cc3a8dc519 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemGetAllocationPropertiesFromHandle.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc @@ -20,26 +20,27 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ -#include -#include "hip_vmm_common.hh" - -#define DATA_SIZE (1 << 13) - /** * @addtogroup hipMemGetAllocationPropertiesFromHandle hipMemGetAllocationPropertiesFromHandle * @{ - * @ingroup MemoryTest + * @ingroup VirtualMemoryManagementTest * `hipError_t hipMemGetAllocationPropertiesFromHandle(hipMemAllocationProp* prop, * hipMemGenericAllocationHandle_t handle)` - * Retrieve the property structure of the given handle. */ +#include + +#include "hip_vmm_common.hh" + +#define DATA_SIZE (1 << 13) + /** * Test Description * ------------------------ * - Functional test to verify the values of hipMemAllocationProp properties. * ------------------------ - * - catch\unit\memory\hipMemGetAllocationPropertiesFromHandle.cc + * - unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -47,7 +48,7 @@ THE SOFTWARE. TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") { hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemGenericAllocationHandle_t handle; hipMemAllocationProp prop = {}; prop.type = hipMemAllocationTypePinned; @@ -58,11 +59,10 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") { size_t granularity = 0; int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); - HIP_CHECK(hipMemGetAllocationGranularity - (&granularity, &prop, hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t mem_size = ((granularity + buffer_size - 1) / granularity) - * granularity; + size_t mem_size = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, mem_size, &prop, 0)); // verify properties has been retrived from handle @@ -78,7 +78,7 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") { * ------------------------ * - Negative Tests. * ------------------------ - * - catch\unit\memory\hipMemGetAllocationPropertiesFromHandle.cc + * - unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -86,7 +86,7 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") { TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_Negative") { hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, 0)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemGenericAllocationHandle_t handle; hipMemAllocationProp prop = {}; prop.type = hipMemAllocationTypePinned; @@ -97,23 +97,21 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_Negative") { size_t granularity = 0; int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); - HIP_CHECK(hipMemGetAllocationGranularity - (&granularity, &prop, hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t mem_size = ((granularity + buffer_size - 1) / granularity) - * granularity; + size_t mem_size = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, mem_size, &prop, 0)); SECTION("Nullptr as prop") { - REQUIRE(hipMemGetAllocationPropertiesFromHandle(nullptr, handle) - == hipErrorInvalidValue); + REQUIRE(hipMemGetAllocationPropertiesFromHandle(nullptr, handle) == hipErrorInvalidValue); } SECTION("null handle") { prop.location.type = hipMemLocationTypeInvalid; - REQUIRE(hipMemGetAllocationPropertiesFromHandle(&prop_temp, nullptr) - == hipErrorInvalidValue); + REQUIRE(hipMemGetAllocationPropertiesFromHandle(&prop_temp, nullptr) == hipErrorInvalidValue); } + HIP_CHECK(hipMemRelease(handle)); } diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc new file mode 100644 index 0000000000..f362e4f049 --- /dev/null +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc @@ -0,0 +1,210 @@ +/* +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. +*/ + +/** + * @addtogroup hipMemExportToShareableHandle hipMemExportToShareableHandle + * @{ + * @ingroup VirtualMemoryManagementTest + * `hipError_t hipMemImportFromShareableHandle(hipMemGenericAllocationHandle_t *handle, + * void *osHandle, + * hipMemAllocationHandleType shHandleType)` - + * Imports an allocation from a requested shareable handle type. + */ + +#include +#include +#include + +#include + +#include "hip_vmm_common.hh" + +/** + * Test Description + * ------------------------ + * - Basic sanity test. + * ------------------------ + * - unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemImportFromShareableHandle_Positive_Basic") { + 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)); + + hipMemGenericAllocationHandle_t imported_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 + * ------------------------ + * - Negative parameters test. + * ------------------------ + * - unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc + * Test requirements + * ------------------------ + * - Host specific (LINUX) + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemImportFromShareableHandle_Negative_Parameters") { + 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)); + + 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, + hipMemHandleTypePosixFileDescriptor), + hipErrorInvalidValue); + } + + SECTION("invalid handleType") { + HIP_CHECK_ERROR( + hipMemImportFromShareableHandle(&imported_handle, shareable_handle, hipMemHandleTypeWin32), + hipErrorNotSupported); + } + + HIP_CHECK(hipMemRelease(handle)); +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/memory/hipMemMap.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc similarity index 86% rename from projects/hip-tests/catch/unit/memory/hipMemMap.cc rename to projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc index d9b2d13a26..7a46f0f3cb 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemMap.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc @@ -19,10 +19,11 @@ LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. */ + /** * @addtogroup hipMemMap hipMemMap * @{ - * @ingroup MemoryTest + * @ingroup VirtualMemoryManagementTest * `hipError_t hipMemMap (void* ptr, * size_t size, * size_t offset, @@ -32,6 +33,7 @@ THE SOFTWARE. */ #include + #include "hip_vmm_common.hh" constexpr int N = (1 << 13); @@ -45,7 +47,7 @@ constexpr int initializer = 0; * vmm address range repeatedly. This test validates physical memory * euse using same vmm range. * ------------------------ - * - catch\unit\memory\hipMemMap.cc + * - unit/virtualMemoryManagement/hipMemMap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -57,24 +59,22 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + 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), C_h(N); // Initialize with data for (size_t idx = 0; idx < N; idx++) { A_h[idx] = idx; - C_h[idx] = idx*idx; + C_h[idx] = idx * idx; } // Allocate a physical memory chunk HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); @@ -94,8 +94,7 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") { 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<<>>(static_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())); @@ -114,7 +113,7 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") { * vmm addresses. This test validates physical memory reuse using * different vmm ranges. * ------------------------ - * - catch\unit\memory\hipMemMap.cc + * - unit/virtualMemoryManagement/hipMemMap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -125,24 +124,22 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + 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), C_h(N); // Initialize with data for (size_t idx = 0; idx < N; idx++) { A_h[idx] = idx; - C_h[idx] = idx*idx; + C_h[idx] = idx * idx; } // Allocate a physical memory chunk HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); @@ -164,8 +161,8 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") { 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])); + square_kernel<<>>( + static_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())); @@ -186,7 +183,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") { * vmm addresses at the same time and check data values integrity * between different VMMs. * ------------------------ - * - catch\unit\memory\hipMemMap.cc + * - unit/virtualMemoryManagement/hipMemMap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -197,17 +194,15 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemory_Map2MultVMMs") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + 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); @@ -253,7 +248,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemory_Map2MultVMMs") { * multiple vmm addresses. This test validates physical memory * reuse using different vmm ranges on multiple devices. * ------------------------ - * - catch\unit\memory\hipMemMap.cc + * - unit/virtualMemoryManagement/hipMemMap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -270,17 +265,15 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_MultiDev") { for (int devX = 0; devX < devicecount; devX++) { hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, devX)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + 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); @@ -328,7 +321,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_MultiDev") { * for single vmm address. This test validates VMM memory reuse * using different physical ranges. * ------------------------ - * - catch\unit\memory\hipMemMap.cc + * - unit/virtualMemoryManagement/hipMemMap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -339,24 +332,22 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle[num_buf]; // Allocate host memory and intialize data 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; + C_h[idx] = idx * idx; } // Allocate a physical memory chunk for (int buf = 0; buf < num_buf; buf++) { @@ -379,8 +370,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<<>>(static_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())); @@ -401,7 +391,7 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") { * can be mapped/unmapped to single vmm address. This test validates VMM * memory reuse using different physical ranges. * ------------------------ - * - catch\unit\memory\hipMemMap.cc + * - unit/virtualMemoryManagement/hipMemMap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -417,17 +407,15 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_MultiGPU") { size_t buffer_size = N * sizeof(int); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; std::vector handle(devicecount); // Allocate host memory and intialize data std::vector A_h(N), B_h(N); @@ -493,7 +481,7 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_MultiGPU") { * - Check if a partial part of a physical chunk can be mapped/unmapped * to a smaller vmm address. * ------------------------ - * - catch\unit\memory\hipMemMap.cc + * - unit/virtualMemoryManagement/hipMemMap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -504,17 +492,15 @@ TEST_CASE("Unit_hipMemMap_MapPartialPhysicalMem") { size_t buffer_size = N * sizeof(int); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + 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); @@ -523,7 +509,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialPhysicalMem") { A_h[idx] = idx; } // Allocate a bigger physical memory chunk of twice size_mem - HIP_CHECK(hipMemCreate(&handle, 2*size_mem, &prop, 0)); + 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)); @@ -549,7 +535,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialPhysicalMem") { * - Check if a partial part of a VMM range can be mapped/unmapped * to a physical address. * ------------------------ - * - catch\unit\memory\hipMemMap.cc + * - unit/virtualMemoryManagement/hipMemMap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -560,17 +546,15 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") { size_t buffer_size = N * sizeof(int); hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + 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); @@ -582,13 +566,13 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") { HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); // Allocate virtual address range of size twice size_mem hipDeviceptr_t ptrA; - HIP_CHECK(hipMemAddressReserve(&ptrA, 2*size_mem, 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrA, 2 * 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(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)); @@ -596,7 +580,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") { HIP_CHECK(hipMemUnmap(ptrA, size_mem)); // Release resources HIP_CHECK(hipMemRelease(handle)); - HIP_CHECK(hipMemAddressFree(ptrA, 2*size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, 2 * size_mem)); } /** @@ -604,7 +588,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") { * ------------------------ * - Negative Argument Tests * ------------------------ - * - catch\unit\memory\hipMemMap.cc + * - unit/virtualMemoryManagement/hipMemMap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -615,16 +599,15 @@ TEST_CASE("Unit_hipMemMap_negative") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; hipDeviceptr_t ptrA; // Allocate physical memory @@ -633,8 +616,7 @@ 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(nullptr, size_mem, 0, handle, 0) == hipErrorInvalidValue); } SECTION("pass zero to size") { @@ -642,9 +624,9 @@ TEST_CASE("Unit_hipMemMap_negative") { } SECTION("pass negative to offset") { - REQUIRE(hipMemMap(&ptrA, size_mem, -1, handle, 0) == - hipErrorInvalidValue); + REQUIRE(hipMemMap(&ptrA, size_mem, -1, handle, 0) == hipErrorInvalidValue); } + HIP_CHECK(hipMemRelease(handle)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); } diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMapArrayAsync.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMapArrayAsync.cc new file mode 100644 index 0000000000..2bc726f82c --- /dev/null +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMapArrayAsync.cc @@ -0,0 +1,110 @@ +/* +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. +*/ + +/** + * @addtogroup hipMemMapArrayAsync hipMemMapArrayAsync + * @{ + * @ingroup VirtualMemoryManagementTest + * `hipError_t hipMemMapArrayAsync(hipArrayMapInfo *mapInfoList, + * unsigned int count, + * hipStream_t stream)` - + * Maps or unmaps subregions of sparse HIP arrays and sparse HIP mipmapped arrays. + */ + +#include +#include +#include + +#include "hip_vmm_common.hh" + +/** + * Test Description + * ------------------------ + * - Basic sanity test. + * ------------------------ + * - unit/virtualMemoryManagement/hipMemMapArrayAsync.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemMapArrayAsync_Positive_Basic") { + HIP_CHECK(hipFree(0)); + + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, 0)); + checkVMMSupported(device); + + CHECK_IMAGE_SUPPORT; + + hipmipmappedArray array; + + HIP_ARRAY3D_DESCRIPTOR desc = {}; + using vec_info = vector_info; + desc.Format = vec_info::format; + desc.NumChannels = vec_info::size; + desc.Width = 1; + desc.Height = 1; + desc.Flags = CUDA_ARRAY3D_SPARSE; + + unsigned int levels = 2; + + HIP_CHECK(hipMipmappedArrayCreate(&array, &desc, levels)); + + hipMemAllocationProp prop = {}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; + prop.allocFlags.usage = CU_MEM_CREATE_USAGE_TILE_POOL; + + size_t granularity; + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityRecommended)); + + hipMemGenericAllocationHandle_t handle; + HIP_CHECK(hipMemCreate(&handle, granularity, &prop, 0)); + + hipArrayMapInfo map_info_list = {}; + map_info_list.resourceType = HIP_RESOURCE_TYPE_MIPMAPPED_ARRAY; + map_info_list.resource.mipmap = array; + map_info_list.subresourceType = hipArraySparseSubresourceTypeSparseLevel; + map_info_list.subresource.sparseLevel.extentWidth = 1; + map_info_list.subresource.sparseLevel.extentHeight = 1; + map_info_list.subresource.sparseLevel.extentDepth = 1; + map_info_list.memOperationType = hipMemOperationTypeMap; + map_info_list.memHandleType = hipMemHandleTypeGeneric; + map_info_list.memHandle.memHandle = handle; + map_info_list.deviceBitMask = 0x1; + + StreamGuard stream(Streams::created); + + HIP_CHECK(hipMemMapArrayAsync(&map_info_list, 1, stream.stream())); + HIP_CHECK(hipStreamSynchronize(stream.stream())); + + map_info_list.memOperationType = hipMemOperationTypeUnmap; + map_info_list.memHandle.memHandle = NULL; + HIP_CHECK(hipMemMapArrayAsync(&map_info_list, 1, stream.stream())); + HIP_CHECK(hipStreamSynchronize(stream.stream())); + + HIP_CHECK(hipMemRelease(handle)); + + HIP_CHECK(hipMipmappedArrayDestroy(array)); +} \ No newline at end of file diff --git a/projects/hip-tests/catch/unit/memory/hipMemRelease.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRelease.cc similarity index 89% rename from projects/hip-tests/catch/unit/memory/hipMemRelease.cc rename to projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRelease.cc index d43647da24..e544710757 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemRelease.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRelease.cc @@ -23,7 +23,7 @@ THE SOFTWARE. /** * @addtogroup hipMemRelease hipMemRelease * @{ - * @ingroup MemoryTest + * @ingroup VirtualMemoryManagementTest * `hipMemRelease(hipMemGenericAllocationHandle_t handle)` - * Release a memory handle representing a memory allocation which was previously * allocated through hipMemCreate. @@ -36,13 +36,11 @@ THE SOFTWARE. * ------------------------ * - Negative Tests * ------------------------ - * - catch\unit\memory\hipMemRelease.cc + * - unit/virtualMemoryManagement/hipMemRelease.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemRelease_negative") { - SECTION("Nullptr to handle") { - REQUIRE(hipMemRelease(nullptr) == hipErrorInvalidValue); - } + SECTION("Nullptr to handle") { REQUIRE(hipMemRelease(nullptr) == hipErrorInvalidValue); } } diff --git a/projects/hip-tests/catch/unit/memory/hipMemRetainAllocationHandle.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc similarity index 76% rename from projects/hip-tests/catch/unit/memory/hipMemRetainAllocationHandle.cc rename to projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc index 29910c61aa..1b4d00dea3 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemRetainAllocationHandle.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc @@ -23,7 +23,7 @@ THE SOFTWARE. /** * @addtogroup hipMemRetainAllocationHandle hipMemRetainAllocationHandle * @{ - * @ingroup MemoryTest + * @ingroup VirtualMemoryManagementTest * `hipError_t hipMemRetainAllocationHandle(hipMemGenericAllocationHandle_t* handle, * void* addr)` - * Returns the allocation handle of the backing memory allocation given the address. @@ -31,6 +31,7 @@ THE SOFTWARE. #include #include + #include "hip_vmm_common.hh" #define DATA_SIZE (1 << 13) @@ -41,41 +42,39 @@ THE SOFTWARE. * - Create a VM mapped to physical memory. Input addr to * hipMemRetainAllocationHandle and validate the handle. * ------------------------ - * - catch\unit\memory\hipMemRetainAllocationHandle.cc + * - unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemRetainAllocationHandle_SetGet") { + HIP_CHECK(hipFree(0)); size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; hipDeviceptr_t ptrA; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); // Allocate virtual address range - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, - 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); // Test hipMemRetainAllocationHandle hipMemGenericAllocationHandle_t gethandle; // Check beginning of VMM ptr - HIP_CHECK(hipMemRetainAllocationHandle(&gethandle, - reinterpret_cast(ptrA))); + HIP_CHECK(hipMemRetainAllocationHandle(&gethandle, reinterpret_cast(ptrA))); REQUIRE(gethandle == handle); HIP_CHECK(hipMemRelease(handle)); HIP_CHECK(hipMemUnmap(ptrA, size_mem)); @@ -87,58 +86,56 @@ TEST_CASE("Unit_hipMemRetainAllocationHandle_SetGet") { * ------------------------ * - Negative Tests * ------------------------ - * - catch\unit\memory\hipMemRetainAllocationHandle.cc + * - unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 */ TEST_CASE("Unit_hipMemRetainAllocationHandle_NegTst") { + HIP_CHECK(hipFree(0)); size_t granularity = 0; constexpr int N = DATA_SIZE; size_t buffer_size = N * sizeof(int); int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; hipDeviceptr_t ptrA; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); // Allocate virtual address range - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, - 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); // Test hipMemRetainAllocationHandle hipMemGenericAllocationHandle_t gethandle; SECTION("nullptr handle") { - REQUIRE(hipMemRetainAllocationHandle(nullptr, - reinterpret_cast(ptrA)) == hipErrorInvalidValue); + REQUIRE(hipMemRetainAllocationHandle(nullptr, reinterpret_cast(ptrA)) == + hipErrorInvalidValue); } SECTION("nullptr Vmm ptr") { - REQUIRE(hipMemRetainAllocationHandle(&gethandle, nullptr) == - hipErrorInvalidValue); + REQUIRE(hipMemRetainAllocationHandle(&gethandle, nullptr) == hipErrorInvalidValue); } SECTION("not mapped address") { hipDeviceptr_t ptrB; HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem, 0, 0, 0)); - REQUIRE(hipMemRetainAllocationHandle(&gethandle, - reinterpret_cast(ptrB)) == hipErrorInvalidValue); + REQUIRE(hipMemRetainAllocationHandle(&gethandle, reinterpret_cast(ptrB)) == + hipErrorInvalidValue); HIP_CHECK(hipMemAddressFree(ptrB, size_mem)); } HIP_CHECK(hipMemRelease(handle)); HIP_CHECK(hipMemUnmap(ptrA, size_mem)); SECTION("unmapped address") { - REQUIRE(hipMemRetainAllocationHandle(&gethandle, - reinterpret_cast(ptrA)) == hipErrorInvalidValue); + REQUIRE(hipMemRetainAllocationHandle(&gethandle, reinterpret_cast(ptrA)) == + hipErrorInvalidValue); } HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); } diff --git a/projects/hip-tests/catch/unit/memory/hipMemSetGetAccess.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc similarity index 83% rename from projects/hip-tests/catch/unit/memory/hipMemSetGetAccess.cc rename to projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc index 20910f9803..dca05c5f1a 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemSetGetAccess.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc @@ -23,7 +23,7 @@ THE SOFTWARE. /** * @addtogroup hipMemSetAccess hipMemSetAccess * @{ - * @ingroup MemoryTest + * @ingroup VirtualMemoryManagementTest * `hipError_t hipMemSetAccess (void* ptr, * size_t size, * const hipMemAccessDesc* desc, @@ -31,18 +31,23 @@ THE SOFTWARE. * Set the access flags for each location specified in desc for the given * virtual address range. */ -#include "hipMallocManagedCommon.hh" -#include "hip_vmm_common.hh" -#include -#include + #ifdef __linux__ #include #include #endif + +#include +#include + +#include "hipMallocManagedCommon.hh" +#include "hip_vmm_common.hh" + #define THREADS_PER_BLOCK 512 #define NUM_OF_BUFFERS 3 #define DATA_SIZE (1 << 13) -#define NEW_DATA_SIZE (2*DATA_SIZE) +#define NEW_DATA_SIZE (2 * DATA_SIZE) + constexpr int initializer = 0; /** @@ -63,7 +68,7 @@ static __global__ void square_kernel(int* Buff) { * Validate that flags = hipMemAccessFlagsProtNone is returned by * hipMemGetAccess() when location is set to device 1. * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -75,23 +80,21 @@ TEST_CASE("Unit_hipMemSetAccess_SetGet") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); // Allocate virtual address range hipDeviceptr_t ptrA; - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, - 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); HIP_CHECK(hipMemRelease(handle)); // Set access @@ -132,7 +135,7 @@ TEST_CASE("Unit_hipMemSetAccess_SetGet") { * flags = hipMemAccessFlagsProtReadWrite is returned by hipMemGetAccess() * when location is set to device 1. * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -150,18 +153,17 @@ TEST_CASE("Unit_hipMemSetAccess_MultDevSetGet") { } HIP_CHECK(hipDeviceGet(&device0, deviceId)); - checkVMMSupported(device0) + checkVMMSupported(device0); HIP_CHECK(hipDeviceGet(&device1, (deviceId + 1))); - checkVMMSupported(device1) + checkVMMSupported(device1); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device0; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); @@ -203,7 +205,7 @@ TEST_CASE("Unit_hipMemSetAccess_MultDevSetGet") { * to device 0. Validate that flags = 3 is returned by hipMemGetAccess() * for entire virtual address range when location is set to device 0. * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -215,23 +217,21 @@ TEST_CASE("Unit_hipMemSetAccess_EntireVMMRangeSetGet") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate physical memory hipMemGenericAllocationHandle_t handle; hipDeviceptr_t ptrA; HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); // Allocate virtual address range - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, - 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); HIP_CHECK(hipMemRelease(handle)); // Set access @@ -250,8 +250,7 @@ TEST_CASE("Unit_hipMemSetAccess_EntireVMMRangeSetGet") { REQUIRE(flags == hipMemAccessFlagsProtReadWrite); uint64_t 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)); @@ -262,7 +261,7 @@ TEST_CASE("Unit_hipMemSetAccess_EntireVMMRangeSetGet") { * ------------------------ * - Negative Tests * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -274,23 +273,21 @@ TEST_CASE("Unit_hipMemGetAccess_NegTst") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate physical memory hipMemGenericAllocationHandle_t handle; hipDeviceptr_t ptrA; HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); // Allocate virtual address range - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, - 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); HIP_CHECK(hipMemRelease(handle)); // Set access @@ -326,7 +323,7 @@ TEST_CASE("Unit_hipMemGetAccess_NegTst") { * address range, launch a kernel to perform operation on the data and * validate the result. * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -340,16 +337,15 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") { HIP_CHECK(hipGetDeviceCount(&devicecount)); for (deviceId = 0; deviceId < devicecount; deviceId++) { HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate physical memory hipDeviceptr_t ptrA; hipMemGenericAllocationHandle_t handle; @@ -371,9 +367,8 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") { } HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), 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 / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0, + static_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())); @@ -389,7 +384,7 @@ TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") { * Access (Read/Write) the virtual pointer directly on host. * Ensure this behavior for all devices on host. * ------------------------ - * - catch\unit\memory\hipMemMap.cc + * - unit/virtualMemoryManagement/hipMemMap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -407,17 +402,15 @@ TEST_CASE("Unit_hipMemSetAccess_AccessDirectlyFromHost") { for (int dev = 0; dev < devicecount; dev++) { hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, dev)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; // Allocate a physical memory chunk HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); @@ -453,7 +446,7 @@ TEST_CASE("Unit_hipMemSetAccess_AccessDirectlyFromHost") { * the property of the range to read only. Check if the memory * range can be read. * ------------------------ - * - catch\unit\memory\hipMemMap.cc + * - unit/virtualMemoryManagement/hipMemMap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -465,17 +458,15 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") { int dev = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, dev)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices HIP_CHECK( - hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + 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 @@ -526,14 +517,12 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") { 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)); + 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)); + REQUIRE(hipErrorInvalidValue == hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); } #endif HIP_CHECK(hipMemUnmap(ptrA, size_mem)); @@ -541,6 +530,7 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") { HIP_CHECK(hipMemRelease(handle)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); } + /** * Test Description * ------------------------ @@ -548,7 +538,7 @@ TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") { * a Virtual Memory chunk and a Unified Memory chunk. Test if data can * be exchanged between these chunks. * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -565,23 +555,21 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2UnifiedMemCpy") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate physical memory hipMemGenericAllocationHandle_t handle; hipDeviceptr_t ptrA, ptrB; HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); // Allocate virtual address range - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, - 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); HIP_CHECK(hipMemRelease(handle)); // Set access @@ -623,7 +611,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2UnifiedMemCpy") { * Memory chunk and a Device Memory chunk. Test if data can be exchanged * between these chunks. * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -635,23 +623,21 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2DevMemCpy") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate physical memory hipMemGenericAllocationHandle_t handle; hipDeviceptr_t ptrA, ptrB; HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); // Allocate virtual address range - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, - 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); HIP_CHECK(hipMemRelease(handle)); // Set access @@ -682,7 +668,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2DevMemCpy") { * Peer Device Memory chunk. Test if data can be exchanged between * these chunks using hipMemcpyDtoD. * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -694,23 +680,21 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") { int deviceId = 0, value = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate physical memory hipMemGenericAllocationHandle_t handle; hipDeviceptr_t ptrA; HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); // Allocate virtual address range - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, - 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); HIP_CHECK(hipMemRelease(handle)); // Set access @@ -737,9 +721,8 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") { break; } HIP_CHECK(hipDeviceGet(&device_other, deviceId)); - HIP_CHECK(hipDeviceGetAttribute(&value, - hipDeviceAttributeVirtualMemoryManagementSupported, - device_other)); + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeVirtualMemoryManagementSupported, + device_other)); if (value == 0) { // Virtual Memory Mgmt is not supported WARN("Machine does not support Virtual Memory Management\n"); @@ -764,7 +747,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") { * a Peer Device Memory chunk. Test if data can be exchanged between * these chunks using hipMemcpyPeer. * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -776,23 +759,21 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy") { int deviceId = 0, value = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate physical memory hipMemGenericAllocationHandle_t handle; hipDeviceptr_t ptrA; HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); // Allocate virtual address range - HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, - 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); HIP_CHECK(hipMemRelease(handle)); // Set access @@ -820,9 +801,8 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy") { break; } HIP_CHECK(hipDeviceGet(&device_other, deviceId)); - HIP_CHECK(hipDeviceGetAttribute(&value, - hipDeviceAttributeVirtualMemoryManagementSupported, - device_other)); + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeVirtualMemoryManagementSupported, + device_other)); if (value == 0) { // Virtual Memory Mgmt is not supported WARN("Machine does not support Virtual Memory Management\n"); @@ -848,7 +828,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy") { * address space in device 0(PtrB). Check if data can be copied from * PtrA -> PtrB using hipMemcpy. * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -860,16 +840,15 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMMemCpy") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; // Allocate physical memory hipMemGenericAllocationHandle_t handle1, handle2; HIP_CHECK(hipMemCreate(&handle1, size_mem, &prop, 0)); @@ -912,7 +891,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMMemCpy") { * address space in device 1(PtrB). Check if data can be copied from * PtrA -> PtrB using hipMemcpyPeer. * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -924,16 +903,15 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") { int deviceId = 0, value = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); @@ -966,9 +944,8 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") { } std::fill(B_h.begin(), B_h.end(), initializer); HIP_CHECK(hipDeviceGet(&device_other, deviceId)); - HIP_CHECK(hipDeviceGetAttribute(&value, - hipDeviceAttributeVirtualMemoryManagementSupported, - device_other)); + HIP_CHECK(hipDeviceGetAttribute(&value, hipDeviceAttributeVirtualMemoryManagementSupported, + device_other)); if (value == 0) { // Virtual Memory Mgmt is not supported WARN("Machine does not support Virtual Memory Management\n"); @@ -980,9 +957,8 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") { prop_loc.location.type = hipMemLocationTypeDevice; prop_loc.location.id = device_other; // Current Devices HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop_loc, - hipMemAllocationGranularityMinimum)); - size_t size_mem_loc = - ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemAllocationGranularityMinimum)); + size_t size_mem_loc = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle_loc; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle_loc, size_mem_loc, &prop_loc, 0)); @@ -1015,7 +991,7 @@ TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") { * chunk of memory and map it to device1. Check if these 2 distinct memory * chunks can be mapped to a single address space. * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -1028,32 +1004,28 @@ TEST_CASE("Unit_hipMemSetAccess_MapPhysChksFromMulDev") { size_t buffer_size = N * sizeof(int), granularity = 0; int deviceId = 0; // Allocate resources for all gpus - hipMemGenericAllocationHandle_t *handle = - static_cast( - malloc(sizeof(hipMemGenericAllocationHandle_t)*numOfBuffers)); + hipMemGenericAllocationHandle_t* handle = static_cast( + malloc(sizeof(hipMemGenericAllocationHandle_t) * numOfBuffers)); REQUIRE(handle != nullptr); - size_t *size_mem = static_cast( - malloc(sizeof(size_t)*numOfBuffers)); + size_t* size_mem = static_cast(malloc(sizeof(size_t) * numOfBuffers)); REQUIRE(size_mem != nullptr); size_t total_mem = 0; // Create memory chunks for (deviceId = 0; deviceId < numOfBuffers; deviceId++) { hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop_loc{}; prop_loc.type = hipMemAllocationTypePinned; prop_loc.location.type = hipMemLocationTypeDevice; prop_loc.location.id = device; // Current Devices HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop_loc, - hipMemAllocationGranularityMinimum)); + hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_mem[deviceId] = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_mem[deviceId] = ((granularity + buffer_size - 1) / granularity) * granularity; total_mem = total_mem + size_mem[deviceId]; // Allocate physical memory chunks - HIP_CHECK(hipMemCreate(&handle[deviceId], size_mem[deviceId], - &prop_loc, 0)); + HIP_CHECK(hipMemCreate(&handle[deviceId], size_mem[deviceId], &prop_loc, 0)); } // Allocate virtual address range for all the memory chunks hipDeviceptr_t ptrA; @@ -1063,9 +1035,9 @@ TEST_CASE("Unit_hipMemSetAccess_MapPhysChksFromMulDev") { hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); uint64_t uiptr = reinterpret_cast(ptrA); - uiptr = uiptr + deviceId*size_mem[deviceId]; - HIP_CHECK(hipMemMap(reinterpret_cast(uiptr), - size_mem[deviceId], 0, handle[deviceId], 0)); + uiptr = uiptr + deviceId * size_mem[deviceId]; + HIP_CHECK( + hipMemMap(reinterpret_cast(uiptr), size_mem[deviceId], 0, handle[deviceId], 0)); HIP_CHECK(hipMemRelease(handle[deviceId])); // Set access hipMemAccessDesc accessDesc_loc = {}; @@ -1073,46 +1045,28 @@ TEST_CASE("Unit_hipMemSetAccess_MapPhysChksFromMulDev") { accessDesc_loc.location.id = device; accessDesc_loc.flags = hipMemAccessFlagsProtReadWrite; // Make the address accessible to deviceId - HIP_CHECK(hipMemSetAccess(reinterpret_cast(uiptr), - size_mem[deviceId], &accessDesc_loc, 1)); + HIP_CHECK( + hipMemSetAccess(reinterpret_cast(uiptr), size_mem[deviceId], &accessDesc_loc, 1)); } - std::vector A_h(numOfBuffers*N), - B_h(numOfBuffers*N); + std::vector A_h(numOfBuffers * N), B_h(numOfBuffers * N); // Fill Data - for (int idx = 0; idx < (numOfBuffers*N); idx++) { - A_h[idx] = idx*idx; + for (int idx = 0; idx < (numOfBuffers * N); idx++) { + A_h[idx] = idx * idx; } - HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), numOfBuffers*buffer_size)); - HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, numOfBuffers*buffer_size)); + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), numOfBuffers * buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, numOfBuffers * buffer_size)); // Validate Results REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); for (deviceId = 0; deviceId < numOfBuffers; deviceId++) { uint64_t uiptr = reinterpret_cast(ptrA); - uiptr = uiptr + deviceId*size_mem[deviceId]; - HIP_CHECK(hipMemUnmap(reinterpret_cast(uiptr), - size_mem[deviceId])); + uiptr = uiptr + deviceId * size_mem[deviceId]; + HIP_CHECK(hipMemUnmap(reinterpret_cast(uiptr), size_mem[deviceId])); } HIP_CHECK(hipMemAddressFree(ptrA, total_mem)); free(handle); free(size_mem); } -/** - * Test Description - * ------------------------ - * - Testing memory resize: Allocate physical memory and map it to virtual - * address range (PtrA). After setting device permission, copy data from - * host to device. Allocate another chunk of memory of a different size. - * Map the new chunk to offset (PtrA + size of old chunk). - * After setting device permission, copy data from host to device at - * offset (PtrA + size of old chunk). Validate both the old data and new - * data after copying back to host. - * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.1 - */ class vmm_resize_class { size_t current_size_tot; size_t current_size_rounded_tot; @@ -1120,18 +1074,16 @@ class vmm_resize_class { std::vector vhandle; std::vector vsize; // allocate initial VMM memory chunk - int allocate_vmm(hipDeviceptr_t *ptr, hipDevice_t device, - size_t size) { + int allocate_vmm(hipDeviceptr_t* ptr, hipDevice_t device, size_t size) { size_t granularity = 0; hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_rounded = - ((granularity + size - 1) / granularity) * granularity; + size_t size_rounded = ((granularity + size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_rounded, &prop, 0)); @@ -1155,20 +1107,20 @@ class vmm_resize_class { } public: - vmm_resize_class(hipDeviceptr_t *ptr, hipDevice_t device, size_t size): - current_size_tot(0), current_size_rounded_tot(0) { + vmm_resize_class(hipDeviceptr_t* ptr, hipDevice_t device, size_t size) + : current_size_tot(0), current_size_rounded_tot(0) { allocate_vmm(ptr, device, size); } // Free all VMM void free_vmm() { - for (hipMemGenericAllocationHandle_t &myhandle : vhandle) { + for (hipMemGenericAllocationHandle_t& myhandle : vhandle) { HIP_CHECK(hipMemRelease(myhandle)); } HIP_CHECK(hipMemUnmap(ptrVmm, current_size_rounded_tot)); HIP_CHECK(hipMemAddressFree(ptrVmm, current_size_rounded_tot)); } // grow memory chunk - int grow_vmm(hipDeviceptr_t *ptr, hipDevice_t device, size_t size) { + int grow_vmm(hipDeviceptr_t* ptr, hipDevice_t device, size_t size) { size_t granularity = 0; if (size <= current_size_tot) { return -1; @@ -1177,13 +1129,12 @@ class vmm_resize_class { prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); // diff size size_t grow_size = (size - current_size_tot); - size_t size_rounded = - ((granularity + grow_size - 1) / granularity) * granularity; + size_t size_rounded = ((granularity + grow_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_rounded, &prop, 0)); @@ -1194,17 +1145,15 @@ class vmm_resize_class { // Unmap and Free the old vmm HIP_CHECK(hipMemUnmap(ptrVmm, current_size_rounded_tot)); HIP_CHECK(hipMemAddressFree(ptrVmm, current_size_rounded_tot)); - HIP_CHECK(hipMemAddressReserve(&ptrVmm, - (size_rounded + current_size_rounded_tot), 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrVmm, (size_rounded + current_size_rounded_tot), 0, 0, 0)); int idx = 0; - for (hipMemGenericAllocationHandle_t &myhandle : vhandle) { + for (hipMemGenericAllocationHandle_t& myhandle : vhandle) { if (idx == 0) { HIP_CHECK(hipMemMap(ptrVmm, vsize[idx], 0, myhandle, 0)); } else { uint64_t uiptr = reinterpret_cast(ptrVmm); - uiptr = uiptr + vsize[idx-1]; - HIP_CHECK(hipMemMap(reinterpret_cast(uiptr), - vsize[idx], 0, myhandle, 0)); + uiptr = uiptr + vsize[idx - 1]; + HIP_CHECK(hipMemMap(reinterpret_cast(uiptr), vsize[idx], 0, myhandle, 0)); } idx++; } @@ -1214,9 +1163,7 @@ class vmm_resize_class { accessDesc.location.id = device; accessDesc.flags = hipMemAccessFlagsProtReadWrite; // Make the address accessible to GPU 0 - HIP_CHECK(hipMemSetAccess(ptrVmm, - (size_rounded + current_size_rounded_tot), - &accessDesc, 1)); + HIP_CHECK(hipMemSetAccess(ptrVmm, (size_rounded + current_size_rounded_tot), &accessDesc, 1)); *ptr = ptrVmm; current_size_tot += size; current_size_rounded_tot += size_rounded; @@ -1224,6 +1171,22 @@ class vmm_resize_class { } }; +/** + * Test Description + * ------------------------ + * - Testing memory resize: Allocate physical memory and map it to virtual + * address range (PtrA). After setting device permission, copy data from + * host to device. Allocate another chunk of memory of a different size. + * Map the new chunk to offset (PtrA + size of old chunk). + * After setting device permission, copy data from host to device at + * offset (PtrA + size of old chunk). Validate both the old data and new + * data after copying back to host. + * ------------------------ + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ TEST_CASE("Unit_hipMemSetAccess_GrowVMM") { hipDeviceptr_t ptr; constexpr int N = DATA_SIZE; @@ -1231,11 +1194,11 @@ TEST_CASE("Unit_hipMemSetAccess_GrowVMM") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); // Create VMM Object of size buffer_size vmm_resize_class resizeobj(&ptr, device, buffer_size); // Inititalize Host Buffer - int *ptrA_h = static_cast(malloc(buffer_size)); + int* ptrA_h = static_cast(malloc(buffer_size)); REQUIRE(ptrA_h != nullptr); for (int idx = 0; idx < N; idx++) { ptrA_h[idx] = idx; @@ -1250,17 +1213,16 @@ TEST_CASE("Unit_hipMemSetAccess_GrowVMM") { return; } free(ptrA_h); - ptrA_h = static_cast(malloc(buffer_size_new - buffer_size)); + ptrA_h = static_cast(malloc(buffer_size_new - buffer_size)); REQUIRE(ptrA_h != nullptr); for (int idx = 0; idx < (Nnew - N); idx++) { ptrA_h[idx] = N + idx; } - int *ptrB_h = static_cast(malloc(buffer_size_new)); + int* ptrB_h = static_cast(malloc(buffer_size_new)); REQUIRE(ptrB_h != nullptr); uint64_t 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++) { @@ -1275,18 +1237,6 @@ TEST_CASE("Unit_hipMemSetAccess_GrowVMM") { resizeobj.free_vmm(); } -/** - * Test Description - * ------------------------ - * - Multithreaded test: Allocate unique virtual memory chunks from - * multiple threads. Transfer data to these chunks from host and execute - * kernel function on these data. Validate the results. - * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.1 - */ std::atomic bTestPassed{1}; #define NUM_THREADS 5 void test_thread(hipDevice_t device) { @@ -1296,14 +1246,14 @@ void test_thread(hipDevice_t device) { // Create VMM Object of size buffer_size vmm_resize_class vmmobj(&ptr, device, buffer_size); // Inititalize Host Buffer - int *ptrA_h = static_cast(malloc(buffer_size)); + int* ptrA_h = static_cast(malloc(buffer_size)); REQUIRE(ptrA_h != nullptr); for (int idx = 0; idx < N; idx++) { ptrA_h[idx] = idx; } // Copy to VMM HIP_CHECK(hipMemcpyHtoD(ptr, ptrA_h, buffer_size)); - int *ptrB_h = static_cast(malloc(buffer_size)); + int* ptrB_h = static_cast(malloc(buffer_size)); REQUIRE(ptrB_h != nullptr); HIP_CHECK(hipMemcpyDtoH(ptrB_h, ptr, buffer_size)); bool bPassed = true; @@ -1323,11 +1273,23 @@ void test_thread(hipDevice_t device) { vmmobj.free_vmm(); } +/** + * Test Description + * ------------------------ + * - Multithreaded test: Allocate unique virtual memory chunks from + * multiple threads. 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_Multithreaded") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); std::thread T[NUM_THREADS]; for (int i = 0; i < NUM_THREADS; i++) { T[i] = std::thread(test_thread, device); @@ -1340,18 +1302,7 @@ TEST_CASE("Unit_hipMemSetAccess_Multithreaded") { } #ifdef __linux__ -/** - * 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. - * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc - * Test requirements - * ------------------------ - * - HIP_VERSION >= 6.1 - */ + bool test_mprocess() { int fd[2]; bool testResult = false; @@ -1369,7 +1320,7 @@ bool test_mprocess() { hipDeviceptr_t ptr; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupportedRetVal(device) + checkVMMSupportedRetVal(device); // Create VMM Object of size buffer_size vmm_resize_class vmmobj(&ptr, device, buffer_size); // Inititalize Host Buffer @@ -1399,7 +1350,7 @@ bool test_mprocess() { hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupportedRetVal(device) + checkVMMSupportedRetVal(device); // Create VMM Object of size buffer_size vmm_resize_class vmmobj(&ptr, device, buffer_size); // Inititalize Host Buffer @@ -1424,9 +1375,20 @@ bool test_mprocess() { return testResult; } -TEST_CASE("Unit_hipMemSetAccess_MultiProc") { - REQUIRE(true == test_mprocess()); -} +/** + * 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 /** @@ -1434,7 +1396,7 @@ TEST_CASE("Unit_hipMemSetAccess_MultiProc") { * ------------------------ * - Negative Tests for hipMemSetAccess() * ------------------------ - * - catch\unit\memory\hipMemSetGetAccess.cc + * - unit/virtualMemoryManagement/hipMemSetGetAccess.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -1446,16 +1408,15 @@ TEST_CASE("Unit_hipMemSetAccess_negative") { int deviceId = 0; hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; // Allocate physical memory HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); @@ -1470,73 +1431,62 @@ TEST_CASE("Unit_hipMemSetAccess_negative") { accessDesc.flags = hipMemAccessFlagsProtReadWrite; SECTION("nullptr to ptrA") { - REQUIRE(hipMemSetAccess(nullptr, size_mem, &accessDesc, 1) == - hipErrorInvalidValue); + REQUIRE(hipMemSetAccess(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); + REQUIRE(hipMemSetAccess(&ptrA, size_mem - 1, &accessDesc, 1) == hipErrorInvalidValue); } SECTION("invalid location type") { accessDesc.location.type = hipMemLocationTypeInvalid; - REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 1) == - hipErrorInvalidValue); + REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); } 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); + REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 1) == hipErrorInvalidValue); } 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(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)); } diff --git a/projects/hip-tests/catch/unit/memory/hipMemUnmap.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc similarity index 83% rename from projects/hip-tests/catch/unit/memory/hipMemUnmap.cc rename to projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc index 1244b794c6..eeadb83099 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemUnmap.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc @@ -23,22 +23,24 @@ THE SOFTWARE. /** * @addtogroup hipMemUnmap hipMemUnmap * @{ - * @ingroup MemoryTest + * @ingroup VirtualMemoryManagementTest * `hipError_t hipMemUnmap (void* ptr, size_t size)` - * Unmap memory allocation of a given address range. */ #include + #include "hip_vmm_common.hh" constexpr int N = (1 << 13); + /** * Test Description * ------------------------ * - Negative Tests * ------------------------ - * - catch\unit\memory\hipMemUnmap.cc + * - unit/virtualMemoryManagement/hipMemUnmap.cc * Test requirements * ------------------------ * - HIP_VERSION >= 6.1 @@ -50,18 +52,17 @@ TEST_CASE("Unit_hipMemUnmap_negative") { hipDevice_t device; HIP_CHECK(hipDeviceGet(&device, deviceId)); - checkVMMSupported(device) + checkVMMSupported(device); hipMemAllocationProp prop{}; prop.type = hipMemAllocationTypePinned; prop.location.type = hipMemLocationTypeDevice; prop.location.id = device; // Current Devices - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, - hipMemAllocationGranularityMinimum)); + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum)); REQUIRE(granularity > 0); - size_t size_mem = - ((granularity + buffer_size - 1) / granularity) * granularity; + size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity; hipMemGenericAllocationHandle_t handle; hipDeviceptr_t ptrA; @@ -70,18 +71,17 @@ TEST_CASE("Unit_hipMemUnmap_negative") { // Allocate virtual address range 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("pass zero to size") { - REQUIRE(hipMemUnmap(ptrA, 0) == hipErrorInvalidValue); - } + SECTION("nullptr to ptrA") { REQUIRE(hipMemUnmap(nullptr, size_mem) == hipErrorInvalidValue); } + + SECTION("pass zero to size") { REQUIRE(hipMemUnmap(ptrA, 0) == hipErrorInvalidValue); } + #if HT_NVIDIA SECTION("unmap a smaller size") { REQUIRE(hipMemUnmap(ptrA, (size_mem - 1)) == hipErrorInvalidValue); } #endif + HIP_CHECK(hipMemRelease(handle)); HIP_CHECK(hipMemUnmap(ptrA, size_mem)); HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); diff --git a/projects/hip-tests/catch/unit/memory/hipMemVmm.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc similarity index 58% rename from projects/hip-tests/catch/unit/memory/hipMemVmm.cc rename to projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc index d9f51f29ac..c2258f057d 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemVmm.cc +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc @@ -1,30 +1,30 @@ /* - Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. - Permission is hereby granted, free of charge, to any person obtaining a copy - of this software and associated documentation files (the "Software"), to deal - in the Software without restriction, including without limitation the rights - to use, copy, modify, merge, publish, distribute, sublicense, and/or sell - copies of the Software, and to permit persons to whom the Software is - furnished to do so, subject to the following conditions: - The above copyright notice and this permission notice shall be included in - all copies or substantial portions of the Software. - THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR - IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE - AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER - LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, - OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN - THE SOFTWARE. - */ +Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ /* Test Case Description: 1) This testcase verifies the basic scenario - supported on all devices */ -#include -#include -#include #include #include #include @@ -32,6 +32,10 @@ #include #include +#include +#include +#include + /* This testcase verifies HIP Mem VMM API basic scenario - supported on all devices */ @@ -42,8 +46,9 @@ TEST_CASE("Unit_hipMemVmm_Basic") { INFO("hipDeviceAttributeVirtualMemoryManagementSupported: " << vmm); if (vmm == 0) { - SUCCEED("GPU 0 doesn't support hipDeviceAttributeVirtualMemoryManagement " - "attribute. Hence skipping the testing with Pass result.\n"); + SUCCEED( + "GPU 0 doesn't support hipDeviceAttributeVirtualMemoryManagement " + "attribute. Hence skipping the testing with Pass result.\n"); return; } @@ -54,7 +59,8 @@ TEST_CASE("Unit_hipMemVmm_Basic") { memAllocationProp.location.id = 0; memAllocationProp.location.type = hipMemLocationTypeDevice; - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &memAllocationProp, hipMemAllocationGranularityRecommended)); + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &memAllocationProp, + hipMemAllocationGranularityRecommended)); size_t size = 4 * 1024; void* reservedAddress{nullptr}; @@ -78,7 +84,7 @@ TEST_CASE("Unit_hipMemVmm_Basic") { HIP_CHECK(hipMemset(reservedAddress, value, size)); HIP_CHECK(hipMemcpy(&values[0], reservedAddress, size, hipMemcpyDeviceToHost)); - for (size_t i=0; i < size; ++i) { + for (size_t i = 0; i < size; ++i) { REQUIRE(values[i] == value); } @@ -87,4 +93,3 @@ TEST_CASE("Unit_hipMemVmm_Basic") { HIP_CHECK(hipMemRelease(gaHandle)); HIP_CHECK(hipMemAddressFree(reservedAddress, size)); } - diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hip_vmm_common.hh b/projects/hip-tests/catch/unit/virtualMemoryManagement/hip_vmm_common.hh new file mode 100644 index 0000000000..a43af62758 --- /dev/null +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/hip_vmm_common.hh @@ -0,0 +1,49 @@ +/* +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