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

SWDEV-499927 - Enable Virtual Memory tests on NV platform (#79)

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