From de348d8b35c5eba19850782d0d27debaf9291035 Mon Sep 17 00:00:00 2001 From: Rupam Chetia Date: Wed, 27 Sep 2023 16:12:19 +0530 Subject: [PATCH] SWDEV-322620 - [catch2][dtest] Adding test for VMM APIs Change-Id: Ia4dea28601b3c2f0163012a7a2ceb14d3a0c369b --- catch/hipTestMain/config/config_amd_linux | 47 + catch/hipTestMain/config/config_amd_windows | 47 + catch/unit/memory/CMakeLists.txt | 12 +- catch/unit/memory/hipMemAddressFree.cc | 76 + catch/unit/memory/hipMemAddressReserve.cc | 152 ++ catch/unit/memory/hipMemCreate.cc | 463 +++++ .../memory/hipMemGetAllocationGranularity.cc | 179 ++ ...hipMemGetAllocationPropertiesFromHandle.cc | 119 ++ catch/unit/memory/hipMemMap.cc | 650 +++++++ catch/unit/memory/hipMemRelease.cc | 48 + .../memory/hipMemRetainAllocationHandle.cc | 144 ++ catch/unit/memory/hipMemSetGetAccess.cc | 1547 +++++++++++++++++ catch/unit/memory/hipMemUnmap.cc | 88 + catch/unit/memory/hip_vmm_common.hh | 59 + 14 files changed, 3630 insertions(+), 1 deletion(-) create mode 100644 catch/unit/memory/hipMemAddressFree.cc create mode 100644 catch/unit/memory/hipMemAddressReserve.cc create mode 100644 catch/unit/memory/hipMemCreate.cc create mode 100644 catch/unit/memory/hipMemGetAllocationGranularity.cc create mode 100644 catch/unit/memory/hipMemGetAllocationPropertiesFromHandle.cc create mode 100644 catch/unit/memory/hipMemMap.cc create mode 100644 catch/unit/memory/hipMemRelease.cc create mode 100644 catch/unit/memory/hipMemRetainAllocationHandle.cc create mode 100644 catch/unit/memory/hipMemSetGetAccess.cc create mode 100644 catch/unit/memory/hipMemUnmap.cc create mode 100644 catch/unit/memory/hip_vmm_common.hh diff --git a/catch/hipTestMain/config/config_amd_linux b/catch/hipTestMain/config/config_amd_linux index 75d3ffb608..f0e9487aa3 100644 --- a/catch/hipTestMain/config/config_amd_linux +++ b/catch/hipTestMain/config/config_amd_linux @@ -134,6 +134,53 @@ "Unit_deviceAllocation_InOneThread_AccessInAllThreads", "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===", "Unit_hiprtc_stdheaders", + "Unit_hipMemAddressFree_negative", + "Unit_hipMemAddressReserve_AlignmentTest", + "Unit_hipMemAddressReserve_Negative", + "Unit_hipMemCreate_BasicAllocateDeAlloc_MultGranularity", + "Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap", + "Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse", + "Unit_hipMemCreate_ChkWithKerLaunch", + "Unit_hipMemCreate_MapNonContiguousChunks", + "Unit_hipMemCreate_ChkWithMemset", + "Unit_hipMemCreate_Negative", + "Unit_hipMemGetAllocationGranularity_MinGranularity", + "Unit_hipMemGetAllocationGranularity_RecommendedGranularity", + "Unit_hipMemGetAllocationGranularity_AllGPUs", + "Unit_hipMemGetAllocationGranularity_NegativeTests", + "Unit_hipMemGetAllocationPropertiesFromHandle_functional", + "Unit_hipMemGetAllocationPropertiesFromHandle_Negative", + "Unit_hipMemMap_SameMemoryReuse", + "Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU", + "Unit_hipMemMap_PhysicalMemory_Map2MultVMMs", + "Unit_hipMemMap_PhysicalMemoryReuse_MultiDev", + "Unit_hipMemMap_VMMMemoryReuse_SingleGPU", + "Unit_hipMemMap_VMMMemoryReuse_MultiGPU", + "Unit_hipMemMap_MapPartialPhysicalMem", + "Unit_hipMemMap_MapPartialVMMMem", + "Unit_hipMemMap_negative", + "Unit_hipMemRelease_negative", + "Unit_hipMemRetainAllocationHandle_SetGet", + "Unit_hipMemRetainAllocationHandle_NegTst", + "Unit_hipMemSetAccess_SetGet", + "Unit_hipMemSetAccess_MultDevSetGet", + "Unit_hipMemSetAccess_EntireVMMRangeSetGet", + "Unit_hipMemGetAccess_NegTst", + "Unit_hipMemSetAccess_FuncTstOnMultDev", + "Unit_hipMemSetAccess_AccessDirectlyFromHost", + "Unit_hipMemSetAccess_ChangeAccessProp", + "Unit_hipMemSetAccess_Vmm2UnifiedMemCpy", + "Unit_hipMemSetAccess_Vmm2DevMemCpy", + "Unit_hipMemSetAccess_Vmm2PeerDevMemCpy", + "Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy", + "Unit_hipMemSetAccess_Vmm2VMMMemCpy", + "Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy", + "Unit_hipMemSetAccess_MapPhysChksFromMulDev", + "Unit_hipMemSetAccess_GrowVMM", + "Unit_hipMemSetAccess_Multithreaded", + "Unit_hipMemSetAccess_MultiProc", + "Unit_hipMemSetAccess_negative", + "Unit_hipMemUnmap_negative", #endif #if defined VEGA20 "=== SWDEV-419112 Below tests fail in stress test on 29/08/23 ===", diff --git a/catch/hipTestMain/config/config_amd_windows b/catch/hipTestMain/config/config_amd_windows index ad4995ec99..0938722c5e 100644 --- a/catch/hipTestMain/config/config_amd_windows +++ b/catch/hipTestMain/config/config_amd_windows @@ -235,6 +235,53 @@ "Unit_Layered2DTexture_Check_DeviceBufferToFromLayered2DArray - float4", "=== Patch which removes the typetraits implementation from std namespace in hiprtc is reverted ===", "Unit_hiprtc_stdheaders", + "Unit_hipMemAddressFree_negative", + "Unit_hipMemAddressReserve_AlignmentTest", + "Unit_hipMemAddressReserve_Negative", + "Unit_hipMemCreate_BasicAllocateDeAlloc_MultGranularity", + "Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap", + "Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse", + "Unit_hipMemCreate_ChkWithKerLaunch", + "Unit_hipMemCreate_MapNonContiguousChunks", + "Unit_hipMemCreate_ChkWithMemset", + "Unit_hipMemCreate_Negative", + "Unit_hipMemGetAllocationGranularity_MinGranularity", + "Unit_hipMemGetAllocationGranularity_RecommendedGranularity", + "Unit_hipMemGetAllocationGranularity_AllGPUs", + "Unit_hipMemGetAllocationGranularity_NegativeTests", + "Unit_hipMemGetAllocationPropertiesFromHandle_functional", + "Unit_hipMemGetAllocationPropertiesFromHandle_Negative", + "Unit_hipMemMap_SameMemoryReuse", + "Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU", + "Unit_hipMemMap_PhysicalMemory_Map2MultVMMs", + "Unit_hipMemMap_PhysicalMemoryReuse_MultiDev", + "Unit_hipMemMap_VMMMemoryReuse_SingleGPU", + "Unit_hipMemMap_VMMMemoryReuse_MultiGPU", + "Unit_hipMemMap_MapPartialPhysicalMem", + "Unit_hipMemMap_MapPartialVMMMem", + "Unit_hipMemMap_negative", + "Unit_hipMemRelease_negative", + "Unit_hipMemRetainAllocationHandle_SetGet", + "Unit_hipMemRetainAllocationHandle_NegTst", + "Unit_hipMemSetAccess_SetGet", + "Unit_hipMemSetAccess_MultDevSetGet", + "Unit_hipMemSetAccess_EntireVMMRangeSetGet", + "Unit_hipMemGetAccess_NegTst", + "Unit_hipMemSetAccess_FuncTstOnMultDev", + "Unit_hipMemSetAccess_AccessDirectlyFromHost", + "Unit_hipMemSetAccess_ChangeAccessProp", + "Unit_hipMemSetAccess_Vmm2UnifiedMemCpy", + "Unit_hipMemSetAccess_Vmm2DevMemCpy", + "Unit_hipMemSetAccess_Vmm2PeerDevMemCpy", + "Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy", + "Unit_hipMemSetAccess_Vmm2VMMMemCpy", + "Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy", + "Unit_hipMemSetAccess_MapPhysChksFromMulDev", + "Unit_hipMemSetAccess_GrowVMM", + "Unit_hipMemSetAccess_Multithreaded", + "Unit_hipMemSetAccess_MultiProc", + "Unit_hipMemSetAccess_negative", + "Unit_hipMemUnmap_negative", #endif "End of json" ] diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 173b78bea1..0ce102ed92 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -79,7 +79,17 @@ if(HIP_PLATFORM MATCHES "amd") hipExtMallocWithFlags.cc hipMallocMngdMultiThread.cc hipArray.cc - hipMemVmm.cc) + hipMemVmm.cc + hipMemCreate.cc + hipMemMap.cc + hipMemGetAllocationGranularity.cc + hipMemSetGetAccess.cc + hipMemRetainAllocationHandle.cc + hipMemUnmap.cc + hipMemAddressFree.cc + hipMemAddressReserve.cc + hipMemRelease.cc + hipMemGetAllocationPropertiesFromHandle.cc) else() set(TEST_SRC ${TEST_SRC} hipGetSymbolSizeAddress.cc) endif() diff --git a/catch/unit/memory/hipMemAddressFree.cc b/catch/unit/memory/hipMemAddressFree.cc new file mode 100644 index 0000000000..6890a26d8c --- /dev/null +++ b/catch/unit/memory/hipMemAddressFree.cc @@ -0,0 +1,76 @@ +/* +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 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. +*/ + +/** + * @addtogroup hipMemAddressFree hipMemAddressFree + * @{ + * @ingroup MemoryTest + * `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) + +/** + * Test Description + * ------------------------ + * - Negative Tests + * ------------------------ + * - catch\unit\memory\hipMemAddressFree.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemAddressFree_negative") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + // Allocate virtual address range + hipDeviceptr_t ptrA; + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); + + SECTION("nullptr to devptr") { + REQUIRE(hipMemAddressFree(nullptr, size_mem) == hipErrorInvalidValue); + } + + SECTION("pass zero to size") { + REQUIRE(hipMemAddressFree(ptrA, 0) == hipErrorInvalidValue); + } + + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} diff --git a/catch/unit/memory/hipMemAddressReserve.cc b/catch/unit/memory/hipMemAddressReserve.cc new file mode 100644 index 0000000000..c6a76fc56c --- /dev/null +++ b/catch/unit/memory/hipMemAddressReserve.cc @@ -0,0 +1,152 @@ +/* +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 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. +*/ + +/** + * @addtogroup hipMemAddressReserve hipMemAddressReserve + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemAddressReserve (void** ptr, + * size_t size, + * size_t alignment, + * void* addr, + * unsigned long long flags)` - + * Reserves an address range. + */ + +#include +#include "hip_vmm_common.hh" + +#define DATA_SIZE (1 << 13) + +/** + * Test Description + * ------------------------ + * - Verify if reserved address returned by hipMemAddressReserve + * for different alignment values are correctly aligned. + * ------------------------ + * - catch\unit\memory\hipMemAddressReserve.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") { + size_t granularity = 0; + constexpr int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + constexpr int initializer = 0; + int deviceId = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + // Allocate virtual address range + hipDeviceptr_t ptrA; + size_t alignmnt = 1; + hipMemGenericAllocationHandle_t handle; + // Allocate physical memory + HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); + // Allocate host memory and intialize data + std::vector A_h(N), B_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + // check for address alignment fron 2 to 1024 + for (int iter = 0; iter < 12; iter++) { + alignmnt = alignmnt * 2; + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, alignmnt, 0, 0)); + REQUIRE((reinterpret_cast(ptrA) % alignmnt) == 0); + std::fill(B_h.begin(), B_h.end(), initializer); + HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + } + HIP_CHECK(hipMemRelease(handle)); +} + +/** + * Test Description + * ------------------------ + * - Negative Tests + * ------------------------ + * - catch\unit\memory\hipMemAddressReserve.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemAddressReserve_Negative") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + // Allocate virtual address range + hipDeviceptr_t ptrA; + SECTION("Nullptr to ptr") { + REQUIRE(hipMemAddressReserve(nullptr, size_mem, 0, 0, 0) == + hipErrorInvalidValue); + } + + SECTION("pass size as 0") { + 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); + } +#endif + SECTION("pass size as non multiple of host page size") { + REQUIRE(hipMemAddressReserve(&ptrA, (size_mem - 1), 0, 0, 0) == + hipErrorMemoryAllocation); + } +} diff --git a/catch/unit/memory/hipMemCreate.cc b/catch/unit/memory/hipMemCreate.cc new file mode 100644 index 0000000000..c6823505ac --- /dev/null +++ b/catch/unit/memory/hipMemCreate.cc @@ -0,0 +1,463 @@ +/* +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 hipMemCreate hipMemCreate + * @{ + * @ingroup MemoryTest + * `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 + +#define THREADS_PER_BLOCK 512 +#define NUM_OF_BUFFERS 3 +#define DATA_SIZE (1 << 13) + +/** + * Test Description + * ------------------------ + * - Allocate physical memories for different multiples of + * granularity and deallocate them. + * ------------------------ + * - catch\unit\memory\hipMemCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemCreate_BasicAllocateDeAlloc_MultGranularity") { + size_t granularity = 0; + int deviceId = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + hipMemGenericAllocationHandle_t handle; + // Allocate physical memory + for (int mul = 1; mul < 64; mul++) { + HIP_CHECK(hipMemCreate(&handle, granularity*mul, &prop, 0)); + HIP_CHECK(hipMemRelease(handle)); + } +} + +/** + * Test Description + * ------------------------ + * - Allocate physical memory and map it to virtual address range. + * After setting device permission, copy data from host to VMM memory + * and back to host. Verify the result. Release handle at end after + * unmapping VMM range. + * ------------------------ + * - catch\unit\memory\hipMemCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + // Allocate host memory and intialize data + std::vector A_h(N), B_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + HIP_CHECK(hipMemRelease(handle)); +} + +/** + * Test Description + * ------------------------ + * - Allocate physical memory and map it to virtual address + * range. After setting device permission, copy data from host + * to VMM memory and back to host. Verify the result. Release + * handle before the VMM range is used. + * ------------------------ + * - catch\unit\memory\hipMemCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + // Allocate host memory and intialize data + std::vector A_h(N), B_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Allocate physical memory and map it to virtual address + * range. After setting 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 + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +#if HT_NVIDIA +// This test is disabled. Will be enabled once VMM feature is fully +// available +TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + std::vector A_h(N), B_h(N), C_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + C_h[idx] = idx*idx; + } + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + // Invoke kernel + hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), + dim3(THREADS_PER_BLOCK), 0, 0, static_cast(ptrA)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} +#endif + +/** + * Test Description + * ------------------------ + * - Allocate multiple non-contiguous physical memory chunks + * and map it to contiguous virtual address range. After setting + * 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 + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +#if HT_NVIDIA +// This test is disabled. Will be enabled once VMM feature is fully +// available +TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") { + size_t granularity = 0; + constexpr int numOfBuffers = NUM_OF_BUFFERS; + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle[NUM_OF_BUFFERS]; + // Allocate 3 physical memory chunks + for (int count = 0; count < numOfBuffers; count++) { + HIP_CHECK(hipMemCreate(&handle[count], size_mem, &prop, 0)); + } + // Allocate virtual address range for all the memory chunks + hipDeviceptr_t ptrA; + HIP_CHECK(hipMemAddressReserve(&ptrA, (numOfBuffers * size_mem), 0, 0, 0)); + for (int idx = 0; idx < numOfBuffers; idx++) { + uint64_t uiptr = reinterpret_cast(ptrA); + uiptr = uiptr + idx * size_mem; + HIP_CHECK(hipMemMap(reinterpret_cast(uiptr), size_mem, 0, + handle[idx], 0)); + HIP_CHECK(hipMemRelease(handle[idx])); + } + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // 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); + // Fill Data + for (size_t idx = 0; idx < (numOfBuffers * N); idx++) { + A_h[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)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, numOfBuffers * buffer_size)); + HIP_CHECK(hipDeviceSynchronize()); + // Validate Results + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); + for (int idx = 0; idx < numOfBuffers; idx++) { + uint64_t uiptr = reinterpret_cast(ptrA); + uiptr = uiptr + idx * size_mem; + HIP_CHECK(hipMemUnmap(reinterpret_cast(uiptr), size_mem)); + } + HIP_CHECK(hipMemAddressFree(ptrA, (numOfBuffers * size_mem))); +} +#endif + +/** + * Test Description + * ------------------------ + * - (Check if the VMM address can be memset) Map a physical chunk + * to the VMM address range. Memset the VMM address range with initial + * value. Validate. + * ------------------------ + * - catch\unit\memory\hipMemCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemCreate_ChkWithMemset") { + size_t granularity = 0; + constexpr int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + constexpr int init_val = 0; + int deviceId = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + std::vector A_h(N); + HIP_CHECK(hipMemset(ptrA, init_val, buffer_size)); + HIP_CHECK(hipMemcpyDtoH(A_h.data(), ptrA, buffer_size)); + for (int idx = 0; idx < N; idx++) { + REQUIRE(A_h[idx] == init_val); + } + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + HIP_CHECK(hipMemRelease(handle)); +} + +/** + * Test Description + * ------------------------ + * - Negative Tests + * ------------------------ + * - catch\unit\memory\hipMemCreate.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemCreate_Negative") { + size_t granularity = 0; + int deviceId = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + 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)); + + SECTION("Nullptr to handle") { + REQUIRE(hipMemCreate(nullptr, granularity, &prop, 0) == + hipErrorInvalidValue); + } + + SECTION("Nullptr to prop") { + REQUIRE(hipMemCreate(&handle, granularity, nullptr, 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); + } + + SECTION("pass location as invalid") { + prop.location.type = hipMemLocationTypeInvalid; + REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == + hipErrorInvalidValue); + } + + SECTION("non multiple of granularity") { + 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); + } + + 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); + } +} diff --git a/catch/unit/memory/hipMemGetAllocationGranularity.cc b/catch/unit/memory/hipMemGetAllocationGranularity.cc new file mode 100644 index 0000000000..5f5821de48 --- /dev/null +++ b/catch/unit/memory/hipMemGetAllocationGranularity.cc @@ -0,0 +1,179 @@ +/* +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 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. +*/ + +/** + * @addtogroup hipMemGetAllocationGranularity hipMemGetAllocationGranularity + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemGetAllocationGranularity (size_t* granularity, + * const hipMemAllocationProp* prop, + * hipMemAllocationGranularity_flags option)` - + * Calculates either the minimal or recommended granularity. + */ + +#include +#include +#include +#include "hip_vmm_common.hh" + +/** + local function to invoke hipMemGetAllocationGranularity. + */ +void getGranularity(size_t *granularity, + hipMemAllocationGranularity_flags option, + int device) { + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(granularity, &prop, option)); +} + +/** + * Test Description + * ------------------------ + * - Functional Test to get granularity size for + * hipMemAllocationGranularityMinimum option. + * ------------------------ + * - catch\unit\memory\hipMemGetAllocationGranularity.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemGetAllocationGranularity_MinGranularity") { + size_t granularity = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, 0)); + checkVMMSupported(device) + getGranularity(&granularity, hipMemAllocationGranularityMinimum, 0); + REQUIRE(granularity > 0); +} + +/** + * Test Description + * ------------------------ + * - Functional Test to get granularity size for + * hipMemAllocationGranularityRecommended option. + * ------------------------ + * - catch\unit\memory\hipMemGetAllocationGranularity.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemGetAllocationGranularity_RecommendedGranularity") { + size_t granularity = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, 0)); + checkVMMSupported(device) + getGranularity(&granularity, hipMemAllocationGranularityRecommended, 0); + REQUIRE(granularity > 0); +} + +/** + * Test Description + * ------------------------ + * - Functional Test to get granularity size for + * hipMemAllocationGranularityMinimum option for all GPUs. + * ------------------------ + * - catch\unit\memory\hipMemGetAllocationGranularity.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemGetAllocationGranularity_AllGPUs") { + 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); + REQUIRE(granularity > 0); + } +} + +/** + * Test Description + * ------------------------ + * - Negative Tests + * ------------------------ + * - catch\unit\memory\hipMemGetAllocationGranularity.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemGetAllocationGranularity_NegativeTests") { + size_t granularity = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, 0)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = 0; // Current Devices + + SECTION("Granularity is nullptr") { + REQUIRE(hipErrorInvalidValue == + hipMemGetAllocationGranularity(nullptr, &prop, + hipMemAllocationGranularityMinimum)); + } + SECTION("Prop is nullptr") { + REQUIRE(hipErrorInvalidValue == + hipMemGetAllocationGranularity(&granularity, nullptr, + hipMemAllocationGranularityMinimum)); + } +#if HT_NVIDIA + SECTION("flag is invalid") { + REQUIRE(hipErrorInvalidValue == + hipMemGetAllocationGranularity(&granularity, &prop, + (hipMemAllocationGranularity_flags)0xff)); + } +#endif + 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)); + } + SECTION("device id < lowest device id") { + prop.location.id = -1; // set to non existing device + REQUIRE(hipErrorInvalidValue == + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + } + SECTION("allocation type as invalid") { + prop.type = hipMemAllocationTypeInvalid; + REQUIRE(hipErrorInvalidValue == + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + } + SECTION("location type as invalid") { + prop.location.type = hipMemLocationTypeInvalid; + REQUIRE(hipErrorInvalidValue == + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + } +} diff --git a/catch/unit/memory/hipMemGetAllocationPropertiesFromHandle.cc b/catch/unit/memory/hipMemGetAllocationPropertiesFromHandle.cc new file mode 100644 index 0000000000..5c348a85f8 --- /dev/null +++ b/catch/unit/memory/hipMemGetAllocationPropertiesFromHandle.cc @@ -0,0 +1,119 @@ +/* +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 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. +*/ + +#include +#include "hip_vmm_common.hh" + +#define DATA_SIZE (1 << 13) + +/** + * @addtogroup hipMemGetAllocationPropertiesFromHandle hipMemGetAllocationPropertiesFromHandle + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemGetAllocationPropertiesFromHandle(hipMemAllocationProp* prop, + * hipMemGenericAllocationHandle_t handle)` - + * Retrieve the property structure of the given handle. + */ + +/** + * Test Description + * ------------------------ + * - Functional test to verify the values of hipMemAllocationProp properties. + * ------------------------ + * - catch\unit\memory\hipMemGetAllocationPropertiesFromHandle.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") { + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, 0)); + checkVMMSupported(device) + hipMemGenericAllocationHandle_t handle; + hipMemAllocationProp prop = {}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; + // create a temp prop structure. + hipMemAllocationProp prop_temp = {}; + size_t granularity = 0; + int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + HIP_CHECK(hipMemGetAllocationGranularity + (&granularity, &prop, hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + 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 + HIP_CHECK(hipMemGetAllocationPropertiesFromHandle(&prop_temp, handle)); + REQUIRE(prop_temp.type == prop.type); + REQUIRE(prop_temp.location.type == prop.location.type); + REQUIRE(prop_temp.location.id == prop.location.id); + HIP_CHECK(hipMemRelease(handle)); +} + +/** + * Test Description + * ------------------------ + * - Negative Tests. + * ------------------------ + * - catch\unit\memory\hipMemGetAllocationPropertiesFromHandle.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_Negative") { + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, 0)); + checkVMMSupported(device) + hipMemGenericAllocationHandle_t handle; + hipMemAllocationProp prop = {}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; + // create a temp prop structure. + hipMemAllocationProp prop_temp = {}; + size_t granularity = 0; + int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + HIP_CHECK(hipMemGetAllocationGranularity + (&granularity, &prop, hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + 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); + } + + SECTION("null handle") { + prop.location.type = hipMemLocationTypeInvalid; + REQUIRE(hipMemGetAllocationPropertiesFromHandle(&prop_temp, nullptr) + == hipErrorInvalidValue); + } + HIP_CHECK(hipMemRelease(handle)); +} diff --git a/catch/unit/memory/hipMemMap.cc b/catch/unit/memory/hipMemMap.cc new file mode 100644 index 0000000000..d9b2d13a26 --- /dev/null +++ b/catch/unit/memory/hipMemMap.cc @@ -0,0 +1,650 @@ +/* +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 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. +*/ +/** + * @addtogroup hipMemMap hipMemMap + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemMap (void* ptr, + * size_t size, + * size_t offset, + * hipMemGenericAllocationHandle_t handle, + * unsigned long long flags)` - + * Maps an allocation handle to a reserved virtual address range. + */ + +#include +#include "hip_vmm_common.hh" + +constexpr int N = (1 << 13); +constexpr int num_buf = 3; +constexpr int initializer = 0; + +/** + * Test Description + * ------------------------ + * - Check if a physical chunk can be mapped/unmapped to same + * vmm address range repeatedly. This test validates physical memory + * euse using same vmm range. + * ------------------------ + * - catch\unit\memory\hipMemMap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemMap_SameMemoryReuse") { + constexpr int iterations = 20; + size_t granularity = 0; + size_t buffer_size = N * sizeof(int); + int deviceId = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate host memory and intialize data + std::vector A_h(N), B_h(N), C_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + C_h[idx] = idx*idx; + } + // Allocate a physical memory chunk + HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); + // Allocate num_buf virtual address ranges + hipDeviceptr_t ptrA; + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + for (int i = 0; i < iterations; i++) { + std::fill(B_h.begin(), B_h.end(), initializer); + HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); + // Set access to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); +#if HT_NVIDIA + 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())); +#endif + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + } + // Release resources + HIP_CHECK(hipMemRelease(handle)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Check if a physical chunk can be mapped/unmapped for multiple + * vmm addresses. This test validates physical memory reuse using + * different vmm ranges. + * ------------------------ + * - catch\unit\memory\hipMemMap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") { + size_t granularity = 0; + size_t buffer_size = N * sizeof(int); + int deviceId = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate host memory and intialize data + std::vector A_h(N), B_h(N), C_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + C_h[idx] = idx*idx; + } + // Allocate a physical memory chunk + HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); + // Allocate num_buf virtual address ranges + hipDeviceptr_t ptrA[num_buf]; + for (int buf = 0; buf < num_buf; buf++) { + HIP_CHECK(hipMemAddressReserve(&ptrA[buf], size_mem, 0, 0, 0)); + } + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + for (int buf = 0; buf < num_buf; buf++) { + std::fill(B_h.begin(), B_h.end(), initializer); + HIP_CHECK(hipMemMap(ptrA[buf], size_mem, 0, handle, 0)); + // Set access to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA[buf], size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemcpyHtoD(ptrA[buf], A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); +#if HT_NVIDIA + square_kernel <<>>( + static_cast(ptrA[buf])); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size)); + HIP_CHECK(hipStreamSynchronize(0)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data())); +#endif + HIP_CHECK(hipMemUnmap(ptrA[buf], size_mem)); + } + // Release resources + HIP_CHECK(hipMemRelease(handle)); + for (int buf = 0; buf < num_buf; buf++) { + HIP_CHECK(hipMemAddressFree(ptrA[buf], size_mem)); + } +} + +/** + * Test Description + * ------------------------ + * - Check if a physical chunk can be mapped to multiple + * vmm addresses at the same time and check data values integrity + * between different VMMs. + * ------------------------ + * - catch\unit\memory\hipMemMap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemMap_PhysicalMemory_Map2MultVMMs") { + size_t granularity = 0; + size_t buffer_size = N * sizeof(int); + int deviceId = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate host memory and intialize data + std::vector A_h(N), B_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + // Allocate a physical memory chunk + HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); + // Allocate num_buf virtual address ranges + hipDeviceptr_t ptrA[num_buf]; + for (int buf = 0; buf < num_buf; buf++) { + HIP_CHECK(hipMemAddressReserve(&ptrA[buf], size_mem, 0, 0, 0)); + } + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + for (int buf = 0; buf < num_buf; buf++) { + HIP_CHECK(hipMemMap(ptrA[buf], size_mem, 0, handle, 0)); + } + // Copy data to VMM via ptrA[0] + HIP_CHECK(hipMemSetAccess(ptrA[0], size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemcpyHtoD(ptrA[0], A_h.data(), buffer_size)); + // Validate the data contained in VMM using ptrA[0], ptrA[1], + // ......, ptrA[num_buf-1] + for (int buf = 0; buf < num_buf; buf++) { + std::fill(B_h.begin(), B_h.end(), initializer); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + } + // Release resources + HIP_CHECK(hipMemRelease(handle)); + for (int buf = 0; buf < num_buf; buf++) { + HIP_CHECK(hipMemAddressFree(ptrA[buf], size_mem)); + } +} + +/** + * Test Description + * ------------------------ + * - Check if a physical chunk can be mapped/unmapped for + * multiple vmm addresses. This test validates physical memory + * reuse using different vmm ranges on multiple devices. + * ------------------------ + * - catch\unit\memory\hipMemMap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_MultiDev") { + int devicecount = 0; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + if (devicecount < 2) { + HipTest::HIP_SKIP_TEST("Machine is Single GPU. Skipping Test.."); + return; + } + size_t granularity = 0; + size_t buffer_size = N * sizeof(int); + for (int devX = 0; devX < devicecount; devX++) { + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, devX)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate host memory and intialize data + std::vector A_h(N), B_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + // Allocate a physical memory chunk + HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); + // Allocate devicecount virtual address ranges + std::vector ptrA(devicecount); + for (int devY = 0; devY < devicecount; devY++) { + HIP_CHECK(hipMemAddressReserve(&ptrA[devY], size_mem, 0, 0, 0)); + } + for (int devY = 0; devY < devicecount; devY++) { + hipDevice_t deviceToTest; + HIP_CHECK(hipDeviceGet(&deviceToTest, devY)); + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = deviceToTest; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + HIP_CHECK(hipSetDevice(devY)); + std::fill(B_h.begin(), B_h.end(), initializer); + HIP_CHECK(hipMemMap(ptrA[devY], size_mem, 0, handle, 0)); + // Set access to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA[devY], size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemcpyHtoD(ptrA[devY], A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[devY], buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipMemUnmap(ptrA[devY], size_mem)); + } + HIP_CHECK(hipSetDevice(0)); // set the device back to 0. + // Release resources + HIP_CHECK(hipMemRelease(handle)); + for (int devY = 0; devY < devicecount; devY++) { + HIP_CHECK(hipMemAddressFree(ptrA[devY], size_mem)); + } + } +} + +/** + * Test Description + * ------------------------ + * - Check if different physical chunk can be mapped/unmapped + * for single vmm address. This test validates VMM memory reuse + * using different physical ranges. + * ------------------------ + * - catch\unit\memory\hipMemMap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") { + size_t granularity = 0; + size_t buffer_size = N * sizeof(int); + int deviceId = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle[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; + } + // Allocate a physical memory chunk + for (int buf = 0; buf < num_buf; buf++) { + HIP_CHECK(hipMemCreate(&handle[buf], size_mem, &prop, 0)); + } + // Allocate num_buf virtual address ranges + hipDeviceptr_t ptrA; + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Map ptrA to physical chunk + for (int buf = 0; buf < num_buf; buf++) { + std::fill(B_h.begin(), B_h.end(), initializer); + HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle[buf], 0)); + // Set access to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); +#if HT_NVIDIA + 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())); +#endif + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + } + // Release resources + for (int buf = 0; buf < num_buf; buf++) { + HIP_CHECK(hipMemRelease(handle[buf])); + } + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Check if different physical chunk allocated in different devices + * can be mapped/unmapped to single vmm address. This test validates VMM + * memory reuse using different physical ranges. + * ------------------------ + * - catch\unit\memory\hipMemMap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_MultiGPU") { + int deviceId = 0, devicecount = 0; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + if (devicecount < 2) { + HipTest::HIP_SKIP_TEST("Machine is Single GPU. Skipping Test.."); + return; + } + size_t granularity = 0; + size_t buffer_size = N * sizeof(int); + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + std::vector handle(devicecount); + // Allocate host memory and intialize data + std::vector A_h(N), B_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + // Allocate a physical memory chunk + for (int dev = 0; dev < devicecount; dev++) { + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, dev)); + prop.location.id = device; + HIP_CHECK(hipMemCreate(&handle[dev], size_mem, &prop, 0)); + } + // Allocate devicecount virtual address ranges + hipDeviceptr_t ptrA; + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); + // Map ptrA to physical chunk + SECTION("Set Access of VMM to Different GPU") { + for (int dev = 0; dev < devicecount; dev++) { + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, dev)); + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + HIP_CHECK(hipSetDevice(dev)); + std::fill(B_h.begin(), B_h.end(), initializer); + HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle[dev], 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)); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + } + } + SECTION("Set Access of VMM to default GPU") { + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + for (int dev = 0; dev < devicecount; dev++) { + std::fill(B_h.begin(), B_h.end(), initializer); + HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle[dev], 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)); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + } + } + HIP_CHECK(hipSetDevice(0)); + // Release resources + for (int dev = 0; dev < devicecount; dev++) { + HIP_CHECK(hipMemRelease(handle[dev])); + } + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Check if a partial part of a physical chunk can be mapped/unmapped + * to a smaller vmm address. + * ------------------------ + * - catch\unit\memory\hipMemMap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemMap_MapPartialPhysicalMem") { + int deviceId = 0; + size_t granularity = 0; + size_t buffer_size = N * sizeof(int); + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate host memory and intialize data + std::vector A_h(N), B_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + // Allocate a bigger physical memory chunk of twice size_mem + HIP_CHECK(hipMemCreate(&handle, 2*size_mem, &prop, 0)); + // Allocate virtual address range of size size_mem + hipDeviceptr_t ptrA; + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + std::fill(B_h.begin(), B_h.end(), initializer); + HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + // Release resources + HIP_CHECK(hipMemRelease(handle)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Check if a partial part of a VMM range can be mapped/unmapped + * to a physical address. + * ------------------------ + * - catch\unit\memory\hipMemMap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") { + int deviceId = 0; + size_t granularity = 0; + size_t buffer_size = N * sizeof(int); + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate host memory and intialize data + std::vector A_h(N), B_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + // Allocate a bigger physical memory chunk of size_mem + 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)); + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + std::fill(B_h.begin(), B_h.end(), initializer); + HIP_CHECK(hipMemMap(ptrA , size_mem, 0, handle, 0)); + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + // Release resources + HIP_CHECK(hipMemRelease(handle)); + HIP_CHECK(hipMemAddressFree(ptrA, 2*size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Negative Argument Tests + * ------------------------ + * - catch\unit\memory\hipMemMap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemMap_negative") { + size_t granularity = 0; + size_t buffer_size = N * sizeof(int); + int deviceId = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + 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)); + + SECTION("nullptr to ptrA") { + REQUIRE(hipMemMap(nullptr, size_mem, 0, handle, 0) == + hipErrorInvalidValue); + } + + SECTION("pass zero to size") { + REQUIRE(hipMemMap(&ptrA, 0, 0, handle, 0) == hipErrorInvalidValue); + } + + SECTION("pass negative to offset") { + REQUIRE(hipMemMap(&ptrA, size_mem, -1, handle, 0) == + hipErrorInvalidValue); + } + HIP_CHECK(hipMemRelease(handle)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} diff --git a/catch/unit/memory/hipMemRelease.cc b/catch/unit/memory/hipMemRelease.cc new file mode 100644 index 0000000000..d43647da24 --- /dev/null +++ b/catch/unit/memory/hipMemRelease.cc @@ -0,0 +1,48 @@ +/* +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 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. +*/ + +/** + * @addtogroup hipMemRelease hipMemRelease + * @{ + * @ingroup MemoryTest + * `hipMemRelease(hipMemGenericAllocationHandle_t handle)` - + * Release a memory handle representing a memory allocation which was previously + * allocated through hipMemCreate. + */ + +#include + +/** + * Test Description + * ------------------------ + * - Negative Tests + * ------------------------ + * - catch\unit\memory\hipMemRelease.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemRelease_negative") { + SECTION("Nullptr to handle") { + REQUIRE(hipMemRelease(nullptr) == hipErrorInvalidValue); + } +} diff --git a/catch/unit/memory/hipMemRetainAllocationHandle.cc b/catch/unit/memory/hipMemRetainAllocationHandle.cc new file mode 100644 index 0000000000..29910c61aa --- /dev/null +++ b/catch/unit/memory/hipMemRetainAllocationHandle.cc @@ -0,0 +1,144 @@ +/* +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 hipMemRetainAllocationHandle hipMemRetainAllocationHandle + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemRetainAllocationHandle(hipMemGenericAllocationHandle_t* handle, + * void* addr)` - + * Returns the allocation handle of the backing memory allocation given the address. + */ + +#include +#include +#include "hip_vmm_common.hh" + +#define DATA_SIZE (1 << 13) + +/** + * Test Description + * ------------------------ + * - Create a VM mapped to physical memory. Input addr to + * hipMemRetainAllocationHandle and validate the handle. + * ------------------------ + * - catch\unit\memory\hipMemRetainAllocationHandle.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemRetainAllocationHandle_SetGet") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + // Test hipMemRetainAllocationHandle + hipMemGenericAllocationHandle_t gethandle; + // Check beginning of VMM ptr + HIP_CHECK(hipMemRetainAllocationHandle(&gethandle, + reinterpret_cast(ptrA))); + REQUIRE(gethandle == handle); + HIP_CHECK(hipMemRelease(handle)); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Negative Tests + * ------------------------ + * - catch\unit\memory\hipMemRetainAllocationHandle.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemRetainAllocationHandle_NegTst") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + // Test hipMemRetainAllocationHandle + hipMemGenericAllocationHandle_t gethandle; + SECTION("nullptr handle") { + REQUIRE(hipMemRetainAllocationHandle(nullptr, + reinterpret_cast(ptrA)) == hipErrorInvalidValue); + } + SECTION("nullptr Vmm ptr") { + 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); + 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); + } + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} diff --git a/catch/unit/memory/hipMemSetGetAccess.cc b/catch/unit/memory/hipMemSetGetAccess.cc new file mode 100644 index 0000000000..220c2c11ba --- /dev/null +++ b/catch/unit/memory/hipMemSetGetAccess.cc @@ -0,0 +1,1547 @@ +/* +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 hipMemSetAccess hipMemSetAccess + * @{ + * @ingroup MemoryTest + * `hipError_t hipMemSetAccess (void* ptr, + * size_t size, + * const hipMemAccessDesc* desc, + * size_t count)` - + * 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 +#define THREADS_PER_BLOCK 512 +#define NUM_OF_BUFFERS 3 +#define DATA_SIZE (1 << 13) +#define NEW_DATA_SIZE (2*DATA_SIZE) +constexpr int initializer = 0; +/** + * Test Description + * ------------------------ + * - Create a VM mapped to physical memory. Set the access of the + * VMM chunk to device 0. Validate that flags = hipMemAccessFlagsProtReadWrite + * is returned by hipMemGetAccess() when location is set to device 0. + * Validate that flags = hipMemAccessFlagsProtNone is returned by + * hipMemGetAccess() when location is set to device 1. + * ------------------------ + * - catch\unit\memory\hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +#if HT_NVIDIA +// This test is disabled. Will be enabled once VMM feature is fully +// available +TEST_CASE("Unit_hipMemSetAccess_SetGet") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + // Validate using hipMemGetAccess() + hipMemLocation location; + location.type = hipMemLocationTypeDevice; + location.id = device; + unsigned long long flags = 0; // NOLINT + HIP_CHECK(hipMemGetAccess(&flags, &location, ptrA)); + REQUIRE(flags == hipMemAccessFlagsProtReadWrite); + int devicecount = 0; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + if (devicecount > 1) { + flags = 0; + HIP_CHECK(hipDeviceGet(&device, 1)); + location.type = hipMemLocationTypeDevice; + location.id = device; + HIP_CHECK(hipMemGetAccess(&flags, &location, ptrA)); + REQUIRE(flags == hipMemAccessFlagsProtNone); + } + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Mult Device Functional Test: Create a VM mapped to physical memory. + * Set the access of the VMM chunk to both device 0 and device 1. + * Validate that flags = hipMemAccessFlagsProtReadWrite is returned by + * hipMemGetAccess() when location is set to device 0. Validate that + * flags = hipMemAccessFlagsProtReadWrite is returned by hipMemGetAccess() + * when location is set to device 1. + * ------------------------ + * - catch\unit\memory\hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemSetAccess_MultDevSetGet") { + size_t granularity = 0; + constexpr int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + int deviceId = 0, value = 0, device_count = 0; + hipDevice_t device0, device1; + HIP_CHECK(hipGetDeviceCount(&device_count)); + if (device_count < 2) { + HipTest::HIP_SKIP_TEST("Need 2 GPUs to run test. Skipping Test.."); + return; + } + + HIP_CHECK(hipDeviceGet(&device0, deviceId)); + checkVMMSupported(device0) + HIP_CHECK(hipDeviceGet(&device1, (deviceId + 1))); + checkVMMSupported(device1) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device0; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc[2]; + accessDesc[0].location.type = hipMemLocationTypeDevice; + accessDesc[0].location.id = device0; + accessDesc[0].flags = hipMemAccessFlagsProtReadWrite; + accessDesc[1].location.type = hipMemLocationTypeDevice; + accessDesc[1].location.id = device1; + accessDesc[1].flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 and 1 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc[0], 2)); + // Validate using hipMemGetAccess() + hipMemLocation location; + location.type = hipMemLocationTypeDevice; + location.id = device0; + unsigned long long flags = 0; // NOLINT + HIP_CHECK(hipMemGetAccess(&flags, &location, ptrA)); + REQUIRE(flags == hipMemAccessFlagsProtReadWrite); + location.type = hipMemLocationTypeDevice; + location.id = device1; + flags = 0; + HIP_CHECK(hipMemGetAccess(&flags, &location, ptrA)); + REQUIRE(flags == hipMemAccessFlagsProtReadWrite); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} +#endif + +/** + * Test Description + * ------------------------ + * - Create a VM mapped to physical memory. Set the access of the VMM chunk + * 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 + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemSetAccess_EntireVMMRangeSetGet") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + // 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + // Validate hipMemGetAccess() + hipMemLocation location; + location.type = hipMemLocationTypeDevice; + location.id = device; + unsigned long long flags = 0; // NOLINT + HIP_CHECK(hipMemGetAccess(&flags, &location, ptrA)); + REQUIRE(flags == hipMemAccessFlagsProtReadWrite); + uint64_t uiptr = reinterpret_cast(ptrA); + uiptr += (size_mem - 1); + HIP_CHECK(hipMemGetAccess(&flags, &location, + reinterpret_cast(uiptr))); + REQUIRE(flags == hipMemAccessFlagsProtReadWrite); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Negative Tests + * ------------------------ + * - catch\unit\memory\hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemGetAccess_NegTst") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + // 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + // Validate hipMemGetAccess() -ve scenarios + hipMemLocation location; + location.type = hipMemLocationTypeDevice; + location.id = device; + unsigned long long flags = 0; // NOLINT + hipError_t status = hipSuccess; + status = hipMemGetAccess(nullptr, &location, ptrA); + REQUIRE(status == hipErrorInvalidValue); + status = hipMemGetAccess(&flags, nullptr, ptrA); + REQUIRE(status == hipErrorInvalidValue); + uint64_t uiptr = reinterpret_cast(ptrA); + uiptr += size_mem; + status = hipMemGetAccess(&flags, &location, reinterpret_cast(uiptr)); + REQUIRE(status == hipErrorInvalidValue); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Test VMM functionality on multiple device. In each device, create + * a VM mapped to physical memory of the device, copy test data to the VM + * address range, launch a kernel to perform operation on the data and + * validate the result. + * ------------------------ + * - catch\unit\memory\hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +#if HT_NVIDIA +// This test is disabled. Will be enabled once VMM feature is fully +// available +TEST_CASE("Unit_hipMemSetAccess_FuncTstOnMultDev") { + size_t granularity = 0; + constexpr int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + int deviceId = 0, devicecount = 0; + hipDevice_t device; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + for (deviceId = 0; deviceId < devicecount; deviceId++) { + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + // Allocate physical memory + hipDeviceptr_t ptrA; + hipMemGenericAllocationHandle_t handle; + HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); + // Allocate virtual address range + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); + HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU deviceId + std::vector A_h(N), B_h(N); + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + for (int idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + 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)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + HIP_CHECK(hipDeviceSynchronize()); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + } +} +#endif + +/** + * Test Description + * ------------------------ + * - Allocate physical memory and map it to a VMM range. + * Access (Read/Write) the virtual pointer directly on host. + * Ensure this behavior for all devices on host. + * ------------------------ + * - catch\unit\memory\hipMemMap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemSetAccess_AccessDirectlyFromHost") { + size_t granularity = 0; + constexpr int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + int devicecount = 0; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + if (devicecount < 2) { + HipTest::HIP_SKIP_TEST("Machine is Single GPU. Skipping Test.."); + return; + } + for (int dev = 0; dev < devicecount; dev++) { + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, dev)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate a physical memory chunk + HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); + // Allocate num_buf virtual address ranges + hipDeviceptr_t ptrA; + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + int* vptr = reinterpret_cast(ptrA); + for (int idx = 0; idx < N; idx++) { + *(vptr + idx) = idx; + } + // validate + for (int idx = 0; idx < N; idx++) { + REQUIRE(*(vptr + idx) == idx); + } + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + // Release resources + HIP_CHECK(hipMemRelease(handle)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + } +} + +/** + * Test Description + * ------------------------ + * - Create a virtual memnory chunk and set the property of + * the range to read/write. Write to the memory chunk. Change + * the property of the range to read only. Check if the memory + * range can be read. + * ------------------------ + * - catch\unit\memory\hipMemMap.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemSetAccess_ChangeAccessProp") { + size_t granularity = 0; + constexpr int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + int dev = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, dev)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK( + hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; // Allocate host memory and intialize data + std::vector A_h(N), B_h(N); + // Initialize with data + for (size_t idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + // Allocate a physical memory chunk + HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0)); + // Allocate num_buf virtual address ranges + hipDeviceptr_t ptrA; + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); + HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0)); + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + SECTION("Change ReadWrite to Read") { + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + // Change property of virtual memory range to read only + accessDesc.flags = hipMemAccessFlagsProtRead; + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + // validate + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + } + SECTION("Change Read to ReadWrite") { + accessDesc.flags = hipMemAccessFlagsProtRead; + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + // Change property of virtual memory range to read only + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + } + SECTION("Change Inaccessible to ReadWrite") { + accessDesc.flags = hipMemAccessFlagsProtNone; + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + // Change property of virtual memory range to read only + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + } +#if HT_NVIDIA + SECTION("Check error while writing on Read-Only memory") { + accessDesc.flags = hipMemAccessFlagsProtRead; + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + REQUIRE(hipErrorInvalidValue == + hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + } + SECTION("Check error while writing on inaccessible memory") { + accessDesc.flags = hipMemAccessFlagsProtNone; + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + REQUIRE(hipErrorInvalidValue == + hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + } +#endif + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + // Release resources + HIP_CHECK(hipMemRelease(handle)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} +/** + * Test Description + * ------------------------ + * - Test Virtual Memory to Unified Memory data transfer. Allocate + * a Virtual Memory chunk and a Unified Memory chunk. Test if data can + * be exchanged between these chunks. + * ------------------------ + * - catch\unit\memory\hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemSetAccess_Vmm2UnifiedMemCpy") { + auto managed = HmmAttrPrint(); + if (managed != 1) { + HipTest::HIP_SKIP_TEST("GPU doesn't support managed memory.Skipping Test.."); + return; + } + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + // 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + int *ptrA_h, *ptrB_h; + HIP_CHECK(hipMallocManaged(&ptrA_h, buffer_size)); + HIP_CHECK(hipMallocManaged(&ptrB_h, buffer_size)); + for (int idx = 0; idx < N; idx++) { + ptrA_h[idx] = idx; + } + HIP_CHECK(hipMemcpyHtoD(ptrA, ptrA_h, buffer_size)); + HIP_CHECK(hipMalloc(&ptrB, buffer_size)); + HIP_CHECK(hipMemcpyDtoD(ptrB, ptrA, buffer_size)); + HIP_CHECK(hipMemcpyDtoH(ptrB_h, ptrB, buffer_size)); + bool bPassed = true; + for (int idx = 0; idx < N; idx++) { + if (ptrB_h[idx] != idx) { + bPassed = false; + break; + } + } + REQUIRE(bPassed == true); + HIP_CHECK(hipFree(ptrB)); + HIP_CHECK(hipFree(ptrA_h)); + HIP_CHECK(hipFree(ptrB_h)); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Test Virtual Memory to Device Memory data transfer. Allocate a Virtual + * Memory chunk and a Device Memory chunk. Test if data can be exchanged + * between these chunks. + * ------------------------ + * - catch\unit\memory\hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemSetAccess_Vmm2DevMemCpy") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + // 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + std::vector A_h(N), B_h(N); + for (int idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMalloc(&ptrB, buffer_size)); + HIP_CHECK(hipMemcpyDtoD(ptrB, ptrA, buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrB, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipFree(ptrB)); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - VM to Device Memory Copy. Allocate a Virtual Memory chunk and a + * Peer Device Memory chunk. Test if data can be exchanged between + * these chunks using hipMemcpyDtoD. + * ------------------------ + * - catch\unit\memory\hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerDevMemCpy") { + size_t granularity = 0; + constexpr int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + int deviceId = 0, value = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + // 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + std::vector A_h(N), B_h(N); + for (int idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + int devicecount = 0; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + // Check Peer Access + for (deviceId = 1; deviceId < devicecount; deviceId++) { + int canAccessPeer = 0; + hipDevice_t device_other; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, deviceId)); + if (0 == canAccessPeer) { + WARN("Machine does not support Peer Access\n"); + break; + } + HIP_CHECK(hipDeviceGet(&device_other, deviceId)); + 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"); + break; + } + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceptr_t dptr_peer; + HIP_CHECK(hipMalloc(&dptr_peer, buffer_size)); + HIP_CHECK(hipMemcpyDtoD(dptr_peer, ptrA, buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), dptr_peer, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipFree(dptr_peer)); + } + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - VM to Device Memory Copy: Allocate a Virtual Memory chunk and + * a Peer Device Memory chunk. Test if data can be exchanged between + * these chunks using hipMemcpyPeer. + * ------------------------ + * - catch\unit\memory\hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemSetAccess_Vmm2PeerPeerMemCpy") { + size_t granularity = 0; + constexpr int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + int deviceId = 0, value = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + // 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + std::vector A_h(N), B_h(N); + for (int idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + int devicecount = 0; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + // Check Peer Access + for (deviceId = 1; deviceId < devicecount; deviceId++) { + std::fill(B_h.begin(), B_h.end(), initializer); + int canAccessPeer = 0; + hipDevice_t device_other; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, deviceId)); + if (0 == canAccessPeer) { + WARN("Machine does not support Peer Access\n"); + break; + } + HIP_CHECK(hipDeviceGet(&device_other, deviceId)); + 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"); + break; + } + HIP_CHECK(hipSetDevice(deviceId)); + hipDeviceptr_t dptr_peer; + HIP_CHECK(hipMalloc(&dptr_peer, buffer_size)); + HIP_CHECK(hipMemcpyPeer(dptr_peer, deviceId, ptrA, 0, buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), dptr_peer, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipFree(dptr_peer)); + } + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - VM to VM copy: Allocate memory and map it to an address space in + * device 0(PtrA). Allocate another chunk of memory and map it to an + * address space in device 0(PtrB). Check if data can be copied from + * PtrA -> PtrB using hipMemcpy. + * ------------------------ + * - catch\unit\memory\hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMMemCpy") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + // Allocate physical memory + hipMemGenericAllocationHandle_t handle1, handle2; + HIP_CHECK(hipMemCreate(&handle1, size_mem, &prop, 0)); + HIP_CHECK(hipMemCreate(&handle2, size_mem, &prop, 0)); + // Allocate virtual address range + hipDeviceptr_t ptrA, ptrB; + HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0)); + HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem, 0, 0, 0)); + HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle1, 0)); + HIP_CHECK(hipMemMap(ptrB, size_mem, 0, handle2, 0)); + HIP_CHECK(hipMemRelease(handle1)); + HIP_CHECK(hipMemRelease(handle2)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the addresses accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + HIP_CHECK(hipMemSetAccess(ptrB, size_mem, &accessDesc, 1)); + std::vector A_h(N), B_h(N); + for (int idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoD(ptrB, ptrA, buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrB, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemUnmap(ptrB, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrB, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Functional Test: Allocate memory and map it to an address space in + * device 0(PtrA). Allocate another chunk of memory and map it to an + * address space in device 1(PtrB). Check if data can be copied from + * PtrA -> PtrB using hipMemcpyPeer. + * ------------------------ + * - catch\unit\memory\hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemSetAccess_Vmm2VMMInterDevMemCpy") { + size_t granularity = 0; + constexpr int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + int deviceId = 0, value = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1)); + std::vector A_h(N), B_h(N); + for (int idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size)); + int devicecount = 0; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + for (deviceId = 1; deviceId < devicecount; deviceId++) { + int canAccessPeer = 0; + hipDevice_t device_other; + HIP_CHECK(hipDeviceCanAccessPeer(&canAccessPeer, 0, deviceId)); + if (0 == canAccessPeer) { + WARN("Machine does not support Peer Access\n"); + break; + } + std::fill(B_h.begin(), B_h.end(), initializer); + HIP_CHECK(hipDeviceGet(&device_other, deviceId)); + 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"); + break; + } + HIP_CHECK(hipSetDevice(deviceId)); + hipMemAllocationProp prop_loc{}; + prop_loc.type = hipMemAllocationTypePinned; + 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; + hipMemGenericAllocationHandle_t handle_loc; + // Allocate physical memory + HIP_CHECK(hipMemCreate(&handle_loc, size_mem_loc, &prop_loc, 0)); + // Allocate virtual address range + hipDeviceptr_t ptrB; + HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem_loc, 0, 0, 0)); + HIP_CHECK(hipMemMap(ptrB, size_mem_loc, 0, handle, 0)); + HIP_CHECK(hipMemRelease(handle_loc)); + // Set access + hipMemAccessDesc accessDesc_loc = {}; + accessDesc_loc.location.type = hipMemLocationTypeDevice; + accessDesc_loc.location.id = device_other; + accessDesc_loc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrB, size_mem_loc, &accessDesc_loc, 1)); + HIP_CHECK(hipMemcpyPeer(ptrB, deviceId, ptrA, 0, buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrB, buffer_size)); + REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data())); + HIP_CHECK(hipMemUnmap(ptrB, size_mem_loc)); + HIP_CHECK(hipMemAddressFree(ptrB, size_mem_loc)); + } + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); +} + +/** + * Test Description + * ------------------------ + * - Allocate a chunk of memory and map it to device0. Allocate another + * 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 + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +#if HT_NVIDIA +// This test is disabled. Will be enabled once VMM feature is fully +// available +TEST_CASE("Unit_hipMemSetAccess_MapPhysChksFromMulDev") { + int devicecount = 0; + HIP_CHECK(hipGetDeviceCount(&devicecount)); + int numOfBuffers = devicecount; + constexpr int N = DATA_SIZE; + 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)); + REQUIRE(handle != nullptr); + 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) + 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)); + REQUIRE(granularity > 0); + 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)); + } + // Allocate virtual address range for all the memory chunks + hipDeviceptr_t ptrA; + HIP_CHECK(hipMemAddressReserve(&ptrA, total_mem, 0, 0, 0)); + // Map the allocated chunks + for (deviceId = 0; deviceId < numOfBuffers; deviceId++) { + 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)); + HIP_CHECK(hipMemRelease(handle[deviceId])); + // Set access + hipMemAccessDesc accessDesc_loc = {}; + accessDesc_loc.location.type = hipMemLocationTypeDevice; + 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)); + } + 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; + } + 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])); + } + HIP_CHECK(hipMemAddressFree(ptrA, total_mem)); + free(handle); + free(size_mem); +} +#endif + +/** + * 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; + hipDeviceptr_t ptrVmm; + std::vector vhandle; + std::vector vsize; + // allocate initial VMM memory chunk + 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)); + REQUIRE(granularity > 0); + size_t size_rounded = + ((granularity + size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate physical memory + HIP_CHECK(hipMemCreate(&handle, size_rounded, &prop, 0)); + // Store the handle for future reference + vhandle.push_back(handle); + vsize.push_back(size_rounded); + // Allocate virtual address range + HIP_CHECK(hipMemAddressReserve(&ptrVmm, size_rounded, 0, 0, 0)); + HIP_CHECK(hipMemMap(ptrVmm, size_rounded, 0, handle, 0)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU device + HIP_CHECK(hipMemSetAccess(ptrVmm, size_rounded, &accessDesc, 1)); + *ptr = ptrVmm; + current_size_tot += size; + current_size_rounded_tot += size_rounded; + return 0; + } + + public: + 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) { + 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) { + size_t granularity = 0; + if (size <= current_size_tot) { + return -1; + } + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + 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; + hipMemGenericAllocationHandle_t handle; + // Allocate physical memory + HIP_CHECK(hipMemCreate(&handle, size_rounded, &prop, 0)); + // Store the handle for future reference + vhandle.push_back(handle); + vsize.push_back(size_rounded); + // Allocate virtual address range + // 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)); + int idx = 0; + 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)); + } + idx++; + } + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + // Make the address accessible to GPU 0 + HIP_CHECK(hipMemSetAccess(ptrVmm, + (size_rounded + current_size_rounded_tot), + &accessDesc, 1)); + *ptr = ptrVmm; + current_size_tot += size; + current_size_rounded_tot += size_rounded; + return 0; + } +}; +#if HT_NVIDIA +// This test is disabled. Will be enabled once VMM feature is fully +// available +TEST_CASE("Unit_hipMemSetAccess_GrowVMM") { + hipDeviceptr_t ptr; + 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) + // 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)); + 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)); + // Resize the VMM + int Nnew = NEW_DATA_SIZE; + size_t buffer_size_new = Nnew * sizeof(int); + if (-1 == resizeobj.grow_vmm(&ptr, device, buffer_size_new)) { + WARN("Virtual Memory Management Grow Failed"); + return; + } + free(ptrA_h); + 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)); + 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(hipMemcpyDtoH(ptrB_h, ptr, buffer_size_new)); + bool bPassed = true; + for (int idx = 0; idx < Nnew; idx++) { + if (ptrB_h[idx] != idx) { + bPassed = false; + break; + } + } + REQUIRE(bPassed == true); + free(ptrB_h); + free(ptrA_h); + resizeobj.free_vmm(); +} +#endif + +/** + * 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) { + hipDeviceptr_t ptr; + constexpr int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + // 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)); + 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)); + REQUIRE(ptrB_h != nullptr); + HIP_CHECK(hipMemcpyDtoH(ptrB_h, ptr, buffer_size)); + bool bPassed = true; + for (int idx = 0; idx < N; idx++) { + if (ptrB_h[idx] != idx) { + bPassed = false; + break; + } + } + if (bPassed) { + bTestPassed.fetch_and(1); + } else { + bTestPassed.fetch_and(0); + } + free(ptrB_h); + free(ptrA_h); + vmmobj.free_vmm(); +} + +TEST_CASE("Unit_hipMemSetAccess_Multithreaded") { + int deviceId = 0; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + std::thread T[NUM_THREADS]; + for (int i = 0; i < NUM_THREADS; i++) { + T[i] = std::thread(test_thread, device); + } + // Wait until all the threads finish their execution + for (int i = 0; i < NUM_THREADS; i++) { + T[i].join(); + } + REQUIRE(1 == bTestPassed.load()); +} + +#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; + pid_t childpid; + int testResultChild = 0; + int deviceId = 0; + constexpr int N = DATA_SIZE; + size_t buffer_size = N * sizeof(int); + // create pipe descriptors + pipe(fd); + // fork process + childpid = fork(); + if (childpid > 0) { // Parent + close(fd[1]); + hipDeviceptr_t ptr; + hipDevice_t device; + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupportedRetVal(device) + // Create VMM Object of size buffer_size + vmm_resize_class vmmobj(&ptr, device, buffer_size); + // Inititalize Host Buffer + std::vector A_h(N), B_h(N); + for (int idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + // Copy to VMM + HIP_CHECK(hipMemcpyHtoD(ptr, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptr, buffer_size)); + bool bPassed = std::equal(B_h.begin(), B_h.end(), A_h.data()); + vmmobj.free_vmm(); + // parent will wait to read the device cnt + read(fd[0], &testResultChild, sizeof(int)); + if (testResultChild == 0) { + testResult = bPassed & false; + } else { + testResult = bPassed & true; + } + // close the read-descriptor + close(fd[0]); + // wait for child exit + wait(NULL); + } else if (!childpid) { // Child + close(fd[0]); + hipDeviceptr_t ptr; + hipDevice_t device; + + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupportedRetVal(device) + // Create VMM Object of size buffer_size + vmm_resize_class vmmobj(&ptr, device, buffer_size); + // Inititalize Host Buffer + std::vector A_h(N), B_h(N); + for (int idx = 0; idx < N; idx++) { + A_h[idx] = idx; + } + // Copy to VMM + HIP_CHECK(hipMemcpyHtoD(ptr, A_h.data(), buffer_size)); + HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptr, buffer_size)); + int result = 0; + if (true == std::equal(B_h.begin(), B_h.end(), A_h.data())) { + result = 1; + } + vmmobj.free_vmm(); + // send the value on the write-descriptor: + write(fd[1], &result, sizeof(int)); + // close the write descriptor: + close(fd[1]); + exit(0); + } + return testResult; +} + +TEST_CASE("Unit_hipMemSetAccess_MultiProc") { + REQUIRE(true == test_mprocess()); +} +#endif + +/** + * Test Description + * ------------------------ + * - Negative Tests for hipMemSetAccess() + * ------------------------ + * - catch\unit\memory\hipMemSetGetAccess.cc + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemSetAccess_negative") { + 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) + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + hipMemGenericAllocationHandle_t handle; + // Allocate 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(hipMemMap(ptrA, size_mem, 0, handle, 0)); + // Set access + hipMemAccessDesc accessDesc = {}; + accessDesc.location.type = hipMemLocationTypeDevice; + accessDesc.location.id = device; + accessDesc.flags = hipMemAccessFlagsProtReadWrite; + + SECTION("nullptr to ptrA") { + REQUIRE(hipMemSetAccess(nullptr, size_mem, &accessDesc, 1) == + hipErrorInvalidValue); + } + + SECTION("pass zero to size") { + REQUIRE(hipMemSetAccess(&ptrA, 0, &accessDesc, 1) == + hipErrorInvalidValue); + } + + SECTION("pass a size greater than reserved size") { + 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); + } + + SECTION("invalid location type") { + accessDesc.location.type = hipMemLocationTypeInvalid; + REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 1) == + hipErrorInvalidValue); + } + + SECTION("invalid id") { + accessDesc.location.id = -1; + 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); + } + + SECTION("invalid flag") { + accessDesc.flags = static_cast(-1); + REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 1) == + hipErrorInvalidValue); + } + + SECTION(" pass zero to count") { + REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 0) == + hipErrorInvalidValue); + } + + SECTION("pass desc as nullptr") { + 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); + } + + HIP_CHECK(hipMemUnmap(ptrA, size_mem)); + SECTION("unmapped virtual memory") { + REQUIRE(hipMemSetAccess(&ptrA, size_mem, &accessDesc, 1) == + hipErrorInvalidValue); + } + HIP_CHECK(hipMemAddressFree(ptrA, size_mem)); + HIP_CHECK(hipMemRelease(handle)); +} diff --git a/catch/unit/memory/hipMemUnmap.cc b/catch/unit/memory/hipMemUnmap.cc new file mode 100644 index 0000000000..1244b794c6 --- /dev/null +++ b/catch/unit/memory/hipMemUnmap.cc @@ -0,0 +1,88 @@ +/* +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 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. +*/ + +/** + * @addtogroup hipMemUnmap hipMemUnmap + * @{ + * @ingroup MemoryTest + * `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 + * Test requirements + * ------------------------ + * - HIP_VERSION >= 6.1 + */ +TEST_CASE("Unit_hipMemUnmap_negative") { + size_t granularity = 0; + size_t buffer_size = N * sizeof(int); + int deviceId = 0; + hipDevice_t device; + + HIP_CHECK(hipDeviceGet(&device, deviceId)); + checkVMMSupported(device) + + hipMemAllocationProp prop{}; + prop.type = hipMemAllocationTypePinned; + prop.location.type = hipMemLocationTypeDevice; + prop.location.id = device; // Current Devices + + HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop, + hipMemAllocationGranularityMinimum)); + REQUIRE(granularity > 0); + size_t size_mem = + ((granularity + buffer_size - 1) / granularity) * granularity; + + hipMemGenericAllocationHandle_t handle; + 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(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); + } +#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/catch/unit/memory/hip_vmm_common.hh b/catch/unit/memory/hip_vmm_common.hh new file mode 100644 index 0000000000..83d209ac97 --- /dev/null +++ b/catch/unit/memory/hip_vmm_common.hh @@ -0,0 +1,59 @@ +/* +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 + +#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;\ + }\ +} + +#if HT_NVIDIA +constexpr int threadsPerBlk = 64; +/** + Kernel to perform Square of input data. + */ +static __global__ void square_kernel(int* Buff) { + int i = threadIdx.x + blockDim.x * blockIdx.x; + int temp = Buff[i] * Buff[i]; + Buff[i] = temp; +} +#endif