Program startup slowness with HIP

[BradWhitlock]

I noticed that several shaping tests with the shaping_driver took about 11 seconds on rzadams vs 2 seconds on rzwhippet.

I dug into this a bit and found that most of that time (~9 seconds) was spent in axom::getDefaultAllocatorID in the course of setting up the sidre::DataStore. That function just calls the Umpire ResourceManager to get the default allocator. ResourceManager queries the HIP runtime to determine the number of GPUs, etc.

I hacked the shaping_driver with a timed block to call hipGetDeviceCount() before the DataStore creation and the 9 seconds of time moved to this first touch of HIP.

Next, I had codex make a short standalone program that did the types of queries that Umpire was doing. The program was way faster, ~245ms instead of 9s.

Next, I had codex start linking that standalone program with additional Axom dependencies until it showed a slowdown in HIP calls. The hipblas library caused the highest jump in program startup time. Maybe there is some static initialization?

I believe Axom is getting this dependency from MFEM, which makes sense since MFEM is needed for the shaping_driver.

There might not be much we can do about this but I thought it was worth recording.

Timings

================================================================
case: baseline-hip
program: hip
command: hip
start: Wed Jul 22 12:10:31 PDT 2026
HIP/hipBLAS startup reproducer
build=hip
[begin] program total
[begin] hipGetDeviceCount (baseline)
[end]   hipGetDeviceCount (baseline) : 254.035 ms
  baseline device count: 4
HIP device count: 4
  device 0: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
  device 1: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
  device 2: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
  device 3: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
[end]   program total : 254.070 ms
real 0.30
user 0.06
sys 0.20
exit_status: 0
end: Wed Jul 22 12:10:31 PDT 2026

================================================================
case: hipblas-linked
program: hip_hipblas
command: hip_hipblas
start: Wed Jul 22 12:10:31 PDT 2026
HIP/hipBLAS startup reproducer
build=hip+hipblas
hipblas_smoke=0
[begin] program total
[begin] hipGetDeviceCount (baseline)
[end]   hipGetDeviceCount (baseline) : 4645.007 ms
  baseline device count: 4
HIP device count: 4
  device 0: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
  device 1: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
  device 2: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
  device 3: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
[end]   program total : 4645.454 ms
real 4.82
user 2.18
sys 2.60
exit_status: 0
end: Wed Jul 22 12:10:36 PDT 2026

================================================================
case: hipblas-smoke
program: hip_hipblas
command: hip_hipblas --hipblas-smoke
start: Wed Jul 22 12:10:36 PDT 2026
HIP/hipBLAS startup reproducer
build=hip+hipblas
hipblas_smoke=1
[begin] program total
[begin] hipblasCreate/Destroy
[end]   hipblasCreate/Destroy : 4520.662 ms
[begin] hipGetDeviceCount (baseline)
[end]   hipGetDeviceCount (baseline) : 0.008 ms
  baseline device count: 4
HIP device count: 4
  device 0: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
  device 1: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
  device 2: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
  device 3: AMD Radeon Graphics, totalGlobalMem=137438953472, gcnArchName=gfx942:sramecc+:xnack-
[end]   program total : 4520.772 ms
real 4.70
user 1.56
sys 3.02
exit_status: 0
end: Wed Jul 22 12:10:41 PDT 2026

hip.cpp

#include <hip/hip_runtime.h>

#if defined(HIP_REPRO_ENABLE_HIPBLAS)
#include <hipblas/hipblas.h>
#endif

#include <chrono>
#include <cstdlib>
#include <iomanip>
#include <iostream>
#include <sstream>
#include <stdexcept>
#include <string>

namespace {

using clock_type = std::chrono::steady_clock;

struct Options
{
  bool run_hipblas_smoke {false};
  std::string program_name {"./hip"};
};

struct StepTimer
{
  explicit StepTimer(std::string label)
    : m_label(std::move(label))
    , m_start(clock_type::now())
  {
    std::cout << "[begin] " << m_label << '\n';
    std::cout.flush();
  }

