Live attach/detach and its unit tests (#53)

projects/rocprofiler-compute/CHANGELOG.md

@@ -5,6 +5,9 @@ Full documentation for ROCm Compute Profiler is available at [https://rocm.docs.
 ## Unreleased
 
 ### Added
+* Live attach/detach feature that allows coupling with a workload process, without controlling its start or end.
+  * Use '--attach-pid' to specify the target process ID.
+  * Use '--attach-duration-msec' to specify time duration.
 
 * Add `rocpd` choice for `--format-rocprof-output` option in profile mode
 

projects/rocprofiler-compute/CMakeLists.txt

@@ -322,6 +322,15 @@ add_test(
     WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
 )
 
+add_test(
+    NAME test_profile_live_attach_detach
+    COMMAND
+        ${Python3_EXECUTABLE} -m pytest -s -m live_attach_detach
+        --junitxml=tests/test_profile_live_attach_detach.xml ${COV_OPTION}
+        ${PROJECT_SOURCE_DIR}/tests/test_profile_general.py
+    WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}
+)
+
 set_tests_properties(
     test_profile_kernel_execution
     test_profile_dispatch

projects/rocprofiler-compute/docs/how-to/live_attach_detach.rst

@@ -0,0 +1,46 @@
+.. meta::
+   :description: ROCm Compute Profiler: using Live Attach Detach
+   :keywords: ROCm Compute Profiler, Attach Detach
+
+***********************************************************
+Using Live Attach/Detach in ROCm Compute Profiler
+***********************************************************
+
+Live Attach/Detach is a new feature of ROCm Compute Profiler that allows coupling with a workload process, without controlling its start or end. The application can already be running before the profiler application is invoked. The profiler simply attaches to the process, collects the required counters, and then detaches—without altering the lifecycle of the workload.
+
+A specific attach is not repeatable, and it can only collect the set of counters that the hardware is capable of capturing in a single run. As such, in the current implementation, you must specify a subset of counter groups that can be collected within one run. This can be done either by using the ``--block`` option (for example, --block 3.1.1 4.1.1 5.1.1) or by providing a predefined set through the use of single pass counter collection ``--set``.
+
+Detachment can be achieved in two ways:
+
+a) By setting the ``--attach-duration-msec`` parameter to a specific duration (in milliseconds). In this case, the detachment occurs automatically after the specified time has elapsed since the ``rocprof`` subprocess started.
+b) By pressing the Enter key after a successful attach within the same profiling terminal session. Upon a successful attach, a confirmation message is displayed in the terminal log of the workload process.
+
+---------------------
+Profiling options
+---------------------
+For using profiling options for PC sampling the configuration needed are:
+
+* ``--attach-pid``: Should be the process ID of the process of workload's application.
+* ``--attach-duration-msec``: (Optional) This is for setting up the synchronized detach, and is optional. Its unit is in milliseconds. When setting up, the detach will happen after this time has elapsed since the ``rocprof`` subprocess started. For example, setting it to 60000 yields 1 minute.
+
+**Sample command:**
+
+.. code-block:: shell
+   $ rocprof-compute profile -n try_live_attach_detach -b 3.1.1 4.1.1 5.1.1 --no-roof -VVV --attach-pid <process id of workload>
+
+   $ rocprof-compute profile -n try_live_attach_detach --set launch_stats --no-roof -VVV --attach-pid <process id of workload>
+
+   $ rocprof-compute profile -n try_live_attach_detach -b 3.1.1 4.1.1 5.1.1 --no-roof -VVV --attach-pid <process id of workload> --attach-duration-msec <time before detach>
+
+   $ rocprof-compute profile -n try_live_attach_detach --set launch_stats --no-roof -VVV --attach-pid <process id of workload> --attach-duration-msec <time before detach>
+
+-----------------------
+Analysis options
+-----------------------
+The analyze options for attach/detach are completely compatible with the non-attach/detach option.
+
+.. note::
+
+  * Live Attach Detach feature is currently in BETA version. To enable Live/Attach Detach, you need to have the correct supported proper version of ROCprofiler-SDK and rocprofiler-register.
+  * To make the Live Attach/Detach feature work, you must use "--block" or a single path to limit the number of counter input files to one. This limitation will be removed in a later version with implementations such as Iteration Multiplexing.
+  * Due to the limitation of ROCprofiler-SDK, the attach can now only happen before Heterogeneous System Architecture (HSA) initialization. HSA initialization happens before the execution of the first HIP kernel call. It only happens once to save all the kernels' function signature, such as the function name and other launch parameters. Attaching after this stage misses all crucial information of the HIP kernel and makes it impossible to store the output. This limitation will be solved in later releases of ROCprofiler-SDK.

projects/rocprofiler-compute/pyproject.toml

