A HIP program to monitor MaxRSS

kernel-loop-test.hip

#include <hip/hip_runtime.h>
#include <iostream>
#include <sys/resource.h>

__global__ void k() {}

size_t getMaxRSS() {
  rusage Rusage;
  if (getrusage(RUSAGE_SELF, &Rusage) != 0)
    exit(3);
  return Rusage.ru_maxrss;
}

int main() {
  constexpr size_t N = 1ull << 63;

  size_t PrevMaxRSS = 0;
  size_t PrevI = 0;
  for (size_t I = 0; I < N; I++) {
    k<<<1, 1>>>();
    hipDeviceSynchronize();

    if (I % 10000 == 0)
      std::cerr << ".";

    auto CurrMaxRSS = getMaxRSS();
    if (CurrMaxRSS > PrevMaxRSS) {
      auto Diff = CurrMaxRSS - PrevMaxRSS;
      double Rate = (I > PrevI) ? (double(Diff) / (I - PrevI)) : Diff;
      std::cerr << "\nMaxRSS=" << CurrMaxRSS << " KB, diff=" << Diff
                << ", rate=" << Rate << " KB/launch\n";
      PrevI = I;
      PrevMaxRSS = CurrMaxRSS;
    }
  }
  std::cerr << "\n";
  return 0;
}