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: 1f96593b4f]
Este cometimento está contido em:
Jonathan R. Madsen
cometido por
GitHub
GitHub
ascendente
ea56f79495
cometimento
4cd076b9cc
@@ -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();
|
||||
|
||||
|
||||
@@ -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
|
||||
{
|
||||
|
||||
@@ -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)
|
||||
|
||||
|
||||
@@ -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)
|
||||
|
||||
@@ -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):
|
||||
|
||||
@@ -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)
|
||||
|
||||
@@ -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)
|
||||
@@ -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 <future>
|
||||
#include <iomanip>
|
||||
#include <iostream>
|
||||
#include <mutex>
|
||||
|
||||
// hip header file
|
||||
#include <hip/hip_runtime.h>
|
||||
|
||||
#include <stdio.h>
|
||||
#include <unistd.h>
|
||||
#include <string>
|
||||
#include <thread>
|
||||
#include <vector>
|
||||
|
||||
namespace
|
||||
{
|
||||
using auto_lock_t = std::unique_lock<std::mutex>;
|
||||
auto print_mutex = std::mutex{};
|
||||
auto global_kern_num = std::atomic<uint64_t>{0};
|
||||
} // namespace
|
||||
|
||||
template <typename T>
|
||||
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<unsigned int>(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<uint64_t>* progress,
|
||||
const std::shared_future<void>& 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<hipStream_t>(nstream, nullptr);
|
||||
auto stream_num = std::vector<int*>(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<hipGraph_t>(nstream);
|
||||
auto execs = std::vector<hipGraphExec_t>(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<uint64_t>(ndevice, ndevice_real);
|
||||
|
||||
auto progress = std::atomic<uint64_t>{0};
|
||||
auto promise = std::promise<void>{};
|
||||
auto future = promise.get_future().share();
|
||||
auto threads = std::vector<std::thread>{};
|
||||
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;
|
||||
}
|
||||
@@ -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}:$<TARGET_FILE:rocprofiler-sdk-json-tool>")
|
||||
else()
|
||||
set(PRELOAD_ENV "LD_PRELOAD=$<TARGET_FILE:rocprofiler-sdk-json-tool>")
|
||||
endif()
|
||||
|
||||
add_test(NAME test-hip-graph-tracing-execute COMMAND $<TARGET_FILE:hip-graph>)
|
||||
|
||||
set(hip-graph-tracing-env
|
||||
"${PRELOAD_ENV}"
|
||||
"ROCPROFILER_TOOL_OUTPUT_FILE=hip-graph-tracing-test.json"
|
||||
"LD_LIBRARY_PATH=$<TARGET_FILE_DIR:rocprofiler::rocprofiler-shared-library>:$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}")
|
||||
@@ -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)
|
||||
@@ -0,0 +1,5 @@
|
||||
|
||||
[pytest]
|
||||
addopts = --durations=20 -rA -s -vv
|
||||
testpaths = validate.py
|
||||
pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages
|
||||
@@ -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)
|
||||
@@ -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):
|
||||
|
||||
@@ -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 (
|
||||
|
||||
Criar uma nova questão referindo esta
Bloquear um utilizador