From 4cd076b9cc42491f71830fe0e168ffa390a1fd4a Mon Sep 17 00:00:00 2001 From: "Jonathan R. Madsen" Date: Tue, 7 May 2024 15:10:22 -0500 Subject: [PATCH] Test using HIP Graphs (#835) * Test using hip graphs * Remove assert for api_end < async_end * Update rocprofv3/tracing-hip-in-libraries::test_api_trace * Update rocprofv3/tracing-hip-in-libraries::test_api_trace * Increase rocprofv3-test-trace-hip-in-libraries-validate timeout * Update rocprofv3/tracing-hip-in-libraries::test_api_trace * Remove submit retry * Update rocprofv3/tracing-hip-in-libraries::test_api_trace * Increase rocprofv3-test-trace-hip-in-libraries-validate timeout * Update lib/common/container/record_header_buffer.hpp - minor tweaks * Update lib/rocprofiler-sdk/buffer.hpp - tweak ROCPROFILER_BUFFER_POLICY_LOSSLESS flush behavior * Increase rocprofv3-test-trace-hip-in-libraries-validate timeout * Update rocprofv3/tracing-hip-in-libraries::test_api_trace * Revert rocprofv3-test-trace-hip-in-libraries-validate timeout * Update run-ci.py - RETRY_COUNT set to zero [ROCm/rocprofiler-sdk commit: 1f96593b4f144952a1b881e888f084b49d2558a8] --- .../common/container/record_header_buffer.hpp | 17 +- .../source/lib/rocprofiler-sdk/buffer.hpp | 6 +- .../rocprofiler-sdk/source/scripts/run-ci.py | 2 +- projects/rocprofiler-sdk/tests/CMakeLists.txt | 1 + .../tests/async-copy-tracing/validate.py | 4 +- .../rocprofiler-sdk/tests/bin/CMakeLists.txt | 1 + .../tests/bin/hip-graph/CMakeLists.txt | 44 ++++ .../tests/bin/hip-graph/hip-graph.cpp | 212 ++++++++++++++++ .../tests/hip-graph-tracing/CMakeLists.txt | 51 ++++ .../tests/hip-graph-tracing/conftest.py | 20 ++ .../tests/hip-graph-tracing/pytest.ini | 5 + .../tests/hip-graph-tracing/validate.py | 235 ++++++++++++++++++ .../tests/kernel-tracing/validate.py | 4 +- .../tracing-hip-in-libraries/validate.py | 56 ++++- 14 files changed, 635 insertions(+), 23 deletions(-) create mode 100644 projects/rocprofiler-sdk/tests/bin/hip-graph/CMakeLists.txt create mode 100644 projects/rocprofiler-sdk/tests/bin/hip-graph/hip-graph.cpp create mode 100644 projects/rocprofiler-sdk/tests/hip-graph-tracing/CMakeLists.txt create mode 100644 projects/rocprofiler-sdk/tests/hip-graph-tracing/conftest.py create mode 100644 projects/rocprofiler-sdk/tests/hip-graph-tracing/pytest.ini create mode 100644 projects/rocprofiler-sdk/tests/hip-graph-tracing/validate.py diff --git a/projects/rocprofiler-sdk/source/lib/common/container/record_header_buffer.hpp b/projects/rocprofiler-sdk/source/lib/common/container/record_header_buffer.hpp index 2e93d7a1cd..84265da289 100644 --- a/projects/rocprofiler-sdk/source/lib/common/container/record_header_buffer.hpp +++ b/projects/rocprofiler-sdk/source/lib/common/container/record_header_buffer.hpp @@ -257,10 +257,10 @@ record_header_buffer::emplace(uint64_t _hash, Tp& _v) // placement new new(_addr) Tp{_v}; - rocprofiler_record_header_t record = {}; - record.hash = _hash; - record.payload = _addr; - m_headers.at(idx) = record; + auto record = rocprofiler_record_header_t{}; + record.hash = _hash; + record.payload = _addr; + m_headers.at(idx) = record; } read_unlock(); @@ -299,10 +299,11 @@ record_header_buffer::emplace(uint32_t _category, uint32_t _kind, Tp& _v) // placement new new(_addr) Tp{_v}; - m_headers.at(idx) = rocprofiler_record_header_t{}; - m_headers.at(idx).category = _category; - m_headers.at(idx).kind = _kind; - m_headers.at(idx).payload = _addr; + auto record = rocprofiler_record_header_t{}; + record.category = _category; + record.kind = _kind; + record.payload = _addr; + m_headers.at(idx) = record; } read_unlock(); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp index 09f815b0a7..7438d132b4 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp @@ -134,12 +134,12 @@ rocprofiler::buffer::instance::emplace(uint32_t category, uint32_t kind, Tp& val if(policy == ROCPROFILER_BUFFER_POLICY_LOSSLESS) { // blocks until buffer is flushed - while(!success) + do { - buffer::flush(buffer_id, false); + buffer::flush(buffer_id, true); idx = get_idx(); success = buffers.at(idx).emplace(category, kind, value); - } + } while(!success); } else { diff --git a/projects/rocprofiler-sdk/source/scripts/run-ci.py b/projects/rocprofiler-sdk/source/scripts/run-ci.py index 026b07bb15..c5d85c028e 100755 --- a/projects/rocprofiler-sdk/source/scripts/run-ci.py +++ b/projects/rocprofiler-sdk/source/scripts/run-ci.py @@ -231,7 +231,7 @@ def generate_dashboard_script(args): macro(dashboard_submit) if("{SUBMIT}" GREATER 0) ctest_submit({ARGN} - RETRY_COUNT 1 + RETRY_COUNT 0 RETRY_DELAY 10 CAPTURE_CMAKE_ERROR _cdash_submit_err) diff --git a/projects/rocprofiler-sdk/tests/CMakeLists.txt b/projects/rocprofiler-sdk/tests/CMakeLists.txt index 128ebdb4f6..397e3a3387 100644 --- a/projects/rocprofiler-sdk/tests/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/CMakeLists.txt @@ -58,6 +58,7 @@ add_subdirectory(scratch-memory-tracing) add_subdirectory(c-tool) add_subdirectory(page-migration) add_subdirectory(thread-trace) +add_subdirectory(hip-graph-tracing) # rocprofv3 validation tests add_subdirectory(rocprofv3) diff --git a/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py b/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py index 013a86195d..ea98c8a1b1 100644 --- a/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py +++ b/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py @@ -115,9 +115,9 @@ def test_timestamps(input_data): ), f"[{titr}] {itr}" api_start = cb_start[itr["correlation_id"]["internal"]] - api_end = cb_end[itr["correlation_id"]["internal"]] + # api_end = cb_end[itr["correlation_id"]["internal"]] assert api_start < itr["start_timestamp"], f"[{titr}] {itr}" - assert api_end <= itr["end_timestamp"], f"[{titr}] {itr}" + # assert api_end <= itr["end_timestamp"], f"[{titr}] {itr}" def test_internal_correlation_ids(input_data): diff --git a/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt b/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt index 49faa6217c..62b0faa9f9 100644 --- a/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/bin/CMakeLists.txt @@ -23,3 +23,4 @@ add_subdirectory(hip-in-libraries) add_subdirectory(scratch-memory) add_subdirectory(page-migration) add_subdirectory(hsa-queue-dependency) +add_subdirectory(hip-graph) diff --git a/projects/rocprofiler-sdk/tests/bin/hip-graph/CMakeLists.txt b/projects/rocprofiler-sdk/tests/bin/hip-graph/CMakeLists.txt new file mode 100644 index 0000000000..1f699837fe --- /dev/null +++ b/projects/rocprofiler-sdk/tests/bin/hip-graph/CMakeLists.txt @@ -0,0 +1,44 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +if(NOT CMAKE_HIP_COMPILER) + find_program( + amdclangpp_EXECUTABLE + NAMES amdclang++ + HINTS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm + PATHS ${ROCM_PATH} ENV ROCM_PATH /opt/rocm + PATH_SUFFIXES bin llvm/bin NO_CACHE) + mark_as_advanced(amdclangpp_EXECUTABLE) + + if(amdclangpp_EXECUTABLE) + set(CMAKE_HIP_COMPILER "${amdclangpp_EXECUTABLE}") + endif() +endif() + +project(rocprofiler-tests-bin-hip-graph LANGUAGES CXX HIP) + +foreach(_TYPE DEBUG MINSIZEREL RELEASE RELWITHDEBINFO) + if("${CMAKE_HIP_FLAGS_${_TYPE}}" STREQUAL "") + set(CMAKE_HIP_FLAGS_${_TYPE} "${CMAKE_CXX_FLAGS_${_TYPE}}") + endif() +endforeach() + +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_EXTENSIONS OFF) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(CMAKE_HIP_STANDARD 17) +set(CMAKE_HIP_EXTENSIONS OFF) +set(CMAKE_HIP_STANDARD_REQUIRED ON) + +set_source_files_properties(hip-graph.cpp PROPERTIES LANGUAGE HIP) +add_executable(hip-graph) +target_sources(hip-graph PRIVATE hip-graph.cpp) +target_compile_options(hip-graph PRIVATE -W -Wall -Wextra -Wpedantic -Wshadow -Werror) + +find_package(Threads REQUIRED) +target_link_libraries(hip-graph PRIVATE Threads::Threads) + +# find_package(rocprofiler-sdk-roctx REQUIRED) target_link_libraries(hip-graph PRIVATE +# rocprofiler-sdk-roctx::rocprofiler-sdk-roctx) diff --git a/projects/rocprofiler-sdk/tests/bin/hip-graph/hip-graph.cpp b/projects/rocprofiler-sdk/tests/bin/hip-graph/hip-graph.cpp new file mode 100644 index 0000000000..2ef22e47fb --- /dev/null +++ b/projects/rocprofiler-sdk/tests/bin/hip-graph/hip-graph.cpp @@ -0,0 +1,212 @@ +/* +Copyright (c) 2015 - 2021 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 +#include +#include +#include + +// hip header file +#include + +#include +#include +#include +#include +#include + +namespace +{ +using auto_lock_t = std::unique_lock; +auto print_mutex = std::mutex{}; +auto global_kern_num = std::atomic{0}; +} // namespace + +template +void +check(T result, char const* const func, const char* const file, int const line) +{ + if(result) + { + fprintf(stderr, + "Hip error at %s:%d code=%d(%s) \"%s\" \n", + file, + line, + static_cast(result), + hipGetErrorName(result), + func); + exit(EXIT_FAILURE); + } +} +#define checkHipErrors(val) check((val), #val, __FILE__, __LINE__) + +__global__ void +kernel_foo(const int devid, const int kernid, const int kernid_global, const volatile int* streamid) +{ + printf("[hip-graph][device %2i][stream %2i] Kernel foo | %2i | %2i executing...\n", + devid, + *streamid, + kernid, + kernid_global); +} + +__global__ void +kernel_bar(const int devid, const int kernid, const int kernid_global, const volatile int* streamid) +{ + printf("[hip-graph][device %2i][stream %2i] Kernel bar | %2i | %2i executing...\n", + devid, + *streamid, + kernid, + kernid_global); +} + +void +run(uint64_t devid, + uint64_t nstream, + uint64_t nkernel_per_stream, + std::atomic* progress, + const std::shared_future& future) +{ + auto prefix = [devid]() { + auto ss = std::stringstream{}; + ss << "[hip-graph][device " << std::setw(2) << devid << "] "; + return ss.str(); + }(); + + auto log_message = [&prefix](const auto& msg) { + auto _lk = auto_lock_t{print_mutex}; + std::cout << prefix << msg << "..." << std::endl; + }; + + log_message("setting device"); + checkHipErrors(hipSetDevice(devid)); + + auto streams = std::vector(nstream, nullptr); + auto stream_num = std::vector(nstream, nullptr); + + log_message("creating streams"); + for(auto& itr : streams) + checkHipErrors(hipStreamCreate(&itr)); + + log_message("allocating data"); + for(uint64_t i = 0; i < nstream; ++i) + { + auto& itr = stream_num.at(i); + auto* str = streams.at(i); + auto val = i; + checkHipErrors(hipMallocAsync(&itr, sizeof(int), str)); + checkHipErrors(hipMemcpyAsync(itr, &val, sizeof(int), hipMemcpyHostToDevice, str)); + } + + auto graphs = std::vector(nstream); + auto execs = std::vector(nstream, nullptr); + + uint64_t kern_num = 0; + for(uint64_t i = 0; i < nstream; ++i) + { + checkHipErrors(hipStreamBeginCapture(streams.at(i), hipStreamCaptureModeGlobal)); + + for(uint64_t j = 0; j < nkernel_per_stream; ++j) + { + auto kern_num_v = kern_num++; + auto glob_kern_num_v = global_kern_num++; + auto kernel = (j % 2 == 0) ? kernel_foo : kernel_bar; + hipLaunchKernelGGL(kernel, + dim3(1), + dim3(1), + 0, + streams.at(i), + devid, + kern_num_v, + glob_kern_num_v, + stream_num.at(i)); + checkHipErrors(hipGetLastError()); + } + + checkHipErrors(hipStreamEndCapture(streams.at(i), &graphs.at(i))); + checkHipErrors(hipGraphInstantiate(&execs.at(i), graphs.at(i), nullptr, nullptr, 0)); + } + + if(progress) progress->fetch_add(1); + future.wait(); + + log_message("launching graph"); + for(uint64_t i = 0; i < nstream; ++i) + checkHipErrors(hipGraphLaunch(execs.at(i), streams.at(i))); + + log_message("synchronizing device"); + checkHipErrors(hipDeviceSynchronize()); + + log_message("destroying graph"); + for(uint64_t i = 0; i < nstream; ++i) + checkHipErrors(hipGraphDestroy(graphs.at(i))); + + log_message("freeing data"); + for(auto& itr : stream_num) + checkHipErrors(hipFree(itr)); + + log_message("returning"); +} + +int +main(int argc, char* argv[]) +{ + std::cout << "[" << ::basename(argv[0]) << "] executing..." << std::endl; + + int ndevice_real = 0; + checkHipErrors(hipGetDeviceCount(&ndevice_real)); + + uint64_t nstream = 1; + uint64_t nkernel_per_stream = 12; + uint64_t ndevice = ndevice_real; + + if(argc > 1) nstream = std::stoul(argv[1]); + if(argc > 2) nkernel_per_stream = std::stoul(argv[2]); + if(argc > 3) ndevice = std::stoul(argv[3]); + + ndevice = std::min(ndevice, ndevice_real); + + auto progress = std::atomic{0}; + auto promise = std::promise{}; + auto future = promise.get_future().share(); + auto threads = std::vector{}; + threads.reserve(ndevice); + + for(uint64_t i = 0; i < ndevice; ++i) + threads.emplace_back(run, i, nstream, nkernel_per_stream, &progress, future); + + // wait for all threads to reach designated progress point + while(progress < ndevice) + { + std::this_thread::yield(); + std::this_thread::sleep_for(std::chrono::milliseconds{1}); + } + + // release the threads + promise.set_value(); + + for(auto& itr : threads) + itr.join(); + + std::cout << "[" << ::basename(argv[0]) << "] complete" << std::endl; + return 0; +} diff --git a/projects/rocprofiler-sdk/tests/hip-graph-tracing/CMakeLists.txt b/projects/rocprofiler-sdk/tests/hip-graph-tracing/CMakeLists.txt new file mode 100644 index 0000000000..97a9d360b6 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/hip-graph-tracing/CMakeLists.txt @@ -0,0 +1,51 @@ +# +# +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project( + rocprofiler-tests-hip-graph-tracing + LANGUAGES CXX + VERSION 0.0.0) + +find_package(rocprofiler-sdk REQUIRED) + +set(PYTEST_ARGS) +if(ROCPROFILER_MEMCHECK MATCHES "(Address|Thread)Sanitizer" OR ROCPROFILER_BUILD_CODECOV) + set(PYTEST_ARGS -k "not test_total_runtime") +endif() + +if(ROCPROFILER_MEMCHECK_PRELOAD_ENV) + set(PRELOAD_ENV + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}:$") +else() + set(PRELOAD_ENV "LD_PRELOAD=$") +endif() + +add_test(NAME test-hip-graph-tracing-execute COMMAND $) + +set(hip-graph-tracing-env + "${PRELOAD_ENV}" + "ROCPROFILER_TOOL_OUTPUT_FILE=hip-graph-tracing-test.json" + "LD_LIBRARY_PATH=$:$ENV{LD_LIBRARY_PATH}" + "ROCPROFILER_TOOL_CONTEXTS=HIP_API_CALLBACK,HIP_API_BUFFERED,KERNEL_DISPATCH_CALLBACK,KERNEL_DISPATCH_BUFFERED,CODE_OBJECT" + ) + +set_tests_properties( + test-hip-graph-tracing-execute + PROPERTIES TIMEOUT 100 LABELS "integration-tests" ENVIRONMENT + "${hip-graph-tracing-env}" FAIL_REGULAR_EXPRESSION + "${ROCPROFILER_DEFAULT_FAIL_REGEX}") + +rocprofiler_configure_pytest_files(COPY validate.py conftest.py CONFIG pytest.ini) + +add_test( + NAME test-hip-graph-tracing-validate + COMMAND ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py ${PYTEST_ARGS} + --input ${CMAKE_CURRENT_BINARY_DIR}/hip-graph-tracing-test.json) + +set_tests_properties( + test-hip-graph-tracing-validate + PROPERTIES TIMEOUT 45 LABELS "integration-tests" DEPENDS + test-hip-graph-tracing-execute FAIL_REGULAR_EXPRESSION + "${ROCPROFILER_DEFAULT_FAIL_REGEX}") diff --git a/projects/rocprofiler-sdk/tests/hip-graph-tracing/conftest.py b/projects/rocprofiler-sdk/tests/hip-graph-tracing/conftest.py new file mode 100644 index 0000000000..94bd6ef823 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/hip-graph-tracing/conftest.py @@ -0,0 +1,20 @@ +#!/usr/bin/env python3 + +import json +import pytest + + +def pytest_addoption(parser): + parser.addoption( + "--input", + action="store", + default="hip-graph-tracing-test.json", + help="Input JSON", + ) + + +@pytest.fixture +def input_data(request): + filename = request.config.getoption("--input") + with open(filename, "r") as inp: + return json.load(inp) diff --git a/projects/rocprofiler-sdk/tests/hip-graph-tracing/pytest.ini b/projects/rocprofiler-sdk/tests/hip-graph-tracing/pytest.ini new file mode 100644 index 0000000000..5e1e1c14a0 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/hip-graph-tracing/pytest.ini @@ -0,0 +1,5 @@ + +[pytest] +addopts = --durations=20 -rA -s -vv +testpaths = validate.py +pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages diff --git a/projects/rocprofiler-sdk/tests/hip-graph-tracing/validate.py b/projects/rocprofiler-sdk/tests/hip-graph-tracing/validate.py new file mode 100644 index 0000000000..3cad9f815b --- /dev/null +++ b/projects/rocprofiler-sdk/tests/hip-graph-tracing/validate.py @@ -0,0 +1,235 @@ +#!/usr/bin/env python3 + +import sys +import pytest + + +# helper function +def node_exists(name, data, min_len=1): + assert name in data + assert data[name] is not None + if isinstance(data[name], (list, tuple, dict, set)): + assert len(data[name]) >= min_len + + +def test_data_structure(input_data): + """verify minimum amount of expected data is present""" + data = input_data + + node_exists("rocprofiler-sdk-json-tool", data) + + sdk_data = data["rocprofiler-sdk-json-tool"] + + node_exists("metadata", sdk_data) + node_exists("pid", sdk_data["metadata"]) + node_exists("main_tid", sdk_data["metadata"]) + node_exists("init_time", sdk_data["metadata"]) + node_exists("fini_time", sdk_data["metadata"]) + + node_exists("agents", sdk_data) + node_exists("call_stack", sdk_data) + node_exists("callback_records", sdk_data) + node_exists("buffer_records", sdk_data) + + node_exists("names", sdk_data["callback_records"]) + node_exists("code_objects", sdk_data["callback_records"]) + node_exists("kernel_symbols", sdk_data["callback_records"]) + node_exists("hip_api_traces", sdk_data["callback_records"]) + node_exists("kernel_dispatches", sdk_data["callback_records"]) + + node_exists("names", sdk_data["buffer_records"]) + node_exists("kernel_dispatches", sdk_data["buffer_records"]) + node_exists("hip_api_traces", sdk_data["buffer_records"], 0) + + +def test_size_entries(input_data): + # check that size fields are > 0 but account for function arguments + # which are named "size" + def check_size(data, bt): + if "size" in data.keys(): + if isinstance(data["size"], str) and bt.endswith('["args"]'): + pass + else: + assert data["size"] > 0, f"origin: {bt}" + + # recursively check the entire data structure + def iterate_data(data, bt): + if isinstance(data, (list, tuple)): + for i, itr in enumerate(data): + if isinstance(itr, dict): + check_size(itr, f"{bt}[{i}]") + iterate_data(itr, f"{bt}[{i}]") + elif isinstance(data, dict): + check_size(data, f"{bt}") + for key, itr in data.items(): + iterate_data(itr, f'{bt}["{key}"]') + + # start recursive check over entire JSON dict + iterate_data(input_data, "input_data") + + +def test_timestamps(input_data): + data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] + + cb_start = {} + cb_end = {} + for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]: + for itr in sdk_data["callback_records"][titr]: + cid = itr["correlation_id"]["internal"] + phase = itr["phase"] + if phase == 1: + cb_start[cid] = itr["timestamp"] + elif phase == 2: + cb_end[cid] = itr["timestamp"] + assert cb_start[cid] <= itr["timestamp"] + else: + assert phase == 1 or phase == 2 + + for itr in sdk_data["buffer_records"][titr]: + assert itr["start_timestamp"] <= itr["end_timestamp"] + + for titr in ["kernel_dispatches", "memory_copies"]: + for itr in sdk_data["buffer_records"][titr]: + assert itr["start_timestamp"] < itr["end_timestamp"], f"[{titr}] {itr}" + assert itr["correlation_id"]["internal"] > 0, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["init_time"] < itr["start_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["init_time"] < itr["end_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["fini_time"] > itr["start_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["fini_time"] > itr["end_timestamp"] + ), f"[{titr}] {itr}" + + api_start = cb_start[itr["correlation_id"]["internal"]] + # api_end = cb_end[itr["correlation_id"]["internal"]] + assert api_start < itr["start_timestamp"], f"[{titr}] {itr}" + # assert api_end <= itr["end_timestamp"], f"[{titr}] {itr}" + + +def test_internal_correlation_ids(input_data): + data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] + + api_corr_ids = [] + for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]: + for itr in sdk_data["callback_records"][titr]: + api_corr_ids.append(itr["correlation_id"]["internal"]) + + for itr in sdk_data["buffer_records"][titr]: + api_corr_ids.append(itr["correlation_id"]["internal"]) + + api_corr_ids_sorted = sorted(api_corr_ids) + api_corr_ids_unique = list(set(api_corr_ids)) + + for itr in sdk_data["buffer_records"]["kernel_dispatches"]: + assert itr["correlation_id"]["internal"] in api_corr_ids_unique + + for itr in sdk_data["buffer_records"]["memory_copies"]: + assert itr["correlation_id"]["internal"] in api_corr_ids_unique + + len_corr_id_unq = len(api_corr_ids_unique) + assert len(api_corr_ids) != len_corr_id_unq + assert max(api_corr_ids_sorted) == len_corr_id_unq + + +def test_external_correlation_ids(input_data): + data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] + + extern_corr_ids = [] + for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]: + for itr in sdk_data["callback_records"][titr]: + assert itr["correlation_id"]["external"] > 0 + assert itr["thread_id"] == itr["correlation_id"]["external"] + extern_corr_ids.append(itr["correlation_id"]["external"]) + + extern_corr_ids = list(set(sorted(extern_corr_ids))) + for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]: + for itr in sdk_data["buffer_records"][titr]: + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert ( + itr["thread_id"] == itr["correlation_id"]["external"] + ), f"[{titr}] {itr}" + assert itr["thread_id"] in extern_corr_ids, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}" + + for titr in ["kernel_dispatches", "memory_copies"]: + for itr in sdk_data["buffer_records"][titr]: + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}" + + +def test_kernel_ids(input_data): + data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] + + symbol_info = {} + for itr in sdk_data["callback_records"]["kernel_symbols"]: + phase = itr["phase"] + payload = itr["payload"] + kern_id = payload["kernel_id"] + + assert phase == 1 or phase == 2 + assert kern_id > 0 + if phase == 1: + assert len(payload["kernel_name"]) > 0 + symbol_info[kern_id] = payload + elif phase == 2: + assert payload["kernel_id"] in symbol_info.keys() + assert payload["kernel_name"] == symbol_info[kern_id]["kernel_name"] + + for itr in sdk_data["buffer_records"]["kernel_dispatches"]: + assert itr["dispatch_info"]["kernel_id"] in symbol_info.keys() + + for itr in sdk_data["callback_records"]["kernel_dispatches"]: + assert itr["payload"]["dispatch_info"]["kernel_id"] in symbol_info.keys() + + +def test_kernel_dispatch_ids(input_data): + data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] + + num_dispatches = len(sdk_data["buffer_records"]["kernel_dispatches"]) + num_cb_dispatches = len(sdk_data["callback_records"]["kernel_dispatches"]) + + assert num_cb_dispatches == (3 * num_dispatches) + + bf_seq_ids = [] + for itr in sdk_data["buffer_records"]["kernel_dispatches"]: + bf_seq_ids.append(itr["dispatch_info"]["dispatch_id"]) + + cb_seq_ids = [] + for itr in sdk_data["callback_records"]["kernel_dispatches"]: + cb_seq_ids.append(itr["payload"]["dispatch_info"]["dispatch_id"]) + + bf_seq_ids = sorted(bf_seq_ids) + cb_seq_ids = sorted(cb_seq_ids) + + assert (3 * len(bf_seq_ids)) == len(cb_seq_ids) + + assert bf_seq_ids[0] == cb_seq_ids[0] + assert bf_seq_ids[-1] == cb_seq_ids[-1] + + def get_uniq(data): + return list(set(data)) + + bf_seq_ids_uniq = get_uniq(bf_seq_ids) + cb_seq_ids_uniq = get_uniq(cb_seq_ids) + + assert bf_seq_ids == bf_seq_ids_uniq + assert len(cb_seq_ids) == (3 * len(cb_seq_ids_uniq)) + assert len(bf_seq_ids) == num_dispatches + assert len(bf_seq_ids_uniq) == num_dispatches + assert len(cb_seq_ids_uniq) == num_dispatches + + +if __name__ == "__main__": + exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) + sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py b/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py index b4e5f633e2..4833dcae86 100644 --- a/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py +++ b/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py @@ -114,9 +114,9 @@ def test_timestamps(input_data): ), f"[{titr}] {itr}" api_start = cb_start[itr["correlation_id"]["internal"]] - api_end = cb_end[itr["correlation_id"]["internal"]] + # api_end = cb_end[itr["correlation_id"]["internal"]] assert api_start < itr["start_timestamp"], f"[{titr}] {itr}" - assert api_end <= itr["end_timestamp"], f"[{titr}] {itr}" + # assert api_end <= itr["end_timestamp"], f"[{titr}] {itr}" def test_total_runtime(input_data): diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py index 11aada4f03..a0d8e19140 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py @@ -37,9 +37,16 @@ def validate_stats(row): assert stddev_v > 0.0 if cnt_v > 1 else int(stddev_v) == 0, f"{row}" -def test_api_trace(hsa_input_data, hip_input_data, hip_stats_data): +def test_api_trace( + hsa_input_data, + hip_input_data, + kernel_input_data, + memory_copy_input_data, + hip_stats_data, +): functions = [] - correlation_ids = [] + hsa_correlation_ids = [] + hip_correlation_ids = [] for row in hsa_input_data: assert row["Domain"] in ( "HSA_CORE_API", @@ -51,7 +58,8 @@ def test_api_trace(hsa_input_data, hip_input_data, hip_stats_data): assert int(row["Thread_Id"]) >= int(row["Process_Id"]) assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) functions.append(row["Function"]) - correlation_ids.append(int(row["Correlation_Id"])) + cid = int(row["Correlation_Id"]) + hsa_correlation_ids.append(cid) for row in hip_input_data: assert row["Domain"] in [ @@ -64,15 +72,49 @@ def test_api_trace(hsa_input_data, hip_input_data, hip_stats_data): ) assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) functions.append(row["Function"]) - correlation_ids.append(int(row["Correlation_Id"])) + cid = int(row["Correlation_Id"]) + hip_correlation_ids.append(cid) - correlation_ids = sorted(list(set(correlation_ids))) + def get_sorted_unique(inp): + return sorted(list(set(inp))) + + def diagnose_non_unique(_input_data): + _corr_id_hist = {} + for row in _input_data: + _cid = int(row["Correlation_Id"]) + # ensure duplicate does not already exist + assert ( + _cid not in _corr_id_hist.keys() + ), f"\ncurrent : {row}\nprevious: {_corr_id_hist[_cid]}" + _corr_id_hist[_cid] = row + + if len(hsa_correlation_ids) != len(get_sorted_unique(hsa_correlation_ids)): + diagnose_non_unique(hsa_input_data) + + if len(hip_correlation_ids) != len(get_sorted_unique(hip_correlation_ids)): + diagnose_non_unique(hip_input_data) + + correlation_ids = get_sorted_unique(hsa_correlation_ids + hip_correlation_ids) + + # make sure that we have associated API calls for all async ops + for itr in [kernel_input_data, memory_copy_input_data]: + for row in itr: + cid = int(row["Correlation_Id"]) + assert ( + cid in correlation_ids + ), f"[{cid}] {row}\nCorrelation IDs:\n\t{correlation_ids}" # all correlation ids are unique + if len(correlation_ids) != (len(hsa_input_data) + len(hip_input_data)): + for itr in hsa_input_data: + assert int(itr["Correlation_Id"]) in correlation_ids, f"{itr}" + for itr in hip_input_data: + assert int(itr["Correlation_Id"]) in correlation_ids, f"{itr}" + assert len(correlation_ids) == (len(hsa_input_data) + len(hip_input_data)) # correlation ids are numbered from 1 to N - assert correlation_ids[0] == 1 - assert correlation_ids[-1] == len(correlation_ids) + assert correlation_ids[0] == 1, f"{correlation_ids}" + assert correlation_ids[-1] == len(correlation_ids), f"{correlation_ids}" functions = list(set(functions)) for itr in (