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

projects/rocprofiler-systems/CHANGELOG.md

@@ -12,6 +12,8 @@ Full documentation for ROCm Systems Profiler is available at [https://rocm.docs.
   the profiler output. Enable collection with `ROCPROFSYS_USE_HIPFILE` and select metrics with
   `ROCPROFSYS_HIPFILE_METRICS`. See
   [hipFile Infinity Storage I/O telemetry](./docs/how-to/hipfile-telemetry.rst).
+- `--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.
 - `rocprof-sys-attach` finds the tool library in the target's own ROCm installation when the
   target cannot see this installation's library (for example, inside a container), so
   `ROCPROF_ATTACH_TOOL_LIBRARY` is no longer needed there.

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,9 +103,16 @@ 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
@@ -117,8 +126,6 @@ run(int rank, int tid, int devid, 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;
@@ -150,19 +157,22 @@ run(int rank, int tid, int devid, 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
@@ -214,7 +224,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);
         }
@@ -223,8 +237,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)
@@ -246,18 +271,30 @@ main(int argc, char** argv)
         devid = rank % ndevice;
         HIP_API_CALL(hipSetDevice(devid));
         printf("[transpose] Rank %i assigned to device %i\n", rank, devid);
-
-        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, devid, _streams.at(i), argc, argv);
-        run(rank, 0, devid, _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)));
+    }
+    if(rank == devid && rank < ndevice)
+    {
+        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, devid, _streams.at(i), argc, argv);
+            run(rank, 0, devid, _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

@@ -9,6 +9,7 @@
 #include <fmt/ranges.h>
 
 #include <algorithm>
+#include <chrono>
 #include <cstdlib>
 #include <cstring>
 #include <exception>
@@ -17,6 +18,7 @@
 #include <optional>
 #include <sstream>
 #include <string>
+#include <thread>
 #include <vector>
 
 #include <rocprofiler-sdk-rocattach/rocattach.h>
@@ -25,32 +27,38 @@ namespace
 {
 struct attach_options
 {
-    int                      pid = -1;
+    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
@@ -189,6 +197,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)
 {
@@ -241,6 +267,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);
@@ -307,8 +339,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/conftest.py

@@ -345,6 +345,7 @@ def pytest_configure(config: pytest.Config) -> None:
         "selective_regions",
         "minimal",
         "rank_filter",
+        "rocattach",
         "pytest_impl",
     ]
     for label in non_functional_markers + generic_functional_markers:
@@ -1709,6 +1710,7 @@ def _subrow(label: str, value) -> str:
         _row("Avail:", rocprof_config.rocprofsys_avail),
         _row("Causal:", rocprof_config.rocprofsys_causal),
         _row("Python:", rocprof_config.rocprofsys_python),
+        _row("Attach:", rocprof_config.rocprofsys_attach),
         "-" * 70,
         "Executables:",
         _row("MPI:", cap.mpiexec_exec),

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,7 @@ class RocprofsysConfig:
     rocprofsys_sample: Path
     rocprofsys_causal: Path
     rocprofsys_avail: Path
+    rocprofsys_attach: Optional[Path]
     rocm_path: Optional[Path]
     rocprofsys_lib_dir: Path
     rocprofsys_bin_dir: Path
@@ -477,6 +480,8 @@ def discover_install_config(
 
     search_paths = [bin_dir]
     sys_execs = _find_rocprofsys_core_executables(search_paths)
+    # rocprof-sys-attach is optional: tests that need it should skip when None.
+    rocprof_attach = _find_executable("rocprof-sys-attach", search_paths)
     rocprofsys_python, rocprofsys_site_packages = _find_rocprofsys_python(
         search_paths, install_dir
     )
@@ -488,6 +493,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=rocprof_attach,
         rocm_path=rocm_path,
         rocprofsys_lib_dir=lib_dir,
         rocprofsys_bin_dir=bin_dir,
@@ -562,6 +568,8 @@ def discover_build_config(
 
     search_paths = [bin_dir]
     sys_execs = _find_rocprofsys_core_executables(search_paths)
+    # rocprof-sys-attach is optional: tests that need it should skip when None.
+    rocprof_attach = _find_executable("rocprof-sys-attach", search_paths)
     rocprofsys_python, rocprofsys_site_packages = _find_rocprofsys_python(
         search_paths, build_dir
     )
@@ -579,6 +587,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=rocprof_attach,
         rocm_path=rocm_path,
         rocprofsys_lib_dir=lib_dir,
         rocprofsys_bin_dir=bin_dir,