[rocprofiler-systems] Add pytest for rocprof-sys-attach (attach/detach/re-attach)

projects/rocprofiler-systems/CHANGELOG.md

@@ -37,6 +37,9 @@ Full documentation for ROCm Systems Profiler is available at [https://rocm.docs.
 - Add `--list-domains` and `--list-operations <domain>` to `rocprof-sys-avail`.
   There new options allow the user to query more information about available
   ROCm domains (used in `ROCPROFSYS_ROCM_DOMAINS`) and their operations.
+- `--detach-after <seconds>` flag on `rocprof-sys-attach` for non-interactive use; the binary sleeps for the given duration after attach instead of waiting for `ENTER` on stdin.
+- Infinite mode for the `transpose` example: passing `0` for `NUM_ITERATION` loops indefinitely until SIGTERM/SIGINT is received, useful as a long-lived target for attach/detach scenarios.
+- Pytest coverage for the `rocprof-sys-attach` (rocattach) codepath in `tests/pytest/test_rocattach.py`, parametrized over 1, 2, and 3 attach sessions against a single transpose target.
 
 ### Changed
 

projects/rocprofiler-systems/examples/transpose/transpose.cpp

@@ -22,9 +22,11 @@ THE SOFTWARE.
 
 #include "hip/hip_runtime.h"
 
+#include <atomic>
 #include <cfloat>
 #include <chrono>
 #include <cmath>
+#include <csignal>
 #include <cstdio>
 #include <cstdlib>
 #include <fstream>
@@ -101,18 +103,23 @@ transpose_a(int* in, int* out, int M, int N)
 
 namespace
 {
-size_t nthreads = 2;
-size_t nitr     = 500;
-size_t nsync    = 10;
+size_t            nthreads = 2;
+size_t            nitr     = 500;  // 0 means run until signaled (infinite mode)
+size_t            nsync    = 10;
+std::atomic<bool> stop_requested{ false };
+
+void
+handle_stop_signal(int /*signum*/) noexcept
+{
+    stop_requested.store(true, std::memory_order_relaxed);
+}
 }  // namespace
 
 void
-run(int rank, int tid, hipStream_t stream, int argc, char** argv)
+run(int rank, int tid, hipStream_t stream, int /*argc*/, char** /*argv*/)
 {
     unsigned int M = 4960 * 2;
     unsigned int N = 4960 * 2;
-    if(argc > 2) nitr = atoll(argv[2]);
-    if(argc > 3) nsync = atoll(argv[3]);
 
     auto_lock_t _lk{ print_lock };
     std::cout << "[" << rank << "][" << tid << "] M: " << M << " N: " << N << std::endl;
@@ -144,19 +151,22 @@ run(int rank, int tid, hipStream_t stream, int argc, char** argv)
     dim3 grid(M / 32, N / 32, 1);
     dim3 block(32, 32, 1);  // transpose_a
 
-    auto t1 = std::chrono::high_resolution_clock::now();
+    size_t completed_itr = 0;
+    auto   t1            = std::chrono::high_resolution_clock::now();
     for(size_t i = 0; i < nitr; ++i)
     {
+        if(stop_requested.load(std::memory_order_relaxed)) break;
         transpose_a<<<grid, block, 0, stream>>>(in, out, M, N);
         check_hip_error();
         if(i % nsync == (nsync - 1)) HIP_API_CALL(hipStreamSynchronize(stream));
+        ++completed_itr;
     }
     auto t2 = std::chrono::high_resolution_clock::now();
     HIP_API_CALL(hipStreamSynchronize(stream));
     HIP_API_CALL(hipMemcpyAsync(out_matrix, out, size, hipMemcpyDeviceToHost, stream));
     double time =
         std::chrono::duration_cast<std::chrono::duration<double>>(t2 - t1).count();
-    float GB = static_cast<float>(size) * nitr * 2 / (1 << 30);
+    float GB = static_cast<float>(size) * completed_itr * 2 / (1 << 30);
 
     print_lock.lock();
     std::cout << "[" << rank << "][" << tid << "] Runtime of transpose is " << time
@@ -208,7 +218,11 @@ main(int argc, char** argv)
         {
             fprintf(stderr,
                     "usage: transpose [NUM_THREADS (%zu)] [NUM_ITERATION (%zu)] "
-                    "[SYNC_EVERY_N_ITERATIONS (%zu)]\n",
+                    "[SYNC_EVERY_N_ITERATIONS (%zu)]\n"
+                    "\n"
+                    "  NUM_ITERATION = 0 enables infinite-workload mode: each round\n"
+                    "  runs the default 500 iterations end-to-end (alloc, kernel\n"
+                    "  loop, verify, free) and rounds repeat until SIGTERM/SIGINT.\n",
                     nthreads, nitr, nsync);
             exit(EXIT_SUCCESS);
         }
@@ -217,8 +231,19 @@ main(int argc, char** argv)
     if(argc > 2) nitr = atoll(argv[2]);
     if(argc > 3) nsync = atoll(argv[3]);
 
+    const bool infinite_workload = (nitr == 0);
+    if(infinite_workload) nitr = 500;
+
+    std::signal(SIGTERM, handle_stop_signal);
+    std::signal(SIGINT, handle_stop_signal);
+
     printf("[transpose] Number of threads: %zu\n", nthreads);
-    printf("[transpose] Number of iterations: %zu\n", nitr);
+    if(infinite_workload)
+        printf("[transpose] Number of iterations per round: %zu (infinite "
+               "rounds until SIGTERM/SIGINT)\n",
+               nitr);
+    else
+        printf("[transpose] Number of iterations: %zu\n", nitr);
     printf("[transpose] Syncing every %zu iterations\n", nsync);
 
 #if defined(USE_MPI)
@@ -241,17 +266,27 @@ main(int argc, char** argv)
     }
     if(rank == devid && rank < ndevice)
     {
-        std::vector<std::thread> _threads{};
-        std::vector<hipStream_t> _streams(nthreads);
-        for(size_t i = 0; i < nthreads; ++i)
-            HIP_API_CALL(hipStreamCreate(&_streams.at(i)));
-        for(size_t i = 1; i < nthreads; ++i)
-            _threads.emplace_back(run, rank, i, _streams.at(i), argc, argv);
-        run(rank, 0, _streams.at(0), argc, argv);
-        for(auto& itr : _threads)
-            itr.join();
-        for(size_t i = 0; i < nthreads; ++i)
-            HIP_API_CALL(hipStreamDestroy(_streams.at(i)));
+        size_t round = 0;
+        do
+        {
+            if(infinite_workload)
+            {
+                print_lock.lock();
+                std::cout << "[transpose] Starting round " << round++ << std::endl;
+                print_lock.unlock();
+            }
+            std::vector<std::thread> _threads{};
+            std::vector<hipStream_t> _streams(nthreads);
+            for(size_t i = 0; i < nthreads; ++i)
+                HIP_API_CALL(hipStreamCreate(&_streams.at(i)));
+            for(size_t i = 1; i < nthreads; ++i)
+                _threads.emplace_back(run, rank, i, _streams.at(i), argc, argv);
+            run(rank, 0, _streams.at(0), argc, argv);
+            for(auto& itr : _threads)
+                itr.join();
+            for(size_t i = 0; i < nthreads; ++i)
+                HIP_API_CALL(hipStreamDestroy(_streams.at(i)));
+        } while(infinite_workload && !stop_requested.load(std::memory_order_relaxed));
     }
     HIP_API_CALL(hipDeviceSynchronize());
     HIP_API_CALL(hipDeviceReset());

projects/rocprofiler-systems/source/bin/rocprof-sys-attach/rocprof-sys-attach.cpp

@@ -8,13 +8,15 @@
 #include <spdlog/fmt/ranges.h>
 
 #include <algorithm>
+#include <chrono>
 #include <cstdlib>
 #include <cstring>
 #include <exception>
 #include <iostream>
 #include <limits>
 #include <sstream>
 #include <string>
+#include <thread>
 #include <vector>
 
 #include <rocprofiler-sdk-rocattach/rocattach.h>
@@ -24,31 +26,37 @@ namespace
 struct attach_options
 {
     int                      pid            = -1;
+    int                      detach_after   = -1;  // seconds; -1 = wait for stdin ENTER
     std::string              output_path    = {};
     std::vector<std::string> profile_format = {};
 };
 
 void
 print_usage(const char* prog_name)
 {
-    std::cout << "Usage: " << prog_name << " -p <pid> [OPTIONS]\n"
-              << "\n"
-              << "Attach to a running process for profiling.\n"
-              << "\n"
-              << "Options:\n"
-              << "  -p <pid>             Process ID to attach to (required)\n"
-              << "  -o, --output PATH    Output path for profiling results\n"
-              << "  -F, --format FORMAT[,FORMAT,...]\n"
-              << "                       Output format(s): perfetto, rocpd\n"
-              << "  -h, --help           Show this help message\n"
-              << "\n"
-              << "Environment variables:\n"
-              << "  ROCPROFSYS_OUTPUT_PATH       Output directory for profiling data\n"
-              << "  ROCPROFSYS_TRACE             Enable perfetto trace output\n"
-              << "  ROCPROFSYS_USE_ROCPD         Enable rocpd database output\n"
-              << "  ROCPROF_ATTACH_TOOL_LIBRARY  Path to the tool library\n"
-              << "\n"
-              << "Once attached, press ENTER to detach from the process.\n";
+    std::cout
+        << "Usage: " << prog_name << " -p <pid> [OPTIONS]\n"
+        << "\n"
+        << "Attach to a running process for profiling.\n"
+        << "\n"
+        << "Options:\n"
+        << "  -p <pid>             Process ID to attach to (required)\n"
+        << "  -o, --output PATH    Output path for profiling results\n"
+        << "  -F, --format FORMAT[,FORMAT,...]\n"
+        << "                       Output format(s): perfetto, rocpd\n"
+        << "  -d, --detach-after SECONDS\n"
+        << "                       Detach automatically after SECONDS "
+           "(non-interactive).\n"
+        << "                       When omitted, waits for ENTER on stdin.\n"
+        << "  -h, --help           Show this help message\n"
+        << "\n"
+        << "Environment variables:\n"
+        << "  ROCPROFSYS_OUTPUT_PATH       Output directory for profiling data\n"
+        << "  ROCPROFSYS_TRACE             Enable perfetto trace output\n"
+        << "  ROCPROFSYS_USE_ROCPD         Enable rocpd database output\n"
+        << "  ROCPROF_ATTACH_TOOL_LIBRARY  Path to the tool library\n"
+        << "\n"
+        << "Once attached, press ENTER to detach from the process (interactive mode).\n";
 }
 
 void
@@ -156,6 +164,24 @@ parse_pid(attach_options& opts, const char* arg)
     }
 }
 
