diff --git a/.github/workflows/reconverse-cuda.yaml b/.github/workflows/reconverse-cuda.yaml index 83e80c983e..a0fd2b4f64 100644 --- a/.github/workflows/reconverse-cuda.yaml +++ b/.github/workflows/reconverse-cuda.yaml @@ -170,6 +170,13 @@ jobs: run: | make -j4 -C netlrts-linux-x86_64-cuda/tests/charm++/cuda/d2dtest \ GPU=cuda CUDATOOLKIT_HOME=/usr OPTS="-g" + - name: gpumigrate compiles + # Device-state migration. Running it needs a GPU, but compiling it + # gates the PUPMode::DEVICE surface that application pup routines see, + # and the nocopydevice entry method the migration payload travels on. + run: | + make -j4 -C netlrts-linux-x86_64-cuda/tests/charm++/cuda/gpumigrate \ + OPTS="-g" - name: legacy buildold path run: | diff --git a/.gitignore b/.gitignore index ae38f6acec..b851194f72 100644 --- a/.gitignore +++ b/.gitignore @@ -14,6 +14,9 @@ configure~ *.def.h *.ci.stamp +# Ignore crash backtrace dumps written beside a crashing binary +*.btr + # Ignore build artifacts config_opts.sh smart-build.log diff --git a/CMakeLists.txt b/CMakeLists.txt index 868a4787e0..1523343eef 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -755,6 +755,11 @@ if(BUILD_CUDA OR BUILD_HIP) message(WARNING "CUPTI library not found. GPU load balancing will not be available.") else() message(STATUS "Found CUPTI: ${CUPTI_LIBRARY}") + # Global, not per-target: the GPU load fields live in LDObjData, which + # ck-core, ck-ldb and hybridapi all see. A per-target definition would + # give them different layouts for the same struct. + add_compile_definitions(HAPI_CUPTI_LB) + set(HAPI_CUPTI_LB ON) endif() add_library(hybridapi ${hybridAPI-cxx-sources}) diff --git a/cmake/converse.cmake b/cmake/converse.cmake index 542fb86451..6419e14771 100644 --- a/cmake/converse.cmake +++ b/cmake/converse.cmake @@ -327,6 +327,9 @@ if(NOT RECONVERSE) # be shadowed. list(APPEND conv-core-h-to-install src/conv-core/converse.h) list(APPEND conv-core-h-to-install src/conv-core/charm-config.h) + # Classic seed-balancer header: ck.h includes it when CMK_USE_SHMEM is on. + # Reconverse builds get reconverse's own cldb.h from cmake/reconverse. + list(APPEND conv-core-h-to-install src/conv-ldb/cldb.h) endif() if(RECONVERSE) # Reconverse provides its own conv-rdma.h, which is the authority on the diff --git a/doc/charm++/manual.rst b/doc/charm++/manual.rst index 19130c4443..4ee92a9c53 100644 --- a/doc/charm++/manual.rst +++ b/doc/charm++/manual.rst @@ -2679,6 +2679,51 @@ infrastructure: options to point to the include and library directories used, respectively. (``+balancer ScotchLB``) +The following centralized balancer targets GPU applications, and is +available on CUDA builds only: + +- **GreedyRefineCentralGPULB**: Balances at two granularities at once. + Objects are assigned across *GPU groups* — the sets of PEs that share a + device — by their measured GPU load, and then across the PEs within a + group by their host load. This matches the usual arrangement in which + each process owns a GPU: which device an object runs on is decided by + its device work, and which PE drives it by its host work. + (``+balancer GreedyRefineCentralGPULB``) + + The GPU load it reads is measured by CUPTI, per object, and normalized + to seconds of whole-device occupancy, so it estimates what an object + would cost on a *different* device rather than only what it cost where + it ran. Measurement follows ``LBTurnInstrumentOn()`` / + ``LBTurnInstrumentOff()``, so an application that instruments a window + around its own ``AtSync`` pays the tracing cost only inside it. See + :numref:`lbFramework` for the instrumentation calls. + + Like the other greedy-refine balancers it accepts a *tolerance*: how far + above the maximum load plain greedy would produce it may go, in exchange + for migrating fewer objects. ``+LBGreedyRefineTolerance 1.1`` allows a + maximum load 10% higher than greedy's, and an object then stays where it + is for as long as its PE remains under that target. + + The default is ``1.1``. Passing ``0`` or less asks the balancer to search + for a tolerance instead of being told one: each PE builds a candidate + assignment from a different parameter pair and the best is chosen, which + costs a reduction per balancing step and is bounded by the PE count. + ``+LBDebug 1`` reports the migration count and the achieved maximum load + against greedy's, separately for the GPU and host dimensions, which is + what says whether the tolerance is buying anything. + +.. note:: + + Migrating a GPU application requires the runtime to be built with + ``CMK_GLOBAL_LOCATION_UPDATE`` (see `Global Location Update`_). A device + zerocopy send is addressed to the PE the sender believes hosts the + target, so without it a send issued around a migration arrives at a PE + that no longer hosts the element, and the receive path aborts rather + than silently reading the wrong buffer (``ckrdmadevice.C``). Note the + restriction that comes with the option: migrations must happen at load + balancing steps, so a GPU application must migrate through ``AtSync()`` + rather than ``migrateMe()``. + In distributed approaches, the strategy executes across multiple PEs, providing scalable computational and communication performance. diff --git a/src/arch/common/conv-mach-cuda.sh b/src/arch/common/conv-mach-cuda.sh index 5cfd46cece..d51cf1bb4b 100644 --- a/src/arch/common/conv-mach-cuda.sh +++ b/src/arch/common/conv-mach-cuda.sh @@ -1,4 +1,9 @@ BUILD_CUDA=1 -CMK_INCDIR="-I$CUDA_DIR/include $CMK_INCDIR " -CMK_LIBDIR="-L$CUDA_DIR/lib64 $CMK_LIBDIR " -CMK_LIBS="-lhybridapi -lcudart -lrt $CMK_LIBS " +CMK_INCDIR="-I$CUDA_DIR/include -I$CUDA_DIR/extras/CUPTI/include $CMK_INCDIR " +CMK_LIBDIR="-L$CUDA_DIR/lib64 -L$CUDA_DIR/extras/CUPTI/lib64 $CMK_LIBDIR " +CMK_LIBS="-lhybridapi -lcudart -lcupti -lrt $CMK_LIBS " +# HAPI_CUPTI_LB turns on per-object GPU load measurement. It has to be a global +# define rather than one confined to hybridapi: the fields it guards live in +# LDObjData, which ck-core and ck-ldb see too, and a partial definition would +# give them different layouts for the same struct. +CMK_DEFS="$CMK_DEFS -DHAPI_CUPTI_LB " diff --git a/src/arch/cuda/hybridAPI/devicemanager.h b/src/arch/cuda/hybridAPI/devicemanager.h index 2351048aa7..45a8cae49c 100644 --- a/src/arch/cuda/hybridAPI/devicemanager.h +++ b/src/arch/cuda/hybridAPI/devicemanager.h @@ -19,8 +19,23 @@ struct DeviceManager { // Buddy allocator for communication buffer buddy::allocator* comm_buffer; + // Device properties needed to estimate how many SMs a kernel occupies while + // it runs, which is what turns a raw kernel timeline into a load. Filled + // lazily by hapiPopulateDeviceProps, because device_managers is not + // populated yet when a DeviceManager is constructed. + int multi_processor_count; + int max_threads_per_sm; + int max_blocks_per_sm; + int max_registers_per_sm; + int max_shared_mem_per_sm; + int warp_size; + bool props_initialized; + DeviceManager(int local_index_, int global_index_) : - local_index(local_index_), global_index(global_index_), comm_buffer(nullptr) { + local_index(local_index_), global_index(global_index_), comm_buffer(nullptr), + multi_processor_count(0), max_threads_per_sm(0), max_blocks_per_sm(0), + max_registers_per_sm(0), max_shared_mem_per_sm(0), warp_size(0), + props_initialized(false) { #if CMK_SMP lock = CmiCreateLock(); #endif diff --git a/src/arch/cuda/hybridAPI/gpumanager.h b/src/arch/cuda/hybridAPI/gpumanager.h index a590da5243..364cf834f8 100644 --- a/src/arch/cuda/hybridAPI/gpumanager.h +++ b/src/arch/cuda/hybridAPI/gpumanager.h @@ -13,6 +13,12 @@ #include #include +#include +#include + +#if CMK_CUDA && CMK_LBDB_ON +#include "lbdb.h" // for LBKernelRecord, LDObjKey and GpuObjectTokenTable +#endif // Local rank of the logical node (process) that the given PE belongs to, // within its physical node: the logical node index modulo the number of @@ -55,7 +61,7 @@ struct hapi_ipc_event_shared { pthread_mutex_t lock; }; -#ifdef HAPI_CUPTI_LB +#if CMK_CUDA && CMK_LBDB_ON struct CuptiBufferItem { uint8_t* buffer; size_t validSize; @@ -183,14 +189,89 @@ struct GPUManager { std::vector hapi_ipc_device_infos; //CUPTI load balancing -#ifdef HAPI_CUPTI_LB - std::unordered_map cupti_correlation_db_;//correlationID -> ObjectID - - std::unordered_map cupti_obj_gpu_times_;//objectID -> accumulated GPU time in ns - +#if CMK_CUDA && CMK_LBDB_ON + // Runtime correlation ID -> process-local full-object token. + std::unordered_map cupti_object_correlation_db_; + + GpuObjectTokenTable cupti_object_tokens_; + std::mutex cupti_object_token_lock_; + + // Full LB object identity -> attributed kernel records. + std::unordered_map, LDObjKeyHash> + cupti_obj_kernel_records_; + + // Kernels that could not be attributed to any object (launched outside a + // migratable entry method, or with no correlation record). They occupy SMs + // and so take part in the sweep-line as contention, but receive no load. + // Kept separate rather than under a sentinel object ID, because 0 is a + // perfectly valid chare element ID. + std::vector cupti_unattributed_kernels_; + + // Full object identity -> SM-utilization-normalized GPU load in seconds. + std::unordered_map cupti_obj_norm_load_; + + // Previous round's count of kernels whose correlations had not been parsed + // yet, used to size the parked vector. + uint32_t cupti_pending_hint_ = 0; + + // Written by CUPTI's buffer-completed callback, which may run on a + // CUPTI-owned thread. That thread exists in non-SMP builds too, where the + // Converse locks compile out, so this needs a real mutex rather than a + // CmiNodeLock. std::queue cupti_buffer_queue_; + std::mutex cupti_queue_lock_; bool cupti_initialized_; + // Whether activity tracing is currently running. Separate from + // cupti_initialized_: the buffer callbacks are registered once, but tracing + // itself is switched on and off as the application asks for it. + // + // Atomic because the entry-method hooks read it on every invocation from + // every PE thread while another thread may be switching tracing on or off. + std::atomic cupti_tracing_active_{false}; + // Serializes hapiCuptiStartTracing/hapiCuptiStopTracing. This state lives in + // the node-wide GPUManager, but the switch is reached per-PE through + // LBDatabase::TurnStatsOn/Off, so every PE thread calls in. Without this, + // several threads enable or disable the same CUPTI activity kinds and flush + // concurrently, which corrupts CUPTI's internal buffer bookkeeping and shows + // up later as heap corruption in an unrelated allocation. Must NOT be the + // same mutex as cupti_queue_lock_: the flush in the stop path invokes the + // buffer-completed callback, which takes that one. + std::mutex cupti_tracing_lock_; + // PEs of this process whose instrumentation is currently on. Tracing runs + // while this is non-zero; see hapiCuptiStartTracing. + int cupti_tracing_users_ = 0; + // Bumped every time CUPTI is detached. Detaching clears CUPTI's + // external-correlation stack for every PE, but the counters that keep the + // entry-method push/pop hooks paired are per-PE, and only the PE that ran the + // detach could reset its own. Each PE compares this against the generation it + // last saw and zeroes its counter when they differ. + uint64_t cupti_generation_ = 0; + // Serializes turning this round's raw CUPTI records into cupti_obj_norm_load_, + // and makes that work happen exactly once per LB round no matter how many PE + // threads ask for it. Doing it with "rank 0 works between two CmiNodeBarrier + // calls" does not hold: both barriers sit behind #if CMK_SMP, which is 0 in + // the multicore build even though a process really does run many PE threads, + // so the barriers vanish and the other ranks read cupti_obj_norm_load_ while + // rank 0 is rebuilding it. A lock also cannot deadlock the way a spin barrier + // can: a PE waiting here waits on a PE that is running, not on one that has + // yet to arrive. + std::mutex cupti_prepare_lock_; + // Set once this round's loads are built; cleared by hapiClearCuptiData. + bool cupti_loads_ready_ = false; + // Epoch the built loads correspond to, so a second caller in the same round + // reads what the first built instead of rebuilding it. + uint64_t cupti_loads_epoch_ = 0; + // Arrival gate for the per-round load build; see hapiCuptiArrive. A load + // balancer's per-PE barrier fires when THAT PE's objects are at AtSync, and + // the first PE to fire would otherwise flush, drain and clear the + // process-wide CUPTI records on the spot. Any PE whose objects were still + // finishing kernels at that instant has them dropped from the round -- + // consistently the last PE of the process, whose objects then read as zero + // GPU load every step. + std::mutex cupti_arrive_lock_; + uint64_t cupti_arrive_epoch_ = 0; + int cupti_arrive_count_ = 0; #endif void init() { @@ -234,6 +315,15 @@ struct GPUManager { hapi_ipc_event_pool_size_pe = -1; hapi_ipc_event_pool_size_total = -1; +#if CMK_CUDA && CMK_LBDB_ON + // CUPTI load balancing + cupti_initialized_ = false; + cupti_tracing_active_.store(false, std::memory_order_relaxed); + cupti_generation_ = 0; + cupti_loads_ready_ = false; + cupti_loads_epoch_ = 0; +#endif + // Allocate host/device buffers array (both user and system-addressed) host_buffers_ = new void*[NUM_BUFFERS*2]; device_buffers_ = new void*[NUM_BUFFERS*2]; diff --git a/src/arch/cuda/hybridAPI/hapi.h b/src/arch/cuda/hybridAPI/hapi.h index 1e3c6017aa..e142f4aa88 100644 --- a/src/arch/cuda/hybridAPI/hapi.h +++ b/src/arch/cuda/hybridAPI/hapi.h @@ -2,11 +2,13 @@ #define __HAPI_H_ #include "hapi_portable.h" -/* HAPI_CUPTI_LB: per-object GPU time attribution (CUPTI activity tracing, - * event-based kernel timing, launch wrappers) feeding GPU-aware load - * balancing. Dormant -- nothing defines it -- until the GPU-LB series - * (plan item 11) enables it together with its ck-ldb counterparts - * (setObjGPUTime and friends, LBHasBalancersRegistered). */ +/* HAPI_CUPTI_LB: per-object GPU time attribution via CUPTI activity tracing, + * feeding GPU-aware load balancing. Defined for a CUDA build in which CUPTI + * was found; without it the entry points below are no-ops and every object + * reads zero GPU load, so a GPU-aware balancer falls back to the host + * dimension. The data structures the attribution produces are declared under + * CMK_CUDA rather than under this macro, so that LDObjData has one layout + * across every translation unit either way. */ /* See hapi_functions.h for the majority of function declarations provided * by the Hybrid API. */ @@ -263,29 +265,50 @@ static inline hapiError_t hapiFreeHost_Pool(void* ptr, bool pool) { return hapiFreeHost(ptr, pool); } -void hapiRecordTime(hapiStream_t stream, hapiEvent_t start); -#ifdef HAPI_CUPTI_LB +#if CMK_CUDA && CMK_LBDB_ON void hapiCuptiInit(); void hapiCuptiFinalize(); + +// Stamp/unstamp the running migratable object onto every kernel launched +// inside the scope. Called around every entry method, so both are no-ops +// unless tracing is running. uint64_t hapiCuptiPushObjCorrelation(); void hapiCuptiPopObjCorrelation(); + +// Epoch a load balancer passes to hapiPrepareCuptiLoads: larger than any +// epoch an application-level sampler will use, so an LB round always rebuilds +// rather than reading a sampler's older loads. +#define HAPI_CUPTI_EPOCH_LB_ROUND UINT64_MAX + +// Flush, parse and normalize the CUPTI records accumulated since the last +// hapiClearCuptiData, once per epoch however many PE threads call it. The +// result is GPUManager::cupti_obj_norm_load_, which every PE of the process +// then reads; see the comment on GPUManager::cupti_prepare_lock_. +void hapiPrepareCuptiLoads(uint64_t epoch = HAPI_CUPTI_EPOCH_LB_ROUND); void hapiProcessCuptiBuffers(); +void hapiNormalizeCuptiLoads(); void hapiClearCuptiData(); -#endif -#ifdef HAPI_CUPTI_LB -#define HAPI_LAUNCH_KERNEL_WRAPPER(call, stream)\ - hapiEvent_t start;\ - hapiEventCreate(&start);\ - hapiEventRecord(start, stream);\ - call;\ - hapiRecordTime(stream, start); -#else -#define HAPI_LAUNCH_KERNEL_WRAPPER(call, stream)\ - call; +// Arrival gate in front of hapiPrepareCuptiLoads for a load-balancing round. +// Returns true to exactly one caller per epoch, once `expected` callers have +// arrived: the records are process-wide, so the drain has to wait for the last +// PE of the process rather than run on the first. +bool hapiCuptiArrive(uint64_t epoch, int expected); + +// Start/stop CUPTI activity tracing. Tracing is the dominant cost of GPU load +// measurement, so an application that instruments a window rather than the +// whole run pays for it only inside that window. Reached per-PE through +// LBDatabase::TurnStatsOn/Off; the process traces while any of its PEs wants +// instrumentation. +void hapiCuptiStartTracing(); +void hapiCuptiStopTracing(); +bool hapiCuptiTracingActive(); #endif -#ifdef HAPI_CUPTI_LB +// Attribute one kernel launch to the running object. The runtime already +// brackets every entry method this way (see CkCallstackPush/Pop), so an +// application needs this only for a launch it makes outside one. +#if CMK_CUDA && CMK_LBDB_ON #define CUPTI_LAUNCH_WRAPPER(call)\ hapiCuptiPushObjCorrelation();\ call;\ diff --git a/src/arch/cuda/hybridAPI/hapi_functions.h b/src/arch/cuda/hybridAPI/hapi_functions.h index 6f4bd5a0b6..0e0f04da7e 100644 --- a/src/arch/cuda/hybridAPI/hapi_functions.h +++ b/src/arch/cuda/hybridAPI/hapi_functions.h @@ -49,11 +49,6 @@ AMPI_CUSTOM_FUNC(void, hapiAddCallback, hapiStream_t, void*, void*) // AMPI_CUSTOM_FUNC(cudaError_t, hapiMemcpyAsync, void*, const void*, size_t, enum cudaMemcpyKind, cudaStream_t) // AMPI_CUSTOM_FUNC(cudaError_t, hapiMemcpy2DAsync, void*, size_t, const void*, size_t, size_t, size_t, enum cudaMemcpyKind, cudaStream_t) -// Kernel launch wrapper -#ifdef HAPI_CUPTI_LB /* pairs with HAPI_LAUNCH_KERNEL_WRAPPER; lands with plan item 11 */ -AMPI_CUSTOM_FUNC(hapiError_t, hapiLaunchKernel, const void*, dim3, dim3, void**, size_t, hapiStream_t) -#endif - // Explicit memory allocations using pinned memory pool. AMPI_CUSTOM_FUNC(hapiError_t, hapiPoolMalloc, void**, size_t) AMPI_CUSTOM_FUNC(hapiError_t, hapiPoolFree, void*) @@ -63,6 +58,9 @@ AMPI_CUSTOM_FUNC(void, hapiErrorDie, hapiError_t, const char*, const char*, int) // Returns the GPU device index this PE is mapped to (set during hapiMapping). AMPI_CUSTOM_FUNC(uint64_t, hapiMyDevice, void) +// SM count of that device, or 0 if it has not been queried yet. A GPU-aware +// balancer needs it to compare devices of different sizes. +AMPI_CUSTOM_FUNC(int, hapiMyDeviceTotalSMs, void) #ifdef HAPI_INSTRUMENT_WRS AMPI_CUSTOM_FUNC(void, hapiInitInstrument, int n_chares, char n_types) diff --git a/src/arch/cuda/hybridAPI/hapi_impl.cpp b/src/arch/cuda/hybridAPI/hapi_impl.cpp index f4c5c35eec..4110f82b45 100644 --- a/src/arch/cuda/hybridAPI/hapi_impl.cpp +++ b/src/arch/cuda/hybridAPI/hapi_impl.cpp @@ -25,13 +25,43 @@ #include "hapi_nvtx.h" #endif +#if CMK_CUDA && CMK_LBDB_ON +// lbdb.h, not LBManager.h: this file needs the LB *types* (LDObjKey, +// GpuObjectTokenTable, LBKernelRecord), which are header-only, and must not +// acquire a link dependency on ck-ldb globals such as _lb_args. See the note on +// cuptiDebugLevel below. +#include "lbdb.h" +#include "cklocation.h" +#include +#include + +// Defined in ck.C. Forward-declared rather than reached through ck.h, which +// this file does not otherwise need. Gives the entry-method correlation hook +// the running element's full LB identity. +CkLocRec* CkActiveLocRec(void); + +// Verbosity for the CUPTI accounting below, from the environment rather than +// from _lb_args.debug(). +// +// This object is only ever pulled into a link that also has libck -- hapi is +// reached from ck-core/init.C -- but nothing enforces that, and when it is +// violated the failure is a wall of undefined references in unrelated +// pure-Converse programs rather than anything pointing here. Reading a global +// that lives in ck-ldb would add one more way to trip it, for two diagnostic +// printfs. The other diagnostic in this file (CHARM_LB_CUPTI_TIME) is already +// environment-driven, so this matches it. +static int cuptiDebugLevel() +{ + static const int level = []() { + const char* s = getenv("CHARM_LB_CUPTI_DEBUG"); + return s != nullptr ? atoi(s) : 0; + }(); + return level; +} + #ifdef HAPI_CUPTI_LB -#if CMK_CUDA #include -#endif -#include "LBManager.h" -#if CMK_CUDA #define CUPTI_SAFE_CALL(call) \ do { \ CUptiResult _status = call; \ @@ -72,11 +102,9 @@ typedef struct hapiEvent { CkCallback cb; void* cb_msg; hapiWorkRequest* wr; // if this is not NULL, buffers and request itself are deallocated - CkMigratable* obj; // pointer to the object whose load we want to set - hapiEvent_t start_ev; // event to record the start time - hapiEvent(hapiEvent_t event_, const CkCallback& cb_, void* cb_msg_, hapiWorkRequest* wr_ = NULL, CkMigratable* obj_ = NULL, hapiEvent_t start_ev_ = NULL) - : event(event_), cb(cb_), cb_msg(cb_msg_), wr(wr_), obj(obj_), start_ev(start_ev_) {} + hapiEvent(hapiEvent_t event_, const CkCallback& cb_, void* cb_msg_, hapiWorkRequest* wr_ = NULL) + : event(event_), cb(cb_), cb_msg(cb_msg_), wr(wr_) {} } hapiEvent; CpvDeclare(std::queue, hapi_event_queue); @@ -141,52 +169,184 @@ static void shmCleanup(); static void ipcHandleCreate(); static void ipcHandleOpen(); +#if CMK_CUDA && CMK_LBDB_ON + +// Sentinel external-correlation ID meaning "no owning migratable object". +// Must not collide with a real chare ID -- 0 is a perfectly valid one, which +// is why this is UINT64_MAX rather than the obvious choice. +static constexpr uint64_t HAPI_CUPTI_NO_OBJECT = + GpuObjectTokenTable::noObjectToken(); + #ifdef HAPI_CUPTI_LB -#if CMK_CUDA static void CUPTIAPI cuptiBufferRequested(uint8_t **buffer, size_t *size, size_t *maxNumRecords) { - *size = 5*1024 * 1024; // 5MB per buffer - *buffer = (uint8_t *)malloc(*size); - *maxNumRecords = 0; + // CUPTI writes activity records straight into this buffer and has no way to + // tell us the allocation failed: hand it NULL and it writes through a null + // pointer, and the NULL comes back through cuptiBufferCompleted to be parsed + // later, so the fault surfaces inside cuptiActivityGetNextRecord with nothing + // left to say where it came from. Step down to a smaller buffer before giving + // up, and if even that fails, say so here. + static const size_t sizes[] = {5*1024*1024, 1024*1024, 256*1024}; + for (size_t s : sizes) { + *buffer = (uint8_t *)malloc(s); + if (*buffer != NULL) { + *size = s; + *maxNumRecords = 0; + return; + } + } + CmiAbort("HAPI: could not allocate a CUPTI activity buffer (tried down to " + "%zu bytes). GPU load instrumentation cannot continue.", sizes[2]); } -//TODO: handle SMP mode static void CUPTIAPI cuptiBufferCompleted(CUcontext ctx, uint32_t streamId, uint8_t *buffer, size_t size, size_t validSize) { GPUManager& gm = CsvAccess(gpu_manager); + std::lock_guard lk(gm.cupti_queue_lock_); gm.cupti_buffer_queue_.push({buffer, validSize}); } + +// Populate DeviceManager with the device attributes needed to compute per-kernel +// SM usage from CUPTI records. Queried once per local device, lazily, because +// device_managers is not populated when the GPUManager is constructed. +static void hapiPopulateDeviceProps(GPUManager& gm) { + for (DeviceManager& dm : gm.device_managers) { + if (dm.props_initialized) continue; + int dev = dm.global_index; + cudaDeviceProp props; + hapiCheck(cudaGetDeviceProperties(&props, dev)); + + dm.multi_processor_count = props.multiProcessorCount; + dm.max_threads_per_sm = props.maxThreadsPerMultiProcessor; +#ifdef cudaDevAttrMaxBlocksPerMultiprocessor + hapiCheck(cudaDeviceGetAttribute(&dm.max_blocks_per_sm, + cudaDevAttrMaxBlocksPerMultiprocessor, dev)); +#else + dm.max_blocks_per_sm = 0; #endif + dm.max_registers_per_sm = props.regsPerMultiprocessor; + dm.max_shared_mem_per_sm = static_cast(props.sharedMemPerMultiprocessor); + dm.warp_size = props.warpSize; + dm.props_initialized = true; + } +} -// Initialize CUPTI activity tracing — called once per process -void hapiCuptiInit() { -#if CMK_CUDA - CmiPrintf("HAPI: Initializing CUPTI...\n"); - hapiDeviceSynchronize(); +// Kept for callers that want tracing up before the balancer asks for it; +// hapiCuptiStartTracing attaches on its own, so this is not needed at startup. +void hapiCuptiInit() { hapiCuptiStartTracing(); } + +// Attaching CUPTI to the process is NOT free even when no activity kind is +// enabled -- measured at ~1.3 ms per step on a 4-PE run, which is most of the +// cost that remains once tracing itself is windowed. So attach here and stop in +// hapiCuptiStopTracing, rather than staying enabled for the whole run. +// Enabling an activity kind is separately what makes records flow. +// +// Per-PE view of the tracing switch. LBDatabase::TurnStatsOn/Off is a PE-local +// switch, but CUPTI tracing is process-wide: a PE that switched its own +// instrumentation off would otherwise switch tracing off for every PE in the +// process, and whichever PE was still finishing its iteration would lose every +// kernel it launched after that instant -- one whole PE per process reading +// zero GPU load at every step. So the process traces while ANY PE wants +// instrumentation: tracing starts with the first PE to switch on and stops with +// the last to switch off, counted per PE so repeated switches do not skew the +// count. +static thread_local bool cupti_pe_tracing = false; + +void hapiCuptiStartTracing() { GPUManager& gm = CsvAccess(gpu_manager); - if (gm.cupti_initialized_) return; + // Every PE thread reaches this through its own LBDatabase::TurnStatsOn, so + // the check and the enable must be one atomic step -- otherwise several + // threads each enable the same activity kinds. + std::lock_guard lk(gm.cupti_tracing_lock_); + if (!cupti_pe_tracing) { + cupti_pe_tracing = true; + gm.cupti_tracing_users_++; + } + if (gm.cupti_tracing_active_.load(std::memory_order_relaxed)) return; + + if (!gm.cupti_initialized_) { + cudaDeviceSynchronize(); + CUPTI_SAFE_CALL( + cuptiActivityRegisterCallbacks(cuptiBufferRequested, cuptiBufferCompleted)); + gm.cupti_initialized_ = true; + } - CUPTI_SAFE_CALL(cuptiActivityRegisterCallbacks(cuptiBufferRequested, cuptiBufferCompleted)); + // RUNTIME must stay enabled alongside the kernel records even though nothing + // consumes its records directly: EXTERNAL_CORRELATION records are only + // emitted for correlation IDs generated by runtime-API tracking, so without + // it every kernel arrives unattributed and the balancer sees zero GPU load. CUPTI_SAFE_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL)); CUPTI_SAFE_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_RUNTIME)); CUPTI_SAFE_CALL(cuptiActivityEnable(CUPTI_ACTIVITY_KIND_EXTERNAL_CORRELATION)); - gm.cupti_initialized_ = true; -#endif + gm.cupti_tracing_active_.store(true, std::memory_order_relaxed); +} + +void hapiCuptiStopTracing() { + GPUManager& gm = CsvAccess(gpu_manager); + std::lock_guard lk(gm.cupti_tracing_lock_); + if (cupti_pe_tracing) { + cupti_pe_tracing = false; + if (gm.cupti_tracing_users_ > 0) gm.cupti_tracing_users_--; + } + // Other PEs of this process still have their instrumentation on: their + // kernels are still being launched and must keep being recorded. + if (gm.cupti_tracing_users_ > 0) return; + if (!gm.cupti_initialized_ || + !gm.cupti_tracing_active_.load(std::memory_order_relaxed)) + return; + + CUPTI_SAFE_CALL(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL)); + CUPTI_SAFE_CALL(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_RUNTIME)); + CUPTI_SAFE_CALL(cuptiActivityDisable(CUPTI_ACTIVITY_KIND_EXTERNAL_CORRELATION)); + + // Clear the flag before flushing so the buffers handed back by the flush are + // the last ones, and no further correlation pushes race with them. Flush on + // the way out so records buffered before the stop are not lost when the + // application switches instrumentation off around its own AtSync. The flush + // drives the buffer-completed callback, which takes cupti_queue_lock_ -- a + // different mutex from the one held here, so this cannot deadlock. + gm.cupti_tracing_active_.store(false, std::memory_order_relaxed); + CUPTI_SAFE_CALL(cuptiActivityFlushAll(CUPTI_ACTIVITY_FLAG_FLUSH_FORCED)); + + // Deliberately NOT detaching with cuptiFinalize(). Staying attached costs + // ~1.3 ms per step even with every kind disabled, and detaching measured + // cheaper (~2.3 ms per window against 1.3 ms per step) -- but it cannot be + // done safely from here. The entry-method hooks check cupti_tracing_active_ + // and then call into CUPTI without holding this lock, so a thread that has + // already passed that check can be inside cuptiActivityPushExternalCorrelationId + // while this one finalizes underneath it, which corrupts CUPTI's allocator + // and surfaces later as heap corruption in unrelated allocations. Reclaiming + // that 1.3 ms needs the hooks made safe against detach first. +} + +bool hapiCuptiTracingActive() { + return CsvAccess(gpu_manager).cupti_tracing_active_.load( + std::memory_order_relaxed); } void hapiCuptiFinalize() { - CmiPrintf("HAPI: Finalizing CUPTI...\n"); - hapiDeviceSynchronize(); // Ensure all activity records are flushed GPUManager& gm = CsvAccess(gpu_manager); - if(gm.cupti_initialized_== false) return; + if (!gm.cupti_initialized_) return; + cudaDeviceSynchronize(); // Ensure all activity records are flushed gm.cupti_initialized_ = false; -#if CMK_CUDA + gm.cupti_tracing_active_.store(false, std::memory_order_relaxed); + ++gm.cupti_generation_; + CUPTI_SAFE_CALL(cuptiFinalize()); -#endif } -#endif + +#else /* !HAPI_CUPTI_LB: no CUPTI in this build */ + +void hapiCuptiInit() {} +void hapiCuptiFinalize() {} +void hapiCuptiStartTracing() {} +void hapiCuptiStopTracing() {} +bool hapiCuptiTracingActive() { return false; } + +#endif /* HAPI_CUPTI_LB */ +#endif /* CMK_CUDA && CMK_LBDB_ON */ #ifndef HAPI_CUDA_CALLBACK #if CSD_NO_SCHEDLOOP @@ -252,104 +412,491 @@ static void hapiInitCsv(char** argv) { // Create and initialize GPU Manager object CsvInitialize(GPUManager, gpu_manager); CsvAccess(gpu_manager).init(); - #ifdef HAPI_CUPTI_LB - if (LBHasBalancersRegistered() && _lb_args.statsOn()) - hapiCuptiInit(); - #endif + // CUPTI is attached lazily by hapiCuptiStartTracing, which the balancer + // reaches through LBDatabase::TurnStatsOn. Attaching here instead would pay + // the attach cost for the whole run even when the application only wants + // instrumentation around its load-balancing steps. } +#if CMK_CUDA && CMK_LBDB_ON #ifdef HAPI_CUPTI_LB +// Find the DeviceManager that matches this kernel's device id. +static DeviceManager* findDeviceManager(GPUManager& gm, uint32_t device_id) { + for (DeviceManager& dm : gm.device_managers) { + if ((uint32_t)dm.global_index == device_id) return &dm; + } + return nullptr; +} + +// Compute the number of SMs this kernel occupies while running. +// Uses the CUDA occupancy model: theoretical max_active_blocks_per_sm is +// limited by (a) max blocks per SM, (b) warp count, (c) register pressure, +// (d) shared memory. Then: +// sms_used = min(num_sms, ceil(total_blocks / max_active_blocks_per_sm)) +static int computeKernelSMs(const DeviceManager& dm, + const CUpti_ActivityKernel4* k) { + if (!dm.props_initialized || dm.multi_processor_count <= 0) return 1; + + uint64_t threads_per_block = + (uint64_t)k->blockX * (uint64_t)k->blockY * (uint64_t)k->blockZ; + uint64_t total_blocks = + (uint64_t)k->gridX * (uint64_t)k->gridY * (uint64_t)k->gridZ; + if (threads_per_block == 0 || total_blocks == 0) return 1; + + // Warp-count limit: maxThreadsPerSM / threadsPerBlock (rounded down). + int limit_warps = + dm.max_threads_per_sm > 0 + ? (int)(dm.max_threads_per_sm / threads_per_block) + : INT_MAX; + if (limit_warps <= 0) limit_warps = 1; + + // Block-count limit (CUDA 11+; 0 means not available -> use a large value). + int limit_blocks = dm.max_blocks_per_sm > 0 ? dm.max_blocks_per_sm : INT_MAX; + + // Register-pressure limit. + uint64_t regs_per_block = (uint64_t)k->registersPerThread * threads_per_block; + int limit_regs = INT_MAX; + if (regs_per_block > 0 && dm.max_registers_per_sm > 0) { + uint64_t r = (uint64_t)dm.max_registers_per_sm / regs_per_block; + limit_regs = r > INT_MAX ? INT_MAX : (int)r; + if (limit_regs <= 0) limit_regs = 1; + } + + // Shared-memory limit. + uint64_t smem_per_block = + (uint64_t)k->staticSharedMemory + (uint64_t)k->dynamicSharedMemory; + int limit_smem = INT_MAX; + if (smem_per_block > 0 && dm.max_shared_mem_per_sm > 0) { + uint64_t s = (uint64_t)dm.max_shared_mem_per_sm / smem_per_block; + limit_smem = s > INT_MAX ? INT_MAX : (int)s; + if (limit_smem <= 0) limit_smem = 1; + } + + int max_active_blocks_per_sm = + std::min(std::min(limit_blocks, limit_warps), + std::min(limit_regs, limit_smem)); + if (max_active_blocks_per_sm < 1) max_active_blocks_per_sm = 1; + + uint64_t sms_needed = + (total_blocks + max_active_blocks_per_sm - 1) / max_active_blocks_per_sm; + int sms_used = (int)std::min(sms_needed, + (uint64_t)dm.multi_processor_count); + if (sms_used < 1) sms_used = 1; + return sms_used; +} + void hapiProcessCuptiBuffers() { - #if CMK_CUDA GPUManager& gm = CsvAccess(gpu_manager); - + hapiPopulateDeviceProps(gm); // lazy: device_managers is ready by now + + // A kernel record can be parsed before the correlation record that names it: + // correlation and kernel records are queued at different points in the + // launch's life and land in buffers that complete independently. Only those + // kernels are parked for a second pass; one whose correlation is already + // known is filed immediately, so the common case never holds two copies of + // every record. + struct PendingKernel { + uint32_t correlation_id; + LBKernelRecord rec; + }; + std::vector pending; + pending.reserve(gm.cupti_pending_hint_); + uint32_t kernel_count = 0; - uint32_t corr_count = 0; + uint32_t object_corr_count = 0; + uint32_t invalid_duration_count = 0; + uint32_t attributed = 0; + uint32_t unattributed = 0; + uint32_t unresolved_token = 0; + uint32_t deferred = 0; + + // Resolve each distinct token once rather than once per kernel record. A + // round holds far more kernels than objects, and this lock is the same one + // every entry method needs, so taking it per record would both dominate the + // drain and stall PEs that are still running. + struct ResolvedToken { + LDObjKey key{}; + bool valid = false; + }; + std::unordered_map resolved_tokens; + + auto fileKernel = [&](const LBKernelRecord& rec, uint64_t object_token) { + if (object_token == HAPI_CUPTI_NO_OBJECT) { + gm.cupti_unattributed_kernels_.push_back(rec); + unattributed++; + return; + } + auto memo = resolved_tokens.find(object_token); + if (memo == resolved_tokens.end()) { + ResolvedToken entry; + { + std::lock_guard token_lock(gm.cupti_object_token_lock_); + entry.valid = gm.cupti_object_tokens_.resolve(object_token, entry.key); + } + memo = resolved_tokens.emplace(object_token, entry).first; + } + if (!memo->second.valid) { + gm.cupti_unattributed_kernels_.push_back(rec); + unattributed++; + unresolved_token++; + return; + } + gm.cupti_obj_kernel_records_[memo->second.key].push_back(rec); + attributed++; + }; + while (true) { - uint32_t record_count = 0; CuptiBufferItem item; // Pop one buffer from the queue - if (gm.cupti_buffer_queue_.empty()) { - break; + { + std::lock_guard lk(gm.cupti_queue_lock_); + if (gm.cupti_buffer_queue_.empty()) break; + item = gm.cupti_buffer_queue_.front(); + gm.cupti_buffer_queue_.pop(); + } + + // A buffer CUPTI never wrote to (or one it handed back empty) has nothing + // to parse, and passing it on would dereference whatever came back. + if (item.buffer == NULL || item.validSize == 0) { + free(item.buffer); + continue; } - item = gm.cupti_buffer_queue_.front(); - gm.cupti_buffer_queue_.pop(); // Parse records in this buffer CUpti_Activity *record = NULL; - // ckout<<"valid size for the CUPTI buffer: "<kind == CUPTI_ACTIVITY_KIND_EXTERNAL_CORRELATION) { CUpti_ActivityExternalCorrelation *corr = (CUpti_ActivityExternalCorrelation *)record; - corr_count++; - if(gm.cupti_correlation_db_.find(corr->correlationId)!=gm.cupti_correlation_db_.end()) - { - //out of order block - uint64_t curr_kernel_time = gm.cupti_correlation_db_[corr->correlationId]; - gm.cupti_obj_gpu_times_[corr->externalId] += curr_kernel_time; - gm.cupti_correlation_db_.erase(corr->correlationId); // Remove correlation ID after processing - } - else - { - gm.cupti_correlation_db_[corr->correlationId] = corr->externalId; + if (corr->externalKind == CUPTI_EXTERNAL_CORRELATION_KIND_UNKNOWN) { + object_corr_count++; + gm.cupti_object_correlation_db_[corr->correlationId] = corr->externalId; } } else if (record->kind == CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL || record->kind == CUPTI_ACTIVITY_KIND_KERNEL) { kernel_count++; CUpti_ActivityKernel4 *kernel = (CUpti_ActivityKernel4 *)record; - uint64_t duration_ns = kernel->end - kernel->start; - // ckout<<"the current kernel's duration is "<correlationId); - if (it != gm.cupti_correlation_db_.end()) { - uint64_t obj_id = it->second; - gm.cupti_obj_gpu_times_[obj_id] += duration_ns; - gm.cupti_correlation_db_.erase(it); // Remove correlation ID after processing - } - else - { - // CmiPrintf("found an out of order entry\n"); - gm.cupti_correlation_db_[kernel->correlationId] = duration_ns; + + DeviceManager* dm = findDeviceManager(gm, kernel->deviceId); + + LBKernelRecord rec{}; + rec.start_ns = kernel->start; + rec.end_ns = kernel->end; + rec.device_id = kernel->deviceId; + rec.sms_used = dm ? computeKernelSMs(*dm, kernel) : 1; + if (rec.end_ns <= rec.start_ns) invalid_duration_count++; + + auto object = gm.cupti_object_correlation_db_.find(kernel->correlationId); + if (object != gm.cupti_object_correlation_db_.end()) { + const uint64_t object_token = object->second; + gm.cupti_object_correlation_db_.erase(object); + fileKernel(rec, object_token); + } else { + pending.push_back({kernel->correlationId, rec}); + deferred++; } } } - - // ckout<<"number of CUPTI records in this buffer: "<second; + gm.cupti_object_correlation_db_.erase(object); + fileKernel(pending_kernel.rec, object_token); + } + + // Size next round's parked vector from this one: the straggler count is a + // property of how CUPTI is batching buffers, which changes slowly. + gm.cupti_pending_hint_ = deferred; + + // Every entry method pushes a correlation ID, and CUPTI emits a record for + // each runtime call made under it -- memcpys, syncs and so on, not just + // kernel launches. Those never match a kernel record. We flush before + // draining, so any kernel that was going to arrive has arrived; whatever is + // left over is non-kernel traffic and would otherwise accumulate without + // bound (entry-method correlation emits on the order of 10^5 records a step). + size_t object_corr_dropped = gm.cupti_object_correlation_db_.size(); + gm.cupti_object_correlation_db_.clear(); + + if (cuptiDebugLevel() > 1) { + CmiPrintf("HAPI[pe=%d]: hapiProcessCuptiBuffers kernels=%u " + "object_correlations=%u attributed=%u unattributed=%u " + "deferred=%u invalid_durations=%u unresolved_tokens=%u " + "objects=%zu object_corr_dropped=%zu\n", + CmiMyPe(), kernel_count, object_corr_count, attributed, + unattributed, deferred, invalid_duration_count, unresolved_token, + gm.cupti_obj_kernel_records_.size(), object_corr_dropped); + } +} + +// Convert this process's raw kernel timeline into one SM-utilization- +// normalized load per object, in seconds of whole-device occupancy. +// +// Per-device sweep-line over all kernel intervals. At each event (kernel start +// or end) the interval's device time is split among the kernels then running, +// in proportion to the SMs each occupies. +// +// Because the result is device-seconds of demand rather than elapsed time, it +// is (to first order) invariant to how the objects happen to be placed right +// now, which is what makes it usable as a load estimate for a *different* +// placement. +// +// This runs in the process that produced the records, before the stats leave +// for the central LB: every PE bound to a device is in the same process, so the +// timeline for that device is already complete here. If several processes share +// one GPU, each sees only its own kernels and the contention between processes +// is not modelled. +void hapiNormalizeCuptiLoads() { GPUManager& gm = CsvAccess(gpu_manager); + gm.cupti_obj_norm_load_.clear(); + + struct SweepKernel { + LDObjKey obj_key; + uint64_t start_ns; + uint64_t end_ns; + int sms_used; + bool attributed; // false => consumes SMs but earns no load + // Whole-device occupancy this kernel earned, filled in by the sweep. + double demand; + }; + std::unordered_map> byDevice; + for (const auto& kv : gm.cupti_obj_kernel_records_) { + for (const LBKernelRecord& k : kv.second) { + if (k.end_ns <= k.start_ns) continue; + byDevice[k.device_id].push_back({kv.first, k.start_ns, k.end_ns, + k.sms_used, true, 0.0}); + } + } + for (const LBKernelRecord& k : gm.cupti_unattributed_kernels_) { + if (k.end_ns <= k.start_ns) continue; + byDevice[k.device_id].push_back({LDObjKey{}, k.start_ns, k.end_ns, + k.sms_used, false, 0.0}); + } - gm.cupti_obj_gpu_times_.clear(); - gm.cupti_correlation_db_.clear(); + size_t total_kernels = 0, devices_normalized = 0; + for (auto& kv : byDevice) { + std::vector& kernels = kv.second; + if (kernels.empty()) continue; + + DeviceManager* dm = findDeviceManager(gm, kv.first); + int total_sms = (dm != nullptr) ? dm->multi_processor_count : 0; + if (total_sms <= 0) continue; // unknown device size -- cannot normalize + total_kernels += kernels.size(); + devices_normalized++; + + // Two events per kernel. END sorts before START on a tie so a kernel + // ending at instant t does not briefly count alongside one starting at t. + struct Event { uint64_t time; int kind; int kidx; }; // kind: 0=END, 1=START + std::vector events; + events.reserve(2 * kernels.size()); + for (int ki = 0; ki < (int)kernels.size(); ++ki) { + events.push_back({kernels[ki].start_ns, 1, ki}); + events.push_back({kernels[ki].end_ns, 0, ki}); + } + std::sort(events.begin(), events.end(), + [](const Event& a, const Event& b) { + if (a.time != b.time) return a.time < b.time; + return a.kind < b.kind; + }); + + // Kernels holding the device right now. Order is irrelevant: the two loops + // below sum sms_used and then credit each kernel its own share, and both + // give the same answer in any order. (An earlier version handed SMs out + // FIFO by submission, which is what the proportional split above replaced.) + // A vector with swap-and-pop erase beats a node-based container here -- + // these sets are small, and every event does one insert or one erase. + std::vector active; + active.reserve(kernels.size()); + + uint64_t t_prev = events.front().time; + for (const auto& ev : events) { + if (ev.time > t_prev && !active.empty()) { + double dt_s = (double)(ev.time - t_prev) / 1.0e9; + + // Split this interval's occupancy in proportion to what each active + // kernel asked for. + // + // Handing SMs out FIFO by submission and stopping at the first kernel + // that finds the pool empty does not work: every kernel behind that + // point earns nothing for the interval, so whether an object is + // credited depends on where its kernels land in the launch order + // relative to its neighbours' -- which reshuffles whenever placement + // changes. That makes the metric unstable in the one direction that + // matters: a device running more objects oversubscribes harder, so more + // of its kernels fall past the cut-off and its objects measure cheaper + // than they are, and a balancer reading that sends the busiest node + // still more work. Proportional sharing is also the better model of the + // hardware: kernels that oversubscribe an SM pool time-share it rather + // than running strictly in submission order. + long want = 0; + for (int ki : active) { + if (kernels[ki].sms_used > 0) want += kernels[ki].sms_used; + } + if (want > 0) { + // Split the interval's DEVICE TIME among the kernels holding the + // device, in proportion to the SMs each asked for. The weights decide + // how concurrent work is divided; they must not decide how much there + // is to divide. + // + // Charging dt * sms_used/total_sms instead -- SM-occupancy seconds -- + // records a kernel that held the device for an interval at low + // occupancy as almost no load. That is the right measure of + // throughput and the wrong one for balancing: an object whose kernels + // tie up the GPU for 10ms costs its PE 10ms whether they fill 4 SMs + // or 80. The error is also not uniform, which is what makes it + // dangerous: an imbalanced distribution gives lightly-loaded objects + // small grids, so a node holding many of them runs many low-occupancy + // kernels -- busy the whole interval, yet reporting almost nothing -- + // and the balancer reads it as idle and sends it still more work. + // + // Weighting by sms_used keeps a big kernel worth more than a small + // one running beside it, while the interval total stays dt: summed + // over a device, attributed demand equals its busy time. + for (int ki : active) { + const int used = kernels[ki].sms_used; + if (used <= 0) continue; + // Computed for unowned kernels too; they are still not credited to + // any object below, but leaving their share uncomputed is what + // makes such a loss invisible to the audit. + kernels[ki].demand += dt_s * ((double)used / (double)want); + } + } + } + if (ev.kind == 1) { + active.push_back(ev.kidx); + } else { + // Swap-and-pop: nothing here depends on the order of the remainder. + auto it = std::find(active.begin(), active.end(), ev.kidx); + if (it != active.end()) { *it = active.back(); active.pop_back(); } + } + t_prev = ev.time; + } + + for (const SweepKernel& k : kernels) { + if (k.demand <= 0.0) continue; + if (!k.attributed) continue; + gm.cupti_obj_norm_load_[k.obj_key] += k.demand; + } + } + + if (cuptiDebugLevel() > 1) { + CmiPrintf("HAPI[pe=%d]: hapiNormalizeCuptiLoads %zu kernels across %zu " + "device(s) -> %zu objects\n", + CmiMyPe(), total_kernels, devices_normalized, + gm.cupti_obj_norm_load_.size()); + } + +} + +void hapiPrepareCuptiLoads(uint64_t epoch) { + GPUManager& gm = CsvAccess(gpu_manager); + std::lock_guard lk(gm.cupti_prepare_lock_); + // Already built for this epoch or a later one. + if (gm.cupti_loads_ready_ && epoch <= gm.cupti_loads_epoch_) return; + + const bool timeIt = (getenv("CHARM_LB_CUPTI_TIME") != nullptr); + const double t0 = timeIt ? CmiWallTimer() : 0.0; + // Only flush while CUPTI is attached: an application driving its own + // instrumentation window may already have switched tracing off, and its stop + // path flushed on the way out, so there is nothing left to pull. + // + // FLUSH_FORCED rather than a plain flush: without it CUPTI hands back only + // the buffers it considers complete, and the kernel records still sitting in + // a partly-filled buffer are read a round late or not at all. + if (hapiCuptiTracingActive()) + CUPTI_SAFE_CALL(cuptiActivityFlushAll(CUPTI_ACTIVITY_FLAG_FLUSH_FORCED)); + const double t1 = timeIt ? CmiWallTimer() : 0.0; + hapiProcessCuptiBuffers(); + const double t2 = timeIt ? CmiWallTimer() : 0.0; + hapiNormalizeCuptiLoads(); + // Close the measurement window exactly where it was read. These records were + // just consumed; anything recorded from here on belongs to the next round. + // + // Dropping them at MigrationDone instead -- after the strategy has run and + // the migrations have executed -- means kernels running through all of that + // are recorded and then discarded, belonging to no round at all, and the + // amount discarded depends on how long that round's LB step took. That is a + // feedback loop between migration count and measured load. + gm.cupti_obj_kernel_records_.clear(); + gm.cupti_unattributed_kernels_.clear(); + if (timeIt) { + const double t3 = CmiWallTimer(); + CmiPrintf("[LBCUPTI pe=%d] flush=%.3fs process=%.3fs normalize=%.3fs total=%.3fs\n", + CmiMyPe(), t1 - t0, t2 - t1, t3 - t2, t3 - t0); + fflush(stdout); + } + + gm.cupti_loads_ready_ = true; + gm.cupti_loads_epoch_ = epoch; +} + +void hapiClearCuptiData() { + GPUManager& gm = CsvAccess(gpu_manager); + // Same lock as hapiPrepareCuptiLoads: this drops the maps that function + // builds and that every PE reads, so it must not run underneath either. + std::lock_guard lk(gm.cupti_prepare_lock_); + gm.cupti_loads_ready_ = false; + gm.cupti_loads_epoch_ = 0; + + gm.cupti_obj_kernel_records_.clear(); + gm.cupti_unattributed_kernels_.clear(); + gm.cupti_obj_norm_load_.clear(); + // The correlation map is drained alongside the CUPTI buffers in + // hapiProcessCuptiBuffers. Do not clear the object-token table: later epochs + // must reuse the same token for the same full LB identity, and the per-PE + // caches in front of it have no invalidation protocol. } +#else /* !HAPI_CUPTI_LB */ + +void hapiProcessCuptiBuffers() {} +void hapiNormalizeCuptiLoads() {} +void hapiPrepareCuptiLoads(uint64_t epoch) {} +void hapiClearCuptiData() {} + +#endif /* HAPI_CUPTI_LB */ + +// Build this round's per-object GPU loads exactly once per round, however many +// PE threads call in. A load balancer's per-PE barrier fires when THAT PE's +// objects are at AtSync, but the records are shared by the whole process, so +// the drain has to wait for the last PE rather than run on the first. +bool hapiCuptiArrive(uint64_t epoch, int expected) { +#ifdef HAPI_CUPTI_LB + GPUManager& gm = CsvAccess(gpu_manager); + std::lock_guard lk(gm.cupti_arrive_lock_); + // A new round resets the count. Rounds are strictly sequential -- a step + // completes on a job-wide barrier before the next can start -- so a stale + // count from a previous epoch can only mean that epoch is over. + if (gm.cupti_arrive_epoch_ != epoch) { + gm.cupti_arrive_epoch_ = epoch; + gm.cupti_arrive_count_ = 0; + } + gm.cupti_arrive_count_++; + if (gm.cupti_arrive_count_ < expected) return false; + gm.cupti_arrive_count_ = 0; + return true; +#else + // Nothing to build, so nothing to wait for: every caller proceeds and reads + // an empty load map. + return true; #endif +} + +#endif /* CMK_CUDA && CMK_LBDB_ON */ // Initialize per-PE variables @@ -665,25 +1212,19 @@ static void hapiMapping(char** argv) { } #ifndef HAPI_CUDA_CALLBACK -void recordEvent(hapiStream_t stream, const CkCallback& cb, void* cb_msg, hapiWorkRequest* wr = NULL, CkMigratable* obj = NULL, hapiEvent_t start_ev = NULL) { - // if(obj!=NULL) - // CmiAbort("non null without HAPI hapi CALLBACK"); +void recordEvent(hapiStream_t stream, const CkCallback& cb, void* cb_msg, hapiWorkRequest* wr = NULL) { // create hapi event / get hapi event from the pool and insert into stream hapiEvent_t ev; auto& hapi_event_pool_local = CpvAccess(hapi_event_pool); if(hapi_event_pool_local.size() == 0) { - #ifdef HAPI_CUPTI_LB - hapiEventCreateWithFlags(&ev, hapiEventDefault); - #else hapiEventCreateWithFlags(&ev, hapiEventDisableTiming); - #endif } else { ev = hapi_event_pool_local.front(); hapi_event_pool_local.pop(); } hapiEventRecord(ev, stream); - hapiEvent hev(ev, cb, cb_msg, wr, obj, start_ev); + hapiEvent hev(ev, cb, cb_msg, wr); // push event information in queue CpvAccess(hapi_event_queue).push(hev); @@ -1641,17 +2182,6 @@ void hapiPollEvents(void* param) { if (hapiEventQuery(hev.event) == hapiSuccess) { queue.pop(); // TODO: investigate possible race condition with charm4py futures - temporarily resolved by popping here -#ifdef HAPI_CUPTI_LB - if (hev.obj) { - // CmiPrintf("should not be printed w/o hapi hapi callback \n"); - float gpu_time; - hapiEventElapsedTime(&gpu_time, hev.start_ev, hev.event); - // hapiEventElapsedTime returns ms, convert to seconds to match wallTime units - double gpu_time_s = gpu_time / 1000.0; - hev.obj->setObjGPUTime(gpu_time_s + hev.obj->getObjGPUTime()); - hapiEventDestroy(hev.start_ev); - } else -#endif // invoke Charm++ callback if one was given hev.cb.send(hev.cb_msg); @@ -1705,64 +2235,113 @@ hapiStream_t hapiGetStream() { return ret; } +#if CMK_CUDA && CMK_LBDB_ON #ifdef HAPI_CUPTI_LB -// Lightweight HAPI, to be invoked after data transfer or kernel execution. -void hapiRecordTime(hapiStream_t stream, hapiEvent_t start) { - Chare* obj = CkActiveObj(); - if (obj && dynamic_cast(obj)) { - - #ifndef HAPI_CUDA_CALLBACK - // record hapi event - recordEvent(stream, CkCallback(), NULL, NULL, dynamic_cast(obj), start); -#else - #error hapi record time with HAPI_CUDA_CALLBACK not supported -#endif - // while there is an ongoing workrequest, quiescence should not be detected - // even if all PEs seem idle - CmiAssert(hapiQdCreate); - hapiQdCreate(1); +// How many external-correlation IDs this PE has actually pushed and not yet +// popped. Tracing is switched on and off from inside entry methods, so a push +// can be skipped while its matching pop still runs (or the reverse). Pairing +// the pop against this count rather than against the tracing flag keeps +// CUPTI's stack balanced across those transitions; without it the pop reports +// CUPTI_ERROR_QUEUE_EMPTY and attribution drifts. Each Charm++ PE is its own +// thread, so thread_local is per-PE. +static thread_local int cupti_pushed_depth = 0; +// The detach generation this PE last observed; see GPUManager::cupti_generation_. +static thread_local uint64_t cupti_seen_generation = 0; + +// This PE's view of the process-wide token table. Every entry method on a +// migratable chare needs its object's token, and the table behind it is shared +// by every PE in the process, so consulting it under the node-wide lock would +// serialize the whole process on one mutex for the length of the run. +// +// The table is append-only for the lifetime of the process, which is what makes +// this cache safe without any invalidation protocol: an entry, once correct, +// stays correct, including across migration (the destination PE simply misses +// once and interns the same token the source PE already has). Anything that +// gains the ability to clear or renumber GpuObjectTokenTable must also +// invalidate these caches. +static thread_local std::unordered_map + cupti_local_object_tokens; + +// Drop this PE's outstanding push count if CUPTI has been detached since we +// last looked -- the stack those pushes referred to no longer exists, so +// popping against it would report CUPTI_ERROR_QUEUE_EMPTY. +static inline void hapiCuptiSyncGeneration(GPUManager& gm) { + if (cupti_seen_generation != gm.cupti_generation_) { + cupti_seen_generation = gm.cupti_generation_; + cupti_pushed_depth = 0; } } -#endif -#ifdef HAPI_CUPTI_LB uint64_t hapiCuptiPushObjCorrelation() { - // printf("seeing CsvAccess(gpu_manager).cupti_initialized_ as %d\n", CsvAccess(gpu_manager).cupti_initialized_); - if (!CsvAccess(gpu_manager).cupti_initialized_) return 0; - - // Get the active Charm++ object - Chare* chare = CkActiveObj(); - if (!chare) - CmiAbort("hapiCuptiPushObjCorrelation call without active object is not possible"); - - CkMigratable* mig = dynamic_cast(chare); - // printf("mig %p\n", mig); - if (!mig) return 0; - - // Use the raw element ID as the external correlation ID - // CmiUInt8 is a 64-bit unique object identifier - uint64_t obj_id = (uint64_t)mig->ckGetID(); -#if CMK_CUDA + GPUManager& gm = CsvAccess(gpu_manager); + // Gated on tracing rather than initialization: this runs on every entry + // method, so when tracing is off the whole body -- the active-object lookup + // and two CUPTI calls -- must be skipped, not just wasted. + if (!gm.cupti_tracing_active_.load(std::memory_order_relaxed)) return 0; + hapiCuptiSyncGeneration(gm); + + // The CUPTI external ID is a process-local token for the complete LB object + // key. Using CkMigratable::ckGetID() here loses the object-manager identity + // and aliases equal element IDs from different chare arrays. + uint64_t object_token = HAPI_CUPTI_NO_OBJECT; + if (CkLocRec* active = CkActiveLocRec()) { + const LDObjHandle& handle = active->getLdHandle(); + LDObjKey key; + key.omID() = handle.omID(); + key.objID() = handle.objID(); + + // Steady state is a PE-local hit: the shared lock is taken only the first + // time this PE runs a given object, so it costs O(objects that ever run + // here) acquisitions rather than one per entry method. + auto cached = cupti_local_object_tokens.find(key); + if (cached != cupti_local_object_tokens.end()) { + object_token = cached->second; + } else { + { + std::lock_guard token_lock(gm.cupti_object_token_lock_); + if (!gm.cupti_object_tokens_.intern(key, object_token)) + CmiAbort("HAPI: exhausted CUPTI object-correlation tokens"); + } + cupti_local_object_tokens.emplace(key, object_token); + } + } + + // Always push, even with the sentinel, so that the matching pop always has + // something to remove; an unbalanced stack would mis-attribute every + // subsequent kernel. CUPTI_SAFE_CALL(cuptiActivityPushExternalCorrelationId( - CUPTI_EXTERNAL_CORRELATION_KIND_UNKNOWN, obj_id)); -#endif - // printf("pushed corr id\n"); + CUPTI_EXTERNAL_CORRELATION_KIND_UNKNOWN, object_token)); + ++cupti_pushed_depth; - return obj_id; + return object_token; } void hapiCuptiPopObjCorrelation() { - if (!CsvAccess(gpu_manager).cupti_initialized_) return; + // Runs the generation check even when tracing is off: a detach may have + // happened between this entry method's push and its pop, and the stale count + // has to be cleared here rather than on the next push. + GPUManager& gm = CsvAccess(gpu_manager); + hapiCuptiSyncGeneration(gm); + + // Pop exactly what was pushed. Checking the tracing flag here instead would + // pop entries this PE never pushed, once tracing is switched on part-way + // through an entry method. + if (cupti_pushed_depth == 0 || !gm.cupti_initialized_) return; + --cupti_pushed_depth; - // printf("popped corr id\n"); uint64_t tag; -#if CMK_CUDA CUPTI_SAFE_CALL(cuptiActivityPopExternalCorrelationId( CUPTI_EXTERNAL_CORRELATION_KIND_UNKNOWN, &tag)); -#endif } -#endif /* HAPI_CUPTI_LB */ + +#else /* !HAPI_CUPTI_LB */ + +uint64_t hapiCuptiPushObjCorrelation() { return 0; } +void hapiCuptiPopObjCorrelation() {} + +#endif /* HAPI_CUPTI_LB */ +#endif /* CMK_CUDA && CMK_LBDB_ON */ // Lightweight HAPI, to be invoked after data transfer or kernel execution. void hapiAddCallback(hapiStream_t stream, const CkCallback& cb, void* cb_msg) { @@ -1850,3 +2429,19 @@ uint64_t hapiMyDevice() { return (static_cast(physical_node_id) << 32) | my_device; } +int hapiMyDeviceTotalSMs() { + GPUManager& gm = CsvAccess(gpu_manager); + int local_id = CpvAccess(my_device_id); + if (local_id < 0 || local_id >= (int)gm.device_managers.size()) return 0; + DeviceManager& dm = gm.device_managers[local_id]; + if (!dm.props_initialized) { + // The lazy population in hapiProcessCuptiBuffers has not run yet (or CUPTI + // is not in this build at all), so ask for just this one attribute. + int count = 0; + hapiCheck(cudaDeviceGetAttribute(&count, cudaDevAttrMultiProcessorCount, + dm.global_index)); + return count; + } + return dm.multi_processor_count; +} + diff --git a/src/ck-core/CMakeLists.txt b/src/ck-core/CMakeLists.txt index d67136c0fe..ffde2b4818 100644 --- a/src/ck-core/CMakeLists.txt +++ b/src/ck-core/CMakeLists.txt @@ -66,6 +66,7 @@ set(ldb-cxx-sources set(ldb-h-sources ../ck-ldb/BaseLB.h ../ck-ldb/CentralLBMsg.h ../ck-ldb/DistBaseLB.h ../ck-ldb/DistributedLB.h ../ck-ldb/HybridBaseLB.h ../ck-ldb/HybridLBMsg.h + ../ck-ldb/GreedyRefineCentralGPULB.h ../ck-ldb/LBComm.h ../ck-ldb/LBDatabase.h ../ck-ldb/LBManager.h ../ck-ldb/LBMachineUtil.h ../ck-ldb/LBOM.h ../ck-ldb/LBObj.h ../ck-ldb/LBSimulation.h @@ -92,8 +93,8 @@ add_dependencies(ckmain ck) # CommonLBs used to be a subset of EveryLB, but they were unified as # part of the LB refactor. However, maintain both targets for # backwards compatibility and for possible future divergence. -add_library(moduleCommonLBs ${ldb-cxx-sources} ${ldb-h-sources} ../ck-ldb/MetisLB.C ../ck-ldb/MetisLB.h) -add_library(moduleEveryLB ${ldb-cxx-sources} ${ldb-h-sources} ../ck-ldb/MetisLB.C ../ck-ldb/MetisLB.h) +add_library(moduleCommonLBs ${ldb-cxx-sources} ${ldb-h-sources} ../ck-ldb/MetisLB.C ../ck-ldb/MetisLB.h ../ck-ldb/GreedyRefineCentralGPULB.C ../ck-ldb/GreedyRefineCentralGPULB.h) +add_library(moduleEveryLB ${ldb-cxx-sources} ${ldb-h-sources} ../ck-ldb/MetisLB.C ../ck-ldb/MetisLB.h ../ck-ldb/GreedyRefineCentralGPULB.C ../ck-ldb/GreedyRefineCentralGPULB.h) configure_file(../ck-ldb/libmoduleCommonLBs.dep ${CMAKE_BINARY_DIR}/lib/ COPYONLY) configure_file(../ck-ldb/libmoduleEveryLB.dep ${CMAKE_BINARY_DIR}/lib/ COPYONLY) add_dependencies(moduleCommonLBs ck ckmetis) @@ -117,6 +118,8 @@ add_library(moduleTreeLB ../ck-ldb/TreeLB.C) add_dependencies(moduleTreeLB ck) add_library(moduleRecBipartLB ../ck-ldb/RecBipartLB.C ../ck-ldb/RecBipartLB.h) add_dependencies(moduleRecBipartLB ck) +add_library(moduleGreedyRefineCentralGPULB ../ck-ldb/GreedyRefineCentralGPULB.C ../ck-ldb/GreedyRefineCentralGPULB.h) +add_dependencies(moduleGreedyRefineCentralGPULB ck) add_library(moduleScotchLB EXCLUDE_FROM_ALL ../ck-ldb/ScotchLB.C ../ck-ldb/ScotchLB.h) configure_file(../ck-ldb/libmoduleScotchLB.dep ${CMAKE_BINARY_DIR}/lib/ COPYONLY) add_dependencies(moduleScotchLB ck) @@ -142,7 +145,7 @@ foreach(filename ${ck-h-sources} ${ldb-h-sources}) endforeach() target_include_directories(ck PRIVATE . ../ck-ldb ../ck-perf ../ck-cp ../util/topomanager ../conv-perf) -if(HAPI_CUPTI_LB AND BUILD_CUDA AND CUPTI_LIBRARY) # dormant until the GPU-LB series (plan item 11) +if(HAPI_CUPTI_LB AND BUILD_CUDA AND CUPTI_LIBRARY) target_include_directories(ck PRIVATE "${CUDA_DIR}/extras/CUPTI/include") target_link_libraries(ck ${CUPTI_LIBRARY}) endif() diff --git a/src/ck-core/ck.C b/src/ck-core/ck.C index 538fa5a3b1..ca4c4e3263 100644 --- a/src/ck-core/ck.C +++ b/src/ck-core/ck.C @@ -603,6 +603,15 @@ inline Chare *_popObj(void) { } } +#if CMK_LBDB_ON && CMK_CUDA +// Declared in hapi.h; forward-declared here to avoid pulling the full HAPI +// header into ck.C. These bracket the entry-method body so that CUPTI kernel +// records can be correlated back to the active migratable object, however the +// application happens to launch its kernels. +extern uint64_t hapiCuptiPushObjCorrelation(); +extern void hapiCuptiPopObjCorrelation(); +#endif + inline void _ckStartTiming(void) { #if CMK_LBDB_ON auto *active = CkActiveLocRec(); @@ -622,6 +631,13 @@ void CkCallstackPush(Chare *obj) { _ckStopTiming(); // suspend timing of the previous obj _pushObj(obj); // push the current object onto the stack _ckStartTiming(); // start timing the current obj +#if CMK_LBDB_ON && CMK_CUDA + // After _pushObj, `obj` is the active chare. Push a CUPTI external + // correlation ID -- a token for its full LB key, or a sentinel when there is + // no migratable owner -- so every kernel launched until the matching pop is + // attributed to this chare. Structurally paired 1:1 with CkCallstackPop. + hapiCuptiPushObjCorrelation(); +#endif } // removes all instances of ( obj ) from the stack @@ -646,6 +662,10 @@ void CkCallstackUnwind(Chare *obj) { // pops ( obj ) from the stack (and manages timing) void CkCallstackPop(Chare *obj) { +#if CMK_LBDB_ON && CMK_CUDA + // Paired 1:1 with the push in CkCallstackPush. + hapiCuptiPopObjCorrelation(); +#endif _ckStopTiming(); // stop timing the current obj auto *popd = _popObj(); // pop it from the stack CkAssertMsg(!popd || popd == obj, "object tracking mismatch"); diff --git a/src/ck-core/cklocation.C b/src/ck-core/cklocation.C index fa618c73af..f8d8033bc7 100644 --- a/src/ck-core/cklocation.C +++ b/src/ck-core/cklocation.C @@ -11,6 +11,12 @@ #include "TopoManager.h" #include "charm++.h" #include "ck.h" +#if CMK_CUDA +// For the staged device migration payload: hapiMalloc/hapiFree hold it, and +// CkDeviceBuffer hands it to the device zerocopy layer. +#include "hapi.h" +#include "ckrdmadevice.h" +#endif #include "cksyncbarrier.h" #include "hilbert.h" #include "partitioning_strategies.h" @@ -1833,6 +1839,8 @@ void CkMigratable::UserSetLBLoad() // user can call this helper function to set obj load (for model-based lb) void CkMigratable::setObjTime(double cputime) { myRec->setObjTime(cputime); } double CkMigratable::getObjTime() { return myRec->getObjTime(); } +void CkMigratable::setObjGPUTime(double gputime) { myRec->setObjGPUTime(gputime); } +double CkMigratable::getObjGPUTime() { return myRec->getObjGPUTime(); } # if CMK_LB_USER_DATA /** @@ -1949,6 +1957,12 @@ void CkMigratable::AtSync(int waitForMigration) this->virtual_pup(ps); if (_lb_psizer_on) setPupSize(ps.size()); +#if CMK_CUDA + // Device bytes are reported separately and exactly: a balancer weighing a + // migration against free device memory needs the real figure, and pupSize + // is an encoded approximation tuned for host state. + setGPUPupSize(ps.gpu_size()); +#endif if (_lb_args.metaLbOn()) myRec->getMetaBalancer()->SetCharePupSize(ps.size()); } @@ -2047,6 +2061,13 @@ void CkMigratable::setMigratable(int migratable) { myRec->setMigratable(migratab void CkMigratable::setPupSize(size_t obj_pup_size) { myRec->setPupSize(obj_pup_size); } +void CkMigratable::setGPUPupSize(size_t obj_gpu_pup_size) +{ +#if CMK_CUDA + myRec->setGPUPupSize(obj_gpu_pup_size); +#endif +} + void CkMigratable::CkAddThreadListeners(CthThread tid, void* msg) { Chare::CkAddThreadListeners(tid, msg); // for trace @@ -2056,6 +2077,8 @@ void CkMigratable::CkAddThreadListeners(CthThread tid, void* msg) #else void CkMigratable::setObjTime(double cputime) {} double CkMigratable::getObjTime() { return 0.0; } +void CkMigratable::setObjGPUTime(double gputime) {} +double CkMigratable::getObjGPUTime() { return 0.0; } # if CMK_LB_USER_DATA void* CkMigratable::getObjUserData(int idx) { return NULL; } @@ -2143,6 +2166,22 @@ double CkLocRec::getObjTime() lbmgr->GetObjLoad(ldHandle, walltime, cputime); return walltime; } +void CkLocRec::setObjGPUTime(double gputime) +{ +#if CMK_CUDA + lbmgr->EstObjGPULoad(ldHandle, gputime); +#endif +} +double CkLocRec::getObjGPUTime() +{ +#if CMK_CUDA + LBRealType gputime; + lbmgr->GetObjGPULoad(ldHandle, gputime); + return gputime; +#else + return 0.0; +#endif +} # if CMK_LB_USER_DATA void* CkLocRec::getObjUserData(int idx) { return lbmgr->GetDBObjUserData(ldHandle, idx); } # endif @@ -2280,6 +2319,13 @@ void CkLocRec::setPupSize(size_t obj_pup_size) lbmgr->setPupSize(ldHandle, obj_pup_size); } +void CkLocRec::setGPUPupSize(size_t obj_gpu_pup_size) +{ +#if CMK_CUDA + lbmgr->setGPUPupSize(ldHandle, obj_gpu_pup_size); +#endif +} + #endif // Call ckDestroy for each record, which deletes the record, and ~CkLocRec() @@ -2983,10 +3029,14 @@ void CkLocMgr::emigrate(CkLocRec* rec, int toPe) // First pass: find size of migration message size_t bufSize; + size_t gpuBufSize = 0; { PUP::sizer p(PUP::er::IS_MIGRATION); pupElementsFor(p, rec, CkElementCreation_migrate); bufSize = p.size(); +#if CMK_CUDA + gpuBufSize = p.gpu_size(); +#endif } #if CMK_ERROR_CHECKING if (bufSize > std::numeric_limits::max()) @@ -2996,6 +3046,22 @@ void CkLocMgr::emigrate(CkLocRec* rec, int toPe) } #endif + void* gpuMsg = nullptr; +#if CMK_CUDA + if (gpuBufSize > 0) + { + if (gpuBufSize > (size_t)std::numeric_limits::max()) + CkAbort("Cannot migrate an object with more than %d bytes of device " + "state!\n", std::numeric_limits::max()); + // A staging copy, deliberately: with the element's device state copied + // out, the element can be destroyed as soon as the pack below finishes, + // and only this buffer has to survive until the destination has read it. + // Handing the destination the element's own buffers instead would mean + // keeping the element alive across the transfer. + hapiCheck(hapiMalloc(&gpuMsg, gpuBufSize)); + } +#endif + // Allocate and pack into message CkArrayElementMigrateMessage* msg = new (bufSize, 0) CkArrayElementMigrateMessage(idx, id, @@ -3005,10 +3071,11 @@ void CkLocMgr::emigrate(CkLocRec* rec, int toPe) false, #endif bufSize, managers.size(), - cache->getEpoch(id) + 1); + cache->getEpoch(id) + 1, + gpuBufSize > 0); { - PUP::toMem p(msg->packData, PUP::er::IS_MIGRATION); + PUP::toMem p(msg->packData, gpuMsg, PUP::er::IS_MIGRATION); p.becomeDeleting(); pupElementsFor(p, rec, CkElementCreation_migrate); if (p.size() != bufSize) @@ -3019,12 +3086,36 @@ void CkLocMgr::emigrate(CkLocRec* rec, int toPe) bufSize, p.size()); CkAbort("Array element's pup routine has a direction mismatch.\n"); } +#if CMK_CUDA + if (p.gpu_size() != gpuBufSize) + { + CkError( + "ERROR! Array element claimed it was %zu device bytes to a " + "sizing PUP::er, but copied %zu device bytes into the packing " + "PUP::er!\n", + gpuBufSize, p.gpu_size()); + CkAbort("Array element's pup routine has a device direction mismatch.\n"); + } +#endif } DEBM((AA "Migrated index size %s to %d \n" AB, idx2str(idx), toPe)); thisProxy[toPe].immigrate(msg); +#if CMK_CUDA + if (gpuBufSize > 0) + { + // Dispatch the device payload from a fresh entry method rather than + // sending it here: the element this migration is tearing down is still + // the running one at this point, and the payload belongs to the runtime + // rather than to it. sendGPUMsg runs once this call has returned and the + // element is gone. + sendGPUBuffers[id] = GPUMigrateData(toPe, gpuBufSize, gpuMsg); + thisProxy[CkMyPe()].sendGPUMsg(id); + } +#endif + duringMigration = true; for (auto itr = managers.begin(); itr != managers.end(); ++itr) { @@ -3064,8 +3155,6 @@ void CkLocMgr::immigrate(CkArrayElementMigrateMessage* msg) { const CkArrayIndex& idx = msg->idx; - PUP::fromMem p(msg->packData, PUP::er::IS_MIGRATION); - if (msg->nManagers < managers.size()) CkAbort("Array element arrived from location with fewer managers!\n"); if (msg->nManagers > managers.size()) @@ -3077,6 +3166,35 @@ void CkLocMgr::immigrate(CkArrayElementMigrateMessage* msg) return; } +#if CMK_CUDA + // The device payload is a separate send and may not be here yet. Hold the + // host message until it lands; immigrateGPU resumes from the other side. + if (msg->hasGPUMsg && receivedDeviceMsgs.find(msg->id) == receivedDeviceMsgs.end()) + { + bufferedHostMigrateMsgs[msg->id] = msg; + return; + } + immigrateWithDevice(msg); +} + +/// Both halves of the migration are in hand: unpack the element, reading its +/// device state out of the landed payload. +void CkLocMgr::immigrateWithDevice(CkArrayElementMigrateMessage* msg) +{ + const CkArrayIndex& idx = msg->idx; + void* gpuData = nullptr; + if (msg->hasGPUMsg) + { + auto it = receivedDeviceMsgs.find(msg->id); + CmiAssert(it != receivedDeviceMsgs.end()); + gpuData = it->second; + receivedDeviceMsgs.erase(it); + } + PUP::fromMem p(msg->packData, gpuData, PUP::er::IS_MIGRATION); +#else + PUP::fromMem p(msg->packData, PUP::er::IS_MIGRATION); +#endif + insertID(idx, msg->id); // Create a record for this element @@ -3108,6 +3226,16 @@ void CkLocMgr::immigrate(CkArrayElementMigrateMessage* msg) CkAbort("Array element's pup routine has a direction mismatch.\n"); } +#if CMK_CUDA + if (msg->hasGPUMsg) + { + // The element's own device buffers now hold copies of everything in the + // landed payload, so the landing buffer is finished with. The sender's + // staged copy is not this side's concern: its source callback releases it. + hapiCheck(hapiFree(gpuData)); + } +#endif + if (!zcRgetsActive) { // Let all the elements know we've arrived @@ -3117,6 +3245,76 @@ void CkLocMgr::immigrate(CkArrayElementMigrateMessage* msg) delete msg; } +#if CMK_CUDA +/// Source callback for a migration payload: the transport has finished reading +/// the staged copy, so it goes back to the device allocator. +/// +/// CkCallback(CkCallbackFn, param) records the PE that built it and routes back +/// there however far the buffer travelled, so this runs on the sending PE -- +/// the same context the old finishGPUSend entry method ran in. `param` is the +/// staged pointer itself, which is why nothing has to be looked up here. +static void gpuMigrateStagedFree(void* param, void* msg) +{ + hapiCheck(hapiFree(param)); +} + +void CkLocMgr::sendGPUMsg(CmiUInt8 id) +{ + auto it = sendGPUBuffers.find(id); + CmiAssert(it != sendGPUBuffers.end()); + const GPUMigrateData gpuData = it->second; // by value: the entry goes now + sendGPUBuffers.erase(it); + + // The staged block has to outlive this send -- the transport may still be + // reading it -- but the zerocopy layer already reports exactly that, through + // the buffer's source callback. Hand it ownership rather than keeping the + // entry alive and waiting for the destination to send an ack back. + thisProxy[gpuData.toPe].immigrateGPU( + id, (int)gpuData.size, + CkDeviceBuffer(gpuData.data, CkCallback(gpuMigrateStagedFree, gpuData.data))); +} + +/// Post variant: supply the device buffer the payload should land in. The +/// transport writes into it and then runs the delivery variant below. +/// +/// The buffer is recorded here under the element's id rather than recovered +/// from the delivery variant's `data` argument: for a non-SDAG entry method +/// the generated dispatch passes the posted pointer only on the zerocopy leg, +/// and the delivery leg -- the one that runs once the transfer has landed -- +/// sees a null. Keying off the id is also what an application would do, since +/// it is the one that knows where it posted. +void CkLocMgr::immigrateGPU(CmiUInt8& id, int& size, char*& data, + CkDeviceBufferPost* post) +{ + void* landing = nullptr; + hapiCheck(hapiMalloc(&landing, (size_t)size)); + data = (char*)landing; + postedDeviceBuffers[id] = landing; +} + +void CkLocMgr::immigrateGPU(CmiUInt8 id, int size, char* data) +{ + // Only now has the transfer landed. The buffer moves from posted to + // received here and not in the post method above, because immigrate() reads + // receivedDeviceMsgs to decide whether the device half is ready -- seeing a + // merely posted buffer there would let it unpack from memory the transport + // has not written yet. + auto posted = postedDeviceBuffers.find(id); + CmiAssert(posted != postedDeviceBuffers.end()); + receivedDeviceMsgs[id] = posted->second; + postedDeviceBuffers.erase(posted); + + // If the host message beat us here it is parked; this was the missing half. + auto host = bufferedHostMigrateMsgs.find(id); + if (host != bufferedHostMigrateMsgs.end()) + { + CkArrayElementMigrateMessage* msg = host->second; + bufferedHostMigrateMsgs.erase(host); + immigrateWithDevice(msg); + } +} +#endif + void CkLocMgr::restore(const CkArrayIndex& idx, CmiUInt8 id, PUP::er& p) { insertID(idx, id); diff --git a/src/ck-core/cklocation.ci b/src/ck-core/cklocation.ci index db8ea2643a..d65b1d4397 100644 --- a/src/ck-core/cklocation.ci +++ b/src/ck-core/cklocation.ci @@ -1,3 +1,5 @@ +#include "conv-mach-opt.h" + module CkLocation { extern module CkMarshall; @@ -14,6 +16,13 @@ module CkLocation { group [migratable] CkLocMgr { entry CkLocMgr(CkArrayOptions opts); entry [expedited] void immigrate(CkArrayElementMigrateMessage *msg); +#ifdef CMK_CUDA + // The device half of a migration. nocopydevice lets the existing device + // zerocopy layer pick the transport (same-process copy, CUDA IPC, or + // device RDMA) rather than this path choosing one. + entry [expedited] void sendGPUMsg(CmiUInt8 id); + entry [expedited] void immigrateGPU(CmiUInt8 id, int size, nocopydevice char data[size]); +#endif entry [expedited] void requestLocation(const CkArrayIndex& idx, int peToTell); entry [expedited] void updateLocation(const CkArrayIndex& idx, const CkLocEntry& e); entry void reclaimRemote(const CkArrayIndex& idx, int deletedOnPe); diff --git a/src/ck-core/cklocation.h b/src/ck-core/cklocation.h index 380db7570a..eb03ea967f 100644 --- a/src/ck-core/cklocation.h +++ b/src/ck-core/cklocation.h @@ -92,13 +92,15 @@ class CkArrayElementMigrateMessage : public CMessage_CkArrayElementMigrateMessag { public: CkArrayElementMigrateMessage(CkArrayIndex idx_, CmiUInt8 id_, bool ignoreArrival_, - int length_, int nManagers_, int epoch_) + int length_, int nManagers_, int epoch_, + bool hasGPUMsg_ = false) : idx(idx_), id(id_), ignoreArrival(ignoreArrival_), length(length_), nManagers(nManagers_), - epoch(epoch_) + epoch(epoch_), + hasGPUMsg(hasGPUMsg_) { } @@ -108,6 +110,10 @@ class CkArrayElementMigrateMessage : public CMessage_CkArrayElementMigrateMessag int length; // Size in bytes of the packed data int nManagers; // Number of associated array managers int epoch; + // Whether a device payload travels alongside this message, on its own + // device-zerocopy send. The two arrive independently and in either order; + // whichever gets there first is buffered until the other lands. + bool hasGPUMsg; char* packData; }; @@ -220,6 +226,24 @@ CkpvExtern(int, CkSaveRestorePrefetch); #include "ckmigratable.h" +#if CMK_CUDA +/// A device migration payload staged on the source, waiting for the +/// destination to read it. The buffer is a copy of the element's device state, +/// not the state itself, so the element can be destroyed as soon as the pack +/// finishes; only this buffer has to outlive the transfer. +class GPUMigrateData +{ +public: + int toPe; + size_t size; + void* data; + + GPUMigrateData() : toPe(-1), size(0), data(nullptr) {} + GPUMigrateData(int toPe_, size_t size_, void* data_) + : toPe(toPe_), size(size_), data(data_) {} +}; +#endif + /********************** CkLocMgr ********************/ /// A tiny class for detecting heap corruption class CkMagicNumber_impl @@ -418,6 +442,28 @@ class CkLocMgr : public IrrGroup // Immigration messages which are waiting for all array managers to be ready std::list pendingImmigrate; +#if CMK_CUDA + // A migration with device state travels as two independent sends: the host + // message through immigrate(), and the device payload through immigrateGPU() + // on the device-zerocopy path. Neither ordering is guaranteed, so whichever + // arrives first waits here for the other. + // + // Source side: staged device payloads awaiting the destination's ack. + std::unordered_map sendGPUBuffers; + // Destination side: a host message that arrived before its device payload, + // and a device payload that arrived before its host message. + std::unordered_map bufferedHostMigrateMsgs; + // Landing buffers handed to the transport, before and after it has written + // them. Only an entry in the second means the device state is readable. + std::unordered_map postedDeviceBuffers; + std::unordered_map receivedDeviceMsgs; + // Source PE of a landed device payload, so the ack can be addressed once + // the unpack has consumed it. + + // Unpack a migration whose host message and device payload are both present. + void immigrateWithDevice(CkArrayElementMigrateMessage* msg); +#endif + // The mapping of index to ID is either done via compression or an explicit map, // depending on if the bounds of this array are compressible into a 64bit ID. CkArrayIndex bounds; @@ -691,6 +737,18 @@ class CkLocMgr : public IrrGroup // Communication: void immigrate(CkArrayElementMigrateMessage* msg); +#if CMK_CUDA + // Send a staged device payload to its destination. Runs as its own entry + // method rather than inline in emigrate so the send is attributed to the + // runtime and not to whichever element last ran on this PE. + void sendGPUMsg(CmiUInt8 id); + // Device-zerocopy receive: the post variant supplies the landing buffer, + // the second runs once the transfer has landed. + void immigrateGPU(CmiUInt8& id, int& size, char*& data, + CkDeviceBufferPost* post); + void immigrateGPU(CmiUInt8 id, int size, char* data); + // Destination's ack: the staged payload has been read and can be released. +#endif void requestLocation(CmiUInt8 id); void requestLocation(const CkArrayIndex& idx); bool requestLocation(const CkArrayIndex& idx, int peToTell); diff --git a/src/ck-core/cklocrec.h b/src/ck-core/cklocrec.h index 8528aafcd2..05272664ff 100644 --- a/src/ck-core/cklocrec.h +++ b/src/ck-core/cklocrec.h @@ -49,6 +49,8 @@ class CkLocRec { void stopTiming(int ignore_running=0); void setObjTime(double cputime); double getObjTime(); + void setObjGPUTime(double gputime); + double getObjGPUTime(); void *getObjUserData(int idx); #else inline void startTiming(int ignore_running=0) { } @@ -70,6 +72,7 @@ class CkLocRec { void recvMigrate(int dest); void setMigratable(int migratable); /// set migratable void setPupSize(size_t obj_pup_size); + void setGPUPupSize(size_t obj_gpu_pup_size); void AsyncMigrate(bool use); bool isAsyncMigrate() { return asyncMigrate; } void ReadyMigrate(bool ready) { readyMigrate = ready; } ///called from user diff --git a/src/ck-core/ckmigratable.h b/src/ck-core/ckmigratable.h index 7cec2884ad..6f78fb8c4a 100644 --- a/src/ck-core/ckmigratable.h +++ b/src/ck-core/ckmigratable.h @@ -103,6 +103,7 @@ class CkMigratable : public Chare { void AtSync(int waitForMigration=1) { ResumeFromSync();} void setMigratable(int migratable) { } void setPupSize(size_t obj_pup_size) { } + void setGPUPupSize(size_t obj_gpu_pup_size) { } public: void ckFinishConstruction(int epoch) { } #endif diff --git a/src/ck-ldb/BaseLB.h b/src/ck-ldb/BaseLB.h index e7b6683d7f..860414b850 100644 --- a/src/ck-ldb/BaseLB.h +++ b/src/ck-ldb/BaseLB.h @@ -53,9 +53,21 @@ class BaseLB: public CBase_BaseLB // double utilization; int pe; // processor id bool available; +#if CMK_CUDA + /// GPU this PE is mapped to, as (physical node id << 32 | device index), + /// so that the same physical device is one value across every process that + /// shares it. -1 when this PE has no GPU. PEs with equal values form the + /// GPU group a GPU-aware strategy balances across. + uint64_t gpu_device_id; + /// SM count of that GPU, 0 if it could not be queried. + int gpu_total_sms; +#endif ProcStats(): n_objs(0), pe_speed(1), total_walltime(0.0), idletime(0.0), #if CMK_LB_CPUTIMER total_cputime(0.0), bg_cputime(0.0), +#endif +#if CMK_CUDA + gpu_device_id((uint64_t)-1), gpu_total_sms(0), #endif bg_walltime(0.0), pe(-1), available(true) {} inline void clearBgLoad() { @@ -78,7 +90,11 @@ class BaseLB: public CBase_BaseLB double dummy; p|dummy; // for old format with utilization } p|available; p|n_objs; - if (_lb_args.lbversion()>=2) p|pe; + if (_lb_args.lbversion()>=2) p|pe; +#if CMK_CUDA + p|gpu_device_id; + p|gpu_total_sms; +#endif } }; diff --git a/src/ck-ldb/CentralLB.C b/src/ck-ldb/CentralLB.C index 23c46a3075..754f186f06 100644 --- a/src/ck-ldb/CentralLB.C +++ b/src/ck-ldb/CentralLB.C @@ -11,6 +11,12 @@ #include "CentralLB.h" #include "LBSimulation.h" +#if CMK_CUDA +#include "hapi.h" +#include "gpumanager.h" +CsvExtern(GPUManager, gpu_manager); +#endif + #define DEBUGF(x) // CmiPrintf x; #define DEBUG(x) // x; @@ -143,9 +149,36 @@ void CentralLB::InvokeLB() MigrationDone(0); return; } - { - thisProxy [CkMyPe()].ProcessAtSync(); - } + +#if CMK_CUDA + // Reduce the kernel timeline to one scalar load per object, while the records + // are still local -- only those scalars are sent to the central LB. + // + // Built by the LAST PE of the process to get here, not the first: InvokeLB + // runs when this PE's own objects are at AtSync, but the records are shared + // by the whole process, and a drain by the first arrival would drop every + // kernel the other PEs' objects were still finishing. The earlier arrivals + // continue from gpuLoadsReady. + if (!hapiCuptiArrive((uint64_t)step(), CkNodeSize(CkMyNode()))) return; + hapiPrepareCuptiLoads(); + const int first = CkNodeFirst(CkMyNode()); + for (int r = 0; r < CkNodeSize(CkMyNode()); r++) + thisProxy[first + r].gpuLoadsReady(); +#else + thisProxy [CkMyPe()].ProcessAtSync(); +#endif +#endif +} + +// The round's GPU loads exist: copy this PE's share out, then carry on into +// the stats path exactly as the non-CUDA build does. +void CentralLB::gpuLoadsReady() +{ +#if CMK_LBDB_ON +#if CMK_CUDA + lbmgr->SetObjGPULoad(CsvAccess(gpu_manager).cupti_obj_norm_load_); +#endif + thisProxy [CkMyPe()].ProcessAtSync(); #endif } @@ -305,6 +338,10 @@ void CentralLB::BuildStatsMsg() #else msg->pe_speed = myspeed; #endif +#if CMK_CUDA + msg->gpu_device_id = hapiMyDevice(); + msg->gpu_total_sms = hapiMyDeviceTotalSMs(); +#endif DEBUGF(("Processor %d Total time (wall,cpu) = %f %f Idle = %f Bg = %f %f\n", CkMyPe(),msg->total_walltime,msg->total_cputime,msg->idletime,msg->bg_walltime,msg->bg_cputime)); @@ -435,6 +472,10 @@ void CentralLB::depositData(CLBStatsMsg *m) procStat.bg_cputime = m->bg_cputime; #endif procStat.pe_speed = m->pe_speed; +#if CMK_CUDA + procStat.gpu_device_id = m->gpu_device_id; + procStat.gpu_total_sms = m->gpu_total_sms; +#endif //procStat.utilization = 1.0; procStat.available = true; @@ -510,6 +551,10 @@ void CentralLB::ReceiveStats(CkMarshalledCLBStatsMessage &&msg) procStat.bg_cputime = m->bg_cputime; #endif procStat.pe_speed = m->pe_speed; +#if CMK_CUDA + procStat.gpu_device_id = m->gpu_device_id; + procStat.gpu_total_sms = m->gpu_total_sms; +#endif //procStat.utilization = 1.0; procStat.available = true; procStat.n_objs = msg_n_objs; @@ -999,11 +1044,27 @@ void CentralLB::ProcessReceiveMigration() future_migrates_expected = 0; for(i=0; i < m->n_moves; i++) { MigrateInfo& move = m->moves[i]; - #if CMK_GLOBAL_LOCATION_UPDATE - UpdateLocation(move); - #endif const int me = CkMyPe(); - if (move.from_pe == me && move.to_pe != me) { + const bool iAmSource = (move.from_pe == me && move.to_pe != me); +#if CMK_GLOBAL_LOCATION_UPDATE + // Every PE but the one the object is leaving learns the new location here, + // before the move is acted on. + // + // The source is the exception. Its entry still says "the element is on me", + // and emigrate() -> CkLocCache::recordEmigration is what turns that into + // "it is on the destination" -- asserting on the way that it was here to + // begin with. Updating it first makes that assert fail on the very first + // migration of any centralized balancer. + // + // The destination is NOT an exception, even though it will learn the + // location for itself when the element lands in createLocal. Between this + // decision and that arrival it would otherwise still believe the element is + // on the source, and address messages there; the source has already let it + // go, so each one takes an extra hop -- which is precisely what + // CkLocMgr::multiHop asserts against in this mode. + if (!iAmSource) UpdateLocation(move); +#endif + if (iAmSource) { #if CMK_DRONE_MODE int to_pe_rank0 = CMK_RANK_0(move.to_pe); if(move.from_pe == to_pe_rank0) continue; @@ -1650,6 +1711,10 @@ CLBStatsMsg::~CLBStatsMsg() { void CLBStatsMsg::pup(PUP::er &p) { p|from_pe; p|pe_speed; +#if CMK_CUDA + p|gpu_device_id; + p|gpu_total_sms; +#endif p|total_walltime; p|idletime; #if defined(TEMP_LDB) diff --git a/src/ck-ldb/CentralLB.ci b/src/ck-ldb/CentralLB.ci index 59694ce685..679e7d6956 100644 --- a/src/ck-ldb/CentralLB.ci +++ b/src/ck-ldb/CentralLB.ci @@ -12,8 +12,9 @@ readonly CkGroupID loadbalancer; initnode void lbinit(void); group [migratable] CentralLB : BaseLB { - entry void CentralLB(const CkLBOptions &); + entry void CentralLB(const CkLBOptions &); entry void ProcessAtSync(void); + entry void gpuLoadsReady(void); entry [reductiontarget] void SendStats(); entry void ReceiveStats(CkMarshalledCLBStatsMessage data); entry void ReceiveStatsViaTree(CkMarshalledCLBStatsMessage data); diff --git a/src/ck-ldb/CentralLB.h b/src/ck-ldb/CentralLB.h index bbc9a5c140..6ae8757245 100644 --- a/src/ck-ldb/CentralLB.h +++ b/src/ck-ldb/CentralLB.h @@ -97,6 +97,9 @@ class CentralLB : public CBase_CentralLB inline void setConcurrent(bool c) { concurrent = c; } void InvokeLB(); // Everything is at the PE barrier + // Second half of InvokeLB on a CUDA build: runs once the process's + // per-object GPU loads for this round exist. + void gpuLoadsReady(); void ProcessAtSync(void); // Receive a message from AtSync to avoid // making projections output look funny void SendStats(); @@ -283,6 +286,12 @@ class CLBStatsMsg { int from_pe; int pe_speed; +#if CMK_CUDA + // GPU this PE is mapped to and that device's SM count; see + // BaseLB::ProcStats, into which these are copied on arrival. + uint64_t gpu_device_id; + int gpu_total_sms; +#endif LBRealType total_walltime; LBRealType idletime; LBRealType bg_walltime; @@ -298,7 +307,11 @@ class CLBStatsMsg { public: CLBStatsMsg(int osz, int csz); - CLBStatsMsg(): from_pe(0), pe_speed(0), total_walltime(0.0), idletime(0.0), + CLBStatsMsg(): from_pe(0), pe_speed(0), +#if CMK_CUDA + gpu_device_id((uint64_t)-1), gpu_total_sms(0), +#endif + total_walltime(0.0), idletime(0.0), bg_walltime(0.0), #if defined(TEMP_LDB) pe_temp(1.0), diff --git a/src/ck-ldb/CommonLBs.ci b/src/ck-ldb/CommonLBs.ci index 436ffa6729..fddf562ddc 100644 --- a/src/ck-ldb/CommonLBs.ci +++ b/src/ck-ldb/CommonLBs.ci @@ -5,6 +5,7 @@ module CommonLBs { extern module DistributedLB; extern module MetisLB; extern module RecBipartLB; + extern module GreedyRefineCentralGPULB; initnode void initCommonLBs(void); }; diff --git a/src/ck-ldb/EveryLB.ci b/src/ck-ldb/EveryLB.ci index a634d9c9e0..bf00ca3a30 100644 --- a/src/ck-ldb/EveryLB.ci +++ b/src/ck-ldb/EveryLB.ci @@ -5,6 +5,7 @@ module EveryLB { extern module DistributedLB; extern module MetisLB; extern module RecBipartLB; + extern module GreedyRefineCentralGPULB; initnode void initEveryLB(void); }; diff --git a/src/ck-ldb/GreedyRefineCentralGPULB.C b/src/ck-ldb/GreedyRefineCentralGPULB.C new file mode 100644 index 0000000000..519958feed --- /dev/null +++ b/src/ck-ldb/GreedyRefineCentralGPULB.C @@ -0,0 +1,830 @@ +/** + * \addtogroup CkLdb +*/ +/*@{*/ + +/** + * Two-level GPU-aware greedy-refine load balancer. + * + * Cross-group (per-GPU/per-process) balancing uses GPU load (gpuTime); + * within-group (per-PE) balancing uses CPU load (wallTime). See the header + * for the full description. Derived from GreedyRefineCentralLB + * (jjgalvez@illinois.edu). +*/ + +#include "charm++.h" +#include "ckgraph.h" +#include "GreedyRefineCentralGPULB.h" + +#include +#include +#include +#include +#if CMK_CUDA +#include +#endif + +extern int quietModeRequested; + +// a solution is feasible if num migrations <= user-specified limit +// LOAD_MIG_BAL is used to control tradeoff between maxload and migrations +// when selecting solutions from the feasible set +#define LOAD_MIG_BAL 1.003 + +using namespace std; + +class GreedyRefineCentralGPULB::Solution { +public: + Solution() {} + Solution(int pe, double maxLoad, int nmoves) : pe(pe), max_load(maxLoad), migrations(nmoves) {} + int pe; // pe who produced this solution + float max_load; + int migrations; + + void pup(PUP::er &p) { + p|pe; + p|max_load; + p|migrations; + } +}; + +// custom heap to allow removal of processors from any position +class GreedyRefineCentralGPULB::PHeap { +public: + PHeap(int numpes) { + Q.reserve(numpes+1); + Q.push_back(NULL); // first element of the array is NULL + } + + void addProcessors(std::vector &procs, bool bgLoadZero, bool insert=true) { + for (int i=0; i < procs.size(); i++) { + GreedyRefineCentralGPULB::GProc &p = procs[i]; + if (p.available) { + p.load = p.bgload; + if (insert) { + Q.push_back(&p); + p.pos = Q.size()-1; + } + } + } + if (!bgLoadZero) buildMinHeap(); + } + + inline GreedyRefineCentralGPULB::GProc *top() const { + CkAssert(Q.size() > 1); + return Q[1]; + } + + inline void push(GreedyRefineCentralGPULB::GProc *p) { + Q.push_back(p); + p->pos = Q.size()-1; + siftUp(p->pos); + } + + inline GreedyRefineCentralGPULB::GProc *pop() { + if (Q.size() == 1) return NULL; + GreedyRefineCentralGPULB::GProc *retval; + if (Q.size() == 2) { + retval = Q[1]; + Q.pop_back(); + return retval; + } + retval = Q[1]; + Q[1] = Q.back(); + Q.pop_back(); + Q[1]->pos = 1; + siftDown(1); + return retval; + } + + // remove processor from any position in the heap + void remove(GreedyRefineCentralGPULB::GProc *p) { + int pos = p->pos; + if ((Q.size() == 2) || (pos == Q.size()-1)) return Q.pop_back(); + if (pos == 1) { pop(); return; } + Q[pos] = Q.back(); + Q.pop_back(); + Q[pos]->pos = pos; + if (Q[pos/2]->load > Q[pos]->load) siftUp(pos); + else siftDown(pos); + } + + inline void clear() { + Q.clear(); + Q.push_back(NULL); + } + +private: + + void min_heapify(int i) { + const int left = 2*i; + const int right = 2*i + 1; + int smallest = i; + if ((left < Q.size()) && (Q[left]->load < Q[smallest]->load)) smallest = left; + if ((right < Q.size()) && (Q[right]->load < Q[smallest]->load)) smallest = right; + if (smallest != i) { + swap(i,smallest); + Q[i]->pos = i; + Q[smallest]->pos = smallest; + min_heapify(smallest); + } + } + + void inline buildMinHeap() { + for (int i=Q.size()/2; i > 0; i--) min_heapify(i); + } + + inline void swap(int pos1, int pos2) { + GreedyRefineCentralGPULB::GProc *t = Q[pos1]; + Q[pos1] = Q[pos2]; + Q[pos2] = t; + } + + void siftUp(int pos) { + if (pos == 1) return; // reached root + int ppos = pos/2; + if (Q[ppos]->load > Q[pos]->load) { + swap(ppos,pos); + Q[ppos]->pos = ppos; + Q[pos]->pos = pos; + siftUp(ppos); + } + } + + inline int minChild(int pos) const { + int c1 = pos*2; + int c2 = pos*2 + 1; + if (c1 >= Q.size()) return -1; + if (c2 >= Q.size()) return c1; + if (Q[c1]->load < Q[c2]->load) return c1; + else return c2; + } + + void siftDown(int pos) { + int cpos = minChild(pos); + if (cpos == -1) return; + if (Q[pos]->load > Q[cpos]->load) { + swap(pos,cpos); + Q[cpos]->pos = cpos; + Q[pos]->pos = pos; + siftDown(cpos); + } + } + + std::vector Q; +}; + +CreateLBFunc_Def(GreedyRefineCentralGPULB, "Two-level GPU-aware greedy refinement-based algorithm") + +GreedyRefineCentralGPULB::GreedyRefineCentralGPULB(const CkLBOptions &opt): CBase_GreedyRefineCentralGPULB(opt), migrationTolerance(1.0) +{ + lbname = "GreedyRefineCentralGPULB"; + if ((CkMyPe() == 0) && !quietModeRequested) + CkPrintf("CharmLB> GreedyRefineCentralGPULB created.\n"); +#if CMK_HIP + // This balancer's cross-GPU stage balances on LDObjData::gpuTime, which is + // filled in from CUPTI activity records -- a CUDA-only facility with no HIP + // equivalent wired up yet. On HIP every object's gpuTime stays zero, so that + // stage sums zeros and only the within-group CPU pass does any work. Warn + // rather than balance silently on a dimension that is not being measured. + // Printed even under quiet mode: it reports degraded behaviour, not a banner. + if (CkMyPe() == 0) + CkPrintf("CharmLB> Warning: GPU load measurement on HIP is not implemented " + "yet. GreedyRefineCentralGPULB will see zero GPU load for every " + "object and balance on CPU load alone.\n"); +#endif + if (_lb_args.percentMovesAllowed() < 100) { + migrationTolerance = float(_lb_args.percentMovesAllowed())/100.0; + } + // A tolerance (the default, 1.1) says what to aim for, so one candidate is + // built. Asking for 0 or less asks for the search instead: every PE builds a + // candidate from a different (A,B) pair and receiveSolutions picks between + // them, at the cost of a reduction per balancing step. + concurrent = (_lb_args.greedyRefineTolerance() <= 0.0); +} + +GreedyRefineCentralGPULB::GreedyRefineCentralGPULB(CkMigrateMessage *m): CBase_GreedyRefineCentralGPULB(m), migrationTolerance(1.0) { + lbname = "GreedyRefineCentralGPULB"; + if (_lb_args.percentMovesAllowed() < 100) + migrationTolerance = float(_lb_args.percentMovesAllowed())/100.0; + concurrent = (_lb_args.greedyRefineTolerance() <= 0.0); +} + +// ------------------------------------------------ + +// regular greedy lb algorithm (CPU-only / non-CUDA path) +double GreedyRefineCentralGPULB::greedyLB(const std::vector &pobjs, + GreedyRefineCentralGPULB::PHeap &procHeap, + const BaseLB::LDStats *stats) const +{ + double max_load = 0; + int nmoves = 0; + for (int i=0; i < pobjs.size(); i++) { + const GreedyRefineCentralGPULB::GObj *obj = pobjs[i]; + GreedyRefineCentralGPULB::GProc *p = procHeap.pop(); // least loaded processor + // update processor load + p->load += (obj->load / p->speed); + procHeap.push(p); + + if (p->id != obj->oldPE) nmoves++; + if (p->load > max_load) max_load = p->load; + } + + if ((CkMyPe() == cur_ld_balancer+1) && (_lb_args.debug() > 1)) { + CkPrintf("[%d] %f : Greedy strategy nmoves=%d, max_load=%f\n", CkMyPe(), + CkWallTimer() - strategyStartTime, nmoves, max_load); + } + return max_load; +} + +// ----------------------------------------------- +#if __DEBUG_GREEDY_REFINE_GPU_ +#include +void GreedyRefineCentralGPULB::dumpObjLoads(std::vector &objs) { + std::ofstream outfile("objloads.txt"); + outfile << objs.size() << std::endl; + for (int i=0; i < objs.size(); i++) { + GreedyRefineCentralGPULB::GObj &obj = objs[i]; + if ((i > 0) && (i % 100 == 0)) outfile << obj.load << std::endl; + else outfile << obj.load << " "; + } + outfile.close(); +} +void GreedyRefineCentralGPULB::dumpProcLoads(std::vector &procs) { + std::ofstream outfile("proc_bg_loads.txt"); + outfile << procs.size() << std::endl; + for (int i=0; i < procs.size(); i++) { + GreedyRefineCentralGPULB::GProc &p = procs[i]; + if ((i > 0) && (i % 100 == 0)) outfile << p.load << std::endl; + else outfile << p.load << " "; + } + outfile.close(); +} +#endif + +double GreedyRefineCentralGPULB::fillData(LDStats *stats, + std::vector &objs, + std::vector &pobjs, + std::vector &procs, + PHeap &procHeap) +{ + const int n_pes = stats->nprocs(); + // Walk every object, not just the migratable ones: the loop below indexes + // stats->objData[i] directly, and non-migratable objects must still be charged + // as background load on their PE. + // + // n_migrateobjs is a *count* of migratable objects (CentralLB.C:485 computes it + // with count_if over the whole array), not an upper bound on their indices -- + // objData is filled in per-PE registration order and is never partitioned, so + // migratable and non-migratable objects interleave. Using the count as a loop + // bound silently truncates objData: the tail is neither balanced nor charged as + // background, so its load vanishes from the objective entirely. Latent whenever + // every chare is migratable (n_migrateobjs == objData.size()), which is why + // pic2d and jacobi2d never exposed it. + const int n_objs = stats->objData.size(); + // most of these variables are just for printing stats when _lb_args.debug() + int unmigratableObjs = 0; + availablePes = 0; totalObjLoad = 0; + double minBGLoad = DBL_MAX; double avgBGLoad = 0; double maxBGLoad = 0; + double minSpeed = DBL_MAX; double maxSpeed = 0; double avgSpeed = 0; + double minOload = DBL_MAX; double maxOload = 0; + + for (int pe=0; pe < n_pes; pe++) { + GreedyRefineCentralGPULB::GProc &p = procs[pe]; + p.id = pe; + p.available = stats->procs[pe].available; + p.speed = stats->procs[pe].pe_speed; + if (p.available) { + availablePes++; + // CMK_CUDA is always defined -- 0 in a non-CUDA build -- so this must + // test its value, not whether it exists. + #if !CMK_CUDA + p.bgload = stats->procs[pe].bg_walltime; + if (p.bgload > maxBGLoad) maxBGLoad = p.bgload; + #else + // bgload is the background GPU load (aggregated per group below). + // cpuLoad is the background CPU/wall load, used for within-group + // PE-level balancing; seed it with this PE's background walltime. + p.bgload = 0.0; + p.bg_walltime = stats->procs[pe].bg_walltime; + p.cpuLoad = stats->procs[pe].bg_walltime; + #endif + if (_lb_args.debug() > 1) { + double &speed = stats->procs[pe].pe_speed; + if (speed < minSpeed) minSpeed = speed; + if (speed > maxSpeed) maxSpeed = speed; + avgSpeed += speed; + } + } + } + if (!availablePes) CkAbort("GreedyRefineCentralGPULB: No available processors\n"); + + for (int i=0; i < n_objs; i++) { + LDObjData &oData = stats->objData[i]; + GreedyRefineCentralGPULB::GObj &obj = objs[i]; + int pe = stats->from_proc[i]; + obj.id = i; + obj.oldPE = pe; + CkAssert(pe >= 0 && pe <= n_pes); + if (pe == n_pes) obj.oldPE = -1; // this can happen in HybridLB if object comes from outside group. mark oldPE as -1 in this situation + if (!oData.migratable) { + CkAssert(pe < n_pes); + unmigratableObjs++; + GreedyRefineCentralGPULB::GProc &p = procs[pe]; + if (!p.available) + CkAbort("GreedyRefineCentralGPULB: nonmigratable object on unavailable processor\n"); +#if CMK_CUDA + // GPU load is background for the PE's GPU group; CPU load is background + // for the PE itself. + p.bgload += oData.gpuTime; + p.cpuLoad += oData.wallTime; + if (p.bgload > maxBGLoad) maxBGLoad = p.bgload; +#else + double nmObjLoad = oData.wallTime; + p.bgload += nmObjLoad; // take non-migratable object load as background load + if (p.bgload > maxBGLoad) maxBGLoad = p.bgload; +#endif + } else { +#if CMK_CUDA + // load -> GPU work (cross-group objective); cpuLoad -> CPU work (intra-group) + obj.load = oData.gpuTime; + obj.cpuLoad = oData.wallTime; +#else + obj.load = oData.wallTime * stats->procs[pe].pe_speed; + obj.cpuLoad = oData.wallTime; +#endif + pobjs.push_back(&obj); + totalObjLoad += obj.load; + if (_lb_args.debug() > 1) { + if (obj.load < minOload) minOload = obj.load; + if (obj.load > maxOload) maxOload = obj.load; + } + } + } + + procHeap.addProcessors(procs, (maxBGLoad <= 0.001), true); + + // ---- print some stats ---- + if ((_lb_args.debug() > 1) && (!concurrent || (CkMyPe() == cur_ld_balancer))) { + for (int pe=0; pe < n_pes; pe++) { + GreedyRefineCentralGPULB::GProc &p = procs[pe]; + if (!p.available) continue; + if (p.bgload < minBGLoad) minBGLoad = p.bgload; + avgBGLoad += p.bgload; + } + CkPrintf("[%d] GreedyRefineCentralGPULB: num pes=%d, num objs=%d (%d migratable, %d background)\n", + CkMyPe(), n_pes, n_objs, n_objs - unmigratableObjs, unmigratableObjs); + CkPrintf("[%d] Unavailable processors=%d, Unmigratable objs=%d\n", CkMyPe(), n_pes - availablePes, unmigratableObjs); + CkPrintf("[%d] min_bgload=%f mean_bgload=%f max_bgload=%f\n", CkMyPe(), minBGLoad, (avgBGLoad / availablePes), maxBGLoad); + CkPrintf("[%d] min_oload=%f mean_oload=%f max_oload=%f\n", CkMyPe(), minOload, (totalObjLoad / (n_objs - unmigratableObjs)), maxOload); + CkPrintf("[%d] min_speed=%f mean_speed=%f max_speed=%f\n", CkMyPe(), minSpeed, (avgSpeed / availablePes), maxSpeed); + } + + return maxBGLoad; +} + +static const float Avals[] = {1.0, 1.005, 1.01, 1.015, 1.02, 1.03, 1.04, 1.05, 1.06, 1.07, 1.08, 1.16, 1.20, 1.30}; +static const float Bvals[] = {FLT_MAX, 1.0, 1.05, 1.1, 1.2, 1.3, 1.4, 1.5, 1.6, 1.7, 1.8, 1.9, 2.0, 2.1, 2.2, 2.3}; +#define Avals_len 14 +#define Bvals_len 16 +#define NUM_SOLUTIONS Avals_len*Bvals_len+1 +static void getGreedyRefineParams(int rank, float &A, float &B) { + if (rank == 0) { A = 0; B = -1; return; } // causes PE0 to run regular greedy + rank--; + int x = rank / Bvals_len; + if (x >= Avals_len) { + A = B = -1; + } else { + A = Avals[x]; + B = Bvals[rank % Bvals_len]; + } +} + +void GreedyRefineCentralGPULB::sendSolution(double maxLoad, int migrations) +{ + // gather results in central PE, who will decide which solution is the best + // only the objective values of the solutions are sent, not the whole solutions + + GreedyRefineCentralGPULB::Solution sol(CkMyPe(), maxLoad, migrations); + size_t buf_size = sizeof(GreedyRefineCentralGPULB::Solution); + void *buffer = malloc(buf_size); + PUP::toMem pd(buffer); + pd|sol; + + CkCallback cb(CkIndex_GreedyRefineCentralGPULB::receiveSolutions((CkReductionMsg*)NULL), thisProxy[cur_ld_balancer]); + contribute(buf_size, buffer, CkReduction::set, cb); + + if ((_lb_args.debug() > 1) && (CkMyPe() == cur_ld_balancer)) { + CkPrintf("[%d] %f : Called gather/reduction\n", CkMyPe(), CkWallTimer() - strategyStartTime); + } + + free(buffer); +} + +void GreedyRefineCentralGPULB::work(LDStats *stats) +{ + strategyStartTime = CkWallTimer(); + // Greedy's max host load, the CPU twin of greedyMaxLoad. Set in the GPU path + // only; the CPU-only path below has a single dimension and uses greedyMaxLoad. + double greedyMaxCpuLoad = 0.0; + // maxCpuLoad itself is scoped to the GPU branch; carry it out for the report. + double maxCpuLoadReported = 0.0; + + // A is the tolerance: the refine target is plain greedy's max load times A + // (M *= A below), and an object stays where it is while its PE is under that. + // B is the stay-put band against the lightest PE; FLT_MAX means the target is + // the only thing holding an object back, which is the documented behaviour. + float A = 1.001, B = FLT_MAX; // Use A=0, B=-1 to imitate regular Greedy (ignore migrations) + if (_lb_args.greedyRefineTolerance() > 0.0) + A = (float)_lb_args.greedyRefineTolerance(); + if (concurrent) { + getGreedyRefineParams(CkMyPe(), A, B); + if (A < 0) { + sendSolution(-1,-1); // send empty response to PE0 + return; + } + } + + const int n_pes = stats->nprocs(); + totalObjs = stats->n_migrateobjs; + + // Sized to every object, because fillData indexes this by LDStats object index. + // totalObjs stays the migratable count -- it bounds how many objects may move + // (migrationsAllowed below), which is a different quantity from how many exist. + std::vector objs(stats->objData.size()); + // will sort pobjs instead of objs (faster swapping). will only contain pointers + // to migratable objects + std::vector pobjs; + pobjs.reserve(totalObjs); + + std::vector procs(n_pes); + PHeap procHeap(n_pes); + + // fill data structures used by algorithm + double maxLoad = fillData(stats, objs, pobjs, procs, procHeap); + + // ------------ apply greedy refine algorithm -------------- + + // Sort by GPU load (GObj::load) — the cross-group objective. + std::sort(pobjs.begin(), pobjs.end(), GreedyRefineCentralGPULB::ObjLoadGreater()); + + int nmoves = 0; + double greedyMaxLoad = 0; + +#if CMK_CUDA + // ---- Two-level GPU-aware path ---- + // + // Outer: balance GPU load (GObj::load) across GPU groups. A GPU group is the + // set of available PEs that share a gpu_device_id (the PEs of one + // process bound to one GPU). + // Inner: within the chosen group, balance CPU load (GObj::cpuLoad / GProc::cpuLoad) + // across the group's PEs. + + struct GPUGrp { + uint64_t gpu_id; + double load; // aggregate GPU load across PEs in this group + std::vector peIds; // indices into procs[] + }; + + std::vector gpuGroups; + std::unordered_map gpuIdToIdx; + + for (int pe = 0; pe < n_pes; pe++) { + GreedyRefineCentralGPULB::GProc &p = procs[pe]; + if (!p.available) continue; + uint64_t devId = stats->procs[pe].gpu_device_id; + + auto it = gpuIdToIdx.find(devId); + if (it == gpuIdToIdx.end()) { + gpuIdToIdx[devId] = gpuGroups.size(); + GPUGrp g; + g.gpu_id = devId; + g.load = p.bgload; // background GPU load + g.peIds.push_back(pe); + gpuGroups.push_back(std::move(g)); + } else { + gpuGroups[it->second].load += p.bgload; + gpuGroups[it->second].peIds.push_back(pe); + } + } + int nGroups = gpuGroups.size(); + + if ((_lb_args.debug() > 1) && (CkMyPe() == cur_ld_balancer)) { + CkPrintf("[%d] GreedyRefineCentralGPULB: %d GPU group(s), %d available PEs, %d migratable objs\n", + CkMyPe(), nGroups, availablePes, (int)pobjs.size()); + for (auto &g : gpuGroups) + CkPrintf("[%d] GPU %llu: %d PEs, bgGpuLoad=%.6f\n", + CkMyPe(), (unsigned long long)g.gpu_id, (int)g.peIds.size(), g.load); + } + + // What a step is actually made of, per PE, so the objective can be checked + // against reality: if the measured step time is not tracked by whichever of + // the CPU and GPU terms the balancer minimises, then balancing that term + // cannot make the step shorter no matter how well it succeeds. + if ((_lb_args.debug() > 1) && (CkMyPe() == cur_ld_balancer)) { + std::vector peCpu(n_pes, 0.0), peGpu(n_pes, 0.0); + std::vector peObjs(n_pes, 0); + for (size_t i = 0; i < stats->objData.size(); i++) { + const int pe = stats->from_proc[i]; + if (pe < 0 || pe >= n_pes) continue; + peCpu[pe] += stats->objData[i].wallTime; + peGpu[pe] += stats->objData[i].gpuTime; + peObjs[pe]++; + } + CkPrintf("[%d] --- step composition (s) ---\n", CkMyPe()); + for (int pe = 0; pe < n_pes; pe++) { + const double tot = stats->procs[pe].total_walltime; + const double idle = stats->procs[pe].idletime; + CkPrintf("[%d] PE %d objs=%d total=%.6f idle=%.6f (%.0f%%) bg=%.6f " + "objCPU=%.6f objGPU=%.6f busy=%.6f\n", + CkMyPe(), pe, peObjs[pe], tot, idle, + tot > 0 ? 100.0 * idle / tot : 0.0, + stats->procs[pe].bg_walltime, peCpu[pe], peGpu[pe], tot - idle); + } + } + + // --- Greedy preprocessing at GPU-group level to establish target M --- + double M = 0; + { + std::vector grpLoad(nGroups); + for (int gi = 0; gi < nGroups; gi++) grpLoad[gi] = gpuGroups[gi].load; + + for (int i = 0; i < (int)pobjs.size(); i++) { + int lightest = 0; + for (int gi = 1; gi < nGroups; gi++) + if (grpLoad[gi] < grpLoad[lightest]) lightest = gi; + grpLoad[lightest] += pobjs[i]->load; + if (grpLoad[lightest] > M) M = grpLoad[lightest]; + } + greedyMaxLoad = M; + } + M *= A; + + // --- The same preprocessing at PE level, for the CPU target Mcpu --- + // + // The inner decision balances host load across the PEs of a group, and until + // now had no ceiling at all: only B held an object in place, so B = FLT_MAX + // meant nothing ever moved whatever the tolerance said. Give it the twin of + // M so the tolerance governs both decisions in the same terms. + // + // On a copy: procs[].cpuLoad is the real background load the assignment pass + // starts from, and unlike the GPU groups it is never reset. + double Mcpu = 0; + { + std::vector peLoad(n_pes); + for (int pe = 0; pe < n_pes; pe++) + peLoad[pe] = procs[pe].available ? procs[pe].cpuLoad : 0.0; + for (int i = 0; i < (int)pobjs.size(); i++) { + int lightest = -1; + for (int pe = 0; pe < n_pes; pe++) { + if (!procs[pe].available) continue; + if (lightest < 0 || peLoad[pe] < peLoad[lightest]) lightest = pe; + } + if (lightest < 0) break; // no available PE: nothing to target + peLoad[lightest] += pobjs[i]->cpuLoad / procs[lightest].speed; + if (peLoad[lightest] > Mcpu) Mcpu = peLoad[lightest]; + } + greedyMaxCpuLoad = Mcpu; + } + Mcpu *= A; + + // Reset GPU group loads back to bg-only for the real assignment pass. + for (int gi = 0; gi < nGroups; gi++) { + gpuGroups[gi].load = 0; + for (int pe : gpuGroups[gi].peIds) + gpuGroups[gi].load += procs[pe].bgload; + } + + // Reverse map: PE index -> GPU group index + std::unordered_map peToGrpIdx; + for (int gi = 0; gi < nGroups; gi++) + for (int pe : gpuGroups[gi].peIds) + peToGrpIdx[pe] = gi; + + // Two dimensions are tracked separately. maxLoad follows the GPU-group + // aggregate and keeps driving the refine target M below; maxCpuLoad follows + // the busiest PE's CPU load. Only their combination is reported as the + // solution objective -- see the note after the assignment loop. + maxLoad = 0; + for (int gi = 0; gi < nGroups; gi++) + if (gpuGroups[gi].load > maxLoad) maxLoad = gpuGroups[gi].load; + + double maxCpuLoad = 0; + for (int pe = 0; pe < n_pes; pe++) + if (procs[pe].available && procs[pe].cpuLoad > maxCpuLoad) + maxCpuLoad = procs[pe].cpuLoad; + + // --- Main assignment loop --- + for (int i = 0; i < (int)pobjs.size(); i++) { + const GreedyRefineCentralGPULB::GObj *obj = pobjs[i]; + double obj_gpu = obj->load; + double obj_cpu = obj->cpuLoad; + + // ----- Outer: choose a GPU group by GPU load ----- + int lightest_gi = 0; + for (int gi = 1; gi < nGroups; gi++) + if (gpuGroups[gi].load < gpuGroups[lightest_gi].load) lightest_gi = gi; + + int chosen_gi = lightest_gi; + if (obj->oldPE >= 0) { + auto curIt = peToGrpIdx.find(obj->oldPE); + if (curIt != peToGrpIdx.end()) { + int cur_gi = curIt->second; + GPUGrp &curGrp = gpuGroups[cur_gi]; + // Keep on current GPU group if it is within B tolerance of the lightest + // group and adding this object keeps it under the target M. + if ((curGrp.load <= (gpuGroups[lightest_gi].load + 0.01) * B) && + (curGrp.load + obj_gpu <= M)) + chosen_gi = cur_gi; + } + } + + GPUGrp &g = gpuGroups[chosen_gi]; + + // ----- Inner: choose a PE within the group by CPU load ----- + int bestPe = g.peIds[0]; + for (int pe : g.peIds) + if (procs[pe].cpuLoad < procs[bestPe].cpuLoad) bestPe = pe; + + // Keep the object on its old PE if that PE is in the chosen group and its + // CPU load is within B tolerance of the lightest PE — reduces migrations. + if (obj->oldPE >= 0 && peToGrpIdx.count(obj->oldPE) && + peToGrpIdx[obj->oldPE] == chosen_gi && + procs[obj->oldPE].cpuLoad <= (procs[bestPe].cpuLoad + 0.01) * B && + procs[obj->oldPE].cpuLoad + obj_cpu / procs[obj->oldPE].speed <= Mcpu) + bestPe = obj->oldPE; + + GreedyRefineCentralGPULB::GProc *p = &procs[bestPe]; + + // Update GPU group aggregate (GPU load) and PE CPU load. + g.load += obj_gpu; + p->cpuLoad += obj_cpu / p->speed; + + if (g.load > maxLoad) { + maxLoad = g.load; + if (maxLoad > M) M = maxLoad; + } + if (p->cpuLoad > maxCpuLoad) { + maxCpuLoad = p->cpuLoad; + if (maxCpuLoad > Mcpu) Mcpu = maxCpuLoad; + } + + if (bestPe != obj->oldPE) { + nmoves++; + stats->to_proc[obj->id] = bestPe; + } + } + + // A step is not over until both resources are done with it, and because + // kernels are launched asynchronously the two overlap, so a group's step time + // is set by whichever dimension is busier: its GPU or its busiest PE. + // Reporting only the GPU term hides every intra-group difference -- with a + // single GPU group that term is the same for every (A,B) candidate, so + // receiveSolutions falls through to its fewest-migrations tie-break and picks + // the do-nothing solution. That is why a one-GPU run never migrated anything. + maxCpuLoadReported = maxCpuLoad; + maxLoad = std::max(maxLoad, maxCpuLoad); + + if ((_lb_args.debug() > 1) && (CkMyPe() == cur_ld_balancer)) { + CkPrintf("[%d] --- Per-GPU-group GPU load after LB ---\n", CkMyPe()); + for (int gi = 0; gi < nGroups; gi++) + CkPrintf("[%d] GPU %llu: aggregate gpuLoad=%.6f\n", + CkMyPe(), (unsigned long long)gpuGroups[gi].gpu_id, gpuGroups[gi].load); + } + +#else + // ---- Original PE-level greedy refine (non-GPU path, CPU only) ---- + + double M = 0; + if (B > 0) { + M = greedyLB(pobjs, procHeap, stats); + greedyMaxLoad = M; + procHeap.addProcessors(procs, (maxLoad <= 0.001), false); + } + + M *= A; + for (int i=0; i < pobjs.size(); i++) { + const GreedyRefineCentralGPULB::GObj *obj = pobjs[i]; + GreedyRefineCentralGPULB::GProc *llp = procHeap.top(); + GreedyRefineCentralGPULB::GProc *prevPe = NULL; + if (obj->oldPE >= 0) prevPe = &(procs[obj->oldPE]); + + GreedyRefineCentralGPULB::GProc *p = llp; + if (prevPe && (prevPe->load <= (llp->load+0.01)*B) && (prevPe->load + obj->load <= M) && (prevPe->available)) + p = prevPe; + + procHeap.remove(p); + p->load += (obj->load / p->speed); + procHeap.push(p); + + if (p->id != obj->oldPE) { + nmoves++; + stats->to_proc[obj->id] = p->id; + } + if (p->load > maxLoad) { + maxLoad = p->load; + if (maxLoad > M) M = maxLoad; + } + } +#endif + // ---------------------------------------------- + + if (concurrent) { + sendSolution(maxLoad, nmoves); + +#if __DEBUG_GREEDY_REFINE_GPU_ + CkCallback cb(CkReductionTarget(GreedyRefineCentralGPULB, receiveTotalTime), thisProxy[cur_ld_balancer]); + contribute(sizeof(double), &strategyStartTime, CkReduction::sum_double, cb); +#endif + } else if (_lb_args.debug() > 0) { + // Single-candidate mode, so this PE's result is the result. Report what the + // tolerance bought: the ratio to greedy says how much max load was given up, + // the migration count says what that saved. + double gpuRatio = 1.0, cpuRatio = 1.0; + if (greedyMaxLoad > 0) gpuRatio = maxLoad / greedyMaxLoad; + if (greedyMaxCpuLoad > 0) cpuRatio = maxCpuLoadReported / greedyMaxCpuLoad; + CkPrintf("CharmLB> %s: after lb, migrations=%d(%.2f%%), tolerance=%.3f, " + "gpu max=%.4f (%.2fx greedy), cpu max=%.4f (%.2fx greedy)\n", + lbname, nmoves, 100.0 * nmoves / double(pobjs.size()), A, + maxLoad, gpuRatio, maxCpuLoadReported, cpuRatio); + } +} + +void GreedyRefineCentralGPULB::receiveTotalTime(double time) +{ + CkPrintf("Avg start time of GreedyRefineCentralGPULB strategy is %f\n", time / CkNumPes()); +} + +// decide which solution among all PEs is best and apply it +void GreedyRefineCentralGPULB::receiveSolutions(CkReductionMsg *msg) +{ + std::vector results(NUM_SOLUTIONS); + + int migrationsAllowed = totalObjs * migrationTolerance; + // feasible solutions are those satistying user's migration constraint + bool feasibleSolutions = false; + float lowest_max_load = FLT_MAX; // lowest max load of all solutions + float lowest_max_load_f = FLT_MAX; // lowest max load of feasible solution set + float highest_max_load = 0; // highest max load of all solutions + int lowestMigrations = INT_MAX; // lowest num migrations of all solutions + const GreedyRefineCentralGPULB::Solution *bestSol = NULL; // best solution + + // first pass. Will record solution with lowest migrations as the best, in case + // there is no feasible solution + CkReduction::setElement *current = (CkReduction::setElement*)msg->getData(); // Get the first element in the set + int numSolutions = 0; + for ( ; current && (numSolutions < NUM_SOLUTIONS); current = current->next()) { + PUP::fromMem pd(¤t->data); + pd|results[numSolutions]; // store result + if (results[numSolutions].migrations >= 0) { // valid result + const GreedyRefineCentralGPULB::Solution &r = results[numSolutions++]; + if ((r.migrations <= migrationsAllowed) && (r.max_load < lowest_max_load_f)) { + lowest_max_load_f = r.max_load; + feasibleSolutions = true; + } + + if ((r.migrations < lowestMigrations) || + ((r.migrations == lowestMigrations) && (r.max_load < bestSol->max_load))) { + lowestMigrations = r.migrations; + bestSol = &r; + } + + if (r.max_load < lowest_max_load) lowest_max_load = r.max_load; + if (r.max_load > highest_max_load) highest_max_load = r.max_load; + } + } + results.resize(numSolutions); // for cases where CkNumPes() < NUM_SOLUTIONS + CkAssert(numSolutions > 0); + + if (feasibleSolutions) { + // second pass, get solution with low max load and migrations from feasible set + int bestMigrations = INT_MAX; // num migrations of best solution + for (int i=0; i < results.size(); i++) { + const GreedyRefineCentralGPULB::Solution &r = results[i]; + if ((r.migrations < bestMigrations && r.max_load <= lowest_max_load_f*LOAD_MIG_BAL) || + (r.migrations == bestMigrations && r.max_load < bestSol->max_load)) { + bestMigrations = r.migrations; + bestSol = &r; + } + } + } + // else: can't satisfy user migration constraint (for this lb step), + // so just use solution with lowest num migrations + + if (_lb_args.debug() > 1) { + CkPrintf("GreedyRefineCentralGPULB: Lowest max_load is %f, worst max_load is %f, lowest migrations=%d\n", + lowest_max_load, highest_max_load, lowestMigrations); + CkPrintf("GreedyRefineCentralGPULB: Got %d solutions at %f\nBest one is from PE %d with max_load=%f, migrations=%d\n", + numSolutions, CkWallTimer(), bestSol->pe, bestSol->max_load, bestSol->migrations); + float A, B; + getGreedyRefineParams(bestSol->pe, A, B); + CkPrintf("Best PE used params A=%f B=%f\n", A, B); + } + + // notify PE that produced the best solution + thisProxy[bestSol->pe].ApplyDecision(); +} + +#include "GreedyRefineCentralGPULB.def.h" + +/*@}*/ diff --git a/src/ck-ldb/GreedyRefineCentralGPULB.ci b/src/ck-ldb/GreedyRefineCentralGPULB.ci new file mode 100644 index 0000000000..4022f7af86 --- /dev/null +++ b/src/ck-ldb/GreedyRefineCentralGPULB.ci @@ -0,0 +1,10 @@ +module GreedyRefineCentralGPULB { + + extern module CentralLB; + initnode void lbinit(void); + group [migratable] GreedyRefineCentralGPULB : CentralLB { + entry void GreedyRefineCentralGPULB(const CkLBOptions &); + entry void receiveSolutions(CkReductionMsg *msg); + entry [reductiontarget] void receiveTotalTime(double time); + }; +}; diff --git a/src/ck-ldb/GreedyRefineCentralGPULB.h b/src/ck-ldb/GreedyRefineCentralGPULB.h new file mode 100644 index 0000000000..892ec2b8c3 --- /dev/null +++ b/src/ck-ldb/GreedyRefineCentralGPULB.h @@ -0,0 +1,115 @@ +/** + * \addtogroup CkLdb +*/ +/*@{*/ + +/** + * Two-level GPU-aware variant of GreedyRefineCentralLB. + * + * Objects are balanced at two granularities: + * - Across GPU groups (the set of PEs that share a gpu_device_id, i.e. the + * PEs of a process bound to one GPU): balanced by GPU load (gpuTime). + * - Within a GPU group, across that group's PEs: balanced by CPU load + * (wallTime). + * + * This matches a setup where each process owns a GPU. The coarse decision of + * which GPU/process an object lands on is driven by the object's GPU work; the + * fine decision of which PE inside that process runs the object's host-side + * code is driven by its CPU work. + * + * Inherits the greedy-refine migration-minimization machinery and the + * concurrent multi-parameter solution search from GreedyRefineCentralLB. + * + * supports processor avail bitvector + * supports nonmigratable attrib +*/ + +#ifndef _GREEDY_REFINE_GPU_LB_H_ +#define _GREEDY_REFINE_GPU_LB_H_ + +#include "CentralLB.h" +#include "GreedyRefineCentralGPULB.decl.h" + +#include + +void CreateGreedyRefineCentralGPULB(); +BaseLB *AllocateGreedyRefineCentralGPULB(); + +#define __DEBUG_GREEDY_REFINE_GPU_ 0 + +class GreedyRefineCentralGPULB : public CBase_GreedyRefineCentralGPULB { +public: + GreedyRefineCentralGPULB(const CkLBOptions &); + GreedyRefineCentralGPULB(CkMigrateMessage *m); + void work(LDStats* stats); + void receiveSolutions(CkReductionMsg *msg); + void receiveTotalTime(double time); + void setMigrationTolerance(float tol) { migrationTolerance = tol; } + +private: + bool QueryBalanceNow(int step) { return true; } + + class GProc { + public: + GProc() : available(true), load(0), bgload(0), cpuLoad(0), bg_walltime(0) {} + int id; + bool available; + int pos; // position in min heap + double load; // GPU load accumulator (used by non-CUDA PE-level path) + double bgload; // background GPU load (group level under CUDA) + double cpuLoad; // CPU/wall load accumulator for within-group balancing + double bg_walltime; + float speed; + }; + + class GObj { + public: + int id; + double load; // GPU load (gpuTime) — drives cross-group assignment + double cpuLoad; // CPU load (wallTime) — drives within-group PE assignment + int oldPE; + }; + + class ObjLoadGreater { + public: + inline bool operator() (const GObj *o1, const GObj *o2) const { +#if CMK_CUDA + // Greedy packing is only as good as its ordering: placing the largest + // items first is what keeps the bins even. Placement here is scored on + // both dimensions, so order by whichever one binds for each object. + // Ordering on GPU time alone leaves the CPU dimension in essentially + // arbitrary order, which packs it worse than not moving anything at all. + return std::max(o1->load, o1->cpuLoad) > std::max(o2->load, o2->cpuLoad); +#else + return (o1->load > o2->load); +#endif + } + }; + + class PHeap; + class Solution; + + double fillData(LDStats *stats, + std::vector &objs, + std::vector &pobjs, + std::vector &procs, + PHeap &procHeap); + + double greedyLB(const std::vector &pobjs, PHeap &procHeap, const BaseLB::LDStats *stats) const; + void sendSolution(double maxLoad, int migrations); + + double strategyStartTime; + double totalObjLoad; + int availablePes; + float migrationTolerance; + int totalObjs; + +#if __DEBUG_GREEDY_REFINE_GPU_ + void dumpObjLoads(std::vector &objs); + void dumpProcLoads(std::vector &procs); +#endif +}; + +#endif + +/*@}*/ diff --git a/src/ck-ldb/LBDatabase.C b/src/ck-ldb/LBDatabase.C index 8c6c1b1d4c..81c1d34647 100644 --- a/src/ck-ldb/LBDatabase.C +++ b/src/ck-ldb/LBDatabase.C @@ -259,6 +259,9 @@ void LBDatabase::ClearLoads(void) obj->data.wallTime = 0.0; #if CMK_LB_CPUTIMER obj->data.cpuTime = 0.0; +#endif +#if CMK_CUDA + obj->data.gpuTime = 0.0; #endif } } @@ -328,3 +331,15 @@ void LBDatabase::EstObjLoad(const LDObjHandle &_h, double cputime) obj->setTiming(cputime); #endif } + +void LBDatabase::EstObjGPULoad(const LDObjHandle &_h, double gputime) +{ +#if CMK_CUDA && CMK_LBDB_ON + LBObj *const obj = LbObj(_h); + + CmiAssert(obj != NULL); + obj->setGPUTiming(gputime); +#else + CmiAbort("LBDatabase::EstObjGPULoad called but CMK_CUDA is not set"); +#endif +} diff --git a/src/ck-ldb/LBDatabase.h b/src/ck-ldb/LBDatabase.h index b344d7c29f..48280d60a2 100644 --- a/src/ck-ldb/LBDatabase.h +++ b/src/ck-ldb/LBDatabase.h @@ -8,6 +8,15 @@ #include "LBComm.h" #include "LBMachineUtil.h" +#if CMK_CUDA && CMK_LBDB_ON +// For hapiCuptiStartTracing/hapiCuptiStopTracing, called from TurnStatsOn/Off. +// Declared here rather than by including hapi.h, which drags the CUDA runtime +// headers into every consumer of LBDatabase.h. +void hapiCuptiStartTracing(); +void hapiCuptiStopTracing(); +#endif + +#include #include class CkSyncBarrier; @@ -67,13 +76,50 @@ friend class LBManager; LbObj(h)->getTime(&walltime, &cputime); }; + inline void GetObjGPULoad(LDObjHandle &h, LBRealType &gputime) { + LbObj(h)->getGPUTime(&gputime); + }; + + // Copy the per-object normalized GPU loads computed by + // hapiNormalizeCuptiLoads into this PE's LB objects. The map is shared + // per-process and every PE reads it concurrently, so this must not mutate it. + inline void SetObjGPULoad( + const std::unordered_map &id_loadMap) + { + for (size_t i = 0; i < objs.size(); i++) { + if (objs[i].obj == nullptr) + continue; + const LDObjHandle &handle = objs[i].obj->GetLDObjHandle(); + LDObjKey key; + key.omID() = handle.omID(); + key.objID() = handle.objID(); + auto it = id_loadMap.find(key); + if (it == id_loadMap.end()) + continue; + objs[i].obj->setGPUTiming(it->second); + } + } + inline void* GetObjUserData(LDObjHandle &h) { return LbObj(h)->getLocalUserData(); } + // GPU activity tracing follows the same switch as CPU instrumentation, so an + // application that calls LBTurnInstrumentOn()/LBTurnInstrumentOff() around + // its own AtSync schedule controls both with one call. Tracing is by far the + // more expensive of the two, so leaving it off between load-balancing steps + // is what makes instrumented runs affordable. inline void TurnStatsOn(void) - {statsAreOn = true; machineUtil.StatsOn();} + {statsAreOn = true; machineUtil.StatsOn(); +#if CMK_CUDA && CMK_LBDB_ON + hapiCuptiStartTracing(); +#endif + } inline void TurnStatsOff(void) - {statsAreOn = false; machineUtil.StatsOff();} + {statsAreOn = false; machineUtil.StatsOff(); +#if CMK_CUDA && CMK_LBDB_ON + hapiCuptiStopTracing(); +#endif + } inline bool StatsOn(void) const { return statsAreOn; }; inline void IdleTime(LBRealType *walltime) { @@ -89,6 +135,7 @@ friend class LBManager; inline void NonMigratable(LDObjHandle h) { LbObj(h)->SetMigratable(false); }; inline void Migratable(LDObjHandle h) { LbObj(h)->SetMigratable(true); }; inline void setPupSize(LDObjHandle h, size_t pup_size) { LbObj(h)->setPupSize(pup_size);}; + inline void setGPUPupSize(LDObjHandle h, size_t gpu_pup_size) { LbObj(h)->setGPUPupSize(gpu_pup_size);}; inline void UseAsyncMigrate(LDObjHandle h, bool flag) { LbObj(h)->UseAsyncMigrate(flag); }; inline int GetCommDataSz(void) { if (commTable) @@ -121,6 +168,7 @@ friend class LBManager; int migratable); void UnregisterObj(LDObjHandle h); void EstObjLoad(const LDObjHandle &h, double cpuload); + void EstObjGPULoad(const LDObjHandle &h, double gpuload); void BackgroundLoad(LBRealType *walltime, LBRealType *cputime); void Send(const LDOMHandle &destOM, const CmiUInt8 &destID, unsigned int bytes, int destObjProc, int force = 0); void MulticastSend(const LDOMHandle &_om, CmiUInt8 *_ids, int _n, unsigned int _b, int _nMsgs=1); diff --git a/src/ck-ldb/LBManager.C b/src/ck-ldb/LBManager.C index 21ccc2ee28..66d15958c9 100644 --- a/src/ck-ldb/LBManager.C +++ b/src/ck-ldb/LBManager.C @@ -315,6 +315,11 @@ void _loadbalancerInit() CmiGetArgIntDesc(argv, "+LBVersion", &_lb_args.lbversion(), "LB database file version number"); CmiGetArgIntDesc(argv, "+LBCentPE", &_lb_args.central_pe(), "CentralLB processor"); + CmiGetArgIntDesc(argv, "+LBPercentMovesAllowed", &_lb_args.percentMovesAllowed(), + "For the greedy-refine balancers, the percentage of chares that can be moved"); + CmiGetArgDoubleDesc(argv, "+LBGreedyRefineTolerance", &_lb_args.greedyRefineTolerance(), + "For the greedy-refine balancers, how far above plain greedy's max load to " + "allow in exchange for fewer migrations (1.1 = 10% higher, the default). 0 or below runs the parameter search instead"); bool _lb_dump_activated = false; if (CmiGetArgIntDesc(argv, "+LBDump", &LBSimulation::dumpStep, "Dump the LB state from this step")) @@ -474,6 +479,7 @@ void LBManager::initnodeFn() _registerCommandLineOpt("+LBPredictorWindow"); _registerCommandLineOpt("+LBVersion"); _registerCommandLineOpt("+LBCentPE"); + _registerCommandLineOpt("+LBPercentMovesAllowed"); _registerCommandLineOpt("+LBDump"); _registerCommandLineOpt("+LBDumpSteps"); _registerCommandLineOpt("+LBDumpFile"); diff --git a/src/ck-ldb/LBManager.h b/src/ck-ldb/LBManager.h index 314513422e..19459bb043 100644 --- a/src/ck-ldb/LBManager.h +++ b/src/ck-ldb/LBManager.h @@ -44,6 +44,22 @@ class CkLBArgs bool _lb_metaLbOn; char* _lb_metaLbModelDir; char* _lb_treeLBFile = (char*)"treelb.json"; + int _lb_percentMovesAllowed; // for the greedy-refine balancers, as a + // percentage of chares that may be moved + // For the greedy-refine balancers: how far above the max load plain greedy + // achieves this balancer may go, in exchange for migrating less. 1.1 allows + // a 10% higher max load. It is the "tolerance" the manual documents. + // + // Defaults to 1.1: allow a max load 10% above what plain greedy achieves, and + // spend that slack on not migrating. Measured on jacobi2d-imbalance, 256 + // chares on 32 PEs, that is 26 moves against 366 for the parameter search it + // replaces as the default, and less than half the time. + // + // 0 or below asks for that search instead: every PE builds a candidate from a + // different (A,B) pair and the best is chosen. It costs a reduction per + // balancing step and is bounded by the PE count, which is why it is no longer + // the default -- but it stays reachable with +LBGreedyRefineTolerance 0. + double _lb_greedyRefineTolerance; public: CkLBArgs() @@ -60,6 +76,8 @@ class CkLBArgs _lb_targetRatio = 1.05; _lb_metaLbOn = false; _lb_metaLbModelDir = nullptr; + _lb_percentMovesAllowed = 100; + _lb_greedyRefineTolerance = 1.1; // 10% over greedy; <=0 searches instead } inline char*& treeLBFile() { return _lb_treeLBFile; } inline double& lbperiod() { return _autoLbPeriod; } @@ -82,6 +100,8 @@ class CkLBArgs inline double& targetRatio() { return _lb_targetRatio; } inline bool& metaLbOn() { return _lb_metaLbOn; } inline char*& metaLbModelDir() { return _lb_metaLbModelDir; } + inline int& percentMovesAllowed() { return _lb_percentMovesAllowed; } + inline double& greedyRefineTolerance() { return _lb_greedyRefineTolerance; } }; extern CkLBArgs _lb_args; @@ -288,6 +308,10 @@ class LBManager : public CBase_LBManager void NonMigratable(LDObjHandle h) { lbdb_obj->NonMigratable(h); } void Migratable(LDObjHandle h) { lbdb_obj->Migratable(h); } void setPupSize(LDObjHandle h, size_t pup_size) { lbdb_obj->setPupSize(h, pup_size); } + void setGPUPupSize(LDObjHandle h, size_t gpu_pup_size) + { + lbdb_obj->setGPUPupSize(h, gpu_pup_size); + } void UseAsyncMigrate(LDObjHandle h, bool flag) { lbdb_obj->UseAsyncMigrate(h, flag); }; int GetObjDataSz(void) { return lbdb_obj->GetObjDataSz(); } int GetCommDataSz(void) { return lbdb_obj->GetCommDataSz(); } @@ -310,6 +334,18 @@ class LBManager : public CBase_LBManager { lbdb_obj->GetObjLoad(h, walltime, cputime); }; + void GetObjGPULoad(LDObjHandle& h, LBRealType& gputime) + { + lbdb_obj->GetObjGPULoad(h, gputime); + }; + // Copy a round's normalized per-object GPU loads (built by + // hapiPrepareCuptiLoads) into this PE's LB objects, just before its stats + // are sent. + void SetObjGPULoad( + const std::unordered_map& id_loadMap) + { + lbdb_obj->SetObjGPULoad(id_loadMap); + } void* GetObjUserData(LDObjHandle& h) { return lbdb_obj->GetObjUserData(h); } void MetaLBCallLBOnChares() { lbdb_obj->MetaLBCallLBOnChares(); } void MetaLBResumeWaitingChares(int lb_period) @@ -338,6 +374,10 @@ class LBManager : public CBase_LBManager { lbdb_obj->EstObjLoad(h, cpuload); } + void EstObjGPULoad(const LDObjHandle& h, double gpuload) + { + lbdb_obj->EstObjGPULoad(h, gpuload); + } void BackgroundLoad(LBRealType* walltime, LBRealType* cputime) { lbdb_obj->BackgroundLoad(walltime, cputime); diff --git a/src/ck-ldb/LBObj.C b/src/ck-ldb/LBObj.C index 36ca87d85f..7c444d4b76 100644 --- a/src/ck-ldb/LBObj.C +++ b/src/ck-ldb/LBObj.C @@ -28,6 +28,10 @@ void LBObj::Clear(void) data.minWall = 1e6; data.maxWall = 0.; #endif +#if CMK_CUDA + data.gpuTime = 0.; + data.gpuPupSize = 0; +#endif } void LBObj::IncrementTime(LBRealType walltime, LBRealType cputime) @@ -42,6 +46,15 @@ void LBObj::IncrementTime(LBRealType walltime, LBRealType cputime) #endif } +void LBObj::IncrementGPUTime(LBRealType gputime) +{ +#if CMK_CUDA + data.gpuTime += gputime; +#else + CmiAbort("LBObj::IncrementGPUTime called but CMK_CUDA is not set"); +#endif +} + #endif /*@}*/ diff --git a/src/ck-ldb/LBObj.h b/src/ck-ldb/LBObj.h index be61c68b02..ff701c3089 100644 --- a/src/ck-ldb/LBObj.h +++ b/src/ck-ldb/LBObj.h @@ -40,6 +40,8 @@ friend class LBDatabase; void Clear(void); void IncrementTime(LBRealType walltime, LBRealType cputime); + void IncrementGPUTime(LBRealType gputime); + inline void StartTimer(void) { startWTime = CkWallTimer(); #if CMK_LB_CPUTIMER @@ -79,12 +81,37 @@ friend class LBDatabase; #endif } + inline void getGPUTime(LBRealType *w) + { +#if CMK_CUDA + *w = data.gpuTime; +#else + CmiAbort("LBObj::getGPUTime called but CMK_CUDA is not set"); +#endif + } + + inline void setGPUTiming(LBRealType gputime) + { +#if CMK_CUDA + data.gpuTime = gputime; +#else + CmiAbort("LBObj::setGPUTiming called but CMK_CUDA is not set"); +#endif + } + inline LDOMHandle &parentOM() { return data.handle.omhandle; } inline const LDObjHandle &GetLDObjHandle() const { return data.handle; } inline void SetMigratable(bool mig) { data.migratable = mig; } inline void setPupSize(size_t obj_pup_size) { data.pupSize = pup_encodeSize(obj_pup_size); } + inline void setGPUPupSize(size_t obj_gpu_pup_size) { +#if CMK_CUDA + data.gpuPupSize = obj_gpu_pup_size; +#else + CmiAbort("LBObj::setGPUPupSize called but CMK_CUDA is not set"); +#endif + } inline void UseAsyncMigrate(bool async) { data.asyncArrival = async; } inline LDObjData &ObjData() { return data; }; inline void lastKnownLoad(LBRealType *w, LBRealType *c) { diff --git a/src/ck-ldb/Make.lb b/src/ck-ldb/Make.lb index 4b53c60d48..d625d22240 100644 --- a/src/ck-ldb/Make.lb +++ b/src/ck-ldb/Make.lb @@ -4,6 +4,7 @@ COMMON_LDBS=\ DistributedLB \ MetisLB \ RecBipartLB \ + GreedyRefineCentralGPULB \ manager.o ALL_LDBS=\ @@ -11,6 +12,7 @@ ALL_LDBS=\ DistributedLB \ MetisLB \ RecBipartLB \ + GreedyRefineCentralGPULB \ manager.o $(L)/libmoduleTreeLB.a: @@ -25,6 +27,9 @@ LBHEADERS += MetisLB.h MetisLB.decl.h $(L)/libmoduleRecBipartLB.a: LBHEADERS += RecBipartLB.h RecBipartLB.decl.h +$(L)/libmoduleGreedyRefineCentralGPULB.a: +LBHEADERS += GreedyRefineCentralGPULB.h GreedyRefineCentralGPULB.decl.h + $(L)/libmoduleScotchLB.a: LBHEADERS += ScotchLB.h ScotchLB.decl.h @@ -38,6 +43,7 @@ ALL_LB_OBJS=EveryLB.o \ DistributedLB.o \ MetisLB.o \ RecBipartLB.o \ + GreedyRefineCentralGPULB.o \ ScotchLB.o \ TempAwareRefineLB.o \ # EveryLB dependecies @@ -46,12 +52,14 @@ EVERYLB_DEPS=EveryLB.o \ DistributedLB.o \ MetisLB.o \ RecBipartLB.o \ + GreedyRefineCentralGPULB.o \ # CommonLBs dependencies COMMONLBS_DEPS=CommonLBs.o \ TreeLB.o \ DistributedLB.o \ MetisLB.o \ RecBipartLB.o \ + GreedyRefineCentralGPULB.o \ manager.o \ $(L)/libmoduleEveryLB.a: $(EVERYLB_DEPS) diff --git a/src/ck-ldb/Makefile_lb.sh b/src/ck-ldb/Makefile_lb.sh index 662052596d..0e6a562463 100755 --- a/src/ck-ldb/Makefile_lb.sh +++ b/src/ck-ldb/Makefile_lb.sh @@ -1,7 +1,7 @@ #!/bin/bash #Typical load balancers -COMMON_LDBS="TreeLB DistributedLB MetisLB RecBipartLB" +COMMON_LDBS="TreeLB DistributedLB MetisLB RecBipartLB GreedyRefineCentralGPULB" #Load balancers for more specialized circumstances SPECIALIZED_LDBS="" #Load balanders which have an external dependency, or require some other kind of intervention diff --git a/src/ck-ldb/lbdb.h b/src/ck-ldb/lbdb.h index 12f330eddd..bf18a3dced 100644 --- a/src/ck-ldb/lbdb.h +++ b/src/ck-ldb/lbdb.h @@ -18,6 +18,7 @@ #endif #include #include +#include #include #include "pup_stl.h" @@ -92,6 +93,68 @@ struct LDObjKey { inline void pup(PUP::er &p); }; +// Hash the complete load-balancing identity. Equality still verifies both +// fields, so hash collisions cannot alias objects from different managers. +struct LDObjKeyHash { + std::size_t operator()(const LDObjKey &key) const noexcept { + uint64_t hash = static_cast(key.objID()); + hash ^= static_cast(key.omID().id.idx) + + UINT64_C(0x9e3779b97f4a7c15) + (hash << 6) + (hash >> 2); + hash ^= hash >> 30; + hash *= UINT64_C(0xbf58476d1ce4e5b9); + hash ^= hash >> 27; + hash *= UINT64_C(0x94d049bb133111eb); + hash ^= hash >> 31; + return static_cast(hash); + } +}; + +#if CMK_CUDA +// CUPTI external-correlation IDs carry only 64 bits, while a Charm++ load- +// balancing identity contains both an object-manager ID and an element ID. +// This table assigns process-local tokens without discarding either part of +// that identity. +// +// Append-only for the lifetime of the process: a token is never reassigned, +// reused for a different identity, or dropped. That is what lets each PE cache +// its own lookups without an invalidation protocol, including across migration +// -- a destination PE simply misses once and interns the same token the source +// already holds. Synchronization is left to the owner, which serializes the +// few paths that reach it. +class GpuObjectTokenTable { +public: + static constexpr uint64_t noObjectToken() { return UINT64_MAX; } + + bool intern(const LDObjKey &key, uint64_t &token) { + auto found = keyToToken_.find(key); + if (found != keyToToken_.end()) { + token = found->second; + return true; + } + if (nextToken_ == noObjectToken()) return false; + + token = nextToken_++; + keyToToken_.emplace(key, token); + tokenToKey_.emplace(token, key); + return true; + } + + bool resolve(uint64_t token, LDObjKey &key) const { + auto found = tokenToKey_.find(token); + if (found == tokenToKey_.end()) return false; + key = found->second; + return true; + } + + std::size_t size() const { return keyToToken_.size(); } + +private: + std::unordered_map keyToToken_; + std::unordered_map tokenToKey_; + uint64_t nextToken_ = 1; +}; +#endif + typedef int LDObjIndex; typedef int LDOMIndex; @@ -154,9 +217,31 @@ class LBObjUserData { void *getData(int idx) { return data.data()+idx; } }; +#if CMK_CUDA +// Per-kernel record captured via CUPTI, used for SM-utilization-normalized +// GPU load attribution. +// +// These are consumed entirely within the process that produced them (see +// hapiNormalizeCuptiLoads) and are never shipped to the central LB: every PE +// sharing a GPU is in the same process, so the whole kernel timeline for a +// device is already local. Only the resulting scalar load per object crosses +// the wire, in LDObjData::gpuTime. +struct LBKernelRecord { + uint64_t start_ns; // CUPTI device-clock timestamp (ns) + uint64_t end_ns; // CUPTI device-clock timestamp (ns) + uint32_t device_id; // CUPTI device id this kernel ran on + int sms_used; // Number of SMs occupied while this kernel was running +}; +#endif + struct LDObjData { LDObjHandle handle; LBRealType wallTime; +#if CMK_CUDA + // SM-utilization-normalized GPU load, in seconds of whole-device occupancy. + // Computed in-process by hapiNormalizeCuptiLoads before the stats are sent. + LBRealType gpuTime; +#endif #if CMK_LB_CPUTIMER LBRealType cpuTime; #endif @@ -171,6 +256,13 @@ struct LDObjData { // An encoded approximation of the amount of data the object would pack; // call pup_decodeSize(pupSize) to get the actual approximate value CmiUInt2 pupSize; +#if CMK_CUDA + // Device bytes the object would pack, exactly rather than encoded: a device + // buffer is sized in megabytes where pupSize's approximation is tuned for + // host state, and a balancer weighing a migration against free device memory + // needs the real figure. + size_t gpuPupSize; +#endif inline const LDOMHandle &omHandle() const { return handle.omhandle; } inline const LDOMid &omID() const { return handle.omhandle.id; } inline const CmiUInt8 &objID() const { return handle.id; } @@ -333,6 +425,9 @@ inline void LBObjUserData::pup(PUP::er &p) { inline void LDObjData::pup(PUP::er &p) { p|handle; p|wallTime; +#if CMK_CUDA + p|gpuTime; +#endif #if CMK_LB_CPUTIMER p|cpuTime; #endif @@ -348,6 +443,9 @@ inline void LDObjData::pup(PUP::er &p) { } #endif p|pupSize; +#if CMK_CUDA + p|gpuPupSize; +#endif } inline bool LDCommDesc::operator==(const LDCommDesc &obj) const { diff --git a/src/scripts/Make.cidepends b/src/scripts/Make.cidepends index 71e3aff443..b54f582d7d 100644 --- a/src/scripts/Make.cidepends +++ b/src/scripts/Make.cidepends @@ -5,6 +5,7 @@ CommonLBs.decl.h CommonLBs.def.h: CommonLBs.ci.stamp DistBaseLB.decl.h DistBaseLB.def.h: DistBaseLB.ci.stamp DistributedLB.decl.h DistributedLB.def.h: DistributedLB.ci.stamp EveryLB.decl.h EveryLB.def.h: EveryLB.ci.stamp +GreedyRefineCentralGPULB.decl.h GreedyRefineCentralGPULB.def.h: GreedyRefineCentralGPULB.ci.stamp HybridBaseLB.decl.h HybridBaseLB.def.h: HybridBaseLB.ci.stamp LBManager.decl.h LBManager.def.h: LBManager.ci.stamp MetaBalancer.decl.h MetaBalancer.def.h: MetaBalancer.ci.stamp diff --git a/src/util/pup.h b/src/util/pup.h index e9b89ddf81..a058646d70 100644 --- a/src/util/pup.h +++ b/src/util/pup.h @@ -133,6 +133,28 @@ typedef enum { dataType_last //<- for setting table lengths, etc. } dataType; +/// Which side of the machine a pupped buffer lives on. DEVICE means `p` is a +/// device pointer, so the bytes travel through the migration message's device +/// region rather than its host region. PUP::ers that have no device region +/// ignore DEVICE buffers entirely -- checkpointing a chare does not capture +/// its device state. +enum class PUPMode { + HOST, + DEVICE +}; + +/// Layout rule for the device side of a migration stream. Every DEVICE-mode +/// buffer starts at an offset that is a multiple of this, measured from the +/// start of the device region -- relative, not absolute, because the region's +/// own base carries no alignment guarantee. The sizer and both memory walkers +/// apply the same rule, which is what lets the receiver reconstruct the +/// sender's layout from its own pup calls alone, with nothing on the wire +/// describing it. +constexpr size_t DEVICE_PUP_ALIGN = 256; +static inline size_t alignDeviceOffset(size_t offset) { + return (offset + (DEVICE_PUP_ALIGN - 1)) & ~(DEVICE_PUP_ALIGN - 1); +} + static inline dataType getXlateDataType(signed char *a) { return Tchar; } #if CMK_SIGNEDCHAR_DIFF_CHAR static inline dataType getXlateDataType(char *a) { return Tchar; } @@ -267,6 +289,18 @@ class er { bytes((void *)a,nItems, sizeof(T), getXlateDataType(a)); } + /// For a buffer that lives on the device. `a` must already point at + /// device memory in BOTH directions -- unlike pup_buffer, this does not + /// allocate, so an unpacking chare allocates its device buffer first and + /// then pups into it. + template + void operator()(T *a,size_t nItems,PUPMode mode) { + if (mode == PUPMode::DEVICE) + bytes_device((void *)a,nItems, sizeof(T), getXlateDataType(a)); + else + bytes((void *)a,nItems, sizeof(T), getXlateDataType(a)); + } + // Standard pup_buffer API that calls malloc for allocation on isUnpacking and free for deallocation on isPacking template void pup_buffer(T *&a, size_t nItems) { @@ -324,6 +358,17 @@ class er { //Generic bottleneck: pack/unpack n items of size itemSize // and data type t from p. Desc describes the data item virtual void bytes(void *p,size_t n,size_t itemSize,dataType t) =0; + /// Device-mode bottleneck. Deliberately a distinct name rather than an + /// overload of bytes(): as an overload it would be hidden in every subclass + /// that overrides the host bottleneck alone, which is all but three of them, + /// and every compiler that warns about a partially overridden virtual would + /// say so in user code that merely includes this header. + /// + /// Defaults to dropping the buffer, which is what every PUP::er without a + /// device region should do: sizing for a checkpoint, writing to disk and + /// converting to text all operate on host memory and cannot dereference `p` + /// at all. Only the migration walkers override it. + virtual void bytes_device(void *p,size_t n,size_t itemSize,dataType t) {} virtual void object(able** a); virtual void pup_buffer(void *&p, size_t n, size_t itemSize, dataType t) = 0; @@ -391,21 +436,29 @@ enum { class sizer : public er { protected: size_t nBytes; + size_t gpuBytes; //Generic bottleneck: n items of size itemSize virtual void bytes(void *p,size_t n,size_t itemSize,dataType t); + virtual void bytes_device(void *p,size_t n,size_t itemSize,dataType t); virtual void pup_buffer(void *&p, size_t n, size_t itemSize, dataType t); virtual void pup_buffer(void *&p, size_t n, size_t itemSize, dataType t, std::function allocate, std::function deallocate); public: //Write data to the given buffer - sizer(const unsigned int purpose = 0) : er(IS_SIZING | purpose), nBytes(0) + sizer(const unsigned int purpose = 0) : er(IS_SIZING | purpose), nBytes(0), + gpuBytes(0) { CmiAssert((purpose & TYPE_MASK) == 0); } //Return the current number of bytes to be packed size_t size(void) const {return nBytes;} + + //Return the number of DEVICE bytes to be packed, including the padding + //DEVICE_PUP_ALIGN inserts between buffers, so it is directly comparable + //with mem::gpu_size(). + size_t gpu_size(void) const {return gpuBytes;} }; template @@ -418,8 +471,15 @@ class mem : public er { //Memory-buffer packers and unpackers protected: myByte *origBuf;//Start of memory buffer myByte *buf;//Memory buffer (stuff gets packed into/out of here) - mem(const unsigned int type, myByte* Nbuf, const unsigned int purpose = 0) - : er(type | purpose), origBuf(Nbuf), buf(Nbuf) + //Device-side counterparts, NULL unless this walker was given a device + //region. A DEVICE-mode buffer pupped through a walker with no device region + //is dropped, the same as for every other PUP::er. + myByte *gpuOrigBuf; + myByte *gpuBuf; + mem(const unsigned int type, myByte* Nbuf, myByte* gpuNbuf, + const unsigned int purpose = 0) + : er(type | purpose), origBuf(Nbuf), buf(Nbuf), + gpuOrigBuf(gpuNbuf), gpuBuf(gpuNbuf) { CmiAssert((purpose & TYPE_MASK) == 0); } @@ -434,6 +494,12 @@ class mem : public er { //Memory-buffer packers and unpackers //Return the current number of buffer bytes used size_t size(void) const {return buf-origBuf;} + //Device bytes consumed so far, including inter-buffer alignment padding, + //so it is directly comparable with sizer::gpu_size(). + size_t gpu_size(void) const { + return (gpuOrigBuf == nullptr) ? 0 : (size_t)(gpuBuf - gpuOrigBuf); + } + inline char* get_current_pointer() const { return reinterpret_cast(buf); } @@ -456,14 +522,20 @@ class toMem : public mem { protected: //Generic bottleneck: pack n items of size itemSize from p. virtual void bytes(void *p,size_t n,size_t itemSize,dataType t); + virtual void bytes_device(void *p,size_t n,size_t itemSize,dataType t); virtual void pup_buffer(void *&p, size_t n, size_t itemSize, dataType t); virtual void pup_buffer(void *&p, size_t n, size_t itemSize, dataType t, std::function allocate, std::function deallocate); public: - //Write data to the given buffer + //Write data to the given buffer, with an optional device region for + //DEVICE-mode buffers. + toMem(void* Nbuf, void* gpuNbuf, const unsigned int purpose = 0) + : mem(IS_PACKING, (myByte*)Nbuf, (myByte*)gpuNbuf, purpose) + { + } toMem(void* Nbuf, const unsigned int purpose = 0) - : mem(IS_PACKING, (myByte*)Nbuf, purpose) + : mem(IS_PACKING, (myByte*)Nbuf, nullptr, purpose) { } }; @@ -480,6 +552,7 @@ class fromMem : public mem { protected: //Generic bottleneck: unpack n items of size itemSize from p. virtual void bytes(void *p,size_t n,size_t itemSize,dataType t); + virtual void bytes_device(void *p,size_t n,size_t itemSize,dataType t); virtual void pup_buffer(void *&p, size_t n, size_t itemSize, dataType t); virtual void pup_buffer(void *&p, size_t n, size_t itemSize, dataType t, std::function allocate, std::function deallocate); @@ -487,9 +560,14 @@ class fromMem : public mem { void pup_buffer_generic(void *&p,size_t n, size_t itemSize, dataType t, std::function allocate, bool isMalloc); public: - //Read data from the given buffer + //Read data from the given buffer, with an optional device region holding + //the DEVICE-mode buffers the sender packed. + fromMem(const void* Nbuf, const void* gpuNbuf, const unsigned int purpose = 0) + : mem(IS_UNPACKING, (myByte*)Nbuf, (myByte*)gpuNbuf, purpose) + { + } fromMem(const void* Nbuf, const unsigned int purpose = 0) - : mem(IS_UNPACKING, (myByte*)Nbuf, purpose) + : mem(IS_UNPACKING, (myByte*)Nbuf, nullptr, purpose) { } }; diff --git a/src/util/pup_util.C b/src/util/pup_util.C index 561b7ecbad..74e7717017 100644 --- a/src/util/pup_util.C +++ b/src/util/pup_util.C @@ -24,6 +24,17 @@ virtual functions are defined here. #include "ckhashtable.h" #include "conv-rdma.h" + +#if CMK_CUDA +// The CUDA runtime directly, NOT hapi.h, for the copies behind PUPMode::DEVICE. +// pup_util.o is part of the Converse-level utility library (LIBCONV_UTIL), which +// is linked into pure-Converse programs that have no libck. hapiCheck expands to +// hapiErrorDie, which lives in hapi_impl.cpp, so referencing it from here drags +// that whole object into every such link -- along with the Charm++ symbols it +// uses (_lb_args, CkActiveLocRec, CkCallback::send), none of which can resolve +// there. Keep this file's dependencies at the Converse level. +#include +#endif #if defined(_WIN32) #include @@ -162,11 +173,59 @@ void PUP::fromMem::bytes(void *p,size_t n,size_t itemSize,dataType t) ((pupCheckRec *)buf)->check(t,n); buf+=sizeof(pupCheckRec); #endif - n*=itemSize; - memcpy(p,(const void *)buf,n); + n*=itemSize; + memcpy(p,(const void *)buf,n); buf+=n; } +/*Device-mode PUP::er's. + * + * The sender's device region layout is never described on the wire: the + * receiver reconstructs it by making the same sequence of pup calls, so the + * sizer, the packer and the unpacker must agree exactly on where each buffer + * starts. DEVICE_PUP_ALIGN is that agreement. + * + * A walker with no device region (gpuOrigBuf == NULL) drops DEVICE buffers, + * matching the base er::bytes default. A chare with device state pupped this + * way therefore migrates its device data but does not checkpoint it. + */ +#if CMK_CUDA +static void pupDeviceCopy(void *dst, const void *src, size_t n) +{ + cudaError_t err = cudaMemcpy(dst, src, n, cudaMemcpyDeviceToDevice); + if (err != cudaSuccess) + CmiAbort("PUP: device-to-device copy of %zu bytes failed: %s", + n, cudaGetErrorString(err)); +} +#endif + +void PUP::sizer::bytes_device(void *p,size_t n,size_t itemSize,dataType t) +{ + gpuBytes = alignDeviceOffset(gpuBytes) + n*itemSize; +} + +void PUP::toMem::bytes_device(void *p,size_t n,size_t itemSize,dataType t) +{ + if (gpuOrigBuf == nullptr) return; + n*=itemSize; +#if CMK_CUDA + gpuBuf = gpuOrigBuf + alignDeviceOffset((size_t)(gpuBuf - gpuOrigBuf)); + pupDeviceCopy((void *)gpuBuf, p, n); + gpuBuf += n; +#endif +} + +void PUP::fromMem::bytes_device(void *p,size_t n,size_t itemSize,dataType t) +{ + if (gpuOrigBuf == nullptr) return; + n*=itemSize; +#if CMK_CUDA + gpuBuf = gpuOrigBuf + alignDeviceOffset((size_t)(gpuBuf - gpuOrigBuf)); + pupDeviceCopy(p, (const void *)gpuBuf, n); + gpuBuf += n; +#endif +} + void PUP::sizer::pup_buffer(void *&p,size_t n, size_t itemSize, dataType t) { #ifdef CK_CHECK_PUP nBytes+=sizeof(pupCheckRec); diff --git a/tests/charm++/cuda/gpumigrate/.gitignore b/tests/charm++/cuda/gpumigrate/.gitignore new file mode 100644 index 0000000000..9c08629e30 --- /dev/null +++ b/tests/charm++/cuda/gpumigrate/.gitignore @@ -0,0 +1 @@ +gpumigrate diff --git a/tests/charm++/cuda/gpumigrate/Makefile b/tests/charm++/cuda/gpumigrate/Makefile new file mode 100644 index 0000000000..12e5ce3d77 --- /dev/null +++ b/tests/charm++/cuda/gpumigrate/Makefile @@ -0,0 +1,41 @@ +# tests/charm++/cuda/gpumigrate -- device-state migration test. +# +# No .cu file: the buffers are filled and read back with plain device copies, +# because what is under test is whether device state survives a migration, not +# whether a kernel ran. +# +# CHARM_DIR defaults to ../../../.. , which is the Charm++ build root when this +# directory is reached through the build tree, the same convention the other +# tests use. Building straight out of the source tree needs it spelled out: +# +# make CHARM_DIR=../../../../multicore-linux-x86_64-cuda + +OPTS = -O2 -g + +CHARM_DIR ?= ../../../.. +CHARMC = $(CHARM_DIR)/bin/charmc $(OPTS) + +TARGET = gpumigrate +all: $(TARGET) + +$(TARGET): $(TARGET).o + $(CHARMC) -language charm++ -module CommonLBs -o $@ $(TARGET).o $(LD_LIBS) + +$(TARGET).decl.h $(TARGET).def.h: $(TARGET).ci + $(CHARMC) $< + +$(TARGET).o: $(TARGET).C $(TARGET).decl.h $(TARGET).def.h + $(CHARMC) -c $< + +clean: + rm -f *.decl.h *.def.h conv-host *.o $(TARGET) charmrun + +# Two PEs in one process: the same-process transport only. IPC needs more than +# one process and RDMA more than one physical node, which a bare `make test` +# cannot arrange -- see README.txt. +# +# RotateLB moves every object one PE onward at each step, which is what makes +# the migration happen at all: no other strategy would move a set of identically +# loaded objects. +test: all + ./$(TARGET) +pe 2 +balancer RotateLB diff --git a/tests/charm++/cuda/gpumigrate/README.txt b/tests/charm++/cuda/gpumigrate/README.txt new file mode 100644 index 0000000000..67b6ecde12 --- /dev/null +++ b/tests/charm++/cuda/gpumigrate/README.txt @@ -0,0 +1,80 @@ +gpumigrate -- device-state migration test +========================================= + +What it covers +-------------- +A chare that keeps state in device memory and pups it with PUPMode::DEVICE +must find that state intact after it migrates. This test builds the smallest +thing that can check that: each element owns two device buffers filled with a +pattern derived from its array index, every element is moved one PE onward +several times, and after each move the buffers are read back and compared +element by element. + +Two buffers rather than one, and the second is deliberately not a multiple of +DEVICE_PUP_ALIGN. The device side of a migration stream carries no description +of its own layout -- the receiver reconstructs it by making the same sequence +of pup calls the sender made -- so a packer and an unpacker that disagree about +where the second buffer begins would go undetected with a single buffer. + +On arrival the buffers are poisoned before the device pup runs. Without that +the test can pass while doing nothing: on the same-process path the source has +just freed buffers of exactly the right size, and cudaMalloc is free to return +that same device memory with the old contents still in it. + +Migration is driven through AtSync with RotateLB, which moves every object one +PE onward at each step. No load-based strategy would move a set of identically +loaded objects, and RotateLB exists for exactly this -- exercising pup routines +and migration paths. + +Anytime migration (ckMigrate/migrateMe) would be the more direct way to move an +element, and is deliberately not used: it is unsupported under +CMK_GLOBAL_LOCATION_UPDATE, which this test requires (below). Going through a +real load-balancing step is also closer to the thing being protected, since +that is how a GPU application's chares actually move. + +Build prerequisite +------------------ +The Charm++ build must set -DCMK_GLOBAL_LOCATION_UPDATE=1, passed through +EXTRA_OPTS at configure time: + + cmake ... -DEXTRA_OPTS="-DCMK_GLOBAL_LOCATION_UPDATE=1" + +Device zerocopy sends are addressed to the PE the sender believes hosts the +target. Without the global update a send issued around a migration arrives at +a PE that no longer hosts it, and CkRdmaDeviceIssueRgets aborts rather than +read the wrong buffer. This test sends no device zerocopy messages of its own, +so that particular abort is not what it is checking -- but it migrates the way +an application in this mode has to, so build it the way such an application is +built. See "Global Location Update" in the manual. + +Running it +---------- + make CHARM_DIR= + make test # 8 blocks, 3 rounds, 2 PEs, +balancer RotateLB + +or directly: + + ./gpumigrate [blocks] [rounds] +pe +balancer RotateLB + +It needs at least 2 PEs; with one there is nowhere to migrate to. It also needs +a balancer that actually moves something -- without +balancer RotateLB the +blocks stay put and the test verifies nothing. + +Transport coverage +------------------ +The device payload travels on the ordinary device zerocopy path, so which +transport carries it is decided by findTransferModeDevice, not by this test: + + same process device-to-device copy `+pe 2` + same physical node, 2+ procs CUDA IPC two processes on one host + different physical nodes device RDMA a 2-node job + +`make test` covers only the first. The other two need a job layout a bare make +cannot arrange, and are what a reviewer with a multi-GPU machine or a two-node +allocation should run: + + # same node, two processes, one GPU each + ./gpumigrate 16 5 +p2 ++ppn 1 + + # two nodes + -n 2 -N 1 ./gpumigrate 16 5 diff --git a/tests/charm++/cuda/gpumigrate/gpumigrate.C b/tests/charm++/cuda/gpumigrate/gpumigrate.C new file mode 100644 index 0000000000..7b572872dc --- /dev/null +++ b/tests/charm++/cuda/gpumigrate/gpumigrate.C @@ -0,0 +1,230 @@ +/** + * Device-state migration test. + * + * Each element owns two device buffers holding patterns derived from its + * index. The test moves every element to the next PE, several times, and after + * each move reads the buffers back and compares them element by element. What + * is under test is PUPMode::DEVICE and the migration path behind it: the + * source packs the device buffers into a staged payload, that payload travels + * on the device zerocopy path, and the destination copies it into the + * element's freshly allocated buffers. + * + * Two buffers, not one, and the second is deliberately not a multiple of + * DEVICE_PUP_ALIGN: a single buffer would not catch a packer and an unpacker + * that disagree about where the next one starts. + * + * Migration is driven through AtSync with RotateLB, which moves every object one + * PE onward at each step -- the strategy the manual recommends for exercising + * pup routines and migration paths. Anytime migration (ckMigrate/migrateMe) is + * deliberately NOT used: it is unsupported under CMK_GLOBAL_LOCATION_UPDATE, + * which a GPU build needs (see "Global Location Update" in the manual). Driving + * the move through a real load-balancing step is also closer to what this test + * exists to protect. + * + * Verification runs in ResumeFromSync, which is reached once the element and its + * device state have both landed, so the test needs no barrier of its own. + * + * Two PEs in one process covers the same-process transport. The IPC and RDMA + * transports need more than one process and more than one physical node + * respectively; see README.txt. + */ +#include "gpumigrate.decl.h" +#include "hapi.h" +#include + +/* readonly */ CProxy_Main mainProxy; +/* readonly */ int numBlocks; +/* readonly */ int numRounds; + +// Big enough that the payload is a real device allocation spanning several +// DEVICE_PUP_ALIGN boundaries, small enough to stay trivial. +#define BUF_LEN 4096 +#define SECOND_LEN 517 // deliberately not a multiple of the device alignment + +class Main : public CBase_Main +{ + CProxy_Block blocks; + int round; + int checked; + +public: + Main(CkArgMsg* m) + { + numBlocks = 4 * CkNumPes(); + numRounds = 3; + if (m->argc > 1) numBlocks = atoi(m->argv[1]); + if (m->argc > 2) numRounds = atoi(m->argv[2]); + delete m; + + if (CkNumPes() < 2) + CkAbort("gpumigrate needs at least 2 PEs: with one PE there is nowhere " + "to migrate to and nothing would be tested."); + + round = 0; + checked = 0; + mainProxy = thisProxy; + CkPrintf("gpumigrate: %d blocks over %d PEs, %d migration round(s), " + "%d + %d device elements per block\n", + numBlocks, CkNumPes(), numRounds, BUF_LEN, SECOND_LEN); + + blocks = CProxy_Block::ckNew(numBlocks); + blocks.check(); // round 0: before anything has moved + } + + void blockChecked() + { + if (++checked < numBlocks) return; + checked = 0; + + if (round == numRounds) + { + CkPrintf("gpumigrate: PASSED -- device state survived %d migration " + "round(s)\n", numRounds); + CkExit(); + return; + } + round++; + CkPrintf("gpumigrate: round %d, migrating every block one PE onward\n", + round); + blocks.step(); + } +}; + +class Block : public CBase_Block +{ + double* d_buf; + int* d_second; + int moves; // how many times this element has migrated + int lastPe; // PE this element was on before the current step + +public: + Block() : d_buf(nullptr), d_second(nullptr), moves(0), lastPe(CkMyPe()) + { + usesAtSync = true; + allocate(); + fill(); + } + + Block(CkMigrateMessage* m) + : CBase_Block(m), d_buf(nullptr), d_second(nullptr), moves(0), lastPe(-1) + { + } + + ~Block() + { + if (d_buf) hapiCheck(hapiFree(d_buf)); + if (d_second) hapiCheck(hapiFree(d_second)); + } + + void allocate() + { + hapiCheck(hapiMalloc((void**)&d_buf, BUF_LEN * sizeof(double))); + hapiCheck(hapiMalloc((void**)&d_second, SECOND_LEN * sizeof(int))); + } + + // The patterns are pure functions of the index, so they are checkable from + // wherever the element lands. + double expectedD(int i) const { return thisIndex * 1000.0 + i; } + int expectedI(int i) const { return thisIndex * 7 + i; } + + void fill() + { + std::vector h(BUF_LEN); + for (int i = 0; i < BUF_LEN; i++) h[i] = expectedD(i); + hapiCheck(cudaMemcpy(d_buf, h.data(), BUF_LEN * sizeof(double), + cudaMemcpyHostToDevice)); + + std::vector h2(SECOND_LEN); + for (int i = 0; i < SECOND_LEN; i++) h2[i] = expectedI(i); + hapiCheck(cudaMemcpy(d_second, h2.data(), SECOND_LEN * sizeof(int), + cudaMemcpyHostToDevice)); + } + + void pup(PUP::er& p) + { + CBase_Block::pup(p); + p | moves; + p | lastPe; + + // The device buffers must exist before they are pupped: PUPMode::DEVICE + // copies into what the pointer names, it does not allocate. + if (p.isUnpacking()) + { + allocate(); + // Poison them first, or the test can pass without the device pup doing + // anything: on the same-process path the source has just freed buffers + // of exactly this size, and cudaMalloc is free to hand the very same + // device memory back with the old contents still in it. + hapiCheck(cudaMemset(d_buf, 0xA5, BUF_LEN * sizeof(double))); + hapiCheck(cudaMemset(d_second, 0xA5, SECOND_LEN * sizeof(int))); + } + + p(d_buf, BUF_LEN, PUP::PUPMode::DEVICE); + p(d_second, SECOND_LEN, PUP::PUPMode::DEVICE); + + // The packing walker for a migration is deleting, and the copies it just + // made are synchronous, so the source's buffers are finished with here. + if (p.isDeleting()) + { + hapiCheck(hapiFree(d_buf)); + d_buf = nullptr; + hapiCheck(hapiFree(d_second)); + d_second = nullptr; + } + } + + void verify() + { + std::vector h(BUF_LEN); + hapiCheck(cudaMemcpy(h.data(), d_buf, BUF_LEN * sizeof(double), + cudaMemcpyDeviceToHost)); + for (int i = 0; i < BUF_LEN; i++) + { + if (h[i] != expectedD(i)) + CkAbort("gpumigrate: block %d on PE %d: d_buf[%d] is %f, expected %f " + "(after %d migration(s))", + thisIndex, CkMyPe(), i, h[i], expectedD(i), moves); + } + + std::vector h2(SECOND_LEN); + hapiCheck(cudaMemcpy(h2.data(), d_second, SECOND_LEN * sizeof(int), + cudaMemcpyDeviceToHost)); + for (int i = 0; i < SECOND_LEN; i++) + { + if (h2[i] != expectedI(i)) + CkAbort("gpumigrate: block %d on PE %d: d_second[%d] is %d, expected " + "%d (after %d migration(s))", + thisIndex, CkMyPe(), i, h2[i], expectedI(i), moves); + } + } + + void check() + { + verify(); + mainProxy.blockChecked(); + } + + // Join the load-balancing step. RotateLB will move this element one PE on. + void step() { AtSync(); } + + // Reached once this element and its device payload have both landed. + void ResumeFromSync() + { + // A step in which nothing actually moved would let every later check pass + // trivially -- the buffers were never packed, sent or unpacked, so of + // course they still read correctly. Catch that here rather than report a + // green run that tested nothing. RotateLB moves every object at every + // step, so on more than one PE the destination is always a different one. + if (CkMyPe() == lastPe) + CkAbort("gpumigrate: block %d did not move at step %d (it is still on " + "PE %d). Run with +balancer RotateLB -- without a balancer that " + "migrates, this test verifies nothing.", + thisIndex, moves + 1, CkMyPe()); + lastPe = CkMyPe(); + moves++; + verify(); + mainProxy.blockChecked(); + } +}; + +#include "gpumigrate.def.h" diff --git a/tests/charm++/cuda/gpumigrate/gpumigrate.ci b/tests/charm++/cuda/gpumigrate/gpumigrate.ci new file mode 100644 index 0000000000..fa692a6dce --- /dev/null +++ b/tests/charm++/cuda/gpumigrate/gpumigrate.ci @@ -0,0 +1,21 @@ +mainmodule gpumigrate { + readonly CProxy_Main mainProxy; + readonly int numBlocks; + readonly int numRounds; + + mainchare Main { + entry Main(CkArgMsg *m); + // One element finished verifying its device state. + entry void blockChecked(); + }; + + array [1D] Block { + entry Block(); + // Read the device buffers back and compare them against what this element + // wrote before it moved. + entry void check(); + // Join the load-balancing step that hands this element to the next PE, + // taking its device state with it. + entry void step(); + }; +};