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

SWDEV-534207 - solve phoenix mem test failures (#218)

Browse Source 
* make sure symbolStatus after hipGetProcAddress() is the expected in the tests that are failing on Phoenix
* fix that calls to hipMemset2DAsync() and hipMemset2DAsync_spt() to set values on submatrices where not taking into account the pitch, causing test failures on Phoenix
* before testing whether hipMemset2DAsync() works, initialize the whole matrix to a known value. This makes sure the test fails even if the uninitialized memory returned after hipMallocPitch() happens to have the expected value
* Sbefore testing whether hipMemset2DAsync_spt() works, initialize the whole matrix to a known value. This makes sure the test fails even if the uninitialized memory returned after hipMallocPitch() happens to have the expected value
* fix up compiler error: dyn_hipMemset2DAsync_ptr should have been dyn_hipMemset2DAsync_spt_ptr
* when Unit_hipMalloc3D_Basic fails due to a potential memory leak, print the values that hipMemGetInfo() returns before and after the allocation/deallocation pair.
* Also print intermediate free memory
* Make sure Unit_hipMalloc3D_Basic allocates at least PalSettings::subAllocationChunkSize_ on Windows. Otherwise hipMemGetInfo() will not report an increase on available memory after hipFree() is called, as the minimum amount of memory that cause a call to the CoarseMemorySubAllocator::Create() is that chunk size
* Fix up previous commit; allocated too much memory
* skip some extent sizes in Unit_hipFreeMipmappedArrayImplicitSyncArray if allocating them would be require more memory than the actual totalGlobalMem of the device
* Do not expect an exact match when comparing the memory available memory before and after hipMalloc3D() + hipFree()
* Do not allocate more memory than the total GPU memory in Unit_hipFreeMipmappedArrayImplicitSyncArray
* fix expected available memory amount comparison in hipMalloc3D - Basic
* use SUCCEED() macro in Unit_hipFreeMipmappedArrayImplicitSyncArray to log more information when there is not enough memory for the mipmapped array to be allocated
* fix formatting
This commit is contained in:
Hernandez, Gerardo
2025-07-18 09:15:19 +01:00
committed by GitHub
parent 0c318b4ca5
commit 4b2ed7653f
4 changed files with 112 additions and 54 deletions
Download Patch File Download Diff File Expand all files Collapse all files
catch/unit/streamperthread/hipGetProcAddressSptApis.cc
+20 -14
View File
@@ -672,22 +672,23 @@ TEST_CASE("Unit_hipGetProcAddress_spt_Memset2D3D") {
void* hipMemset2DAsync_spt_ptr = nullptr;
void* hipMemset3D_spt_ptr = nullptr;
void* hipMemset3DAsync_spt_ptr = nullptr;
hipDriverProcAddressQueryResult symbolStatus = HIP_GET_PROC_ADDRESS_SYMBOL_NOT_FOUND;
int currentHipVersion = 0;
HIP_CHECK(hipRuntimeGetVersion(&currentHipVersion));
HIP_CHECK(hipGetProcAddress("hipMemset2D_spt",
&hipMemset2D_spt_ptr,
currentHipVersion, 0, nullptr));
HIP_CHECK(hipGetProcAddress("hipMemset2DAsync_spt",
&hipMemset2DAsync_spt_ptr,
currentHipVersion, 0, nullptr));
HIP_CHECK(hipGetProcAddress("hipMemset3D_spt",
&hipMemset3D_spt_ptr,
currentHipVersion, 0, nullptr));
HIP_CHECK(hipGetProcAddress("hipMemset3DAsync_spt",
&hipMemset3DAsync_spt_ptr,
currentHipVersion, 0, nullptr));
HIP_CHECK(hipGetProcAddress("hipMemset2D_spt", &hipMemset2D_spt_ptr, currentHipVersion, 0,
&symbolStatus));
REQUIRE(symbolStatus == HIP_GET_PROC_ADDRESS_SUCCESS);
HIP_CHECK(hipGetProcAddress("hipMemset2DAsync_spt", &hipMemset2DAsync_spt_ptr, currentHipVersion,
0, &symbolStatus));
REQUIRE(symbolStatus == HIP_GET_PROC_ADDRESS_SUCCESS);
HIP_CHECK(hipGetProcAddress("hipMemset3D_spt", &hipMemset3D_spt_ptr, currentHipVersion, 0,
&symbolStatus));
REQUIRE(symbolStatus == HIP_GET_PROC_ADDRESS_SUCCESS);
HIP_CHECK(hipGetProcAddress("hipMemset3DAsync_spt", &hipMemset3DAsync_spt_ptr, currentHipVersion,
0, &symbolStatus));
REQUIRE(symbolStatus == HIP_GET_PROC_ADDRESS_SUCCESS);
hipError_t (*dyn_hipMemset2D_spt_ptr)(void *, size_t, int, size_t, size_t) =
reinterpret_cast<hipError_t (*)(void *, size_t, int, size_t, size_t)>
@@ -747,12 +748,17 @@ TEST_CASE("Unit_hipGetProcAddress_spt_Memset2D3D") {
HIP_CHECK(hipStreamCreate(&stream[s]));
}
// set the whole matrix first to something different than 'value'
HIP_CHECK(dyn_hipMemset2DAsync_spt_ptr(devMem, pitch, 5, width, height, 0));
HIP_CHECK(hipStreamSynchronize(0));
for ( int s = 0; s < Ns; s++ ) {
int startIndex = s * (N/Ns);
HIP_CHECK(dyn_hipMemset2DAsync_spt_ptr(devMem + startIndex, pitch/Ns,
value, width/Ns, height/Ns,
int row = startIndex / width;
HIP_CHECK(dyn_hipMemset2DAsync_spt_ptr(devMem + row * pitch, pitch, value, width, height / Ns,
stream[s]));
}
for ( int s = 0; s < Ns; s++ ) {
HIP_CHECK(hipStreamSynchronize(stream[s]));
}