From 3e49440495ce6945bcfb59aef35da6ed90ea16a4 Mon Sep 17 00:00:00 2001 From: marandje Date: Tue, 23 Dec 2025 10:27:42 +0100 Subject: [PATCH] SWDEV-555178 - Calculate phys mem offset for remap range (#1879) --- projects/clr/rocclr/device/pal/palvirtual.cpp | 13 ++- .../hipTestMain/config/config_amd_windows | 2 - .../hip-tests/catch/unit/memory/hipMemVmm.cc | 4 +- .../virtualMemoryManagement/CMakeLists.txt | 3 +- .../virtualMemoryManagement/hipMemVmm_old.cc | 91 ------------------- 5 files changed, 11 insertions(+), 102 deletions(-) delete mode 100644 projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc diff --git a/projects/clr/rocclr/device/pal/palvirtual.cpp b/projects/clr/rocclr/device/pal/palvirtual.cpp index 38342c0f64..3946d5b729 100644 --- a/projects/clr/rocclr/device/pal/palvirtual.cpp +++ b/projects/clr/rocclr/device/pal/palvirtual.cpp @@ -2319,12 +2319,18 @@ void VirtualGPU::submitVirtualMap(amd::VirtualMapCommand& vcmd) { // Create a view, since original base obj will map the whole memory and multimap cases wont work. amd::Memory* vaddr_sub_obj = nullptr; + Pal::IGpuMemory* phymem_igpu_mem = nullptr; size_t vaddr_offset = 0; + size_t phys_offset = 0; if (phys_mem_obj != nullptr) { constexpr bool kParent = false; vaddr_sub_obj = phys_mem_obj->getContext().devices()[0]->CreateVirtualBuffer( phys_mem_obj->getContext(), const_cast(vcmd.ptr()), vcmd.size(), phys_mem_obj->getUserData().deviceId, phys_mem_obj->getUserData().locationType, kParent); + + pal::Memory* phys_pal_mem = dev().getGpuMemory(phys_mem_obj); + phymem_igpu_mem = phys_pal_mem->iMem(); + phys_offset = phys_pal_mem->offset(); } else { vaddr_sub_obj = amd::MemObjMap::FindMemObj(vcmd.ptr()); } @@ -2335,11 +2341,8 @@ void VirtualGPU::submitVirtualMap(amd::VirtualMapCommand& vcmd) { // The imem() in the backend is shared between base and sub/view object. pal::Memory* vaddr_pal_mem = dev().getGpuMemory(vaddr_base_obj); - Pal::IGpuMemory* phymem_igpu_mem = - (phys_mem_obj == nullptr) ? nullptr : dev().getGpuMemory(phys_mem_obj)->iMem(); - Pal::VirtualMemoryRemapRange range{vaddr_pal_mem->iMem(), vaddr_offset, - phymem_igpu_mem, 0, + phymem_igpu_mem, phys_offset, vcmd.size(), Pal::VirtualGpuMemAccessMode::NoAccess}; // Wait for previous operations before unmap @@ -2364,7 +2367,7 @@ void VirtualGPU::submitVirtualMap(amd::VirtualMapCommand& vcmd) { phys_mem_obj->getUserData().vaddr_mem_obj = vaddr_sub_obj; } else { // assert the vaddr_mem_obj is mapped and needs to be removed - amd::Memory* vaddr_sub_obj = amd::MemObjMap::FindMemObj(vcmd.ptr()); + vaddr_sub_obj = amd::MemObjMap::FindMemObj(vcmd.ptr()); assert(vaddr_sub_obj != nullptr); assert(vcmd.ptr() == vaddr_sub_obj->getSvmPtr()); diff --git a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows index d03d9322cf..162ec6080a 100644 --- a/projects/hip-tests/catch/hipTestMain/config/config_amd_windows +++ b/projects/hip-tests/catch/hipTestMain/config/config_amd_windows @@ -177,7 +177,6 @@ "Unit_hipStreamCreateWithPriority_MulthreadNonblockingflag", "SWDEV-396617 ExecMemcpyNodeSetParamsFromSymbol fails in direction", "SWDEV-396616 hipMemMap returns invalid error", - "Unit_hipMemVmm_Basic", "SWDEV-396615 mGPUs not considered correctly", "Unit_hipManagedKeyword_MultiGpu", "Disabling test tracked SWDEV-391555", @@ -699,7 +698,6 @@ "====================================================", "Unit_hipEventRecord", "========Rock_Window_Failures_on_gfx1151===========================================", - "Unit_hipMemVmm_Uncached", "Unit_Uuid_FntlTstsFor_SetEnv_HIP_VISIBLE_DEVICES", "Unit_UUID_setEnv_Thread", "Unit_hipTexRefSetArray_Positive", diff --git a/projects/hip-tests/catch/unit/memory/hipMemVmm.cc b/projects/hip-tests/catch/unit/memory/hipMemVmm.cc index 1c1fcb025b..2d9b74a4df 100644 --- a/projects/hip-tests/catch/unit/memory/hipMemVmm.cc +++ b/projects/hip-tests/catch/unit/memory/hipMemVmm.cc @@ -58,9 +58,9 @@ TEST_CASE("Unit_hipMemVmm_Basic") { HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &memAllocationProp, hipMemAllocationGranularityRecommended)); - size_t size = 4 * 1024; + size_t size = granularity * 4; void* reservedAddress{nullptr}; - HIP_CHECK(hipMemAddressReserve(&reservedAddress, size, granularity, nullptr, 0)); + HIP_CHECK(hipMemAddressReserve(&reservedAddress, size, 0, nullptr, 0)); hipMemGenericAllocationHandle_t gaHandle{nullptr}; HIP_CHECK(hipMemCreate(&gaHandle, size, &memAllocationProp, 0)); diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt b/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt index 3bf34bb852..6796f11006 100644 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt +++ b/projects/hip-tests/catch/unit/virtualMemoryManagement/CMakeLists.txt @@ -55,8 +55,7 @@ set(TEST_SRC hipMemGetAllocationPropertiesFromHandle.cc hipMemMap.cc hipMemRelease.cc - hipMemUnmap.cc - hipMemVmm_old.cc) + hipMemUnmap.cc) hip_add_exe_to_target(NAME VirtualMemoryManagementTest TEST_SRC ${TEST_SRC} diff --git a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc b/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc deleted file mode 100644 index cc22d56723..0000000000 --- a/projects/hip-tests/catch/unit/virtualMemoryManagement/hipMemVmm_old.cc +++ /dev/null @@ -1,91 +0,0 @@ -/* -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 - -#include -#include -#include - -/* - 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); - - if (vmm == 0) { - SUCCEED( - "GPU 0 doesn't support hipDeviceAttributeVirtualMemoryManagement " - "attribute. Hence skipping the testing with Pass result.\n"); - return; - } - - size_t granularity = 0; - - hipMemAllocationProp memAllocationProp{}; - memAllocationProp.type = hipMemAllocationTypePinned; - memAllocationProp.location.id = 0; - memAllocationProp.location.type = hipMemLocationTypeDevice; - - HIP_CHECK(hipMemGetAllocationGranularity(&granularity, &memAllocationProp, - hipMemAllocationGranularityRecommended)); - - size_t size = granularity; - void* reservedAddress{nullptr}; - HIP_CHECK(hipMemAddressReserve(&reservedAddress, size, 0, nullptr, 0)); - - hipMemGenericAllocationHandle_t gaHandle; - HIP_CHECK(hipMemCreate(&gaHandle, size, &memAllocationProp, 0)); - - HIP_CHECK(hipMemMap(reservedAddress, size, 0, gaHandle, 0)); - - hipDevice_t device; - HIP_CHECK(hipDeviceGet(&device, 0)); - hipMemAccessDesc desc; - desc.location.type = hipMemLocationTypeDevice; - desc.location.id = device; - desc.flags = hipMemAccessFlagsProtReadWrite; - std::vector values(size); - const char value = 1; - - HIP_CHECK(hipMemSetAccess(reservedAddress, size, &desc, 1)); - HIP_CHECK(hipMemset(reservedAddress, value, size)); - HIP_CHECK(hipMemcpy(&values[0], reservedAddress, size, hipMemcpyDeviceToHost)); - - for (size_t i = 0; i < size; ++i) { - REQUIRE(values[i] == value); - } - - HIP_CHECK(hipMemUnmap(reservedAddress, size)); - - HIP_CHECK(hipMemRelease(gaHandle)); - HIP_CHECK(hipMemAddressFree(reservedAddress, size)); - CTX_DESTROY(); -}