Madsen, Jonathan
8f891cdcc8
* [SDK][rocprofv3] HIP API buffer records with args (ext)
- New buffer tracing domain(s) for HIP APIs which include the arguments and the return value in the buffer records
- Update HIP stream support for extended HIP buffer tracing
- Update rocprofv3 tool library and output library to use extended HIP buffer tracing recods
* Update stream.cpp
- handle hipStream_t address being reused for a new stream
* Update doxygen docs for rocprofiler_iterate_buffer_tracing_record_args
* Update rocprofv3 tool.cpp
- configure buffer tracing services with HIP_*_API_EXT variants
- tweak logging level for hip_stream_display_callback
* Fix validation tests
- add HIP_RUNTIME_API_EXT and HIP_COMPILER_API_EXT to valid domain names
* Serialization support for buffer tracing args
* Disable stream service for __hipPopCallConfiguration
- this is interpreted as a stream create but it doesn't create a stream
* Fix execute_buffer_record_emplace for HIP extended contexts
* Add uint64_t_retval to rocprofiler_hip_api_retval_t union
- reading in hipError_t_retval during serialization of pointer return value causes undefined behavior
* Fix compilation warning about unused but set parameter
- in hip/stream.cpp
* Add synchronization for async_copy_data
* Fix compilation error
* Fix compilation error
---------
Co-authored-by: Jonathan R. Madsen <jonathanrmadsen@gmail.com>
[ROCm/rocprofiler-sdk commit: e33dff7ad0]
133 baris
4.5 KiB
Python
133 baris
4.5 KiB
Python
#!/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
|
|
import re
|
|
|
|
kernel_trace_list = sorted(["addition_kernel", "subtract_kernel"])
|
|
kernel_counter_list = ["addition_kernel"]
|
|
|
|
|
|
def unique(lst):
|
|
return list(set(lst))
|
|
|
|
|
|
def test_counter_collection_json_data(json_data):
|
|
data = json_data["rocprofiler-sdk-tool"]
|
|
|
|
def get_kernel_name(kernel_id):
|
|
return data["kernel_symbols"][kernel_id]["formatted_kernel_name"]
|
|
|
|
counter_collection_data = data["callback_records"]["counter_collection"]
|
|
|
|
for counter in counter_collection_data:
|
|
kernel_name = get_kernel_name(counter.dispatch_data.dispatch_info.kernel_id)
|
|
assert kernel_name in kernel_counter_list
|
|
|
|
|
|
def test_kernel_trace_json(json_data):
|
|
data = json_data["rocprofiler-sdk-tool"]
|
|
|
|
def get_kernel_name(kernel_id):
|
|
return data["kernel_symbols"][kernel_id]["formatted_kernel_name"]
|
|
|
|
def get_kind_name(kind_id):
|
|
return data["strings"]["buffer_records"][kind_id]["kind"]
|
|
|
|
kernel_dispatch_data = data["buffer_records"]["kernel_dispatch"]
|
|
kernels = []
|
|
assert len(kernel_dispatch_data) == 2
|
|
for dispatch in kernel_dispatch_data:
|
|
dispatch_info = dispatch["dispatch_info"]
|
|
kernel_name = get_kernel_name(dispatch_info["kernel_id"])
|
|
|
|
assert get_kind_name(dispatch["kind"]) == "KERNEL_DISPATCH"
|
|
assert dispatch["correlation_id"]["internal"] > 0
|
|
assert dispatch_info["agent_id"]["handle"] > 0
|
|
assert dispatch_info["queue_id"]["handle"] > 0
|
|
assert dispatch_info["kernel_id"] > 0
|
|
if not re.search(r"__amd_rocclr_.*", kernel_name):
|
|
kernels.append(kernel_name)
|
|
|
|
assert dispatch_info["workgroup_size"]["x"] == 64
|
|
assert dispatch_info["workgroup_size"]["y"] == 1
|
|
assert dispatch_info["workgroup_size"]["z"] == 1
|
|
assert dispatch_info["grid_size"]["x"] == 1024
|
|
assert dispatch_info["grid_size"]["y"] == 1024
|
|
assert dispatch_info["grid_size"]["z"] == 1
|
|
assert dispatch["end_timestamp"] >= dispatch["start_timestamp"]
|
|
|
|
assert kernels == kernel_trace_list
|
|
|
|
|
|
def test_hip_api_trace_json(json_data):
|
|
data = json_data["rocprofiler-sdk-tool"]
|
|
|
|
def get_operation_name(kind_id, op_id):
|
|
return data["strings"]["buffer_records"][kind_id]["operations"][op_id]
|
|
|
|
def get_kind_name(kind_id):
|
|
return data["strings"]["buffer_records"][kind_id]["kind"]
|
|
|
|
valid_domain_names = (
|
|
"HIP_RUNTIME_API",
|
|
"HIP_RUNTIME_API_EXT",
|
|
)
|
|
|
|
hip_api_data = data["buffer_records"]["hip_api"]
|
|
|
|
functions = []
|
|
for api in hip_api_data:
|
|
kind = get_kind_name(api["kind"])
|
|
assert kind in valid_domain_names
|
|
assert api["end_timestamp"] >= api["start_timestamp"]
|
|
functions.append(get_operation_name(api["kind"], api["operation"]))
|
|
|
|
expected_functions = (
|
|
[
|
|
"hipGetDeviceCount",
|
|
"hipSetDevice",
|
|
"hipDeviceSynchronize",
|
|
"hipStreamCreateWithFlags",
|
|
]
|
|
+ (["hipHostMalloc"] * 3)
|
|
+ (["hipMallocAsync"] * 3)
|
|
+ (["hipMemcpyAsync"] * 2)
|
|
+ [
|
|
"hipStreamSynchronize",
|
|
"hipDeviceSynchronize",
|
|
"hipLaunchKernel",
|
|
"hipGetLastError",
|
|
"hipLaunchKernel",
|
|
"hipGetLastError",
|
|
]
|
|
)
|
|
assert functions == expected_functions
|
|
|
|
|
|
if __name__ == "__main__":
|
|
exit_code = pytest.main(["-x", __file__] + sys.argv[1:])
|
|
sys.exit(exit_code)
|