AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

EXSWHTEC-380 - Implement tests for Virtual Memory Management API functions #448

Browse Source 
Change-Id: Ic766be69fddd0309f7ad4093465494cc14c7c70b


[ROCm/hip-tests commit: 00433d4f87]
This commit is contained in:
Mirza Halilcevic
2024-02-02 12:29:18 +05:30
committed by Rakesh Roy
parent ad4d04fa88
commit d439c6b986
23 changed files with 1043 additions and 787 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip-tests/catch/hipTestMain/config/config_nvidia_linux.json
+3 -1
View File
@@ -92,6 +92,8 @@
"Unit_hipMemcpy_Positive_Synchronization_Behavior",
"=== Below tests fail in external CI for PR https://github.com/ROCm-Developer-Tools/hip-tests/pull/210 ===",
"Unit_StaticAssert_Positive_Basic",
"Unit_StaticAssert_Negative_Basic"
"Unit_StaticAssert_Negative_Basic",
"Unit_hipMemImportFromShareableHandle_Positive_MultiProc",
"Unit_hipMemMapArrayAsync_Positive_Basic"
]
}
projects/hip-tests/catch/hipTestMain/config/config_nvidia_windows.json
+2 -1
View File
@@ -44,6 +44,7 @@
"Performance_hipMemsetD32",
"Performance_hipMemsetD32Async",
"Unit_hipMemcpyParam2D_Positive_Synchronization_Behavior",
"Unit_hipMemcpy_Positive_Synchronization_Behavior"
"Unit_hipMemcpy_Positive_Synchronization_Behavior",
"Unit_hipMemMapArrayAsync_Positive_Basic"
]
}
projects/hip-tests/catch/unit/memory/hipMallocManagedCommon.hh → projects/hip-tests/catch/include/hipMallocManagedCommon.hh
View File
projects/hip-tests/catch/include/hip_test_defgroups.hh
+6 -184
View File
@@ -103,16 +103,9 @@ THE SOFTWARE.
*/
/**
* @defgroup KernelTest Kernel Functions Management
* @{
* This section describes the various kernel functions invocation.
* @}
*/
/**
* @defgroup AtomicsTest Device Atomics
* @defgroup KernelTest Kernel Functions Management
* @{
* This section describes tests for the Device Atomic APIs.
* This section describes the various kernel functions invocation.
* @}
*/
@@ -140,7 +133,8 @@ THE SOFTWARE.
/**
* @defgroup PeerToPeerTest PeerToPeer Device Memory Access
* @{
* This section describes tests for the PeerToPeer device memory access functions of HIP runtime API.
* This section describes tests for the PeerToPeer device memory access functions of HIP runtime
* API.
* @warning PeerToPeer support is experimental.
* @}
*/
@@ -200,124 +194,6 @@ THE SOFTWARE.
* @}
*/
/**
* @defgroup AtomicsTest Device Atomics
* @{
* This section describes tests for the Device Atomic APIs.
*/
/**
* @addtogroup atomicAdd atomicAdd
* @{
* @ingroup AtomicsTest
*/
/**
* Test Description
* ------------------------
* - Compiles atomicAdd with invalid parameters.
* - Compiles the source with specialized Python tool.
* -# Utilizes sub-process to invoke compilation of faulty source.
* -# Performs post-processing of compiler output and counts errors.
* Test source
* ------------------------
* - unit/atomics/CMakeLists.txt
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_atomicAdd_Negative_Parameters") {}
/**
* End doxygen group atomicAdd.
* @}
*/
/**
* @addtogroup atomicSub atomicSub
* @{
* @ingroup AtomicsTest
*/
/**
* Test Description
* ------------------------
* - Compiles atomicSub with invalid parameters.
* - Compiles the source with specialized Python tool.
* -# Utilizes sub-process to invoke compilation of faulty source.
* -# Performs post-processing of compiler output and counts errors.
* Test source
* ------------------------
* - unit/atomics/CMakeLists.txt
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_atomicSub_Negative_Parameters") {}
/**
* End doxygen group atomicSub.
* @}
*/
/**
* @addtogroup atomicInc atomicInc
* @{
* @ingroup AtomicsTest
*/
/**
* Test Description
* ------------------------
* - Compiles atomicInc with invalid parameters.
* - Compiles the source with specialized Python tool.
* -# Utilizes sub-process to invoke compilation of faulty source.
* -# Performs post-processing of compiler output and counts errors.
* Test source
* ------------------------
* - unit/atomics/CMakeLists.txt
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_atomicInc_Negative_Parameters") {}
/**
* End doxygen group atomicInc.
* @}
*/
/**
* @addtogroup atomicDec atomicDec
* @{
* @ingroup AtomicsTest
*/
/**
* Test Description
* ------------------------
* - Compiles atomicDec with invalid parameters.
* - Compiles the source with specialized Python tool.
* -# Utilizes sub-process to invoke compilation of faulty source.
* -# Performs post-processing of compiler output and counts errors.
* Test source
* ------------------------
* - unit/atomics/CMakeLists.txt
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_atomicDec_Negative_Parameters") {}
/**
* End doxygen group atomicDec.
* @}
*/
/**
* End doxygen group AtomicsTest.
* @defgroup MathTest Math Device Functions
* @{
* This section describes tests for device math functions of HIP runtime API.
* @}
*/
/**
* @defgroup MathTest Math Device Functions
* @{
@@ -347,62 +223,8 @@ TEST_CASE("Unit_atomicDec_Negative_Parameters") {}
*/
/**
* @defgroup DeviceLanguageTest Device Language
* @defgroup VirtualMemoryManagementTest Virtual Memory Management APIs
* @{
* This section describes tests for the Device Language API.
*/
/**
* @addtogroup launch_bounds launch_bounds
* @{
* @ingroup DeviceLanguageTest
*/
/**
* Test Description
* ------------------------
* - Validates handling of invalid arguments:
* -# Compiles kernels that are not created appropriately:
* - Maximum number of threads is 0
* - Maximum number of threads is not integer value
* - Mimimum number of warps is not integer value
* -# Expected output: compiler error
* Test source
* ------------------------
* - unit/launch_bounds/CMakeLists.txt
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Kernel_Launch_bounds_Negative_Parameters_CompilerError") {}
/**
* Test Description
* ------------------------
* - Validates handling of invalid arguments:
* -# Compiles kernels that are not created appropriately:
* - Maximum number of threads is negative
* - Mimimum number of warps is negative
* - Validates handling of invalid arguments:
* -# Expected output: parse error
* Test source
* ------------------------
* - unit/launch_bounds/CMakeLists.txt
* Test requirements
* ------------------------
* - HIP_VERSION >= 5.2
*/
TEST_CASE("Unit_Kernel_Launch_bounds_Negative_Parameters_ParseError") {}
/**
* End doxygen group launch_bounds.
* This section describes the virtual memory management types & functions of HIP runtime API.
* @}
*/
/**
* End doxygen group DeviceLanguageTest.
* @}
* @defgroup VectorTypeTest Vector types
* @{
* This section describes tests for the Vector type functions and operators.
*/
projects/hip-tests/catch/unit/CMakeLists.txt
+1
View File
@@ -52,6 +52,7 @@ add_subdirectory(p2p)
add_subdirectory(gcc)
add_subdirectory(syncthreads)
add_subdirectory(threadfence)
add_subdirectory(virtualMemoryManagement)
if(HIP_PLATFORM STREQUAL "amd")
add_subdirectory(callback)
projects/hip-tests/catch/unit/memory/CMakeLists.txt
+2 -1
View File
@@ -94,7 +94,8 @@ if(HIP_PLATFORM MATCHES "amd")
hipMemAddressFree.cc
hipMemAddressReserve.cc
hipMemRelease.cc
hipMemGetAllocationPropertiesFromHandle.cc)
hipMemGetAllocationPropertiesFromHandle.cc
hipArray.cc)
else()
set(TEST_SRC ${TEST_SRC} hipGetSymbolSizeAddress.cc)
endif()
projects/hip-tests/catch/unit/memory/hip_vmm_common.hh
-51
View File
@@ -1,51 +0,0 @@
/*
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
#include "hip_test_context.hh"
#define checkVMMSupported(device) {\
int value = 0;\
hipDeviceAttribute_t\
attr = hipDeviceAttributeVirtualMemoryManagementSupported;\
HIP_CHECK(\
hipDeviceGetAttribute(&value, attr, device));\
if (value == 0) {\
HipTest::HIP_SKIP_TEST("Machine does not support VMM. Skipping Test..");\
return;\
}\
}
#define checkVMMSupportedRetVal(device) {\
int value = 0;\
hipDeviceAttribute_t\
attr = hipDeviceAttributeVirtualMemoryManagementSupported;\
HIP_CHECK(\
hipDeviceGetAttribute(&value, attr, device));\
if (value == 0) {\
HipTest::HIP_SKIP_TEST("Machine does not support VMM. Skipping Test..");\
return true;\
}\
}
constexpr int threadsPerBlk = 64;
projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt
+55
View File
@@ -0,0 +1,55 @@
# Copyright (c) 2023 Advanced Micro Devices, Inc. All Rights Reserved.
#
# Permission is hereby granted, free of charge, to any person obtaining a copy
# of this software and associated documentation files (the "Software"), to deal
# in the Software without restriction, including without limitation the rights
# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
# copies of the Software, and to permit persons to whom the Software is
# furnished to do so, subject to the following conditions:
#
# The above copyright notice and this permission notice shall be included in all
# copies or substantial portions of the Software.
#
# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
# SOFTWARE.
set(TEST_SRC
hipMemGetAllocationGranularity.cc
hipMemRetainAllocationHandle.cc)
if(HIP_PLATFORM MATCHES "nvidia")
set(TEST_SRC
${TEST_SRC}
hipMemMapArrayAsync.cc)
if(UNIX) # Disabled on AMD due to defect EXSWHTEC-375
set(TEST_SRC
${TEST_SRC}
hipMemExportToShareableHandle.cc
hipMemImportFromShareableHandle.cc)
endif()
endif()
if(HIP_PLATFORM MATCHES "amd")
set(TEST_SRC
${TEST_SRC}
hipMemAddressFree.cc
hipMemAddressReserve.cc
hipMemCreate.cc
hipMemSetGetAccess.cc
hipMemGetAllocationPropertiesFromHandle.cc
hipMemMap.cc
hipMemRelease.cc
hipMemUnmap.cc
hipMemVmm_old.cc)
endif()
hip_add_exe_to_target(NAME VirtualMemoryManagementTest
TEST_SRC ${TEST_SRC}
TEST_TARGET_NAME build_tests COMMON_SHARED_SRC ${COMMON_SHARED_SRC})
projects/hip-tests/catch/unit/memory/hipMemAddressFree.cc → projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressFree.cc
+8 -10
View File
@@ -23,12 +23,13 @@ THE SOFTWARE.
/**
* @addtogroup hipMemAddressFree hipMemAddressFree
* @{
* @ingroup MemoryTest
* @ingroup VirtualMemoryManagementTest
* `hipError_t hipMemAddressFree (void* devPtr, size_t size)` -
* Frees an address range reservation made via hipMemAddressReserve.
*/
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
#define DATA_SIZE (1 << 13)
@@ -38,7 +39,7 @@ THE SOFTWARE.
* ------------------------
* - Negative Tests
* ------------------------
* - catch\unit\memory\hipMemAddressFree.cc
* - unit/virtualMemoryManagement/hipMemAddressFree.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -50,16 +51,15 @@ TEST_CASE("Unit_hipMemAddressFree_negative") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate virtual address range
hipDeviceptr_t ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -68,9 +68,7 @@ TEST_CASE("Unit_hipMemAddressFree_negative") {
REQUIRE(hipMemAddressFree(nullptr, size_mem) == hipErrorInvalidValue);
}
SECTION("pass zero to size") {
REQUIRE(hipMemAddressFree(ptrA, 0) == hipErrorInvalidValue);
}
SECTION("pass zero to size") { REQUIRE(hipMemAddressFree(ptrA, 0) == hipErrorInvalidValue); }
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
}
projects/hip-tests/catch/unit/memory/hipMemAddressReserve.cc → projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemAddressReserve.cc
+19 -21
View File
@@ -23,7 +23,7 @@ THE SOFTWARE.
/**
* @addtogroup hipMemAddressReserve hipMemAddressReserve
* @{
* @ingroup MemoryTest
* @ingroup VirtualMemoryManagementTest
* `hipError_t hipMemAddressReserve (void** ptr,
* size_t size,
* size_t alignment,
@@ -33,6 +33,7 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
#define DATA_SIZE (1 << 13)
@@ -43,7 +44,7 @@ THE SOFTWARE.
* - Verify if reserved address returned by hipMemAddressReserve
* for different alignment values are correctly aligned.
* ------------------------
* - catch\unit\memory\hipMemAddressReserve.cc
* - unit/virtualMemoryManagement/hipMemAddressReserve.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -56,16 +57,15 @@ TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate virtual address range
hipDeviceptr_t ptrA;
size_t alignmnt = 1;
@@ -106,7 +106,7 @@ TEST_CASE("Unit_hipMemAddressReserve_AlignmentTest") {
* ------------------------
* - Negative Tests
* ------------------------
* - catch\unit\memory\hipMemAddressReserve.cc
* - unit/virtualMemoryManagement/hipMemAddressReserve.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -118,35 +118,33 @@ TEST_CASE("Unit_hipMemAddressReserve_Negative") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate virtual address range
hipDeviceptr_t ptrA;
SECTION("Nullptr to ptr") {
REQUIRE(hipMemAddressReserve(nullptr, size_mem, 0, 0, 0) ==
hipErrorInvalidValue);
REQUIRE(hipMemAddressReserve(nullptr, size_mem, 0, 0, 0) == hipErrorInvalidValue);
}
SECTION("pass size as 0") {
REQUIRE(hipMemAddressReserve(&ptrA, 0, 0, 0, 0) ==
hipErrorMemoryAllocation);
REQUIRE(hipMemAddressReserve(&ptrA, 0, 0, 0, 0) == hipErrorMemoryAllocation);
}
#if HT_NVIDIA
SECTION("pass non power of two for alignment") {
REQUIRE(hipMemAddressReserve(&ptrA, size_mem, 3, 0, 0) ==
hipErrorMemoryAllocation);
REQUIRE(hipMemAddressReserve(&ptrA, size_mem, 3, 0, 0) == hipErrorMemoryAllocation);
}
#endif
SECTION("pass size as non multiple of host page size") {
REQUIRE(hipMemAddressReserve(&ptrA, (size_mem - 1), 0, 0, 0) ==
hipErrorMemoryAllocation);
REQUIRE(hipMemAddressReserve(&ptrA, (size_mem - 1), 0, 0, 0) == hipErrorMemoryAllocation);
}
}
projects/hip-tests/catch/unit/memory/hipMemCreate.cc → projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemCreate.cc
+48 -68
View File
@@ -19,20 +19,23 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
* @addtogroup hipMemCreate hipMemCreate
* @{
* @ingroup MemoryTest
* @ingroup VirtualMemoryManagementTest
* `hipError_t hipMemCreate (hipMemGenericAllocationHandle_t* handle,
* size_t size,
* const hipMemAllocationProp* prop,
* unsigned long long flags)` -
* Creates a memory allocation described by the properties and size.
*/
#include "hip_vmm_common.hh"
#include <hip_test_kernels.hh>
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
#define THREADS_PER_BLOCK 512
#define NUM_OF_BUFFERS 3
#define DATA_SIZE (1 << 13)
@@ -52,7 +55,7 @@ static __global__ void square_kernel(int* Buff) {
* - Allocate physical memories for different multiples of
* granularity and deallocate them.
* ------------------------
* - catch\unit\memory\hipMemCreate.cc
* - unit/virtualMemoryManagement/hipMemCreate.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -62,19 +65,18 @@ TEST_CASE("Unit_hipMemCreate_BasicAllocateDeAlloc_MultGranularity") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
hipMemGenericAllocationHandle_t handle;
// Allocate physical memory
for (int mul = 1; mul < 64; mul++) {
HIP_CHECK(hipMemCreate(&handle, granularity*mul, &prop, 0));
HIP_CHECK(hipMemCreate(&handle, granularity * mul, &prop, 0));
HIP_CHECK(hipMemRelease(handle));
}
}
@@ -87,7 +89,7 @@ TEST_CASE("Unit_hipMemCreate_BasicAllocateDeAlloc_MultGranularity") {
* and back to host. Verify the result. Release handle at end after
* unmapping VMM range.
* ------------------------
* - catch\unit\memory\hipMemCreate.cc
* - unit/virtualMemoryManagement/hipMemCreate.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -99,17 +101,15 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
@@ -146,7 +146,7 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPostUnmap") {
* to VMM memory and back to host. Verify the result. Release
* handle before the VMM range is used.
* ------------------------
* - catch\unit\memory\hipMemCreate.cc
* - unit/virtualMemoryManagement/hipMemCreate.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -158,17 +158,15 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
@@ -205,12 +203,11 @@ TEST_CASE("Unit_hipMemCreate_ChkDev2HstMemcpy_ReleaseHdlPreUse") {
* to device, launch kernel to square the data, copy data back
* to host. Verify the result.
* ------------------------
* - catch\unit\memory\hipMemCreate.cc
* - unit/virtualMemoryManagement/hipMemCreate.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") {
size_t granularity = 0;
constexpr int N = DATA_SIZE;
@@ -218,17 +215,15 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
@@ -248,12 +243,12 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") {
// Initialize with data
for (size_t idx = 0; idx < N; idx++) {
A_h[idx] = idx;
C_h[idx] = idx*idx;
C_h[idx] = idx * idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
// Invoke kernel
hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK),
dim3(THREADS_PER_BLOCK), 0, 0, static_cast<int*>(ptrA));
hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0,
static_cast<int*>(ptrA));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipDeviceSynchronize());
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
@@ -269,7 +264,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithKerLaunch") {
* device permission, copy data from host to device, launch kernel
* to square the data, copy data back to host. Verify the result.
* ------------------------
* - catch\unit\memory\hipMemCreate.cc
* - unit/virtualMemoryManagement/hipMemCreate.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -282,17 +277,15 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle[NUM_OF_BUFFERS];
// Allocate 3 physical memory chunks
for (int count = 0; count < numOfBuffers; count++) {
@@ -304,8 +297,7 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") {
for (int idx = 0; idx < numOfBuffers; idx++) {
uint64_t uiptr = reinterpret_cast<uint64_t>(ptrA);
uiptr = uiptr + idx * size_mem;
HIP_CHECK(hipMemMap(reinterpret_cast<void*>(uiptr), size_mem, 0,
handle[idx], 0));
HIP_CHECK(hipMemMap(reinterpret_cast<void*>(uiptr), size_mem, 0, handle[idx], 0));
HIP_CHECK(hipMemRelease(handle[idx]));
}
hipMemAccessDesc accessDesc = {};
@@ -315,16 +307,16 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") {
// Make the address accessible to GPU 0
HIP_CHECK(hipMemSetAccess(ptrA, (numOfBuffers * size_mem), &accessDesc, 1));
std::vector<int> A_h(numOfBuffers * size_mem), B_h(numOfBuffers * size_mem),
C_h(numOfBuffers * size_mem);
C_h(numOfBuffers * size_mem);
// Fill Data
for (size_t idx = 0; idx < (numOfBuffers * N); idx++) {
A_h[idx] = idx;
C_h[idx] = idx*idx;
C_h[idx] = idx * idx;
}
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), numOfBuffers * buffer_size));
// Launch square kernel
hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK),
dim3(THREADS_PER_BLOCK), 0, 0, static_cast<int*>(ptrA));
hipLaunchKernelGGL(square_kernel, dim3(N / THREADS_PER_BLOCK), dim3(THREADS_PER_BLOCK), 0, 0,
static_cast<int*>(ptrA));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, numOfBuffers * buffer_size));
HIP_CHECK(hipDeviceSynchronize());
// Validate Results
@@ -344,7 +336,7 @@ TEST_CASE("Unit_hipMemCreate_MapNonContiguousChunks") {
* to the VMM address range. Memset the VMM address range with initial
* value. Validate.
* ------------------------
* - catch\unit\memory\hipMemCreate.cc
* - unit/virtualMemoryManagement/hipMemCreate.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -357,17 +349,15 @@ TEST_CASE("Unit_hipMemCreate_ChkWithMemset") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
@@ -398,7 +388,7 @@ TEST_CASE("Unit_hipMemCreate_ChkWithMemset") {
* ------------------------
* - Negative Tests
* ------------------------
* - catch\unit\memory\hipMemCreate.cc
* - unit/virtualMemoryManagement/hipMemCreate.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -408,58 +398,48 @@ TEST_CASE("Unit_hipMemCreate_Negative") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemGenericAllocationHandle_t handle;
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Device
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
SECTION("Nullptr to handle") {
REQUIRE(hipMemCreate(nullptr, granularity, &prop, 0) ==
hipErrorInvalidValue);
REQUIRE(hipMemCreate(nullptr, granularity, &prop, 0) == hipErrorInvalidValue);
}
SECTION("Nullptr to prop") {
REQUIRE(hipMemCreate(&handle, granularity, nullptr, 0) ==
hipErrorInvalidValue);
REQUIRE(hipMemCreate(&handle, granularity, nullptr, 0) == hipErrorInvalidValue);
}
SECTION("pass size as 0") {
REQUIRE(hipMemCreate(&handle, 0, &prop, 0) == hipErrorInvalidValue);
}
SECTION("pass size as 0") { REQUIRE(hipMemCreate(&handle, 0, &prop, 0) == hipErrorInvalidValue); }
SECTION("Pass prop type as invalid") {
prop.type = hipMemAllocationTypeInvalid;
REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) ==
hipErrorInvalidValue);
REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidValue);
}
SECTION("pass location as invalid") {
prop.location.type = hipMemLocationTypeInvalid;
REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) ==
hipErrorInvalidValue);
REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidValue);
}
SECTION("non multiple of granularity") {
REQUIRE(hipMemCreate(&handle, (granularity - 1), &prop, 0) ==
hipErrorInvalidValue);
REQUIRE(hipMemCreate(&handle, (granularity - 1), &prop, 0) == hipErrorInvalidValue);
}
SECTION("pass location id as -1") {
prop.location.id = -1; // set to non existing device
REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) ==
hipErrorInvalidValue);
REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidValue);
}
SECTION("pass location id as > highest device number") {
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
prop.location.id = numDevices; // set to non existing device
REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) ==
hipErrorInvalidValue);
REQUIRE(hipMemCreate(&handle, granularity, &prop, 0) == hipErrorInvalidValue);
}
}
projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc
+145
View File
@@ -0,0 +1,145 @@
/*
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
* @addtogroup hipMemExportToShareableHandle hipMemExportToShareableHandle
* @{
* @ingroup VirtualMemoryManagementTest
* `hipError_t hipMemExportToShareableHandle(void *shareableHandle,
* hipMemGenericAllocationHandle_t handle,
* hipMemAllocationHandleType handleType,
* unsigned long long flags)` -
* Exports an allocation to a requested shareable handle type.
*/
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
/**
* Test Description
* ------------------------
* - Basic sanity test.
* ------------------------
* - unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc
* Test requirements
* ------------------------
* - Host specific (LINUX)
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemExportToShareableHandle_Positive_Basic") {
HIP_CHECK(hipFree(0));
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device);
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
size_t granularity;
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, granularity * 2, &prop, 0));
void* shareable_handle = nullptr;
HIP_CHECK(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 0));
REQUIRE(shareable_handle != nullptr);
HIP_CHECK(hipMemRelease(handle));
}
/**
* Test Description
* ------------------------
* - Negative parameters test.
* ------------------------
* - unit/virtualMemoryManagement/hipMemExportToShareableHandle.cc
* Test requirements
* ------------------------
* - Host specific (LINUX)
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemExportToShareableHandle_Negative_Parameters") {
HIP_CHECK(hipFree(0));
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device);
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
size_t granularity;
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, granularity * 2, &prop, 0));
void* shareable_handle = nullptr;
SECTION("shareableHandle == nullptr") {
HIP_CHECK_ERROR(
hipMemExportToShareableHandle(nullptr, handle, hipMemHandleTypePosixFileDescriptor, 0),
hipErrorInvalidValue);
}
#if HT_AMD
SECTION("handle == nullptr") {
HIP_CHECK_ERROR(hipMemExportToShareableHandle(&shareable_handle, nullptr,
hipMemHandleTypePosixFileDescriptor, 0),
hipErrorInvalidValue);
}
#endif
SECTION("invalid handleType") {
HIP_CHECK_ERROR(
hipMemExportToShareableHandle(&shareable_handle, handle, hipMemHandleTypeWin32, 0),
hipErrorInvalidValue);
}
SECTION("non-zero flags") {
HIP_CHECK_ERROR(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 1),
hipErrorInvalidValue);
}
HIP_CHECK(hipMemRelease(handle));
#if HT_AMD // segfaults on NVIDIA
SECTION("released handle") {
HIP_CHECK_ERROR(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 0),
hipErrorInvalidValue);
}
#endif
}
projects/hip-tests/catch/unit/memory/hipMemGetAllocationGranularity.cc → projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc
+38 -33
View File
@@ -23,7 +23,7 @@ THE SOFTWARE.
/**
* @addtogroup hipMemGetAllocationGranularity hipMemGetAllocationGranularity
* @{
* @ingroup MemoryTest
* @ingroup VirtualMemoryManagementTest
* `hipError_t hipMemGetAllocationGranularity (size_t* granularity,
* const hipMemAllocationProp* prop,
* hipMemAllocationGranularity_flags option)` -
@@ -33,14 +33,13 @@ THE SOFTWARE.
#include <hip_test_checkers.hh>
#include <hip_test_kernels.hh>
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
/**
local function to invoke hipMemGetAllocationGranularity.
*/
void getGranularity(size_t *granularity,
hipMemAllocationGranularity_flags option,
int device) {
void getGranularity(size_t* granularity, hipMemAllocationGranularity_flags option, int device) {
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
@@ -54,16 +53,17 @@ void getGranularity(size_t *granularity,
* - Functional Test to get granularity size for
* hipMemAllocationGranularityMinimum option.
* ------------------------
* - catch\unit\memory\hipMemGetAllocationGranularity.cc
* - unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemGetAllocationGranularity_MinGranularity") {
HIP_CHECK(hipFree(0));
size_t granularity = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device)
checkVMMSupported(device);
getGranularity(&granularity, hipMemAllocationGranularityMinimum, 0);
REQUIRE(granularity > 0);
}
@@ -74,16 +74,17 @@ TEST_CASE("Unit_hipMemGetAllocationGranularity_MinGranularity") {
* - Functional Test to get granularity size for
* hipMemAllocationGranularityRecommended option.
* ------------------------
* - catch\unit\memory\hipMemGetAllocationGranularity.cc
* - unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemGetAllocationGranularity_RecommendedGranularity") {
HIP_CHECK(hipFree(0));
size_t granularity = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device)
checkVMMSupported(device);
getGranularity(&granularity, hipMemAllocationGranularityRecommended, 0);
REQUIRE(granularity > 0);
}
@@ -94,21 +95,21 @@ TEST_CASE("Unit_hipMemGetAllocationGranularity_RecommendedGranularity") {
* - Functional Test to get granularity size for
* hipMemAllocationGranularityMinimum option for all GPUs.
* ------------------------
* - catch\unit\memory\hipMemGetAllocationGranularity.cc
* - unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemGetAllocationGranularity_AllGPUs") {
HIP_CHECK(hipFree(0));
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
for (int dev = 0; dev < numDevices; dev++) {
size_t granularity = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, dev));
checkVMMSupported(device)
getGranularity(&granularity, hipMemAllocationGranularityRecommended,
dev);
checkVMMSupported(device);
getGranularity(&granularity, hipMemAllocationGranularityRecommended, dev);
REQUIRE(granularity > 0);
}
}
@@ -118,16 +119,17 @@ TEST_CASE("Unit_hipMemGetAllocationGranularity_AllGPUs") {
* ------------------------
* - Negative Tests
* ------------------------
* - catch\unit\memory\hipMemGetAllocationGranularity.cc
* - unit/virtualMemoryManagement/hipMemGetAllocationGranularity.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemGetAllocationGranularity_NegativeTests") {
HIP_CHECK(hipFree(0));
size_t granularity = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
@@ -135,45 +137,48 @@ TEST_CASE("Unit_hipMemGetAllocationGranularity_NegativeTests") {
SECTION("Granularity is nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipMemGetAllocationGranularity(nullptr, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(nullptr, &prop, hipMemAllocationGranularityMinimum));
}
#if HT_AMD // segfaults on NVIDIA
SECTION("Prop is nullptr") {
REQUIRE(hipErrorInvalidValue ==
hipMemGetAllocationGranularity(&granularity, nullptr,
hipMemAllocationGranularityMinimum));
REQUIRE(
hipErrorInvalidValue ==
hipMemGetAllocationGranularity(&granularity, nullptr, hipMemAllocationGranularityMinimum));
}
#endif
#if HT_NVIDIA
SECTION("flag is invalid") {
REQUIRE(hipErrorInvalidValue ==
hipMemGetAllocationGranularity(&granularity, &prop,
(hipMemAllocationGranularity_flags)0xff));
hipMemGetAllocationGranularity(&granularity, &prop,
(hipMemAllocationGranularity_flags)0xff));
}
#endif
#if HT_AMD // succeeds on NVIDIA
SECTION("device id > highest device id") {
int numDevices = 0;
HIP_CHECK(hipGetDeviceCount(&numDevices));
prop.location.id = numDevices; // set to non existing device
REQUIRE(hipErrorInvalidValue ==
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
REQUIRE(
hipErrorInvalidValue ==
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
}
SECTION("device id < lowest device id") {
prop.location.id = -1; // set to non existing device
REQUIRE(hipErrorInvalidValue ==
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
REQUIRE(
hipErrorInvalidValue ==
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
}
SECTION("allocation type as invalid") {
prop.type = hipMemAllocationTypeInvalid;
REQUIRE(hipErrorInvalidValue ==
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
REQUIRE(
hipErrorInvalidValue ==
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
}
SECTION("location type as invalid") {
prop.location.type = hipMemLocationTypeInvalid;
REQUIRE(hipErrorInvalidValue ==
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
REQUIRE(
hipErrorInvalidValue ==
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
}
#endif
}
projects/hip-tests/catch/unit/memory/hipMemGetAllocationPropertiesFromHandle.cc → projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc
+20 -22
View File
@@ -20,26 +20,27 @@ OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
#define DATA_SIZE (1 << 13)
/**
* @addtogroup hipMemGetAllocationPropertiesFromHandle hipMemGetAllocationPropertiesFromHandle
* @{
* @ingroup MemoryTest
* @ingroup VirtualMemoryManagementTest
* `hipError_t hipMemGetAllocationPropertiesFromHandle(hipMemAllocationProp* prop,
* hipMemGenericAllocationHandle_t handle)` -
* Retrieve the property structure of the given handle.
*/
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
#define DATA_SIZE (1 << 13)
/**
* Test Description
* ------------------------
* - Functional test to verify the values of hipMemAllocationProp properties.
* ------------------------
* - catch\unit\memory\hipMemGetAllocationPropertiesFromHandle.cc
* - unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -47,7 +48,7 @@ THE SOFTWARE.
TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") {
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemGenericAllocationHandle_t handle;
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
@@ -58,11 +59,10 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") {
size_t granularity = 0;
int N = DATA_SIZE;
size_t buffer_size = N * sizeof(int);
HIP_CHECK(hipMemGetAllocationGranularity
(&granularity, &prop, hipMemAllocationGranularityMinimum));
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t mem_size = ((granularity + buffer_size - 1) / granularity)
* granularity;
size_t mem_size = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, mem_size, &prop, 0));
// verify properties has been retrived from handle
@@ -78,7 +78,7 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") {
* ------------------------
* - Negative Tests.
* ------------------------
* - catch\unit\memory\hipMemGetAllocationPropertiesFromHandle.cc
* - unit/virtualMemoryManagement/hipMemGetAllocationPropertiesFromHandle.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -86,7 +86,7 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_functional") {
TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_Negative") {
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemGenericAllocationHandle_t handle;
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
@@ -97,23 +97,21 @@ TEST_CASE("Unit_hipMemGetAllocationPropertiesFromHandle_Negative") {
size_t granularity = 0;
int N = DATA_SIZE;
size_t buffer_size = N * sizeof(int);
HIP_CHECK(hipMemGetAllocationGranularity
(&granularity, &prop, hipMemAllocationGranularityMinimum));
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t mem_size = ((granularity + buffer_size - 1) / granularity)
* granularity;
size_t mem_size = ((granularity + buffer_size - 1) / granularity) * granularity;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, mem_size, &prop, 0));
SECTION("Nullptr as prop") {
REQUIRE(hipMemGetAllocationPropertiesFromHandle(nullptr, handle)
== hipErrorInvalidValue);
REQUIRE(hipMemGetAllocationPropertiesFromHandle(nullptr, handle) == hipErrorInvalidValue);
}
SECTION("null handle") {
prop.location.type = hipMemLocationTypeInvalid;
REQUIRE(hipMemGetAllocationPropertiesFromHandle(&prop_temp, nullptr)
== hipErrorInvalidValue);
REQUIRE(hipMemGetAllocationPropertiesFromHandle(&prop_temp, nullptr) == hipErrorInvalidValue);
}
HIP_CHECK(hipMemRelease(handle));
}
projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc
+210
View File
@@ -0,0 +1,210 @@
/*
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
* @addtogroup hipMemExportToShareableHandle hipMemExportToShareableHandle
* @{
* @ingroup VirtualMemoryManagementTest
* `hipError_t hipMemImportFromShareableHandle(hipMemGenericAllocationHandle_t *handle,
* void *osHandle,
* hipMemAllocationHandleType shHandleType)` -
* Imports an allocation from a requested shareable handle type.
*/
#include <unistd.h>
#include <sys/types.h>
#include <sys/wait.h>
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
/**
* Test Description
* ------------------------
* - Basic sanity test.
* ------------------------
* - unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc
* Test requirements
* ------------------------
* - Host specific (LINUX)
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemImportFromShareableHandle_Positive_Basic") {
HIP_CHECK(hipFree(0));
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device);
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
size_t granularity;
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, granularity * 2, &prop, 0));
void* shareable_handle = nullptr;
HIP_CHECK(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 0));
hipMemGenericAllocationHandle_t imported_handle;
HIP_CHECK(hipMemImportFromShareableHandle(&imported_handle, shareable_handle,
hipMemHandleTypePosixFileDescriptor));
HIP_CHECK(hipMemRelease(handle));
}
/**
* Test Description
* ------------------------
* - Basic multiprocess sanity test.
* ------------------------
* - unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc
* Test requirements
* ------------------------
* - Host specific (LINUX)
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemImportFromShareableHandle_Positive_MultiProc") {
int fd[2];
REQUIRE(pipe(fd) == 0);
auto pid = fork();
REQUIRE(pid >= 0);
if (pid == 0) { // child
REQUIRE(close(fd[1]) == 0);
void* shareable_handle = nullptr;
REQUIRE(read(fd[0], &shareable_handle, sizeof(shareable_handle)) >= 0);
REQUIRE(close(fd[0]) == 0);
REQUIRE(shareable_handle != nullptr);
HIP_CHECK(hipFree(0));
hipMemGenericAllocationHandle_t imported_handle;
HIP_CHECK(hipMemImportFromShareableHandle(&imported_handle, shareable_handle,
hipMemHandleTypePosixFileDescriptor));
exit(0);
} else { // parent
REQUIRE(close(fd[0]) == 0);
HIP_CHECK(hipFree(0));
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device);
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
size_t granularity;
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, granularity * 2, &prop, 0));
void* shareable_handle = nullptr;
HIP_CHECK(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 0));
REQUIRE(write(fd[1], &shareable_handle, sizeof(shareable_handle)) >= 0);
REQUIRE(close(fd[1]) == 0);
REQUIRE(wait(NULL) >= 0);
HIP_CHECK(hipMemRelease(handle));
}
}
/**
* Test Description
* ------------------------
* - Negative parameters test.
* ------------------------
* - unit/virtualMemoryManagement/hipMemImportFromShareableHandle.cc
* Test requirements
* ------------------------
* - Host specific (LINUX)
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemImportFromShareableHandle_Negative_Parameters") {
HIP_CHECK(hipFree(0));
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device);
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.requestedHandleTypes = hipMemHandleTypePosixFileDescriptor;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
size_t granularity;
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, granularity * 2, &prop, 0));
void* shareable_handle = nullptr;
HIP_CHECK(hipMemExportToShareableHandle(&shareable_handle, handle,
hipMemHandleTypePosixFileDescriptor, 0));
hipMemGenericAllocationHandle_t imported_handle;
#if HT_AMD
SECTION("handle == nullptr") {
HIP_CHECK_ERROR(hipMemImportFromShareableHandle(nullptr, shareable_handle,
hipMemHandleTypePosixFileDescriptor),
hipErrorInvalidValue);
}
#endif
SECTION("shareableHandle == nullptr") {
HIP_CHECK_ERROR(hipMemImportFromShareableHandle(&imported_handle, nullptr,
hipMemHandleTypePosixFileDescriptor),
hipErrorInvalidValue);
}
SECTION("invalid handleType") {
HIP_CHECK_ERROR(
hipMemImportFromShareableHandle(&imported_handle, shareable_handle, hipMemHandleTypeWin32),
hipErrorNotSupported);
}
HIP_CHECK(hipMemRelease(handle));
}
projects/hip-tests/catch/unit/memory/hipMemMap.cc → projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMap.cc
+54 -72
View File
@@ -19,10 +19,11 @@ LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
* @addtogroup hipMemMap hipMemMap
* @{
* @ingroup MemoryTest
* @ingroup VirtualMemoryManagementTest
* `hipError_t hipMemMap (void* ptr,
* size_t size,
* size_t offset,
@@ -32,6 +33,7 @@ THE SOFTWARE.
*/
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
constexpr int N = (1 << 13);
@@ -45,7 +47,7 @@ constexpr int initializer = 0;
* vmm address range repeatedly. This test validates physical memory
* euse using same vmm range.
* ------------------------
* - catch\unit\memory\hipMemMap.cc
* - unit/virtualMemoryManagement/hipMemMap.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -57,24 +59,22 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
// Allocate host memory and intialize data
std::vector<int> A_h(N), B_h(N), C_h(N);
// Initialize with data
for (size_t idx = 0; idx < N; idx++) {
A_h[idx] = idx;
C_h[idx] = idx*idx;
C_h[idx] = idx * idx;
}
// Allocate a physical memory chunk
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
@@ -94,8 +94,7 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") {
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
#if HT_NVIDIA
square_kernel <<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(
static_cast<int*>(ptrA));
square_kernel<<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(static_cast<int*>(ptrA));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipStreamSynchronize(0));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
@@ -114,7 +113,7 @@ TEST_CASE("Unit_hipMemMap_SameMemoryReuse") {
* vmm addresses. This test validates physical memory reuse using
* different vmm ranges.
* ------------------------
* - catch\unit\memory\hipMemMap.cc
* - unit/virtualMemoryManagement/hipMemMap.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -125,24 +124,22 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
// Allocate host memory and intialize data
std::vector<int> A_h(N), B_h(N), C_h(N);
// Initialize with data
for (size_t idx = 0; idx < N; idx++) {
A_h[idx] = idx;
C_h[idx] = idx*idx;
C_h[idx] = idx * idx;
}
// Allocate a physical memory chunk
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
@@ -164,8 +161,8 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") {
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
#if HT_NVIDIA
square_kernel <<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(
static_cast<int*>(ptrA[buf]));
square_kernel<<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(
static_cast<int*>(ptrA[buf]));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA[buf], buffer_size));
HIP_CHECK(hipStreamSynchronize(0));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
@@ -186,7 +183,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_SingleGPU") {
* vmm addresses at the same time and check data values integrity
* between different VMMs.
* ------------------------
* - catch\unit\memory\hipMemMap.cc
* - unit/virtualMemoryManagement/hipMemMap.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -197,17 +194,15 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemory_Map2MultVMMs") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
// Allocate host memory and intialize data
std::vector<int> A_h(N), B_h(N);
@@ -253,7 +248,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemory_Map2MultVMMs") {
* multiple vmm addresses. This test validates physical memory
* reuse using different vmm ranges on multiple devices.
* ------------------------
* - catch\unit\memory\hipMemMap.cc
* - unit/virtualMemoryManagement/hipMemMap.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -270,17 +265,15 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_MultiDev") {
for (int devX = 0; devX < devicecount; devX++) {
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, devX));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
// Allocate host memory and intialize data
std::vector<int> A_h(N), B_h(N);
@@ -328,7 +321,7 @@ TEST_CASE("Unit_hipMemMap_PhysicalMemoryReuse_MultiDev") {
* for single vmm address. This test validates VMM memory reuse
* using different physical ranges.
* ------------------------
* - catch\unit\memory\hipMemMap.cc
* - unit/virtualMemoryManagement/hipMemMap.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -339,24 +332,22 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle[num_buf];
// Allocate host memory and intialize data
std::vector<int> A_h(N), B_h(N), C_h(N);
// Initialize with data
for (size_t idx = 0; idx < N; idx++) {
A_h[idx] = idx;
C_h[idx] = idx*idx;
C_h[idx] = idx * idx;
}
// Allocate a physical memory chunk
for (int buf = 0; buf < num_buf; buf++) {
@@ -379,8 +370,7 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") {
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), A_h.data()));
#if HT_NVIDIA
square_kernel <<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(
static_cast<int*>(ptrA));
square_kernel<<<dim3(N / threadsPerBlk), dim3(threadsPerBlk), 0, 0>>>(static_cast<int*>(ptrA));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
HIP_CHECK(hipStreamSynchronize(0));
REQUIRE(true == std::equal(B_h.begin(), B_h.end(), C_h.data()));
@@ -401,7 +391,7 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_SingleGPU") {
* can be mapped/unmapped to single vmm address. This test validates VMM
* memory reuse using different physical ranges.
* ------------------------
* - catch\unit\memory\hipMemMap.cc
* - unit/virtualMemoryManagement/hipMemMap.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -417,17 +407,15 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_MultiGPU") {
size_t buffer_size = N * sizeof(int);
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
std::vector<hipMemGenericAllocationHandle_t> handle(devicecount);
// Allocate host memory and intialize data
std::vector<int> A_h(N), B_h(N);
@@ -493,7 +481,7 @@ TEST_CASE("Unit_hipMemMap_VMMMemoryReuse_MultiGPU") {
* - Check if a partial part of a physical chunk can be mapped/unmapped
* to a smaller vmm address.
* ------------------------
* - catch\unit\memory\hipMemMap.cc
* - unit/virtualMemoryManagement/hipMemMap.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -504,17 +492,15 @@ TEST_CASE("Unit_hipMemMap_MapPartialPhysicalMem") {
size_t buffer_size = N * sizeof(int);
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
// Allocate host memory and intialize data
std::vector<int> A_h(N), B_h(N);
@@ -523,7 +509,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialPhysicalMem") {
A_h[idx] = idx;
}
// Allocate a bigger physical memory chunk of twice size_mem
HIP_CHECK(hipMemCreate(&handle, 2*size_mem, &prop, 0));
HIP_CHECK(hipMemCreate(&handle, 2 * size_mem, &prop, 0));
// Allocate virtual address range of size size_mem
hipDeviceptr_t ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
@@ -549,7 +535,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialPhysicalMem") {
* - Check if a partial part of a VMM range can be mapped/unmapped
* to a physical address.
* ------------------------
* - catch\unit\memory\hipMemMap.cc
* - unit/virtualMemoryManagement/hipMemMap.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -560,17 +546,15 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") {
size_t buffer_size = N * sizeof(int);
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
// Allocate host memory and intialize data
std::vector<int> A_h(N), B_h(N);
@@ -582,13 +566,13 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") {
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range of size twice size_mem
hipDeviceptr_t ptrA;
HIP_CHECK(hipMemAddressReserve(&ptrA, 2*size_mem, 0, 0, 0));
HIP_CHECK(hipMemAddressReserve(&ptrA, 2 * size_mem, 0, 0, 0));
hipMemAccessDesc accessDesc = {};
accessDesc.location.type = hipMemLocationTypeDevice;
accessDesc.location.id = device;
accessDesc.flags = hipMemAccessFlagsProtReadWrite;
std::fill(B_h.begin(), B_h.end(), initializer);
HIP_CHECK(hipMemMap(ptrA , size_mem, 0, handle, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
HIP_CHECK(hipMemSetAccess(ptrA, size_mem, &accessDesc, 1));
HIP_CHECK(hipMemcpyHtoD(ptrA, A_h.data(), buffer_size));
HIP_CHECK(hipMemcpyDtoH(B_h.data(), ptrA, buffer_size));
@@ -596,7 +580,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") {
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
// Release resources
HIP_CHECK(hipMemRelease(handle));
HIP_CHECK(hipMemAddressFree(ptrA, 2*size_mem));
HIP_CHECK(hipMemAddressFree(ptrA, 2 * size_mem));
}
/**
@@ -604,7 +588,7 @@ TEST_CASE("Unit_hipMemMap_MapPartialVMMMem") {
* ------------------------
* - Negative Argument Tests
* ------------------------
* - catch\unit\memory\hipMemMap.cc
* - unit/virtualMemoryManagement/hipMemMap.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -615,16 +599,15 @@ TEST_CASE("Unit_hipMemMap_negative") {
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
// Allocate physical memory
@@ -633,8 +616,7 @@ TEST_CASE("Unit_hipMemMap_negative") {
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
SECTION("nullptr to ptrA") {
REQUIRE(hipMemMap(nullptr, size_mem, 0, handle, 0) ==
hipErrorInvalidValue);
REQUIRE(hipMemMap(nullptr, size_mem, 0, handle, 0) == hipErrorInvalidValue);
}
SECTION("pass zero to size") {
@@ -642,9 +624,9 @@ TEST_CASE("Unit_hipMemMap_negative") {
}
SECTION("pass negative to offset") {
REQUIRE(hipMemMap(&ptrA, size_mem, -1, handle, 0) ==
hipErrorInvalidValue);
REQUIRE(hipMemMap(&ptrA, size_mem, -1, handle, 0) == hipErrorInvalidValue);
}
HIP_CHECK(hipMemRelease(handle));
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
}
projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemMapArrayAsync.cc
+110
View File
@@ -0,0 +1,110 @@
/*
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/**
* @addtogroup hipMemMapArrayAsync hipMemMapArrayAsync
* @{
* @ingroup VirtualMemoryManagementTest
* `hipError_t hipMemMapArrayAsync(hipArrayMapInfo *mapInfoList,
* unsigned int count,
* hipStream_t stream)` -
* Maps or unmaps subregions of sparse HIP arrays and sparse HIP mipmapped arrays.
*/
#include <hip_array_common.hh>
#include <hip_test_common.hh>
#include <resource_guards.hh>
#include "hip_vmm_common.hh"
/**
* Test Description
* ------------------------
* - Basic sanity test.
* ------------------------
* - unit/virtualMemoryManagement/hipMemMapArrayAsync.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemMapArrayAsync_Positive_Basic") {
HIP_CHECK(hipFree(0));
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, 0));
checkVMMSupported(device);
CHECK_IMAGE_SUPPORT;
hipmipmappedArray array;
HIP_ARRAY3D_DESCRIPTOR desc = {};
using vec_info = vector_info<float>;
desc.Format = vec_info::format;
desc.NumChannels = vec_info::size;
desc.Width = 1;
desc.Height = 1;
desc.Flags = CUDA_ARRAY3D_SPARSE;
unsigned int levels = 2;
HIP_CHECK(hipMipmappedArrayCreate(&array, &desc, levels));
hipMemAllocationProp prop = {};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device;
prop.allocFlags.usage = CU_MEM_CREATE_USAGE_TILE_POOL;
size_t granularity;
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityRecommended));
hipMemGenericAllocationHandle_t handle;
HIP_CHECK(hipMemCreate(&handle, granularity, &prop, 0));
hipArrayMapInfo map_info_list = {};
map_info_list.resourceType = HIP_RESOURCE_TYPE_MIPMAPPED_ARRAY;
map_info_list.resource.mipmap = array;
map_info_list.subresourceType = hipArraySparseSubresourceTypeSparseLevel;
map_info_list.subresource.sparseLevel.extentWidth = 1;
map_info_list.subresource.sparseLevel.extentHeight = 1;
map_info_list.subresource.sparseLevel.extentDepth = 1;
map_info_list.memOperationType = hipMemOperationTypeMap;
map_info_list.memHandleType = hipMemHandleTypeGeneric;
map_info_list.memHandle.memHandle = handle;
map_info_list.deviceBitMask = 0x1;
StreamGuard stream(Streams::created);
HIP_CHECK(hipMemMapArrayAsync(&map_info_list, 1, stream.stream()));
HIP_CHECK(hipStreamSynchronize(stream.stream()));
map_info_list.memOperationType = hipMemOperationTypeUnmap;
map_info_list.memHandle.memHandle = NULL;
HIP_CHECK(hipMemMapArrayAsync(&map_info_list, 1, stream.stream()));
HIP_CHECK(hipStreamSynchronize(stream.stream()));
HIP_CHECK(hipMemRelease(handle));
HIP_CHECK(hipMipmappedArrayDestroy(array));
}
projects/hip-tests/catch/unit/memory/hipMemRelease.cc → projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRelease.cc
+3 -5
View File
@@ -23,7 +23,7 @@ THE SOFTWARE.
/**
* @addtogroup hipMemRelease hipMemRelease
* @{
* @ingroup MemoryTest
* @ingroup VirtualMemoryManagementTest
* `hipMemRelease(hipMemGenericAllocationHandle_t handle)` -
* Release a memory handle representing a memory allocation which was previously
* allocated through hipMemCreate.
@@ -36,13 +36,11 @@ THE SOFTWARE.
* ------------------------
* - Negative Tests
* ------------------------
* - catch\unit\memory\hipMemRelease.cc
* - unit/virtualMemoryManagement/hipMemRelease.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemRelease_negative") {
SECTION("Nullptr to handle") {
REQUIRE(hipMemRelease(nullptr) == hipErrorInvalidValue);
}
SECTION("Nullptr to handle") { REQUIRE(hipMemRelease(nullptr) == hipErrorInvalidValue); }
}
projects/hip-tests/catch/unit/memory/hipMemRetainAllocationHandle.cc → projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc
+24 -27
View File
@@ -23,7 +23,7 @@ THE SOFTWARE.
/**
* @addtogroup hipMemRetainAllocationHandle hipMemRetainAllocationHandle
* @{
* @ingroup MemoryTest
* @ingroup VirtualMemoryManagementTest
* `hipError_t hipMemRetainAllocationHandle(hipMemGenericAllocationHandle_t* handle,
* void* addr)` -
* Returns the allocation handle of the backing memory allocation given the address.
@@ -31,6 +31,7 @@ THE SOFTWARE.
#include <hip_test_kernels.hh>
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
#define DATA_SIZE (1 << 13)
@@ -41,41 +42,39 @@ THE SOFTWARE.
* - Create a VM mapped to physical memory. Input addr to
* hipMemRetainAllocationHandle and validate the handle.
* ------------------------
* - catch\unit\memory\hipMemRetainAllocationHandle.cc
* - unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemRetainAllocationHandle_SetGet") {
HIP_CHECK(hipFree(0));
size_t granularity = 0;
constexpr int N = DATA_SIZE;
size_t buffer_size = N * sizeof(int);
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem,
0, 0, 0));
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
// Test hipMemRetainAllocationHandle
hipMemGenericAllocationHandle_t gethandle;
// Check beginning of VMM ptr
HIP_CHECK(hipMemRetainAllocationHandle(&gethandle,
reinterpret_cast<void*>(ptrA)));
HIP_CHECK(hipMemRetainAllocationHandle(&gethandle, reinterpret_cast<void*>(ptrA)));
REQUIRE(gethandle == handle);
HIP_CHECK(hipMemRelease(handle));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
@@ -87,58 +86,56 @@ TEST_CASE("Unit_hipMemRetainAllocationHandle_SetGet") {
* ------------------------
* - Negative Tests
* ------------------------
* - catch\unit\memory\hipMemRetainAllocationHandle.cc
* - unit/virtualMemoryManagement/hipMemRetainAllocationHandle.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
*/
TEST_CASE("Unit_hipMemRetainAllocationHandle_NegTst") {
HIP_CHECK(hipFree(0));
size_t granularity = 0;
constexpr int N = DATA_SIZE;
size_t buffer_size = N * sizeof(int);
int deviceId = 0;
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
// Allocate physical memory
HIP_CHECK(hipMemCreate(&handle, size_mem, &prop, 0));
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem,
0, 0, 0));
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
// Test hipMemRetainAllocationHandle
hipMemGenericAllocationHandle_t gethandle;
SECTION("nullptr handle") {
REQUIRE(hipMemRetainAllocationHandle(nullptr,
reinterpret_cast<void*>(ptrA)) == hipErrorInvalidValue);
REQUIRE(hipMemRetainAllocationHandle(nullptr, reinterpret_cast<void*>(ptrA)) ==
hipErrorInvalidValue);
}
SECTION("nullptr Vmm ptr") {
REQUIRE(hipMemRetainAllocationHandle(&gethandle, nullptr) ==
hipErrorInvalidValue);
REQUIRE(hipMemRetainAllocationHandle(&gethandle, nullptr) == hipErrorInvalidValue);
}
SECTION("not mapped address") {
hipDeviceptr_t ptrB;
HIP_CHECK(hipMemAddressReserve(&ptrB, size_mem, 0, 0, 0));
REQUIRE(hipMemRetainAllocationHandle(&gethandle,
reinterpret_cast<void*>(ptrB)) == hipErrorInvalidValue);
REQUIRE(hipMemRetainAllocationHandle(&gethandle, reinterpret_cast<void*>(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<void*>(ptrA)) == hipErrorInvalidValue);
REQUIRE(hipMemRetainAllocationHandle(&gethandle, reinterpret_cast<void*>(ptrA)) ==
hipErrorInvalidValue);
}
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
}
projects/hip-tests/catch/unit/memory/hipMemSetGetAccess.cc → projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemSetGetAccess.cc
+203 -253
View File
File diff suppressed because it is too large Load Diff
projects/hip-tests/catch/unit/memory/hipMemUnmap.cc → projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemUnmap.cc
+13 -13
View File
@@ -23,22 +23,24 @@ THE SOFTWARE.
/**
* @addtogroup hipMemUnmap hipMemUnmap
* @{
* @ingroup MemoryTest
* @ingroup VirtualMemoryManagementTest
* `hipError_t hipMemUnmap (void* ptr, size_t size)` -
* Unmap memory allocation of a given address range.
*/
#include <hip_test_common.hh>
#include "hip_vmm_common.hh"
constexpr int N = (1 << 13);
/**
* Test Description
* ------------------------
* - Negative Tests
* ------------------------
* - catch\unit\memory\hipMemUnmap.cc
* - unit/virtualMemoryManagement/hipMemUnmap.cc
* Test requirements
* ------------------------
* - HIP_VERSION >= 6.1
@@ -50,18 +52,17 @@ TEST_CASE("Unit_hipMemUnmap_negative") {
hipDevice_t device;
HIP_CHECK(hipDeviceGet(&device, deviceId));
checkVMMSupported(device)
checkVMMSupported(device);
hipMemAllocationProp prop{};
prop.type = hipMemAllocationTypePinned;
prop.location.type = hipMemLocationTypeDevice;
prop.location.id = device; // Current Devices
HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &prop,
hipMemAllocationGranularityMinimum));
HIP_CHECK(
hipMemGetAllocationGranularity(&granularity, &prop, hipMemAllocationGranularityMinimum));
REQUIRE(granularity > 0);
size_t size_mem =
((granularity + buffer_size - 1) / granularity) * granularity;
size_t size_mem = ((granularity + buffer_size - 1) / granularity) * granularity;
hipMemGenericAllocationHandle_t handle;
hipDeviceptr_t ptrA;
@@ -70,18 +71,17 @@ TEST_CASE("Unit_hipMemUnmap_negative") {
// Allocate virtual address range
HIP_CHECK(hipMemAddressReserve(&ptrA, size_mem, 0, 0, 0));
HIP_CHECK(hipMemMap(ptrA, size_mem, 0, handle, 0));
SECTION("nullptr to ptrA") {
REQUIRE(hipMemUnmap(nullptr, size_mem) == hipErrorInvalidValue);
}
SECTION("pass zero to size") {
REQUIRE(hipMemUnmap(ptrA, 0) == hipErrorInvalidValue);
}
SECTION("nullptr to ptrA") { REQUIRE(hipMemUnmap(nullptr, size_mem) == hipErrorInvalidValue); }
SECTION("pass zero to size") { REQUIRE(hipMemUnmap(ptrA, 0) == hipErrorInvalidValue); }
#if HT_NVIDIA
SECTION("unmap a smaller size") {
REQUIRE(hipMemUnmap(ptrA, (size_mem - 1)) == hipErrorInvalidValue);
}
#endif
HIP_CHECK(hipMemRelease(handle));
HIP_CHECK(hipMemUnmap(ptrA, size_mem));
HIP_CHECK(hipMemAddressFree(ptrA, size_mem));
projects/hip-tests/catch/unit/memory/hipMemVmm.cc → projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc
+30 -25
View File
@@ -1,30 +1,30 @@
/*
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
Copyright (c) 2022 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
/* Test Case Description:
1) This testcase verifies the basic scenario - supported on
all devices
*/
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
#include <cstdio>
#include <cstdint>
#include <algorithm>
@@ -32,6 +32,10 @@
#include <chrono>
#include <vector>
#include <hip_test_common.hh>
#include <hip_test_kernels.hh>
#include <hip_test_checkers.hh>
/*
This testcase verifies HIP Mem VMM API basic scenario - supported on all devices
*/
@@ -42,8 +46,9 @@ TEST_CASE("Unit_hipMemVmm_Basic") {
INFO("hipDeviceAttributeVirtualMemoryManagementSupported: " << vmm);
if (vmm == 0) {
SUCCEED("GPU 0 doesn't support hipDeviceAttributeVirtualMemoryManagement "
"attribute. Hence skipping the testing with Pass result.\n");
SUCCEED(
"GPU 0 doesn't support hipDeviceAttributeVirtualMemoryManagement "
"attribute. Hence skipping the testing with Pass result.\n");
return;
}
@@ -54,7 +59,8 @@ TEST_CASE("Unit_hipMemVmm_Basic") {
memAllocationProp.location.id = 0;
memAllocationProp.location.type = hipMemLocationTypeDevice;
HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &memAllocationProp, hipMemAllocationGranularityRecommended));
HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &memAllocationProp,
hipMemAllocationGranularityRecommended));
size_t size = 4 * 1024;
void* reservedAddress{nullptr};
@@ -78,7 +84,7 @@ TEST_CASE("Unit_hipMemVmm_Basic") {
HIP_CHECK(hipMemset(reservedAddress, value, size));
HIP_CHECK(hipMemcpy(&values[0], reservedAddress, size, hipMemcpyDeviceToHost));
for (size_t i=0; i < size; ++i) {
for (size_t i = 0; i < size; ++i) {
REQUIRE(values[i] == value);
}
@@ -87,4 +93,3 @@ TEST_CASE("Unit_hipMemVmm_Basic") {
HIP_CHECK(hipMemRelease(gaHandle));
HIP_CHECK(hipMemAddressFree(reservedAddress, size));
}
projects/hip-tests/catch/unit/virtualMemoryManagement/hip_vmm_common.hh
+49
View File
@@ -0,0 +1,49 @@
/*
Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved.
Permission is hereby granted, free of charge, to any person obtaining a copy
of this software and associated documentation files (the "Software"), to deal
in the Software without restriction, including without limitation the rights
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
copies of the Software, and to permit persons to whom the Software is
furnished to do so, subject to the following conditions:
The above copyright notice and this permission notice shall be included in
all copies or substantial portions of the Software.
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
THE SOFTWARE.
*/
#pragma once
#include "hip_test_context.hh"
#define checkVMMSupported(device) \
{ \
int value = 0; \
hipDeviceAttribute_t attr = hipDeviceAttributeVirtualMemoryManagementSupported; \
HIP_CHECK(hipDeviceGetAttribute(&value, attr, device)); \
if (value == 0) { \
HipTest::HIP_SKIP_TEST("Machine does not support VMM. Skipping Test.."); \
return; \
} \
}
#define checkVMMSupportedRetVal(device) \
{ \
int value = 0; \
hipDeviceAttribute_t attr = hipDeviceAttributeVirtualMemoryManagementSupported; \
HIP_CHECK(hipDeviceGetAttribute(&value, attr, device)); \
if (value == 0) { \
HipTest::HIP_SKIP_TEST("Machine does not support VMM. Skipping Test.."); \
return true; \
} \
}
constexpr int threadsPerBlk = 64;