@@ -107,6 +107,7 @@ markers = [
     "sets_func",
     "sets_perf",
     "pc_sampling",
+    "live_attach_detach",
     "roofline",
     "path",
     "sci_notion",

projects/rocprofiler-compute/sample/dynamic_shared/dynamic_shared.hip

@@ -0,0 +1,174 @@
+// MIT License
+//
+// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#include "example_utils.hpp"
+
+#include <hip/hip_runtime.h>
+
+#include <iostream>
+#include <vector>
+
+#include <cstddef>
+#include <cstdlib>
+#include <thread>
+
+/// \brief A simple matrix transpose kernel that using dynamic shared memory.
+/// - The number of rows in the input and output matrices is equal, and given by the \p width parameter.
+/// - Each thread in the grid is responsible for one element of the input and output matrices.
+/// - Because the transposition is computed in shared memory, which cannot be accessed between different
+/// blocks, the matrix has to be processed by a single block.
+__global__ void matrix_transpose_kernel(float* out, const float* in, const unsigned int width)
+{
+    // Declare that this kernel is using dynamic shared memory to store a number of floats.
+    // The unsized array type indicates that the total amount of memory that is going
+    // to be used here is not known ahead of time, and will be computed at runtime and
+    // passed to the kernel launch function.
+    extern __shared__ float shared_matrix_memory[];
+
+    // Compute the row and column index of the element this thread is going to process.
+    const unsigned int x = blockDim.x * blockIdx.x + threadIdx.x;
+    const unsigned int y = blockDim.y * blockIdx.y + threadIdx.y;
+
+    // Perform the transpose by reading an element of the input matrix from global memory and
+    // by storing it in the tranposed index in shared memory.
+    shared_matrix_memory[y * width + x] = in[x * width + y];
+
+    // Synchronization is required to make sure that all threads have written
+    // their part of the input matrix to the shared memory, before the values
+    // are read by another thread.
+    __syncthreads();
+
+    // Copy the transposed matrix from shared memory to the output array, which
+    // is in global memory.
+    out[y * width + x] = shared_matrix_memory[y * width + x];
+}
+
+// CPU implementation of matrix transpose
+std::vector<float> matrix_transpose_reference(const std::vector<float>& input,
+                                              const unsigned int        width)
+{
+    std::vector<float> output(width * width);
+    for(unsigned int j = 0; j < width; j++)
+    {
+        for(unsigned int i = 0; i < width; i++)
+        {
+            output[i * width + j] = input[j * width + i];
+        }
+    }
+    return output;
+}
+
+int main()
+{
+    // Number of rows and columns in the transposed square matrix.
+    constexpr unsigned int width = 4;
+
+    // Number of threads in each kernel block along the X dimension.
+    // Because each thread will process exactly one element, this value
+    // is equal to the width of the matrix.
+    constexpr unsigned int threads_per_block_x = width;
+
+    // Number of threads in each kernel block along the Y dimension.
+    // Because each thread will process exactly one element, this value
+    // is equal to the width of the matrix.
+    constexpr unsigned int threads_per_block_y = width;
+
+    // Total element count of the transposed matrix.
+    constexpr unsigned int size = width * width;
+
+    // Total size (in bytes) of the transposed matrix.
+    constexpr size_t size_bytes = sizeof(float) * size;
+
+    // Total amount of shared memory that each block is going to use.
+    // Exactly one matrix will be stored in shared memory.
+    constexpr size_t shared_memory_bytes = size_bytes;
+    std::cout << "Run transpose continuously" << std::endl;
+
+    // Set a timer to 30 seconds for rocprofv3 preparation
+    std::this_thread::sleep_for(std::chrono::seconds(30));
+    
+    while (true) 
+    {
+        std::this_thread::sleep_for(std::chrono::seconds(5));
+        // Allocate host vectors.
+        std::vector<float> h_matrix(size);
+        std::vector<float> h_transposed_matrix(size);
+
+        // Set up input data.
+        for(unsigned int i = 0; i < size; i++)
+        {
+            h_matrix[i] = i * 10.0f;
+        }
+
+        // Allocate device memory for the input and output matrices.
+        float* d_matrix{};
+        float* d_transposed_matrix{};
+        HIP_CHECK(hipMalloc(&d_matrix, size_bytes));
+        HIP_CHECK(hipMalloc(&d_transposed_matrix, size_bytes));
+
+        // Transfer the input matrix to the device memory.
+        HIP_CHECK(hipMemcpy(d_matrix, h_matrix.data(), size_bytes, hipMemcpyHostToDevice));
+
+        // Lauching kernel from host.
+        matrix_transpose_kernel<<<dim3(width / threads_per_block_x, width / threads_per_block_y),
+                                dim3(threads_per_block_x, threads_per_block_y),
+                                shared_memory_bytes,
+                                hipStreamDefault>>>(d_transposed_matrix, d_matrix, width);
+
+        // Check if the kernel launch was successful.
+        HIP_CHECK(hipGetLastError());
+
+        // Transfer the result back to the host.
+        HIP_CHECK(hipMemcpy(h_transposed_matrix.data(),
+                            d_transposed_matrix,
+                            size_bytes,
+                            hipMemcpyDeviceToHost));
+
+        // Free the resources on the device.
+        HIP_CHECK(hipFree(d_matrix));
+        HIP_CHECK(hipFree(d_transposed_matrix));
+
+        // Perform the reference (CPU) calculation.
+        std::vector<float> ref_transposed_matrix = matrix_transpose_reference(h_matrix, width);
+
+        // Check the results' validity.
+        constexpr float eps = 1.0E-6f;
+        unsigned int    errors{};
+        for(unsigned int i = 0; i < size; i++)
+        {
+            if(std::fabs(h_transposed_matrix[i] - ref_transposed_matrix[i]) > eps)
+            {
+                errors++;
+            }
+        }
+
+        if(errors != 0)
+        {
+            std::cout << "Validation failed. Errors: " << errors << std::endl;
+            return error_exit_code;
+        }
+        else
+        {
+            std::cout << "Validation passed." << std::endl;
+        }
+    }
+}

projects/rocprofiler-compute/sample/dynamic_shared/example_utils.hpp

@@ -0,0 +1,300 @@
+// MIT License
+//
+// Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
+//
+// Permission is hereby granted, free of charge, to any person obtaining a copy
+// of this software and associated documentation files (the "Software"), to deal
+// in the Software without restriction, including without limitation the rights
+// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+// copies of the Software, and to permit persons to whom the Software is
+// furnished to do so, subject to the following conditions:
+//
+// The above copyright notice and this permission notice shall be included in all
+// copies or substantial portions of the Software.
+//
+// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+// SOFTWARE.
+
+#ifndef COMMON_EXAMPLE_UTILS_HPP
+#define COMMON_EXAMPLE_UTILS_HPP
+
+// Compiling HIP on Windows includes windows.h, and this triggers many silly warnings.
+#include <cstdint>
+#if defined(_WIN32) && defined(__NVCC__)
+    #pragma nv_diag_suppress 108 // signed bit field of length 1
+    #pragma nv_diag_suppress 174 // expression has no effect
+    #pragma nv_diag_suppress 1835 // attribute "dllimport" does not apply here
+#endif
+
+// rocPRIM adds a #warning about printf on NAVI.
+#ifdef __clang__
+    #pragma clang diagnostic ignored "-W#warnings"
+#endif
+
+#include <algorithm>
+#include <cassert>
+#include <chrono>
+#include <iomanip>
+#include <iostream>
+#include <iterator>
+#include <sstream>
+#include <string>
+#include <type_traits>
+#include <vector>
+
+#include <hip/hip_runtime.h>
+
+constexpr int error_exit_code = -1;
+
+/// \brief Checks if the provided error code is \p hipSuccess and if not,
+/// prints an error message to the standard error output and terminates the program
+/// with an error code.
+#define HIP_CHECK(condition)                                                                \
+    {                                                                                       \
+        const hipError_t error = condition;                                                 \
+        if(error != hipSuccess)                                                             \
+        {                                                                                   \
+            std::cerr << "An error encountered: \"" << hipGetErrorString(error) << "\" at " \
+                      << __FILE__ << ':' << __LINE__ << std::endl;                          \
+            std::exit(error_exit_code);                                                     \
+        }                                                                                   \
+    }
+
+/// \brief Formats a range of elements to a pretty string.
+/// \tparam BidirectionalIterator - must implement the BidirectionalIterator concept and
+/// must be dereferencable in host code. Its value type must be formattable to
+/// \p std::ostream.
+template<class BidirectionalIterator>
+inline std::string format_range(const BidirectionalIterator begin, const BidirectionalIterator end)
+{
+    std::stringstream sstream;
+    sstream << "[ ";
+    for(auto it = begin; it != end; ++it)
+    {
+        sstream << *it;
+        if(it != std::prev(end))
+        {
+            sstream << ", ";
+        }
+    }
+    sstream << " ]";
+    return sstream.str();
+}
+
+/// \brief Formats a range of pairs to a pretty string. The length of the two ranges must match.
+/// \tparam BidirectionalIteratorT - must implement the BidirectionalIterator concept and
+/// must be dereferencable in host code. Its value type must be formattable to \p std::ostream.
+/// \tparam BidirectionalIteratorU - must implement the BidirectionalIterator concept and
+/// must be dereferencable in host code. Its value type must be formattable to \p std::ostream.
+template<class BidirectionalIteratorT, typename BidirectionalIteratorU>
+inline std::string format_pairs(const BidirectionalIteratorT begin_a,
+                                const BidirectionalIteratorT end_a,
+                                const BidirectionalIteratorU begin_b,
+                                const BidirectionalIteratorU end_b)
+{
+    (void)end_b;
+    assert(std::distance(begin_a, end_a) == std::distance(begin_b, end_b));
+
+    std::stringstream sstream;
+    sstream << "[ ";
+    auto it_a = begin_a;
+    auto it_b = begin_b;
+    for(; it_a < end_a; ++it_a, ++it_b)
+    {
+        sstream << "(" << *it_a << ", " << *it_b << ")";
+
+        if(it_a != std::prev(end_a))
+        {
+            sstream << ", ";
+        }
+    }
+    sstream << " ]";
+    return sstream.str();
+}
+
+/// \brief A function to parse a string for an int. If the string is a valid integer then return true
+/// else if it has non-numeric character then return false.
+inline bool parse_int_string(const std::string& str, int& out)
+{
+    try
+    {
+        size_t end;
+        int    value = std::stoi(str, &end);
+        if(end == str.size())
+        {
+            out = value;
+            return true;
+        }
+        return false;
+    }
+    catch(const std::exception&)
+    {
+        return false;
+    }
+}
+
+/// \brief A class to measures time between intervals
+class HostClock
+{
+private:
+    std::chrono::steady_clock::time_point start_time;
+    std::chrono::steady_clock::duration   elapsed_time;
+
+public:
+    HostClock()
+    {
+        this->reset_timer();
+    }
+
+    inline void reset_timer()
+    {
+        this->elapsed_time = std::chrono::steady_clock::duration(0);
+    }
+
+    inline void start_timer()
+    {
+        this->start_time = std::chrono::steady_clock::now();
+    }
+
+    inline void stop_timer()
+    {
+        const auto end_time = std::chrono::steady_clock::now();
+        this->elapsed_time += end_time - this->start_time;
+    }
+
+    /// @brief Returns time elapsed in Seconds
+    /// @return type double that contains the elapsed time in Seconds
+    inline double get_elapsed_time() const
+    {
+        return std::chrono::duration_cast<std::chrono::duration<double>>(this->elapsed_time)
+            .count();
+    }
+};
+
+/// \brief Returns <tt>ceil(dividend / divisor)</tt>, where \p dividend is an integer and
+/// \p divisor is an unsigned integer.
+template<typename T,
+         typename U,
+         std::enable_if_t<std::is_integral<T>::value && std::is_unsigned<U>::value, int> = 0>
+__host__ __device__ constexpr auto ceiling_div(const T& dividend, const U& divisor)
+{
+    return (dividend + divisor - 1) / divisor;
+}
+
+/// \brief Report validation results.
+inline int report_validation_result(int errors)
+{
+    if(errors)
+    {
+        std::cout << "Validation failed. Errors: " << errors << std::endl;
+        return error_exit_code;
+    }
+
+    std::cout << "Validation passed." << std::endl;
+    return 0;
+}
+
+/// \brief Generate an identity matrix.
+/// The identity matrix is a $m \times n$ matrix with ones in the main diagonal and zeros elsewhere.
+template<typename T>
+void generate_identity_matrix(T* A, int m, int n, size_t lda)
+{
+    for(int i = 0; i < m; ++i)
+    {
+        for(int j = 0; j < n; ++j)
+        {
+            A[i + j * lda] = T(i == j);
+        }
+    }
+}
+
+/// \brief Multiply an $A$ matrix ($m \times k$) with a $B$ matrix ($k \times n$) as:
+/// $C := \alpha \cdot A \cdot B + \beta \cdot C$
+template<typename T>
+void multiply_matrices(T        alpha,
+                       T        beta,
+                       int      m,
+                       int      n,
+                       int      k,
+                       const T* A,
+                       int      stride1_a,
+                       int      stride2_a,
+                       const T* B,
+                       int      stride1_b,
+                       int      stride2_b,
+                       T*       C,
+                       int      stride_c)
+{
+    for(int i1 = 0; i1 < m; ++i1)
+    {
+        for(int i2 = 0; i2 < n; ++i2)
+        {
+            T t = T(0.0);
+            for(int i3 = 0; i3 < k; ++i3)
+            {
+                t += A[i1 * stride1_a + i3 * stride2_a] * B[i3 * stride1_b + i2 * stride2_b];
+            }
+            C[i1 + i2 * stride_c] = beta * C[i1 + i2 * stride_c] + alpha * t;
+        }
+    }
+}
+
+/// \brief Prints an {1,2,3}-dimensional array. The last dimension (fastest-index) specified in
+/// \p n will be printed horizontally.
+///
+/// By default a row-major layout of the data is assumed. When printing data in column-major
+/// layout, the \p column_major parameter must be set to \p true for a correct interpretation
+/// of the dimensions' sizes.
+template<class Tdata, class Tsize>
+void print_nd_data(const std::vector<Tdata>& data,
+                   std::vector<Tsize>        np,
+                   const int                 column_width = 4,
+                   const bool                column_major = false)
+{
+    if(column_major)
+    {
+        std::reverse(np.begin(), np.end());
+    }
+    const std::vector<Tsize> n(np);
+    // Note: we want to print the last dimension horizontally (on the x-axis)!
+    int size_x = n[n.size() - 1];
+    int size_y = n.size() > 1 ? n[n.size() - 2] : 1;
+    int size_z = n.size() > 2 ? n[n.size() - 3] : 1;
+    for(int z = 0; z < size_z; ++z)
+    {
+        for(int y = 0; y < size_y; ++y)
+        {
+            for(int x = 0; x < size_x; ++x)
+            {
+                auto index = (z * size_y + y) * size_x + x;
+                std::cout << std::setfill(' ') << std::setw(column_width) << data[index] << " ";
+            }
+            std::cout << "\n";
+        }
+        if(z != size_z - 1)
+        {
+            std::cout << "\n";
+        }
+    }
+    std::cout << std::flush;
+}
+
+/// \brief Returns a string from the double \p value with specified \p precision .
+inline std::string
+    double_precision(const double value, const int precision, const bool fixed = false)
+{
+    std::stringstream ss;
+    if(fixed)
+    {
+        ss << std::fixed;
+    }
+    ss << std::setprecision(precision) << value;
+    return ss.str();
+}
+
+#endif // COMMON_EXAMPLE_UTILS_HPP

projects/rocprofiler-compute/src/argparser.py

@@ -146,6 +146,28 @@ Examples:
     profile_group.add_argument(
         "--target", type=str, default=None, help=argparse.SUPPRESS
     )
+    profile_group.add_argument(
+        "--attach-pid",
+        type=str,
+        dest="attach_pid",
+        metavar="",
+        default=None,
+        required=False,
+        help="\t\t\tProcess id to be attached for profiling.",
+    )
+    profile_group.add_argument(
+        "--attach-duration-msec",
+        type=str,
+        dest="attach_duration_msec",
+        metavar="",
+        default=None,
+        required=False,
+        help=(
+            "\t\t\tWhen --attach-pid is used, it specifies the attach duration "
+            "in milliseconds. If not set, detachment occurs when "
+            '"Enter" key is pressed.'
+        ),
+    )
     profile_group.add_argument(
         "-p",
         "--path",

projects/rocprofiler-compute/src/rocprof_compute_profile/profiler_base.py

@@ -109,7 +109,7 @@ class RocProfCompute_Base:
                     "Please verify."
                 )
             args.remaining = " ".join(args.remaining)
-        else:
+        elif not args.attach_pid:
             console_error(
                 "Profiling command required. Pass application executable after -- "
                 "at the end of options.\n"
@@ -357,6 +357,9 @@ class RocProfCompute_Base:
         args = self.get_args()
         console_debug("profiling", f"pre-processing using {self.__profiler} profiler")
 
+        if args.attach_pid:
+            args.remaining = ""
+
         self._filter_blocks = self._soc.profiling_setup()
 
         # Write profiling configuration as yaml file
@@ -472,6 +475,22 @@ class RocProfCompute_Base:
             ):
                 options = self.get_profiler_options(str(fname), self._soc)
                 start_time = time.time()
+
+                # Only 1-run case is permitted for attach/detach
+                if (isinstance(options, list) and "--pid" in options) or (
+                    isinstance(options, dict)
+                    and (options.get("ROCPROF_ATTACH_PID") is not None)
+                ):
+                    if total_runs > 1:
+                        console_error(
+                            f"Cannot attach process for profiling as the requested "
+                            f"performance counters exceed the collection capacity of "
+                            f"single pass counter collection. The current setup of "
+                            f"requested counter blocks needs {total_runs} number of "
+                            f'passes. Please use "--block" or "--set" '
+                            f"to adjust or reduce the requested performance metrics!"
+                        )
+
                 run_prof(
                     fname=str(fname),
                     profiler_options=options,

projects/rocprofiler-compute/src/rocprof_compute_profile/profiler_rocprof_v3.py

@@ -72,6 +72,14 @@ class rocprof_v3_profiler(RocProfCompute_Base):
             args.format_rocprof_output,
         ]
 
+        if args.attach_pid:
+            profiling_options.append("--pid")
+            profiling_options.append(args.attach_pid)
+
+            if args.attach_duration_msec:
+                profiling_options.append("--attach-duration-msec")
+                profiling_options.append(args.attach_duration_msec)
+
         # Kernel filtering
         if args.kernel:
             profiling_options.extend(["--kernel-include-regex", "|".join(args.kernel)])
@@ -94,8 +102,9 @@ class rocprof_v3_profiler(RocProfCompute_Base):
                 f"[{','.join(dispatch)}]",
             ])
 
