Merge pull request #35 from rkebichi/rkebichi-trace-compare
Rkebichi trace compare
[ROCm/roctracer commit: 6dbfe515dc]
이 커밋은 다음에 포함됨:
eshcherb
GitHub
@@ -0,0 +1,135 @@
|
||||
#!/usr/bin/python
|
||||
|
||||
#Copyright (c) 2015-present 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, os, re
|
||||
import filecmp
|
||||
import argparse
|
||||
|
||||
events_count = {}
|
||||
events_order = {}
|
||||
trace2level = {}
|
||||
trace2level_filename = 'test/tests_trace_cmp_levels.txt'
|
||||
|
||||
def parse_trace_levels(filename):
|
||||
f = open(filename)
|
||||
trace2level = {}
|
||||
for line in f:
|
||||
item = line.split(' ', 1)
|
||||
trace2level[item[0]] = eval(item[1])
|
||||
return trace2level
|
||||
|
||||
# check trace againt golden reference and returns 0 for match, 1 for mismatch
|
||||
def check_trace_status(tracename):
|
||||
trace2level = parse_trace_levels(trace2level_filename)
|
||||
|
||||
trace = tracename + '.txt'
|
||||
rtrace = tracename + '_r.txt'
|
||||
if os.path.basename(tracename) in trace2level:
|
||||
trace_level = trace2level[os.path.basename(tracename)]
|
||||
print 'Trace comparison for ' + os.path.basename(tracename) + ' is at level ' + str(trace_level)
|
||||
else:
|
||||
print 'Trace ' + os.path.basename(tracename) + ' not found in ' + trace2level_filename + ', defaulting to level 0'
|
||||
return 0
|
||||
|
||||
if trace_level == 1:
|
||||
cnt_r = gen_events_info(rtrace,'cnt')
|
||||
cnt = gen_events_info(trace,'cnt')
|
||||
if cnt_r == cnt:
|
||||
return 0
|
||||
else:
|
||||
return 1
|
||||
elif trace_level == 2:
|
||||
cnt_r = gen_events_info(rtrace,'or')
|
||||
cnt = gen_events_info(trace,'or')
|
||||
if cnt_r == cnt:
|
||||
return 0
|
||||
else:
|
||||
return 1
|
||||
elif trace_level == 3:
|
||||
if filecmp.cmp(trace,rtrace):
|
||||
return 0
|
||||
else:
|
||||
return 1
|
||||
|
||||
#Parses roctracer trace file for regression purpose
|
||||
#and generates events count per event (when cnt is on) or events order per tid (when order is on)
|
||||
def gen_events_info(tracefile, metric):
|
||||
events_count = {}
|
||||
events_order = {}
|
||||
res=''
|
||||
with open(tracefile) as f:
|
||||
for line in f:
|
||||
event_pattern_s = re.compile(r'# START \((\d+)\) #############################')
|
||||
ms = event_pattern_s.match(line)
|
||||
if ms:
|
||||
start_id = ms.group(1)
|
||||
continue
|
||||
event_pattern = re.compile(r'.*<(\w+)\s+id\(\d+\)\s+.*tid\((\d+)\)>')
|
||||
# event_pattern extracts event(grp1) and tid (grp2) from a line like this:
|
||||
# <hsaKmtGetVersion id(2) correlation_id(0) on-enter pid(26224) tid(26224)>
|
||||
m = event_pattern.match(line)
|
||||
if m:
|
||||
event = m.group(1)
|
||||
tid = m.group(2)
|
||||
event_pattern2 = re.compile(r'\d+:\d+\s+\d+:(\d+)\s+(\w+)')
|
||||
# event_pattern2 extracts tid (grp1) and event (grp2) from a line like this:
|
||||
# 1822810364769411:1822810364771941 116477:116477 hsa_agent_get_info(<agent 0x8990e0>, 17, 0x7ffeac015fec) = 0
|
||||
m2 = event_pattern2.match(line)
|
||||
if m2:
|
||||
event = m2.group(2)
|
||||
tid = m2.group(1)
|
||||
event_pattern3 = re.compile(r'<rocTX "(.*)">')
|
||||
# event_pattern2 extracts rocTX event like:
|
||||
# <rocTX "before hipLaunchKernel">
|
||||
# <rocTX "hipLaunchKernel">
|
||||
m3 = event_pattern3.match(line)
|
||||
if m3:
|
||||
event = m3.group(1)
|
||||
tid = start_id
|
||||
if metric == 'cnt' and (m or m2 or m3):
|
||||
if event in events_count:
|
||||
events_count[event] = events_count[event] + 1
|
||||
else:
|
||||
events_count[event] = 1
|
||||
if metric == 'or' and (m or m2 or m3):
|
||||
if tid in events_order.keys():
|
||||
events_order[tid].append(event)
|
||||
else:
|
||||
events_order[tid] = [event]
|
||||
if metric == 'cnt':
|
||||
for event,count in events_count.items():
|
||||
res = res + event + " : count " + str(count) + '\n'
|
||||
if metric == 'or':
|
||||
for tid in sorted (events_order.keys()) :
|
||||
#res = res + 'Events for tid ' + tid + ' are:\n' + str(events_order[tid]) + '\n'
|
||||
res = res + str(events_order[tid])
|
||||
return res
|
||||
|
||||
|
||||
parser = argparse.ArgumentParser(description='check_trace.py: check a trace aainst golden ref. Returns 0 for success, 1 for failure')
|
||||
requiredNamed = parser.add_argument_group('Required arguments')
|
||||
requiredNamed.add_argument('-in', metavar='file', help='Name of trace to be checked', required=True)
|
||||
args = vars(parser.parse_args())
|
||||
|
||||
if __name__ == '__main__':
|
||||
sys.exit(check_trace_status(args['in']))
|
||||
|
||||
@@ -40,6 +40,10 @@ add_custom_target( mytest
|
||||
COMMAND sh -xc "cp ${TEST_DIR}/MatrixTranspose_test/MatrixTranspose ${PROJECT_BINARY_DIR}/test/MatrixTranspose_mgpu"
|
||||
COMMAND C_TEST=1 ${TEST_CFLAGS} make -C "${TEST_DIR}/MatrixTranspose_test"
|
||||
COMMAND sh -xc "cp ${TEST_DIR}/MatrixTranspose_test/MatrixTranspose ${PROJECT_BINARY_DIR}/test/MatrixTranspose_ctest"
|
||||
# copy traces
|
||||
COMMAND sh -xc "cp ${TEST_DIR}/MatrixTranspose*/*_trace.txt ${PROJECT_BINARY_DIR}/test/"
|
||||
COMMAND sh -xc "cp ${TEST_DIR}/*_trace.txt ${PROJECT_BINARY_DIR}/test/"
|
||||
COMMAND sh -xc "cp ${TEST_DIR}/tests_trace_cmp_levels.txt ${PROJECT_BINARY_DIR}/test/"
|
||||
)
|
||||
|
||||
## Util sources
|
||||
@@ -59,4 +63,4 @@ set ( TEST_DIR ${HSA_TEST_DIR} )
|
||||
add_subdirectory ( ${TEST_DIR} ${PROJECT_BINARY_DIR}/test/hsa )
|
||||
|
||||
## copying run script
|
||||
execute_process ( COMMAND sh -xc "cp ${RUN_SCRIPT} ${PROJECT_BINARY_DIR}" )
|
||||
execute_process ( COMMAND sh -xc "cp ${RUN_SCRIPT} ${PROJECT_BINARY_DIR}" )
|
||||
@@ -0,0 +1,503 @@
|
||||
standalone C test: "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_ctest"
|
||||
# INIT #############################
|
||||
# START (99) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (98) #############################
|
||||
PASSED!
|
||||
# START (97) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (96) #############################
|
||||
PASSED!
|
||||
# START (95) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (94) #############################
|
||||
PASSED!
|
||||
# START (93) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (92) #############################
|
||||
PASSED!
|
||||
# START (91) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (90) #############################
|
||||
PASSED!
|
||||
# START (89) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (88) #############################
|
||||
PASSED!
|
||||
# START (87) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (86) #############################
|
||||
PASSED!
|
||||
# START (85) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (84) #############################
|
||||
PASSED!
|
||||
# START (83) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (82) #############################
|
||||
PASSED!
|
||||
# START (81) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (80) #############################
|
||||
PASSED!
|
||||
# START (79) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (78) #############################
|
||||
PASSED!
|
||||
# START (77) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (76) #############################
|
||||
PASSED!
|
||||
# START (75) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (74) #############################
|
||||
PASSED!
|
||||
# START (73) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (72) #############################
|
||||
PASSED!
|
||||
# START (71) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (70) #############################
|
||||
PASSED!
|
||||
# START (69) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (68) #############################
|
||||
PASSED!
|
||||
# START (67) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (66) #############################
|
||||
PASSED!
|
||||
# START (65) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (64) #############################
|
||||
PASSED!
|
||||
# START (63) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (62) #############################
|
||||
PASSED!
|
||||
# START (61) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (60) #############################
|
||||
PASSED!
|
||||
# START (59) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (58) #############################
|
||||
PASSED!
|
||||
# START (57) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (56) #############################
|
||||
PASSED!
|
||||
# START (55) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (54) #############################
|
||||
PASSED!
|
||||
# START (53) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (52) #############################
|
||||
PASSED!
|
||||
# START (51) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (50) #############################
|
||||
PASSED!
|
||||
# START (49) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (48) #############################
|
||||
PASSED!
|
||||
# START (47) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (46) #############################
|
||||
PASSED!
|
||||
# START (45) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (44) #############################
|
||||
PASSED!
|
||||
# START (43) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (42) #############################
|
||||
PASSED!
|
||||
# START (41) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (40) #############################
|
||||
PASSED!
|
||||
# START (39) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (38) #############################
|
||||
PASSED!
|
||||
# START (37) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (36) #############################
|
||||
PASSED!
|
||||
# START (35) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (34) #############################
|
||||
PASSED!
|
||||
# START (33) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (32) #############################
|
||||
PASSED!
|
||||
# START (31) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (30) #############################
|
||||
PASSED!
|
||||
# START (29) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (28) #############################
|
||||
PASSED!
|
||||
# START (27) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (26) #############################
|
||||
PASSED!
|
||||
# START (25) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (24) #############################
|
||||
PASSED!
|
||||
# START (23) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (22) #############################
|
||||
PASSED!
|
||||
# START (21) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (20) #############################
|
||||
PASSED!
|
||||
# START (19) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (18) #############################
|
||||
PASSED!
|
||||
# START (17) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (16) #############################
|
||||
PASSED!
|
||||
# START (15) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (14) #############################
|
||||
PASSED!
|
||||
# START (13) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (12) #############################
|
||||
PASSED!
|
||||
# START (11) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (10) #############################
|
||||
PASSED!
|
||||
# START (9) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (8) #############################
|
||||
PASSED!
|
||||
# START (7) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (6) #############################
|
||||
PASSED!
|
||||
# START (5) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (4) #############################
|
||||
PASSED!
|
||||
# START (3) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (2) #############################
|
||||
PASSED!
|
||||
# START (1) #############################
|
||||
<rocTX "before hipLaunchKernel">
|
||||
<rocTX "hipLaunchKernel">
|
||||
<rocTX "after hipLaunchKernel">
|
||||
<rocTX "hipMemcpy">
|
||||
<rocTX "(null)">
|
||||
<rocTX "(null)">
|
||||
PASSED!
|
||||
# START (0) #############################
|
||||
PASSED!
|
||||
# STOP #############################
|
||||
@@ -0,0 +1,466 @@
|
||||
tool period test: "ROCP_CTRL_RATE=10:100000:1000000 ./test/MatrixTranspose"
|
||||
ROCTracer (pid=116438):
|
||||
HIP-trace()
|
||||
ROCTracer: trace control: delay(10us), length(100000us), rate(1000000us)
|
||||
Device name Vega 10 XT [Radeon RX Vega 64]
|
||||
## Iteration (99) #################
|
||||
PASSED!
|
||||
## Iteration (98) #################
|
||||
PASSED!
|
||||
## Iteration (97) #################
|
||||
PASSED!
|
||||
## Iteration (96) #################
|
||||
PASSED!
|
||||
## Iteration (95) #################
|
||||
PASSED!
|
||||
## Iteration (94) #################
|
||||
PASSED!
|
||||
## Iteration (93) #################
|
||||
PASSED!
|
||||
## Iteration (92) #################
|
||||
PASSED!
|
||||
## Iteration (91) #################
|
||||
PASSED!
|
||||
## Iteration (90) #################
|
||||
PASSED!
|
||||
## Iteration (89) #################
|
||||
PASSED!
|
||||
## Iteration (88) #################
|
||||
PASSED!
|
||||
## Iteration (87) #################
|
||||
PASSED!
|
||||
## Iteration (86) #################
|
||||
PASSED!
|
||||
## Iteration (85) #################
|
||||
PASSED!
|
||||
## Iteration (84) #################
|
||||
PASSED!
|
||||
## Iteration (83) #################
|
||||
PASSED!
|
||||
## Iteration (82) #################
|
||||
PASSED!
|
||||
## Iteration (81) #################
|
||||
PASSED!
|
||||
## Iteration (80) #################
|
||||
PASSED!
|
||||
## Iteration (79) #################
|
||||
PASSED!
|
||||
## Iteration (78) #################
|
||||
PASSED!
|
||||
## Iteration (77) #################
|
||||
PASSED!
|
||||
## Iteration (76) #################
|
||||
PASSED!
|
||||
## Iteration (75) #################
|
||||
PASSED!
|
||||
## Iteration (74) #################
|
||||
PASSED!
|
||||
## Iteration (73) #################
|
||||
PASSED!
|
||||
## Iteration (72) #################
|
||||
PASSED!
|
||||
## Iteration (71) #################
|
||||
PASSED!
|
||||
## Iteration (70) #################
|
||||
PASSED!
|
||||
## Iteration (69) #################
|
||||
PASSED!
|
||||
## Iteration (68) #################
|
||||
PASSED!
|
||||
## Iteration (67) #################
|
||||
PASSED!
|
||||
## Iteration (66) #################
|
||||
PASSED!
|
||||
## Iteration (65) #################
|
||||
PASSED!
|
||||
## Iteration (64) #################
|
||||
PASSED!
|
||||
## Iteration (63) #################
|
||||
PASSED!
|
||||
## Iteration (62) #################
|
||||
PASSED!
|
||||
## Iteration (61) #################
|
||||
PASSED!
|
||||
## Iteration (60) #################
|
||||
PASSED!
|
||||
## Iteration (59) #################
|
||||
PASSED!
|
||||
## Iteration (58) #################
|
||||
PASSED!
|
||||
## Iteration (57) #################
|
||||
PASSED!
|
||||
## Iteration (56) #################
|
||||
PASSED!
|
||||
## Iteration (55) #################
|
||||
PASSED!
|
||||
## Iteration (54) #################
|
||||
PASSED!
|
||||
## Iteration (53) #################
|
||||
PASSED!
|
||||
## Iteration (52) #################
|
||||
PASSED!
|
||||
## Iteration (51) #################
|
||||
PASSED!
|
||||
## Iteration (50) #################
|
||||
PASSED!
|
||||
## Iteration (49) #################
|
||||
PASSED!
|
||||
## Iteration (48) #################
|
||||
PASSED!
|
||||
## Iteration (47) #################
|
||||
PASSED!
|
||||
## Iteration (46) #################
|
||||
PASSED!
|
||||
## Iteration (45) #################
|
||||
PASSED!
|
||||
## Iteration (44) #################
|
||||
PASSED!
|
||||
## Iteration (43) #################
|
||||
PASSED!
|
||||
## Iteration (42) #################
|
||||
PASSED!
|
||||
## Iteration (41) #################
|
||||
PASSED!
|
||||
## Iteration (40) #################
|
||||
PASSED!
|
||||
## Iteration (39) #################
|
||||
PASSED!
|
||||
## Iteration (38) #################
|
||||
PASSED!
|
||||
## Iteration (37) #################
|
||||
PASSED!
|
||||
## Iteration (36) #################
|
||||
PASSED!
|
||||
## Iteration (35) #################
|
||||
PASSED!
|
||||
## Iteration (34) #################
|
||||
PASSED!
|
||||
## Iteration (33) #################
|
||||
PASSED!
|
||||
## Iteration (32) #################
|
||||
PASSED!
|
||||
## Iteration (31) #################
|
||||
PASSED!
|
||||
## Iteration (30) #################
|
||||
PASSED!
|
||||
## Iteration (29) #################
|
||||
PASSED!
|
||||
## Iteration (28) #################
|
||||
PASSED!
|
||||
## Iteration (27) #################
|
||||
PASSED!
|
||||
## Iteration (26) #################
|
||||
PASSED!
|
||||
## Iteration (25) #################
|
||||
PASSED!
|
||||
## Iteration (24) #################
|
||||
PASSED!
|
||||
## Iteration (23) #################
|
||||
PASSED!
|
||||
## Iteration (22) #################
|
||||
PASSED!
|
||||
## Iteration (21) #################
|
||||
PASSED!
|
||||
## Iteration (20) #################
|
||||
PASSED!
|
||||
## Iteration (19) #################
|
||||
PASSED!
|
||||
## Iteration (18) #################
|
||||
PASSED!
|
||||
## Iteration (17) #################
|
||||
PASSED!
|
||||
## Iteration (16) #################
|
||||
PASSED!
|
||||
## Iteration (15) #################
|
||||
PASSED!
|
||||
## Iteration (14) #################
|
||||
PASSED!
|
||||
## Iteration (13) #################
|
||||
PASSED!
|
||||
## Iteration (12) #################
|
||||
PASSED!
|
||||
## Iteration (11) #################
|
||||
PASSED!
|
||||
## Iteration (10) #################
|
||||
PASSED!
|
||||
## Iteration (9) #################
|
||||
PASSED!
|
||||
## Iteration (8) #################
|
||||
PASSED!
|
||||
## Iteration (7) #################
|
||||
PASSED!
|
||||
## Iteration (6) #################
|
||||
PASSED!
|
||||
## Iteration (5) #################
|
||||
PASSED!
|
||||
## Iteration (4) #################
|
||||
PASSED!
|
||||
## Iteration (3) #################
|
||||
PASSED!
|
||||
## Iteration (2) #################
|
||||
PASSED!
|
||||
## Iteration (1) #################
|
||||
PASSED!
|
||||
## Iteration (0) #################
|
||||
PASSED!
|
||||
1822770162190096:1822770162201236 116438:116438 hipGetDeviceProperties()
|
||||
1822770164732632:1822770164861753 116438:116438 hipMalloc(ptr(0x7fdfe4c00000) size(0x400000))
|
||||
1822770164862903:1822770164922593 116438:116438 hipMalloc(ptr(0x7fdfe4600000) size(0x400000))
|
||||
1822770164934104:1822770173240357 116438:116438 hipMemcpy(dst(0x7fdfe4c00000) src(0x20b7940) size(0x400000) kind(1))
|
||||
1822770173243647:1822770173243648 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770618561423:1822770618561424 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770636292136:1822770636292137 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770636318946:1822770636318947 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770649027427:1822770649027428 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770649055147:1822770649055148 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770661858819:1822770661858820 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770661886030:1822770661886031 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770674690761:1822770674690762 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770674719182:1822770674719183 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770687446783:1822770687446784 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770687474763:1822770687474764 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770700179044:1822770700179045 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770700207174:1822770700207175 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770713015755:1822770713015756 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770713043885:1822770713043886 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770725806637:1822770725806638 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770725833717:1822770725833718 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770738590118:1822770738590119 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770738618148:1822770738618149 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770751343939:1822770751343940 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770751372549:1822770751372550 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770764097681:1822770764097682 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770764136631:1822770764136632 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770776839992:1822770776839993 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770776867912:1822770776867913 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770789778774:1822770789778775 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770789807094:1822770789807095 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770802678836:1822770802678837 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770802707166:1822770802707167 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770815520828:1822770815520829 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770815548698:1822770815548699 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770828292289:1822770828292290 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770828320769:1822770828320770 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770841038531:1822770841038532 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770841065971:1822770841065972 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770853790232:1822770853790233 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770853820062:1822770853820063 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770866522263:1822770866522264 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770866550493:1822770866550494 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770879265854:1822770879265855 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770879294035:1822770879294036 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770892077246:1822770892077247 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770892105206:1822770892105207 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770904861617:1822770904861618 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770904892457:1822770904892458 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770917627368:1822770917627369 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770917655519:1822770917655520 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770930409670:1822770930409671 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770930438360:1822770930438361 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770943189151:1822770943189152 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770943217081:1822770943217082 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770956024813:1822770956024814 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770956052314:1822770956052315 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770968829485:1822770968829486 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770968857945:1822770968857946 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770981593826:1822770981593827 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770981621566:1822770981621567 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770994423678:1822770994423679 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822770994452568:1822770994452569 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771007194429:1822771007194430 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771007221919:1822771007221920 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771020222042:1822771020222043 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771020249972:1822771020249973 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771032985534:1822771032985535 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771033013724:1822771033013725 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771045714144:1822771045714145 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771045741935:1822771045741936 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771058596217:1822771058596218 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771058624617:1822771058624618 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771071441309:1822771071441310 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771071469609:1822771071469610 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771084319850:1822771084319851 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771084348361:1822771084348362 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771097128132:1822771097128133 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771097156612:1822771097156613 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771109883183:1822771109883184 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771109921563:1822771109921564 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771122729855:1822771122729856 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771122757155:1822771122757156 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771135568637:1822771135568638 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771135597657:1822771135597658 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771148375219:1822771148375220 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771148403629:1822771148403630 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771159180137:1822771161235110 116438:116438 hipMemcpy(dst(0x7fdfe4c00000) src(0x20b7940) size(0x400000) kind(1))
|
||||
1822771161236660:1822771161236661 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771161248381:1822771161264431 116438:116438 hipModuleLaunchKernel(kernel(matrixTranspose(float*, float*, int)) stream((nil)))
|
||||
1822771161285751:1822771161285752 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771161287281:1822771165514938 116438:116438 hipMemcpy(dst(0x24b7950) src(0x7fdfe4600000) size(0x400000) kind(2))
|
||||
1822771171918848:1822771173984782 116438:116438 hipMemcpy(dst(0x7fdfe4c00000) src(0x20b7940) size(0x400000) kind(1))
|
||||
1822771173986152:1822771173986153 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771173996942:1822771174012792 116438:116438 hipModuleLaunchKernel(kernel(matrixTranspose(float*, float*, int)) stream((nil)))
|
||||
1822771174014272:1822771174014273 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771174015382:1822771178279619 116438:116438 hipMemcpy(dst(0x24b7950) src(0x7fdfe4600000) size(0x400000) kind(2))
|
||||
1822771184673329:1822771186731413 116438:116438 hipMemcpy(dst(0x7fdfe4c00000) src(0x20b7940) size(0x400000) kind(1))
|
||||
1822771186732863:1822771186732864 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771186744233:1822771186760743 116438:116438 hipModuleLaunchKernel(kernel(matrixTranspose(float*, float*, int)) stream((nil)))
|
||||
1822771186761783:1822771186761784 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771186762873:1822771191066460 116438:116438 hipMemcpy(dst(0x24b7950) src(0x7fdfe4600000) size(0x400000) kind(2))
|
||||
1822771197562312:1822771199606165 116438:116438 hipMemcpy(dst(0x7fdfe4c00000) src(0x20b7940) size(0x400000) kind(1))
|
||||
1822771199607685:1822771199607686 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771199618535:1822771199634505 116438:116438 hipModuleLaunchKernel(kernel(matrixTranspose(float*, float*, int)) stream((nil)))
|
||||
1822771199635835:1822771199635836 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771199636915:1822771203924882 116438:116438 hipMemcpy(dst(0x24b7950) src(0x7fdfe4600000) size(0x400000) kind(2))
|
||||
1822771210390814:1822771212448587 116438:116438 hipMemcpy(dst(0x7fdfe4c00000) src(0x20b7940) size(0x400000) kind(1))
|
||||
1822771212450127:1822771212450128 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771212462917:1822771212479457 116438:116438 hipModuleLaunchKernel(kernel(matrixTranspose(float*, float*, int)) stream((nil)))
|
||||
1822771212480477:1822771212480478 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771212481567:1822771216792744 116438:116438 hipMemcpy(dst(0x24b7950) src(0x7fdfe4600000) size(0x400000) kind(2))
|
||||
1822771223198475:1822771225253498 116438:116438 hipMemcpy(dst(0x7fdfe4c00000) src(0x20b7940) size(0x400000) kind(1))
|
||||
1822771225255288:1822771225255289 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771225266888:1822771225282819 116438:116438 hipModuleLaunchKernel(kernel(matrixTranspose(float*, float*, int)) stream((nil)))
|
||||
1822771225283809:1822771225283810 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771225284909:1822771229566596 116438:116438 hipMemcpy(dst(0x24b7950) src(0x7fdfe4600000) size(0x400000) kind(2))
|
||||
1822771236041977:1822771238101610 116438:116438 hipMemcpy(dst(0x7fdfe4c00000) src(0x20b7940) size(0x400000) kind(1))
|
||||
1822771238103080:1822771238103081 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771238113820:1822771238129850 116438:116438 hipModuleLaunchKernel(kernel(matrixTranspose(float*, float*, int)) stream((nil)))
|
||||
1822771238130840:1822771238130841 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771238131950:1822771242369628 116438:116438 hipMemcpy(dst(0x24b7950) src(0x7fdfe4600000) size(0x400000) kind(2))
|
||||
1822771248769078:1822771250830021 116438:116438 hipMemcpy(dst(0x7fdfe4c00000) src(0x20b7940) size(0x400000) kind(1))
|
||||
1822771250831601:1822771250831602 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771250842911:1822771250858631 116438:116438 hipModuleLaunchKernel(kernel(matrixTranspose(float*, float*, int)) stream((nil)))
|
||||
1822771250859621:1822771250859622 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771250860701:1822771255145869 116438:116438 hipMemcpy(dst(0x24b7950) src(0x7fdfe4600000) size(0x400000) kind(2))
|
||||
1822771263595032:1822771263595033 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771263622582:1822771263622583 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771276404694:1822771276404695 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771276431864:1822771276431865 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771289132475:1822771289132476 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771289160596:1822771289160597 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771301911557:1822771301911558 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771301950127:1822771301950128 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771314764229:1822771314764230 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771314791509:1822771314791510 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771327616470:1822771327616471 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771327644471:1822771327644472 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771340467792:1822771340467793 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771340495392:1822771340495393 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771353284764:1822771353284765 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771353314034:1822771353314035 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771366027805:1822771366027806 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771366054975:1822771366054976 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771378747056:1822771378747057 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771378775616:1822771378775617 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771391520358:1822771391520359 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771391548328:1822771391548329 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771404259849:1822771404259850 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771404287629:1822771404287630 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771417078650:1822771417078651 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771417105521:1822771417105522 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771429921622:1822771429921623 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771429948912:1822771429948913 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771442703863:1822771442703864 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771442732114:1822771442732115 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771455546175:1822771455546176 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771455574765:1822771455574766 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771468309596:1822771468309597 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771468337556:1822771468337557 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771481124388:1822771481124389 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771481151899:1822771481151900 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771493874879:1822771493874880 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771493904120:1822771493904121 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771506687371:1822771506687372 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771506715871:1822771506715872 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771519464533:1822771519464534 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771519492293:1822771519492294 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771532279194:1822771532279195 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771532307234:1822771532307235 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771545016015:1822771545016016 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771545044205:1822771545044206 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771557801846:1822771557801847 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771557829627:1822771557829628 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771570644168:1822771570644169 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771570672029:1822771570672030 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771583412800:1822771583412801 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771583440560:1822771583440561 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771596207231:1822771596207232 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771596235072:1822771596235073 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771609169614:1822771609169615 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771609198584:1822771609198585 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771621935145:1822771621935146 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771621963045:1822771621963046 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771634755077:1822771634755078 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771634782347:1822771634782348 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771647514558:1822771647514559 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771647554188:1822771647554189 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771660410781:1822771660410782 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771660438721:1822771660438722 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771673205902:1822771673205903 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771673233972:1822771673233973 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771686004863:1822771686004864 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771686032954:1822771686032955 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771698793305:1822771698793306 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771698821805:1822771698821806 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771711576346:1822771711576347 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771711604186:1822771711604187 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771724367387:1822771724367388 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771724394878:1822771724394879 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771737130739:1822771737130740 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771737160079:1822771737160080 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771749894700:1822771749894701 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771749922980:1822771749922981 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771762658122:1822771762658123 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771762686292:1822771762686293 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771775443853:1822771775443854 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771775471433:1822771775471434 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771788218125:1822771788218126 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771788245825:1822771788245826 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771801042946:1822771801042947 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771801071416:1822771801071417 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771813839688:1822771813839689 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771813867578:1822771813867579 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771826878471:1822771826878472 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771826906261:1822771826906262 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771839641572:1822771839641573 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771839669053:1822771839669054 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771852458764:1822771852458765 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771852485704:1822771852485705 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771865247615:1822771865247616 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771865275756:1822771865275757 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771878163408:1822771878163409 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771878191438:1822771878191439 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822771891031999:1822771891032000 116438:116438 MARK(name(before HIP LaunchKernel))
|
||||
1822771891060780:1822771891060781 116438:116438 MARK(name(after HIP LaunchKernel))
|
||||
1822770164959094:1822770173233007 0:0 hcMemcpyHostToDevice:4
|
||||
1822771159196627:1822771161231010 0:0 hcMemcpyHostToDevice:89
|
||||
1822771161271486:1822771162462508 0:0 hcCommandKernel:91
|
||||
1822771162546779:1822771165512668 0:0 hcMemcpyDeviceToHost:93
|
||||
1822771171936428:1822771173981162 0:0 hcMemcpyHostToDevice:94
|
||||
1822771174019952:1822771175216159 0:0 hcCommandKernel:96
|
||||
1822771162467378:1822771162499379 0:0 hcCommandMarker:93
|
||||
1822771175297180:1822771178277099 0:0 hcMemcpyDeviceToHost:98
|
||||
1822771184690200:1822771186728003 0:0 hcMemcpyHostToDevice:99
|
||||
1822771186767974:1822771187965367 0:0 hcCommandKernel:101
|
||||
1822771175221178:1822771175252439 0:0 hcCommandMarker:98
|
||||
1822771188052511:1822771191063960 0:0 hcMemcpyDeviceToHost:103
|
||||
1822771197579732:1822771199602785 0:0 hcMemcpyHostToDevice:104
|
||||
1822771199641385:1822771200839370 0:0 hcCommandKernel:106
|
||||
1822771187970434:1822771188001695 0:0 hcCommandMarker:103
|
||||
1822771200920833:1822771203922192 0:0 hcMemcpyDeviceToHost:108
|
||||
1822771210408554:1822771212444577 0:0 hcMemcpyHostToDevice:109
|
||||
1822771212486611:1822771213670374 0:0 hcCommandKernel:111
|
||||
1822771200844255:1822771200875960 0:0 hcCommandMarker:108
|
||||
1822771213752345:1822771216790154 0:0 hcMemcpyDeviceToHost:113
|
||||
1822771223214605:1822771225250368 0:0 hcMemcpyHostToDevice:114
|
||||
1822771225289765:1822771226493528 0:0 hcCommandKernel:116
|
||||
1822771213675410:1822771213707115 0:0 hcCommandMarker:113
|
||||
1822771226574837:1822771229564026 0:0 hcMemcpyDeviceToHost:118
|
||||
1822771236058777:1822771238097360 0:0 hcMemcpyHostToDevice:119
|
||||
1822771238136777:1822771239325132 0:0 hcCommandKernel:121
|
||||
1822771226498570:1822771226530127 0:0 hcCommandMarker:118
|
||||
1822771239406108:1822771242367398 0:0 hcMemcpyDeviceToHost:123
|
||||
1822771248787248:1822771250826071 0:0 hcMemcpyHostToDevice:124
|
||||
1822771250866138:1822771252059234 0:0 hcCommandKernel:126
|
||||
1822771239330157:1822771239361418 0:0 hcCommandMarker:123
|
||||
1822771252142610:1822771255143279 0:0 hcMemcpyDeviceToHost:128
|
||||
1822771252064420:1822771252095680 0:0 hcCommandMarker:128
|
||||
파일 크기가 너무 크기때문에 변경 상태를 표시하지 않습니다.
Diff 로드
파일 크기가 너무 크기때문에 변경 상태를 표시하지 않습니다.
Diff 로드
파일 크기가 너무 크기때문에 변경 상태를 표시하지 않습니다.
Diff 로드
@@ -207,6 +207,9 @@ int main() {
|
||||
#include <inc/roctracer_kfd.h>
|
||||
#include <inc/roctracer_roctx.h>
|
||||
|
||||
#include <unistd.h>
|
||||
#include <sys/syscall.h> /* For SYS_xxx definitions */
|
||||
|
||||
// Macro to check ROC-tracer calls status
|
||||
#define ROCTRACER_CALL(call) \
|
||||
do { \
|
||||
@@ -217,6 +220,10 @@ int main() {
|
||||
} \
|
||||
} while (0)
|
||||
|
||||
static inline uint32_t GetTid() { return syscall(__NR_gettid); }
|
||||
static inline uint32_t GetPid() { return syscall(__NR_getpid); }
|
||||
|
||||
|
||||
// Runtime API callback function
|
||||
void api_callback(
|
||||
uint32_t domain,
|
||||
@@ -237,7 +244,8 @@ void api_callback(
|
||||
roctracer_op_string(ACTIVITY_DOMAIN_KFD_API, cid, 0),
|
||||
cid,
|
||||
data->correlation_id,
|
||||
(data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit");
|
||||
(data->phase == ACTIVITY_API_PHASE_ENTER) ? "on-enter" : "on-exit",
|
||||
);
|
||||
return;
|
||||
}
|
||||
const hip_api_data_t* data = (const hip_api_data_t*)(callback_data);
|
||||
@@ -249,23 +257,39 @@ void api_callback(
|
||||
if (data->phase == ACTIVITY_API_PHASE_ENTER) {
|
||||
switch (cid) {
|
||||
case HIP_API_ID_hipMemcpy:
|
||||
fprintf(stdout, "dst(%p) src(%p) size(0x%x) kind(%u)",
|
||||
fprintf(stdout, "<%s id(%u)\tcorrelation_id(%lu) %s>\n dst(%p) src(%p) size(0x%x) kind(%u)\n",
|
||||
roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0),
|
||||
cid,
|
||||
data->correlation_id,
|
||||
"on-enter",
|
||||
data->args.hipMemcpy.dst,
|
||||
data->args.hipMemcpy.src,
|
||||
(uint32_t)(data->args.hipMemcpy.sizeBytes),
|
||||
(uint32_t)(data->args.hipMemcpy.kind));
|
||||
break;
|
||||
case HIP_API_ID_hipMalloc:
|
||||
fprintf(stdout, "ptr(%p) size(0x%x)",
|
||||
fprintf(stdout, "<%s id(%u)\tcorrelation_id(%lu) %s>\n ptr(%p) size(0x%x)\n",
|
||||
roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0),
|
||||
cid,
|
||||
data->correlation_id,
|
||||
"on-enter",
|
||||
data->args.hipMalloc.ptr,
|
||||
(uint32_t)(data->args.hipMalloc.size));
|
||||
break;
|
||||
case HIP_API_ID_hipFree:
|
||||
fprintf(stdout, "ptr(%p)",
|
||||
fprintf(stdout, "<%s id(%u)\tcorrelation_id(%lu) %s>\n ptr(%p)\n",
|
||||
roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0),
|
||||
cid,
|
||||
data->correlation_id,
|
||||
"on-enter",
|
||||
data->args.hipFree.ptr);
|
||||
break;
|
||||
case HIP_API_ID_hipModuleLaunchKernel:
|
||||
fprintf(stdout, "kernel(\"%s\") stream(%p)",
|
||||
fprintf(stdout, "<%s id(%u)\tcorrelation_id(%lu) %s>\n kernel(\"%s\") stream(%p)\n",
|
||||
roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0),
|
||||
cid,
|
||||
data->correlation_id,
|
||||
"on-enter",
|
||||
hipKernelNameRef(data->args.hipModuleLaunchKernel.f),
|
||||
data->args.hipModuleLaunchKernel.stream);
|
||||
break;
|
||||
@@ -275,16 +299,20 @@ void api_callback(
|
||||
} else {
|
||||
switch (cid) {
|
||||
case HIP_API_ID_hipMalloc:
|
||||
fprintf(stdout, "*ptr(0x%p)",
|
||||
fprintf(stdout, "<%s id(%u)\tcorrelation_id(%lu) %s>\n *ptr(0x%p)\n",
|
||||
roctracer_op_string(ACTIVITY_DOMAIN_HIP_API, cid, 0),
|
||||
cid,
|
||||
data->correlation_id,
|
||||
"on-exit",
|
||||
*(data->args.hipMalloc.ptr));
|
||||
break;
|
||||
default:
|
||||
break;
|
||||
}
|
||||
}
|
||||
fprintf(stdout, "\n"); fflush(stdout);
|
||||
//fprintf(stdout, "\n");
|
||||
fflush(stdout);
|
||||
}
|
||||
|
||||
// Activity tracing callback
|
||||
// hipMalloc id(3) correlation_id(1): begin_ns(1525888652762640464) end_ns(1525888652762877067)
|
||||
void activity_callback(const char* begin, const char* end, void* arg) {
|
||||
@@ -305,7 +333,7 @@ void activity_callback(const char* begin, const char* end, void* arg) {
|
||||
record->thread_id
|
||||
);
|
||||
} else if (record->domain == ACTIVITY_DOMAIN_HCC_OPS) {
|
||||
fprintf(stdout, " device_id(%d) queue_id(%lu)",
|
||||
fprintf(stdout, " device_id(%d) queue_id(%lu)\n",
|
||||
record->device_id,
|
||||
record->queue_id
|
||||
);
|
||||
@@ -317,14 +345,13 @@ void activity_callback(const char* begin, const char* end, void* arg) {
|
||||
record->pc_sample.pc
|
||||
);
|
||||
} else if (record->domain == ACTIVITY_DOMAIN_EXT_API) {
|
||||
fprintf(stdout, " external_id(%lu)",
|
||||
fprintf(stdout, " external_id(%lu)\n",
|
||||
record->external_id
|
||||
);
|
||||
} else {
|
||||
fprintf(stderr, "Bad domain %d\n", record->domain);
|
||||
fprintf(stderr, "Bad domain %d\n\n", record->domain);
|
||||
abort();
|
||||
}
|
||||
fprintf(stdout, "\n");
|
||||
fflush(stdout);
|
||||
ROCTRACER_CALL(roctracer_next_record(record, &record));
|
||||
}
|
||||
|
||||
파일 크기가 너무 크기때문에 변경 상태를 표시하지 않습니다.
Diff 로드
@@ -0,0 +1,77 @@
|
||||
tool HSA test input: "./test/hsa/ctrl"
|
||||
ROCTracer (pid=116477): input from "input.xml"
|
||||
HSA-trace( hsa_agent_get_info hsa_amd_memory_pool_allocate)
|
||||
HSA-activity-trace()
|
||||
> GPU agents :
|
||||
> agent[0] :
|
||||
>> Name : gfx900
|
||||
>> APU : 0
|
||||
>> HSAIL profile : 0
|
||||
>> Max Wave Size : 64
|
||||
>> Max Queue Size : 131072
|
||||
>> CU number : 64
|
||||
>> Waves per CU : 40
|
||||
>> SIMDs per CU : 4
|
||||
>> SE number : 4
|
||||
>> Shader Arrays per SE : 1
|
||||
TestHsa::Initialize :
|
||||
> Using agent[0] : gfx900
|
||||
TestHsa::setup :
|
||||
Code object filename: gfx9_DummyKernel.hsaco
|
||||
TestHsa::run :
|
||||
> Executing kernel: "DummyKernel"
|
||||
> Waiting on kernel dispatch signal, que_idx=0
|
||||
> DONE, que_idx=0
|
||||
Test : Passed
|
||||
Time taken for Setup by DummyKernel : 0.000853027
|
||||
Time taken for Dispatch by DummyKernel : 2.39258e-05
|
||||
Time taken in Total by DummyKernel : 0.000876953
|
||||
TestHsa::Initialize :
|
||||
> Using agent[0] : gfx900
|
||||
TestHsa::setup :
|
||||
Code object filename: gfx9_SimpleConvolution.hsaco
|
||||
SimpleConvolution::init :
|
||||
> Input[0] :
|
||||
> 15 201 51 89 92 34 96 66 11 225 161 96 81 211 108 124 202 244 182 90 215 92 98 20 44 225 55 247 202 0 45 218 202 97 51 39 131 147 105 143 116 11 239 198 222 92 67 169 81 250 3 40 86 101 60 131 70 116 123 17 117 168 236 64
|
||||
> Mask :
|
||||
> 0 0.2 0
|
||||
> 0.2 0.2 0.2
|
||||
> 0 0.2 0
|
||||
TestHsa::run :
|
||||
> Executing kernel: "SimpleConvolution"
|
||||
> Waiting on kernel dispatch signal, que_idx=1
|
||||
> DONE, que_idx=1
|
||||
> Output[0] :
|
||||
> 45 60 89 75 79 86 45 43 104 82 144 105 99 90 109 124 123 146 149 124 120 87 43 36 88 91 113 103 98 53 68 104 113 106 76 90 90 122 82 92 102 124 95 149 112 102 69 82 146 116 103 62 50 96 99 87 84 110 88 81 61 105 134 71
|
||||
Test : Passed
|
||||
Time taken for Setup by SimpleConvolution : 0.000602051
|
||||
Time taken for Dispatch by SimpleConvolution : 4.00391e-05
|
||||
Time taken in Total by SimpleConvolution : 0.00064209
|
||||
1822810967909998:1822810967921554 async-copy0
|
||||
1822810364769411:1822810364771941 116477:116477 hsa_agent_get_info(<agent 0x8990e0>, 17, 0x7ffeac015fec) = 0
|
||||
1822810364778011:1822810364778611 116477:116477 hsa_agent_get_info(<agent 0x86f770>, 17, 0x7ffeac015fec) = 0
|
||||
1822810364779111:1822810364782691 116477:116477 hsa_agent_get_info(<agent 0x86f770>, 0, 0x8aac04) = 0
|
||||
1822810364784511:1822810364785051 116477:116477 hsa_agent_get_info(<agent 0x86f770>, 6, 0x8aac44) = 0
|
||||
1822810364785491:1822810364785841 116477:116477 hsa_agent_get_info(<agent 0x86f770>, 14, 0x8aac48) = 0
|
||||
1822810364786281:1822810364786611 116477:116477 hsa_agent_get_info(<agent 0x86f770>, 4, 0x8aac4c) = 0
|
||||
1822810364787051:1822810364787401 116477:116477 hsa_agent_get_info(<agent 0x86f770>, 40962, 0x8aac68) = 0
|
||||
1822810364787841:1822810364788181 116477:116477 hsa_agent_get_info(<agent 0x86f770>, 40970, 0x8aac6c) = 0
|
||||
1822810364788621:1822810364788951 116477:116477 hsa_agent_get_info(<agent 0x86f770>, 40971, 0x8aac70) = 0
|
||||
1822810364789381:1822810364789711 116477:116477 hsa_agent_get_info(<agent 0x86f770>, 40972, 0x8aac74) = 0
|
||||
1822810364790151:1822810364790471 116477:116477 hsa_agent_get_info(<agent 0x86f770>, 40973, 0x8aac78) = 0
|
||||
1822810966872900:1822810966889591 116477:116485 hsa_amd_memory_pool_allocate(<amd_memory_pool>, <uint64_t 0x4000>, <uint32_t 0x0>, 0x7fe988df8c68) = 0
|
||||
1822810966960311:1822810966961441 116477:116485 hsa_amd_memory_pool_allocate(<amd_memory_pool>, <uint64_t 0x4000>, <uint32_t 0x0>, 0x7fe988df8c70) = 0
|
||||
1822810966961951:1822810966973001 116477:116485 hsa_amd_memory_pool_allocate(<amd_memory_pool>, <uint64_t 0x1000>, <uint32_t 0x0>, 0x7fe988df8c68) = 0
|
||||
1822810967025981:1822810967035781 116477:116485 hsa_amd_memory_pool_allocate(<amd_memory_pool>, <uint64_t 0x1000>, <uint32_t 0x0>, 0x7fe988df8c68) = 0
|
||||
1822810967298203:1822810967317743 116477:116485 hsa_amd_memory_pool_allocate(<amd_memory_pool>, <uint64_t 0x4000>, <uint32_t 0x0>, 0x7fe988df8c68) = 0
|
||||
1822810967384054:1822810967384814 116477:116485 hsa_agent_get_info(<agent 0x8990e0>, 17, 0x7fe988df87bc) = 0
|
||||
1822810967387674:1822810967388054 116477:116485 hsa_agent_get_info(<agent 0x86f770>, 17, 0x7fe988df87bc) = 0
|
||||
1822810967388514:1822810967391744 116477:116485 hsa_agent_get_info(<agent 0x86f770>, 0, 0x7fe930006e04) = 0
|
||||
1822810967394494:1822810967394794 116477:116485 hsa_agent_get_info(<agent 0x86f770>, 6, 0x7fe930006e44) = 0
|
||||
1822810967395194:1822810967395514 116477:116485 hsa_agent_get_info(<agent 0x86f770>, 14, 0x7fe930006e48) = 0
|
||||
1822810967395924:1822810967396214 116477:116485 hsa_agent_get_info(<agent 0x86f770>, 4, 0x7fe930006e4c) = 0
|
||||
1822810967396604:1822810967396904 116477:116485 hsa_agent_get_info(<agent 0x86f770>, 40962, 0x7fe930006e68) = 0
|
||||
1822810967397294:1822810967397594 116477:116485 hsa_agent_get_info(<agent 0x86f770>, 40970, 0x7fe930006e6c) = 0
|
||||
1822810967397984:1822810967398284 116477:116485 hsa_agent_get_info(<agent 0x86f770>, 40971, 0x7fe930006e70) = 0
|
||||
1822810967398674:1822810967398974 116477:116485 hsa_agent_get_info(<agent 0x86f770>, 40972, 0x7fe930006e74) = 0
|
||||
1822810967399364:1822810967399654 116477:116485 hsa_agent_get_info(<agent 0x86f770>, 40973, 0x7fe930006e78) = 0
|
||||
@@ -0,0 +1,163 @@
|
||||
tool HSA test: "./test/hsa/ctrl"
|
||||
ROCTracer (pid=116451):
|
||||
HSA-trace()
|
||||
HSA-activity-trace()
|
||||
> GPU agents :
|
||||
> agent[0] :
|
||||
>> Name : gfx900
|
||||
>> APU : 0
|
||||
>> HSAIL profile : 0
|
||||
>> Max Wave Size : 64
|
||||
>> Max Queue Size : 131072
|
||||
>> CU number : 64
|
||||
>> Waves per CU : 40
|
||||
>> SIMDs per CU : 4
|
||||
>> SE number : 4
|
||||
>> Shader Arrays per SE : 1
|
||||
TestHsa::Initialize :
|
||||
> Using agent[0] : gfx900
|
||||
TestHsa::setup :
|
||||
Code object filename: gfx9_DummyKernel.hsaco
|
||||
TestHsa::run :
|
||||
> Executing kernel: "DummyKernel"
|
||||
> Waiting on kernel dispatch signal, que_idx=0
|
||||
> DONE, que_idx=0
|
||||
Test : Passed
|
||||
Time taken for Setup by DummyKernel : 0.000803955
|
||||
Time taken for Dispatch by DummyKernel : 2.68555e-05
|
||||
Time taken in Total by DummyKernel : 0.000830811
|
||||
TestHsa::Initialize :
|
||||
> Using agent[0] : gfx900
|
||||
TestHsa::setup :
|
||||
Code object filename: gfx9_SimpleConvolution.hsaco
|
||||
SimpleConvolution::init :
|
||||
> Input[0] :
|
||||
> 15 201 51 89 92 34 96 66 11 225 161 96 81 211 108 124 202 244 182 90 215 92 98 20 44 225 55 247 202 0 45 218 202 97 51 39 131 147 105 143 116 11 239 198 222 92 67 169 81 250 3 40 86 101 60 131 70 116 123 17 117 168 236 64
|
||||
> Mask :
|
||||
> 0 0.2 0
|
||||
> 0.2 0.2 0.2
|
||||
> 0 0.2 0
|
||||
TestHsa::run :
|
||||
> Executing kernel: "SimpleConvolution"
|
||||
> Waiting on kernel dispatch signal, que_idx=1
|
||||
> DONE, que_idx=1
|
||||
> Output[0] :
|
||||
> 45 60 89 75 79 86 45 43 104 82 144 105 99 90 109 124 123 146 149 124 120 87 43 36 88 91 113 103 98 53 68 104 113 106 76 90 90 122 82 92 102 124 95 149 112 102 69 82 146 116 103 62 50 96 99 87 84 110 88 81 61 105 134 71
|
||||
Test : Passed
|
||||
Time taken for Setup by SimpleConvolution : 0.00060498
|
||||
Time taken for Dispatch by SimpleConvolution : 4.29687e-05
|
||||
Time taken in Total by SimpleConvolution : 0.000647949
|
||||
1822787198892285:1822787198903100 async-copy0
|
||||
1822786593183215:1822786593187325 116451:116451 hsa_amd_profiling_async_copy_enable(<bool 0x1>) = 0
|
||||
1822786593205695:1822786593206405 116451:116451 hsa_agent_get_info(<agent 0x14290e0>, 17, 0x7ffd08170f5c) = 0
|
||||
1822786593211235:1822786593212485 116451:116451 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 0, 0x7ffd08170df8) = 0
|
||||
1822786593213335:1822786593213925 116451:116451 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 1, 0x7ffd08170dfc) = 0
|
||||
1822786593214745:1822786593215305 116451:116451 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 0, 0x7ffd08170df8) = 0
|
||||
1822786593216065:1822786593216635 116451:116451 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 1, 0x7ffd08170dfc) = 0
|
||||
1822786593216065:1822786593217435 116451:116451 hsa_amd_agent_iterate_memory_pools(<agent 0x14290e0>, 1, 0x143cfc0) = 1
|
||||
1822786593219035:1822786593219605 116451:116451 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 0, 0x7ffd08170df8) = 0
|
||||
1822786593220385:1822786593220945 116451:116451 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 1, 0x7ffd08170dfc) = 0
|
||||
1822786593220385:1822786593221715 116451:116451 hsa_amd_agent_iterate_memory_pools(<agent 0x14290e0>, 1, 0x143cfd0) = 1
|
||||
1822786593224085:1822786593224895 116451:116451 hsa_agent_get_info(<agent 0x13ff770>, 17, 0x7ffd08170f5c) = 0
|
||||
1822786593225835:1822786593231426 116451:116451 hsa_agent_get_info(<agent 0x13ff770>, 0, 0x143d094) = 0
|
||||
1822786593233936:1822786593234536 116451:116451 hsa_agent_get_info(<agent 0x13ff770>, 6, 0x143d0d4) = 0
|
||||
1822786593235326:1822786593235926 116451:116451 hsa_agent_get_info(<agent 0x13ff770>, 14, 0x143d0d8) = 0
|
||||
1822786593236696:1822786593237276 116451:116451 hsa_agent_get_info(<agent 0x13ff770>, 4, 0x143d0dc) = 0
|
||||
1822786593238046:1822786593238646 116451:116451 hsa_agent_get_info(<agent 0x13ff770>, 40962, 0x143d0f8) = 0
|
||||
1822786593239416:1822786593240136 116451:116451 hsa_agent_get_info(<agent 0x13ff770>, 40970, 0x143d0fc) = 0
|
||||
1822786593240946:1822786593241526 116451:116451 hsa_agent_get_info(<agent 0x13ff770>, 40971, 0x143d100) = 0
|
||||
1822786593242296:1822786593242886 116451:116451 hsa_agent_get_info(<agent 0x13ff770>, 40972, 0x143d104) = 0
|
||||
1822786593243666:1822786593244246 116451:116451 hsa_agent_get_info(<agent 0x13ff770>, 40973, 0x143d108) = 0
|
||||
1822786593245736:1822786593246326 116451:116451 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 0, 0x7ffd08170de8) = 0
|
||||
1822786593247096:1822786593247676 116451:116451 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 1, 0x7ffd08170dec) = 0
|
||||
1822786593247096:1822786593248456 116451:116451 hsa_amd_agent_iterate_memory_pools(<agent 0x13ff770>, 1, 0x143d0e8) = 1
|
||||
1822786593247096:1822786593250036 116451:116451 hsa_iterate_agents(1, 0x13fd320) = 0
|
||||
1822786593251826:1822786593416117 116451:116451 hsa_system_get_extension_table(<uint16_t 0x202>, <uint16_t 0x1>, <uint16_t 0x0>, 0x13fd3b8) = 0
|
||||
1822786593417707:1822786593420377 116451:116451 hsa_system_get_extension_table(<uint16_t 0x201>, <uint16_t 0x1>, <uint16_t 0x0>, 0x13fd408) = 0
|
||||
1822786593422447:1822786593423047 116451:116451 hsa_system_get_info(3, 0x7ffd08171128) = 0
|
||||
1822786593658778:1822786608352362 116451:116459 hsa_queue_create(<agent 0x13ff770>, <uint32_t 0x80>, <uint32_t 0x0>, 0, 0, <uint32_t 0xffffffff>, <uint32_t 0xffffffff>, <queue 0x1102d0e9006a>) = 0
|
||||
1822786902531377:1822786902535757 116451:116459 hsa_signal_create(1, <uint32_t 0x0>, 0, 0x7fa610001f20) = 0
|
||||
1822786902571037:1822786902584537 116451:116459 hsa_code_object_reader_create_from_file(5, 0x7fa665ffcbe0) = 0
|
||||
1822786902586167:1822786902597097 116451:116459 hsa_executable_create_alt(1, 0,
|
||||
1822786902598597:1822786902965610 116451:116459 hsa_executable_load_agent_code_object(<executable>, <agent 0x13ff770>, <code_object_reader>,
|
||||
1822786902967780:1822786903353172 116451:116459 hsa_executable_freeze(<executable>, ) = 0
|
||||
1822786903355162:1822786903356442 116451:116459 hsa_executable_get_symbol(<executable>,
|
||||
1822786903362512:1822786903363052 116451:116459 hsa_executable_symbol_get_info(<executable_symbol>, 13, 0x7fa665ffcc44) = 0
|
||||
1822786903363492:1822786903363782 116451:116459 hsa_executable_symbol_get_info(<executable_symbol>, 14, 0x7fa665ffcc48) = 0
|
||||
1822786903364172:1822786903364492 116451:116459 hsa_executable_symbol_get_info(<executable_symbol>, 22, 0x7fa665ffcc58) = 0
|
||||
1822786903365702:1822786903366022 116451:116459 hsa_signal_store_relaxed(<signal 0x7fa7a2d2cc80>, 1) = void
|
||||
1822786903370912:1822786903371252 116451:116459 hsa_queue_load_write_index_relaxed(0x7fa7a2ce1000) = 0
|
||||
1822786903372632:1822786903372942 116451:116459 hsa_queue_store_write_index_relaxed(0x7fa7a2ce1000, <uint64_t 0x1>) = void
|
||||
1822786903374072:1822786903374512 116451:116459 hsa_queue_load_read_index_relaxed(0x7fa7a2ce1000) = 0
|
||||
1822786903375092:1822786903375562 116451:116459 hsa_signal_store_relaxed(<signal 0x7fa7a2d2cd80>, 0) = void
|
||||
1822786903379112:1822786903393132 116451:116459 hsa_signal_wait_scacquire(<signal 0x7fa7a2d2cc80>, 2, 1, <uint64_t 0xffffffffffffffff>, 0) = 0
|
||||
1822786903425033:1822786903474693 116451:116459 hsa_executable_destroy(<executable>) = 0
|
||||
1822786903475683:1822786903476973 116451:116459 hsa_signal_destroy(<signal 0x7fa7a2d2cc80>) = 0
|
||||
1822787197612807:1822787197613567 116451:116459 hsa_signal_create(1, <uint32_t 0x0>, 0, 0x7fa610007340) = 0
|
||||
1822787197622077:1822787197628137 116451:116459 hsa_code_object_reader_create_from_file(7, 0x7fa665ffcbe0) = 0
|
||||
1822787197628677:1822787197629627 116451:116459 hsa_executable_create_alt(1, 0,
|
||||
1822787197630087:1822787197791048 116451:116459 hsa_executable_load_agent_code_object(<executable>, <agent 0x13ff770>, <code_object_reader>,
|
||||
1822787197792018:1822787197833668 116451:116459 hsa_executable_freeze(<executable>, ) = 0
|
||||
1822787197834218:1822787197839559 116451:116459 hsa_executable_get_symbol(<executable>,
|
||||
1822787197841769:1822787197856699 116451:116459 hsa_amd_memory_pool_allocate(<amd_memory_pool>, <uint64_t 0x4000>, <uint32_t 0x0>, 0x7fa665ffcc68) = 0
|
||||
1822787197858079:1822787197924589 116451:116459 hsa_amd_agents_allow_access(<uint32_t 0x1>, 0x7fa665ffcc70, 0, 0x7fa7a2cc4000) = 0
|
||||
1822787197926269:1822787197927179 116451:116459 hsa_amd_memory_pool_allocate(<amd_memory_pool>, <uint64_t 0x4000>, <uint32_t 0x0>, 0x7fa665ffcc70) = 0
|
||||
1822787197927659:1822787197939219 116451:116459 hsa_amd_memory_pool_allocate(<amd_memory_pool>, <uint64_t 0x1000>, <uint32_t 0x0>, 0x7fa665ffcc68) = 0
|
||||
1822787197939719:1822787197992339 116451:116459 hsa_amd_agents_allow_access(<uint32_t 0x1>, 0x7fa665ffcc70, 0, 0x7fa7a2cd6000) = 0
|
||||
1822787197993039:1822787197993399 116451:116459 hsa_executable_symbol_get_info(<executable_symbol>, 11, 0x7fa665ffcca0) = 0
|
||||
1822787197995009:1822787198005070 116451:116459 hsa_amd_memory_pool_allocate(<amd_memory_pool>, <uint64_t 0x1000>, <uint32_t 0x0>, 0x7fa665ffcc68) = 0
|
||||
1822787198005560:1822787198054310 116451:116459 hsa_amd_agents_allow_access(<uint32_t 0x1>, 0x7fa665ffcc70, 0, 0x7fa7a2cd4000) = 0
|
||||
1822787198222651:1822787198223061 116451:116459 hsa_executable_symbol_get_info(<executable_symbol>, 13, 0x7fa665ffcc44) = 0
|
||||
1822787198223531:1822787198223831 116451:116459 hsa_executable_symbol_get_info(<executable_symbol>, 14, 0x7fa665ffcc48) = 0
|
||||
1822787198224221:1822787198224531 116451:116459 hsa_executable_symbol_get_info(<executable_symbol>, 22, 0x7fa665ffcc58) = 0
|
||||
1822787198224991:1822787198225351 116451:116459 hsa_signal_store_relaxed(<signal 0x7fa7a2d2cc80>, 1) = void
|
||||
1822787198228821:1822787198229191 116451:116459 hsa_queue_load_write_index_relaxed(0x7fa7a2ce1000) = 1
|
||||
1822787198229651:1822787198229961 116451:116459 hsa_queue_store_write_index_relaxed(0x7fa7a2ce1000, <uint64_t 0x2>) = void
|
||||
1822787198230421:1822787198230851 116451:116459 hsa_queue_load_read_index_relaxed(0x7fa7a2ce1000) = 1
|
||||
1822787198231381:1822787198231871 116451:116459 hsa_signal_store_relaxed(<signal 0x7fa7a2d2cd80>, 1) = void
|
||||
1822787198234591:1822787198268311 116451:116459 hsa_signal_wait_scacquire(<signal 0x7fa7a2d2cc80>, 2, 1, <uint64_t 0xffffffffffffffff>, 0) = 0
|
||||
1822787198271101:1822787198289651 116451:116459 hsa_amd_memory_pool_allocate(<amd_memory_pool>, <uint64_t 0x4000>, <uint32_t 0x0>, 0x7fa665ffcc68) = 0
|
||||
1822787198290261:1822787198347092 116451:116459 hsa_amd_agents_allow_access(<uint32_t 0x1>, 0x7fa665ffcc70, 0, 0x7fa7a2c78000) = 0
|
||||
1822787198347792:1822787198349312 116451:116459 hsa_signal_create(1, <uint32_t 0x0>, 0, 0x7fa665ffcc40) = 0
|
||||
1822787198356932:1822787198357642 116451:116459 hsa_agent_get_info(<agent 0x14290e0>, 17, 0x7fa665ffc7bc) = 0
|
||||
1822787198359432:1822787198360132 116451:116459 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 0, 0x7fa665ffc658) = 0
|
||||
1822787198360542:1822787198360842 116451:116459 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 1, 0x7fa665ffc65c) = 0
|
||||
1822787198361262:1822787198361552 116451:116459 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 0, 0x7fa665ffc658) = 0
|
||||
1822787198361942:1822787198362232 116451:116459 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 1, 0x7fa665ffc65c) = 0
|
||||
1822787198361942:1822787198362632 116451:116459 hsa_amd_agent_iterate_memory_pools(<agent 0x14290e0>, 1, 0x7fa610006d60) = 1
|
||||
1822787198363522:1822787198363812 116451:116459 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 0, 0x7fa665ffc658) = 0
|
||||
1822787198364212:1822787198364502 116451:116459 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 1, 0x7fa665ffc65c) = 0
|
||||
1822787198364212:1822787198364892 116451:116459 hsa_amd_agent_iterate_memory_pools(<agent 0x14290e0>, 1, 0x7fa610006d70) = 1
|
||||
1822787198365732:1822787198366072 116451:116459 hsa_agent_get_info(<agent 0x13ff770>, 17, 0x7fa665ffc7bc) = 0
|
||||
1822787198366522:1822787198369532 116451:116459 hsa_agent_get_info(<agent 0x13ff770>, 0, 0x7fa610006e34) = 0
|
||||
1822787198372162:1822787198372472 116451:116459 hsa_agent_get_info(<agent 0x13ff770>, 6, 0x7fa610006e74) = 0
|
||||
1822787198372872:1822787198373172 116451:116459 hsa_agent_get_info(<agent 0x13ff770>, 14, 0x7fa610006e78) = 0
|
||||
1822787198373572:1822787198373862 116451:116459 hsa_agent_get_info(<agent 0x13ff770>, 4, 0x7fa610006e7c) = 0
|
||||
1822787198374262:1822787198374572 116451:116459 hsa_agent_get_info(<agent 0x13ff770>, 40962, 0x7fa610006e98) = 0
|
||||
1822787198375892:1822787198376202 116451:116459 hsa_agent_get_info(<agent 0x13ff770>, 40970, 0x7fa610006e9c) = 0
|
||||
1822787198376692:1822787198376992 116451:116459 hsa_agent_get_info(<agent 0x13ff770>, 40971, 0x7fa610006ea0) = 0
|
||||
1822787198377392:1822787198377692 116451:116459 hsa_agent_get_info(<agent 0x13ff770>, 40972, 0x7fa610006ea4) = 0
|
||||
1822787198378092:1822787198378392 116451:116459 hsa_agent_get_info(<agent 0x13ff770>, 40973, 0x7fa610006ea8) = 0
|
||||
1822787198379162:1822787198379462 116451:116459 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 0, 0x7fa665ffc648) = 0
|
||||
1822787198379852:1822787198380152 116451:116459 hsa_amd_memory_pool_get_info(<amd_memory_pool>, 1, 0x7fa665ffc64c) = 0
|
||||
1822787198379852:1822787198380542 116451:116459 hsa_amd_agent_iterate_memory_pools(<agent 0x13ff770>, 1, 0x7fa610006e88) = 1
|
||||
1822787198379852:1822787198381142 116451:116459 hsa_iterate_agents(1, 0x7fa610006980) = 0
|
||||
1822787198381692:1822787198397632 116451:116459 hsa_system_get_major_extension_table(<uint16_t 0x202>, <uint16_t 0x1>, <uint64_t 0x50>, 0x7fa610006a18) = 0
|
||||
1822787198398112:1822787198398412 116451:116459 hsa_system_get_major_extension_table(<uint16_t 0x201>, <uint16_t 0x1>, <uint64_t 0x18>, 0x7fa610006a68) = 0
|
||||
1822787198399032:1822787198399332 116451:116459 hsa_system_get_info(3, 0x7fa665ffc988) = 0
|
||||
1822787198399782:1822787198400322 116451:116459 hsa_system_get_info(2, 0x7fa665ffc9d0) = 0
|
||||
1822787198400762:1822787198401922 116451:116459 hsa_signal_create(1, <uint32_t 0x0>, 0, 0x7fa68f797050) = 0
|
||||
1822787198402472:1822787198405132 116451:116459 hsa_amd_signal_async_handler(<signal 0x7fa7a2d2cb00>, 2, 1, 1, 0x7fa68f797010) = 0
|
||||
1822787198405752:1822787198837475 116451:116459 hsa_amd_memory_async_copy(0x7fa7a2c78000, <agent 0x14290e0>, 0x7fa688203000, <agent 0x13ff770>, <uint64_t 0x4000>, <uint32_t 0x0>, 0, <signal 0x7fa7a2d2cb00>) = 0
|
||||
1822787198931945:1822787198941165 116451:116454 hsa_amd_profiling_get_async_copy_time(<signal 0x7fa7a2d2cb00>, 0x7fa7a092bc40) = 0
|
||||
1822787198944925:1822787198946035 116451:116454 hsa_system_get_info(2, 0x7fa7a092bc38) = 0
|
||||
1822787198947005:1822787198948155 116451:116454 hsa_signal_load_relaxed(<signal 0x7fa7a2d2cb80>) = 1
|
||||
1822787198838365:1822787198950945 116451:116459 hsa_signal_wait_scacquire(<signal 0x7fa7a2d2cb80>, 2, 1, <uint64_t 0xffffffffffffffff>, 0) = 0
|
||||
1822787198949035:1822787198951965 116451:116454 hsa_signal_store_screlease(<signal 0x7fa7a2d2cb80>, 0) = void
|
||||
1822787198952535:1822787198953995 116451:116459 hsa_signal_destroy(<signal 0x7fa7a2d2cb80>) = 0
|
||||
1822787198953835:1822787198955305 116451:116454 hsa_signal_destroy(<signal 0x7fa7a2d2cb00>) = 0
|
||||
1822787199054396:1822787199072966 116451:116459 hsa_memory_free(0x7fa7a2c78000) = 0
|
||||
1822787199096536:1822787199113536 116451:116459 hsa_executable_destroy(<executable>) = 0
|
||||
1822787199114426:1822787199116066 116451:116459 hsa_signal_destroy(<signal 0x7fa7a2d2cc80>) = 0
|
||||
1822787199118006:1822787200339674 116451:116459 hsa_queue_destroy(<queue 0x7fa7a2ce1000>) = 0
|
||||
|
||||
+22
-13
@@ -35,8 +35,8 @@ fi
|
||||
|
||||
# debugger
|
||||
debugger=""
|
||||
if [ -n "$2" ] ; then
|
||||
debugger=$2
|
||||
if [ -n "$3" ] ; then
|
||||
debugger=$3
|
||||
fi
|
||||
|
||||
# test check routin
|
||||
@@ -46,18 +46,27 @@ test_number=0
|
||||
xeval_test() {
|
||||
test_number=$test_number
|
||||
}
|
||||
|
||||
eval_test() {
|
||||
label=$1
|
||||
cmdline=$2
|
||||
rtrace="${3}_r.txt"
|
||||
if [ $test_filter = -1 -o $test_filter = $test_number ] ; then
|
||||
echo "$label: \"$cmdline\""
|
||||
test_runnum=$((test_runnum + 1))
|
||||
eval "$debugger $cmdline"
|
||||
eval "$debugger $cmdline" | tee $rtrace
|
||||
if [ $? != 0 ] ; then
|
||||
echo "$label: FAILED"
|
||||
test_status=$(($test_status + 1))
|
||||
else
|
||||
echo "$label: PASSED"
|
||||
echo "Comparing traces: ../script/check_trace.py -in $3"
|
||||
eval "../script/check_trace.py -in $3"
|
||||
if [ $? != 0 ] ; then
|
||||
echo "$label: FAILED trace comparison"
|
||||
test_status=$(($test_status + 1))
|
||||
else
|
||||
echo "$label: PASSED trace comparison"
|
||||
fi
|
||||
fi
|
||||
fi
|
||||
test_number=$((test_number + 1))
|
||||
@@ -65,9 +74,9 @@ eval_test() {
|
||||
|
||||
# Standalone test
|
||||
# rocTrecer is used explicitely by test
|
||||
eval_test "standalone C test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_ctest"
|
||||
eval_test "standalone HIP test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_test"
|
||||
eval_test "standalone HIP MGPU test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_mgpu"
|
||||
eval_test "standalone C test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_ctest" "test/MatrixTranspose_ctest_trace"
|
||||
eval_test "standalone HIP test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_test" "test/MatrixTranspose_test_trace"
|
||||
eval_test "standalone HIP MGPU test" "LD_PRELOAD=libkfdwrapper64.so ./test/MatrixTranspose_mgpu" "test/MatrixTranspose_mgpu_trace"
|
||||
|
||||
# Tool test
|
||||
# rocTracer/tool is loaded by HSA runtime
|
||||
@@ -75,13 +84,13 @@ export HSA_TOOLS_LIB="test/libtracer_tool.so"
|
||||
|
||||
# SYS test
|
||||
export ROCTRACER_DOMAIN="sys:roctx"
|
||||
eval_test "tool SYS test" ./test/MatrixTranspose
|
||||
eval_test "tool SYS test" ./test/MatrixTranspose "test/MatrixTranspose_sys_trace"
|
||||
export ROCTRACER_DOMAIN="sys:hsa:roctx"
|
||||
eval_test "tool SYS/HSA test" ./test/MatrixTranspose
|
||||
eval_test "tool SYS/HSA test" ./test/MatrixTranspose "test/MatrixTranspose_sys_hsa_trace"
|
||||
# Tracing control <delay:length:rate>
|
||||
export ROCTRACER_DOMAIN="hip"
|
||||
eval_test "tool period test" "ROCP_CTRL_RATE=10:100000:1000000 ./test/MatrixTranspose"
|
||||
eval_test "tool flushing test" "ROCP_FLUSH_RATE=100000 ./test/MatrixTranspose"
|
||||
eval_test "tool period test" "ROCP_CTRL_RATE=10:100000:1000000 ./test/MatrixTranspose" "test/MatrixTranspose_hip_trace"
|
||||
eval_test "tool flushing test" "ROCP_FLUSH_RATE=100000 ./test/MatrixTranspose" "test/MatrixTranspose_hip_trace_flush"
|
||||
|
||||
# HSA test
|
||||
export ROCTRACER_DOMAIN="hsa"
|
||||
@@ -98,11 +107,11 @@ export ROCP_AGENTS=1
|
||||
# each thread creates a queue pre GPU agent
|
||||
export ROCP_THRS=1
|
||||
|
||||
eval_test "tool HSA test" ./test/hsa/ctrl
|
||||
eval_test "tool HSA test" ./test/hsa/ctrl "test/ctrl_hsa_trace"
|
||||
|
||||
echo "<trace name=\"HSA\"><parameters api=\"hsa_agent_get_info, hsa_amd_memory_pool_allocate\"></parameters></trace>" > input.xml
|
||||
export ROCP_INPUT=input.xml
|
||||
eval_test "tool HSA test input" ./test/hsa/ctrl
|
||||
eval_test "tool HSA test input" ./test/hsa/ctrl "test/ctrl_hsa_input_trace"
|
||||
|
||||
#valgrind --leak-check=full $tbin
|
||||
#valgrind --tool=massif $tbin
|
||||
|
||||
@@ -0,0 +1,8 @@
|
||||
ctrl_hsa_input_trace 0
|
||||
ctrl_hsa_trace 1
|
||||
MatrixTranspose_ctest_trace 3
|
||||
MatrixTranspose_hip_trace 1
|
||||
MatrixTranspose_mgpu_trace 1
|
||||
MatrixTranspose_sys_hsa_trace 2
|
||||
MatrixTranspose_sys_trace 1
|
||||
MatrixTranspose_test_trace 1
|
||||
새 이슈에서 참조
사용자 차단