diff --git a/catch/unit/memory/CMakeLists.txt b/catch/unit/memory/CMakeLists.txt index 29f966d9a4..2283d5e558 100644 --- a/catch/unit/memory/CMakeLists.txt +++ b/catch/unit/memory/CMakeLists.txt @@ -82,18 +82,7 @@ if(HIP_PLATFORM MATCHES "amd") hipMemCoherencyTst.cc hipExtMallocWithFlags.cc hipMallocMngdMultiThread.cc - hipArray.cc hipMemVmm.cc - hipMemCreate.cc - hipMemMap.cc - hipMemGetAllocationGranularity.cc - hipMemSetGetAccess.cc - hipMemRetainAllocationHandle.cc - hipMemUnmap.cc - hipMemAddressFree.cc - hipMemAddressReserve.cc - hipMemRelease.cc - hipMemGetAllocationPropertiesFromHandle.cc hipArray.cc hipMemcpyDeviceToDeviceNoCU.cc) if(UNIX) diff --git a/catch/unit/memory/hipMemAddressFree.cc b/catch/unit/memory/hipMemAddressFree.cc deleted file mode 100644 index 6890a26d8c..0000000000 --- a/catch/unit/memory/hipMemAddressFree.cc +++ /dev/null @@ -1,76 +0,0 @@ -/* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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 deleted file mode 100644 index c6a76fc56c..0000000000 --- a/catch/unit/memory/hipMemAddressReserve.cc +++ /dev/null @@ -1,152 +0,0 @@ -/* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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 deleted file mode 100644 index 83f6ff7c01..0000000000 --- a/catch/unit/memory/hipMemCreate.cc +++ /dev/null @@ -1,465 +0,0 @@ -/* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ -/** - * @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) - -/** - Kernel to perform Square of input data. - */ -static __global__ void square_kernel(int* Buff) { - int i = threadIdx.x + blockDim.x * blockIdx.x; - int temp = Buff[i] * Buff[i]; - Buff[i] = temp; -} - -/** - * Test Description - * ------------------------ - * - 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 - */ - -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)); -} - -/** - * 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 - */ -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))); -} - -/** - * 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 deleted file mode 100644 index 5f5821de48..0000000000 --- a/catch/unit/memory/hipMemGetAllocationGranularity.cc +++ /dev/null @@ -1,179 +0,0 @@ -/* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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 deleted file mode 100644 index 5c348a85f8..0000000000 --- a/catch/unit/memory/hipMemGetAllocationPropertiesFromHandle.cc +++ /dev/null @@ -1,119 +0,0 @@ -/* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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 deleted file mode 100644 index d9b2d13a26..0000000000 --- a/catch/unit/memory/hipMemMap.cc +++ /dev/null @@ -1,650 +0,0 @@ -/* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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 deleted file mode 100644 index d43647da24..0000000000 --- a/catch/unit/memory/hipMemRelease.cc +++ /dev/null @@ -1,48 +0,0 @@ -/* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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 deleted file mode 100644 index 29910c61aa..0000000000 --- a/catch/unit/memory/hipMemRetainAllocationHandle.cc +++ /dev/null @@ -1,144 +0,0 @@ -/* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/** - * @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 deleted file mode 100644 index 20910f9803..0000000000 --- a/catch/unit/memory/hipMemSetGetAccess.cc +++ /dev/null @@ -1,1542 +0,0 @@ -/* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -/** - * @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; - -/** - Kernel to perform Square of input data. - */ -static __global__ void square_kernel(int* Buff) { - int i = threadIdx.x + blockDim.x * blockIdx.x; - int temp = Buff[i] * Buff[i]; - Buff[i] = temp; -} - -/** - * Test Description - * ------------------------ - * - 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 - */ -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, 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)); -} - -/** - * 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 - */ -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)); - } -} - -/** - * 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 - */ -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); -} - -/** - * 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; - } -}; - -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(); -} - -/** - * 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 deleted file mode 100644 index 1244b794c6..0000000000 --- a/catch/unit/memory/hipMemUnmap.cc +++ /dev/null @@ -1,88 +0,0 @@ -/* -Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. - -Permission is hereby granted, free of charge, to any person obtaining a copy -of this software and associated documentation files (the "Software"), to deal -in the Software without restriction, including without limitation the rights -to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -copies of the Software, and to permit persons to whom the Software is -furnished to do so, subject to the following conditions: - -The above copyright notice and this permission notice shall be included in -all copies or substantial portions of the Software. - -THE SOFTWARE IS PROVIDED "AS IS", WITHOUT 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)); -}