diff --git a/inc/rocprofiler.h b/inc/rocprofiler.h index 26c332d4d4..5c76d5312a 100644 --- a/inc/rocprofiler.h +++ b/inc/rocprofiler.h @@ -471,7 +471,7 @@ typedef enum { /** * Represents a ATT tracing record (Not available yet) */ - ROCPROFILER_ATT_RECORD = 2, + ROCPROFILER_ATT_TRACER_RECORD = 2, /** * Represents a PC sampling record */ @@ -1046,6 +1046,66 @@ typedef struct { } rocprofiler_record_spm_t; +/** + * struct to store the trace data from a shader engine. + */ +typedef struct { + void* buffer_ptr; + uint32_t buffer_size; +} rocprofiler_record_se_att_data_t; + + /** + * ATT tracing record structure. + * This will represent all the information reported by the + * ATT tracer such as the kernel and its thread trace data. + * This record can be flushed to the user using + * ::rocmtools_buffer_callback_t + */ +typedef struct { + /** + * ROCMtool General Record base header to identify the id and kind of every + * record + */ + rocprofiler_record_header_t header; + /** + * Kernel Identifier to be used by the user to get the kernel info using + * ::rocmtools_query_kernel_info + */ + rocprofiler_kernel_id_t kernel_id; + /** + * Agent Identifier to be used by the user to get the Agent Information using + * ::rocmtools_query_agent_info + */ + rocprofiler_agent_id_t gpu_id; + /** + * Queue Identifier to be used by the user to get the Queue Information using + * ::rocmtools_query_agent_info + */ + rocprofiler_queue_id_t queue_id; + /** + * kernel properties, including the grid size, work group size, + * registers count, wave size and completion signal + */ + rocprofiler_kernel_properties_t kernel_properties; + /** + * Thread id + */ + rocprofiler_thread_id_t thread_id; + /** + * Queue Index - packet index in the queue + */ + rocprofiler_queue_index_t queue_idx; + /** + * ATT data output from each shader engine. + */ + rocprofiler_record_se_att_data_t* shader_engine_data; + /** + * The count of the shader engine ATT data + */ + uint64_t shader_engine_data_count; +} rocprofiler_record_att_tracer_t; + + /** @} */ @@ -1721,7 +1781,7 @@ typedef enum { /** * ATT Tracing. (Not Yet Supported) */ - ROCPROFILER_ATT_TRACE = 4, + ROCPROFILER_ATT_TRACE_COLLECTION = 4, /** * SPM collection. (Not Yet Supported) */ @@ -1764,6 +1824,31 @@ typedef enum { typedef const char* rocprofiler_hip_function_name_t; typedef const char* rocprofiler_hsa_function_name_t; +// ATT tracing parameter names +typedef enum { + ROCPROFILER_ATT_COMPUTE_UNIT_TARGET = 0, + ROCPROFILER_ATT_VM_ID_MASK = 1, + ROCPROFILER_ATT_MASK = 2, + ROCPROFILER_ATT_TOKEN_MASK = 3, + ROCPROFILER_ATT_TOKEN_MASK2 = 4, + ROCPROFILER_ATT_SE_MASK = 5, + ROCPROFILER_ATT_SAMPLE_RATE = 6, + ROCPROFILER_ATT_PERF_MASK = 240, + ROCPROFILER_ATT_PERF_CTRL = 241, + ROCPROFILER_ATT_PERFCOUNTER = 242, + ROCPROFILER_ATT_PERFCOUNTER_NAME = 243, + ROCPROFILER_ATT_MAXVALUE +} rocprofiler_att_parameter_name_t; + +// att tracing parameters object +typedef struct { + rocprofiler_att_parameter_name_t parameter_name; + union { + uint32_t value; + const char* counter_name; + }; +} rocprofiler_att_parameter_t; + /** * Filter Data Type * filter data will be used to report required and optional filters for the @@ -1822,6 +1907,10 @@ typedef union { * Counters to profile */ const char** counters_names; + /** + * att parameters + */ + rocprofiler_att_parameter_t* att_parameters; /** * spm counters parameters */ diff --git a/plugin/CMakeLists.txt b/plugin/CMakeLists.txt index 0726b5bfb8..1576b72a9b 100644 --- a/plugin/CMakeLists.txt +++ b/plugin/CMakeLists.txt @@ -23,3 +23,4 @@ add_subdirectory(file) add_subdirectory(perfetto) add_subdirectory(ctf) +add_subdirectory(att) diff --git a/plugin/att/CMakeLists.txt b/plugin/att/CMakeLists.txt new file mode 100644 index 0000000000..8da375ffef --- /dev/null +++ b/plugin/att/CMakeLists.txt @@ -0,0 +1,56 @@ +# ############################################################################### +# # Copyright (c) 2022 Advanced Micro Devices, Inc. +# # +# # 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. +# ############################################################################### + +# Building att plugin library +file(GLOB ROCMTOOLS_UTIL_SRC_FILES ${PROJECT_SOURCE_DIR}/src/utils/helper.cpp) +file(GLOB FILE_SOURCES att.cpp) +add_library(att_plugin SHARED ${FILE_SOURCES} ${ROCMTOOLS_UTIL_SRC_FILES}) + +set_target_properties(att_plugin PROPERTIES + CXX_VISIBILITY_PRESET hidden + LINK_DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/../exportmap + LIBRARY_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR}) + +target_compile_definitions(att_plugin + PRIVATE HIP_PROF_HIP_API_STRING=1 __HIP_PLATFORM_HCC__=1) + +target_include_directories(att_plugin PRIVATE ${PROJECT_SOURCE_DIR}/inc ${PROJECT_SOURCE_DIR} ${CMAKE_CURRENT_SOURCE_DIR}) +target_link_options(att_plugin PRIVATE -Wl,--version-script=${CMAKE_CURRENT_SOURCE_DIR}/../exportmap -Wl,--no-undefined) +target_link_libraries(att_plugin PRIVATE ${ROCPROFILER_TARGET} systemd hsa-runtime64::hsa-runtime64 stdc++fs) + +install(TARGETS att_plugin LIBRARY + DESTINATION ${CMAKE_INSTALL_LIBDIR}/${PROJECT_NAME} + COMPONENT runtime) + +configure_file(att.py att/att.py COPYONLY) +configure_file(trace_view.py att/trace_view.py COPYONLY) +#configure_file(t.db att/t.db COPYONLY) +configure_file(ui/index.html att/ui/index.html COPYONLY) +configure_file(ui/logo.svg att/ui/logo.svg COPYONLY) +configure_file(ui/styles.css att/ui/styles.css COPYONLY) +#configure_file(ui/trace.json att/ui/trace.json COPYONLY) +install(DIRECTORY + ${CMAKE_CURRENT_BINARY_DIR}/att + DESTINATION ${CMAKE_INSTALL_LIBEXECDIR}/rocprofiler + USE_SOURCE_PERMISSIONS + COMPONENT runtime) + diff --git a/plugin/att/att.cpp b/plugin/att/att.cpp new file mode 100644 index 0000000000..7e3d4dc4c5 --- /dev/null +++ b/plugin/att/att.cpp @@ -0,0 +1,189 @@ +/* Copyright (c) 2022 Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include "rocprofiler.h" +#include "rocprofiler_plugin.h" +#include "../utils.h" + +namespace { + +class att_plugin_t { + public: + att_plugin_t() {} + + std::mutex writing_lock; + bool is_valid_{true}; + + inline bool att_file_exists(const std::string& name) { + struct stat buffer; + return stat(name.c_str(), &buffer) == 0; + } + + bool IsValid() const { return is_valid_; } + + void FlushATTRecord(const rocprofiler_record_att_tracer_t* att_tracer_record, + rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { + std::lock_guard lock(writing_lock); + + if (!att_tracer_record) { + printf("No att data buffer received\n"); + return; + } + + size_t name_length; + CHECK_ROCMTOOLS(rocprofiler_query_kernel_info_size(ROCPROFILER_KERNEL_NAME, + att_tracer_record->kernel_id, &name_length)); + const char* kernel_name_c = static_cast(malloc(name_length * sizeof(char))); + CHECK_ROCMTOOLS(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, + att_tracer_record->kernel_id, &kernel_name_c)); + + std::string name_demangled = rocmtools::truncate_name(rocmtools::cxx_demangle(kernel_name_c)); + + // Get the number of shader engine traces + int se_num = att_tracer_record->shader_engine_data_count; + + // Find if this filename already exists. If so, increment vname. + int file_iteration = -1; + bool bIncrementVersion = true; + while(bIncrementVersion) { + file_iteration += 1; + std::string fss = name_demangled+"_v"+std::to_string(file_iteration); + bIncrementVersion = att_file_exists(fss + "_kernel.txt"); + } + + std::string fname = name_demangled+"_v"+std::to_string(file_iteration)+"_kernel.txt"; + std::ofstream(fname.c_str()) << name_demangled << ": " << kernel_name_c << '\n'; + + // iterate over each shader engine att trace + for (int i = 0; i < se_num; i++) { + if (!att_tracer_record->shader_engine_data && + !att_tracer_record->shader_engine_data[i].buffer_ptr) + continue; + printf("--------------collecting data for shader_engine %d---------------\n", i); + rocprofiler_record_se_att_data_t* se_att_trace = &att_tracer_record->shader_engine_data[i]; + uint32_t size = se_att_trace->buffer_size; + const char* data_buffer_ptr = reinterpret_cast(se_att_trace->buffer_ptr); + + // dump data in binary format + std::ostringstream oss; + oss << name_demangled << "_v" << file_iteration << "_se" << i << ".att"; + std::ofstream out(oss.str().c_str(), std::ios::binary); + if (out.is_open()) { + out.write((char*)data_buffer_ptr, size); + out.close(); + } else { + std::cerr << "\t" << __FUNCTION__ << " Failed to open file: " << oss.str().c_str() << '\n'; + } + } + } + + int WriteBufferRecords(const rocprofiler_record_header_t* begin, + const rocprofiler_record_header_t* end, rocprofiler_session_id_t session_id, + rocprofiler_buffer_id_t buffer_id) { + while (begin < end) { + if (!begin) return 0; + switch (begin->kind) { + case ROCPROFILER_PROFILER_RECORD: + case ROCPROFILER_TRACER_RECORD: + case ROCPROFILER_PC_SAMPLING_RECORD: + case ROCPROFILER_SPM_RECORD: + printf("Invalid record Kind: %d", begin->kind); + break; + + case ROCPROFILER_ATT_TRACER_RECORD: { + rocprofiler_record_att_tracer_t* att_record = const_cast( + reinterpret_cast(begin)); + FlushATTRecord(att_record, session_id, buffer_id); + break; + } + } + rocprofiler_next_record(begin, &begin, session_id, buffer_id); + } + + return 0; + } + + private: +}; + +att_plugin_t* att_plugin = nullptr; + +} // namespace + +ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_version, + uint32_t rocprofiler_minor_version) { + if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR || + rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR) + return -1; + + if (att_plugin != nullptr) return -1; + + att_plugin = new att_plugin_t(); + if (att_plugin->IsValid()) return 0; + + // The plugin failed to initialied, destroy it and return an error. + delete att_plugin; + att_plugin = nullptr; + return -1; +} + +ROCPROFILER_EXPORT void rocprofiler_plugin_finalize() { + if (!att_plugin) return; + delete att_plugin; + att_plugin = nullptr; +} + +ROCPROFILER_EXPORT int rocprofiler_plugin_write_buffer_records(const rocprofiler_record_header_t* begin, + const rocprofiler_record_header_t* end, + rocprofiler_session_id_t session_id, + rocprofiler_buffer_id_t buffer_id) { + if (!att_plugin || !att_plugin->IsValid()) return -1; + return att_plugin->WriteBufferRecords(begin, end, session_id, buffer_id); +} + +ROCPROFILER_EXPORT int rocprofiler_plugin_write_record(rocprofiler_record_tracer_t record, + rocprofiler_session_id_t session_id) { + if (!att_plugin || !att_plugin->IsValid()) return -1; + if (record.header.id.handle == 0) return 0; + return 0; +} diff --git a/plugin/att/att.py b/plugin/att/att.py new file mode 100755 index 0000000000..a913a147c1 --- /dev/null +++ b/plugin/att/att.py @@ -0,0 +1,523 @@ +#!/usr/bin/env python3 +import sys +if sys.version_info[0] < 3: + raise Exception("Must be using Python 3") + +import os +import argparse +from pathlib import Path +from struct import * +from ctypes import * +import ctypes +from copy import deepcopy +from trace_view import view_trace +import sys +import glob +import numpy as np +import matplotlib.pyplot as plt +import json + +class PerfEvent(ctypes.Structure): + _fields_ = [ + ('time', c_uint64), + ('event0', c_uint16), + ('event1', c_uint16), + ('event2', c_uint16), + ('event3', c_uint16), + ('cu', c_uint8), + ('bank', c_uint8), + ] + def toTuple(self): + return (int(self.time), int(self.event0), int(self.event1), + int(self.event2), int(self.event3), int(self.cu), int(self.bank)) + + +class CodeWrapped(ctypes.Structure): + """ Matches CodeWrapped on the python side """ + _fields_ = [('line', ctypes.c_char_p), + ('loc', ctypes.c_char_p), + ('value', ctypes.c_int), + ('to_line', ctypes.c_int), + ('index', ctypes.c_int), + ('line_num', ctypes.c_int)] + + +class KvPair(ctypes.Structure): + """ Matches pair = (key, value) on the python side """ + _fields_ = [('key', ctypes.c_int), + ('value', ctypes.c_int)] + + +class ReturnAssemblyInfo(ctypes.Structure): + """ Matches ReturnAssemblyInfo on the python side """ + _fields_ = [('code', POINTER(CodeWrapped)), + ('jumps', POINTER(KvPair)), + ('code_len', ctypes.c_int), + ('jumps_len', ctypes.c_int)] + + +class Wave(ctypes.Structure): + _fields_ = [ + ('simd', ctypes.c_uint64), + ('wave_id', ctypes.c_uint64), + ('begin_time', ctypes.c_uint64), # Begin and end cycle + ('end_time', ctypes.c_uint64), + + # total VMEM/FLAT/LDS/SMEM instructions issued + # total issued memory instructions + ('num_mem_instrs', ctypes.c_uint64), + # total issued instructions (compute + memory) + ('num_issued_instrs', ctypes.c_uint64), + ('num_valu_instrs', ctypes.c_uint64), + ('num_valu_stalls', ctypes.c_uint64), + # VMEM Pipeline: instrs and stalls + ('num_vmem_instrs', ctypes.c_uint64), + ('num_vmem_stalls', ctypes.c_uint64), + # FLAT instrs and stalls + ('num_flat_instrs', ctypes.c_uint64), + ('num_flat_stalls', ctypes.c_uint64), + + # LDS instr and stalls + ('num_lds_instrs', ctypes.c_uint64), + ('num_lds_stalls', ctypes.c_uint64), + + # SCA instrs stalls + ('num_salu_instrs', ctypes.c_uint64), + ('num_smem_instrs', ctypes.c_uint64), + ('num_salu_stalls', ctypes.c_uint64), + ('num_smem_stalls', ctypes.c_uint64), + + # Branch + ('num_branch_instrs', ctypes.c_uint64), + ('num_branch_taken_instrs', ctypes.c_uint64), + ('num_branch_stalls', ctypes.c_uint64), + + ('timeline_string', ctypes.c_char_p), + ('instructions_string', ctypes.c_char_p)] + + +class ReturnInfo(ctypes.Structure): + _fields_ = [('num_waves', ctypes.c_uint64), + ('wavedata', POINTER(Wave)), + ('num_events', ctypes.c_uint64), + ('perfevents', POINTER(PerfEvent))] + +try: # For build dir + path_to_parser = os.path.abspath('/usr/lib/hsa-amd-aqlprofile/librocprofv2_att.so') + SO = CDLL(path_to_parser) +except: # For installed dir + path_to_parser = os.path.abspath('/usr/local/lib/hsa-amd-aqlprofile/librocprofv2_att.so') + SO = CDLL(path_to_parser) + +SO.AnalyseBinary.restype = ReturnInfo +SO.AnalyseBinary.argtypes = [ctypes.c_char_p, ctypes.c_int, ctypes.c_bool] +SO.wrapped_parse_binary.argtypes = [ctypes.c_char_p, ctypes.c_char_p] +SO.wrapped_parse_binary.restype = ReturnAssemblyInfo + +def parse_binary(filename, kernel=None): + if kernel is None or kernel == '': + kernel = ctypes.c_char_p(0) + print('Parsing all kernels') + else: + with open(glob.glob(kernel)[0], 'r') as file: + kernel = file.readlines() + print('Parsing kernel:', kernel[0].split(': ')[0]) + kernel = kernel[0].split(': ')[1].split('.kd')[0] + kernel = str(kernel).encode('utf-8') + info = SO.wrapped_parse_binary(str(filename).encode('utf-8'), kernel) + + code = [] + for k in range(info.code_len): + code_entry = info.code[k] + + # copy string memory from C++ + line = deepcopy(code_entry.line.decode("utf-8")) + loc = deepcopy(code_entry.loc.decode("utf-8")) + + # Transform empty entries back to python's None + to_line = int(code_entry.to_line) if (code_entry.to_line >= 0) else None + loc = loc if len(loc) > 0 else None + + code.append((line, int(code_entry.value), to_line, loc, + int(code_entry.index), int(code_entry.line_num))) + + jumps = {} + for k in range(info.jumps_len): + jumps[info.jumps[k].key] = info.jumps[k].value + + return code, jumps + + +def getWaves(filename, target_cu, verbose): + info = SO.AnalyseBinary(filename.encode('utf-8'), target_cu, verbose) + + waves = [info.wavedata[k] for k in range(info.num_waves)] + events = [deepcopy(info.perfevents[k]) for k in range(info.num_events)] + + for wave in waves: + wave.timeline = deepcopy(wave.timeline_string.decode("utf-8")) + wave.instructions = deepcopy(wave.instructions_string.decode("utf-8")) + + return waves, events + + +def persist(output_ui, trace_file, SIMD): + trace = Path(trace_file).name + simds, waves = [], [] + begin_time, end_time, timeline, instructions = [], [], [], [] + mem_ins, issued_ins, valu_ins, valu_stalls = [], [], [], [] + vmem_ins, vmem_stalls, flat_ins, flat_stalls = [], [], [], [] + lds_ins, lds_stalls, salu_ins, salu_stalls = [], [], [], [] + smem_ins, smem_stalls, br_ins, br_taken_ins, br_stalls = [], [], [], [], [] + + for wave in SIMD: + simds.append(wave.simd) + waves.append(wave.wave_id) + begin_time.append(wave.begin_time) + end_time.append(wave.end_time) + mem_ins.append(wave.num_mem_instrs) + issued_ins.append(wave.num_issued_instrs) + valu_ins.append(wave.num_valu_instrs) + valu_stalls.append(wave.num_valu_stalls) + vmem_ins.append(wave.num_vmem_instrs) + vmem_stalls.append(wave.num_vmem_stalls) + flat_ins.append(wave.num_flat_instrs) + flat_stalls.append(wave.num_flat_stalls) + lds_ins.append(wave.num_lds_instrs) + lds_stalls.append(wave.num_lds_stalls) + salu_ins.append(wave.num_salu_instrs) + salu_stalls.append(wave.num_salu_stalls) + smem_ins.append(wave.num_smem_instrs) + smem_stalls.append(wave.num_smem_stalls) + br_ins.append(wave.num_branch_instrs) + br_taken_ins.append(wave.num_branch_taken_instrs) + br_stalls.append(wave.num_branch_stalls) + timeline.append(wave.timeline) + instructions.append(wave.instructions) + + #df = pd.DataFrame({ + df = { + 'name': [trace for _ in range(len(begin_time))], + 'id': [i for i in range(len(begin_time))], + 'simd': simds, + 'wave_slot': waves, + 'begin_time': begin_time, + 'end_time': end_time, + 'mem_ins': mem_ins, + 'issued_ins': issued_ins, + 'valu_ins': valu_ins, + 'valu_stalls': valu_stalls, + 'vmem_ins': vmem_ins, + 'vmem_stalls': vmem_stalls, + 'flat_ins': flat_ins, + 'flat_stalls': flat_stalls, + 'lds_ins': lds_ins, + 'lds_stalls': lds_stalls, + 'salu_ins': salu_ins, + 'salu_stalls': salu_stalls, + 'smem_ins': smem_ins, + 'smem_stalls': smem_stalls, + 'br_ins': br_ins, + 'br_taken_ins': br_taken_ins, + 'br_stalls': br_stalls, + 'timeline': timeline, + 'instructions': instructions, + }#) + #[print(d) for c, d in df.iterrows()]; quit() + return df + + +def mem_max(array): + mem_dict = {} + for SE in array: + for wave in SE: + for inst in wave: + try: + mem_dict[inst[0]][0] = max(mem_dict[inst[0]][0], inst[1]) + except: + mem_dict[inst[0]] = inst[1:] + assert(mem_dict[inst[0]][1] == inst[2]) + + return mem_dict + +def lgk(count): + return 'lgkmcnt({0})'.format(count) +def vmc(count): + return 'vmcnt({0})'.format(count) +def both_cnt(count): + return lgk(count)+' '+vmc(count) + +def insert_waitcnt(flight_count, assembly_code): + flight_count = mem_max(flight_count) + for key in sorted(flight_count): + line_n = key + issue_amount, waitcnt_amount, = flight_count[key] + if 'vmcnt' in assembly_code[line_n] and 'lgkmcnt' in assembly_code[line_n]: + counter_type = both_cnt + elif 'vmcnt' in assembly_code[line_n]: + counter_type = vmc + elif 'lgkmcnt' in assembly_code[line_n]: + counter_type = lgk + else: + print('Error: Line mismatch') + exit(-1) + + for count in range(waitcnt_amount+1, issue_amount): + print('Inserted line: '+str(line_n)) + as_index = line_n - count/(issue_amount+1) + assembly_code[as_index] = \ + '\ts_waitcnt {0}\t\t; Timing analysis.'.format(counter_type(count)) + as_index += 0.5/(issue_amount+1) + assembly_code[as_index] = '\ts_nop 0\t\t\t\t\t\t; Counters: '+str(issue_amount) + + return assembly_code + + +def Copy_Files(output_ui): + curpath = os.path.dirname(os.path.abspath(__file__)) + outpath = output_ui+'/ui/' + + os.makedirs(outpath, exist_ok=True) + os.system('cp '+curpath+'/ui/* '+outpath) + + +def get_delta_time(events): + for begin in range(len(events)): + tg_cu = events[begin].cu + for e in range(begin+1,len(events)): + if events[e].cu == tg_cu: + return events[e].time-events[begin].time + return 1 + + +def num_cus(EVENTS): + cus = 0 + for events in EVENTS: + for e in events: + cus = max(cus, e.cu) + return cus+1 + + +def draw_wave_metrics(selections, normalize): + global PIC_SAVE_FOLDER + global EVENTS + global EVENT_NAMES + + #event_names = ['Busy CUs', 'Occupancy', 'Eligible waves', 'Waves waiting'] + with open(PIC_SAVE_FOLDER+'counters.json', 'w') as f: + f.write(json.dumps({"counters": EVENT_NAMES})) + + plt.figure(figsize=(15,3)) + + delta_time = int(0.5+np.mean([get_delta_time(events) for events in EVENTS])) + maxtime = np.max([np.max([e.time for e in events]) for events in EVENTS])+1 + event_timeline = np.zeros((16, maxtime), dtype=np.int32) + print('Delta:', delta_time) + print('Max_cycles:', maxtime) + + kernsize = 2*(delta_time//8)+1 + trim = max(maxtime//5000,1) + cycles = 4*np.arange(maxtime)[::trim] + + kernel = np.asarray([np.exp(-abs(k/kernsize)**2) for k in range(-kernsize*3,kernsize*3+1)]) + kernel /= np.sum(kernel)*len(EVENTS)*delta_time#*5.12 # SEslots/100% + + for events in EVENTS: + for e in range(len(events)-1): + bk = events[e].bank*4 + start = events[e].time + end = start+delta_time + event_timeline[bk:bk+4, start:end] += np.asarray(events[e].toTuple()[1:5])[:, None] + start = events[-1].time + event_timeline[bk:bk+4, start:start+delta_time] += \ + np.asarray(events[-1].toTuple()[1:5])[:, None] + + event_timeline = [np.convolve(e, kernel)[3*kernsize:-3*kernsize] for e in event_timeline] + + if normalize: + event_timeline = [100*e/max(e.max(), 1E-5) for e in event_timeline] + #event_timeline[0] = np.clip(event_timeline[0]*8.5, 0, 100) #CC rate + + colors = ['blue', 'green', 'gray', 'red', 'orange', 'cyan', 'black', 'darkviolet', + 'yellow', 'darkred', 'pink', 'lime', 'gold', 'tan', 'aqua', 'olive'] + [plt.plot(cycles, e[::trim], '-', label=n, color=c) + for e, n, c, sel in zip(event_timeline, EVENT_NAMES, colors, selections) if sel] + + plt.legend() + if normalize: + plt.ylabel('As % of maximum') + else: + plt.ylabel('Value') + plt.subplots_adjust(left=0.05, right=1, top=1, bottom=0.07) + plt.savefig(PIC_SAVE_FOLDER+'timeline.png', dpi=150) + #plt.show() + + +def draw_wave_states(selections, normalize): + global TIMELINES + global PIC_SAVE_FOLDER + plot_indices = [1, 2, 3, 4] + STATES = [['Empty', 'Idle', 'Exec', 'Wait', 'Stall'][k] for k in plot_indices] + colors = [['gray', 'orange', 'green', 'red', 'blue'][k] for k in plot_indices] + + plt.figure(figsize=(15,3)) + + maxtime = max([np.max((TIMELINES[k]!=0)*np.arange(0,TIMELINES[k].size)) for k in plot_indices]) + timelines = [deepcopy(TIMELINES[k][:maxtime]) for k in plot_indices] + timelines = [np.pad(t, [0, maxtime-t.size]) for t in timelines] + + if normalize: + timelines = np.array(timelines) / np.maximum(np.sum(timelines,0)*1E-2,1E-7) + + kernsize = maxtime//120+3 + trim = max(maxtime//5000,1) + cycles = np.arange(timelines[0].size)[::trim] + + kernel = np.asarray([np.exp(-abs(10*k/kernsize)) for k in range(-kernsize//2,kernsize//2+1)]) + kernel /= np.sum(kernel) + + timelines = [np.convolve(time, kernel)[kernsize//2:-kernsize//2][::trim] for time in timelines] + + with open(PIC_SAVE_FOLDER+'counters.json', 'w') as f: + f.write(json.dumps({"counters": STATES})) + + [plt.plot(cycles, t, label='State '+s, linewidth=1.1, color=c) + for t, s, c, sel in zip(timelines, STATES, colors, selections) if sel] + + plt.legend() + if normalize: + plt.ylabel('Waves state %') + else: + plt.ylabel('Waves state total') + plt.ylim(-1) + plt.xlim(-maxtime//200, maxtime+maxtime//200) + plt.subplots_adjust(left=0.05, right=1, top=1, bottom=0.07) + plt.savefig(PIC_SAVE_FOLDER+'timeline.png', dpi=150) + + +def GeneratePIC(selections=[True for k in range(4)], normalize=True): + if len(EVENTS) > 0 and np.sum([len(e) for e in EVENTS]) > 32: + draw_wave_metrics(selections, normalize) + else: + draw_wave_states(selections, normalize) + + +if __name__ == "__main__": + parser = argparse.ArgumentParser() + parser.add_argument("assembly_code", help="Path of the assembly code") + parser.add_argument("--trace_file", help="Filter for trace files", default=None, type=str) + parser.add_argument("-o", "--output_ui", help="Output Folder", default='/dev/shm/attplugin/') + parser.add_argument("-k", "--att_kernel", help="Kernel file", type=str, default='*_kernel.txt') + parser.add_argument("-w", "--wave_id", help="wave id") + parser.add_argument("-p", "--ports", help="Server and websocket ports, default: 8000,18000") + parser.add_argument("--target_cu", help="Collected target CU id{0-15}", type=int, default=None) + parser.add_argument("-v", "--verbose", action="store_true") + parser.add_argument("-g", "--genasm", + help="Generate post-processed asm file at this path", type=str, default="") + args = parser.parse_args() + + global EVENT_NAMES + with open(os.getenv("COUNTERS_PATH"), 'r') as f: + lines = [l.split('//')[0] for l in f.readlines()] + + EVENT_NAMES = [] + clean = lambda x: x.split('=')[1].split(' ')[0].split('\n')[0] + + for line in lines: + if 'PERFCOUNTER_ID=' in line: + EVENT_NAMES += ['id: '+clean(line)] + elif args.target_cu is None and 'att: TARGET_CU' in line: + args.target_cu = int(clean(line)) + print('Target CU set to:', args.target_cu) + for line in lines: + if 'PERFCOUNTER=' in line: + EVENT_NAMES += [clean(line).split('SQ_')[1].lower()] + + if args.target_cu is None: + args.target_cu = 1 + + # Assembly parsing + path = Path(args.assembly_code) + if not path.is_file(): + print("Invalid assembly_code('{0}')!".format(args.assembly_code)) + sys.exit(1) + + att_kernel = glob.glob(args.att_kernel) + + if len(att_kernel) == 0: + print('Could not find att output kernel:', args.att_kernel) + exit(1) + elif len(att_kernel) > 1: + print('Found multiple kernel matching given filters:') + for n, k in enumerate(att_kernel): + print('\t', n, '->', k) + + bValid = False + while bValid == False: + try: + args.att_kernel = att_kernel[int(input("Please select number: "))] + bValid = True + except KeyboardInterrupt: + exit(0) + except: + print('Invalid option.') + else: + args.att_kernel = att_kernel[0] + + print('Att kernel:', args.att_kernel) + code, jumps = parse_binary(args.assembly_code, args.att_kernel) + + # Trace Parsing + if args.trace_file is None: + filenames = glob.glob(args.att_kernel.split('_kernel.txt')[0]+'*.att') + assert(len(filenames) > 0) + else: + filenames = glob.glob(args.trace_file) + + print('Trace filenames:', filenames) + + Copy_Files(args.output_ui) + DBFILES = [] + global TIMELINES + global EVENTS + TIMELINES = [np.zeros(int(1E4),dtype=np.int32) for k in range(5)] + EVENTS = [] + for name in filenames: + SIMD, perfevents = getWaves(name, args.target_cu, args.verbose) + EVENTS.append(perfevents) + DBFILES.append( persist(args.output_ui, name, SIMD) ) + for wave in SIMD: + time_acc = 0 + tuples1 = wave.timeline.split('(') + tuples2 = [t.split(')')[0].split(',') for t in tuples1 if t != ''] + tuples3 = [(int(t[0]),int(t[1])) for t in tuples2] + + for state in tuples3: + if state[1] > 1E7: + print('Warning: Time limit reached for ',state[0], state[1]) + break + + if time_acc+state[1] > TIMELINES[state[0]].size: + TIMELINES[state[0]] = np.hstack([ + TIMELINES[state[0]], + np.zeros_like(TIMELINES[state[0]]) + ]) + TIMELINES[state[0]][time_acc:time_acc+state[1]] += 1 + time_acc += state[1] + + if args.genasm and len(args.genasm) > 0: + flight_count = view_trace(args, 0, code, jumps, DBFILES, filenames, True, None) + + with open(args.assembly_code, 'r') as file: + lines = file.readlines() + assembly_code = {l+1.0: lines[l][:-1] for l in range(len(lines))} + assembly_code = insert_waitcnt(flight_count, assembly_code) + + with open(args.genasm, 'w') as file: + keys = sorted(assembly_code.keys()) + for k in keys: + file.write(assembly_code[k]+'\n') + else: + global PIC_SAVE_FOLDER + PIC_SAVE_FOLDER = args.output_ui+"/ui/" + view_trace(args, 0, code, jumps, DBFILES, filenames, False, GeneratePIC) diff --git a/plugin/att/trace_view.py b/plugin/att/trace_view.py new file mode 100755 index 0000000000..6b651971a2 --- /dev/null +++ b/plugin/att/trace_view.py @@ -0,0 +1,465 @@ +#!/usr/bin/env python3 +import sys +if sys.version_info[0] < 3: + raise Exception("Must be using Python 3") + +import os +import sys +import time +import socket +from pathlib import Path +from struct import * +from collections import defaultdict +import json +import time +import webbrowser +import http.server +import socketserver +import socket +import asyncio +import websockets +from multiprocessing import * +from copy import deepcopy + +PORT, WebSocketPort = 8000, 18000 +SP = '\u00A0' + +RS_TRACE_DEBUG = "RS_TRACE_DEBUG" in os.environ +if RS_TRACE_DEBUG: + LOG = open('./att_viewer.log', 'w') + + +def get_ip(): + s = socket.socket(socket.AF_INET, socket.SOCK_DGRAM) + s.settimeout(0) + try: + hostname = socket.gethostname() + IPAddr = socket.gethostbyname(hostname) + s.connect(({IPAddr}, 1)) + except Exception: + IPAddr = '127.0.0.1' + finally: + return IPAddr + + +IPAddr = get_ip() + + +def debug_log(msg, last=False): + if RS_TRACE_DEBUG: + LOG.write(msg) + + if last: + LOG.close() + + +def try_match_swapped(insts, code, i, line): + return insts[i+1][1] == code[line][1] and insts[i][1] == code[line+1][1] + + +def stitch(insts, code, jumps): + result, i, line, loopCount, N = [], 0, 0, defaultdict(int), len(insts) + + SMEM_INST = [] + VMEM_INST = [] + FLAT_INST = [] + NUM_SMEM = 0 + NUM_VMEM = 0 + NUM_FLAT = 0 + + mem_unroll = [] + flight_count = [] + + ordering = 0 # 0 for in order, 1 for ongoing flats and 2 for SMEM + while i < N: + inst = insts[i] + if line >= len(code): + break + + as_line = code[line] + if inst[1] == as_line[1]: + if line in jumps: + loopCount[line-1] += 1 # label is the previous line + matched, next = True, line + 1 + + num_inflight = NUM_FLAT + NUM_SMEM + NUM_VMEM + + if inst[1] == 1 or inst[1] == 5: # SMEM, LDS + ordering = 2 if inst[1] == 1 else ordering + SMEM_INST.append([line, num_inflight]) + NUM_SMEM += 1 + elif inst[1] == 3 or (inst[1] == 4 and 'global_' in as_line[0]): # VMEM R/W + VMEM_INST.append([line, num_inflight]) + NUM_VMEM += 1 + elif inst[1] == 4: # FLAT + ordering = max(ordering, 1) + FLAT_INST.append([line, num_inflight]) + NUM_FLAT += 1 + elif inst[1] == 9 and 'waitcnt' in as_line[0]: + + if 'lgkmcnt' in as_line[0]: + wait_N = int(as_line[0].split('lgkmcnt(')[1].split(')')[0]) + flight_count.append([as_line[-1], num_inflight, wait_N]) + if wait_N == 0: + ordering = 0 + if ordering == 0: + offset = len(SMEM_INST)-wait_N + mem_unroll.append( [line, SMEM_INST[:offset]+FLAT_INST] ) + SMEM_INST = SMEM_INST[offset:] + FLAT_INST = [] + NUM_FLAT = 0 + NUM_SMEM = 0 + else: + NUM_SMEM = min(max(wait_N-NUM_FLAT, 0), NUM_SMEM) + NUM_FLAT = min(max(wait_N-NUM_SMEM, 0), NUM_FLAT) + + if 'vmcnt' in as_line[0]: + wait_N = int(as_line[0].split('vmcnt(')[1].split(')')[0]) + flight_count.append([as_line[-1], num_inflight, wait_N]) + if wait_N == 0 and ordering != 2: + ordering = 0 + if ordering == 0: + offset = len(VMEM_INST)-wait_N + mem_unroll.append( [line, VMEM_INST[:offset]+FLAT_INST] ) + VMEM_INST = VMEM_INST[offset:] + FLAT_INST = [] + NUM_FLAT = 0 + NUM_VMEM = 0 + else: + NUM_VMEM = min(max(wait_N-NUM_FLAT, 0), NUM_VMEM) + NUM_FLAT = min(max(wait_N-NUM_VMEM, 0), NUM_FLAT) + + + elif inst[1] == 7 and as_line[1] == 10: # jump + matched, next = True, as_line[2] + elif inst[1] == 8 and as_line[1] == 10: # next + matched, next = True, line + 1 + else: + # instructions with almost same timestamp swapped + # if i+1 < N and line+1 < len(code) and inst[0] == insts[i+1][0]: + matched = False + next = line + 1 + if i+1 < N and line+1 < len(code): + if try_match_swapped(insts, code, i, line): + #print('Swap:', code[line]) + #print('For:', code[line+1]) + #result += (*(insts[i+1]), line), + #result += (*inst, line+1), + #i, next = i+2, line+2 + temp = insts[i] + insts[i] = insts[i+1] + insts[i+1] = temp + next = line + #else: + # print('Could not parse tokens:', insts[i], as_line) + + if matched: + new_res = inst + (line,) + result.append(new_res) + i += 1 + line = next + + N = max(N, 1) + if len(result) != N: + print('Warning - Stitching rate: {'+str(len(result) * 100 / N)+'% matched') + + return result, loopCount, mem_unroll, flight_count + + +def extract_tuple(content, num): + vals = content.split(',') + assert (len(vals) == num) + last_val = vals[-1][:-1] if vals[-1].endswith(')') else vals[-1] + vals = [vals[0][1:]] + vals[1:-1] + [last_val] + return tuple(int(val) for val in vals) + + +def get_top_n(stitched): + TOP_N = 10 + by_line_num = defaultdict(lambda: [0, 0, 0]) + for (_, _, s2i, run_time, line_num) in stitched: + entry = by_line_num[line_num] + entry[0] += 1 + entry[1] += s2i + entry[2] += run_time + top_n = sorted( + [(line_num, v[0], v[1], v[2]) + for (line_num, v) in by_line_num.items()], + key=lambda x: x[2] + x[3], + reverse=True) + return top_n[:TOP_N] + + +def rjust_html(s, n): + s = str(s) + return SP * (n-len(s)) + s if len(s) < n else s + + +def rjust_html_format(msg, n1, inst, n2, n3, stall): + return str(rjust_html(msg,n1)) + str(rjust_html(inst,n2)) + str(SP*n3) + str(stall) + + +def wave_info(df, id): + issued_ins, mem_ins = df['issued_ins'][id], df['mem_ins'][id] + valu_ins, valu_stalls = df['valu_ins'][id], df['valu_stalls'][id] + salu_ins, salu_stalls = df['salu_ins'][id], df['salu_stalls'][id] + vmem_ins, vmem_stalls = df['vmem_ins'][id], df['vmem_stalls'][id] + smem_ins, smem_stalls = df['smem_ins'][id], df['smem_stalls'][id] + flat_ins, flat_stalls = df['flat_ins'][id], df['flat_stalls'][id] + lds_ins, lds_stalls = df['lds_ins'][id], df['lds_stalls'][id] + br_ins, br_stalls = df['br_ins'][id], df['br_stalls'][id] + + return 'Issued:' + str(rjust_html(issued_ins,8)) + str(SP*2) + 'Mem:' + str(mem_ins) \ + + "-" * 26 + rjust_html_format("VALU:",6,valu_ins,8,4,valu_stalls) \ + + rjust_html_format("SALU:",6,salu_ins,8,4,salu_stalls) \ + + rjust_html_format("VMEM:",6,vmem_ins,8,4,vmem_stalls) \ + + rjust_html_format("SMEM:",6,smem_ins,8,4,smem_stalls) \ + + rjust_html_format("FLAT:",6,flat_ins,8,4,flat_stalls) \ + + rjust_html_format("LDS:",6,lds_ins,8,4,lds_stalls) \ + + rjust_html_format("BR:",6,br_ins,8,4,br_stalls) + + +def extract_waves(waves): + result, slot2seq = [], {} + for id in waves['id']: + row = {key: waves[key][id] for key in waves.keys()} + + insts, timeline = [], [] + for x in row['instructions'].split('),'): + if len(x) > 0: + insts.append(extract_tuple(x, 4)) + for x in row['timeline'].split('),'): + if len(x) > 0: + timeline.append(extract_tuple(x, 2)) + + # aggregate per wave slot + if (row['simd'], row['wave_slot']) in slot2seq: + slot = result[slot2seq[(row['simd'], row['wave_slot'])]] + last_end_time = slot[2][-1][-1] + slot[2] += (row['id'], row['begin_time'], row['end_time']), + slot[3] += insts + # filler between waves + slot[4] += (0, row['begin_time'] - last_end_time), + slot[4] += timeline + else: + slot2seq[row['simd'], row['wave_slot']] = len(result) + result.append([row['simd'], row['wave_slot'], + [(row['id'], row['begin_time'], row['end_time'])], + insts, + timeline]) + + return result + + +def extract_data(df, output_ui, se_number, code, jumps): + if len(df['id']) == 0 or len(df['instructions']) == 0 or len(df['timeline']) == 0: + return None + + cu_waves = extract_waves(df) + all_filenames = [] + flight_count = [] + + for wave_id in df['id']: + insts, timeline = [], [] + if len(df['instructions'][wave_id]) == 0 or len(df['timeline'][wave_id]) == 0: + continue + + for x in df['instructions'][wave_id].split('),'): + insts.append(extract_tuple(x, 4)) + for x in df['timeline'][wave_id].split('),'): + timeline.append(extract_tuple(x, 2)) + + stitched, loopCount, mem_unroll, count = stitch(insts, code, jumps) + flight_count.append(count) + + wave_entry = { + "id": int(df['id'][wave_id]), + "simd": int(df['simd'][wave_id]), + "slot": int(df['wave_slot'][wave_id]), + "begin": int(df['begin_time'][wave_id]), + "end": int(df['end_time'][wave_id]), + "info": wave_info(df, wave_id), + "instructions": stitched, + "timeline": timeline, + "code": code, + "waitcnt": mem_unroll + } + data_obj = { + "name": 'SE'.format(se_number), + "kernel": code[0][0], + "duration": sum(dur for (_, dur) in timeline), + "wave": wave_entry, + "simd_waves": [], + "cu_waves": cu_waves, + "loop_count": loopCount, + "top_n": get_top_n(stitched), + "websocket_port": WebSocketPort, + "generation_time": time.ctime() + } + if len(data_obj["cu_waves"]) == 0: + continue + + OUT = output_ui+'/ui/se'+str(se_number)+'_sm'+str(df['simd'][wave_id])+\ + '_wv'+str(df['wave_slot'][wave_id])+'.json' + + with open(OUT, 'w') as f: + f.write(json.dumps(data_obj)) + all_filenames.append(OUT.split('/')[-1]) + + return flight_count, all_filenames + +def open_browser(): + time.sleep(0.1) + webbrowser.open_new_tab('http://{0}:{1}'.format(IPAddr, PORT)) + + +class NoCacheHTTPRequestHandler(http.server.SimpleHTTPRequestHandler): + def end_headers(self): + self.send_my_headers() + http.server.SimpleHTTPRequestHandler.end_headers(self) + + def send_my_headers(self): + self.send_header("Cache-Control", "no-cache, no-store, must-revalidate") + self.send_header("Pragma", "no-cache") + self.send_header("Expires", "0") + + def do_GET(self): + global PICTURE_CALLBACK + if 'timeline.png?' in self.path: + selections = [int(s)!=0 for s in self.path.split('timeline.png?')[1]] + print('Sel:', selections) + PICTURE_CALLBACK(selections[1:], selections[0]) + http.server.SimpleHTTPRequestHandler.do_GET(self) + +class RocTCPServer(socketserver.TCPServer): + def server_bind(self): + self.socket.setsockopt(socket.SOL_SOCKET, socket.SO_REUSEADDR, 1) + self.socket.bind(self.server_address) + + +def run_server(): + global RS_HOME + Handler = NoCacheHTTPRequestHandler + + os.chdir(RS_HOME+'/ui') + try: + with RocTCPServer((IPAddr, PORT), Handler) as httpd: + httpd.serve_forever() + except KeyboardInterrupt: + pass + + +def fix_space(line): + line = line.replace(' ', SP) + line = line.replace('\t', SP*4) + return line + + +def WebSocketserver(websocket, path): + data = websocket.recv() + print(354, data) + cpp, ln, _ = data.split(':') + ln = int(ln) + HL, EMP = 'highlight', '' + content = None + print("loading...") + try: + f = open(cpp, 'r', errors='replace') + content = ''.join('
  • '+str(i).ljust(5)+fix_space(l)+'
    • ' + for i, l in enumerate(f.readlines(), 1)) + except FileNotFoundError: + content = cpp + ' not found!' + websocket.send(content) + + +def run_websocket(): + start_server = websockets.serve(WebSocketserver, IPAddr, WebSocketPort) + try: + asyncio.get_event_loop().run_until_complete(start_server) + asyncio.get_event_loop().run_forever() + except KeyboardInterrupt: + pass + + +def assign_ports(ports): + ps = [int(port) for port in ports.split(',')] + if ps[0] <= 5000 or ps[1] <= 5000: + print('Need to have port values > 5000') + sys.exit(1) + elif ps[0] == ps[1]: + print('Can not use the same port for both web server and websocket server: '+ps[0]) + sys.exit(1) + global IPAddr, PORT, WebSocketPort + PORT, WebSocketPort = ps[0], ps[1] + + +def view_trace(args, wait, code, jumps, dbnames, att_filenames, bReturnLoc, pic_callback): + global PICTURE_CALLBACK + PICTURE_CALLBACK = pic_callback + pic_thread = Process(target=pic_callback) + pic_thread.start() + + assert(len(dbnames) > 0) + global RS_HOME + output_ui = args.output_ui + RS_HOME = output_ui + + att_filenames = [Path(f).name for f in att_filenames] + se_numbers = [int(a.split('_se')[1].split('.att')[0]) for a in att_filenames] + flight_count = [] + simd_wave_filenames = {} + + for se_number, dbname in zip(se_numbers, dbnames): + if len(dbname['id']) == 0: + continue + + count, wv_filenames = extract_data(dbname, output_ui, se_number, code, jumps) + + if count is not None: + flight_count.append(count) + simd_wave_filenames[se_number] = wv_filenames + + if bReturnLoc: + return flight_count + + for key in simd_wave_filenames.keys(): + wv_array = [[ + int(s.split('_sm')[1].split('_wv')[0]), + int(s.split('_wv')[1][0]), + s + ] for s in simd_wave_filenames[key]] + + wv_dict = {} + for wv in wv_array: + try: + wv_dict[wv[0]][wv[1]] = wv[2] + except: + try: + wv_dict[wv[0]] = {wv[1]: wv[2]} + except: + exit(-1) + + simd_wave_filenames[key] = wv_dict + + with open(output_ui+'/ui/filenames.json', 'w') as f: + f.write(json.dumps({"filenames": simd_wave_filenames})) + + if args.ports: + assign_ports(args.ports) + print('serving at ports: {0},{1}'.format(PORT, WebSocketPort)) + + if wait == 0: + try: + PROCS = [Process(target=run_server), + Process(target=open_browser), + Process(target=run_websocket)] + if pic_thread is not None: + pic_thread.join() + + for p in PROCS: + p.start() + for p in PROCS: + p.join() + except KeyboardInterrupt: + print("Exitting.") diff --git a/plugin/att/ui/index.html b/plugin/att/ui/index.html new file mode 100644 index 0000000000..b1dbbc7703 --- /dev/null +++ b/plugin/att/ui/index.html @@ -0,0 +1,1047 @@ + + + + + + MI Trace Viewer + + + +
    +
    + +
    + +
    + +
    + +
    + +
    +
    + +
    +
    +
    +
    +
    +
    +
    +
    +
    + +
    +
    +
    +
    +
    +
    +
      +
      + + + + + + diff --git a/plugin/att/ui/logo.svg b/plugin/att/ui/logo.svg new file mode 100644 index 0000000000..340c490f1b --- /dev/null +++ b/plugin/att/ui/logo.svg @@ -0,0 +1,132 @@ + + + + + + + + + + + MI Trace View + + diff --git a/plugin/att/ui/styles.css b/plugin/att/ui/styles.css new file mode 100644 index 0000000000..1949a608d9 --- /dev/null +++ b/plugin/att/ui/styles.css @@ -0,0 +1,106 @@ +h2,h3,h4 { + text-align: center; + font-family: Calibri, Candara, Optima, Arial, 'Trebuchet MS', sans-serif; +} + +ul, ol { + list-style-type: none; + padding: 0; + margin: 20px 5px; + /*font-family: Calibri, Candara, Segoe, "Segoe UI", Optima, Arial, sans-serif; font-size: 15px; font-style: normal; font-variant: normal; */ + font-family: 'Courier New', monospace; font-size: 15px; font-style: normal; +} + +a { + color: royalblue; + cursor: pointer; +} + +div#flexbox { + display: flex; + flex-wrap: wrap; +} + +div#flexbox > div { + flex: 50%; +} + +div#minimap { + position: absolute; + width: 350px; + top: 340px; + display: flex; + justify-content: right; + right: 5px; +} + +div#wave, div#cu_wave { + overflow:scroll; + overflow-y:hidden; +} + +div#ma_code { + height: calc(100vh - 180px); + overflow: auto; + overflow-x: hidden; + display: block; + margin: 20px; +} + +nav { + padding-right: 8px; + position: absolute; + top: 1rem; + right: 1rem; + background: #FFFDE3; + z-index: 5; + margin-left: 10px; + width:330px; +} + +.highlight { + background-color: lightgray; + font-weight: bold; +} + +.clickable { + cursor: pointer; + font-weight: 180; +} + +ul li { + border-bottom: 1px dotted black; +} + +.tooltip { + display: none; + position: relative; + left: 10px; + z-index: 100; + background: #333; + border: 1px solid #c0c0c0; + opacity: 0.8; + color: white; +} +li:hover .tooltip { + display: inline-block; /*block;*/ +} +.loop { + margin-left: 30px; + background: lightseagreen; +} + +.btn { + border: 1px solid black; + background-color: #D7D7D7; + color: black; + padding: 3px 4px; + font-size: 14px; + cursor: pointer; + border-style: ridge; + border-radius: 6px; +} + +.btn:hover { + color: blue; +} \ No newline at end of file diff --git a/plugin/file/file.cpp b/plugin/file/file.cpp index 7e6da7dad3..f21c430927 100644 --- a/plugin/file/file.cpp +++ b/plugin/file/file.cpp @@ -403,6 +403,9 @@ class file_plugin_t { FlushTracerRecord(*tracer_record, session_id, buffer_id); break; } + case ROCPROFILER_ATT_TRACER_RECORD: { + break; + } case ROCPROFILER_PC_SAMPLING_RECORD: { const rocprofiler_record_pc_sample_t *pc_sampling_record = reinterpret_cast(begin); diff --git a/rocprofv2 b/rocprofv2 index 19e82a7f0a..078652cb88 100755 --- a/rocprofv2 +++ b/rocprofv2 @@ -11,7 +11,7 @@ elif [[ $ROCPROFV2_DIR == *"/rocprofiler"* ]]; then fi usage() { - echo -e "ROCProfiler Run Script Usage:" + echo -e "ROCProfilerV2 Run Script Usage:" echo -e "-h | --help For showing this message" echo -e "--list-counters For showing all available counters for the current GPUs" if [ $RUN_FROM_BUILD == 1 ]; then @@ -21,8 +21,8 @@ usage() { echo -e "-ct | --clean-build-test For Running the tests after a clean build" echo -e "-mt | --mem-test For Running the Memory Leak tests. This run requires building using -acb | --asan-clean-build option" echo -e "-acb | --asan-clean-build For compiling with ASAN library attached" - echo -e "--install For installing rocprofiler without clean build in the default installation folder (review build.sh to know more about the default paths)" - echo -e "--clean-install For installing rocprofiler with new clean build in the default installation folder (review build.sh to know more about the default paths)" + echo -e "--install For installing ROCProfilerV2 without clean build in the default installation folder (review build.sh to know more about the default paths)" + echo -e "--clean-install For installing ROCProfilerV2 with new clean build in the default installation folder (review build.sh to know more about the default paths)" fi echo -e "--hip-api For Collecting HIP API Traces" echo -e "--hip-activity For Collecting HSA API Activities Traces" @@ -31,11 +31,18 @@ usage() { echo -e "--roctx-trace For Collecting ROCTx Traces" echo -e "--kernel-trace For Collecting Kernel dispatch Traces" echo -e "--sys-trace For Collecting HIP and HSA APIs and their Activities Traces along ROCTX and Kernel Dispatch traces" - echo -e "--plugin PLUGIN_NAME For enabling a plugin (file/perfetto)" + echo -e "--plugin PLUGIN_NAME For enabling a plugin (file/perfetto/att)" echo -e "-i | --input For adding counters file path (every line in the text file represents a counter)" echo -e "-o | --output-file For the output file name" echo -e "-d | --output-directory For adding output path where the output files will be saved" echo -e "-fi | --flush-interval For adding a flush interval in milliseconds, every \"flush interval\" the buffers will be flushed" + echo -e "\n###ATT Plugin options: ###" + if [ $RUN_FROM_BUILD == 1 ]; then + ATT_PATH=$ROCPROFV2_DIR/build/plugin/att/att/att.py + else + ATT_PATH=$ROCPROFV2_DIR/../libexec/rocprofiler/att/att.py + fi + eval "python3 $ATT_PATH --help" exit 1 } @@ -199,6 +206,35 @@ while [ 1 ] ; do usage exit 1 fi + if [ "$2" = "att" ] ; then + #if [ $RUN_FROM_BUILD == 1 ]; then + # export ROCMTOOLS_ATT_PARSER=build/libatt_parser.so + #else + # export ROCMTOOLS_ATT_PARSER=$ROCPROFV2_DIR/../lib/rocprofiler/libatt_parser.so + #fi + ATT_ARGV=$3 + shift + + ATT_OPTIONS="Not done" + while [ "$ATT_OPTIONS" = "Not done" ]; do + if [[ "$3" = "--trace_file" ]]; then + ATT_ARGV="$ATT_ARGV $3 \"$4\"" + shift + shift + elif [[ "$3" = "--genasm" || "$3" = "--target_cu" || "$3" = "-o" || "$3" == "-k" || "$3" == "--att_kernel" ]]; then + ATT_ARGV="$ATT_ARGV $3 $4" + shift + shift + else + ATT_OPTIONS="Done" + fi + done + if [ $RUN_FROM_BUILD == 1 ]; then + ATT_PATH=$ROCPROFV2_DIR/build/plugin/att/att/att.py + else + ATT_PATH=$ROCPROFV2_DIR/../libexec/rocprofiler/att/att.py + fi + fi shift shift elif [[ "$1" = "-"* || "$1" = "--"* ]] ; then @@ -214,7 +250,10 @@ PMC_LINES=() if [ -n "$COUNTERS_PATH" ]; then input=$COUNTERS_PATH while IFS= read -r line || [[ -n "$line" ]]; do + # if in att mode, only add the first line + if [[ ! -n "$PMC_LINES" ]] || [[ ! -n "$ATT_ARGV" ]]; then PMC_LINES+=( "$line" ) + fi done < $input fi @@ -236,7 +275,7 @@ if [ -n "$PMC_LINES" ]; then LD_PRELOAD=$LD_PRELOAD:$ROCM_DIR/lib/librocprofiler_tool.so $* fi done -else +elif [ ! -n "$ATT_ARGV" ]; then if [ $RUN_FROM_BUILD == 1 ]; then LD_PRELOAD=$LD_PRELOAD:$ROCM_DIR/build/librocprofiler_tool.so $* else @@ -244,4 +283,10 @@ else fi fi - +if [ -n "$ATT_PATH" ]; then + if [ -n "$ATT_ARGV" ]; then + eval "python3 $ATT_PATH $ATT_ARGV" + elif [ ! -n "$PMC_LINES" ]; then + echo "ATT File is required!" + fi +fi diff --git a/samples/att.txt b/samples/att.txt new file mode 100644 index 0000000000..29cf693cb1 --- /dev/null +++ b/samples/att.txt @@ -0,0 +1 @@ +att: TARGET_CU=1 diff --git a/src/api/CMakeLists.txt b/src/api/CMakeLists.txt index b57e566fe7..75e27785c5 100644 --- a/src/api/CMakeLists.txt +++ b/src/api/CMakeLists.txt @@ -150,6 +150,7 @@ file(GLOB ROCPROFILER_SRC_FILES ${CMAKE_CURRENT_SOURCE_DIR}/*.cpp) file(GLOB ROCPROFILER_PROFILER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/profiler/profiler.cpp) file(GLOB ROCPROFILER_TRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/*.cpp) file(GLOB ROCPROFILER_ROCTRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/src/*.cpp) +file(GLOB ROCMTOOLS_ATT_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/att/att.cpp) file(GLOB ROCMTOOL_CLASS_SRC_FILES ${CMAKE_CURRENT_SOURCE_DIR}/rocmtool.cpp) file(GLOB ROCPROFILER_SPM_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/spm/spm.cpp) @@ -189,6 +190,7 @@ add_library(${ROCPROFILER_TARGET} SHARED ${ROCPROFILER_SRC_FILES} ${ROCMTOOL_CLASS_SRC_FILES} ${ROCPROFILER_PROFILER_SRC_FILES} + ${ROCMTOOLS_ATT_SRC_FILES} ${CORE_HARDWARE_SRC_FILES} ${CORE_HSA_SRC_FILES} ${ROCPROFILER_SPM_SRC_FILES} diff --git a/src/core/hsa/packets/packets_generator.cpp b/src/core/hsa/packets/packets_generator.cpp index 290dc99bfb..58ba619879 100644 --- a/src/core/hsa/packets/packets_generator.cpp +++ b/src/core/hsa/packets/packets_generator.cpp @@ -486,4 +486,146 @@ hsa_ven_amd_aqlprofile_profile_t* InitializeDeviceProfilingAqlPackets( return profile; } +// ATT +uint32_t g_output_buffer_size = 0x8000000; // 128M x 16 = 2GB +bool g_output_buffer_local = true; + +// Allocate system memory accessible by both CPU and GPU +uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool) { + hsa_status_t status = HSA_STATUS_ERROR; + uint8_t* buffer = NULL; + size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK; + // if (!cpu_agents_.empty()) { + status = hsa_amd_memory_pool_allocate(*cpu_pool, size, 0, reinterpret_cast(&buffer)); + // Both the CPU and GPU can access the memory + if (status == HSA_STATUS_SUCCESS) { + hsa_agent_t ag_list[1] = {gpu_agent}; + status = hsa_amd_agents_allow_access(1, ag_list, NULL, buffer); + } + // } + uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL; + return ptr; +} + +// Allocate memory for use by a kernel of specified size +uint8_t* AllocateLocalMemory(size_t size, hsa_amd_memory_pool_t* gpu_pool) { + hsa_status_t status = HSA_STATUS_ERROR; + uint8_t* buffer = NULL; + size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK; + status = hsa_amd_memory_pool_allocate(*gpu_pool, size, 0, reinterpret_cast(&buffer)); + uint8_t* ptr = (status == HSA_STATUS_SUCCESS) ? buffer : NULL; + return ptr; +} + +hsa_status_t Allocate(hsa_agent_t gpu_agent, hsa_ven_amd_aqlprofile_profile_t* profile, + hsa_amd_memory_pool_t* cpu_pool, hsa_amd_memory_pool_t* gpu_pool) { + profile->command_buffer.ptr = + AllocateSysMemory(gpu_agent, profile->command_buffer.size, cpu_pool); + profile->output_buffer.size = g_output_buffer_size; + profile->output_buffer.ptr = (g_output_buffer_local) + ? AllocateLocalMemory(profile->output_buffer.size, gpu_pool) + : AllocateSysMemory(gpu_agent, profile->output_buffer.size, cpu_pool); + return (profile->command_buffer.ptr && profile->output_buffer.ptr) ? HSA_STATUS_SUCCESS + : HSA_STATUS_ERROR; +} + +bool AllocateMemoryPools(hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, + hsa_amd_memory_pool_t* cpu_pool, hsa_amd_memory_pool_t* gpu_pool) { + hsa_status_t status = hsa_amd_agent_iterate_memory_pools(cpu_agent, FindStandardPool, cpu_pool); + CHECK_HSA_STATUS("hsa_amd_agent_iterate_memory_pools(cpu_pool)", status); + + status = hsa_amd_agent_iterate_memory_pools(gpu_agent, FindStandardPool, gpu_pool); + CHECK_HSA_STATUS("hsa_amd_agent_iterate_memory_pools(gpu_pool)", status); + + return true; +} + +// map between gpu agent handle and att_memory_pools_t +typedef std::map att_mem_pools_map_t; + +att_mem_pools_map_t* agent_att_mem_pools_map = nullptr; +std::atomic att_map_init{false}; + +att_mem_pools_map_t* GetAttMemPoolsMap() { + if (!att_map_init.load(std::memory_order_relaxed)) { + agent_att_mem_pools_map = new att_mem_pools_map_t(); + att_map_init.exchange(true, std::memory_order_release); + } + + return agent_att_mem_pools_map; +} + + +att_memory_pools_t* GetAttMemPools(hsa_agent_t gpu_agent) { + auto it = GetAttMemPoolsMap()->find(gpu_agent.handle); + if (it != GetAttMemPoolsMap()->end()) { + return it->second; + } + printf("Error: att_memory_pools_t instance not found for given gpu agent handle: %lu\n", + gpu_agent.handle); + + return nullptr; +} + +// Generate start and stop packets for collecting ATT traces +// Also generate and return the profile object which has the PM4 +// command buffer and the output buffer for retrieving the traces +hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets( + hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, + std::vector& att_params, packet_t* start_packet, + packet_t* stop_packet) { + att_memory_pools_t* att_mem_pools = NULL; + auto it = GetAttMemPoolsMap()->find(gpu_agent.handle); + if (it == GetAttMemPoolsMap()->end()) { + att_mem_pools = new att_memory_pools_t; + + // Allocate memory pools for cpu and gpu + AllocateMemoryPools(cpu_agent, gpu_agent, &att_mem_pools->cpu_mem_pool, + &att_mem_pools->gpu_mem_pool); + + GetAttMemPoolsMap()->emplace(gpu_agent.handle, att_mem_pools); + } else + att_mem_pools = it->second; + +#pragma GCC diagnostic push +#pragma GCC diagnostic ignored "-Wconversion-null" + // Preparing the profile structure to get the packets + hsa_ven_amd_aqlprofile_profile_t* profile = + new hsa_ven_amd_aqlprofile_profile_t{gpu_agent, + HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_TRACE, + nullptr, + 0, + &att_params[0], + (uint32_t)att_params.size(), + NULL, + NULL}; +#pragma GCC diagnostic pop + + // Check the profile buffer sizes + hsa_status_t status = hsa_ven_amd_aqlprofile_start(profile, NULL); + if (status != HSA_STATUS_SUCCESS) printf("Error: aqlprofile_start(NULL)"); + // // Double output buffer size if concurrent + // if (is_concurrent) profile.output_buffer.size *= 2; + + // TODO: create a separate class for memory allocations + // Maintain pools per device + // handle allocation and resource cleanup + + + // Allocate command and output buffers + // command buffer -> from CPU memory pool + // output buffer -> from GPU memory pool + status = + Allocate(gpu_agent, profile, &att_mem_pools->cpu_mem_pool, &att_mem_pools->gpu_mem_pool); + if (status != HSA_STATUS_SUCCESS) printf("Error: Allocate()"); + + // Generate start/stop/read profiling packets + status = hsa_ven_amd_aqlprofile_start(profile, start_packet); + if (status != HSA_STATUS_SUCCESS) printf("Error: aqlprofile_start"); + status = hsa_ven_amd_aqlprofile_stop(profile, stop_packet); + if (status != HSA_STATUS_SUCCESS) printf("Error: aqlprofile_stop"); + if (status == HSA_STATUS_ERROR) return nullptr; + return profile; +} + } // namespace Packet diff --git a/src/core/hsa/packets/packets_generator.h b/src/core/hsa/packets/packets_generator.h index d6d6cffb5e..05c92dcd70 100644 --- a/src/core/hsa/packets/packets_generator.h +++ b/src/core/hsa/packets/packets_generator.h @@ -55,5 +55,25 @@ hsa_amd_memory_pool_t& GetCommandPool(); hsa_amd_memory_pool_t& GetOutputPool(); +hsa_ven_amd_aqlprofile_profile_t* GenerateATTPackets( + hsa_agent_t cpu_agent, hsa_agent_t gpu_agent, + std::vector& att_params, packet_t* start_packet, + packet_t* stop_packet); + + +uint8_t* AllocateSysMemory(hsa_agent_t gpu_agent, size_t size, hsa_amd_memory_pool_t* cpu_pool); + +void get_command_buffer_map(std::map ); +void get_outbuffer_map(std::map ); +void initialize_pools(hsa_agent_t cpu_agent); + +typedef struct { + hsa_amd_memory_pool_t cpu_mem_pool; + hsa_amd_memory_pool_t gpu_mem_pool; +} att_memory_pools_t; + +att_memory_pools_t* GetAttMemPools(hsa_agent_t gpu_agent); + + } // namespace Packet #endif // SRC_CORE_HSA_PACKETS_PACKETS_GENERATOR_H_ diff --git a/src/core/hsa/queues/queue.cpp b/src/core/hsa/queues/queue.cpp index bad30c7e7c..0f12fdf283 100644 --- a/src/core/hsa/queues/queue.cpp +++ b/src/core/hsa/queues/queue.cpp @@ -34,6 +34,8 @@ #include "src/utils/helper.h" #define __NR_gettid 186 +#define MAX_ATT_PROFILES 16 + std::mutex sessions_pending_signal_lock; namespace rocmtools { @@ -52,6 +54,8 @@ static inline bool IsEventMatch(const hsa_ven_amd_aqlprofile_event_t& event1, (event1.counter_id == event2.counter_id); } +typedef std::vector att_trace_callback_data_t; + static std::mutex ksymbol_map_lock; static std::map* ksymbols; static std::atomic ksymbols_flag{true}; @@ -234,6 +238,17 @@ hsa_status_t pmcCallback(hsa_ven_amd_aqlprofile_info_type_t info_type, return status; } +hsa_status_t attTraceDataCallback(hsa_ven_amd_aqlprofile_info_type_t info_type, + hsa_ven_amd_aqlprofile_info_data_t* info_data, void* data) { + hsa_status_t status = HSA_STATUS_SUCCESS; + att_trace_callback_data_t* passed_data = reinterpret_cast(data); + passed_data->push_back(*info_data); + // TODO: clear output buffers after copying + // either copy here or in AddattRecord + + return status; +} + void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_signal_t& pending) { rocmtools::metrics::GetCounterData(pending.profile, pending.context->results_list); rocmtools::metrics::GetMetricsData(pending.context->results_map, pending.context->metrics_list); @@ -260,6 +275,47 @@ void AddRecordCounters(rocprofiler_record_profiler_t* record, const pending_sign record->counters_count = rocprofiler_record_counters_instances_count_t{counters_vec.size()}; } +void AddAttRecord(rocprofiler_record_att_tracer_t* record, hsa_agent_t gpu_agent, + att_pending_signal_t& pending) { + att_trace_callback_data_t data; + hsa_ven_amd_aqlprofile_iterate_data(pending.profile, attTraceDataCallback, &data); + + // Get CPU and GPU memory pools + Packet::att_memory_pools_t* att_mem_pools = Packet::GetAttMemPools(gpu_agent); + + // Allocate memory for shader_engine_data + record->shader_engine_data = static_cast( + malloc(data.size() * sizeof(rocprofiler_record_se_att_data_t))); + + att_trace_callback_data_t::iterator trace_data_it; + + uint32_t se_index = 0; + // iterate over the trace data collected from each shader engine + for (trace_data_it = data.begin(); trace_data_it != data.end(); trace_data_it++) { + const void* data_ptr = trace_data_it->trace_data.ptr; + const uint32_t data_size = trace_data_it->trace_data.size; + // fprintf(arg->file, " SE(%u) size(%u)\n", data.sample_id, data_size); + + void* buffer = NULL; + if (data_size != 0) { + // Allocate buffer on CPU to copy out trace data + buffer = Packet::AllocateSysMemory(gpu_agent, data_size, &att_mem_pools->cpu_mem_pool); + if (buffer == NULL) fatal("Trace data buffer allocation failed"); + + auto status = + rocmtools::hsa_support::GetCoreApiTable().hsa_memory_copy_fn(buffer, data_ptr, data_size); + if (status != HSA_STATUS_SUCCESS) fatal("Trace data memcopy to host failed"); + + record->shader_engine_data[se_index].buffer_ptr = buffer; + record->shader_engine_data[se_index].buffer_size = data_size; + ++se_index; + + // TODO: clear output buffers after copying + } + } + record->shader_engine_data_count = data.size(); +} + // static const size_t MEM_PAGE_BYTES = 0x1000; // static const size_t MEM_PAGE_MASK = MEM_PAGE_BYTES - 1; // static std::mutex begin_signal_lock; @@ -418,6 +474,67 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) { return false; } +bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) { + // TODO: finish implementation to iterate trace data and add it to rocmtools record + // and generic buffer + + auto queue_info_session = static_cast(data); + if (!queue_info_session || !GetROCMToolObj() || + !GetROCMToolObj()->GetSession(queue_info_session->session_id) || + !GetROCMToolObj()->GetSession(queue_info_session->session_id)->GetAttTracer()) + return true; + rocmtools::Session* session = GetROCMToolObj()->GetSession(queue_info_session->session_id); + rocmtools::att::AttTracer* att_tracer = session->GetAttTracer(); + std::vector& pending_signals = + const_cast&>( + att_tracer->GetPendingSignals(queue_info_session->writer_id)); + + if (!pending_signals.empty()) { + for (auto it = pending_signals.begin(); it != pending_signals.end(); + it = pending_signals.erase(it)) { + auto& pending = *it; + std::lock_guard lock(session->GetSessionLock()); + if (hsa_support::GetCoreApiTable().hsa_signal_load_relaxed_fn(pending.signal)) return true; + rocprofiler_record_att_tracer_t record{}; + record.kernel_id = rocprofiler_kernel_id_t{pending.kernel_descriptor}; + record.gpu_id = rocprofiler_agent_id_t{ + (uint64_t)hsa_support::GetAgentInfo(queue_info_session->agent.handle).getIndex()}; + record.kernel_properties = pending.kernel_properties; + record.thread_id = rocprofiler_thread_id_t{pending.thread_id}; + record.queue_idx = rocprofiler_queue_index_t{pending.queue_index}; + record.queue_id = rocprofiler_queue_id_t{queue_info_session->queue_id}; + if (/*pending.counters_count > 0 && */ pending.profile) { + AddAttRecord(&record, queue_info_session->agent, pending); + } + record.header = {ROCPROFILER_ATT_TRACER_RECORD, + rocprofiler_record_id_t{GetROCMToolObj()->GetUniqueRecordId()}}; + + if (pending.session_id.handle == 0) { + pending.session_id = GetROCMToolObj()->GetCurrentSessionId(); + } + if (session->FindBuffer(pending.buffer_id)) { + Memory::GenericBuffer* buffer = session->GetBuffer(pending.buffer_id); + record.header.id = rocprofiler_record_id_t{GetROCMToolObj()->GetUniqueRecordId()}; + buffer->AddRecord(record); + } + hsa_status_t status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn( + (pending.profile->output_buffer.ptr)); + if (status != HSA_STATUS_SUCCESS) { + printf("Error: Couldn't free output buffer memory\n"); + } + status = rocmtools::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn( + (pending.profile->command_buffer.ptr)); + if (status != HSA_STATUS_SUCCESS) { + printf("Error: Couldn't free command buffer memory\n"); + } + delete pending.profile; + } + } + delete queue_info_session; + + return false; +} + void CreateBarrierPacket(const hsa_signal_t& packet_completion_signal, std::vector* transformed_packets) { hsa_barrier_and_packet_t barrier{0}; @@ -439,6 +556,12 @@ void SignalAsyncHandler(const hsa_signal_t& signal, void* data) { if (status != HSA_STATUS_SUCCESS) fatal("hsa_amd_signal_async_handler failed"); } +void signalAsyncHandlerATT(const hsa_signal_t& signal, void* data) { + hsa_status_t status = hsa_support::GetAmdExtTable().hsa_amd_signal_async_handler_fn( + signal, HSA_SIGNAL_CONDITION_EQ, 0, AsyncSignalHandlerATT, data); + if (status != HSA_STATUS_SUCCESS) fatal("hsa_amd_signal_async_handler failed"); +} + void CreateSignal(uint32_t attribute, hsa_signal_t* signal) { hsa_status_t status = hsa_support::GetAmdExtTable().hsa_amd_signal_create_fn(1, 0, nullptr, attribute, signal); @@ -459,6 +582,7 @@ template constexpr Integral bit_extract(Integral x, int firs return (x >> first) & bit_mask(0, last - first); } +static int KernelInterceptCount = 0; std::atomic WRITER_ID{0}; /** * @brief This function is a queue write interceptor. It intercepts the @@ -487,9 +611,12 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt bool is_counter_collection_mode = false; bool is_timestamp_collection_mode = false; + bool is_att_collection_mode = false; bool is_pc_sampling_collection_mode = false; - + std::vector att_parameters_data; uint32_t replay_mode_count = 0; + std::vector kernel_profile_names; + std::vector att_counters_names; rocmtools::Session* session = nullptr; @@ -509,6 +636,17 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt session->GetFilterIdWithKind(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION); rocmtools::Filter* filter = session->GetFilter(filter_id); buffer_id = filter->GetBufferId(); + } else if (session && session->FindFilterWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) { + rocprofiler_filter_id_t filter_id = + session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION); + rocmtools::Filter* filter = session->GetFilter(filter_id); + att_parameters_data = filter->GetAttParametersData(); + is_att_collection_mode = true; + buffer_id = session->GetFilter(session->GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) + ->GetBufferId(); + + att_counters_names = filter->GetCounterData(); + kernel_profile_names = std::get>(filter->GetProperty(ROCPROFILER_FILTER_KERNEL_NAMES)); } else if (session && session->FindFilterWithKind(ROCPROFILER_PC_SAMPLING_COLLECTION)) { is_pc_sampling_collection_mode = true; } @@ -527,10 +665,12 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt std::vector>* profiles = nullptr; + // Searching accross all the packets given during this write for (size_t i = 0; i < pkt_count; ++i) { auto& original_packet = static_cast(packets)[i]; + // +Skip kernel dispatch IDs not wanted // Skip packets other than kernel dispatch packets. if (bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) != @@ -659,6 +799,185 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt } /* Write the transformed packets to the hardware queue. */ writer(&transformed_packets[0], transformed_packets.size()); + } else if (session_id.handle > 0 && pkt_count > 0 && is_att_collection_mode && session) { + // att start + // Getting Queue Data and Information + auto& queue_info = *static_cast(data); + std::lock_guard lk(queue_info.qw_mutex); + Agent::AgentInfo* agentInfo = &(hsa_support::GetAgentInfo(queue_info.GetGPUAgent().handle)); + + if (agentInfo->getName().substr(0, 4) != "gfx9") { + printf("ATT collection is only supported for gfx9 at the moment!\n"); + exit(1); + } + + // Preparing att Packets + Packet::packet_t start_packet{}; + Packet::packet_t stop_packet{}; + hsa_ven_amd_aqlprofile_profile_t* profile = nullptr; + + if (att_parameters_data.size() > 0 && is_att_collection_mode) { + // TODO sauverma: convert att_parameters_data to pass to generateattPackets + std::vector att_params; + int num_att_counters = 0; + + for (rocprofiler_att_parameter_t& param : att_parameters_data) { + att_params.push_back({ + static_cast(int(param.parameter_name)), + param.value + }); + num_att_counters += param.parameter_name == ROCPROFILER_ATT_PERFCOUNTER; + } + + if (att_counters_names.size() > 0) { + MetricsDict* metrics_dict_ = MetricsDict::Create(agentInfo); + + for (const std::string& counter_name : att_counters_names) { + const Metric* metric = metrics_dict_->Get(counter_name); + const BaseMetric* base = dynamic_cast(metric); + if (!base) { + printf("Invalid base metric value: %s\n", counter_name.c_str()); + exit(1); + } + std::vector counters; + base->GetCounters(counters); + hsa_ven_amd_aqlprofile_event_t event = counters[0]->event; + if (event.block_name != HSA_VEN_AMD_AQLPROFILE_BLOCK_NAME_SQ) { + printf("Only events from the SQ block can be selected for ATT."); + exit(1); + } + att_params.push_back({ + static_cast(int(ROCPROFILER_ATT_PERFCOUNTER)), + event.counter_id | (event.counter_id ? (0xF<<24) : 0) + }); + num_att_counters += 1; + } + + hsa_ven_amd_aqlprofile_parameter_t zero_perf = { + static_cast(int(ROCPROFILER_ATT_PERFCOUNTER)), 0}; + + // Fill other perfcounters with 0's + for(; num_att_counters<16; num_att_counters++) att_params.push_back(zero_perf); + } + + // Get the PM4 Packets using packets_generator + profile = Packet::GenerateATTPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(), + att_params, &start_packet, &stop_packet); + } + + + // Searching across all the packets given during this write + for (size_t i = 0; i < pkt_count; ++i) { + auto& original_packet = static_cast(packets)[i]; + + // Skip packets other than kernel dispatch packets. + if (bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE, + HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1) != + HSA_PACKET_TYPE_KERNEL_DISPATCH) { + transformed_packets.emplace_back(packets_arr[i]); + continue; + } + + auto& kdispatch = static_cast(packets)[i]; + uint64_t kernel_object = kdispatch.kernel_object; + bool b_profile_this_object = false; + + // Try to match the mangled kernel name with given matches in input.txt + try { + std::lock_guard lock(ksymbol_map_lock); + assert(ksymbols); + const std::string& kernel_name = ksymbols->at(kernel_object); + + // We want to initiate att profiling only if a match exists + for(const std::string& kernel_matches : kernel_profile_names) { + if (kernel_name.find(kernel_matches) != std::string::npos) { + b_profile_this_object = true; + break; + } + } + if (!b_profile_this_object) printf("Skipping: %s\n", kernel_name.c_str()); + } catch (...) { + printf("Warning: Unknown name for object %lu\n", kernel_object); + } + + // If no match was found or intercept count > maximum desired profiles, skip this kernel. + if (!b_profile_this_object || KernelInterceptCount >= MAX_ATT_PROFILES) { + printf("Skipping: %lu\n", kernel_object); + transformed_packets.emplace_back(packets_arr[i]); + continue; + } + KernelInterceptCount += 1; + + uint32_t writer_id = WRITER_ID.fetch_add(1, std::memory_order_release); + + + if (att_parameters_data.size() > 0 && is_att_collection_mode && profile) { + // Adding start packet and its barrier with a dummy signal + hsa_signal_t dummy_signal{}; + dummy_signal.handle = 0; + start_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + AddVendorSpecificPacket(&start_packet, &transformed_packets, dummy_signal); + CreateBarrierPacket(start_packet.completion_signal, &transformed_packets); + } + + auto& packet = transformed_packets.emplace_back(packets_arr[i]); + auto& dispatch_packet = reinterpret_cast(packet); + + CreateSignal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &packet.completion_signal); + // Adding the dispatch packet newly created signal to the pending signals + // list to be processed by the signal interrupt + rocprofiler_kernel_properties_t kernel_properties = + set_kernel_properties(dispatch_packet, queue_info.GetGPUAgent()); + if (session && profile) { + session->GetAttTracer()->AddPendingSignals( + writer_id, dispatch_packet.kernel_object, dispatch_packet.completion_signal, session_id, + buffer_id, profile, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index); + } else { + session->GetAttTracer()->AddPendingSignals( + writer_id, dispatch_packet.kernel_object, dispatch_packet.completion_signal, session_id, + buffer_id, nullptr, kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index); + } + + // Make a copy of the original packet, adding its signal to a barrier packet + if (original_packet.completion_signal.handle) { + hsa_barrier_and_packet_t barrier{0}; + barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; + Packet::packet_t* __attribute__((__may_alias__)) pkt = + (reinterpret_cast(&barrier)); + transformed_packets.emplace_back(*pkt).completion_signal = + original_packet.completion_signal; + } + + // Adding a barrier packet with the original packet's completion signal. + hsa_signal_t interrupt_signal; + CreateSignal(0, &interrupt_signal); + + // Adding Stop PM4 Packets + if (att_parameters_data.size() > 0 && is_att_collection_mode && profile) { + stop_packet.header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE; + AddVendorSpecificPacket(&stop_packet, &transformed_packets, interrupt_signal); + + // Added Interrupt Signal with barrier and provided handler for it + CreateBarrierPacket(interrupt_signal, &transformed_packets); + } else { + hsa_barrier_and_packet_t barrier{0}; + barrier.header = HSA_PACKET_TYPE_BARRIER_AND << HSA_PACKET_HEADER_TYPE; + barrier.completion_signal = interrupt_signal; + Packet::packet_t* __attribute__((__may_alias__)) pkt = + (reinterpret_cast(&barrier)); + transformed_packets.emplace_back(*pkt); + } + + // Creating Async Handler to be called every time the interrupt signal is + // marked complete + signalAsyncHandlerATT( + interrupt_signal, + new queue_info_session_t{queue_info.GetGPUAgent(), session_id, queue_info.GetQueueID(), + writer_id, interrupt_signal}); + } + /* Write the transformed packets to the hardware queue. */ + writer(&transformed_packets[0], transformed_packets.size()); + // ATT end } else { /* Write the original packets to the hardware queue if no profiling session * is active */ diff --git a/src/core/memory/generic_buffer.cpp b/src/core/memory/generic_buffer.cpp index 585425eaca..798f64a91d 100644 --- a/src/core/memory/generic_buffer.cpp +++ b/src/core/memory/generic_buffer.cpp @@ -43,7 +43,8 @@ GenericBuffer::GenericBuffer(rocprofiler_session_id_t session_id, rocprofiler_bu // pointer moves to the other buffer. Each buffer should be large enough to // hold at least 2 activity records, as record pairs may be written when // external correlation ids are used. - const size_t allocation_size = 2 * std::max(2 * sizeof(rocprofiler_record_header_t), buffer_size); + const size_t allocation_size = + 2 * std::max(2 * sizeof(rocprofiler_record_header_t), buffer_size); pool_begin_ = nullptr; AllocateMemory(&pool_begin_, allocation_size); assert(pool_begin_ != nullptr && "pool allocator failed"); @@ -213,6 +214,12 @@ bool GetNextRecord(const rocprofiler_record_header_t* record, // break; *next = reinterpret_cast(tracer_record + 1); } + case ROCPROFILER_ATT_TRACER_RECORD: { + const rocprofiler_record_att_tracer_t* att_tracer_record = + reinterpret_cast(record); + *next = reinterpret_cast(att_tracer_record + 1); + break; + } default: const rocprofiler_record_tracer_t* tracer_record = reinterpret_cast(record); diff --git a/src/core/session/att/att.cpp b/src/core/session/att/att.cpp new file mode 100644 index 0000000000..ffa6a133c3 --- /dev/null +++ b/src/core/session/att/att.cpp @@ -0,0 +1,57 @@ +/* Copyright (c) 2022 Advanced Micro Devices, Inc. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. */ + +#include "att.h" +#include + +namespace rocmtools { + +namespace att { + +AttTracer::AttTracer(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t filter_id, + rocprofiler_session_id_t session_id) + : buffer_id_(buffer_id), filter_id_(filter_id), session_id_(session_id) {} +AttTracer::~AttTracer() {} + +void AttTracer::AddPendingSignals(uint32_t writer_id, uint64_t kernel_object, + const hsa_signal_t& completion_signal, + rocprofiler_session_id_t session_id, + rocprofiler_buffer_id_t buffer_id, + hsa_ven_amd_aqlprofile_profile_t* profile, + rocprofiler_kernel_properties_t kernel_properties, + uint32_t thread_id, uint64_t queue_index) { + std::lock_guard lock(sessions_pending_signals_lock_); + if (sessions_pending_signals_.find(writer_id) == sessions_pending_signals_.end()) + sessions_pending_signals_.emplace(writer_id, std::vector()); + sessions_pending_signals_.at(writer_id).emplace_back( + att_pending_signal_t{kernel_object, completion_signal, session_id_, buffer_id, profile, + kernel_properties, thread_id, queue_index}); +} + +const std::vector& AttTracer::GetPendingSignals(uint32_t writer_id) { + std::lock_guard lock(sessions_pending_signals_lock_); + assert(sessions_pending_signals_.find(writer_id) != sessions_pending_signals_.end() && + "writer_id is not found in the pending_signals"); + return sessions_pending_signals_.at(writer_id); +} + +} // namespace att + +} // namespace rocmtools diff --git a/src/core/session/att/att.h b/src/core/session/att/att.h new file mode 100644 index 0000000000..85360db7b4 --- /dev/null +++ b/src/core/session/att/att.h @@ -0,0 +1,76 @@ +/* Copyright (c) 2022 Advanced Micro Devices, Inc. + + 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. */ + +#ifndef SRC_CORE_SESSION_ATT_ATT_H_ +#define SRC_CORE_SESSION_ATT_ATT_H_ + +#include + +#include +#include +#include +#include + +#include "inc/rocprofiler.h" + +namespace rocmtools { + +typedef struct { + uint64_t kernel_descriptor; + hsa_signal_t signal; + rocprofiler_session_id_t session_id; + rocprofiler_buffer_id_t buffer_id; + hsa_ven_amd_aqlprofile_profile_t* profile; + rocprofiler_kernel_properties_t kernel_properties; + uint32_t thread_id; + uint64_t queue_index; +} att_pending_signal_t; + +namespace att { + +class AttTracer { + public: + AttTracer(rocprofiler_buffer_id_t buffer_id, rocprofiler_filter_id_t filter_id, + rocprofiler_session_id_t session_id); + ~AttTracer(); + + void AddPendingSignals(uint32_t writer_id, uint64_t kernel_object, + const hsa_signal_t& completion_signal, rocprofiler_session_id_t session_id, + rocprofiler_buffer_id_t buffer_id, hsa_ven_amd_aqlprofile_profile_t* profile, + rocprofiler_kernel_properties_t kernel_properties, uint32_t thread_id, + uint64_t queue_index); + + const std::vector& GetPendingSignals(uint32_t writer_id); + + private: + rocprofiler_buffer_id_t buffer_id_; + rocprofiler_filter_id_t filter_id_; + rocprofiler_session_id_t session_id_; + + std::mutex sessions_pending_signals_lock_; + std::map> sessions_pending_signals_; +}; + +} // namespace att + +} // namespace rocmtools + + +#endif // SRC_CORE_SESSION_ATT_ATT_H_ diff --git a/src/core/session/filter.cpp b/src/core/session/filter.cpp index b659739c13..68ade3459e 100644 --- a/src/core/session/filter.cpp +++ b/src/core/session/filter.cpp @@ -35,12 +35,25 @@ Filter::Filter(rocprofiler_filter_id_t id, rocprofiler_filter_kind_t filter_kind } case ROCPROFILER_COUNTERS_COLLECTION: { profiler_counter_names_.clear(); - for (uint32_t j = 0; j < data_count; j++) + for (uint32_t j = 0; j < data_count; j++) { profiler_counter_names_.emplace_back(filter_data.counters_names[j]); + } break; } - case ROCPROFILER_PC_SAMPLING_COLLECTION: - case ROCPROFILER_ATT_TRACE: { + case ROCPROFILER_PC_SAMPLING_COLLECTION:{ + break; + } + case ROCPROFILER_ATT_TRACE_COLLECTION: { + att_parameters_.clear(); + profiler_counter_names_.clear(); + + for (uint32_t j = 0; j < data_count; j++) { + if (filter_data.att_parameters[j].parameter_name != ROCPROFILER_ATT_PERFCOUNTER_NAME) { + att_parameters_.emplace_back(filter_data.att_parameters[j]); + } else { + profiler_counter_names_.emplace_back(filter_data.att_parameters[j].counter_name); + } + } break; } case ROCPROFILER_SPM_COLLECTION: { @@ -49,8 +62,9 @@ Filter::Filter(rocprofiler_filter_id_t id, rocprofiler_filter_kind_t filter_kind } case ROCPROFILER_API_TRACE: { tracer_apis_.clear(); - for (uint32_t j = 0; j < data_count; j++) - tracer_apis_.emplace_back(filter_data.trace_apis[j]); + for (uint32_t j = 0; j < data_count; j++){ + tracer_apis_.emplace_back(filter_data.trace_apis[j]); + } break; } default: { @@ -73,7 +87,7 @@ rocprofiler_filter_kind_t Filter::GetKind() { return kind_; } std::mutex counter_data_lock; std::vector Filter::GetCounterData() { - if (kind_ == ROCPROFILER_COUNTERS_COLLECTION) { + if (kind_ == ROCPROFILER_COUNTERS_COLLECTION || kind_ == ROCPROFILER_ATT_TRACE_COLLECTION) { std::lock_guard lock(counter_data_lock); return profiler_counter_names_; } @@ -90,6 +104,16 @@ std::vector Filter::GetTraceData() { "Error: ROCMtools filter specified is not supported for " "profiler mode!\n"); } + +std::vector Filter::GetAttParametersData() { + if (kind_ == ROCPROFILER_ATT_TRACE_COLLECTION) { + return att_parameters_; + } + fatal( + "Error: ROCMtools filter specified is not supported for " + "ATT tracing mode!\n"); +} + rocprofiler_spm_parameter_t* Filter::GetSpmParameterData() { if (kind_ == ROCPROFILER_SPM_COLLECTION) { return spm_parameter_; @@ -143,7 +167,8 @@ void Filter::SetProperty(rocprofiler_filter_property_t property) { } case ROCPROFILER_FILTER_KERNEL_NAMES: { if (kind_ == ROCPROFILER_COUNTERS_COLLECTION || - kind_ == ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION) { + kind_ == ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION || + kind_ == ROCPROFILER_ATT_TRACE_COLLECTION) { kernel_names_.clear(); for (uint32_t j = 0; j < property.data_count; j++) kernel_names_.emplace_back(property.name_regex[j]); @@ -166,23 +191,29 @@ std::variant, uint32_t*> Filter::GetProperty( switch (kind) { case ROCPROFILER_FILTER_GPU_NAME: { property = agent_names_; + break; } case ROCPROFILER_FILTER_RANGE: { property = static_cast(dispatch_range_); + break; } case ROCPROFILER_FILTER_KERNEL_NAMES: { property = kernel_names_; + break; } case ROCPROFILER_FILTER_HSA_TRACER_API_FUNCTIONS: { property = hsa_tracer_api_calls_; + break; } case ROCPROFILER_FILTER_HIP_TRACER_API_FUNCTIONS: { property = hip_tracer_api_calls_; + break; } default: fatal( "Error: ROCMtools filter specified is not supported for the given " "kind!"); + break; } return property; } diff --git a/src/core/session/filter.h b/src/core/session/filter.h index 98091b8889..8f3621ffde 100644 --- a/src/core/session/filter.h +++ b/src/core/session/filter.h @@ -47,7 +47,7 @@ class Filter { std::vector GetCounterData(); std::vector GetTraceData(); - + std::vector GetAttParametersData(); void SetCallback(rocprofiler_sync_callback_t& callback); rocprofiler_sync_callback_t& GetCallback(); @@ -71,6 +71,7 @@ class Filter { std::vector profiler_counter_names_; // Counter Names to collect std::vector tracer_apis_; // ROCTX/HIP/HSA API rocprofiler_spm_parameter_t* spm_parameter_; // spm parameter + std::vector att_parameters_; // ATT Parameters rocprofiler_sync_callback_t callback_; }; diff --git a/src/core/session/session.cpp b/src/core/session/session.cpp index ce441b96dd..e3d80a25a6 100644 --- a/src/core/session/session.cpp +++ b/src/core/session/session.cpp @@ -54,6 +54,10 @@ Session::~Session() { // delete tracer_; // tracer_started_.exchange(false, std::memory_order_release); // } + if (att_tracer_started_.load(std::memory_order_release)) { + delete att_tracer_; + att_tracer_started_.exchange(false, std::memory_order_release); + } // { // std::lock_guard lock(filters_lock_); // buffers_.clear(); @@ -99,6 +103,13 @@ void Session::Start() { GetFilter(GetFilterIdWithKind(ROCPROFILER_COUNTERS_COLLECTION))->GetId(), session_id_); profiler_started_.exchange(true, std::memory_order_release); } + if (FindFilterWithKind(ROCPROFILER_ATT_TRACE_COLLECTION)) { + if (att_tracer_started_.load(std::memory_order_release)) delete att_tracer_; + att_tracer_ = new att::AttTracer( + GetFilter(GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION))->GetBufferId(), + GetFilter(GetFilterIdWithKind(ROCPROFILER_ATT_TRACE_COLLECTION))->GetId(), session_id_); + att_tracer_started_.exchange(true, std::memory_order_release); + } if (FindFilterWithKind(ROCPROFILER_SPM_COLLECTION)) { if (spm_started_.load(std::memory_order_release)) delete spmcounter_; @@ -176,6 +187,7 @@ rocprofiler_session_id_t Session::GetId() { return session_id_; } bool Session::IsActive() { return is_active_; } profiler::Profiler* Session::GetProfiler() { return profiler_; } +att::AttTracer* Session::GetAttTracer() { return att_tracer_; } tracer::Tracer* Session::GetTracer() { return tracer_; } spm::SpmCounters* Session::GetSpmCounter() { return spmcounter_; } pc_sampler::PCSampler* Session::GetPCSampler() { return pc_sampler_; } diff --git a/src/core/session/session.h b/src/core/session/session.h index 2ee6b195f0..561e16a4a2 100644 --- a/src/core/session/session.h +++ b/src/core/session/session.h @@ -37,6 +37,7 @@ #include "src/core/session/filter.h" #include "profiler/profiler.h" #include "tracer/tracer.h" +#include "att/att.h" #include "spm/spm.h" #include "src/pcsampler/session/pc_sampler.h" @@ -58,6 +59,7 @@ class Session { profiler::Profiler* GetProfiler(); tracer::Tracer* GetTracer(); + att::AttTracer* GetAttTracer(); spm::SpmCounters* GetSpmCounter(); pc_sampler::PCSampler* GetPCSampler(); @@ -104,6 +106,8 @@ class Session { std::atomic profiler_started_{false}; std::atomic tracer_started_{false}; + std::atomic att_tracer_started_{false}; + att::AttTracer* att_tracer_; std::atomic spm_started_{false}; profiler::Profiler* profiler_; diff --git a/src/tools/tool.cpp b/src/tools/tool.cpp index 45db8dba2c..8cd8fe8b40 100644 --- a/src/tools/tool.cpp +++ b/src/tools/tool.cpp @@ -172,6 +172,134 @@ std::vector GetCounterNames() { return counters; } +typedef std::tuple< + std::vector>, + std::vector, std::vector +> att_parsed_input_t; + +att_parsed_input_t GetATTParams() { + std::vector> parameters; + std::vector kernel_names; + std::vector counters_names; + const char* path = getenv("COUNTERS_PATH"); + + // List of parameters the user can set. Maxvalue is unused. + std::unordered_map ATT_PARAM_NAMES{}; + + ATT_PARAM_NAMES["att: TARGET_CU"] = ROCPROFILER_ATT_COMPUTE_UNIT_TARGET; + ATT_PARAM_NAMES["SIMD_MASK"] = ROCPROFILER_ATT_MAXVALUE; + ATT_PARAM_NAMES["PERFCOUNTER_ID"] = ROCPROFILER_ATT_PERFCOUNTER; + ATT_PARAM_NAMES["PERFCOUNTER"] = ROCPROFILER_ATT_PERFCOUNTER_NAME; + ATT_PARAM_NAMES["PERFCOUNTERS_COL_PERIOD"] = ROCPROFILER_ATT_MAXVALUE; + ATT_PARAM_NAMES["KERNEL"] = ROCPROFILER_ATT_MAXVALUE; + ATT_PARAM_NAMES["REDUCED_MEMORY"] = ROCPROFILER_ATT_MAXVALUE; + /* + ATT_PARAM_NAMES["ATT_MASK"] = ROCMTOOLS_ATT_MASK; + ATT_PARAM_NAMES["TOKEN_MASK"] = ROCMTOOLS_ATT_TOKEN_MASK; + ATT_PARAM_NAMES["TOKEN_MASK2"] = ROCMTOOLS_ATT_TOKEN_MASK2; + ATT_PARAM_NAMES["SE_MASK"] = ROCMTOOLS_ATT_SE_MASK; + ATT_PARAM_NAMES["PERF_MASK"] = ROCMTOOLS_ATT_PERF_MASK; + ATT_PARAM_NAMES["PERF_CTRL"] = ROCMTOOLS_ATT_PERF_CTRL; +*/ + + // Default values used for token generation. + std::unordered_map default_params = { + {"ATT_MASK", 0x3F01}, + {"TOKEN_MASK", 0x344B}, + {"TOKEN_MASK2", 0x3FFFF} + }; + + bool started_att_counters = false; + + if (!path) return {parameters, kernel_names, counters_names}; + + std::string line; + std::ifstream trace_file(path); + if (!trace_file.is_open()) { + std::cout << "Unable to open att trace file." << std::endl; + return {parameters, kernel_names, counters_names}; + } + + while (getline(trace_file, line)) { + if (line.find("//") != std::string::npos) + line = line.substr(0, line.find("//")); // Remove comments + + auto pos = line.find('='); + if (pos == std::string::npos) + continue; + + std::string param_name = line.substr(0, pos); + uint32_t param_value; + + if (param_name == "att: TARGET_CU") started_att_counters = true; + if (!started_att_counters) continue; + + if (param_name == "KERNEL") { + kernel_names.push_back(line.substr(pos+1)); + continue; + } else if (param_name == "PERFCOUNTER") { + counters_names.push_back(line.substr(pos+1)); + continue; + } else { // param_value is a number + try { + auto hexa_pos = line.find("0x", pos); // Is it hex? + if (hexa_pos != std::string::npos) + param_value = stoi(line.substr(hexa_pos+2), 0, 16); // hexadecimal + else + param_value = stoi(line.substr(pos+1), 0, 10); // decimal + } catch(...) { + printf("Error: Invalid parameter value %s - (%s)\n", + line.substr(pos+1, line.size()).c_str(), line.c_str()); + exit(1); + } + } + + if (param_name == "PERFCOUNTERS_COL_PERIOD") { + default_params["TOKEN_MASK"] |= 0x4000; + param_value = ((param_value & 0x1F) << 8) | 0x007F; + parameters.push_back(std::make_pair(ROCPROFILER_ATT_PERF_CTRL, param_value)); + continue; + } else if (param_name == "SIMD_MASK") { + default_params["ATT_MASK"] &= ~0xF00; + default_params["ATT_MASK"] |= (param_value<<8) & 0xF00; + continue; + } else if (param_name == "att: TARGET_CU") { + default_params["ATT_MASK"] &= ~0xF; + default_params["ATT_MASK"] |= param_value & 0xF; + } else if (param_name == "PERFCOUNTER_ID") { + param_value = param_value | (param_value ? (0xF<<24) : 0); + } else if (param_name == "REDUCED_MEMORY") { + default_params["TOKEN_MASK2"] = 0; + continue; + } + + if (ATT_PARAM_NAMES.find(param_name) != ATT_PARAM_NAMES.end()) { + parameters.push_back(std::make_pair(ATT_PARAM_NAMES[param_name], param_value)); + try { default_params.erase(param_name); } catch(...) {}; + } else { + printf("Error: Invalid parameter name: %s - (%s)\nList of available params:\n", + param_name.c_str(), line.c_str()); + for (auto& name : ATT_PARAM_NAMES) printf("%s\n", name.first.c_str()); + } + } + trace_file.close(); + + if (!started_att_counters) return {parameters, kernel_names, counters_names}; + + ATT_PARAM_NAMES["ATT_MASK"] = ROCPROFILER_ATT_MASK; + ATT_PARAM_NAMES["TOKEN_MASK"] = ROCPROFILER_ATT_TOKEN_MASK; + ATT_PARAM_NAMES["TOKEN_MASK2"] = ROCPROFILER_ATT_TOKEN_MASK2; + + for (auto& param : default_params) + parameters.push_back(std::make_pair(ATT_PARAM_NAMES[param.first], param.second)); + + // If no kernel names were provided, collect them all. + // Empty string always returns true for "str.find()". + if (kernel_names.size() == 0) kernel_names.push_back(""); + + return {parameters, kernel_names, counters_names}; +} + void finish() { if (amd_sys_handler.load(std::memory_order_release)) { amd_sys_handler.exchange(false, std::memory_order_release); @@ -335,19 +463,34 @@ ROCPROFILER_EXPORT bool OnLoad(HsaApiTable* table, uint64_t runtime_version, std::vector counters_; if (counters.size() > 0) { - printf("ROCMTools: Collecting the following counters:\n"); + printf("ROCProfilerV2: Collecting the following counters:\n"); for (size_t i = 0; i < counters.size(); i++) { counters_.emplace_back(counters.at(i).c_str()); printf("- %s\n", counters_.back()); } } + // ATT Parameters + std::vector parameters; + std::vector> params; + std::vector kernel_names; + std::vector att_counters_names; + std::tie(params, kernel_names, att_counters_names) = GetATTParams(); + + for (auto& kv_pair : params) + parameters.emplace_back(rocprofiler_att_parameter_t{kv_pair.first, kv_pair.second}); + for (std::string& name : att_counters_names) { + rocprofiler_att_parameter_t param; + param.parameter_name = ROCPROFILER_ATT_PERFCOUNTER_NAME; + param.counter_name = name.c_str(); + parameters.emplace_back(param); + } CHECK_ROCMTOOLS(rocprofiler_create_session(ROCPROFILER_KERNEL_REPLAY_MODE, &session_id)); bool want_pc_sampling = getenv("ROCPROFILER_PC_SAMPLING"); std::vector filters_requested; - if ((counters.size() == 0 && (apis_requested.size() == 0 || getenv("ROCPROFILER_KERNEL_TRACE"))) + if (((counters.size() == 0 && parameters.size() == 0) && (apis_requested.size() == 0 || getenv("ROCPROFILER_KERNEL_TRACE"))) || want_pc_sampling /* PC sampling needs a profiler, even it's doing nothing */) filters_requested.emplace_back(ROCPROFILER_DISPATCH_TIMESTAMPS_COLLECTION); @@ -356,6 +499,7 @@ ROCPROFILER_EXPORT bool OnLoad(HsaApiTable* table, uint64_t runtime_version, } if (counters.size() > 0) filters_requested.emplace_back(ROCPROFILER_COUNTERS_COLLECTION); if (apis_requested.size() > 0) filters_requested.emplace_back(ROCPROFILER_API_TRACE); + if (parameters.size() > 0) filters_requested.emplace_back(ROCPROFILER_ATT_TRACE_COLLECTION); rocprofiler_buffer_id_t buffer_id; CHECK_ROCMTOOLS(rocprofiler_create_buffer( @@ -364,7 +508,7 @@ ROCPROFILER_EXPORT bool OnLoad(HsaApiTable* table, uint64_t runtime_version, rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id) { if (plugin) plugin->write_buffer_records(record, end_record, session_id, buffer_id); }, - 0x9999, &buffer_id)); + 1<<20, &buffer_id)); buffer_ids.emplace_back(buffer_id); rocprofiler_buffer_id_t buffer_id_1; @@ -374,7 +518,7 @@ ROCPROFILER_EXPORT bool OnLoad(HsaApiTable* table, uint64_t runtime_version, rocprofiler_session_id_t session_id, rocprofiler_buffer_id_t buffer_id_1) { if (plugin) plugin->write_buffer_records(record, end_record, session_id, buffer_id_1); }, - 0x9999, &buffer_id_1)); + 1<<20, &buffer_id_1)); buffer_ids.emplace_back(buffer_id_1); for (rocprofiler_filter_kind_t filter_kind : filters_requested) { @@ -412,6 +556,26 @@ ROCPROFILER_EXPORT bool OnLoad(HsaApiTable* table, uint64_t runtime_version, filter_ids.emplace_back(filter_id); break; } + case ROCPROFILER_ATT_TRACE_COLLECTION: { + printf("Enabling ATT Tracing\n"); + rocprofiler_filter_id_t filter_id; + + std::vector kernel_names_c; + for (auto& name : kernel_names) kernel_names_c.push_back(name.data()); + + rocprofiler_filter_property_t property = {}; + property.kind = ROCPROFILER_FILTER_KERNEL_NAMES; + property.data_count = kernel_names_c.size(); + property.name_regex = kernel_names_c.data(); + + CHECK_ROCMTOOLS( + rocprofiler_create_filter(session_id, ROCPROFILER_ATT_TRACE_COLLECTION, + rocprofiler_filter_data_t{.att_parameters = ¶meters[0]}, + parameters.size(), &filter_id, property)); + CHECK_ROCMTOOLS(rocprofiler_set_filter_buffer(session_id, filter_id, buffer_id_1)); + filter_ids.emplace_back(filter_id); + break; + } case ROCPROFILER_PC_SAMPLING_COLLECTION: { puts("Enabling PC sampling"); rocprofiler_filter_id_t filter_id; diff --git a/tests/featuretests/profiler/CMakeLists.txt b/tests/featuretests/profiler/CMakeLists.txt index 96c2baeae6..f03e5bfe8e 100644 --- a/tests/featuretests/profiler/CMakeLists.txt +++ b/tests/featuretests/profiler/CMakeLists.txt @@ -185,6 +185,12 @@ target_link_libraries ( profiler_api_test ${ROCPROFILER_TARGET} ${ROCPROFILER_TA target_include_directories(profiler_api_test PRIVATE ${TEST_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH} ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) add_dependencies(tests profiler_api_test) +set_source_files_properties(discretetests/api/att_test.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) +hip_add_executable ( profiler_api_att_test discretetests/api/att_test.cpp utils/test_utils.cpp ) +target_link_libraries ( profiler_api_att_test ${ROCPROFILER_TARGET} ${ROCPROFILER_TARGET} hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs ) +target_include_directories(profiler_api_att_test PRIVATE ${TEST_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH} ${PROJECT_SOURCE_DIR} ${PROJECT_SOURCE_DIR}/tests/featuretests/profiler) +add_dependencies(tests profiler_api_att_test) + set_source_files_properties(discretetests/api/spm_test.cpp PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1) hip_add_executable ( profiler_api_spm_test discretetests/api/spm_test.cpp utils/test_utils.cpp ) target_link_libraries ( profiler_api_spm_test ${ROCPROFILER_TARGET} ${ROCPROFILER_TARGET} hsa-runtime64::hsa-runtime64 Threads::Threads dl stdc++fs ) diff --git a/tests/featuretests/profiler/discretetests/api/att_test.cpp b/tests/featuretests/profiler/discretetests/api/att_test.cpp new file mode 100644 index 0000000000..2b5838a784 --- /dev/null +++ b/tests/featuretests/profiler/discretetests/api/att_test.cpp @@ -0,0 +1,259 @@ +/****************************************************************************** +Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*******************************************************************************/ +#include +#include + +#include +#include +#include +#include +#include + +#include "utils/test_utils.h" + +#ifdef NDEBUG +#define HIP_ASSERT(x) x +#else +#define HIP_ASSERT(x) (assert((x)==hipSuccess)) +#endif + + +#define WIDTH 1024 +#define HEIGHT 1024 + +#define NUM (WIDTH*HEIGHT) + +#define THREADS_PER_BLOCK_X 16 +#define THREADS_PER_BLOCK_Y 16 +#define THREADS_PER_BLOCK_Z 1 + + + +/** \mainpage ROC Profiler API Test + * + * \section introduction Introduction + * + * The goal of this test is to test ROCmTools APIs to collect ATT traces. + * + * A simple vectoradd_float kernel is launched and the trace results are printed + * as console output + */ + + +// function to check att tracing API status +auto CheckApi = [](rocprofiler_status_t status) { + if (status != ROCPROFILER_STATUS_SUCCESS) { + std::cout << "ROCmTools API Error" << std::endl; + } + assert(status == ROCPROFILER_STATUS_SUCCESS); +}; + + +// callback function to dump att tracing data +void FlushCallback(const rocprofiler_record_header_t* record, + const rocprofiler_record_header_t* end_record, rocprofiler_session_id_t session_id, + rocprofiler_buffer_id_t buffer_id) { + while (record < end_record) { + if (!record) break; + else if (record->kind == ROCPROFILER_ATT_TRACER_RECORD){ + const rocprofiler_record_att_tracer_t* att_tracer_record = + reinterpret_cast(record); + size_t name_length; + CheckApi(rocprofiler_query_kernel_info_size(ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, + &name_length)); + const char* kernel_name_c = static_cast(malloc(name_length * sizeof(char))); + CheckApi(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, + &kernel_name_c)); + int gpu_index = att_tracer_record->gpu_id.handle; + printf( + "Kernel Info:\n\tGPU Index: %d\n\tKernel Name: %s\n", + gpu_index, kernel_name_c); + + // Get the number of shader engine traces + int se_num = att_tracer_record->shader_engine_data_count; + + // iterate over each shader engine att trace + for (int i = 0; i < se_num; i++){ + + printf("\n\n-------------- shader_engine %d --------------\n\n", i); + rocprofiler_record_se_att_data_t* se_att_trace = &att_tracer_record->shader_engine_data[i]; + uint32_t size = se_att_trace->buffer_size; + const unsigned short* data_buffer_ptr = reinterpret_cast(se_att_trace->buffer_ptr); + + // Print the buffer in terms of shorts (16 bits) + for (uint32_t j = 0; j < (size / sizeof(short)); j++) + printf("%04x\n", data_buffer_ptr[j]); + + } + + } + CheckApi(rocprofiler_next_record(record, &record, session_id, buffer_id)); + } +} + + + + +__global__ void +vectoradd_float(float* __restrict__ a, const float* __restrict__ b, const float* __restrict__ c, int width, int height) + + { + + int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x; + int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y; + + int i = y * width + x; + if ( i < (width * height)) { + a[i] = b[i] + c[i]; + } + + + + } + +int LaunchVectorAddKernel() { + + float* hostA; + float* hostB; + float* hostC; + + float* deviceA; + float* deviceB; + float* deviceC; + + hipDeviceProp_t devProp; + hipGetDeviceProperties(&devProp, 0); + std::cout << " System minor " << devProp.minor << std::endl; + std::cout << " System major " << devProp.major << std::endl; + std::cout << " agent prop name " << devProp.name << std::endl; + + + + std::cout << "hip Device prop succeeded " << std::endl ; + + + int i; + int errors; + + hostA = (float*)malloc(NUM * sizeof(float)); + hostB = (float*)malloc(NUM * sizeof(float)); + hostC = (float*)malloc(NUM * sizeof(float)); + + // initialize the input data + for (i = 0; i < NUM; i++) { + hostB[i] = (float)i; + hostC[i] = (float)i*100.0f; + } + + HIP_ASSERT(hipMalloc((void**)&deviceA, NUM * sizeof(float))); + HIP_ASSERT(hipMalloc((void**)&deviceB, NUM * sizeof(float))); + HIP_ASSERT(hipMalloc((void**)&deviceC, NUM * sizeof(float))); + + HIP_ASSERT(hipMemcpy(deviceB, hostB, NUM*sizeof(float), hipMemcpyHostToDevice)); + HIP_ASSERT(hipMemcpy(deviceC, hostC, NUM*sizeof(float), hipMemcpyHostToDevice)); + + + hipLaunchKernelGGL(vectoradd_float, + dim3(WIDTH/THREADS_PER_BLOCK_X, HEIGHT/THREADS_PER_BLOCK_Y), + dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), + 0, 0, + deviceA ,deviceB ,deviceC ,WIDTH ,HEIGHT); + + + HIP_ASSERT(hipMemcpy(hostA, deviceA, NUM*sizeof(float), hipMemcpyDeviceToHost)); + + // verify the results + errors = 0; + for (i = 0; i < NUM; i++) { + if (hostA[i] != (hostB[i] + hostC[i])) { + errors++; + } + } + if (errors!=0) { + printf("FAILED: %d errors\n",errors); + } else { + printf ("PASSED!\n"); + } + + HIP_ASSERT(hipFree(deviceA)); + HIP_ASSERT(hipFree(deviceB)); + HIP_ASSERT(hipFree(deviceC)); + + free(hostA); + free(hostB); + free(hostC); + + //hipResetDefaultAccelerator(); + + return errors; +} + + +int main(int argc, char** argv) { + + // inititalize rocmtools + CheckApi(rocprofiler_initialize()); + + // Att trace collection parameters + rocprofiler_session_id_t session_id; + std::vector parameters; + parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_COMPUTE_UNIT_TARGET, 0}); + parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_MASK, 0x0F00}); + parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_TOKEN_MASK, 0x344B}); + parameters.emplace_back(rocprofiler_att_parameter_t{ROCPROFILER_ATT_TOKEN_MASK2, 0xFFFF}); + + // create a session + CheckApi(rocprofiler_create_session(ROCPROFILER_KERNEL_REPLAY_MODE, &session_id)); + + // create a buffer to hold att trace records for each kernel launch + rocprofiler_buffer_id_t buffer_id; + CheckApi(rocprofiler_create_buffer(session_id, FlushCallback, 0x9999, &buffer_id)); + + // create a filter for collecting att traces + rocprofiler_filter_id_t filter_id; + rocprofiler_filter_property_t property = {}; + CheckApi(rocprofiler_create_filter(session_id, ROCPROFILER_ATT_TRACE_COLLECTION, + rocprofiler_filter_data_t{.att_parameters = ¶meters[0]}, + parameters.size(), &filter_id, property)); + + // set buffer for the filter + CheckApi(rocprofiler_set_filter_buffer(session_id, filter_id, buffer_id)); + + // activating att tracing session + CheckApi(rocprofiler_start_session(session_id)); + + // Launch a kernel + LaunchVectorAddKernel(); + + // deactivate att tracing session + CheckApi(rocprofiler_terminate_session(session_id)); + + // dump att tracing data + CheckApi(rocprofiler_flush_data(session_id, buffer_id)); + + // destroy session + CheckApi(rocprofiler_destroy_session(session_id)); + + // finalize att tracing by destroying rocmtool object + CheckApi(rocprofiler_finalize()); + return 0; +} diff --git a/tests/unittests/core/CMakeLists.txt b/tests/unittests/core/CMakeLists.txt index e4a97c3b7c..5efc625421 100644 --- a/tests/unittests/core/CMakeLists.txt +++ b/tests/unittests/core/CMakeLists.txt @@ -43,10 +43,11 @@ file(GLOB ROCPROFILER_SRC_API_FILES ${PROJECT_SOURCE_DIR}/src/api/*.cpp) file(GLOB ROCPROFILER_SRC_PROFILER_FILES ${PROJECT_SOURCE_DIR}/src/core/session/profiler/profiler.cpp) file(GLOB ROCPROFILER_TRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/*.cpp) file(GLOB ROCPROFILER_ROCTRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/src/*.cpp) +file(GLOB ROCPROFILER_ATT_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/att/att.cpp) file(GLOB ROCPROFILER_SRC_CLASS_FILES ${CMAKE_CURRENT_SOURCE_DIR}/rocmtool.cpp) file(GLOB ROCPROFILER_SPM_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/spm/spm.cpp) -set(ROCPROFILER_SRC_FILES ${ROCPROFILER_SRC_API_FILES} ${ROCPROFILER_SRC_CLASS_FILES} ${ROCPROFILER_SRC_PROFILER_FILES}) +set(ROCPROFILER_SRC_FILES ${ROCPROFILER_SRC_API_FILES} ${ROCPROFILER_SRC_CLASS_FILES} ${ROCPROFILER_SRC_PROFILER_FILES} ${ROCPROFILER_ATT_SRC_FILES}) set(CORE_HSA_DIR ${PROJECT_SOURCE_DIR}/src/core/hsa) file(GLOB CORE_HSA_SRC_FILES ${CORE_HSA_DIR}/*.cpp) diff --git a/tests/unittests/profiler/CMakeLists.txt b/tests/unittests/profiler/CMakeLists.txt index 8399df41cc..7e08ace63e 100644 --- a/tests/unittests/profiler/CMakeLists.txt +++ b/tests/unittests/profiler/CMakeLists.txt @@ -31,11 +31,12 @@ file(GLOB CORE_COUNTERS_SRC_FILES ${PROJECT_BINARY_DIR}/src/api/*_counter.cpp) file(GLOB ROCPROFILER_SRC_PROFILER_FILES ${PROJECT_SOURCE_DIR}/src/core/session/profiler/profiler.cpp) file(GLOB ROCPROFILER_TRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/*.cpp) file(GLOB ROCPROFILER_ROCTRACER_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/tracer/src/*.cpp) +file(GLOB ROCPROFILER_ATT_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/att/att.cpp) file(GLOB ROCPROFILER_SRC_CLASS_FILES ${CMAKE_CURRENT_SOURCE_DIR}/rocmtool.cpp) file(GLOB ROCPROFILER_SPM_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/spm/spm.cpp) file(GLOB ROCPROFILER_SRC_API_FILES ${PROJECT_SOURCE_DIR}/src/api/*.cpp) -set(ROCPROFILER_SRC_FILES ${ROCPROFILER_SRC_API_FILES} ${ROCPROFILER_SRC_CLASS_FILES} ${ROCPROFILER_SRC_PROFILER_FILES}) +set(ROCPROFILER_SRC_FILES ${ROCPROFILER_SRC_API_FILES} ${ROCPROFILER_SRC_CLASS_FILES} ${ROCPROFILER_SRC_PROFILER_FILES} ${ROCPROFILER_ATT_SRC_FILES}) set(CORE_HSA_DIR ${PROJECT_SOURCE_DIR}/src/core/hsa) file(GLOB CORE_HSA_SRC_FILES ${CORE_HSA_DIR}/*.cpp)