charmplusplus · anant37289 · Dec 12, 2025 · Dec 24, 2025 · Jan 8, 2026 · Jan 9, 2026
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -208,6 +208,7 @@ option(BUILD_SHARED        "Build Charm++ dynamic libraries" OFF)
 
 # Other options
 option(BUILD_CUDA          "Build with CUDA support" OFF)
+option(BUILD_HIP           "Build with HIP support" OFF)
 option(PXSHM               "Build with PXSHM" OFF)
 
 # LRTS PMI options
@@ -511,7 +512,7 @@ if(EXISTS ${CMAKE_SOURCE_DIR}/src/arch/${NETWORK}/gdir_link)
   file(STRINGS src/arch/${NETWORK}/gdir_link GDIR)
 elseif(${NETWORK} MATCHES "gni-")
   set(GDIR "gni")
-elseif(${NETWORK} MATCHES "mpi-cray")
+elseif(${NETWORK} MATCHES "`mpi`-cray")
   set(GDIR "mpi")
 elseif(${NETWORK} MATCHES "ofi-cray")
   set(GDIR "ofi")
@@ -533,8 +534,12 @@ else()
   set(CMK_BUILD_CHARMRUN 1)
 endif()
 
+set(CHARMRUN_ELASTIC_DIR src/arch/common)
+set(CHARMRUN_HAPI_DIR src/arch/common)
+
 include(cmake/detect-features.cmake)
 include(cmake/ci-files.cmake)
+add_custom_target(ci-generated DEPENDS ${all-ci-outputs})
 
 
 if(${TARGET} STREQUAL "all-test")
@@ -660,6 +665,8 @@ configure_file(src/arch/common/cc-msvc.sh                    include/ COPYONLY)
 configure_file(src/arch/common/conv-mach-craype.sh           include/ COPYONLY)
 configure_file(src/arch/common/conv-mach-cuda.sh             include/ COPYONLY)
 configure_file(src/arch/common/conv-mach-cuda.h              include/ COPYONLY)
+configure_file(src/arch/common/conv-mach-hip.sh              include/ COPYONLY)
+configure_file(src/arch/common/conv-mach-hip.h               include/ COPYONLY)
 configure_file(src/arch/common/conv-mach-darwin.sh           include/ COPYONLY)
 configure_file(src/arch/common/conv-mach-flang.h             include/ COPYONLY)
 configure_file(src/arch/common/conv-mach-flang.sh            include/ COPYONLY)
@@ -688,29 +695,110 @@ configure_file(src/arch/common/conv-mach-tsan.h              include/ COPYONLY)
 configure_file(src/arch/common/conv-mach-tsan.sh             include/ COPYONLY)
 configure_file(src/scripts/conv-config.sh                    include/ COPYONLY)
 configure_file(src/arch/${VDIR}/conv-mach.sh                 include/ COPYONLY)
+configure_file(src/util/ckrescale.h                         include/ COPYONLY)
+
+add_library(ckrescale src/util/ckrescale.C)
 
 set(CUDA_DIR "")
