From 31296c81fce9162da8d15482361caab6901027f3 Mon Sep 17 00:00:00 2001 From: Adrien Taberner Date: Thu, 13 Feb 2025 17:08:39 +0100 Subject: [PATCH 1/7] MemInfo initial commit --- cmake/modules/FindKokkos.cmake | 57 +++++++++++++++++++++++++ memInfo/CMakeLists.txt | 19 +++++++++ memInfo/src/CMakeLists.txt | 3 ++ memInfo/src/cexa_MemInfo.hpp | 69 +++++++++++++++++++++++++++++++ memInfo/unit_test/CMakeLists.txt | 16 +++++++ memInfo/unit_test/TestMemInfo.cpp | 46 +++++++++++++++++++++ 6 files changed, 210 insertions(+) create mode 100644 cmake/modules/FindKokkos.cmake create mode 100644 memInfo/CMakeLists.txt create mode 100644 memInfo/src/CMakeLists.txt create mode 100644 memInfo/src/cexa_MemInfo.hpp create mode 100644 memInfo/unit_test/CMakeLists.txt create mode 100644 memInfo/unit_test/TestMemInfo.cpp diff --git a/cmake/modules/FindKokkos.cmake b/cmake/modules/FindKokkos.cmake new file mode 100644 index 0000000..2603556 --- /dev/null +++ b/cmake/modules/FindKokkos.cmake @@ -0,0 +1,57 @@ +# set CMAKE_BUILD_TYPE if not defined +if(NOT CMAKE_BUILD_TYPE) + set(default_build_type "RelWithDebInfo") + message(STATUS "Setting build type to '${default_build_type}' as none was specified.") + set( + CMAKE_BUILD_TYPE + "${default_build_type}" + CACHE + STRING + "Choose the type of build, options are: Debug, Release, RelWithDebInfo and MinSizeRel." + FORCE + ) +endif() + +# find Kokkos as an already existing target +if(TARGET Kokkos::kokkos) + return() +endif() + +# find Kokkos as installed +find_package(Kokkos CONFIG) +if(Kokkos_FOUND) + message(STATUS "Kokkos provided as installed: ${Kokkos_DIR} (version \"${Kokkos_VERSION}\")") + + return() +endif() + +# find Kokkos as an existing source directory +set( + CexaExperimental_KOKKOS_SOURCE_DIR + "${CMAKE_CURRENT_SOURCE_DIR}/../../../vendor/kokkos" + CACHE + PATH + "Path to the local source directory of Kokkos" +) +if(EXISTS "${CexaExperimental_KOKKOS_SOURCE_DIR}/CMakeLists.txt") + message(STATUS "Kokkos provided as a source directory: ${CexaExperimental_KOKKOS_SOURCE_DIR}") + + add_subdirectory("${CexaExperimental_KOKKOS_SOURCE_DIR}" kokkos) + set(Kokkos_FOUND True) + + return() +endif() + +# download Kokkos from release and find it +message(STATUS "Kokkos downloaded: ${CexaExperimental_KOKKOS_SOURCE_DIR}") + +include(FetchContent) + +FetchContent_Declare( + kokkos + DOWNLOAD_EXTRACT_TIMESTAMP ON + URL https://github.com/kokkos/kokkos/releases/download/4.5.01/kokkos-4.5.01.zip + SOURCE_DIR ${CexaExperimental_KOKKOS_SOURCE_DIR} +) +FetchContent_MakeAvailable(kokkos) +set(Kokkos_FOUND True) diff --git a/memInfo/CMakeLists.txt b/memInfo/CMakeLists.txt new file mode 100644 index 0000000..a39479e --- /dev/null +++ b/memInfo/CMakeLists.txt @@ -0,0 +1,19 @@ +cmake_minimum_required(VERSION 3.16) + +project(cexa-experimental-meminfo LANGUAGES CXX) +set(CMAKE_BUILD_TYPE "RelwithDebInfo") + +if (NOT ${Kokkos_DIR} STREQUAL "") + add_subdirectory(${Kokkos_DIR}) + include_directories(${Kokkos_INCLUDE_DIRS_RET}) +endif() + +list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/../cmake/modules") + +include(CTest) + +find_package(Kokkos REQUIRED) + +add_subdirectory(src) +add_subdirectory(unit_test) + diff --git a/memInfo/src/CMakeLists.txt b/memInfo/src/CMakeLists.txt new file mode 100644 index 0000000..fbd9863 --- /dev/null +++ b/memInfo/src/CMakeLists.txt @@ -0,0 +1,3 @@ +add_library(memInfo INTERFACE) +target_include_directories(memInfo INTERFACE ${CMAKE_CURRENT_SOURCE_DIR}) +target_link_libraries(memInfo INTERFACE Kokkos::kokkos) diff --git a/memInfo/src/cexa_MemInfo.hpp b/memInfo/src/cexa_MemInfo.hpp new file mode 100644 index 0000000..776a853 --- /dev/null +++ b/memInfo/src/cexa_MemInfo.hpp @@ -0,0 +1,69 @@ +#include +#include + +#ifdef _WIN32 +#include +#else +#include +#endif + +namespace Kokkos { +namespace Experimental { + +template +void MemGetInfo(size_t* free, size_t* total) { + *free = 0; + *total = 0; +} + +// Single node memory info +template <> +void MemGetInfo(size_t* free, size_t* total) { +#ifdef _WIN32 + MEMORYSTATUSEX statex; + statex.dwLength = sizeof(statex); + if (GlobalMemoryStatusEx(&statex) != 0) { + *free = statex.ullAvailPhys; + *total = statex.ullTotalPhys; + } +#else + struct sysinfo info; + if (sysinfo(&info) == 0) { + *free = info.freeram * info.mem_unit; + *total = info.totalram * info.mem_unit; + } +#endif +} + +#if defined(KOKKOS_ENABLE_CUDA) +template <> +void MemGetInfo(size_t* free, size_t* total) { + KOKKOS_IMPL_CUDA_SAFE_CALL(cudaMemGetInfo(free, total)); +} +#endif + +#if defined(KOKKOS_ENABLE_HIP) +template <> +void MemGetInfo(size_t* free, size_t* total) { + KOKKOS_IMPL_HIP_SAFE_CALL(hipMemGetInfo(free, total)); +} +#endif + +#if defined(KOKKOS_ENABLE_SYCL) +template <> +void MemGetInfo(size_t* free, size_t* total) { + std::vector devices = Kokkos::Impl : get_sycl_devices(); + for (auto& dev : devices) { + if (dev.is_gpu()) { + *total += dev.get_info(); + // https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/doc/extensions/supported/sycl_ext_intel_device_info.md#free-global-memory + if (dev.has(aspect::ext_intel_free_memory)) { + *free += dev.get_info(); + } + } + } +} +#endif + +} // namespace Experimental +} // namespace Kokkos diff --git a/memInfo/unit_test/CMakeLists.txt b/memInfo/unit_test/CMakeLists.txt new file mode 100644 index 0000000..eacd131 --- /dev/null +++ b/memInfo/unit_test/CMakeLists.txt @@ -0,0 +1,16 @@ + +include(FetchContent) + +FetchContent_Declare( + googletest + URL https://github.com/google/googletest/archive/03597a01ee50ed33e9dfd640b249b4be3799d395.zip + DOWNLOAD_EXTRACT_TIMESTAMP true + ) +FetchContent_MakeAvailable(googletest) + +enable_testing() +add_executable(MemInfoTest TestMemInfo.cpp) +target_link_libraries(MemInfoTest GTest::gtest_main Kokkos::kokkos memInfo) + +include(GoogleTest) +gtest_discover_tests(MemInfoTest) diff --git a/memInfo/unit_test/TestMemInfo.cpp b/memInfo/unit_test/TestMemInfo.cpp new file mode 100644 index 0000000..90a101b --- /dev/null +++ b/memInfo/unit_test/TestMemInfo.cpp @@ -0,0 +1,46 @@ +#include +#include +#include +#include +#include + +template +void testMemInfo() { + size_t step1_dram = 0ul; + size_t step2_dram = 0ul; + size_t step2_freeDram = 0ul; + size_t step1_freeDram = 0ul; + Kokkos::Experimental::MemGetInfo(&step1_freeDram, &step1_dram); + { + Kokkos::View data("1 GiB", 1024*1024*1024); + Kokkos::Experimental::MemGetInfo(&step2_freeDram, &step2_dram); + } + // Same total memory before and after allocating 1 GiB + EXPECT_EQ(step1_dram, step2_dram); + // Check that free memory is less after allocating 1 GiB + EXPECT_LT(step2_freeDram, step1_freeDram); +} + +TEST(TestMemInfo, testHost) { + testMemInfo(); +} + +#if defined(KOKKOS_ENABLE_CUDA) +TEST(TestMemInfo, testCuda) { + testMemInfo(); +} +#endif + +#if defined(KOKKOS_ENABLE_HIP) +TEST(TestMemInfo, testHip) { + testMemInfo(); +} +#endif + +int main(int argc, char *argv[]) { + Kokkos::initialize(argc, argv); + ::testing::InitGoogleTest(&argc, argv); + int result = RUN_ALL_TESTS(); + Kokkos::finalize(); + return result; +} \ No newline at end of file From 8fed5f764ebb95f2fd70271afd3a32ce903f1790 Mon Sep 17 00:00:00 2001 From: Adrien Taberner Date: Wed, 19 Feb 2025 09:55:04 +0100 Subject: [PATCH 2/7] Add Readme & sharedSpace --- memInfo/README.md | 9 +++++++++ memInfo/src/cexa_MemInfo.hpp | 8 ++++++++ memInfo/unit_test/TestMemInfo.cpp | 6 ++++++ 3 files changed, 23 insertions(+) create mode 100644 memInfo/README.md diff --git a/memInfo/README.md b/memInfo/README.md new file mode 100644 index 0000000..13cbe63 --- /dev/null +++ b/memInfo/README.md @@ -0,0 +1,9 @@ +### memInfo +This is a wrapper for the device _memGetInfo_ function, but it's also available for unified memory space and host space. + +### Usage +``` +Kokkos::Experimental::memInfo(&free, &total); +``` +- `total`: Amount of RAM on the system (HBM/DRAM) +- `free`: Available memory \ No newline at end of file diff --git a/memInfo/src/cexa_MemInfo.hpp b/memInfo/src/cexa_MemInfo.hpp index 776a853..6177cc5 100644 --- a/memInfo/src/cexa_MemInfo.hpp +++ b/memInfo/src/cexa_MemInfo.hpp @@ -40,6 +40,10 @@ template <> void MemGetInfo(size_t* free, size_t* total) { KOKKOS_IMPL_CUDA_SAFE_CALL(cudaMemGetInfo(free, total)); } +template <> +void MemGetInfo(size_t* free, size_t* total) { + MemGetInfo(free, total); +} #endif #if defined(KOKKOS_ENABLE_HIP) @@ -47,6 +51,10 @@ template <> void MemGetInfo(size_t* free, size_t* total) { KOKKOS_IMPL_HIP_SAFE_CALL(hipMemGetInfo(free, total)); } +template <> +void MemGetInfo(size_t* free, size_t* total) { + MemGetInfo(free, total); +} #endif #if defined(KOKKOS_ENABLE_SYCL) diff --git a/memInfo/unit_test/TestMemInfo.cpp b/memInfo/unit_test/TestMemInfo.cpp index 90a101b..468fd35 100644 --- a/memInfo/unit_test/TestMemInfo.cpp +++ b/memInfo/unit_test/TestMemInfo.cpp @@ -29,12 +29,18 @@ TEST(TestMemInfo, testHost) { TEST(TestMemInfo, testCuda) { testMemInfo(); } +TEST(TestMemInfo, testCudaUVM) { + testMemInfo(); +} #endif #if defined(KOKKOS_ENABLE_HIP) TEST(TestMemInfo, testHip) { testMemInfo(); } +TEST(TestMemInfo, testHipManaged) { + testMemInfo(); +} #endif int main(int argc, char *argv[]) { From 443dbcc66373d5bc14474f972a8687d066d6c8c6 Mon Sep 17 00:00:00 2001 From: Adrien Taberner Date: Wed, 26 Feb 2025 10:17:21 +0100 Subject: [PATCH 3/7] Correct syntax --- memInfo/CMakeLists.txt | 2 +- memInfo/README.md | 2 +- memInfo/src/cexa_MemInfo.hpp | 79 ++++++++++++++++++++++++++++--- memInfo/unit_test/TestMemInfo.cpp | 76 ++++++++++++++++++++--------- 4 files changed, 128 insertions(+), 31 deletions(-) diff --git a/memInfo/CMakeLists.txt b/memInfo/CMakeLists.txt index a39479e..8444c07 100644 --- a/memInfo/CMakeLists.txt +++ b/memInfo/CMakeLists.txt @@ -1,7 +1,7 @@ cmake_minimum_required(VERSION 3.16) project(cexa-experimental-meminfo LANGUAGES CXX) -set(CMAKE_BUILD_TYPE "RelwithDebInfo") +set(CMAKE_BUILD_TYPE "RelWithDebInfo") if (NOT ${Kokkos_DIR} STREQUAL "") add_subdirectory(${Kokkos_DIR}) diff --git a/memInfo/README.md b/memInfo/README.md index 13cbe63..55ad3e3 100644 --- a/memInfo/README.md +++ b/memInfo/README.md @@ -3,7 +3,7 @@ This is a wrapper for the device _memGetInfo_ function, but it's also available ### Usage ``` -Kokkos::Experimental::memInfo(&free, &total); +Kokkos::Experimental::memInfo(&free, &total); ``` - `total`: Amount of RAM on the system (HBM/DRAM) - `free`: Available memory \ No newline at end of file diff --git a/memInfo/src/cexa_MemInfo.hpp b/memInfo/src/cexa_MemInfo.hpp index 6177cc5..3c75322 100644 --- a/memInfo/src/cexa_MemInfo.hpp +++ b/memInfo/src/cexa_MemInfo.hpp @@ -1,5 +1,7 @@ #include #include +#include +#include #ifdef _WIN32 #include @@ -10,10 +12,68 @@ namespace Kokkos { namespace Experimental { -template +// On some systems, overcommit is disabled, and the kernel will not allow memory +// allocation beyond the commit limit. This means that allocations that only touch +// a small amount of memory will be accounted for at their full allocation size. +bool is_overcommit_limit_set() { + std::ifstream overcommit_file("/proc/sys/vm/overcommit_memory"); + int overcommit_value = 0; + + if (overcommit_file.is_open()) { + overcommit_file >> overcommit_value; + overcommit_file.close(); + } else { + return false; + } + return (overcommit_value == 2); +} + +size_t get_committed_as() { + std::ifstream meminfo("/proc/meminfo"); + size_t committed_as = 0; + char line[256]; + + if (meminfo.is_open()) { + while (meminfo.getline(line, 256)) { + if (strncmp(line, "Committed_AS:", 13) == 0) { + std::istringstream iss(line); + iss.ignore(256, ':'); + iss >> committed_as; + committed_as = (iss.fail()) ? 0 : committed_as * 1024; + meminfo.close(); + break; + } + } + meminfo.close(); + } + return committed_as; +} + +size_t get_commitLimit() { + std::ifstream meminfo("/proc/meminfo"); + size_t commitLimit = 0; + char line[256]; + + if (meminfo.is_open()) { + while (meminfo.getline(line, 256)) { + if (strncmp(line, "CommitLimit:", 12) == 0) { + std::istringstream iss(line); + iss.ignore(256, ':'); + iss >> commitLimit; + commitLimit = (iss.fail()) ? 0 : commitLimit * 1024; + meminfo.close(); + break; + } + } + meminfo.close(); + } + return commitLimit; +} + +template void MemGetInfo(size_t* free, size_t* total) { - *free = 0; - *total = 0; + using MemorySpace = typename Space::memory_space; + MemGetInfo(free, total); } // Single node memory info @@ -27,10 +87,17 @@ void MemGetInfo(size_t* free, size_t* total) { *total = statex.ullTotalPhys; } #else + static bool overcommit_limit = is_overcommit_limit_set(); struct sysinfo info; + if (overcommit_limit) { + *total = get_commitLimit(); + *free = *total - get_committed_as(); + return; + } if (sysinfo(&info) == 0) { *free = info.freeram * info.mem_unit; *total = info.totalram * info.mem_unit; + return; } #endif } @@ -60,13 +127,13 @@ void MemGetInfo(size_t* free, size_t* total) { #if defined(KOKKOS_ENABLE_SYCL) template <> void MemGetInfo(size_t* free, size_t* total) { - std::vector devices = Kokkos::Impl : get_sycl_devices(); + std::vector devices = Kokkos::Impl::get_sycl_devices(); for (auto& dev : devices) { if (dev.is_gpu()) { *total += dev.get_info(); // https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/doc/extensions/supported/sycl_ext_intel_device_info.md#free-global-memory - if (dev.has(aspect::ext_intel_free_memory)) { - *free += dev.get_info(); + if (dev.has(sycl::aspect::ext_intel_free_memory)) { + *free += dev.get_info(); } } } diff --git a/memInfo/unit_test/TestMemInfo.cpp b/memInfo/unit_test/TestMemInfo.cpp index 468fd35..67b2e43 100644 --- a/memInfo/unit_test/TestMemInfo.cpp +++ b/memInfo/unit_test/TestMemInfo.cpp @@ -1,52 +1,82 @@ -#include -#include #include #include #include +#include -template +template void testMemInfo() { - size_t step1_dram = 0ul; - size_t step2_dram = 0ul; - size_t step2_freeDram = 0ul; - size_t step1_freeDram = 0ul; - Kokkos::Experimental::MemGetInfo(&step1_freeDram, &step1_dram); + size_t step1_total = 0ul; + size_t step2_total = 0ul; + size_t step2_free = 0ul; + size_t step1_free = 0ul; + + if constexpr (std::is_same::value) { + Kokkos::Experimental::MemGetInfo(&step1_free, &step1_total); + } else { + Kokkos::Experimental::MemGetInfo(&step1_free, &step1_total); + } + volatile double k = 0.0; { - Kokkos::View data("1 GiB", 1024*1024*1024); - Kokkos::Experimental::MemGetInfo(&step2_freeDram, &step2_dram); + // Allocate 128 MiB of memory + Kokkos::View data("data test", 128, 1024*1024); + if constexpr (std::is_same::value) { + Kokkos::Experimental::MemGetInfo(&step2_free, &step2_total); + } else { + Kokkos::Experimental::MemGetInfo(&step2_free, &step2_total); + } } - // Same total memory before and after allocating 1 GiB - EXPECT_EQ(step1_dram, step2_dram); - // Check that free memory is less after allocating 1 GiB - EXPECT_LT(step2_freeDram, step1_freeDram); + // Same total memory before and after aloccation + EXPECT_EQ(step1_total, step2_total); + // Check that free memory is less after allocation + EXPECT_LT(step2_free, step1_free); } -TEST(TestMemInfo, testHost) { +TEST(MemInfo, HostSpace) { + Kokkos::initialize(); testMemInfo(); + Kokkos::finalize(); +} +TEST(MemInfo, HostSpaceUninitialized) { + size_t free = 0; + size_t total = 0; + Kokkos::Experimental::MemGetInfo(&free, &total); + EXPECT_GT(free, 0); + EXPECT_GT(total, 0); +} + +TEST(MemInfo, DefaultSpace) { + Kokkos::initialize(); + testMemInfo<>(); + Kokkos::finalize(); +} +TEST(MemInfo, DefaultSpaceUninitialized) { + size_t free = 0; + size_t total = 0; + Kokkos::Experimental::MemGetInfo(&free, &total); + EXPECT_GT(free, 0); + EXPECT_GT(total, 0); } #if defined(KOKKOS_ENABLE_CUDA) -TEST(TestMemInfo, testCuda) { +TEST(MemInfo, CudaSpace) { + Kokkos::initialize(); testMemInfo(); -} -TEST(TestMemInfo, testCudaUVM) { testMemInfo(); + Kokkos::finalize(); } #endif #if defined(KOKKOS_ENABLE_HIP) -TEST(TestMemInfo, testHip) { +TEST(MemInfo, HIPSpace) { + Kokkos::initialize(); testMemInfo(); -} -TEST(TestMemInfo, testHipManaged) { testMemInfo(); + Kokkos::finalize(); } #endif int main(int argc, char *argv[]) { - Kokkos::initialize(argc, argv); ::testing::InitGoogleTest(&argc, argv); int result = RUN_ALL_TESTS(); - Kokkos::finalize(); return result; } \ No newline at end of file From 5ce64a912cbf4171e4c09a5657eeda37ab7f5385 Mon Sep 17 00:00:00 2001 From: Adrien Taberner Date: Mon, 3 Mar 2025 14:45:21 +0100 Subject: [PATCH 4/7] No call to close() and char* -> string --- memInfo/src/cexa_MemInfo.hpp | 17 ++++++----------- memInfo/unit_test/TestMemInfo.cpp | 3 +-- 2 files changed, 7 insertions(+), 13 deletions(-) diff --git a/memInfo/src/cexa_MemInfo.hpp b/memInfo/src/cexa_MemInfo.hpp index 3c75322..99ed9b4 100644 --- a/memInfo/src/cexa_MemInfo.hpp +++ b/memInfo/src/cexa_MemInfo.hpp @@ -21,7 +21,6 @@ bool is_overcommit_limit_set() { if (overcommit_file.is_open()) { overcommit_file >> overcommit_value; - overcommit_file.close(); } else { return false; } @@ -31,20 +30,18 @@ bool is_overcommit_limit_set() { size_t get_committed_as() { std::ifstream meminfo("/proc/meminfo"); size_t committed_as = 0; - char line[256]; + std::string line; if (meminfo.is_open()) { - while (meminfo.getline(line, 256)) { - if (strncmp(line, "Committed_AS:", 13) == 0) { + while (std::getline(meminfo, line)) { + if (line.find("Committed_AS:") != std::string::npos ){ std::istringstream iss(line); iss.ignore(256, ':'); iss >> committed_as; committed_as = (iss.fail()) ? 0 : committed_as * 1024; - meminfo.close(); break; } } - meminfo.close(); } return committed_as; } @@ -52,20 +49,18 @@ size_t get_committed_as() { size_t get_commitLimit() { std::ifstream meminfo("/proc/meminfo"); size_t commitLimit = 0; - char line[256]; + std::string line; if (meminfo.is_open()) { - while (meminfo.getline(line, 256)) { - if (strncmp(line, "CommitLimit:", 12) == 0) { + while (std::getline(meminfo, line)) { + if (line.find("CommitLimit:") != std::string::npos ){ std::istringstream iss(line); iss.ignore(256, ':'); iss >> commitLimit; commitLimit = (iss.fail()) ? 0 : commitLimit * 1024; - meminfo.close(); break; } } - meminfo.close(); } return commitLimit; } diff --git a/memInfo/unit_test/TestMemInfo.cpp b/memInfo/unit_test/TestMemInfo.cpp index 67b2e43..6619c73 100644 --- a/memInfo/unit_test/TestMemInfo.cpp +++ b/memInfo/unit_test/TestMemInfo.cpp @@ -15,7 +15,6 @@ void testMemInfo() { } else { Kokkos::Experimental::MemGetInfo(&step1_free, &step1_total); } - volatile double k = 0.0; { // Allocate 128 MiB of memory Kokkos::View data("data test", 128, 1024*1024); @@ -79,4 +78,4 @@ int main(int argc, char *argv[]) { ::testing::InitGoogleTest(&argc, argv); int result = RUN_ALL_TESTS(); return result; -} \ No newline at end of file +} From 7678738071124926c4412e88ce27b1b8563cfd88 Mon Sep 17 00:00:00 2001 From: Adrien Taberner Date: Tue, 4 Mar 2025 09:53:56 +0100 Subject: [PATCH 5/7] Combine the two functions into one Add constants --- memInfo/src/cexa_MemInfo.hpp | 60 +++++++++++++++--------------------- 1 file changed, 25 insertions(+), 35 deletions(-) diff --git a/memInfo/src/cexa_MemInfo.hpp b/memInfo/src/cexa_MemInfo.hpp index 99ed9b4..f2b0246 100644 --- a/memInfo/src/cexa_MemInfo.hpp +++ b/memInfo/src/cexa_MemInfo.hpp @@ -12,11 +12,20 @@ namespace Kokkos { namespace Experimental { -// On some systems, overcommit is disabled, and the kernel will not allow memory -// allocation beyond the commit limit. This means that allocations that only touch -// a small amount of memory will be accounted for at their full allocation size. +namespace { + constexpr int OVERCOMMIT_DISABLED = 2; + constexpr char COMMITTED_AS_KEY[] = "Committed_AS:"; + constexpr char COMMIT_LIMIT_KEY[] = "CommitLimit:"; + constexpr char MEMINFO_PATH[] = "/proc/meminfo"; + constexpr char OVERCOMMIT_PATH[] = "/proc/sys/vm/overcommit_memory"; +} + +// On some systems, overcommit is disabled, and the kernel does not allow +// memory allocation beyond the commit limit. This means that allocations +// that touch only a small amount of memory are still counted at their full size. +// man proc_sys_vm bool is_overcommit_limit_set() { - std::ifstream overcommit_file("/proc/sys/vm/overcommit_memory"); + std::ifstream overcommit_file(OVERCOMMIT_PATH); int overcommit_value = 0; if (overcommit_file.is_open()) { @@ -24,48 +33,29 @@ bool is_overcommit_limit_set() { } else { return false; } - return (overcommit_value == 2); -} - -size_t get_committed_as() { - std::ifstream meminfo("/proc/meminfo"); - size_t committed_as = 0; - std::string line; - - if (meminfo.is_open()) { - while (std::getline(meminfo, line)) { - if (line.find("Committed_AS:") != std::string::npos ){ - std::istringstream iss(line); - iss.ignore(256, ':'); - iss >> committed_as; - committed_as = (iss.fail()) ? 0 : committed_as * 1024; - break; - } - } - } - return committed_as; + return (overcommit_value == OVERCOMMIT_DISABLED); } -size_t get_commitLimit() { - std::ifstream meminfo("/proc/meminfo"); - size_t commitLimit = 0; +size_t get_meminfo_value(const char* key) { + std::ifstream meminfo(MEMINFO_PATH); + size_t value = 0; std::string line; if (meminfo.is_open()) { while (std::getline(meminfo, line)) { - if (line.find("CommitLimit:") != std::string::npos ){ + if (line.find(key) != std::string::npos) { std::istringstream iss(line); iss.ignore(256, ':'); - iss >> commitLimit; - commitLimit = (iss.fail()) ? 0 : commitLimit * 1024; + iss >> value; + value = (iss.fail()) ? 0 : value * 1024; break; } } } - return commitLimit; + return value; } -template +template void MemGetInfo(size_t* free, size_t* total) { using MemorySpace = typename Space::memory_space; MemGetInfo(free, total); @@ -83,12 +73,12 @@ void MemGetInfo(size_t* free, size_t* total) { } #else static bool overcommit_limit = is_overcommit_limit_set(); - struct sysinfo info; if (overcommit_limit) { - *total = get_commitLimit(); - *free = *total - get_committed_as(); + *total = get_meminfo_value(COMMIT_LIMIT_KEY); + *free = *total - get_meminfo_value(COMMITTED_AS_KEY); return; } + struct sysinfo info; if (sysinfo(&info) == 0) { *free = info.freeram * info.mem_unit; *total = info.totalram * info.mem_unit; From 45f0022fea73effaf70814a9d490b5cf38d9730c Mon Sep 17 00:00:00 2001 From: Adrien Taberner Date: Wed, 12 Mar 2025 16:31:56 +0100 Subject: [PATCH 6/7] Requested changes Signed-off-by: Adrien Taberner --- memInfo/CMakeLists.txt | 7 ++- memInfo/src/cexa_MemInfo.hpp | 36 +++++++------ memInfo/unit_test/CMakeLists.txt | 7 +-- memInfo/unit_test/TestMemInfo.cpp | 88 ++++++++++++------------------- 4 files changed, 60 insertions(+), 78 deletions(-) diff --git a/memInfo/CMakeLists.txt b/memInfo/CMakeLists.txt index 8444c07..eb20152 100644 --- a/memInfo/CMakeLists.txt +++ b/memInfo/CMakeLists.txt @@ -6,14 +6,13 @@ set(CMAKE_BUILD_TYPE "RelWithDebInfo") if (NOT ${Kokkos_DIR} STREQUAL "") add_subdirectory(${Kokkos_DIR}) include_directories(${Kokkos_INCLUDE_DIRS_RET}) +else() + list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/../cmake/modules") + find_package(Kokkos REQUIRED) endif() -list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/../cmake/modules") - include(CTest) -find_package(Kokkos REQUIRED) - add_subdirectory(src) add_subdirectory(unit_test) diff --git a/memInfo/src/cexa_MemInfo.hpp b/memInfo/src/cexa_MemInfo.hpp index f2b0246..f8b67ae 100644 --- a/memInfo/src/cexa_MemInfo.hpp +++ b/memInfo/src/cexa_MemInfo.hpp @@ -1,30 +1,34 @@ +#ifndef KOKKOS_MEMINFO_HPP +#define KOKKOS_MEMINFO_HPP + #include -#include #include #include #ifdef _WIN32 #include -#else -#include #endif +#include + namespace Kokkos { namespace Experimental { namespace { constexpr int OVERCOMMIT_DISABLED = 2; - constexpr char COMMITTED_AS_KEY[] = "Committed_AS:"; - constexpr char COMMIT_LIMIT_KEY[] = "CommitLimit:"; - constexpr char MEMINFO_PATH[] = "/proc/meminfo"; - constexpr char OVERCOMMIT_PATH[] = "/proc/sys/vm/overcommit_memory"; + constexpr char MEM_FREE_KEY[] = "MemFree:"; + constexpr char MEM_TOTAL_KEY[] = "MemTotal:"; + constexpr char COMMITTED_AS_KEY[] = "Committed_AS:"; + constexpr char COMMIT_LIMIT_KEY[] = "CommitLimit:"; + constexpr char MEMINFO_PATH[] = "/proc/meminfo"; + constexpr char OVERCOMMIT_PATH[] = "/proc/sys/vm/overcommit_memory"; } // On some systems, overcommit is disabled, and the kernel does not allow // memory allocation beyond the commit limit. This means that allocations // that touch only a small amount of memory are still counted at their full size. // man proc_sys_vm -bool is_overcommit_limit_set() { +bool is_overcommit_disabled() { std::ifstream overcommit_file(OVERCOMMIT_PATH); int overcommit_value = 0; @@ -72,17 +76,13 @@ void MemGetInfo(size_t* free, size_t* total) { *total = statex.ullTotalPhys; } #else - static bool overcommit_limit = is_overcommit_limit_set(); - if (overcommit_limit) { + static bool overcommit_disabled = is_overcommit_disabled(); + if (overcommit_disabled) { *total = get_meminfo_value(COMMIT_LIMIT_KEY); *free = *total - get_meminfo_value(COMMITTED_AS_KEY); - return; - } - struct sysinfo info; - if (sysinfo(&info) == 0) { - *free = info.freeram * info.mem_unit; - *total = info.totalram * info.mem_unit; - return; + } else { + *total = get_meminfo_value(MEM_TOTAL_KEY); + *free = get_meminfo_value(MEM_FREE_KEY); } #endif } @@ -127,3 +127,5 @@ void MemGetInfo(size_t* free, size_t* total) { } // namespace Experimental } // namespace Kokkos + +#endif // KOKKOS_MEMINFO_HPP diff --git a/memInfo/unit_test/CMakeLists.txt b/memInfo/unit_test/CMakeLists.txt index eacd131..a63a25c 100644 --- a/memInfo/unit_test/CMakeLists.txt +++ b/memInfo/unit_test/CMakeLists.txt @@ -3,9 +3,10 @@ include(FetchContent) FetchContent_Declare( googletest - URL https://github.com/google/googletest/archive/03597a01ee50ed33e9dfd640b249b4be3799d395.zip - DOWNLOAD_EXTRACT_TIMESTAMP true - ) + GIT_REPOSITORY https://github.com/google/googletest.git + GIT_TAG 6910c9d9165801d8827d628cb72eb7ea9dd538c5 + DOWNLOAD_EXTRACT_TIMESTAMP true) + FetchContent_MakeAvailable(googletest) enable_testing() diff --git a/memInfo/unit_test/TestMemInfo.cpp b/memInfo/unit_test/TestMemInfo.cpp index 6619c73..c33cc2b 100644 --- a/memInfo/unit_test/TestMemInfo.cpp +++ b/memInfo/unit_test/TestMemInfo.cpp @@ -1,81 +1,61 @@ -#include #include -#include + #include +#include + +#include +#include -template +template void testMemInfo() { - size_t step1_total = 0ul; - size_t step2_total = 0ul; - size_t step2_free = 0ul; - size_t step1_free = 0ul; + std::size_t step1_total = 0ul; + std::size_t step2_total = 0ul; + std::size_t step1_free = 0ul; + std::size_t step2_free = 0ul; - if constexpr (std::is_same::value) { - Kokkos::Experimental::MemGetInfo(&step1_free, &step1_total); - } else { - Kokkos::Experimental::MemGetInfo(&step1_free, &step1_total); - } - { - // Allocate 128 MiB of memory - Kokkos::View data("data test", 128, 1024*1024); - if constexpr (std::is_same::value) { - Kokkos::Experimental::MemGetInfo(&step2_free, &step2_total); - } else { - Kokkos::Experimental::MemGetInfo(&step2_free, &step2_total); - } - } + Kokkos::Experimental::MemGetInfo(&step1_free, &step1_total); + // Allocate 64 MiB of memory + Kokkos::View data("data test", 1024, 8192); + Kokkos::fence(); + Kokkos::Experimental::MemGetInfo(&step2_free, &step2_total); // Same total memory before and after aloccation EXPECT_EQ(step1_total, step2_total); // Check that free memory is less after allocation EXPECT_LT(step2_free, step1_free); } -TEST(MemInfo, HostSpace) { - Kokkos::initialize(); - testMemInfo(); - Kokkos::finalize(); -} -TEST(MemInfo, HostSpaceUninitialized) { - size_t free = 0; - size_t total = 0; - Kokkos::Experimental::MemGetInfo(&free, &total); - EXPECT_GT(free, 0); - EXPECT_GT(total, 0); -} +#define TEST_SPACE(Space) \ + TEST(MemInfo, Space) { \ + testMemInfo(); \ + } TEST(MemInfo, DefaultSpace) { - Kokkos::initialize(); testMemInfo<>(); - Kokkos::finalize(); -} -TEST(MemInfo, DefaultSpaceUninitialized) { - size_t free = 0; - size_t total = 0; - Kokkos::Experimental::MemGetInfo(&free, &total); - EXPECT_GT(free, 0); - EXPECT_GT(total, 0); } +TEST_SPACE(HostSpace) + #if defined(KOKKOS_ENABLE_CUDA) -TEST(MemInfo, CudaSpace) { - Kokkos::initialize(); - testMemInfo(); - testMemInfo(); - Kokkos::finalize(); -} + TEST_SPACE(CudaSpace) + TEST_SPACE(SharedSpace) #endif #if defined(KOKKOS_ENABLE_HIP) -TEST(MemInfo, HIPSpace) { - Kokkos::initialize(); - testMemInfo(); - testMemInfo(); - Kokkos::finalize(); -} + TEST_SPACE(HIPSpace) + TEST_SPACE(SharedSpace) +#endif + +#if defined(KOKKOS_ENABLE_SYCL) + TEST_SPACE(SYCLDeviceUSMSpace) + TEST_SPACE(SharedSpace) #endif int main(int argc, char *argv[]) { + Kokkos::initialize(argc, argv); ::testing::InitGoogleTest(&argc, argv); int result = RUN_ALL_TESTS(); + Kokkos::finalize(); return result; } + +#undef TEST_SPACE From ad519788de814a8ca261113d8b54a175a73506a6 Mon Sep 17 00:00:00 2001 From: Adrien Taberner Date: Sun, 16 Mar 2025 17:30:57 +0100 Subject: [PATCH 7/7] cgroup v1 for slurm Signed-off-by: Adrien Taberner --- memInfo/CMakeLists.txt | 8 +- memInfo/src/cexa_MemInfo.hpp | 77 +++--------- memInfo/src/unixMemInfo.hpp | 194 ++++++++++++++++++++++++++++++ memInfo/unit_test/TestMemInfo.cpp | 31 +++-- 4 files changed, 238 insertions(+), 72 deletions(-) create mode 100644 memInfo/src/unixMemInfo.hpp diff --git a/memInfo/CMakeLists.txt b/memInfo/CMakeLists.txt index eb20152..bb85bbe 100644 --- a/memInfo/CMakeLists.txt +++ b/memInfo/CMakeLists.txt @@ -4,11 +4,11 @@ project(cexa-experimental-meminfo LANGUAGES CXX) set(CMAKE_BUILD_TYPE "RelWithDebInfo") if (NOT ${Kokkos_DIR} STREQUAL "") - add_subdirectory(${Kokkos_DIR}) - include_directories(${Kokkos_INCLUDE_DIRS_RET}) + add_subdirectory(${Kokkos_DIR}) + include_directories(${Kokkos_INCLUDE_DIRS_RET}) else() - list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/../cmake/modules") - find_package(Kokkos REQUIRED) + list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/../cmake/modules") + find_package(Kokkos REQUIRED) endif() include(CTest) diff --git a/memInfo/src/cexa_MemInfo.hpp b/memInfo/src/cexa_MemInfo.hpp index f8b67ae..e6fdbe4 100644 --- a/memInfo/src/cexa_MemInfo.hpp +++ b/memInfo/src/cexa_MemInfo.hpp @@ -1,64 +1,21 @@ #ifndef KOKKOS_MEMINFO_HPP #define KOKKOS_MEMINFO_HPP -#include -#include -#include - #ifdef _WIN32 #include +#else +#include #endif +#include +#include +#include + #include namespace Kokkos { namespace Experimental { -namespace { - constexpr int OVERCOMMIT_DISABLED = 2; - constexpr char MEM_FREE_KEY[] = "MemFree:"; - constexpr char MEM_TOTAL_KEY[] = "MemTotal:"; - constexpr char COMMITTED_AS_KEY[] = "Committed_AS:"; - constexpr char COMMIT_LIMIT_KEY[] = "CommitLimit:"; - constexpr char MEMINFO_PATH[] = "/proc/meminfo"; - constexpr char OVERCOMMIT_PATH[] = "/proc/sys/vm/overcommit_memory"; -} - -// On some systems, overcommit is disabled, and the kernel does not allow -// memory allocation beyond the commit limit. This means that allocations -// that touch only a small amount of memory are still counted at their full size. -// man proc_sys_vm -bool is_overcommit_disabled() { - std::ifstream overcommit_file(OVERCOMMIT_PATH); - int overcommit_value = 0; - - if (overcommit_file.is_open()) { - overcommit_file >> overcommit_value; - } else { - return false; - } - return (overcommit_value == OVERCOMMIT_DISABLED); -} - -size_t get_meminfo_value(const char* key) { - std::ifstream meminfo(MEMINFO_PATH); - size_t value = 0; - std::string line; - - if (meminfo.is_open()) { - while (std::getline(meminfo, line)) { - if (line.find(key) != std::string::npos) { - std::istringstream iss(line); - iss.ignore(256, ':'); - iss >> value; - value = (iss.fail()) ? 0 : value * 1024; - break; - } - } - } - return value; -} - template void MemGetInfo(size_t* free, size_t* total) { using MemorySpace = typename Space::memory_space; @@ -66,26 +23,18 @@ void MemGetInfo(size_t* free, size_t* total) { } // Single node memory info +#ifdef _WIN32 template <> void MemGetInfo(size_t* free, size_t* total) { -#ifdef _WIN32 MEMORYSTATUSEX statex; statex.dwLength = sizeof(statex); if (GlobalMemoryStatusEx(&statex) != 0) { *free = statex.ullAvailPhys; *total = statex.ullTotalPhys; } -#else - static bool overcommit_disabled = is_overcommit_disabled(); - if (overcommit_disabled) { - *total = get_meminfo_value(COMMIT_LIMIT_KEY); - *free = *total - get_meminfo_value(COMMITTED_AS_KEY); - } else { - *total = get_meminfo_value(MEM_TOTAL_KEY); - *free = get_meminfo_value(MEM_FREE_KEY); - } -#endif + return; } +#endif #if defined(KOKKOS_ENABLE_CUDA) template <> @@ -96,6 +45,10 @@ template <> void MemGetInfo(size_t* free, size_t* total) { MemGetInfo(free, total); } +template <> +void MemGetInfo(size_t* free, size_t* total) { + MemGetInfo(free, total); +} #endif #if defined(KOKKOS_ENABLE_HIP) @@ -107,6 +60,10 @@ template <> void MemGetInfo(size_t* free, size_t* total) { MemGetInfo(free, total); } +template <> +void MemGetInfo(size_t* free, size_t* total) { + MemGetInfo(free, total); +} #endif #if defined(KOKKOS_ENABLE_SYCL) diff --git a/memInfo/src/unixMemInfo.hpp b/memInfo/src/unixMemInfo.hpp new file mode 100644 index 0000000..cb37f02 --- /dev/null +++ b/memInfo/src/unixMemInfo.hpp @@ -0,0 +1,194 @@ +#ifndef KOKKOS_UNIX_MEMINFO_HPP +#define KOKKOS_UNIX_MEMINFO_HPP + +#include + +#include +#include +#include + +#include + +namespace Kokkos { +namespace Experimental { + +namespace { + constexpr size_t NO_LIMIT = 1ull << 50; // No limit (like PAGE_COUNT_MAX) + constexpr int OVERCOMMIT_DISABLED = 2; + // Memory info keys + constexpr char MEM_FREE_KEY[] = "MemFree:"; + constexpr char MEM_TOTAL_KEY[] = "MemTotal:"; + constexpr char COMMITTED_AS_KEY[] = "Committed_AS:"; + constexpr char COMMIT_LIMIT_KEY[] = "CommitLimit:"; + // Cgroup v1 memory info + constexpr char CGROUP_PROCS[] = "cgroup.procs"; + constexpr char MEM_LIMIT_BYTES[] = "memory.limit_in_bytes"; + constexpr char MEM_USAGE_BYTES[] = "memory.usage_in_bytes"; + // Paths + constexpr char MEMINFO_PATH[] = "/proc/meminfo"; + constexpr char OVERCOMMIT_PATH[] = "/proc/sys/vm/overcommit_memory"; + constexpr char CGROUP_MEM_PATH[] = "/sys/fs/cgroup/memory"; +} + +template +void MemGetInfo(size_t* free, size_t* total); + +// On some systems, overcommit is disabled, and the kernel does not allow +// memory allocation beyond the commit limit. This means that allocations +// that touch only a small amount of memory are still counted at their full size. +// man proc_sys_vm +bool is_overcommit_disabled() { + std::ifstream overcommit_file(OVERCOMMIT_PATH); + int overcommit_value = 0; + + if (overcommit_file.is_open()) { + overcommit_file >> overcommit_value; + overcommit_value = (overcommit_file.fail()) ? 0 : overcommit_value; + return (overcommit_value == OVERCOMMIT_DISABLED); + } + return false; +} + +size_t get_meminfo_value(const char* key) { + std::ifstream meminfo(MEMINFO_PATH); + size_t value = 0; + std::string line; + + if (meminfo.is_open()) { + while (std::getline(meminfo, line)) { + if (line.find(key) != std::string::npos) { + std::istringstream iss(line); + iss.ignore(256, ':'); + iss >> value; + value = (iss.fail()) ? 0 : value * 1024; + break; + } + } + } + return value; +} + +// Extract a value from a cgroup file +size_t get_cgroup_value(const char* path) { + std::ifstream cgroup_file(path); + size_t value = 0; + + if (cgroup_file.is_open()) { + cgroup_file >> value; + value = (cgroup_file.fail()) ? 0 : value; + } + return value; +} + +// Check if a process is in the cgroup.procs file +bool is_pid_in_cgroup_procs(const char* cgroup_procs_path, const pid_t pid) { + std::ifstream cgroup_procs(cgroup_procs_path); + pid_t proc_id = 0; + + if (cgroup_procs.is_open()) { + while (cgroup_procs >> proc_id) { + if (proc_id == pid) { + return true; + } + } + } + return false; +} + +// Find out if memory controller is enabled (cgroup v1) +// Check in /proc//cgroup +bool is_cgroup_mem_control_enabled() { + const pid_t pid = getpid(); + std::ifstream cgroup_memory_limit("/proc/" + std::to_string(pid) + "/cgroup"); + std::string line; + + if (cgroup_memory_limit.is_open()) { + while (std::getline(cgroup_memory_limit, line)) { + if (line.find("memory") != std::string::npos) { + return true; + } + } + } + return false; +} + +// Find the cgroup memory path for the current process +// Verify if the process is in the cgroup.procs file +std::string find_cgroup_memory_path() { + const pid_t pid = getpid(); + std::ifstream cgroup_memory_limit("/proc/" + std::to_string(pid) + "/cgroup"); + std::string cgroup_path; + + if (cgroup_memory_limit.is_open()) { + std::string line; + while (std::getline(cgroup_memory_limit, line)) { + if (line.find(":memory:") != std::string::npos) { + const size_t pos = line.find_last_of(':'); + if (pos != std::string::npos) { + std::string cgroup_path = line.substr(pos + 1); + if (is_pid_in_cgroup_procs((CGROUP_MEM_PATH + cgroup_path + "/" + CGROUP_PROCS).c_str(), pid)) { + return CGROUP_MEM_PATH + cgroup_path; + } + const size_t last_slash = cgroup_path.find_last_of('/'); + if (last_slash != std::string::npos) { + std::string parent_path = cgroup_path.substr(0, last_slash); + if (is_pid_in_cgroup_procs((CGROUP_MEM_PATH + parent_path + "/" + CGROUP_PROCS).c_str(), pid)) { + return CGROUP_MEM_PATH + parent_path; + } + } + } + } + } + } + + // Fallback to the default path + return CGROUP_MEM_PATH + cgroup_path; +} + +// Single node memory info +template <> +void MemGetInfo(size_t* free, size_t* total) { + static bool overcommit_disabled = is_overcommit_disabled(); + bool cgroup_mem_enable = is_cgroup_mem_control_enabled(); + + // Cgroup memory info + if (cgroup_mem_enable) { + std::string cgroup_mem_path = find_cgroup_memory_path(); + const size_t mem_limit = get_cgroup_value((cgroup_mem_path + "/" + MEM_LIMIT_BYTES).c_str()); + if (mem_limit == 0 || mem_limit > NO_LIMIT) { + if (overcommit_disabled) { + *total = get_meminfo_value(COMMIT_LIMIT_KEY); + } else { + *total = get_meminfo_value(MEM_TOTAL_KEY); + } + } else { + *total = mem_limit; + } + const size_t mem_usage = get_cgroup_value((cgroup_mem_path + "/" + MEM_USAGE_BYTES).c_str()); + if (mem_usage == 0 || mem_usage > NO_LIMIT ) { + if (overcommit_disabled) { + *free = *total - get_meminfo_value(COMMITTED_AS_KEY); + } else { + *free = get_meminfo_value(MEM_FREE_KEY); + } + } else { + *free = *total - mem_usage; + } + return; + } + + // System memory info + if (overcommit_disabled) { + *total = get_meminfo_value(COMMIT_LIMIT_KEY); + *free = *total - get_meminfo_value(COMMITTED_AS_KEY); + } else { + *total = get_meminfo_value(MEM_TOTAL_KEY); + *free = get_meminfo_value(MEM_FREE_KEY); + } + return; +} + +} // namespace Experimental +} // namespace Kokkos + +#endif // KOKKOS_UNIX_MEMINFO_HPP \ No newline at end of file diff --git a/memInfo/unit_test/TestMemInfo.cpp b/memInfo/unit_test/TestMemInfo.cpp index c33cc2b..ee6c577 100644 --- a/memInfo/unit_test/TestMemInfo.cpp +++ b/memInfo/unit_test/TestMemInfo.cpp @@ -8,20 +8,33 @@ template void testMemInfo() { - std::size_t step1_total = 0ul; - std::size_t step2_total = 0ul; - std::size_t step1_free = 0ul; - std::size_t step2_free = 0ul; + using memory_space = typename std::conditional< + Kokkos::is_memory_space::value, + Space, + typename Space::memory_space + >::type; + + std::size_t step1_total = 0ull; + std::size_t step2_total = 0ull; + std::size_t step1_free = 0ull; + std::size_t step2_free = 0ull; + + Kokkos::Experimental::MemGetInfo(&step1_free, &step1_total); - Kokkos::Experimental::MemGetInfo(&step1_free, &step1_total); // Allocate 64 MiB of memory - Kokkos::View data("data test", 1024, 8192); + Kokkos::View data("data test", 1024*8192); Kokkos::fence(); - Kokkos::Experimental::MemGetInfo(&step2_free, &step2_total); + Kokkos::Experimental::MemGetInfo(&step2_free, &step2_total); + // Same total memory before and after aloccation EXPECT_EQ(step1_total, step2_total); // Check that free memory is less after allocation - EXPECT_LT(step2_free, step1_free); + const bool is_shared_space = std::is_same::value; + if (is_shared_space) { + EXPECT_LE(step2_free, step1_free); + } else { + EXPECT_LT(step2_free, step1_free); + } } #define TEST_SPACE(Space) \ @@ -38,11 +51,13 @@ TEST_SPACE(HostSpace) #if defined(KOKKOS_ENABLE_CUDA) TEST_SPACE(CudaSpace) TEST_SPACE(SharedSpace) + TEST_SPACE(SharedHostPinnedSpace) #endif #if defined(KOKKOS_ENABLE_HIP) TEST_SPACE(HIPSpace) TEST_SPACE(SharedSpace) + TEST_SPACE(SharedHostPinnedSpace) #endif #if defined(KOKKOS_ENABLE_SYCL)