+void
+parse_detach_after(attach_options& opts, const char* arg)
+{
+    try
+    {
+        opts.detach_after = std::stoi(arg);
+        if(opts.detach_after < 0)
+        {
+            LOG_ERROR("--detach-after must be non-negative.");
+            std::exit(EXIT_FAILURE);
+        }
+    } catch(const std::exception&)
+    {
+        LOG_ERROR("Invalid --detach-after value '{}'.", arg);
+        std::exit(EXIT_FAILURE);
+    }
+}
+
 void
 parse_formats(attach_options& opts, const char* arg)
 {
@@ -208,6 +234,12 @@ parse_args(int argc, char* argv[])
             continue;
         }
 
+        if(is_option(arg, "-d", "--detach-after"))
+        {
+            parse_detach_after(opts, consume_arg(i, argc, argv, "-d/--detach-after"));
+            continue;
+        }
+
         LOG_ERROR("Unknown option '{}'.", arg);
         print_usage(argv[0]);
         std::exit(EXIT_FAILURE);
@@ -265,8 +297,17 @@ main(int argc, char* argv[])
         return EXIT_FAILURE;
     }
 
-    LOG_INFO("Attached to process {}. Press ENTER to detach.", pid);
-    std::cin.ignore(std::numeric_limits<std::streamsize>::max(), '\n');
+    if(opts.detach_after >= 0)
+    {
+        LOG_INFO("Attached to process {}. Detaching after {} second(s).", pid,
+                 opts.detach_after);
+        std::this_thread::sleep_for(std::chrono::seconds(opts.detach_after));
+    }
+    else
+    {
+        LOG_INFO("Attached to process {}. Press ENTER to detach.", pid);
+        std::cin.ignore(std::numeric_limits<std::streamsize>::max(), '\n');
+    }
 
     result = rocattach_detach(pid);
     if(result != ROCATTACH_STATUS_SUCCESS)