  ~StepTimer()
  {
    const auto end = clock_type::now();
    const auto ms = std::chrono::duration<double, std::milli>(end - m_start).count();
    std::cout << "[end]   " << m_label << " : " << std::fixed << std::setprecision(3) << ms << " ms\n";
    std::cout.flush();
  }

private:
  std::string m_label;
  clock_type::time_point m_start;
};

void check_hip(hipError_t err, const std::string& what)
{
  if(err != hipSuccess) {
    std::ostringstream oss;
    oss << what << " failed: " << hipGetErrorString(err) << " (" << static_cast<int>(err) << ")";
    throw std::runtime_error(oss.str());
  }
}

#if defined(HIP_REPRO_ENABLE_HIPBLAS)
void check_hipblas(hipblasStatus_t err, const std::string& what)
{
  if(err != HIPBLAS_STATUS_SUCCESS) {
    std::ostringstream oss;
    oss << what << " failed: hipBLAS status " << static_cast<int>(err);
    throw std::runtime_error(oss.str());
  }
}

void run_hipblas_smoke()
{
  StepTimer step {"hipblasCreate/Destroy"};
  hipblasHandle_t handle {nullptr};
  check_hipblas(hipblasCreate(&handle), "hipblasCreate");
  check_hipblas(hipblasDestroy(handle), "hipblasDestroy");
}
#endif

Options parse_args(int argc, char** argv)
{
  Options opts;
  if(argc > 0 && argv[0] != nullptr) {
    opts.program_name = argv[0];
  }

  for(int i = 1; i < argc; ++i) {
    const std::string arg {argv[i]};
    if(arg == "--help" || arg == "-h") {
      std::cout << "Usage: " << opts.program_name << " [options]\n"
                << "  --help           Show this message.\n";
#if defined(HIP_REPRO_ENABLE_HIPBLAS)
      std::cout << "  --hipblas-smoke  Time hipblasCreate/Destroy before hipGetDeviceCount.\n";
#endif
      std::exit(0);
    }
#if defined(HIP_REPRO_ENABLE_HIPBLAS)
    else if(arg == "--hipblas-smoke") {
      opts.run_hipblas_smoke = true;
    }
#endif
    else {
      std::ostringstream oss;
      oss << "Unknown argument: " << arg;
      throw std::runtime_error(oss.str());
    }
  }

  return opts;
}

int run_baseline_probe()
{
  int device_count {0};
  {
    StepTimer step {"hipGetDeviceCount (baseline)"};
    const hipError_t err = hipGetDeviceCount(&device_count);
    if(err == hipErrorNoDevice) {
      std::cout << "  baseline device count: 0\n";
      std::cout << "  HIP runtime reports no ROCm-capable device on this node.\n";
      std::cout.flush();
      return 0;
    }
    check_hip(err, "hipGetDeviceCount");
  }

  std::cout << "  baseline device count: " << device_count << '\n';
  std::cout << "HIP device count: " << device_count << '\n';
  for(int device = 0; device < device_count; ++device) {
    hipDeviceProp_t props {};
    check_hip(hipGetDeviceProperties(&props, device), "hipGetDeviceProperties");
    std::cout << "  device " << device << ": " << props.name
              << ", totalGlobalMem=" << props.totalGlobalMem
              << ", gcnArchName=" << props.gcnArchName << '\n';
  }
  std::cout.flush();

  return device_count;
}

}  // namespace

int main(int argc, char** argv)
{
  try {
    const Options opts = parse_args(argc, argv);

    std::cout << "HIP/hipBLAS startup reproducer\n";
    std::cout << "build=hip";
#if defined(HIP_REPRO_ENABLE_HIPBLAS)
    std::cout << "+hipblas";
#endif
    std::cout << '\n';
#if defined(HIP_REPRO_ENABLE_HIPBLAS)
    std::cout << "hipblas_smoke=" << opts.run_hipblas_smoke << '\n';
#endif
    std::cout.flush();

    StepTimer total {"program total"};

#if defined(HIP_REPRO_ENABLE_HIPBLAS)
    if(opts.run_hipblas_smoke) {
      run_hipblas_smoke();
    }
#endif

    (void)run_baseline_probe();
    return 0;
  } catch(const std::exception& ex) {
    std::cerr << "ERROR: " << ex.what() << '\n';
    return 1;
  }
}

