tests: add hip graph bubbles coverage

Author: anushe@amd.com

projects/rocprofiler-sdk/tests/CMakeLists.txt

@@ -88,6 +88,7 @@ add_subdirectory(rocdecode)
 add_subdirectory(rocjpeg)
 add_subdirectory(hip-host-tracing)
 add_subdirectory(code-object-multi-threaded)
+add_subdirectory(hip-graph-bubbles)
 
 # rocpd validation tests
 add_subdirectory(rocpd)

projects/rocprofiler-sdk/tests/bin/hip-graph-bubbles/hip-graph-bubbles.cpp

@@ -24,6 +24,7 @@ THE SOFTWARE.
 #include <rocprofiler-sdk-roctx/roctx.h>
 
 #include <chrono>
+#include <cstdlib>
 #include <iomanip>
 #include <iostream>
 #include <vector>
@@ -47,19 +48,79 @@ simpleKernel(int* data, int value)
     data[idx] = value + idx;
 }
 
+namespace
+{
+struct config
+{
+    int num_kernels       = 2000;
+    int num_iterations    = 200;
+    int array_size        = 256;
+    int progress_interval = 50;
+};
+
+void
+print_usage(const char* argv0)
+{
+    std::cerr << "Usage: " << argv0
+              << " [num_kernels] [num_iterations] [array_size] [progress_interval]"
+              << std::endl;
+}
+
 int
-main()
+parse_positive_integer(const char* arg_name, const char* value)
 {
-    const int NUM_KERNELS    = 2000;
-    const int NUM_ITERATIONS = 200;
-    const int ARRAY_SIZE     = 256;
+    char* end = nullptr;
+    auto  ret = std::strtol(value, &end, 10);
+    if(end == value || (end != nullptr && *end != '\0') || ret <= 0)
+    {
+        std::cerr << "Invalid " << arg_name << ": " << value << std::endl;
+        exit(EXIT_FAILURE);
+    }
+
+    return static_cast<int>(ret);
+}
 
-    std::cout << "Creating HIP graph with " << NUM_KERNELS << " kernel launches" << std::endl;
-    std::cout << "Will execute graph " << NUM_ITERATIONS << " times" << std::endl;
+config
+parse_args(int argc, char** argv)
+{
+    config cfg{};
+
+    if(argc > 5)
+    {
+        print_usage(argv[0]);
+        exit(EXIT_FAILURE);
+    }
+
+    if(argc > 1) cfg.num_kernels = parse_positive_integer("num_kernels", argv[1]);
+    if(argc > 2) cfg.num_iterations = parse_positive_integer("num_iterations", argv[2]);
+    if(argc > 3) cfg.array_size = parse_positive_integer("array_size", argv[3]);
+    if(argc > 4)
+        cfg.progress_interval = parse_positive_integer("progress_interval", argv[4]);
+
+    if(cfg.array_size < 256)
+    {
+        std::cerr << "array_size must be at least 256, got " << cfg.array_size << std::endl;
+        exit(EXIT_FAILURE);
+    }
+
+    return cfg;
+}
+}  // namespace
+
+int
+main(int argc, char** argv)
+{
+    const auto cfg = parse_args(argc, argv);
+
+    std::cout << "Creating HIP graph with " << cfg.num_kernels << " kernel launches"
+              << std::endl;
+    std::cout << "Will execute graph " << cfg.num_iterations << " times" << std::endl;
+    std::cout << "Array size: " << cfg.array_size << std::endl;
+    std::cout << "Progress interval: " << cfg.progress_interval << std::endl;
 
     // Allocate device memory
     int* d_data;
-    HIP_CHECK(hipMalloc(&d_data, ARRAY_SIZE * sizeof(int)));
+    HIP_CHECK(hipMalloc(&d_data, cfg.array_size * sizeof(int)));
 
     // Create graph
     hipGraph_t graph;
@@ -76,7 +137,7 @@ main()
     dim3 blockSize(256);
     dim3 gridSize(1);
 
-    for(int i = 0; i < NUM_KERNELS; i++)
+    for(int i = 0; i < cfg.num_kernels; i++)
     {
         hipLaunchKernelGGL(simpleKernel, gridSize, blockSize, 0, stream, d_data, i);
     }
@@ -95,13 +156,13 @@ main()
     auto start = std::chrono::high_resolution_clock::now();
 
     // Execute the graph multiple times
-    for(int iter = 0; iter < NUM_ITERATIONS; iter++)
+    for(int iter = 0; iter < cfg.num_iterations; iter++)
     {
         roctxRangePush("graph_launch");
         HIP_CHECK(hipGraphLaunch(graphExec, stream));
         roctxRangePop();
 
-        if((iter + 1) % 50 == 0)
+        if((iter + 1) % cfg.progress_interval == 0 || (iter + 1) == cfg.num_iterations)
         {
             std::cout << "Completed " << (iter + 1) << " iterations" << std::endl;
         }
@@ -126,9 +187,10 @@ main()
     std::cout << std::fixed << std::setprecision(4);
     std::cout << "\n=== Timing Results ===" << std::endl;
     std::cout << "Total execution time: " << elapsed.count() << " seconds" << std::endl;
-    std::cout << "Total kernel launches: " << (NUM_KERNELS * NUM_ITERATIONS) << std::endl;
-    std::cout << "Average time per iteration: " << (elapsed.count() / NUM_ITERATIONS) << " seconds"
+    std::cout << "Total kernel launches: " << (cfg.num_kernels * cfg.num_iterations)
               << std::endl;
+    std::cout << "Average time per iteration: " << (elapsed.count() / cfg.num_iterations)
+              << " seconds" << std::endl;
     std::cout << "======================" << std::endl;
 
     // Cleanup

projects/rocprofiler-sdk/tests/hip-graph-bubbles/CMakeLists.txt

@@ -0,0 +1,33 @@
+#
+# Standalone HIP graph bubbles integration test
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+project(
+    rocprofiler-sdk-tests-hip-graph-bubbles
+    LANGUAGES CXX
+    VERSION 0.0.0)
+
+set(HIP_GRAPH_BUBBLES_STANDALONE_ENVIRONMENT)
+if(DEFINED ENV{ROCM_PATH} AND EXISTS "$ENV{ROCM_PATH}/lib/rocm_sysdeps/lib")
+    list(
+        APPEND HIP_GRAPH_BUBBLES_STANDALONE_ENVIRONMENT
+        "LD_LIBRARY_PATH=$ENV{ROCM_PATH}/lib/rocm_sysdeps/lib:$ENV{LD_LIBRARY_PATH}")
+endif()
+
+function(add_hip_graph_bubbles_standalone_test SUFFIX NUM_KERNELS NUM_ITERATIONS TIMEOUT_SECS)
+    math(EXPR EXPECTED_DISPATCH_COUNT "${NUM_KERNELS} * ${NUM_ITERATIONS}")
+
+    rocprofiler_add_integration_execute_test(
+        hip-graph-bubbles-standalone-${SUFFIX}
+        TARGET hip-graph-bubbles
+        ARGS ${NUM_KERNELS} ${NUM_ITERATIONS} 256 25
+        TIMEOUT ${TIMEOUT_SECS}
+        LABELS "integration-tests;standalone"
+        ENVIRONMENT ${HIP_GRAPH_BUBBLES_STANDALONE_ENVIRONMENT}
+        PASS_REGULAR_EXPRESSION
+            "Total kernel launches: ${EXPECTED_DISPATCH_COUNT};Test completed successfully")
+endfunction()
+
+add_hip_graph_bubbles_standalone_test(k256-i20 256 20 30)
+add_hip_graph_bubbles_standalone_test(k2000-i200 2000 200 60)

projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt

@@ -24,6 +24,7 @@ enable_testing()
 include(CTest)
 
 add_subdirectory(tracing)
+add_subdirectory(hip-graph-bubbles)
 add_subdirectory(tracing-plus-counter-collection)
 add_subdirectory(tracing-hip-in-libraries)
 add_subdirectory(counter-collection)

projects/rocprofiler-sdk/tests/rocprofv3/hip-graph-bubbles/CMakeLists.txt

@@ -0,0 +1,58 @@
+#
+# rocprofv3 HIP graph bubble regression tests
+#
+cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR)
+
+project(
+    rocprofiler-sdk-tests-rocprofv3-hip-graph-bubbles
+    LANGUAGES CXX
+    VERSION 0.0.0)
+
+find_package(rocprofiler-sdk REQUIRED)
+
+set(HIP_GRAPH_BUBBLES_ENVIRONMENT)
+if(DEFINED ENV{ROCM_PATH} AND EXISTS "$ENV{ROCM_PATH}/lib/rocm_sysdeps/lib")
+    list(
+        APPEND HIP_GRAPH_BUBBLES_ENVIRONMENT
+        "LD_LIBRARY_PATH=$ENV{ROCM_PATH}/lib/rocm_sysdeps/lib:$ENV{LD_LIBRARY_PATH}")
+endif()
+
+function(add_hip_graph_bubbles_kernel_trace_test SUFFIX NUM_KERNELS NUM_ITERATIONS TIMEOUT_SECS)
+    set(TEST_NAME "rocprofv3-test-hip-graph-bubbles-${SUFFIX}")
+    set(OUTPUT_DIR "${CMAKE_CURRENT_BINARY_DIR}/hip-graph-bubbles-${SUFFIX}")
+    math(EXPR EXPECTED_DISPATCH_COUNT "${NUM_KERNELS} * ${NUM_ITERATIONS}")
+
+    rocprofiler_add_integration_execute_test(
+        ${TEST_NAME}
+        COMMAND
+            $<TARGET_FILE:rocprofiler-sdk::rocprofv3> --kernel-trace -d ${OUTPUT_DIR} -o
+            out --output-format csv -- $<TARGET_FILE:hip-graph-bubbles> ${NUM_KERNELS}
+            ${NUM_ITERATIONS} 256 25
+        DEPENDS hip-graph-bubbles
+        TIMEOUT ${TIMEOUT_SECS}
+        LABELS "integration-tests"
+        ENVIRONMENT ${HIP_GRAPH_BUBBLES_ENVIRONMENT}
+        FIXTURES_SETUP ${TEST_NAME}
+        FAIL_REGULAR_EXPRESSION "KERNEL_DISPATCH|CODE_OBJECT")
+
+    rocprofiler_add_integration_validate_test(
+        ${TEST_NAME}
+        TEST_PATHS validate.py
+        COPY conftest.py
+        CONFIG pytest.ini
+        ARGS --kernel-input
+             ${CMAKE_CURRENT_BINARY_DIR}/hip-graph-bubbles-${SUFFIX}/out_kernel_trace.csv
+             --expected-dispatch-count
+             ${EXPECTED_DISPATCH_COUNT}
+             --expected-kernels
+             ${NUM_KERNELS}
+             --expected-iterations
+             ${NUM_ITERATIONS}
+        TIMEOUT ${TIMEOUT_SECS}
+        LABELS "integration-tests"
+        FIXTURES_REQUIRED ${TEST_NAME})
+endfunction()
+
+add_hip_graph_bubbles_kernel_trace_test(k256-i20 256 20 60)
+add_hip_graph_bubbles_kernel_trace_test(k256-i200 256 200 120)
+add_hip_graph_bubbles_kernel_trace_test(k2000-i200 2000 200 180)

projects/rocprofiler-sdk/tests/rocprofv3/hip-graph-bubbles/conftest.py

@@ -0,0 +1,79 @@
+#!/usr/bin/env python3
+
+# MIT License
+#
+# Copyright (c) 2024-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.
+
+import csv
+import pytest
+
+
+def pytest_addoption(parser):
+    parser.addoption(
+        "--kernel-input",
+        action="store",
+        help="Path to kernel tracing CSV file.",
+    )
+    parser.addoption(
+        "--expected-dispatch-count",
+        action="store",
+        type=int,
+        help="Expected number of kernel dispatch records.",
+    )
+    parser.addoption(
+        "--expected-kernels",
+        action="store",
+        type=int,
+        help="Expected kernels per graph launch.",
+    )
+    parser.addoption(
+        "--expected-iterations",
+        action="store",
+        type=int,
+        help="Expected number of graph launches.",
+    )
+
+
+@pytest.fixture
+def kernel_input_data(request):
+    filename = request.config.getoption("--kernel-input")
+    data = []
+    with open(filename, "r") as inp:
+        reader = csv.DictReader(inp)
+        for row in reader:
+            data.append(row)
+
+    return data
+
+
+@pytest.fixture
+def expected_dispatch_count(request):
+    return request.config.getoption("--expected-dispatch-count")
+
+
+@pytest.fixture
+def expected_kernels(request):
+    return request.config.getoption("--expected-kernels")
+
+
+@pytest.fixture
+def expected_iterations(request):
+    return request.config.getoption("--expected-iterations")

projects/rocprofiler-sdk/tests/rocprofv3/hip-graph-bubbles/pytest.ini

@@ -0,0 +1,4 @@
+[pytest]
+addopts = --durations=20 -rA -s
+testpaths = validate.py
+pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages

projects/rocprofiler-sdk/tests/rocprofv3/hip-graph-bubbles/validate.py

@@ -0,0 +1,66 @@
+#!/usr/bin/env python3
+
+# MIT License
+#
+# Copyright (c) 2024-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.
+
+import sys
+import pytest
+
+
+def test_kernel_trace_row_count(
+    kernel_input_data,
+    expected_dispatch_count,
+    expected_kernels,
+    expected_iterations,
+):
+    assert expected_dispatch_count == expected_kernels * expected_iterations
+    assert len(kernel_input_data) == expected_dispatch_count
+
+
+def test_kernel_trace_dispatch_ids(kernel_input_data, expected_dispatch_count):
+    dispatch_ids = [int(row["Dispatch_Id"]) for row in kernel_input_data]
+
+    assert len(dispatch_ids) == expected_dispatch_count
+    assert len(set(dispatch_ids)) == expected_dispatch_count
+
+
+def test_kernel_trace_fields(kernel_input_data, expected_dispatch_count):
+    assert len(kernel_input_data) == expected_dispatch_count
+
+    for row in kernel_input_data:
+        assert row["Kind"] == "KERNEL_DISPATCH"
+        assert int(row["Agent_Id"].split(" ")[-1]) >= 0
+        assert int(row["Queue_Id"]) > 0
+        assert int(row["Kernel_Id"]) > 0
+        assert int(row["Correlation_Id"]) > 0
+        assert int(row["Workgroup_Size_X"]) == 256
+        assert int(row["Workgroup_Size_Y"]) == 1
+        assert int(row["Workgroup_Size_Z"]) == 1
+        assert int(row["Grid_Size_X"]) == 256
+        assert int(row["Grid_Size_Y"]) == 1
+        assert int(row["Grid_Size_Z"]) == 1
+        assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"])
+
+
+if __name__ == "__main__":
+    exit_code = pytest.main(["-x", __file__] + sys.argv[1:])
+    sys.exit(exit_code)