projects/rocprofiler-systems/tests/pytest/rocprofsys/config.py

@@ -28,6 +28,8 @@ class RocprofsysConfig:
         - rocprofsys_sample: Path to rocprof-sys-sample executable
         - rocprofsys_causal: Path to rocprof-sys-causal executable
         - rocprofsys_avail: Path to rocprof-sys-avail executable
+        - rocprofsys_attach: Path to rocprof-sys-attach executable (optional;
+            None if the binary was not built)
         - rocm_path: Path to ROCm installation directory
         - rocprofsys_lib_dir: Path to rocprofsys library directory
         - rocprofsys_bin_dir: Path to rocprofsys binary directory
@@ -52,6 +54,8 @@ class RocprofsysConfig:
     rocprofsys_sample: Path
     rocprofsys_causal: Path
     rocprofsys_avail: Path
+    rocprofsys_attach: Optional[Path]
+    rocm_path: Path
     rocm_path: Optional[Path]
     rocprofsys_lib_dir: Path
     rocprofsys_bin_dir: Path
@@ -403,6 +407,7 @@ def _find_rocprofsys_core_executables(
     rocprof_run = _find_executable("rocprof-sys-run", search_paths)
     rocprof_causal = _find_executable("rocprof-sys-causal", search_paths)
     rocprof_avail = _find_executable("rocprof-sys-avail", search_paths)
+    rocprof_attach = _find_executable("rocprof-sys-attach", search_paths)
 
     required_executables = {
         "rocprof-sys-instrument": rocprof_instrument,
@@ -419,7 +424,8 @@ def _find_rocprofsys_core_executables(
             f"Searched in: {search_paths}"
         )
 
-    return required_executables
+    # rocprof-sys-attach is optional: tests that need it should skip when None.
+    return {**required_executables, "rocprof-sys-attach": rocprof_attach}
 
 
 def _find_rocprofsys_python(
@@ -567,6 +573,7 @@ def discover_install_config(
         rocprofsys_sample=sys_execs["rocprof-sys-sample"],
         rocprofsys_causal=sys_execs["rocprof-sys-causal"],
         rocprofsys_avail=sys_execs["rocprof-sys-avail"],
+        rocprofsys_attach=sys_execs["rocprof-sys-attach"],
         rocm_path=rocm_path,
         rocprofsys_lib_dir=lib_dir,
         rocprofsys_bin_dir=bin_dir,
@@ -669,6 +676,7 @@ def discover_build_config(
         rocprofsys_sample=sys_execs["rocprof-sys-sample"],
         rocprofsys_causal=sys_execs["rocprof-sys-causal"],
         rocprofsys_avail=sys_execs["rocprof-sys-avail"],
+        rocprofsys_attach=sys_execs["rocprof-sys-attach"],
         rocm_path=rocm_path,
         rocprofsys_lib_dir=lib_dir,
         rocprofsys_bin_dir=bin_dir,