Makefile

ROCM_PATH ?= /opt/rocm-6.4.3
CXX := $(ROCM_PATH)/llvm/bin/amdclang++
SRC := hip.cpp

OFFLOAD_ARCH ?= gfx90a gfx942
ARCH_FLAGS := $(foreach arch,$(OFFLOAD_ARCH),--offload-arch=$(arch))

CXXFLAGS ?= -O2 -g -std=c++17
CPPFLAGS ?= -I$(ROCM_PATH)/include
LDFLAGS ?= -L$(ROCM_PATH)/lib -Wl,-rpath,$(ROCM_PATH)/lib
LDLIBS ?= -lamdhip64
HIPFLAGS ?= -x hip --hip-link
NO_AS_NEEDED := -Wl,--no-as-needed
AS_NEEDED := -Wl,--as-needed

.PHONY: all clean list-targets

all: hip hip_hipblas

list-targets:
	@printf '%s\n' hip hip_hipblas

hip: $(SRC)
	$(CXX) $(CPPFLAGS) $(CXXFLAGS) $(ARCH_FLAGS) $(HIPFLAGS) $< -o $@ $(LDFLAGS) $(LDLIBS)

hip_hipblas: $(SRC)
	$(CXX) $(CPPFLAGS) -DHIP_REPRO_ENABLE_HIPBLAS $(CXXFLAGS) $(ARCH_FLAGS) $(HIPFLAGS) $< -o $@ $(LDFLAGS) $(LDLIBS) $(NO_AS_NEEDED) -lhipblas $(AS_NEEDED)

clean:
	rm -f hip hip_hipblas

Run script

#!/usr/bin/env bash

set -u

SCRIPT_DIR="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
cd "${SCRIPT_DIR}"

TIMESTAMP="$(date +%Y%m%d_%H%M%S)"
OUTFILE="${SCRIPT_DIR}/hipblas_reproducer_${TIMESTAMP}.log"

{
  echo "hipblas reproducer run"
  echo "date: $(date)"
  echo "host: $(hostname)"
  echo "cwd: ${SCRIPT_DIR}"
  echo "outfile: ${OUTFILE}"
  echo
  echo "environment:"
  env | grep -E '^(ROCR|HIP|HSA|ROCP|CALI|CALIPER|UMPIRE)=' || true
  echo
} > "${OUTFILE}"

run_one() {
  local label="$1"
  shift
  local prog="$1"
  shift || true

  local -a args=("$@")
  local -a cmd=("${SCRIPT_DIR}/${prog}" "${args[@]}")

  {
    echo "================================================================"
    echo "case: ${label}"
    echo "program: ${prog}"
    echo "command: ${cmd[*]}"
    echo "start: $(date)"
  } >> "${OUTFILE}"

  if [[ ! -x "${SCRIPT_DIR}/${prog}" ]]; then
    {
      echo "status: missing executable"
      echo
    } >> "${OUTFILE}"
    return 0
  fi

  /usr/bin/time -p "${cmd[@]}" >> "${OUTFILE}" 2>&1
  local status=$?

  {
    echo "exit_status: ${status}"
    echo "end: $(date)"
    echo
  } >> "${OUTFILE}"
}

run_one "baseline-hip" hip
run_one "hipblas-linked" hip_hipblas
run_one "hipblas-smoke" hip_hipblas --hipblas-smoke

echo "wrote ${OUTFILE}"

[BradWhitlock]

Update: I tried rebuilding this test program with ROCM_PATH ?= /opt/rocm-7.2.1 in the Makefile and that produced binaries that did not experience the slow startup in multiple runs. Maybe we just need to update our compiler (and library set).

[rhornung67]

That's interesting. ROCm 7.14 will be available soon (there is a beta version on our machines now). IIRC, AMD folks are encouraging codes to move to that. Should we wait until codes are on 7.14?