-        profiling_options.append("--")
-        profiling_options.extend(app_cmd)
+        if not args.attach_pid:
+            profiling_options.append("--")
+            profiling_options.extend(app_cmd)
         return profiling_options
 
     # -----------------------

projects/rocprofiler-compute/src/rocprof_compute_profile/profiler_rocprofiler_sdk.py

@@ -57,9 +57,14 @@ class rocprofiler_sdk_profiler(RocProfCompute_Base):
         rocprofiler_sdk_tool_path = str(
             rocm_libdir / "rocprofiler-sdk" / "librocprofiler-sdk-tool.so"
         )
+        rocm_dir = Path(args.rocprofiler_sdk_library_path).parent.parent
+        rocprofiler_attach_tool_path = str(
+            rocm_dir / "lib" / "rocprofiler-sdk" / "librocprofv3-attach.so"
+        )
         ld_preload = [
             rocprofiler_sdk_tool_path,
             args.rocprofiler_sdk_library_path,
+            rocprofiler_attach_tool_path,
         ]
         options = {
             "ROCPROFILER_LIBRARY_CTOR": "1",
@@ -71,6 +76,17 @@ class rocprofiler_sdk_profiler(RocProfCompute_Base):
             "ROCPROF_OUTPUT_PATH": f"{args.path}/out/pmc_1",
         }
 