-if(BUILD_CUDA)
+set(HIP_DIR "")
+if(BUILD_CUDA OR BUILD_HIP)
 
   file(GLOB_RECURSE hybridAPI-h-sources ${CMAKE_SOURCE_DIR}/src/arch/cuda/*.h)
   file(GLOB_RECURSE hybridAPI-cxx-sources ${CMAKE_SOURCE_DIR}/src/arch/cuda/*.cpp)
   foreach(file ${hybridAPI-h-sources})
     configure_file(${file} include/ COPYONLY)
   endforeach()
 
-  if(CMAKE_VERSION VERSION_GREATER 3.17 OR CMAKE_VERSION VERSION_EQUAL 3.17)
-    find_package(CUDAToolkit REQUIRED)
-    set(CMAKE_CUDA_COMPILER "${CUDAToolkit_NVCC_EXECUTABLE}")
-    enable_language(CUDA)
-    set(CUDA_DIR "${CUDAToolkit_TARGET_DIR}")
-  else()
-    find_package(CUDA REQUIRED)
-    set(CUDA_DIR "${CUDA_TOOLKIT_ROOT_DIR}")
+  if (BUILD_CUDA)
+    if(CMAKE_VERSION VERSION_GREATER 3.17 OR CMAKE_VERSION VERSION_EQUAL 3.17)
+      find_package(CUDAToolkit REQUIRED)
+      set(CMAKE_CUDA_COMPILER "${CUDAToolkit_NVCC_EXECUTABLE}")
+      enable_language(CUDA)
+      set(CUDA_DIR "${CUDAToolkit_TARGET_DIR}")
+    else()
+      find_package(CUDA REQUIRED)
+      set(CUDA_DIR "${CUDA_TOOLKIT_ROOT_DIR}")
+    endif()
+
+    # Find CUPTI library and include directory
+    find_library(CUPTI_LIBRARY cupti
+      HINTS "${CUDA_DIR}/extras/CUPTI/lib64"
+            "${CUDA_DIR}/lib64"
+    )
+    set(CUPTI_INCLUDE_DIR "${CUDA_DIR}/extras/CUPTI/include")
+    if(NOT CUPTI_LIBRARY)
+      message(WARNING "CUPTI library not found. GPU load balancing will not be available.")
+    else()
+      message(STATUS "Found CUPTI: ${CUPTI_LIBRARY}")
+    endif()
+
+    add_library(hybridapi ${hybridAPI-cxx-sources} $<TARGET_OBJECTS:ckrescale>)
+    add_dependencies(hybridapi ci-generated)
+
+    if(CUPTI_LIBRARY)
+      target_include_directories(hybridapi PRIVATE "${CUPTI_INCLUDE_DIR}")
+      target_link_libraries(hybridapi ${CUPTI_LIBRARY})
+    endif()
+
+    if(TRACING)
+      target_compile_definitions(hybridapi PRIVATE HAPI_TRACE)
+    endif()
   endif()
-  add_library(cudahybridapi ${hybridAPI-cxx-sources})
-  if(TRACING)
-    target_compile_definitions(cudahybridapi PRIVATE HAPI_TRACE)
+
+  if (BUILD_HIP)
+    add_compile_definitions(__HIP_PLATFORM_AMD__)
+    # Modern ROCm/HIP detection
+    if(NOT DEFINED ROCM_PATH)
+      if(NOT DEFINED ENV{ROCM_PATH})
+        set(ROCM_PATH "/opt/rocm" CACHE PATH "Path to ROCm installation")
+      else()
+        set(ROCM_PATH $ENV{ROCM_PATH} CACHE PATH "Path to ROCm installation")
+      endif()
+    endif()
+
+    # Find hipcc wrapper for reference
+    find_program(HIP_HIPCC_EXECUTABLE
+      NAMES hipcc
+      PATHS "${ROCM_PATH}/bin" "${ROCM_PATH}/hip/bin"
+      NO_DEFAULT_PATH
+    )
+
+    if(NOT HIP_HIPCC_EXECUTABLE)
+      message(FATAL_ERROR "Could not find hipcc. Please set ROCM_PATH to your ROCm installation directory.")
+    endif()
+
+    # Find the actual clang compiler used by ROCm (required by CMake)
+    find_program(CMAKE_HIP_COMPILER
+      NAMES clang++
+      PATHS "${ROCM_PATH}/llvm/bin" "${ROCM_PATH}/bin"
+      NO_DEFAULT_PATH
+    )
+
+    if(NOT CMAKE_HIP_COMPILER)
+      message(FATAL_ERROR "Could not find ROCm clang++ compiler in ${ROCM_PATH}")
+    endif()
+
+    set(HIP_DIR "${ROCM_PATH}")
+    set(CMAKE_HIP_ARCHITECTURES "gfx803;gfx900;gfx906;gfx908;gfx90a;gfx1030;gfx1100" CACHE STRING "HIP architectures")
+
+    # Enable HIP language support
+    enable_language(HIP)
+
+    add_library(hybridapi ${hybridAPI-cxx-sources})
+    add_dependencies(hybridapi ci-generated)
+    target_include_directories(hybridapi PRIVATE "${ROCM_PATH}/include")
+
+    if(TRACING)
+      target_compile_definitions(hybridapi PRIVATE HAPI_TRACE)
+    endif()
   endif()
+
+  # hapi_memory_daemon - standalone executable for shrink/expand GPU memory management
+  if(BUILD_CUDA)
+    add_executable(hapi_memory_daemon src/arch/cuda/hybridAPI/hapi_memory_daemon.cpp)
+    add_dependencies(hapi_memory_daemon hybridapi ck converse ckqt moduleNDMeshStreamer ckmain modulecompletion conv-static)
+  endif()
+
 endif()
 
 if(EXISTS ${CMAKE_SOURCE_DIR}/src/arch/${VDIR}/conv-mach-cxi.sh)
@@ -911,6 +999,12 @@ if(${CMK_BUILD_CHARMRUN})
   add_dependencies(charmrun create_symlinks)
 else()
   configure_file(${CHARMRUN_DIR}/charmrun ${CMAKE_BINARY_DIR}/bin COPYONLY)
+  if(EXISTS ${CMAKE_SOURCE_DIR}/${CHARMRUN_ELASTIC_DIR}/charmrun_elastic)
+    configure_file(${CHARMRUN_ELASTIC_DIR}/charmrun_elastic ${CMAKE_BINARY_DIR}/bin COPYONLY)
+  endif()
+  if(EXISTS ${CMAKE_SOURCE_DIR}/${CHARMRUN_HAPI_DIR}/charmrun_hapi)
+    configure_file(${CHARMRUN_HAPI_DIR}/charmrun_hapi ${CMAKE_BINARY_DIR}/bin COPYONLY)
+  endif()
 endif()
 configure_file(src/scripts/testrun bin/ COPYONLY)
 
@@ -1000,7 +1094,11 @@ if(${TARGET} STREQUAL "charm4py")
   endif()
 
   if (${BUILD_CUDA})
-    target_link_libraries(charm cudart cudahybridapi)
+    target_link_libraries(charm cudart hybridapi)
+  endif()
+
+  if (${BUILD_HIP})
+    target_link_libraries(charm hiprtc hybridapi)
   endif()
 
   if(${TRACING})
@@ -1019,9 +1117,12 @@ else()
   if(RECONVERSE)
     target_link_libraries(ckhello PRIVATE reconverse)
   endif()
-  add_dependencies(ckhello ck ckqt conv-static
+  add_dependencies(ckhello ck ckqt ckrescale conv-static
     converse ckmain
     moduleNDMeshStreamer modulecompletion)
+  if(BUILD_CUDA OR BUILD_HIP)
+    add_dependencies(ckhello hybridapi)
+  endif()
 endif()
 
 # Create conv-mach-opt.sh
@@ -1067,7 +1168,7 @@ foreach(l BUILDOPTS CMK_AMPI_WITH_ROMIO CMK_BUILD_PYTHON CMK_CAN_LINK_FORTRAN
         CXX_NO_AS_NEEDED LDXX_WHOLE_ARCHIVE_PRE LDXX_WHOLE_ARCHIVE_POST
         CMK_MACOSX CMK_POST_EXE CMK_SHARED_SUF CMK_USER_SUFFIX OPTS_LD
         CMK_COMPILER_KNOWS_FVISIBILITY CMK_LINKER_KNOWS_UNDEFINED
-        CMK_SUPPORTS_MEMORY_ISOMALLOC CUDA_DIR CMK_USER_DISABLED_TLS CMK_CXI)
+        CMK_SUPPORTS_MEMORY_ISOMALLOC CUDA_DIR HIP_DIR CMK_USER_DISABLED_TLS CMK_CXI)
     file(APPEND ${optfile_sh} "${l}=\"${${l}}\"\n" )
 endforeach(l)
 
@@ -1104,7 +1205,7 @@ endif()
 set(optfile_mak ${CMAKE_BINARY_DIR}/include/conv-mach-opt.mak)
 
 file(WRITE  ${optfile_mak} "# Build-time options header for Makefiles, automatically generated by cmake.\n")
-foreach(l CUDA_DIR BUILD_CUDA CMK_AMPI_WITH_ROMIO CMK_MACOSX CMK_BUILD_PYTHON
+foreach(l CUDA_DIR HIP_DIR BUILD_CUDA BUILD_HIP CMK_AMPI_WITH_ROMIO CMK_MACOSX CMK_BUILD_PYTHON
         CMK_CHARMDEBUG CMK_COMPILER CMK_GDIR CMK_HAS_MALLOC_HOOK CMK_HAS_MMAP CMK_LIBJPEG
         CMK_LUSTREAPI CMK_MULTICORE CMK_NO_BUILD_SHARED CMK_NO_PARTITIONS CMK_SHARED_SUF
         CMK_SMP CMK_SUPPORTS_FSGLOBALS CMK_SUPPORTS_PIPGLOBALS CMK_SUPPORTS_PIEGLOBALS
@@ -1117,7 +1218,8 @@ endforeach(l)
 
 # Add options
 set(CUDA ${BUILD_CUDA}) # need CUDA to match conv-mach file name
-foreach(opt SMP OMP TCP PTHREADS SYNCFT PXSHM PERSISTENT OOC CUDA PAPI CXI)
+set(HIP ${BUILD_HIP})   # need HIP to match conv-mach file name
+foreach(opt SMP OMP TCP PTHREADS SYNCFT PXSHM PERSISTENT OOC CUDA HIP PAPI CXI)
     if(${opt})
         string(TOLOWER ${opt} optl)
         file(APPEND ${optfile_sh} ". ${CMAKE_BINARY_DIR}/include/conv-mach-${optl}.sh\n")

diff --git a/Dockerfile b/Dockerfile
@@ -0,0 +1,25 @@
+FROM mpioperator/openmpi
+
+RUN apt update && apt install -y build-essential zlib1g-dev ca-certificates cmake git
+
+RUN apt update \
+    && apt install -y --no-install-recommends \
+        g++ \
+        gfortran \
+        libopenmpi-dev \
+    && rm -rf /var/lib/apt/lists/*
+
+#RUN git clone https://github.com/charmplusplus/charm.git
+RUN mkdir /home/mpiuser/charm
+COPY . /home/mpiuser/charm
+RUN cd charm && git checkout shrinkexpand-mpi && ./build charm++ mpi-linux-x86_64 --enable-shrinkexpand -j8 --force --with-production
+
+RUN cd charm/examples/charm++/shrink_expand && make clean && make
+RUN cd charm/examples/charm++/shrink_expand/jacobi2d-iter && make clean && make
+RUN cd charm/examples/charm++/shrink_expand/startup && make clean && make
+RUN mkdir /app
+RUN cp charm/examples/charm++/shrink_expand/jacobi2d-iter/charmrun /app/
+RUN cp charm/examples/charm++/shrink_expand/jacobi2d-iter/charmrun_elastic /app/
+RUN cp charm/examples/charm++/shrink_expand/jacobi2d-iter/jacobi2d /app/
+RUN cp charm/examples/charm++/shrink_expand/startup/startup /app/
+RUN chmod 777 /app
diff --git a/benchmarks/charm++/cuda/gpudirect/latency/latency.C b/benchmarks/charm++/cuda/gpudirect/latency/latency.C
@@ -125,6 +125,7 @@ public:
   int peer;
 
   double start_time;
+  double allocTime;
   double* times;
 
   char* h_local_data;
@@ -211,6 +212,13 @@ public:
       cudaStreamSynchronize(stream);
       thisProxy[peer].receiveReg(size, h_local_data);
     } else {
+      double allocStart = CkWallTimer();
+      char* d_local_data_new;
+      hapiCheck(cudaMalloc(&d_local_data_new, max_size));
+      hapiCheck(cudaFree(d_local_data));
+      d_local_data = d_local_data_new;
+      send_buffer = CkDeviceBuffer(d_local_data_new);
+      allocTime = CkWallTimer() - allocStart;
       thisProxy[peer].receiveZC(size, send_buffer);
     }
   }
@@ -230,7 +238,7 @@ public:
     // Inform the runtime where the incoming data should be stored
     // and which CUDA stream should be used for the transfer
     data = d_remote_data;
-    devicePost[0].cuda_stream = stream; // Not used with UCX
+    devicePost[0].hapi_stream = stream; // Not used with UCX
   }
 
   // Second receive (regular entry method), invoked once the data transfers complete
@@ -247,7 +255,7 @@ public:
     } else {
       // PE 0: received pong
       if (iter > warmup_iters) {
-        times[iter-warmup_iters-1] = (CkWallTimer() - start_time) / 2.0;
+        times[iter-warmup_iters-1] = (CkWallTimer() - start_time) / 2.0 - allocTime;
       }
 
       // Start next iteration or end test for current size

diff --git a/buildcmake b/buildcmake
@@ -103,6 +103,7 @@ opt_ccs=0
 opt_charmdebug=0
 opt_controlpoint=0
 opt_cuda=0
+opt_hip=0
 opt_destination=""
 opt_disabletls=0
 opt_install_prefix=""
@@ -181,6 +182,9 @@ function parse_platform_compilers() {
     cuda)
       opt_cuda=1
       ;;
+    hip)
+      opt_hip=1
+      ;;
     cxi)
       opt_cxi=1
       ;;
@@ -681,6 +685,7 @@ CC=$opt_CC CXX=$opt_CXX FC=$opt_FC cmake "$my_srcdir" \
   -DCHARMDEBUG="$opt_charmdebug" \
   -DCONTROLPOINT="$opt_controlpoint" \
   -DBUILD_CUDA="$opt_cuda" \
+  -DBUILD_HIP="$opt_hip" \
   -DDISABLE_TLS="$opt_disabletls" \
   -DDRONE_MODE="$opt_drone_mode" \
   -DENABLE_FORTRAN=$opt_enable_fortran \

diff --git a/cmake/converse.cmake b/cmake/converse.cmake
@@ -240,7 +240,8 @@ add_library(charm_cxx_utils STATIC
 
 add_library(topomanager STATIC
     ${tmgr-cxx-sources}
-    ${tmgr-h-sources})
+    ${tmgr-h-sources}
+    $<TARGET_OBJECTS:ckrescale>)
 
 target_include_directories(topomanager PUBLIC
     src/util/topomanager
@@ -253,7 +254,7 @@ target_include_directories(topomanager PUBLIC
 #     charm_cxx_utils
 # )
 add_custom_target(converse)
-add_dependencies(converse reconverse topomanager charm_cxx_utils)
+add_dependencies(converse reconverse topomanager charm_cxx_utils ckrescale)
 
 #file(MAKE_DIRECTORY ${CMAKE_BINARY_DIR}/include/comm_backend)
 

diff --git a/cmake/fetch_reconverse/CMakeLists.txt b/cmake/fetch_reconverse/CMakeLists.txt
@@ -27,5 +27,6 @@ FetchContent_MakeAvailable(reconverse)
 set(BUILD_SHARED_LIBS ${_save_BUILD_SHARED_LIBS} CACHE INTERNAL "")
 
 configure_file(${reconverse_SOURCE_DIR}/include/converse.h ${CMAKE_BINARY_DIR}/include/ COPYONLY)
+configure_file(${reconverse_SOURCE_DIR}/include/conv-rdma.h ${CMAKE_BINARY_DIR}/include/ COPYONLY)
 configure_file(${reconverse_SOURCE_DIR}/src/cldb.h ${CMAKE_BINARY_DIR}/include/ COPYONLY)
 configure_file(${reconverse_SOURCE_DIR}/include/charm-config.h ${CMAKE_BINARY_DIR}/include/ COPYONLY)
diff --git a/doc/charm++/manual.rst b/doc/charm++/manual.rst
@@ -9314,7 +9314,8 @@ This entry method should be invoked on the sender by wrapping the
 source buffer with ``CkDeviceBuffer``, whose constructor takes a pointer
 to the source buffer, a Charm++ callback to be invoked once the transfer
 completes (optional), and a CUDA stream associated with the transfer
-(which is only used internally in the CUDA memcpy and IPC based implementation and is also optional):
+(which is only used internally in the CUDA memcpy and IPC based implementation and is also optional).
+The user guarantees that the GPU buffer won't be modified until the callback is called:
 
 .. code-block:: c++
 

diff --git a/examples/ampi/Cjacobi3D/Makefile b/examples/ampi/Cjacobi3D/Makefile
@@ -1,7 +1,7 @@
 -include ../../common.mk
 -include ../../../include/conv-mach-opt.mak
 CHARMBASE=../../../
-CHARMC=../../../bin/ampicxx $(OPTS)
+CHARMC=../../../netlrts-linux-x86_64/bin/ampicxx $(OPTS)
 TOKENS=6
 
 -include $(CHARMBASE)/include/conv-mach-opt.mak
@@ -12,6 +12,7 @@ AMPI_TARGETS := \
   jacobi \
   jacobi.pup \
   jacobi-get \
+  jacobi.pie
 
 ifeq (1,$(CMK_SUPPORTS_TLSGLOBALS))
   AMPI_TARGETS += jacobi.tls
@@ -47,6 +48,10 @@ jacobi.tls: jacobi.C
 	$(CHARMC) -c -tlsglobals jacobi.C -o jacobi.tls.o
 	$(CHARMC) -o jacobi.tls jacobi.tls.o -tlsglobals
 
+jacobi.pie: jacobi-pie.C
+	$(CHARMC) -c -pieglobals jacobi-pie.C -o jacobi.pie.o
+	$(CHARMC) -o jacobi.pie jacobi.pie.o -pieglobals
+
 jacobi.rose: jacobi.C
 	$(CHARMC) -roseomptlsglobals -o jacobi.rose.o -c $<
 	$(CHARMC) -roseomptlsglobals -o $@ jacobi.rose.o
@@ -93,5 +98,5 @@ endif
 
 
 clean:
-	rm -f *.o jacobi *~ moduleinit.C charmrun conv-host jacobi-cpp jacobi.iso jacobi-get jacobi.tls ampirun
+	rm -f *.o jacobi *~ moduleinit.C charmrun conv-host jacobi-cpp jacobi.iso jacobi-get jacobi.tls jacobi.pie ampirun
 	rm -rf 40 80 120