hipGraphs acts as a separate activity group as they have their own memory copies and their kernel dispatches and more activities, currently hip api that roctracer depend on doesn't have a way to represent this type, so we are temporarily suppressing the issue up till we have a complete support for the hipGraphs
Change-Id: I5d889be05c9414530672c781e5a712d572ea4104
[ROCm/rocprofiler commit: bbbd177dd7]
merge_traces.sh was depending on tblextr.py, and the location of tblextr.py changed from bin to libexec/rocprofiler
Change-Id: I76f1a0f41ffcff26e6a1119dacfb8b2ebd56d17d
[ROCm/rocprofiler commit: 7e8df53c44]
Some trace entries report the stream argument as 'stream=' and some report it as 'hStream='.
Change-Id: I9e7c7af162db72ddadf6f2a3fe4f44a04bf77ed6
[ROCm/rocprofiler commit: 3bb853c56d]
Enabling the new methodology of ROCP_STATS_OPT of getting HIP activities while the application is running
Change-Id: I19d09e2f2abe16e6528240bf55d6048645aaa09f
[ROCm/rocprofiler commit: 6f97e15d55]
This issue happens when we have slices that have overlapping
timestamp values (i.e. start_ts + dur = start_ts of next slice)
and the flow event references that shared timestamp value.
The google chrome parser sorts all the events
and just taking the most recent one.
Rocprof is outputting flow attachment points as start+dur;
the end ts of the slice that can overlap with the next slice
if no gap between them.
Solution proposal by Michael Steffen (Michael.Steffen@amd.com)
Change-Id: Ifcde142144033c9012b01a78c95f047384f972dd
[ROCm/rocprofiler commit: 93a79eb99a]
Using libroctracer_tool.so instead of libtracer_tool.so and fixing the paths for the tracer tool
Change-Id: I9b3ca885f3ca5385b106d5376894b1b4054f9c1d
[ROCm/rocprofiler commit: 0be6306cae]
In a future change, the tracer API library (libroctracer64.so) will be
automatically registered as a tool library. Until then, explicitly
register it by adding it to the HSA_TOOLS_LIB environment variable.
Change-Id: I44d78ac38608e6da5edf04b498a73485f5609d06
[ROCm/rocprofiler commit: 0123aa61fa]
HIP/HSA traces were asked to access range_data list, however, it was not initialized because roctx tracing was not enabled, moved lists initialization before roctx check
Change-Id: I9942876445cb1b2f69c6bb0d8986d6d9234f1441
[ROCm/rocprofiler commit: 1f925b3f16]
To enable this feature use the --roctx-rename rocprof option. This
implementation records all messages received in roctxPush calls and
use them to replace corresponding kernel names.
Tested with the following HIP program:
\#include <hip/hip_runtime.h>
\#include <roctracer/roctx.h>
__global__ void
ThisIsALongKernelName ()
{
}
int
main (int argc, char* argv[])
{
hipSetDevice (0);
// Not in a roctx range.
ThisIsALongKernelName<<<1, 1>>> ();
roctxRangePush ("A");
// In a simple first level roctx range.
ThisIsALongKernelName<<<1, 1>>> ();
roctxRangePop ();
roctxRangePush ("B");
roctxRangePush ("C");
// In a nested roctx range.
ThisIsALongKernelName<<<1, 1>>> ();
roctxRangePop ();
roctxRangePop ();
roctxRangePush ("D");
roctxRangePush ("E");
roctxRangePop ();
// In a first level roctx range, but after a nested range.
ThisIsALongKernelName<<<1, 1>>> ();
roctxRangePop ();
hipDeviceSynchronize ();
return 0;
}
Change-Id: I629312234468daff8b017caa5cb0773707d98cce
[ROCm/rocprofiler commit: 1078a088e9]
In a previous change the key for the var_table in tblextr.py script has been changed from one value to a tuple without changing the usage of the var_table in the rest of the script
Change-Id: I38964f61afad5323d1ca9b64d538cec426298842
[ROCm/rocprofiler commit: 46c4e5045a]
The Post-Processing script was depending HSA API call for async mem copies to correlate it with the HSA Async Memcpy Activity, now if user decided to include input file with filtering HSA Api calls without adding HSA Memcpy, then all the correlation data will be dropped and the Async activity will be reported with the information given from the HSA async activity result file
Change-Id: I5123a5acab9b35a4c25793e7953fdfb74929c999
[ROCm/rocprofiler commit: fd4767d954]
'merge_traces script from rocprof fails to include GPU / HSA / ROCTX activity in merged trace' change was missing tuple addition to the second for loop causing issues on gfx908 and gfx906 | change NO: 628475
Change-Id: Ic0b6140d4372eb109fdf7bdc8d58c0d84239196d
[ROCm/rocprofiler commit: 7a9692766d]
Changed var_pattern in tblextr.py to include pattern like "name[0]"
Change-Id: Ibe1c512595cfbdcaca8fa5bddceb3f6a570caf43
[ROCm/rocprofiler commit: ff43ca1542]
Changed derived metrics to double from int64.
Fixed standalone test due to int64 to float change
Fixed intercept test due to int64 to float change.
Change-Id: I49631c187406ae9dd94a869b3bb13772012e8cdf
[ROCm/rocprofiler commit: f9017cbdc5]
Removed the old code for trying to locate libhsakmt.so.1 as it is replaced by libhsakmt.a static library
Change-Id: Icc5a0f6ead285e2406e6e83614e536184e3a2663
[ROCm/rocprofiler commit: 14b62557d0]
Add roctx_trace to the list of files that need to be merged when
aggregating results from multiple runs.
Change-Id: I5810be9e9220765ed8e8a84eca854131e97e61b1
[ROCm/rocprofiler commit: 4ab94c410a]
Added Support for launch kernel functions to fill_api_db
Added support for hipMemcpyToSymbol in add_memcpy
Added support for hsa_amd_memory_pool_allocate to be counted as source of allocations
Change-Id: I68806106324b19ca6f09d413df37c27582be2f51
[ROCm/rocprofiler commit: 804e063eda]
When building the json data flow, from_us_list has (timestamp, stream_id, thread_id).
stream_id used to be interpreted as from_tid and tid as to_tid. But that's not correct.
stream_id is always a destination and tid is the initiator (source).
Change-Id: I2f5bb86a387b4003b17271c90bdf9de4b59a79bf
[ROCm/rocprofiler commit: 244dadcb85]
Marker events inside hcc_ops_trace.txt are from barriers so they are not meant to be stored in ops_patch_data map.
Added support for hipMemset events which are a kind of memory copy.
Change-Id: I213fe959bcd35ff0371613ba5bffd95bc53e06b5
[ROCm/rocprofiler commit: caa5f32300]
recordid cannot be just a counter. The code removed was doing
just that i.e. incrementing a counter. Recordid has to come
from recvals data structure. That code was left there since
a while when Evgeny and Rachida were trying to prototype this feature.
I am not sure why it was not spotted before.
Change-Id: Ia867066dcfca083fcd4111f2aefc2fec88c26314
[ROCm/rocprofiler commit: 4ba91a972c]
1st issue was that one of the ostream ops failed to print the
content of the struct.
2nd issue: get_ptr_type was called with args being src/dest
pointers while it should be the agents pointers for src/dest.
3rd issue: memcopies map used (recordid, procid, is_async)
as a key but this is not enough as some copies share same key,
so I added begin/end timestamps as a way to distinguish between them.
Change-Id: I7c6e80e74e30ea572f21612aaf0cf7efec6e91e6
[ROCm/rocprofiler commit: 761bd6a86b]
gobhardw
Sriraksha Nagaraj
Ammar ELWazir
Giovanni LB
Konstantin Zhuravlyov
Laurent Morichetti
Kiumars Sabeti
Ranjith Ramakrishnan
Saurabh Verma
Chun Yang
rachida
Christophe Paquot
Ashutosh Misra