+        if args.attach_pid:
+            options.update({
+                "ROCPROF_ATTACH_TOOL_LIBRARY": rocprofiler_attach_tool_path,
+                "ROCPROF_ATTACH_PID": args.attach_pid,
+            })
+
+            if args.attach_duration_msec:
+                options.update({
+                    "ROCPROF_ATTACH_DURATION": args.attach_duration_msec,
+                })
+
         if args.kokkos_trace:
             # NOTE: --kokkos-trace feature is incomplete and is disabled for now.
             console_error(
@@ -101,7 +117,8 @@ class rocprofiler_sdk_profiler(RocProfCompute_Base):
 
         if dispatch:
             options["ROCPROF_KERNEL_FILTER_RANGE"] = f"[{','.join(dispatch)}]"
-        options["APP_CMD"] = app_cmd
+        if not args.attach_pid:
+            options["APP_CMD"] = app_cmd
         return options
 
     # -----------------------

projects/rocprofiler-compute/src/utils/utils.py

@@ -24,6 +24,7 @@
 ##############################################################################
 
 import argparse
+import ctypes
 import glob
 import io
 import json
@@ -31,15 +32,19 @@ import locale
 import logging
 import os
 import re
+import select
 import selectors
 import shlex
 import shutil
 import subprocess
+import sys
 import tempfile
+import threading
 import time
 import uuid
+from contextlib import contextmanager
 from pathlib import Path
-from typing import Any, Optional, Union, cast
+from typing import Any, Dict, Generator, Optional, Union, cast
 
 import pandas as pd
 import yaml
@@ -292,22 +297,33 @@ def capture_subprocess_output(
     # Start subprocess
     # bufsize = 1 means output is line buffered
     # universal_newlines = True is required for line buffering
+    sanitized_env = (
+        None
+        if new_env is None
+        else {
+            k: ":".join(str(i) for i in v) if isinstance(v, list) else str(v)
+            for k, v in new_env.items()
+        }
+    )
+
     process = (
         subprocess.Popen(
             subprocess_args,
             bufsize=1,
+            stdin=subprocess.PIPE,
             stdout=subprocess.PIPE,
             stderr=subprocess.STDOUT,
             universal_newlines=True,
         )
-        if new_env == None
+        if sanitized_env == None
         else subprocess.Popen(
             subprocess_args,
             bufsize=1,
+            stdin=subprocess.PIPE,
             stdout=subprocess.PIPE,
             stderr=subprocess.STDOUT,
             universal_newlines=True,
-            env=new_env,
+            env=sanitized_env,
         )
     )
 
@@ -319,6 +335,8 @@ def capture_subprocess_output(
             # Because the process' output is line buffered, there's only ever one
             # line to read when this function is called
             line = stream.readline()
+            if not line:
+                return
             buf.write(line)
             if enable_logging:
                 if profileMode:
@@ -334,6 +352,43 @@ def capture_subprocess_output(
     if process.stdout is not None:
         selector.register(process.stdout, selectors.EVENT_READ, handle_output)
 
+    def forward_input() -> None:
+        """
+        Forward the keyboard input from the terminal to the inside subprocess
+        """
+
+        try:
+            sys.stdin.fileno()
+        except (io.UnsupportedOperation, AttributeError):
+            # Stdin can't be used in select; skip input forwarding
+            return
+
+        if sys.stdin.isatty():
+            for line in sys.stdin:
+                if process.poll() is not None:
+                    break
+                process.stdin.write(line)
+                process.stdin.flush()
+        else:
+            while process.poll() is None:
+                try:
+                    rlist, _, _ = select.select([sys.stdin], [], [], 0.1)
+                except (io.UnsupportedOperation, AttributeError):
+                    break
+                if rlist:
+                    line = sys.stdin.readline()
+                    if not line:
+                        break
+                    process.stdin.write(line)
+                    process.stdin.flush()
+        try:
+            process.stdin.close()
+        except Exception:
+            console_warning("forward_input: the stdin did not close properly!")
+
+    input_thread = threading.Thread(target=forward_input, daemon=True)
+    input_thread.start()
+
     # Loop until subprocess is terminated
     while process.poll() is None:
         # Wait for events and handle them with their registered callbacks
@@ -342,6 +397,8 @@ def capture_subprocess_output(
             callback = key.data
             callback(key.fileobj, mask)
 
+    input_thread.join(timeout=1)
+
     # Get process return code
     return_code = process.wait()
     selector.close()
@@ -693,6 +750,13 @@ def run_prof(
     fbase = fpath.stem
     console_debug(f"pmc file: {fpath.name}")
 
+    is_mode_live_attach = (
+        isinstance(profiler_options, list) and "--pid" in profiler_options
+    ) or (
+        isinstance(profiler_options, dict)
+        and profiler_options.get("ROCPROF_ATTACH_PID") is not None
+    )
+
     # standard rocprof options
     if rocprof_cmd == "rocprofiler-sdk":
         options = cast(dict[str, Union[str, list[str]]], profiler_options)
@@ -707,6 +771,11 @@ def run_prof(
             options["ROCPROF_AGENT_INDEX"] = "absolute"
         else:
             options = ["-A", "absolute"] + options
+    else:
+        if is_mode_live_attach:
+            console_error(
+                "The live attach/detach only supports rocprofv3 or rocprofiler-sdk"
+            )
 
     new_env = os.environ.copy()
 
@@ -754,14 +823,66 @@ def run_prof(
     time_1 = time.time()
 
     if rocprof_cmd == "rocprofiler-sdk":
-        app_cmd = options.pop("APP_CMD")
+        app_cmd = options.pop("APP_CMD") if "APP_CMD" in options else None
         for key, value in options.items():
             new_env[key] = value
         console_debug(f"rocprof sdk env vars: {new_env}")
-        console_debug(f"rocprof sdk user provided command: {app_cmd}")
-        success, output = capture_subprocess_output(
-            app_cmd, new_env=new_env, profileMode=True
-        )
+
+        if is_mode_live_attach:
+
+            @contextmanager
+            def temporary_env(env_vars: Dict[str, str]) -> Generator[None, None, None]:
+                """
+                Temporarily change the environment variable of this application.
+                """
+                original_env = os.environ.copy()
+                os.environ.update({k: str(v) for k, v in env_vars.items()})
+                try:
+                    yield
+                finally:
+                    os.environ.clear()
+                    os.environ.update(original_env)
+
+            with temporary_env(new_env):
+                libname = options["ROCPROF_ATTACH_TOOL_LIBRARY"]
+                c_lib = ctypes.CDLL(libname)
+                if c_lib is None:
+                    console_error(f"Error opening {libname}")
+                c_lib.attach.argtypes = [ctypes.c_uint]
+
+                pid = options["ROCPROF_ATTACH_PID"]
+                if pid is None:
+                    console_error(
+                        "Mode of attach/detach must have setup for process ID"
+                    )
+
+                c_lib.attach(int(pid))
+                duration = os.environ.get("ROCPROF_ATTACH_DURATION", None)
+                if duration is None:
+                    console_log(
+                        f"\033[93mAttach to process with ID {pid} is successful, "
+                        "Press Enter to detach...\033[0m"
+                    )
+                    input()
+                else:
+                    console_log(
+                        f"\033[93mAttach to process with ID {pid} is successful, "
+                        f"detach will happen in {duration} milliseconds...\033[0m"
+                    )
+                    time.sleep(int(duration) / 1000)
+                c_lib.detach()
+
+        else:
+            if app_cmd is None:
+                console_error(
+                    "APP_CMD, the workload's execuatble must be provided "
+                    "when not in live attach mode"
+                )
+
+            console_debug(f"rocprof sdk user provided command: {app_cmd}")
+            success, output = capture_subprocess_output(
+                app_cmd, new_env=new_env, profileMode=True
+            )
     else:
         # print in readable format using shlex
         console_debug(f"rocprof command: {shlex.join([rocprof_cmd] + options)}")
@@ -780,7 +901,7 @@ def run_prof(
     if new_env.get("ROCPROFILER_METRICS_PATH"):
         shutil.rmtree(new_env["ROCPROFILER_METRICS_PATH"], ignore_errors=True)
 
-    if not success:
+    if (not is_mode_live_attach) and (not success):
         if loglevel > logging.INFO:
             for line in output.splitlines():
                 console_error(line, exit=False)

projects/rocprofiler-compute/tests/CMakeLists.txt

@@ -49,3 +49,12 @@ set_target_properties(
     mat_mul_max
     PROPERTIES RUNTIME_OUTPUT_DIRECTORY ${CMAKE_SOURCE_DIR}/tests
 )
+
+set(DYNAMIC_SHARED_SOURCES ../sample/dynamic_shared/dynamic_shared.hip)
+set_source_files_properties(${DYNAMIC_SHARED_SOURCES} PROPERTIES LANGUAGE HIP)
+add_executable(hip_dynamic_shared ${DYNAMIC_SHARED_SOURCES})
+target_include_directories(hip_dynamic_shared PRIVATE ../sample/dynamic_shared)
+set_target_properties(
+    hip_dynamic_shared
+    PROPERTIES RUNTIME_OUTPUT_DIRECTORY ${CMAKE_SOURCE_DIR}/tests
+)

projects/rocprofiler-compute/tests/conftest.py

@@ -59,6 +59,7 @@ def binary_handler_profile_rocprof_compute(request):
         check_success=True,
         roof=False,
         app_name="app_1",
+        attach_detach_para=None,
     ):
         if request.config.getoption("--rocprofiler-sdk-library-path"):
             options.extend(
@@ -77,11 +78,25 @@ def binary_handler_profile_rocprof_compute(request):
             ]
             if not roof:
                 baseline_opts.append("--no-roof")
+
+            command_rocprof_compute = baseline_opts + options + ["--path", workload_dir]
+            if not attach_detach_para:
+                command_rocprof_compute = (
+                    command_rocprof_compute + ["--"] + config[app_name]
+                )
+            else:
+                command_rocprof_compute = command_rocprof_compute + [
+                    "--attach-pid",
+                    str(attach_detach_para["attach_pid"]),
+                ]
+                if attach_detach_para["attach-duration-msec"]:
+                    command_rocprof_compute = command_rocprof_compute + [
+                        "--attach-duration-msec",
+                        str(attach_detach_para["attach-duration-msec"]),
+                    ]
+
             process = subprocess.run(
-                baseline_opts
-                + options
-                + ["--path", workload_dir, "--"]
-                + config[app_name],
+                command_rocprof_compute,
                 text=True,
             )
             # verify run status
@@ -92,13 +107,27 @@ def binary_handler_profile_rocprof_compute(request):
             baseline_opts = ["rocprof-compute", "profile", "-n", app_name, "-VVV"]
             if not roof:
                 baseline_opts.append("--no-roof")
+
+            command_rocprof_compute = baseline_opts + options + ["--path", workload_dir]
+            if not attach_detach_para:
+                command_rocprof_compute = (
+                    command_rocprof_compute + ["--"] + config[app_name]
+                )
+            else:
+                command_rocprof_compute = command_rocprof_compute + [
+                    "--attach-pid",
+                    str(attach_detach_para["attach_pid"]),
+                ]
+                if attach_detach_para["attach-duration-msec"]:
+                    command_rocprof_compute = command_rocprof_compute + [
+                        "--attach-duration-msec",
+                        str(attach_detach_para["attach-duration-msec"]),
+                    ]
+
             with pytest.raises(SystemExit) as e:
                 with patch(
                     "sys.argv",
-                    baseline_opts
-                    + options
-                    + ["--path", workload_dir, "--"]
-                    + config[app_name],
+                    command_rocprof_compute,
                 ):
                     rocprof_compute.main()
             # verify run status

projects/rocprofiler-compute/tests/test_profile_general.py

@@ -60,7 +60,9 @@ CHIP_IDS = {
 config = {}
 config["kernel_name_1"] = "vecCopy"
 config["app_1"] = ["./tests/vcopy", "-n", "1048576", "-b", "256", "-i", "3"]
+config["app_occupancy"] = ["./tests/occupancy"]
 config["app_mat_mul_max"] = ["./tests/mat_mul_max"]
+config["app_hip_dynamic_shared"] = ["./tests/hip_dynamic_shared"]
 config["cleanup"] = True
 config["COUNTER_LOGGING"] = False
 config["METRIC_COMPARE"] = False
@@ -1869,6 +1871,133 @@ def test_pc_sampling_stochastic(binary_handler_profile_rocprof_compute):
     test_utils.clean_output_dir(config["cleanup"], workload_dir)
 
 
+@pytest.mark.live_attach_detach
+def test_live_attach_detach_block(binary_handler_profile_rocprof_compute):
+    if not using_v3():
+        assert True
+        return
+
+    options = ["--block", "3.1.1", "4.1.1", "5.1.1"]
+    workload_dir = test_utils.get_output_dir()
+    process_workload = subprocess.Popen(config["app_hip_dynamic_shared"])
+
+    # set the time to detach here to 1 mins, which is 60000 msec
+    time_to_detach = "60000"
+
+    attach_detach = dict()
+    attach_detach["attach_pid"] = process_workload.pid
+    attach_detach["attach-duration-msec"] = time_to_detach
+
+    _ = binary_handler_profile_rocprof_compute(
+        config,
+        workload_dir,
+        options,
+        check_success=True,
+        roof=False,
+        app_name="app_hip_dynamic_shared",
+        attach_detach_para=attach_detach,
+    )
+
+    # kill the process of the workload at thsi point if it's still running
+    if process_workload.poll() is None:
+        print(
+            f"rocprof-compute has detached and finished, "
+            f"killing workload process (pid={process_workload.pid})..."
+        )
+        process_workload.kill()
+        process_workload.wait()
+
+    file_dict = test_utils.check_csv_files(workload_dir, 1, num_kernels)
+    validate(
+        inspect.stack()[0][3],
+        workload_dir,
+        file_dict,
+    )
+
+    assert test_utils.check_file_pattern(
+        "- 3.1.1", f"{workload_dir}/profiling_config.yaml"
+    )
+    assert test_utils.check_file_pattern(
+        "- 4.1.1", f"{workload_dir}/profiling_config.yaml"
+    )
+    assert test_utils.check_file_pattern(
+        "- 5.1.1", f"{workload_dir}/profiling_config.yaml"
+    )
+    test_utils.clean_output_dir(config["cleanup"], workload_dir)
+
+
+@pytest.mark.live_attach_detach
+def test_live_attach_detach_singlepath_launch_stats(
+    binary_handler_profile_rocprof_compute,
+):
+    if not using_v3():
+        assert True
+        return
+
+    options = ["--set", "launch_stats"]
+    workload_dir = test_utils.get_output_dir()
+    process_workload = subprocess.Popen(config["app_hip_dynamic_shared"])
+
+    # set the time to detach here to 1 mins, which is 60000 msec
+    time_to_detach = "60000"
+
+    attach_detach = dict()
+    attach_detach["attach_pid"] = process_workload.pid
+    attach_detach["attach-duration-msec"] = time_to_detach
+
+    _ = binary_handler_profile_rocprof_compute(
+        config,
+        workload_dir,
+        options,
+        check_success=True,
+        roof=False,
+        app_name="app_hip_dynamic_shared",
+        attach_detach_para=attach_detach,
+    )
+
+    # kill the process of the workload at thsi point if it's still running
+    if process_workload.poll() is None:
+        print(
+            f"rocprof-compute has detached and finished, "
+            f"killing workload process (pid={process_workload.pid})..."
+        )
+        process_workload.kill()
+        process_workload.wait()
+
+    file_dict = test_utils.check_csv_files(workload_dir, 1, num_kernels)
+    validate(
+        inspect.stack()[0][3],
+        workload_dir,
+        file_dict,
+    )
+
+    assert test_utils.check_file_pattern(
+        "- 7.1.0", f"{workload_dir}/profiling_config.yaml"
+    )
+    assert test_utils.check_file_pattern(
+        "- 7.1.1", f"{workload_dir}/profiling_config.yaml"
+    )
+    assert test_utils.check_file_pattern(
+        "- 7.1.2", f"{workload_dir}/profiling_config.yaml"
+    )
+    assert test_utils.check_file_pattern(
+        "- 7.1.5", f"{workload_dir}/profiling_config.yaml"
+    )
+    assert test_utils.check_file_pattern(
+        "- 7.1.6", f"{workload_dir}/profiling_config.yaml"
+    )
+    assert test_utils.check_file_pattern(
+        "- 7.1.7", f"{workload_dir}/profiling_config.yaml"
+    )
+    assert test_utils.check_file_pattern(
+        "- 7.1.8", f"{workload_dir}/profiling_config.yaml"
+    )
+    assert test_utils.check_file_pattern(
+        "- 7.1.9", f"{workload_dir}/profiling_config.yaml"
+    )
+    test_utils.clean_output_dir(config["cleanup"], workload_dir)
+
+
 @pytest.mark.sets_func
 class TestSetsIntegration:
     def test_memory_throughput_set(self, binary_handler_profile_rocprof_compute):