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

Use finegrain allocator by default (#140)

Browse Source 
* Use FineGrained allocator for heap by default, consolidate all types of
allocators under saner cmake controls

Co-authored-by: Yiltan <ytemucin@amd.com>

* Uncached may not be only for debug

Need to include the rocshmem config otherwise produce an inconsistent
build with different allocators used in different files

* Undo this pr adding presumably useless hip_host_allocator_noncoherent

* Rename HEAP_IS_COHERENT/USE_COHERENT_HEAP to USE_HDP_FLUSH as the former
was misleading

* Remove unused __roc_inv()

---------

Co-authored-by: Yiltan <ytemucin@amd.com>
This commit is contained in:
Aurelien Bouteiller
2025-06-13 15:26:26 -04:00
committed by GitHub
parent bcc14b1a34
commit 41fd9e2d57
13 changed files with 67 additions and 66 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/assembly.hpp
+1 -16
View File
@@ -165,23 +165,8 @@ NOWARN(-Wdeprecated-volatile,
)
// clang-format on
__device__ __forceinline__ void __roc_inv() {
#if defined USE_COHERENT_HEAP
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
#endif
#if defined(__gfx90a__)
// asm volatile("buffer_wbinvl1;");
#endif
#if defined(__gfx942__)
// asm volatile("buffer_inv sc0 sc1;");
#endif
#endif
}
__device__ __forceinline__ void __roc_flush() {
#if defined USE_COHERENT_HEAP
#if not defined USE_HDP_FLUSH
#if defined(__gfx906__)
#endif
#if defined(__gfx908__)
src/hdp_policy.hpp
+4 -4
View File
@@ -182,16 +182,16 @@ class NoHdpPolicy {
/*
* Select which one of our HDP policies to use at compile time.
*/
#ifdef USE_COHERENT_HEAP
typedef NoHdpPolicy HdpPolicy;
#else
#if defined USE_HDP_FLUSH
// Only when we are using the IB conduit, we have to use a polling thread to
// flush the HDP cache on the GPU's behalf.
#ifdef USE_HOST_SIDE_HDP_FLUSH
#if defined USE_HDP_FLUSH_HOST_SIDE
typedef HdpHostSideFlushRocmPolicy HdpPolicy;
#else
typedef HdpDeviceSideFlushRocmPolicy HdpPolicy;
#endif
#else
typedef NoHdpPolicy HdpPolicy;
#endif
} // namespace rocshmem
src/host/host.cpp
+6 -6
View File
@@ -118,16 +118,16 @@ __host__ HostInterface::HostInterface(HdpPolicy* hdp_policy,
new HostContextWindowInfo(host_comm_world_, heap);
}
#if !defined(USE_COHERENT_HEAP) && !defined(USE_SINGLE_NODE)
#if defined(USE_HDP_FLUSH) && !defined(USE_SINGLE_NODE)
// The single node implementation needs a different path since
// the HDP flush pointers are allocated on the symmetric heap
// and we need to wait for other initialization to happen before
// calling `get_hdp_flush_ptr`.
create_hdp_window();
#endif // defined(USE_COHERENT_HEAP) && !defined(USE_SINGLE_NODE)
#endif // defined(USE_HDP_FLUSH) && !defined(USE_SINGLE_NODE)
}
#ifndef USE_COHERENT_HEAP
#if defined USE_HDP_FLUSH
__host__ void HostInterface::create_hdp_window() {
MPI_Win_create(hdp_policy_->get_hdp_flush_ptr(),
sizeof(unsigned int), /* size of window */
@@ -142,14 +142,14 @@ __host__ void HostInterface::create_hdp_window() {
*/
MPI_Win_lock_all(MPI_MODE_NOCHECK, hdp_win);
}
#endif // USE_COHERENT_HEAP
#endif // USE_HDP_FLUSH
__host__ HostInterface::~HostInterface() {
#ifndef USE_COHERENT_HEAP
#if defined USE_HDP_FLUSH
MPI_Win_unlock_all(hdp_win);
MPI_Win_free(&hdp_win);
#endif // USE_COHERENT_HEAP
#endif // USE_HDP_FLUSH
/* Detroy the pool of contexts */
for (int ctx_i = 0; ctx_i < max_num_ctxs_; ctx_i++) {
src/host/host.hpp
+8 -8
View File
@@ -248,16 +248,16 @@ class HostInterface {
template <typename T>
__host__ int test(T *ivars, int cmp, T val, WindowInfo* window_info);
#ifndef USE_COHERENT_HEAP
#if defined USE_HDP_FLUSH
__host__ void create_hdp_window();
#endif // USE_COHERENT_HEAP
#endif // USE_HDP_FLUSH
private:
/**************************************************************************
**************************** INTERNAL METHODS ****************************
*************************************************************************/
__host__ void flush_remote_hdps() {
#ifndef USE_COHERENT_HEAP
#if defined USE_HDP_FLUSH
unsigned flush_val{HdpPolicy::HDP_FLUSH_VAL};
for (size_t i{0}; i < num_pes_; i++) {
if (i == my_pe_) {
@@ -266,15 +266,15 @@ class HostInterface {
MPI_Put(&flush_val, 1, MPI_UNSIGNED, i, 0, 1, MPI_UNSIGNED, hdp_win);
}
MPI_Win_flush_all(hdp_win);
#endif // USE_COHERENT_HEAP
#endif // USE_HDP_FLUSH
}
__host__ void flush_remote_hdp(int pe) {
#ifndef USE_COHERENT_HEAP
#if defined USE_HDP_FLUSH
unsigned flush_val{HdpPolicy::HDP_FLUSH_VAL};
MPI_Put(&flush_val, 1, MPI_UNSIGNED, pe, 0, 1, MPI_UNSIGNED, hdp_win);
MPI_Win_flush(pe, hdp_win);
#endif // USE_COHERENT_HEAP
#endif // USE_HDP_FLUSH
}
__host__ void initiate_put(void* dest, const void* source, size_t nelems,
@@ -333,12 +333,12 @@ class HostInterface {
*/
int num_pes_{0};
#ifndef USE_COHERENT_HEAP
#if defined USE_HDP_FLUSH
/**
* @brief MPI window for hdp flushing
*/
MPI_Win hdp_win;
#endif // USE_COHERENT_HEAP
#endif // USE_HDP_FLUSH
/**
* @brief Max number of contexts for the application
src/memory/heap_type.hpp
+17 -8
View File
@@ -40,16 +40,25 @@
namespace rocshmem {
#if defined USE_MANAGED_HEAP
using HEAP_T = HeapMemory<HIPAllocatorManaged>;
#elif defined USE_COHERENT_HEAP
// Compilation error 'HEAP_T redefined' indicates that user had more than one
// USE_HEAP_* ON when configuring. Use ccmake to select only one.
#if defined USE_HEAP_DEVICE_COARSEGRAIN
using HEAP_T = HeapMemory<HIPAllocator>;
#elif defined USE_HOST_HEAP
using HEAP_T = HeapMemory<HostAllocator>;
#elif defined USE_HIP_HOST_HEAP
#endif
#if defined USE_HEAP_DEVICE_FINEGRAIN
using HEAP_T = HeapMemory<HIPAllocatorFinegrained>;
#endif
#if defined USE_HEAP_DEVICE_UNCACHED
using HEAP_T = HeapMemory<HIPAllocatorUncached>;
#endif
#if defined USE_HEAP_MANAGED
using HEAP_T = HeapMemory<HIPAllocatorManaged>;
#endif
#if defined USE_HEAP_HOST_HIP
using HEAP_T = HeapMemory<HIPHostAllocator>;
#else
using HEAP_T = HeapMemory<HIPDefaultFinegrainedAllocator>;
#endif
#if defined USE_HEAP_HOST
using HEAP_T = HeapMemory<HostAllocator>;
#endif
} // namespace rocshmem
src/memory/hip_allocator.hpp
+5 -1
View File
@@ -36,12 +36,15 @@
#include <cstdlib>
#include <limits>
#include "rocshmem_config.h" // NOLINT(build/include_subdir)
#include "memory_allocator.hpp"
// `hipDeviceMallocUncached` was introduced at ROCm 5.5
#if (HIP_VERSION_MAJOR > 5) || \
(HIP_VERSION_MAJOR == 5 && HIP_VERSION_MINOR >= 5)
#define HIP_SUPPORTS_MALLOC_UNCACHED
#elif defined USE_HEAP_DEVICE_UNCACHED
#error "USE_HEAP_DEVICE_UNCACHED unsupported in this HIP version"
#endif
namespace rocshmem {
@@ -57,13 +60,14 @@ class HIPAllocatorFinegrained : public MemoryAllocator {
hipDeviceMallocFinegrained) {}
};
#ifdef HIP_SUPPORTS_MALLOC_UNCACHED
#if defined HIP_SUPPORTS_MALLOC_UNCACHED
class HIPAllocatorUncached : public MemoryAllocator {
public:
HIPAllocatorUncached()
: MemoryAllocator(hipExtMallocWithFlags, hipFree,
hipDeviceMallocUncached) {}
};
// The default fine-grained coherence allocator is the uncached allocator
using HIPDefaultFinegrainedAllocator = HIPAllocatorUncached;
#else