From a5555ac45f8feefc7619d99722c8473851c0910b Mon Sep 17 00:00:00 2001 From: Giovanni LB Date: Thu, 13 Jul 2023 03:46:03 -0300 Subject: [PATCH] SWDEV-396239: Automatic ISA dump from ATT Change-Id: Ia66346c0048b779487157961ead08d7567c9153a --- CHANGELOG.md | 2 + README.md | 86 +++--- include/rocprofiler/v2/rocprofiler.h | 189 ++++++++++++- plugin/att/CMakeLists.txt | 6 +- plugin/att/att.cpp | 115 +++++--- plugin/att/att.py | 59 +++-- plugin/att/att_to_csv.py | 18 ++ plugin/att/code_printing.cpp | 226 ++++++++++++++++ plugin/att/code_printing.hpp | 56 ++++ plugin/att/disassembly.cpp | 228 ++++++++++++++++ plugin/att/disassembly.hpp | 59 +++++ plugin/att/drawing.py | 2 + plugin/att/stitch.py | 153 ++++++++--- plugin/att/trace_view.py | 3 +- src/api/CMakeLists.txt | 11 +- src/api/exportmap | 5 + src/api/rocprofilerv2.cpp | 54 ++++ src/core/hsa/hsa_support.cpp | 11 + src/core/hsa/queues/queue.cpp | 84 ++++-- src/core/isa_capture/code_object_track.cpp | 248 ++++++++++++++++++ src/core/isa_capture/code_object_track.hpp | 133 ++++++++++ src/core/session/att/att.cpp | 2 - src/tools/tool.cpp | 4 +- .../featuretests/profiler/profiler_gtest.cpp | 140 ++++++++++ tests-v2/unittests/CMakeLists.txt | 11 +- 25 files changed, 1723 insertions(+), 182 deletions(-) create mode 100755 plugin/att/att_to_csv.py create mode 100644 plugin/att/code_printing.cpp create mode 100644 plugin/att/code_printing.hpp create mode 100644 plugin/att/disassembly.cpp create mode 100644 plugin/att/disassembly.hpp create mode 100644 src/core/isa_capture/code_object_track.cpp create mode 100644 src/core/isa_capture/code_object_track.hpp diff --git a/CHANGELOG.md b/CHANGELOG.md index 908c83eb6d..0f47812af6 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -299,3 +299,5 @@ Example for file plugin output: ### Added - Updated supported GPU architectures in README with profiler versions +- Automatic ISA dumping for ATT +- CSV mode for ATT diff --git a/README.md b/README.md index 3b836c2bb3..af9c0962c8 100644 --- a/README.md +++ b/README.md @@ -306,56 +306,32 @@ Usage: - ATT (Advanced thread tracer) plugin: advanced hardware traces data in binary format. Please refer ATT section. Tool used to collect fine-grained hardware metrics. Provides ISA-level instruction hotspot analysis via hardware tracing. -*** -Note: ATT(Advanced Thread Trace) needs some preparation before running. -*** + - Install plugin package. See Plugin Support section for installation + - Run the following to view the trace. Att-specific options must come right after the assembly file -- Make sure to generate the assembly file for application by executing the following before compiling your HIP Application. This can be achieved globally by following environment variable + ```bash + rocprofv2 -i input.txt --plugin att --mode network + ``` - ```bash - export HIPCC_COMPILE_FLAGS_APPEND="--save-temps -g" - ``` - - Similarly, the --save-temps -g flags can be added per file for better ISA generation control. - -- Install plugin package. See Plugin Support section for installation -- Run the following to view the trace. Att-specific options must come right after the assembly file - - ```bash - rocprofv2 -i input.txt --plugin att --mode network - ``` - -- Example for vectoradd on navi31. - - *** - Note: Special attention to gfx1100.s==navi31 in the ISA file name. - Use gfx1030 for navi21, gfx90a for MI200 and gfx940 for MI300 - *** - - ```bash - hipcc -g --save-temps vectoradd_hip.cpp -o vectoradd_hip.exe - rocprofv2 -i input.txt --plugin att vectoradd_hip-hip-amdgcn-amd-amdhsa-gfx1100.s --mode network ./vectoradd_hip.exe - ``` - - Then open the browser at `http://localhost:8000` - The ISA can also be obtained from llvm/roc objdump, however, annotations will be different - - - app_assembly_file_relative_path - AMDGCN ISA file with .s extension generated in 1st step + - app_assembly_file: + On ROCm 6.0, ATT enables automatic capture of the ISA during kernel execution, and does not require recompiling. It is recommeneded to leave at "auto". - app_relative_path - Path for the running application + Path for the running application - ATT plugin optional parameters - --depth [n]: How many waves per slot to parse (maximum). - --mpi [proc]: Parse with this many mpi processes, for greater analysis speed. Does not change results. Requires mpi4py. - --att_kernel "filename": Kernel filename to use (instead of ATT asking which one to use). - --trace_file "files": glob (wildcards allowed) of traces files to parse. Requires quotes for use with wildcards. - - --mode [network, file, off (default)] + - --mode [network, file, csv, off (default)] - network Opens the server with the browser UI. att needs 2 ports available (e.g. 8000, 18000). There is an option (default: --ports "8000,18000") to change these. In case rocprofv2 is running on a different machine, use port forwarding `ssh -L 8000:localhost:8000 ` so the browser can be used locally. For docker, use --network=host --ipc=host -p8000:8000 -p18000:18000 - file Dumps the analyzed json files to disk for vieweing at a later time. Run python3 httpserver.py from within the generated ui/ folder to view the trace, similarly to network mode. The folder can be copied to another machine, and will run without rocm. + - csv + Dumps the analyzed assembly into a CSV format, with the hitcount and total cycles cost. + Use rocprofv2's -o option to specify output file name (default "att_output.csv"). - off Runs trace collection but not analysis, so it can be analyzed at a later time. Run rocprofv2 ATT [network, file] with the same parameters, removing the application binary, to analyze previously generated traces. - input.txt @@ -385,23 +361,47 @@ Note: ATT(Advanced Thread Trace) needs some preparation before running. - PERFCOUNTERS_COL_PERIOD=0x3 // Multiplier period for counter collection [0~31]. 0=fastest (usually once every 16 cycles). GFX9 only. Counters will be shown in a graph over time in the browser UI. - PERFCOUNTER=counter_name // Add a SQ counter to be collected with ATT; period defined by PERFCOUNTERS_COL_PERIOD. GFX9 only. - BUFFER_SIZE=[size] // Sets size of the ATT buffer collection, per dispatch, in megabytes (shared among all shader engines). + - ISA_CAPTURE_MODE=[0,1,2] // Set capture mode during kernel dispatch. + - 0 = capture symbols only. + - 1 = capture symbols for file:// and make a copy of memory:// + - 2 = Copy file:// and memory:// + + - Example for vectoradd. + + ```bash + # -g add debugging symbols to the binary. Required only for tracking disassembly back to c++. + hipcc -g --save-temps vectoradd_hip.cpp -o vectoradd_hip.exe + # "auto" means to use the automatically captured ISA. "csv" dumps result to "my.csv". + rocprofv2 -i input.txt -o my.csv --plugin att auto --mode csv ./vectoradd_hip.exe + ``` + Instruction latencies will be in my.csv + + - Example with symbolic ISA (as in ROCm 5.7 or previous). + ```bash + hipcc -g --save-temps vectoradd_hip.cpp -o vectoradd_hip.exe + # A custom ISA can be used such as vectoradd_hip-hip-amdgcn-amd-amdhsa-gfx1100.s + # Special attention to the correct architecture for the ISA, such as "gfx1100" (navi31). + rocprofv2 -i input.txt --plugin att vectoradd_hip-hip-amdgcn-amd-amdhsa-gfx1100.s --mode network ./vectoradd_hip.exe + ``` + + Then open the browser at `http://localhost:8000` *** Note: For MPI or long running applications, we recommend to run collection, and later run the parser with already collected data: - Run only collection: The assembly file is not used. Use mpirun [...] rocprofv2 [...] if needed. + Run only collection: The assembly file is not used. Use mpirun [...] rocprofv2 [...] if needed. ```bash - rocprofv2 -i input.txt --plugin att none ./vectoradd_hip.exe + # Run only data collection, not the parser + rocprofv2 -i input.txt --plugin att auto --mode off ./vectoradd_hip.exe ``` - - Remove the binary/application: Only runs the parser. - + Remove the binary/application from the command line. ```bash - rocprofv2 -i input.txt --plugin att vectoradd_hip-hip-amdgcn-amd-amdhsa-gfx1100.s --mode network + # Only runs the parser on previously collected data. + rocprofv2 -i input.txt --plugin att auto --mode network ``` *** - + ### Flush Interval Flush interval can be used to control the interval time in milliseconds between the buffers flush for the tool. However, if the buffers are full the flush will be called on its own. This can be used as in the next example: diff --git a/include/rocprofiler/v2/rocprofiler.h b/include/rocprofiler/v2/rocprofiler.h index 4fc5d2e4b2..0e373969ed 100644 --- a/include/rocprofiler/v2/rocprofiler.h +++ b/include/rocprofiler/v2/rocprofiler.h @@ -1173,6 +1173,75 @@ typedef struct { uint32_t buffer_size; } rocprofiler_record_se_att_data_t; +/** + * struct to store the filepaths and their addresses for intercepted code objects +*/ +typedef struct { + /** + * File path (file://, memory://) of the code object + */ + const char* filepath; + /** + * Addr where codeobj is loaded + */ + uint64_t base_address; + /** + * If a copy of the codeobj is made, contains the data. Nullptr otherwise. + */ + const char* data; + /** + * If a copy of the codeobj is made, contains the size of the data. 0 otherwise. + */ + uint64_t size; + /** + * Timestamp for the time point this codeobj was loaded. + */ + rocprofiler_timestamp_t clock_start; + /** + * Timestamp for the time point this codeobj was unloaded. + * If the obj is still loaded by the time the record was generated, this value is 0 + */ + rocprofiler_timestamp_t clock_end; +} rocprofiler_intercepted_codeobj_t; + +/** + * Enum defines how code object is captured for ATT and PC Sampling +*/ +typedef enum { + /** + * Capture file and memory paths for the loaded code object + */ + ROCPROFILER_CAPTURE_SYMBOLS_ONLY = 0, + /** + * Capture symbols for file:// and memory:// type objects, + * and generate a copy of all kernel code for objects under memory:// + */ + ROCPROFILER_CAPTURE_COPY_MEMORY = 1, + /** + * Capture symbols and all kernel code for file:// and memory:// type objects + */ + ROCPROFILER_CAPTURE_COPY_FILE_AND_MEMORY = 2 +} rocprofiler_codeobj_capture_mode_t; + +/** + * struct to store the filepaths and their addresses for intercepted code objects +*/ +typedef struct { + /** + * List of symbols + */ + const rocprofiler_intercepted_codeobj_t* symbols; + /** + * Number of symbols + */ + uint64_t count; + /** + * Userdata space for custom capture. + * For ATT records, it is the address of the kernel being launched. + */ + uint64_t userdata; +} rocprofiler_codeobj_symbols_t; + /** * ATT tracing record structure. * This will represent all the information reported by the @@ -1214,6 +1283,10 @@ typedef struct { * Queue Index - packet index in the queue */ rocprofiler_queue_index_t queue_idx; + /** + * Writer ID for counting how many kernels + */ + uint64_t writer_id; /** * ATT data output from each shader engine. */ @@ -1222,6 +1295,10 @@ typedef struct { * The count of the shader engine ATT data */ uint64_t shader_engine_data_count; + /** + * Filepaths for the intercepted code objects at the time of kernel dispatch + */ + rocprofiler_codeobj_symbols_t intercept_list; } rocprofiler_record_att_tracer_t; @@ -1650,7 +1727,7 @@ typedef enum { */ ROCPROFILER_PC_SAMPLING_COLLECTION = 3, /** - * ATT Tracing. (Not Yet Supported) + * ATT Tracing. */ ROCPROFILER_ATT_TRACE_COLLECTION = 4, /** @@ -1707,16 +1784,50 @@ typedef const char* rocprofiler_hsa_function_name_t; * ATT parameters to be used by for collection */ typedef enum { - ROCPROFILER_ATT_COMPUTE_UNIT = 0, //! Select the target compute unit (wgp) for profiling. - ROCPROFILER_ATT_VMID_MASK = 1, //! VMID Mask. - ROCPROFILER_ATT_SE_MASK = 5, //! Shader engine mask for selection. - ROCPROFILER_ATT_SIMD_SELECT = 8, //! Set SIMD Mask (GFX9) or SIMD ID for collection (Navi) - ROCPROFILER_ATT_OCCUPANCY = 9, //! Set true for occupancy collection only. - ROCPROFILER_ATT_BUFFER_SIZE = 10, //! ATT collection max data size, in MB. Shared among shader engines. - ROCPROFILER_ATT_PERF_MASK = 240, //! Mask of which compute units to generate perfcounters. GFX9 only. - ROCPROFILER_ATT_PERF_CTRL = 241, //! Select collection period for perfcounters. GFX9 only. - ROCPROFILER_ATT_PERFCOUNTER = 242, //! Select perfcounter ID (SQ block) for collection. GFX9 only. - ROCPROFILER_ATT_PERFCOUNTER_NAME = 243, //! Select perfcounter name (SQ block) for collection. GFX9 only. + /** + * Select the target compute unit (wgp) for profiling. + */ + ROCPROFILER_ATT_COMPUTE_UNIT = 0, + /** + * VMID Mask + */ + ROCPROFILER_ATT_VMID_MASK = 1, + /** + * Shader engine mask for selection. + */ + ROCPROFILER_ATT_SE_MASK = 5, + /** + * Set SIMD Mask (GFX9) or SIMD ID for collection (Navi) + */ + ROCPROFILER_ATT_SIMD_SELECT = 8, + /** + * Set true for occupancy collection only. + */ + ROCPROFILER_ATT_OCCUPANCY = 9, + /** + * ATT collection max data size, in MB. Shared among shader engines. + */ + ROCPROFILER_ATT_BUFFER_SIZE = 10, + /** + * Set ISA capture during ATT collection (rocprofiler_codeobj_capture_mode_t) + */ + ROCPROFILER_ATT_CAPTURE_MODE = 11, + /** + * Mask of which compute units to generate perfcounters. GFX9 only. + */ + ROCPROFILER_ATT_PERF_MASK = 240, + /** + * Select collection period for perfcounters. GFX9 only. + */ + ROCPROFILER_ATT_PERF_CTRL = 241, + /** + * Select perfcounter ID (SQ block) for collection. GFX9 only. + */ + ROCPROFILER_ATT_PERFCOUNTER = 242, + /** + * Select perfcounter name (SQ block) for collection. GFX9 only. + */ + ROCPROFILER_ATT_PERFCOUNTER_NAME = 243, ROCPROFILER_ATT_MAXVALUE, ROCPROFILER_ATT_MASK = 2, //! Deprecated @@ -2230,6 +2341,62 @@ ROCPROFILER_API rocprofiler_status_t rocprofiler_device_profiling_session_stop( ROCPROFILER_API rocprofiler_status_t rocprofiler_device_profiling_session_destroy( rocprofiler_session_id_t session_id) ROCPROFILER_VERSION_9_0; +/** + * Creates a codeobj capture record, returned in ID. + * \param[out] id contains a handle for the created record. + * \param[in] mode Set to capture symbols only, + * make a copy of codeobj under memory:// + * or copy all codeobj. + * \param[in] userdata userdata to be returned in the record. For ATT records, is the kernel addr. + * \retval ::ROCPROFILER_STATUS_SUCCESS The function has been executed successfully. + */ +ROCPROFILER_API rocprofiler_status_t +rocprofiler_codeobj_capture_create( + rocprofiler_record_id_t* id, + rocprofiler_codeobj_capture_mode_t mode, + uint64_t userdata +); + +/** + * API to get the captured codeobj. + * Each call invalidates the previous pointer for the same ID. + * \param[in] id record handle. + * \param[out] capture captured code objects. + * \retval ::ROCPROFILER_STATUS_SUCCESS The function has been executed successfully. + * \retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS invalid ID. + */ +ROCPROFILER_API rocprofiler_status_t +rocprofiler_codeobj_capture_get(rocprofiler_record_id_t id, + rocprofiler_codeobj_symbols_t* capture); + +/** + * API to delete a record. + * Invalidates the pointer returned from rocprofiler_codeobj_capture_get. + * \param[in] id record handle. + * \retval ::ROCPROFILER_STATUS_SUCCESS The function has been executed successfully. + */ +ROCPROFILER_API rocprofiler_status_t +rocprofiler_codeobj_capture_free(rocprofiler_record_id_t id); + +/** + * Records the current loaded codeobjs and any following loads until stop() is called. + * \param[in] id record handle. + * \retval ::ROCPROFILER_STATUS_SUCCESS The function has been executed successfully. + * \retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS invalid ID. + */ +ROCPROFILER_API rocprofiler_status_t +rocprofiler_codeobj_capture_start(rocprofiler_record_id_t id); + +/** + * Stops recording of future codeobjs, until start() is called again. + * Calling stop() immediately after a start() snapshots the current state of loaded codeobjs. + * \param[in] id record handle. + * \retval ::ROCPROFILER_STATUS_SUCCESS The function has been executed successfully. + * \retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS invalid ID. + */ +ROCPROFILER_API rocprofiler_status_t +rocprofiler_codeobj_capture_stop(rocprofiler_record_id_t id); + /** @} */ #ifdef __cplusplus diff --git a/plugin/att/CMakeLists.txt b/plugin/att/CMakeLists.txt index bb8892fc39..f4af4f6795 100644 --- a/plugin/att/CMakeLists.txt +++ b/plugin/att/CMakeLists.txt @@ -26,7 +26,7 @@ set(ENV{ROCPROFV2_ATT_LIB_PATH} $ROCPROFV2_ATT) # Building att plugin library file(GLOB ROCPROFILER_UTIL_SRC_FILES ${PROJECT_SOURCE_DIR}/src/utils/helper.cpp) -file(GLOB FILE_SOURCES att.cpp) +file(GLOB FILE_SOURCES att.cpp disassembly.cpp code_printing.cpp) add_library(att_plugin SHARED ${FILE_SOURCES} ${ROCPROFILER_UTIL_SRC_FILES}) set_target_properties( @@ -44,8 +44,7 @@ target_include_directories(att_plugin PRIVATE ${PROJECT_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-v2 hsa-runtime64::hsa-runtime64 - stdc++fs) +target_link_libraries(att_plugin PRIVATE rocprofiler-v2 hsa-runtime64::hsa-runtime64 stdc++fs dw elf amd_comgr) install(TARGETS att_plugin LIBRARY DESTINATION ${CMAKE_INSTALL_LIBDIR}/${PROJECT_NAME} COMPONENT asan) @@ -56,6 +55,7 @@ configure_file(att.py att/att.py COPYONLY) configure_file(trace_view.py att/trace_view.py COPYONLY) configure_file(stitch.py att/stitch.py COPYONLY) configure_file(drawing.py att/drawing.py COPYONLY) +configure_file(att_to_csv.py att/att_to_csv.py 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) diff --git a/plugin/att/att.cpp b/plugin/att/att.cpp index f5aadba747..a11b3a8a12 100644 --- a/plugin/att/att.cpp +++ b/plugin/att/att.cpp @@ -44,8 +44,11 @@ #include "rocprofiler.h" #include "rocprofiler_plugin.h" #include "../utils.h" +#include "code_printing.hpp" + #define ATT_FILENAME_MAXBYTES 90 +#define TEST_INVALID_KERNEL size_t(-1) namespace { @@ -74,24 +77,33 @@ class att_plugin_t { 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) { + int 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; + if (!att_tracer_record) return ROCPROFILER_STATUS_ERROR; + + std::string kernel_name_mangled{}; + // Found problem with rocprofiler API for invalid kernel_ids; + if (att_tracer_record->kernel_id.handle != TEST_INVALID_KERNEL) { + size_t name_length; + CHECK_ROCPROFILER(rocprofiler_query_kernel_info_size( + ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, &name_length)); + const char* kernel_name_c = nullptr; + CHECK_ROCPROFILER(rocprofiler_query_kernel_info( + ROCPROFILER_KERNEL_NAME, att_tracer_record->kernel_id, &kernel_name_c)); + + assert(kernel_name_c && "Rocprofv2 returned an invalid kernel name"); + + kernel_name_mangled = std::string(kernel_name_c); + free(const_cast(kernel_name_c)); + } else { // Temporary. Adding a valid string. + kernel_name_mangled = "test_kernel"; } - size_t name_length; - CHECK_ROCPROFILER(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_ROCPROFILER(rocprofiler_query_kernel_info(ROCPROFILER_KERNEL_NAME, - att_tracer_record->kernel_id, &kernel_name_c)); std::string name_demangled = - rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_c)); + rocprofiler::truncate_name(rocprofiler::cxx_demangle(kernel_name_mangled)); if (name_demangled.size() > ATT_FILENAME_MAXBYTES) // Limit filename size name_demangled = name_demangled.substr(0, ATT_FILENAME_MAXBYTES); @@ -110,11 +122,11 @@ class att_plugin_t { file_iteration += 1; outfilepath += std::to_string(file_iteration); - auto dispatch_id = att_tracer_record->header.id.handle; + auto writer_id = att_tracer_record->writer_id; std::string fname = outfilepath + "_kernel.txt"; - std::ofstream(fname.c_str()) << name_demangled << " dispatch[" << dispatch_id << "] GPU[" - << att_tracer_record->gpu_id.handle << "]: " << kernel_name_c + std::ofstream(fname.c_str()) << name_demangled << " dispatch[" << writer_id << "] GPU[" + << att_tracer_record->gpu_id.handle << "]: " << kernel_name_mangled << '\n'; // iterate over each shader engine att trace @@ -125,29 +137,69 @@ class att_plugin_t { 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]; - const char* data_buffer_ptr = reinterpret_cast(se_att_trace->buffer_ptr); + char* data_buffer_ptr = reinterpret_cast(se_att_trace->buffer_ptr); // dump data in binary format std::ofstream out(outfilepath + "_se" + std::to_string(i) + ".att", std::ios::binary); - if (out.is_open()) - out.write((char*)data_buffer_ptr, se_att_trace->buffer_size); - else + if (!out.is_open()) { std::cerr << "ATT Failed to open file: " << outfilepath << "_se" << i << ".att\n"; + return ROCPROFILER_STATUS_ERROR; + } + out.write(data_buffer_ptr, se_att_trace->buffer_size); } + + std::ofstream isafile(outfilepath + "_isa.s"); + if (!isafile.is_open()) { + std::cerr << "Could not open ISA file: " << outfilepath << "_isa.s" << std::endl; + return ROCPROFILER_STATUS_ERROR; + } + uint64_t kernel_begin_addr = att_tracer_record->intercept_list.userdata; + isafile << " " << kernel_name_mangled << '\n'; + + for (size_t i = 0; i < att_tracer_record->intercept_list.count; i++) { + const rocprofiler_intercepted_codeobj_t& symbol = + att_tracer_record->intercept_list.symbols[i]; + + std::unique_ptr binary; + std::unique_ptr decoder; + + if (symbol.data && symbol.size) { + decoder = std::make_unique(symbol.data, symbol.size); + } else if (std::string(symbol.filepath).find("file://") != std::string::npos) { + binary = std::make_unique(symbol.filepath); + decoder = + std::make_unique(binary->buffer.data(), binary->buffer.size()); + } else { + continue; + } + + for (auto& instance : decoder->instructions) { + uint64_t addr = instance.address + symbol.base_address; + + if (kernel_begin_addr == addr) + isafile << "; Begin " << kernel_name_mangled << '\n'; + else if (decoder->m_symbol_map.find(instance.address) != decoder->m_symbol_map.end()) + isafile << "; Begin " << decoder->m_symbol_map[instance.address].first << '\n'; + if (instance.cpp_reference) isafile << "; " << instance.cpp_reference << '\n'; + isafile << instance.instruction << " // " << std::hex << addr << '\n'; + } + } + + return ROCPROFILER_STATUS_SUCCESS; } 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; + if (!begin) return ROCPROFILER_STATUS_ERROR; switch (begin->kind) { case ROCPROFILER_PROFILER_RECORD: case ROCPROFILER_TRACER_RECORD: case ROCPROFILER_PC_SAMPLING_RECORD: case ROCPROFILER_SPM_RECORD: case ROCPROFILER_COUNTERS_SAMPLER_RECORD: - printf("Invalid record Kind: %d", begin->kind); + rocprofiler::warning("Invalid record Kind: %d\n", begin->kind); break; case ROCPROFILER_ATT_TRACER_RECORD: { @@ -158,10 +210,11 @@ class att_plugin_t { break; } } - rocprofiler_next_record(begin, &begin, session_id, buffer_id); + int status = rocprofiler_next_record(begin, &begin, session_id, buffer_id); + if (status != ROCPROFILER_STATUS_SUCCESS) return status; } - return 0; + return ROCPROFILER_STATUS_SUCCESS; } private: @@ -176,17 +229,17 @@ ROCPROFILER_EXPORT int rocprofiler_plugin_initialize(uint32_t rocprofiler_major_ void* data) { if (rocprofiler_major_version != ROCPROFILER_VERSION_MAJOR || rocprofiler_minor_version < ROCPROFILER_VERSION_MINOR) - return -1; + return ROCPROFILER_STATUS_ERROR; - if (att_plugin != nullptr) return -1; + if (att_plugin != nullptr) return ROCPROFILER_STATUS_ERROR; att_plugin = new att_plugin_t(); - if (att_plugin->IsValid()) return 0; + if (att_plugin->IsValid()) return ROCPROFILER_STATUS_SUCCESS; // The plugin failed to initialied, destroy it and return an error. delete att_plugin; att_plugin = nullptr; - return -1; + return ROCPROFILER_STATUS_ERROR; } ROCPROFILER_EXPORT void rocprofiler_plugin_finalize() { @@ -198,12 +251,12 @@ ROCPROFILER_EXPORT void rocprofiler_plugin_finalize() { 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; + if (!att_plugin || !att_plugin->IsValid()) return ROCPROFILER_STATUS_ERROR; return att_plugin->WriteBufferRecords(begin, end, session_id, buffer_id); } ROCPROFILER_EXPORT int rocprofiler_plugin_write_record(rocprofiler_record_tracer_t record) { - if (!att_plugin || !att_plugin->IsValid()) return -1; - if (record.header.id.handle == 0) return 0; - return 0; + if (!att_plugin || !att_plugin->IsValid()) return ROCPROFILER_STATUS_ERROR; + if (record.header.id.handle == 0) return ROCPROFILER_STATUS_SUCCESS; + return ROCPROFILER_STATUS_SUCCESS; } diff --git a/plugin/att/att.py b/plugin/att/att.py index 3d35ed011f..cf1778f980 100755 --- a/plugin/att/att.py +++ b/plugin/att/att.py @@ -40,25 +40,26 @@ class PerfEvent(ctypes.Structure): 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)] + ('loc', ctypes.c_char_p), + ('to_line', ctypes.c_int), + ('value', ctypes.c_int), + ('index', ctypes.c_int), + ('line_num', ctypes.c_int), + ('addr', ctypes.c_int64)] class KvPair(ctypes.Structure): """ Matches pair = (key, value) on the python side """ _fields_ = [('key', ctypes.c_int), - ('value', 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)] + ('jumps', POINTER(KvPair)), + ('code_len', ctypes.c_int), + ('jumps_len', ctypes.c_int)] class Wave(ctypes.Structure): @@ -155,8 +156,9 @@ def parse_binary(filename, kernel=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), 0, 0]) # hitcount + cycles + # asm, inst_type, addr, loc, index, line_num, hitcount, cycles + code.append([line, int(code_entry.value), to_line, loc, int(code_entry.index), + int(code_entry.line_num), int(code_entry.addr), 0, 0]) jumps = {} for k in range(info.jumps_len): @@ -187,9 +189,9 @@ def getWaves_binary(name, shader_engine_data_dict, target_cu, depth): shader_engine_data_dict[name] = (waves_python, events, occupancy, flags) -def getWaves_stitch(SIMD, code, jumps, flags, latency_map, hitcount_map): +def getWaves_stitch(SIMD, code, jumps, flags, latency_map, hitcount_map, bIsAuto): for pwave in SIMD: - pwave.instructions = stitch(pwave.instructions, code, jumps, flags) + pwave.instructions = stitch(pwave.instructions, code, jumps, flags, bIsAuto) for inst in pwave.instructions[0]: hitcount_map[inst[-1]] += 1 @@ -363,7 +365,10 @@ if __name__ == "__main__": network: Open att server over the network.''', type=str, default="off") args = parser.parse_args() - if args.mode.lower() == 'file': + CSV_MODE = False + if args.mode.lower() == 'csv': + CSV_MODE = True + elif args.mode.lower() == 'file': args.dumpfiles = True elif args.mode.lower() == 'network': args.dumpfiles = False @@ -387,12 +392,6 @@ if __name__ == "__main__": 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: @@ -418,6 +417,16 @@ if __name__ == "__main__": else: args.att_kernel = att_kernel[0] + # Assembly parsing + bIsAuto = False + if args.assembly_code.lower().strip() == 'auto': + args.assembly_code = args.att_kernel.split('_kernel.txt')[0]+'_isa.s' + bIsAuto = True + path = Path(args.assembly_code) + if not path.is_file(): + print("Invalid assembly_code('{0}')!".format(args.assembly_code)) + sys.exit(1) + # Trace Parsing if args.trace_file is None: filenames = glob.glob(args.att_kernel.split('_kernel.txt')[0]+'_*.att') @@ -431,7 +440,7 @@ if __name__ == "__main__": code = jumps = None if mpi_root: print('Att kernel:', args.att_kernel) - code, jumps = parse_binary(args.assembly_code, args.att_kernel) + code, jumps = parse_binary(args.assembly_code, None if bIsAuto else args.att_kernel) DBFILES = [] TIMELINES = [np.zeros(int(1E4),dtype=np.int16) for k in range(5)] @@ -460,7 +469,7 @@ if __name__ == "__main__": if np.sum([0]+[len(s.instructions) for s in SIMD]) == 0: print("No waves from", name) continue - getWaves_stitch(SIMD, code, jumps, gfxv, latency_map, hitcount_map) + getWaves_stitch(SIMD, code, jumps, gfxv, latency_map, hitcount_map, bIsAuto) analysed_filenames.append(name) EVENTS.append(perfevents) @@ -524,6 +533,12 @@ if __name__ == "__main__": code[k][-2] = int(hitcount_map[k]) code[k][-1] = int(latency_map[k]) + if CSV_MODE: + if mpi_root: + from att_to_csv import dump_csv + dump_csv(code) + quit() + gc.collect() print("Min time:", min_event_time) diff --git a/plugin/att/att_to_csv.py b/plugin/att/att_to_csv.py new file mode 100755 index 0000000000..fbd12aa0b8 --- /dev/null +++ b/plugin/att/att_to_csv.py @@ -0,0 +1,18 @@ +#!/usr/bin/env python3 + +import numpy as np +import csv +import os + +def dump_csv(code): + outpath = os.getenv("OUT_FILE_NAME") + if outpath is None: + outpath = "att_output.csv" + if ".csv" not in outpath: + outpath += ".csv" + + with open(outpath, 'w') as f: + writer = csv.writer(f) + writer.writerow(['Line', 'Instruction', 'Hitcount', 'Cycles', 'Addr', 'C++ Reference']) + [writer.writerow([m[5], m[0], m[7], m[8], hex(m[6]), m[3]]) for m in code] + #[writer.writerow(m) for m in code] diff --git a/plugin/att/code_printing.cpp b/plugin/att/code_printing.cpp new file mode 100644 index 0000000000..aff5259a99 --- /dev/null +++ b/plugin/att/code_printing.cpp @@ -0,0 +1,226 @@ +/* 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 "code_printing.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +#include +#include +#include +#include + +#include +#include "../utils.h" +#include +#include +#include + +code_object_decoder_t::code_object_decoder_t(const char* codeobj_data, uint64_t codeobj_size) { + buffer = std::vector{}; + buffer.resize(codeobj_size); + std::memcpy(buffer.data(), codeobj_data, codeobj_size); + + m_fd = -1; +#if defined(_GNU_SOURCE) && defined(MFD_ALLOW_SEALING) && defined(MFD_CLOEXEC) + m_fd = ::memfd_create(m_uri.c_str(), MFD_ALLOW_SEALING | MFD_CLOEXEC); +#endif + if (m_fd == -1) // If fail, attempt under /tmp + m_fd = ::open("/tmp", O_TMPFILE | O_RDWR, 0666); + + if (m_fd == -1) { + printf("could not create a temporary file for code object\n"); + return; + } + + if (size_t size = ::write(m_fd, buffer.data(), buffer.size()); size != buffer.size()) { + printf("could not write to the temporary file\n"); + return; + } + ::lseek(m_fd, 0, SEEK_SET); + fsync(m_fd); + + m_line_number_map = {}; + + std::unique_ptr dbg(dwarf_begin(m_fd, DWARF_C_READ), + [](Dwarf* dbg) { dwarf_end(dbg); }); + + /*if (!dbg) { + rocprofiler::warning("Error opening Dwarf!\n"); + return; + } */ + + if (dbg) { + Dwarf_Off cu_offset{0}, next_offset; + size_t header_size; + + while (!dwarf_nextcu(dbg.get(), cu_offset, &next_offset, &header_size, nullptr, nullptr, + nullptr)) { + Dwarf_Die die; + if (!dwarf_offdie(dbg.get(), cu_offset + header_size, &die)) continue; + + Dwarf_Lines* lines; + size_t line_count; + if (dwarf_getsrclines(&die, &lines, &line_count)) continue; + + for (size_t i = 0; i < line_count; ++i) { + Dwarf_Addr addr; + int line_number; + + if (Dwarf_Line* line = dwarf_onesrcline(lines, i)) + if (!dwarf_lineaddr(line, &addr) && !dwarf_lineno(line, &line_number) && line_number) { + m_line_number_map.emplace( + addr, std::make_pair(dwarf_linesrc(line, nullptr, nullptr), line_number)); + } + } + cu_offset = next_offset; + } + // load_symbol_map(); + } + disassemble_kernels(); +} + + +code_object_decoder_t::~code_object_decoder_t() { + if (m_fd) ::close(m_fd); +} + +std::optional code_object_decoder_t::find_symbol( + uint64_t address) { + /* Load the symbol table. */ + if (auto it = m_symbol_map.upper_bound(address); it != m_symbol_map.begin()) { + if (auto&& [symbol_value, symbol] = *std::prev(it); address < (symbol_value + symbol.second)) { + std::string symbol_name = symbol.first; + + if (int status; auto* demangled_name = + abi::__cxa_demangle(symbol_name.c_str(), nullptr, nullptr, &status)) { + symbol_name = demangled_name; + free(demangled_name); + } + return symbol_info_t{std::move(symbol_name), symbol_value, symbol.second}; + } + } + return {}; +} + +/* +void code_object_decoder_t::load_symbol_map() { + std::unique_ptr elf ( + elf_begin(m_fd, ELF_C_READ, nullptr), + [](Elf *elf){ elf_end(elf); }); + + if (!elf) { + rocprofiler::warning("Error opening ELF!\n"); + return; + } + + Elf64_Ehdr *ehdr = elf64_getehdr(elf.get()); + if (!ehdr) { + printf("elf64_getehdr failed\n"); + return; + } + + // Slurp the symbol table. + Elf_Scn *scn = nullptr; + while ((scn = elf_nextscn(elf.get(), scn)) != nullptr) { + GElf_Shdr shdr_mem; + GElf_Shdr *shdr = gelf_getshdr(scn, &shdr_mem); + if (shdr->sh_type != SHT_SYMTAB && shdr->sh_type != SHT_DYNSYM) { + continue; + } + + Elf_Data *data = elf_getdata(scn, nullptr); + if (!data) continue; + + size_t symbol_count = data->d_size / gelf_fsize(elf.get(), ELF_T_SYM, 1, EV_CURRENT); + for (size_t j = 0; j < symbol_count; ++j) { + GElf_Sym sym_mem; + GElf_Sym *sym = gelf_getsym(data, j, &sym_mem); + + if (GELF_ST_TYPE(sym->st_info) != STT_FUNC || sym->st_shndx == SHN_UNDEF) continue; + + std::string symbol_name{ elf_strptr(elf.get(), shdr->sh_link, sym->st_name) }; + auto symbol_pair = std::make_pair(symbol_name, sym->st_size); + + auto [it, success] = m_symbol_map.emplace(sym->st_value, symbol_pair); + + // If there already was a symbol defined at this address, but this + // new symbol covers a larger address range, replace the old symbol + // with this new one. + if (!success && sym->st_size > it->second.second) it->second = symbol_pair; + } + } +} */ + +void code_object_decoder_t::disassemble_kernel(uint64_t addr) { + auto symbol = find_symbol(addr); + + if (!symbol) { + std::cerr << "No symbol found at address 0x" << std::hex << addr << std::endl; + return; + } + + // if (symbol->m_name.find("__amd_rocclr_") == 0) + // return; + + std::cout << "Dumping ISA for " << symbol->m_name << std::endl; + + uint64_t end_addr = addr + symbol->m_size; + while (addr < end_addr) { + char* cpp_line = nullptr; + auto it = m_line_number_map.find(addr); + if (it != m_line_number_map.end()) { + const std::string& file_name = it->second.first; + size_t line_number = it->second.second; + + std::string cpp = file_name + ':' + std::to_string(line_number); + cpp_line = (char*)calloc(cpp.size() + 4, sizeof(char)); + std::memcpy(cpp_line, cpp.data(), cpp.size() * sizeof(char)); + } + + size_t size = disassembly->ReadInstruction(addr, cpp_line); + addr += size; + } +} + +void code_object_decoder_t::disassemble_kernels() { + disassembly = std::make_unique(*this); + + // if (m_symbol_map.begin() == m_symbol_map.end()) + m_symbol_map = disassembly->GetKernelMap(); + + for (auto& [k, v] : m_symbol_map) disassemble_kernel(k); +} diff --git a/plugin/att/code_printing.hpp b/plugin/att/code_printing.hpp new file mode 100644 index 0000000000..050b219dff --- /dev/null +++ b/plugin/att/code_printing.hpp @@ -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. */ + +#pragma once + +#include +#include +#include +#include +#include "rocprofiler.h" +#include +#include "disassembly.hpp" + +class code_object_decoder_t { + struct symbol_info_t { + const std::string m_name; + uint64_t m_value; + uint64_t m_size; + }; + + public: + // void load_symbol_map(); + std::optional find_symbol(uint64_t address); + + code_object_decoder_t(const char* codeobj_data, uint64_t codeobj_size); + ~code_object_decoder_t(); + + void disassemble_kernel(uint64_t addr); + void disassemble_kernels(); + int m_fd; + + std::map> m_line_number_map; + std::map> m_symbol_map; + + std::string m_uri; + std::vector buffer; + std::vector instructions; + std::unique_ptr disassembly; +}; diff --git a/plugin/att/disassembly.cpp b/plugin/att/disassembly.cpp new file mode 100644 index 0000000000..d564719ed2 --- /dev/null +++ b/plugin/att/disassembly.cpp @@ -0,0 +1,228 @@ +/* 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. */ + +#if !defined(_GNU_SOURCE) || !defined(_XOPEN_SOURCE) +#define _XOPEN_SOURCE 700 +#endif + +#include "code_printing.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include +#include +#include + +#include +#include +#include +#include +#include + +#include +#include "../utils.h" +#include +#include + +#define CHECK_COMGR(call) \ + if (amd_comgr_status_s status = call) { \ + const char* reason = ""; \ + amd_comgr_status_string(status, &reason); \ + std::cerr << __LINE__ << " failed: " << reason << std::endl; \ + return; \ + } + +CodeObjectBinary::CodeObjectBinary(const std::string& uri) : m_uri(uri) { + const std::string protocol_delim{"://"}; + + size_t protocol_end = m_uri.find(protocol_delim); + std::string protocol = m_uri.substr(0, protocol_end); + protocol_end += protocol_delim.length(); + + std::transform(protocol.begin(), protocol.end(), protocol.begin(), + [](unsigned char c) { return std::tolower(c); }); + + std::string path; + size_t path_end = m_uri.find_first_of("#?", protocol_end); + if (path_end != std::string::npos) { + path = m_uri.substr(protocol_end, path_end++ - protocol_end); + } else { + path = m_uri.substr(protocol_end); + } + + /* %-decode the string. */ + std::string decoded_path; + decoded_path.reserve(path.length()); + for (size_t i = 0; i < path.length(); ++i) + if (path[i] == '%' && std::isxdigit(path[i + 1]) && std::isxdigit(path[i + 2])) { + decoded_path += std::stoi(path.substr(i + 1, 2), 0, 16); + i += 2; + } else { + decoded_path += path[i]; + } + + /* Tokenize the query/fragment. */ + std::vector tokens; + size_t pos, last = path_end; + while ((pos = m_uri.find('&', last)) != std::string::npos) { + tokens.emplace_back(m_uri.substr(last, pos - last)); + last = pos + 1; + } + if (last != std::string::npos) { + tokens.emplace_back(m_uri.substr(last)); + } + + /* Create a tag-value map from the tokenized query/fragment. */ + std::unordered_map params; + std::for_each(tokens.begin(), tokens.end(), [&](std::string& token) { + size_t delim = token.find('='); + if (delim != std::string::npos) { + params.emplace(token.substr(0, delim), token.substr(delim + 1)); + } + }); + + buffer = std::vector{}; + try { + size_t offset{0}, size{0}; + + if (auto offset_it = params.find("offset"); offset_it != params.end()) { + offset = std::stoul(offset_it->second, nullptr, 0); + } + + if (auto size_it = params.find("size"); size_it != params.end()) { + if (!(size = std::stoul(size_it->second, nullptr, 0))) return; + } + + if (protocol != "file") { + printf("\"%s\" protocol not supported\n", protocol.c_str()); + return; + } + + std::ifstream file(decoded_path, std::ios::in | std::ios::binary); + if (!file) { + printf("could not open `%s'\n", decoded_path.c_str()); + return; + } + + if (!size) { + file.ignore(std::numeric_limits::max()); + size_t bytes = file.gcount(); + file.clear(); + + if (bytes < offset) { + printf("invalid uri `%s' (file size < offset)\n", decoded_path.c_str()); + return; + } + size = bytes - offset; + } + + file.seekg(offset, std::ios_base::beg); + buffer.resize(size); + file.read(&buffer[0], size); + } catch (...) { + } +} + +DisassemblyInstance::DisassemblyInstance(code_object_decoder_t& decoder) + : buffer(reinterpret_cast(decoder.buffer.data())), + size(decoder.buffer.size()), + instructions(decoder.instructions) { + amd_comgr_create_data(AMD_COMGR_DATA_KIND_RELOCATABLE, &data); + amd_comgr_set_data(data, size, decoder.buffer.data()); + + char isa_name[128]; + size_t isa_size = sizeof(isa_name); + amd_comgr_get_data_isa_name(data, &isa_size, isa_name); + + CHECK_COMGR(amd_comgr_create_disassembly_info( + isa_name, //"amdgcn-amd-amdhsa--gfx1100", + &DisassemblyInstance::memory_callback, &DisassemblyInstance::inst_callback, + [](uint64_t address, void* user_data) {}, &info)); +} + +amd_comgr_status_t DisassemblyInstance::symbol_callback(amd_comgr_symbol_t symbol, + void* user_data) { + amd_comgr_symbol_type_t type; + amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_TYPE, &type); + + if (type != AMD_COMGR_SYMBOL_TYPE_FUNC && type != AMD_COMGR_SYMBOL_TYPE_AMDGPU_HSA_KERNEL) + return AMD_COMGR_STATUS_SUCCESS; + + uint64_t addr; + amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_VALUE, &addr); + + uint64_t mem_size; + amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_SIZE, &mem_size); + + uint64_t name_size; + amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_NAME_LENGTH, &name_size); + + std::string name; + name.resize(name_size); + + amd_comgr_symbol_get_info(symbol, AMD_COMGR_SYMBOL_INFO_NAME, name.data()); + + static_cast(user_data)->symbol_map[addr] = {name, mem_size}; + return AMD_COMGR_STATUS_SUCCESS; +} + +std::map>& DisassemblyInstance::GetKernelMap() { + symbol_map = std::map>{}; + amd_comgr_iterate_symbols(data, &DisassemblyInstance::symbol_callback, this); + return symbol_map; +} + +DisassemblyInstance::~DisassemblyInstance() { + amd_comgr_release_data(data); + CHECK_COMGR(amd_comgr_destroy_disassembly_info(info)); +} + +uint64_t DisassemblyInstance::ReadInstruction(uint64_t addr, const char* cpp_line) { + uint64_t size_read; + amd_comgr_disassemble_instruction(info, buffer + addr, (void*)this, &size_read); + instructions.back().address = addr; + instructions.back().cpp_reference = cpp_line; + return size_read; +} + +uint64_t DisassemblyInstance::memory_callback(uint64_t from, char* to, uint64_t size, + void* user_data) { + DisassemblyInstance& instance = *static_cast(user_data); + size_t copysize = std::min((int64_t)size, instance.buffer + instance.size - (int64_t)from); + std::memcpy(to, (char*)from, copysize); + return copysize; +} + +void DisassemblyInstance::inst_callback(const char* instruction, void* user_data) { + DisassemblyInstance& instance = *static_cast(user_data); + instance.instructions.push_back({strdup(instruction), nullptr, 0}); +} diff --git a/plugin/att/disassembly.hpp b/plugin/att/disassembly.hpp new file mode 100644 index 0000000000..2344c35bbe --- /dev/null +++ b/plugin/att/disassembly.hpp @@ -0,0 +1,59 @@ +/* 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. */ + +#pragma once + +#include +#include +#include +#include + +typedef struct { + const char* instruction; + const char* cpp_reference; + uint64_t address; +} instruction_instance_t; + +class CodeObjectBinary { + public: + CodeObjectBinary(const std::string& uri); + std::string m_uri; + std::vector buffer; +}; + +class DisassemblyInstance { + public: + DisassemblyInstance(class code_object_decoder_t& decoder); + ~DisassemblyInstance(); + + uint64_t ReadInstruction(uint64_t addr, const char* cpp_line); + std::map>& GetKernelMap(); + + static uint64_t memory_callback(uint64_t from, char* to, uint64_t size, void* user_data); + static void inst_callback(const char* instruction, void* user_data); + static amd_comgr_status_t symbol_callback(amd_comgr_symbol_t symbol, void* user_data); + + int64_t buffer; + int64_t size; + std::vector& instructions; + amd_comgr_disassembly_info_t info; + amd_comgr_data_t data; + std::map> symbol_map; +}; diff --git a/plugin/att/drawing.py b/plugin/att/drawing.py index f9549f6fbe..631762609f 100644 --- a/plugin/att/drawing.py +++ b/plugin/att/drawing.py @@ -120,6 +120,7 @@ def draw_wave_states(selections, normalize, TIMELINES): plt.figure(figsize=(15,4)) maxtime = max([np.max((TIMELINES[k]!=0)*np.arange(0,TIMELINES[k].size)) for k in plot_indices]) + maxtime = max(maxtime, 1) timelines = [deepcopy(TIMELINES[k][:maxtime]) for k in plot_indices] timelines = [np.pad(t, [0, maxtime-t.size]) for t in timelines] @@ -161,6 +162,7 @@ def draw_occupancy(selections, normalize, OCCUPANCY, shadernames): OCCUPANCY = [occ for occ in OCCUPANCY if len(occ) > 0] maxtime = 1 + delta = 1 for name, occ in zip(shadernames, OCCUPANCY): occ_values = [0] occ_times = [0] diff --git a/plugin/att/stitch.py b/plugin/att/stitch.py index 242996a404..b471edb4df 100644 --- a/plugin/att/stitch.py +++ b/plugin/att/stitch.py @@ -6,7 +6,7 @@ if sys.version_info[0] < 3: from collections import defaultdict from copy import deepcopy -MAX_STITCHED_TOKENS = 10000000 +MAX_STITCHED_TOKENS = 100000000 MAX_FAILED_STITCHES = 256 STACK_SIZE_LIMIT = 64 @@ -25,6 +25,7 @@ GETPC = 11 SETPC = 12 SWAPPC = 13 LANEIO = 14 +PCINFO = 15 DONT_KNOW = 100 WaveInstCategory = { @@ -46,10 +47,11 @@ WaveInstCategory = { SETPC: "SETPC", SWAPPC: "SWAPPC", LANEIO: "LANEIO", + PCINFO: "PCINFO", DONT_KNOW: "DONT_KNOW", } - +# Keeps track of register states for hipcc-generated assembly class RegisterWatchList: def __init__(self, labels): self.registers = {'v'+str(k): [[] for m in range(64)] for k in range(64)} @@ -84,7 +86,7 @@ class RegisterWatchList: except: pass - def swappc(self, line, line_num): + def swappc(self, line, line_num, inst_num): try: tokens = self.tokenize(line) dst = tokens[1] @@ -97,10 +99,9 @@ class RegisterWatchList: except: return 0 - def setpc(self, line): + def setpc(self, line, inst_num): try: src = line.split(' ')[1].strip() - #print('Going to:', self.registers[self.range(src)[0]], src) popped = self.registers[self.range(src)[0]][-1] self.registers[self.range(src)[0]] = self.registers[self.range(src)[0]][:-1] return popped @@ -140,16 +141,57 @@ class RegisterWatchList: except Exception as e: pass +# Translates PC values to instructions, for auto captured ISA +class PCTranslator: + def __init__(self, code, insts): + self.code = code + self.insts = insts + self.addrmap = {code[m][-3] : m for m in range(len(code))} + + def try_translate(self, tok): + pass + def range(self, r): + pass + def tokenize(self, line): + pass + def getpc(self, line, next_line): + pass + def swappc(self, line, line_num, inst_index): + try: + loc = self.addrmap[self.insts[inst_index+1][2]] + #print('Jumping to:', loc, self.code[loc]) + return loc + except: + print('SWAPPC: Could not find addr', self.insts[inst_index+1][2], 'for', line) + return -1 + def setpc(self, line, inst_index): + try: + loc = self.addrmap[self.insts[inst_index+1][2]] + #print('Jumping to:', loc, self.code[loc]) + return loc + except: + print('SETPC: Could not find addr', self.insts[inst_index+1][2], 'for', line) + return -1 + def scratch(self, line): + pass + def move(self, line): + pass + def updatelane(self, line): + pass + +# Matches tokens in reverse order def try_match_swapped(insts, code, i, line): return insts[i+1][1] == code[line][1] and insts[i][1] == code[line+1][1] FORK_NAMES = 1 +# A successful parsed instruction class CachedInst: def __init__(self, inst, as_line): self.inst_type = inst self.as_line = as_line self.forks = None +# A branch of the parsing tree class Fork: def __init__(self): global FORK_NAMES @@ -159,19 +201,21 @@ class Fork: FORK_NAMES += 1 #print('Created new fork: ', self.name) -def move_down_fork(fork, insts, i): #def move_down_fork(fork : Fork, insts : list, i : int): +# Try to match sequence "insts" with the branch "fork", starting at position "i" +def move_down_fork(fork, insts, i): #(fork : Fork, insts : list, i : int): N = min(len(insts), len(fork.insts)) while i < N: if insts[i][1] == fork.insts[i].inst_type: i += 1 - elif i= len(cur_fork.insts): return False, cur_fork @@ -204,7 +247,6 @@ def fromDict(insts): bMatchFork = False for fork in last_inst.forks: if fork.insts[0].inst_type == insts[0][1]: - #print('Found match fork', fork.name) cur_fork = fork bMatchFork = True break @@ -217,10 +259,30 @@ def fromDict(insts): return False, cur_fork -def stitch(insts, raw_code, jumps, gfxv): +def stitch(insts, raw_code, jumps, gfxv, bIsAuto): bGFX9 = gfxv == 'vega' + + # Try from cached result from a previous wave that have already been parsed + dict_sucess, current_fork = fromDict(insts) + if dict_sucess: + result, loopCount, mem_unroll, flight_count, maxline, pcsequence = current_fork.data + # Check if the sequence of measured PC values are equal for cached and new wave + if len(pcsequence) > 0: + pcs = [r[2] for r in insts if r[1] == PCINFO] + if len(pcs) != len(pcsequence): + dict_sucess = False + for pc1, pc2 in zip(pcs, pcsequence): + if pc1 != pc2: + dict_sucess = False + + # If successful, use resulting assembly from cache + if dict_sucess: + result = [r+(asm[-1],) for r, asm in zip(insts, result)] + return result, loopCount, mem_unroll, flight_count, maxline, len(result) + result, i, line, loopCount, N = [], 0, 0, defaultdict(int), len(insts) + SMEM_INST = [] # scalar memory VLMEM_INST = [] # vector memory load VSMEM_INST = [] # vector memory store @@ -236,8 +298,14 @@ def stitch(insts, raw_code, jumps, gfxv): labels = {} jump_map = [0] + + # Clean the code and remove comments code = [raw_code[0]] for c in raw_code[1:]: + if bIsAuto and '; Begin ' == c[0][:len('; Begin ')]: + if '; Begin ' in c[0]: + line = len(code) + print('Begin at:', line, c) c = list(c) c[0] = c[0].split(';')[0].split('//')[0].strip() @@ -254,25 +322,22 @@ def stitch(insts, raw_code, jumps, gfxv): jumps = {jump_map[j]+1: j for j in jumps} + # Checks if we have guaranteed ordering in memory operations smem_ordering = 0 vlmem_ordering = 0 vsmem_ordering = 0 - watchlist = RegisterWatchList(labels=labels) - num_failed_stitches = 0 loops = 0 maxline = 0 - dict_sucess, current_fork = fromDict(insts) - if dict_sucess: - result, loopCount, mem_unroll, flight_count, maxline = current_fork.data - result = [r+(asm[-1],) for r, asm in zip(insts, result)] - return result, loopCount, mem_unroll, flight_count, maxline, len(insts) + watchlist = RegisterWatchList(labels=labels) if not bIsAuto else PCTranslator(code, insts) + pcsequence = [] while i < N: loops += 1 - if line >= len(code) or loops > MAX_STITCHED_TOKENS or num_failed_stitches > MAX_FAILED_STITCHES: + if line >= len(code) or loops > MAX_STITCHED_TOKENS \ + or num_failed_stitches > MAX_FAILED_STITCHES: break maxline = max(reverse_map[line], maxline) @@ -282,23 +347,37 @@ def stitch(insts, raw_code, jumps, gfxv): matched = True next = line+1 - if '_mov_' in as_line[0]: - watchlist.move(as_line[0]) - elif 'scratch_' in as_line[0]: - watchlist.scratch(as_line[0]) + if not bIsAuto: + if '_mov_' in as_line[0]: + watchlist.move(as_line[0]) + elif 'scratch_' in as_line[0]: + watchlist.scratch(as_line[0]) if as_line[1] == GETPC: - watchlist.getpc(as_line[0], code[line+1][0]) - matched = inst[1] in [SALU, JUMP] + try: + watchlist.getpc(as_line[0], code[line+1][0]) + matched = inst[1] in [SALU, JUMP] + except: + matched = False elif as_line[1] == LANEIO: watchlist.updatelane(as_line[0]) matched = inst[1] == VALU elif as_line[1] == SETPC: - next = watchlist.setpc(as_line[0]) + next = watchlist.setpc(as_line[0], i) matched = inst[1] in [SALU, JUMP] + if bIsAuto: + matched = next >= 0 + i += 1 + result.append((insts[i][0], PCINFO, 0, 0, 0)) + pcsequence.append(insts[i][2]) elif as_line[1] == SWAPPC: - next = watchlist.swappc(as_line[0], line) + next = watchlist.swappc(as_line[0], line, i) matched = inst[1] in [SALU, JUMP] + if bIsAuto: + matched = next >= 0 + i += 1 + result.append((insts[i][0], PCINFO, 0, 0, 0)) + pcsequence.append(insts[i][2]) elif inst[1] == as_line[1]: if line in jumps: loopCount[jumps[line]-1] += 1 @@ -366,10 +445,10 @@ def stitch(insts, raw_code, jumps, gfxv): if 'vscnt' in as_line[0] or (bGFX9 and 'vmcnt' in as_line[0]): try: - wait_N = int(as_line[0].split('vmcnt(')[1].split(')')[0]) + wait_N = int(as_line[0].split('vscnt(')[1].split(')')[0]) except: try: - wait_N = int(as_line[0].split('vscnt(')[1].split(')')[0]) + wait_N = int(as_line[0].split('vmcnt(')[1].split(')')[0]) except: wait_N = 0 flight_count.append([as_line[5], num_inflight, wait_N]) @@ -407,10 +486,13 @@ def stitch(insts, raw_code, jumps, gfxv): if skipped_immed > 0 and 's_waitcnt ' in as_line[0]: matched = True skipped_immed -= 1 - else: + elif 'scratch_' not in as_line[0]: print('Parsing terminated at:', as_line) break + #print(matched, as_line) + #print([WaveInstCategory[insts[i+k][1]] for k in range(10) if i+k < len(insts)]) + if matched: result.append(inst + (reverse_map[line],)) i += 1 @@ -425,8 +507,8 @@ def stitch(insts, raw_code, jumps, gfxv): line = next N = max(N, 1) - if len(result) != N: - print('Warning - Stitching rate: '+str(len(result) * 100 / N)+'% matched') + if i != N: + print('Warning - Stitching rate: '+str(i * 100 / N)+'% matched') print('Leftovers:', [WaveInstCategory[insts[i+k][1]] for k in range(20) if i+k < len(insts)]) try: print(line, code[line]) @@ -440,5 +522,6 @@ def stitch(insts, raw_code, jumps, gfxv): line += 1 current_fork.insts = [CachedInst(inst[1], inst[-1]) for inst in result] - current_fork.data = result, loopCount, mem_unroll, flight_count, maxline - return result, loopCount, mem_unroll, flight_count, maxline, len(insts) + current_fork.data = result, loopCount, mem_unroll, flight_count, maxline, pcsequence + result = [r for r in result if r[1] != PCINFO] + return result, loopCount, mem_unroll, flight_count, maxline, len(result) if i == N else N diff --git a/plugin/att/trace_view.py b/plugin/att/trace_view.py index ac61713b73..0539b33deb 100755 --- a/plugin/att/trace_view.py +++ b/plugin/att/trace_view.py @@ -278,7 +278,8 @@ def view_trace(args, code, dbnames, att_filenames, bReturnLoc, OCCUPANCY, bDumpO simd_wave_filenames[se_number] = wv_filenames if mpi_root: - JSON_GLOBAL_DICTIONARY['code.json'] = Readable({"code": code[:allse_maxline+16], "top_n": get_top_n(code[:allse_maxline+16])}) + code_sel = [c[:-3]+c[-2:] for c in code[:allse_maxline+16]] + JSON_GLOBAL_DICTIONARY['code.json'] = Readable({"code": code_sel, "top_n": get_top_n(code_sel)}) for key in simd_wave_filenames.keys(): wv_array = [[ diff --git a/src/api/CMakeLists.txt b/src/api/CMakeLists.txt index 7ba1c9632d..f994227ec1 100644 --- a/src/api/CMakeLists.txt +++ b/src/api/CMakeLists.txt @@ -184,11 +184,15 @@ 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) +set(ROCPROFILER_ATT_SRC_FILES + ${PROJECT_SOURCE_DIR}/src/core/session/att/att.cpp) file(GLOB ROCPROFILER_CLASS_SRC_FILES ${CMAKE_CURRENT_SOURCE_DIR}/rocprofiler_singleton.cpp) file(GLOB ROCPROFILER_SPM_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/spm/spm.cpp) +set(CORE_ISA_CAPTURE_DIR ${PROJECT_SOURCE_DIR}/src/core/isa_capture) +file(GLOB CORE_ISA_CAPTURE_SRC_FILES ${CORE_ISA_CAPTURE_DIR}/*.cpp) + set(CORE_HARDWARE_DIR ${PROJECT_SOURCE_DIR}/src/core/hardware) file(GLOB CORE_HARDWARE_SRC_FILES ${CORE_HARDWARE_DIR}/*.cpp) @@ -265,6 +269,7 @@ add_library( ${ROCPROFILER_CLASS_SRC_FILES} ${ROCPROFILER_PROFILER_SRC_FILES} ${ROCPROFILER_ATT_SRC_FILES} + ${CORE_ISA_CAPTURE_SRC_FILES} ${CORE_HARDWARE_SRC_FILES} ${CORE_HSA_SRC_FILES} ${ROCPROFILER_SPM_SRC_FILES} @@ -326,6 +331,8 @@ if(ASAN) stdc++ stdc++fs amd_comgr + dw + elf ${PCIACCESS_LIBRARIES}) else() target_link_options( @@ -343,6 +350,8 @@ else() stdc++ stdc++fs amd_comgr + dw + elf ${PCIACCESS_LIBRARIES}) endif() diff --git a/src/api/exportmap b/src/api/exportmap index 0977a8d25f..d4cc7b8c53 100644 --- a/src/api/exportmap +++ b/src/api/exportmap @@ -38,5 +38,10 @@ global: OnLoad; rocprofiler_device_profiling_session_poll; rocprofiler_device_profiling_session_stop; rocprofiler_device_profiling_session_destroy; + rocprofiler_codeobj_capture_get; + rocprofiler_codeobj_capture_create; + rocprofiler_codeobj_capture_free; + rocprofiler_codeobj_capture_start; + rocprofiler_codeobj_capture_stop; local: *; }; \ No newline at end of file diff --git a/src/api/rocprofilerv2.cpp b/src/api/rocprofilerv2.cpp index d3aa4b40c4..9c473faf39 100644 --- a/src/api/rocprofilerv2.cpp +++ b/src/api/rocprofilerv2.cpp @@ -8,6 +8,7 @@ #include "src/core/hsa/hsa_support.h" #include "src/api/rocprofiler_singleton.h" #include "src/utils/helper.h" +#include "src/core/isa_capture/code_object_track.hpp" // TODO(aelwazir): change that to adapt with our own Exception // What about outside exceptions and callbacks exceptions!! @@ -543,6 +544,59 @@ rocprofiler_device_profiling_session_destroy(rocprofiler_session_id_t session_id API_METHOD_SUFFIX } +ROCPROFILER_API rocprofiler_status_t +rocprofiler_codeobj_capture_get(rocprofiler_record_id_t id, + rocprofiler_codeobj_symbols_t* symbols) { + API_METHOD_PREFIX + try { + *symbols = codeobj_record::get_capture(id); + } catch(const std::out_of_range& e) { + return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS; + } + API_METHOD_SUFFIX +} + +ROCPROFILER_API rocprofiler_status_t +rocprofiler_codeobj_capture_create( + rocprofiler_record_id_t* id, + rocprofiler_codeobj_capture_mode_t mode, + uint64_t userdata +) { + API_METHOD_PREFIX + id->handle = rocprofiler::GetROCProfilerSingleton()->GetUniqueRecordId(); + codeobj_record::make_capture(*id, mode, userdata); + API_METHOD_SUFFIX +} + +ROCPROFILER_API rocprofiler_status_t +rocprofiler_codeobj_capture_free(rocprofiler_record_id_t id) { + API_METHOD_PREFIX + codeobj_record::free_capture(id); + API_METHOD_SUFFIX +} + +ROCPROFILER_API rocprofiler_status_t +rocprofiler_codeobj_capture_start(rocprofiler_record_id_t id) { + API_METHOD_PREFIX + try { + codeobj_record::start_capture(id); + } catch(const std::out_of_range& e) { + return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS; + } + API_METHOD_SUFFIX +} + +ROCPROFILER_API rocprofiler_status_t +rocprofiler_codeobj_capture_stop(rocprofiler_record_id_t id) { + API_METHOD_PREFIX + try { + codeobj_record::stop_capture(id); + } catch(const std::out_of_range& e) { + return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENTS; + } + API_METHOD_SUFFIX +} + static bool started{false}; diff --git a/src/core/hsa/hsa_support.cpp b/src/core/hsa/hsa_support.cpp index 6cb5d21a16..828d9c7c33 100644 --- a/src/core/hsa/hsa_support.cpp +++ b/src/core/hsa/hsa_support.cpp @@ -50,6 +50,7 @@ #include "src/core/hsa/queues/queue.h" #include "src/api/rocprofiler_singleton.h" +#include "src/core/isa_capture/code_object_track.hpp" #include namespace fs = std::experimental::filesystem; @@ -464,6 +465,16 @@ hsa_status_t CodeObjectCallback(hsa_executable_t executable, data.codeobj.unload = *static_cast(arg) ? 1 : 0; ReportActivity(ACTIVITY_DOMAIN_HSA_EVT, HSA_EVT_ID_CODEOBJ, &data); + if (data.codeobj.unload) + codeobj_capture_instance::Unload(data.codeobj.load_base); + else + codeobj_capture_instance::Load( + data.codeobj.load_base, + uri_str, + data.codeobj.memory_base, + data.codeobj.memory_size + ); + hsa_executable_iterate_agent_symbols(executable, data.codeobj.agent, hsa_executable_iteration_callback, &(data.codeobj.unload)); diff --git a/src/core/hsa/queues/queue.cpp b/src/core/hsa/queues/queue.cpp index c173dcab16..202f5b1b43 100644 --- a/src/core/hsa/queues/queue.cpp +++ b/src/core/hsa/queues/queue.cpp @@ -33,6 +33,7 @@ #include "src/core/hsa/packets/packets_generator.h" #include "src/core/hsa/hsa_support.h" #include "src/utils/helper.h" +#include "src/core/isa_capture/code_object_track.hpp" #define CHECK_HSA_STATUS(msg, status) \ do { \ @@ -493,17 +494,20 @@ bool AsyncSignalHandler(hsa_signal_value_t signal_value, void* data) { } bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) { - // TODO: finish implementation to iterate trace data and add it to rocprofiler record - // and generic buffer auto queue_info_session = static_cast(data); - if (!queue_info_session || !GetROCProfilerSingleton() || - !GetROCProfilerSingleton()->GetSession(queue_info_session->session_id) || - !GetROCProfilerSingleton()->GetSession(queue_info_session->session_id)->GetAttTracer()) + if (!queue_info_session || !GetROCProfilerSingleton()) return true; + rocprofiler::Session* session = GetROCProfilerSingleton()->GetSession(queue_info_session->session_id); + if (!session) return true; + + std::lock_guard lock(session->GetSessionLock()); + rocprofiler::att::AttTracer* att_tracer = session->GetAttTracer(); + if (!session->GetAttTracer()) return true; + std::vector& pending_signals = const_cast&>( att_tracer->GetPendingSignals(queue_info_session->writer_id)); @@ -512,9 +516,9 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) { 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.new_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)queue_info_session->gpu_index}; @@ -522,13 +526,18 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) { 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}; + record.writer_id = queue_info_session->writer_id; + if (/*pending.counters_count > 0 && */ pending.profile) { AddAttRecord(&record, queue_info_session->agent, pending); } - // July/01/2023 -> Changed this to writer ID so we can correlate to dispatches - // kernel_id already has the descriptor. + // July/01/2023 -> Changed this to queue_info_session->writer_id + // so we can correlate to dispatches. kernel_id already has the descriptor. record.header = {ROCPROFILER_ATT_TRACER_RECORD, - rocprofiler_record_id_t{queue_info_session->writer_id}}; + rocprofiler_record_id_t{pending.kernel_descriptor}}; + + record.intercept_list = codeobj_record::get_capture(record.header.id); + std::atomic_thread_fence(std::memory_order_release); if (pending.session_id.handle == 0) { pending.session_id = GetROCProfilerSingleton()->GetCurrentSessionId(); @@ -536,7 +545,10 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) { if (session->FindBuffer(pending.buffer_id)) { Memory::GenericBuffer* buffer = session->GetBuffer(pending.buffer_id); buffer->AddRecord(record); + buffer->Flush(); } + codeobj_record::free_capture(record.header.id); + hsa_status_t status = rocprofiler::hsa_support::GetAmdExtTable().hsa_amd_memory_pool_free_fn( (pending.profile->output_buffer.ptr)); CHECK_HSA_STATUS("Error: Couldn't free output buffer memory", status); @@ -548,6 +560,7 @@ bool AsyncSignalHandlerATT(hsa_signal_value_t /* signal */, void* data) { } delete queue_info_session; + std::atomic_thread_fence(std::memory_order_seq_cst); return false; } @@ -721,17 +734,25 @@ std::pair, bool> GetAllowedProfilesList(const void* packets, i return {can_profile_packet, b_can_profile_anypacket}; } -hsa_ven_amd_aqlprofile_profile_t* ProcessATTParams(Packet::packet_t& start_packet, - Packet::packet_t& stop_packet, Queue& queue_info, - Agent::AgentInfo& agentInfo) { +std::pair +ProcessATTParams( + Packet::packet_t& start_packet, + Packet::packet_t& stop_packet, + Queue& queue_info, + Agent::AgentInfo& agentInfo +) { std::vector att_params; int num_att_counters = 0; uint32_t att_buffer_size = DEFAULT_ATT_BUFFER_SIZE; + rocprofiler_codeobj_capture_mode_t capture_mode = ROCPROFILER_CAPTURE_SYMBOLS_ONLY; for (rocprofiler_att_parameter_t& param : att_parameters_data) { switch (param.parameter_name) { case ROCPROFILER_ATT_PERFCOUNTER_NAME: break; + case ROCPROFILER_ATT_CAPTURE_MODE: + capture_mode = static_cast(param.value); + break; case ROCPROFILER_ATT_BUFFER_SIZE: att_buffer_size = std::max(96l << 10l, std::min(int64_t(param.value) << 20l, (1l << 32l) - (3l << 20))); @@ -773,8 +794,8 @@ hsa_ven_amd_aqlprofile_profile_t* ProcessATTParams(Packet::packet_t& start_packe for (; num_att_counters < 16; num_att_counters++) att_params.push_back(zero_perf); } // Get the PM4 Packets using packets_generator - return Packet::GenerateATTPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(), att_params, - &start_packet, &stop_packet, att_buffer_size); + return {Packet::GenerateATTPackets(queue_info.GetCPUAgent(), queue_info.GetGPUAgent(), + att_params, &start_packet, &stop_packet, att_buffer_size), capture_mode}; } /** @@ -786,6 +807,7 @@ hsa_ven_amd_aqlprofile_profile_t* ProcessATTParams(Packet::packet_t& start_packe */ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt_index, void* data, hsa_amd_queue_intercept_packet_writer writer) { + static const char* env_MAX_ATT_PROFILES = getenv("ROCPROFILER_MAX_ATT_PROFILES"); static int MAX_ATT_PROFILES = env_MAX_ATT_PROFILES ? atoi(env_MAX_ATT_PROFILES) : 1; @@ -952,9 +974,15 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt Packet::packet_t start_packet{}; Packet::packet_t stop_packet{}; hsa_ven_amd_aqlprofile_profile_t* profile = nullptr; + rocprofiler_codeobj_capture_mode_t capture_mode = ROCPROFILER_CAPTURE_SYMBOLS_ONLY; - if (att_parameters_data.size() > 0 && is_att_collection_mode) - profile = ProcessATTParams(start_packet, stop_packet, queue_info, agentInfo); + if (att_parameters_data.size() > 0) { + std::tie(profile, capture_mode) = ProcessATTParams(start_packet, + stop_packet, + queue_info, + agentInfo + ); + } // Searching across all the packets given during this write for (size_t i = 0; i < pkt_count; ++i) { @@ -976,7 +1004,7 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt KernelInterceptCount += 1; writer_id = WRITER_ID.fetch_add(1, std::memory_order_release); - if (att_parameters_data.size() > 0 && is_att_collection_mode && profile) { + if (att_parameters_data.size() > 0 && profile) { // Adding start packet and its barrier with a dummy signal hsa_signal_t dummy_signal{}; dummy_signal.handle = 0; @@ -996,17 +1024,17 @@ void WriteInterceptor(const void* packets, uint64_t pkt_count, uint64_t user_pkt uint64_t record_id = GetROCProfilerSingleton()->GetUniqueRecordId(); AddKernelNameWithDispatchID(GetKernelNameFromKsymbols(dispatch_packet.kernel_object), record_id); - if (session && profile) { - session->GetAttTracer()->AddPendingSignals( - writer_id, record_id, original_packet.completion_signal, - dispatch_packet.completion_signal, session_id_snapshot, buffer_id, profile, - kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index); - } else { - session->GetAttTracer()->AddPendingSignals( - writer_id, record_id, original_packet.completion_signal, - dispatch_packet.completion_signal, session_id_snapshot, buffer_id, nullptr, - kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index); - } + + session->GetAttTracer()->AddPendingSignals( + writer_id, record_id, original_packet.completion_signal, + dispatch_packet.completion_signal, session_id_snapshot, buffer_id, profile, + kernel_properties, (uint32_t)syscall(__NR_gettid), user_pkt_index); + + uint64_t off = dispatch_packet.kernel_object + + GetKernelCode(dispatch_packet.kernel_object)->kernel_code_entry_byte_offset; + codeobj_record::make_capture(rocprofiler_record_id_t{record_id}, capture_mode, off); + codeobj_record::start_capture(rocprofiler_record_id_t{record_id}); + codeobj_record::stop_capture(rocprofiler_record_id_t{record_id}); // Make a copy of the original packet, adding its signal to a barrier packet if (original_packet.completion_signal.handle) { diff --git a/src/core/isa_capture/code_object_track.cpp b/src/core/isa_capture/code_object_track.cpp new file mode 100644 index 0000000000..0ea0a24eb6 --- /dev/null +++ b/src/core/isa_capture/code_object_track.cpp @@ -0,0 +1,248 @@ +/* Copyright (c) 2023 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 "src/utils/helper.h" +#include "src/api/rocprofiler_singleton.h" +#include "src/core/isa_capture/code_object_track.hpp" +#include + +std::mutex codeobj_capture_instance::mutex; +std::mutex codeobj_record::mutex; + +std::unordered_map codeobj_capture_instance::codeobjs{}; +std::unordered_map codeobj_record::record_id_map{}; +std::unordered_set codeobj_record::listeners; + +// Codeobj Record +codeobj_record::codeobj_record(rocprofiler_codeobj_capture_mode_t mode) : capture_mode(mode){}; + +void codeobj_record::start_capture() { + listeners.insert(this); + for (auto& [addr, capture] : codeobj_capture_instance::codeobjs) this->addcapture(capture); +} + +void codeobj_record::addcapture(CodeobjPtr& capture) { + if (captures.find(capture) != captures.end()) return; + capture->setmode(capture_mode); + captures.insert(capture); +} + +void codeobj_record::stop_capture() { + try { + listeners.erase(this); + } catch (...) { + }; +} + +// Codeobj Capture +void codeobj_capture_instance::Load(uint64_t addr, const std::string& URI, uint64_t mem_addr, + uint64_t mem_size) { + std::lock_guard lock(mutex); + codeobjs[addr] = std::make_shared( + addr, URI, mem_addr, mem_size, rocprofiler::GetCurrentTimestamp().value); + std::atomic_thread_fence(std::memory_order_release); // Fencing the state of the map + { + std::lock_guard lock(codeobj_record::mutex); + for (auto* listen : codeobj_record::listeners) listen->addcapture(codeobjs.at(addr)); + } +} + +void codeobj_capture_instance::Unload(uint64_t addr) { + std::lock_guard lock(mutex); + codeobjs.at(addr)->end_time = rocprofiler::GetCurrentTimestamp().value; + codeobjs.erase(addr); +} + +void codeobj_capture_instance::copyCodeobjFromFile(uint64_t offset, uint64_t size, + const std::string& decoded_path) { + std::ifstream file(decoded_path, std::ios::in | std::ios::binary); + if (!file) { + printf("could not open `%s'\n", decoded_path.c_str()); + return; + } + + if (!size) { + file.ignore(std::numeric_limits::max()); + size_t bytes = file.gcount(); + file.clear(); + + if (bytes < offset) { + printf("invalid uri `%s' (file size < offset)\n", decoded_path.c_str()); + return; + } + size = bytes - offset; + } + + file.seekg(offset, std::ios_base::beg); + buffer.resize(size); + file.read(&buffer[0], size); +} + +void codeobj_capture_instance::copyCodeobjFromMemory(uint64_t mem_addr, uint64_t mem_size) { + buffer.resize(mem_size); + std::memcpy(buffer.data(), (uint64_t*)mem_addr, mem_size); +} + +std::pair codeobj_capture_instance::parse_uri() { + const std::string protocol_delim{"://"}; + + size_t protocol_end = URI.find(protocol_delim); + protocol = URI.substr(0, protocol_end); + protocol_end += protocol_delim.length(); + + std::transform(protocol.begin(), protocol.end(), protocol.begin(), + [](unsigned char c) { return std::tolower(c); }); + + std::string path; + size_t path_end = URI.find_first_of("#?", protocol_end); + if (path_end != std::string::npos) { + path = URI.substr(protocol_end, path_end++ - protocol_end); + } else { + path = URI.substr(protocol_end); + } + + /* %-decode the string. */ + decoded_path = std::string{}; + decoded_path.reserve(path.length()); + for (size_t i = 0; i < path.length(); ++i) { + if (path[i] == '%' && std::isxdigit(path[i + 1]) && std::isxdigit(path[i + 2])) { + decoded_path += std::stoi(path.substr(i + 1, 2), 0, 16); + i += 2; + } else { + decoded_path += path[i]; + } + } + + /* Tokenize the query/fragment. */ + std::vector tokens; + size_t pos, last = path_end; + while ((pos = URI.find('&', last)) != std::string::npos) { + tokens.emplace_back(URI.substr(last, pos - last)); + last = pos + 1; + } + if (last != std::string::npos) tokens.emplace_back(URI.substr(last)); + + /* Create a tag-value map from the tokenized query/fragment. */ + std::unordered_map params; + std::for_each(tokens.begin(), tokens.end(), [&](std::string& token) { + size_t delim = token.find('='); + if (delim != std::string::npos) { + params.emplace(token.substr(0, delim), token.substr(delim + 1)); + } + }); + + size_t offset{0}, size{0}; + + if (auto offset_it = params.find("offset"); offset_it != params.end()) + offset = std::stoul(offset_it->second, nullptr, 0); + + if (auto size_it = params.find("size"); size_it != params.end()) { + if (!(size = std::stoul(size_it->second, nullptr, 0))) throw std::exception(); + } + + return {offset, size}; +} + +codeobj_capture_instance::codeobj_capture_instance(uint64_t _addr, const std::string& _uri, + uint64_t mem_addr, uint64_t mem_size, + uint64_t start_time) + : addr(_addr), start_time(start_time), URI(_uri), mem_addr(mem_addr), mem_size(mem_size) { + reset(ROCPROFILER_CAPTURE_SYMBOLS_ONLY); +}; + +void codeobj_capture_instance::setmode(rocprofiler_codeobj_capture_mode_t mode) { + // Only reset when needed & check if codeobj was not unloaded + if (end_time == 0 && static_cast(mode) > static_cast(capture_mode)) + reset(mode); +} + +void codeobj_capture_instance::reset(rocprofiler_codeobj_capture_mode_t mode) { + capture_mode = mode; + + size_t offset, size; + try { + std::tie(offset, size) = parse_uri(); + } catch (...) { + return; + } + + buffer = std::vector{}; + + if (mode == ROCPROFILER_CAPTURE_SYMBOLS_ONLY) return; + + if (protocol == "file") { + if (mode == ROCPROFILER_CAPTURE_COPY_FILE_AND_MEMORY) + copyCodeobjFromFile(offset, size, decoded_path); + } else if (protocol == "memory") { + copyCodeobjFromMemory(mem_addr, mem_size); + } else { + printf("\"%s\" protocol not supported\n", protocol.c_str()); + return; + } +} + +// Public static funcs +void codeobj_record::make_capture(rocprofiler_record_id_t id, + rocprofiler_codeobj_capture_mode_t mode, uint64_t userdata) { + std::lock_guard lock(mutex); + record_id_map[id.handle] = {userdata, std::unique_ptr{new codeobj_record(mode)}}; +} + +void codeobj_record::free_capture(rocprofiler_record_id_t id) { + std::lock_guard lock(mutex); + record_id_map.erase(id.handle); +} + +void codeobj_record::start_capture(rocprofiler_record_id_t id) { + std::lock_guard lock(mutex); + record_id_map.at(id.handle).second->start_capture(); +} + +void codeobj_record::stop_capture(rocprofiler_record_id_t id) { + std::lock_guard lock(mutex); + record_id_map.at(id.handle).second->stop_capture(); +} + +rocprofiler_codeobj_symbols_t codeobj_record::get_capture(rocprofiler_record_id_t id) { + std::atomic_thread_fence(std::memory_order_acquire); // Fencing the state of the map + auto& pair = record_id_map.at(id.handle); + return pair.second->get(pair.first); +} diff --git a/src/core/isa_capture/code_object_track.hpp b/src/core/isa_capture/code_object_track.hpp new file mode 100644 index 0000000000..339de2bf20 --- /dev/null +++ b/src/core/isa_capture/code_object_track.hpp @@ -0,0 +1,133 @@ +/* Copyright (c) 2023 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. */ + +#pragma once + +#include "src/utils/helper.h" +#include +#include +#include +#include +#include +#include + +/** + * A class to keep track of currently loaded code objects. + * Only the public static methods are thread-safe and expected to be used. + */ +class codeobj_capture_instance { + public: + codeobj_capture_instance(uint64_t _addr, const std::string& _uri, uint64_t mem_addr, + uint64_t mem_size, uint64_t start_time); + + void setmode(rocprofiler_codeobj_capture_mode_t mode); + + rocprofiler_intercepted_codeobj_t get() const { + const char* buf_ptr = buffer.size() ? buffer.data() : nullptr; + return {URI.c_str(), addr, buf_ptr, buffer.size(), start_time, end_time}; + }; + + const uint64_t addr; + const uint64_t start_time; + + static void Load(uint64_t addr, const std::string& URI, uint64_t mem_addr, uint64_t mem_size); + static void Unload(uint64_t addr); + + static std::unordered_map> codeobjs; + + private: + void reset(rocprofiler_codeobj_capture_mode_t mode); + + std::pair parse_uri(); + void DecodePath(); + void copyCodeobjFromFile(uint64_t offset, uint64_t size, const std::string& decoded_path); + void copyCodeobjFromMemory(uint64_t mem_addr, uint64_t mem_size); + + std::string URI; + std::string decoded_path; + std::string protocol; + std::vector buffer; + + uint64_t mem_addr; + uint64_t mem_size; + uint64_t end_time = 0; + rocprofiler_codeobj_capture_mode_t capture_mode; + + // Address -> codeobj + static std::mutex mutex; +}; + +typedef std::shared_ptr CodeobjPtr; + +template <> struct std::hash { + // addr is typically 2^12-byte aligned. Taking last 44 bits of time == cycle time of many hours. + uint64_t operator()(const CodeobjPtr& p) const { + return (p->addr >> 12) ^ (p->start_time << 20); + }; +}; + +template <> struct std::equal_to { + bool operator()(const CodeobjPtr& a, const CodeobjPtr& b) const { + return (a->addr == b->addr) & (a->start_time == b->start_time); + }; +}; + +/** + * A class to keep track of the history of loaded code objets. + * Only the public static methods are thread-safe and expected to be used. + */ +class codeobj_record { + public: + codeobj_record(rocprofiler_codeobj_capture_mode_t mode); + ~codeobj_record() { + if (listeners.find(this) != listeners.end()) stop_capture(); + }; + + void addcapture(CodeobjPtr& capture); + + public: + static void make_capture(rocprofiler_record_id_t id, rocprofiler_codeobj_capture_mode_t mode, + uint64_t userdata); + static void free_capture(rocprofiler_record_id_t id); + static void start_capture(rocprofiler_record_id_t id); + static void stop_capture(rocprofiler_record_id_t id); + static rocprofiler_codeobj_symbols_t get_capture(rocprofiler_record_id_t id); + + static std::unordered_set listeners; + static std::mutex mutex; + + private: + rocprofiler_codeobj_symbols_t get(uint64_t userdata) { + persist.clear(); + for (auto& capt : captures) persist.push_back(capt->get()); + return rocprofiler_codeobj_symbols_t{persist.data(), persist.size(), userdata}; + }; + + void start_capture(); + void stop_capture(); + + rocprofiler_codeobj_capture_mode_t capture_mode; + std::vector persist; + std::unordered_set captures; + + // Record_id -> codeobj + using RecordInstance = std::pair>; + static std::unordered_map record_id_map; +}; diff --git a/src/core/session/att/att.cpp b/src/core/session/att/att.cpp index 180d244b96..6b4aeafe2d 100644 --- a/src/core/session/att/att.cpp +++ b/src/core/session/att/att.cpp @@ -41,11 +41,9 @@ void AttTracer::AddPendingSignals( sessions_pending_signals_.at(writer_id).emplace_back(att_pending_signal_t{ kernel_object, original_completion_signal, new_completion_signal, session_id_, buffer_id, profile, kernel_properties, thread_id, queue_index}); - std::atomic_thread_fence(std::memory_order_release); } const std::vector& AttTracer::GetPendingSignals(uint32_t writer_id) { - std::atomic_thread_fence(std::memory_order_acquire); 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"); diff --git a/src/tools/tool.cpp b/src/tools/tool.cpp index ec1cf83c20..fdb8e084a7 100644 --- a/src/tools/tool.cpp +++ b/src/tools/tool.cpp @@ -287,12 +287,14 @@ att_parsed_input_t GetATTParams() { ATT_PARAM_NAMES["KERNEL"] = ROCPROFILER_ATT_MAXVALUE; ATT_PARAM_NAMES["BUFFER_SIZE"] = ROCPROFILER_ATT_BUFFER_SIZE; + ATT_PARAM_NAMES["ISA_CAPTURE_MODE"] = ROCPROFILER_ATT_CAPTURE_MODE; // Default values used for token generation. std::unordered_map default_params = { {"SE_MASK", 0x111111}, // One every 4 SEs, by default {"SIMD_SELECT", 0x3}, // 0x3 works for both gfx9 and Navi - {"BUFFER_SIZE", 0x40000000} // 2^30 == 1GB + {"BUFFER_SIZE", 0x40000000}, // 2^30 == 1GB + {"ISA_CAPTURE_MODE", static_cast(ROCPROFILER_CAPTURE_SYMBOLS_ONLY)} }; std::ifstream trace_file(path); diff --git a/tests-v2/featuretests/profiler/profiler_gtest.cpp b/tests-v2/featuretests/profiler/profiler_gtest.cpp index 76cb2a52bb..56d56b6d56 100644 --- a/tests-v2/featuretests/profiler/profiler_gtest.cpp +++ b/tests-v2/featuretests/profiler/profiler_gtest.cpp @@ -1150,6 +1150,146 @@ TEST(ProfilerMPTest, WhenRunningMultiProcessTestItPasses) { ASSERT_TRUE(1); } } + +/* + * ################################################### + * ############ Codeobj capture tests ################ + * ################################################### + */ + +class CodeobjTest : public ::testing::Test { +public: + virtual void SetUp(const char* app_name) {}; + virtual void TearDown(){}; + static 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) {}; + + void SetupRocprofiler() { + int result = ROCPROFILER_STATUS_ERROR; + + result = rocprofiler_initialize(); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + result = rocprofiler_create_session(ROCPROFILER_NONE_REPLAY_MODE, &session_id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + result = rocprofiler_codeobj_capture_create(&id, ROCPROFILER_CAPTURE_SYMBOLS_ONLY, 0); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + } + + void TearDownRocprofiler() { + int result = ROCPROFILER_STATUS_ERROR; + + result = rocprofiler_codeobj_capture_free(id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + result = rocprofiler_destroy_session(session_id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + result = rocprofiler_finalize(); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + } + + rocprofiler_session_id_t session_id; + rocprofiler_record_id_t id; +}; + +TEST_F(CodeobjTest, WhenRunningProfilerWithCodeobjCapture) { + int result = ROCPROFILER_STATUS_ERROR; + SetupRocprofiler(); + + result = rocprofiler_codeobj_capture_start(id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // Launch a kernel + LaunchVectorAddKernel(); + + result = rocprofiler_codeobj_capture_stop(id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + rocprofiler_codeobj_symbols_t capture; + + rocprofiler_codeobj_capture_get(id, &capture); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + EXPECT_GE(capture.count, 1); + bool bCaptured_itself = false; + + for (int i=0; i<(int)capture.count; i++) { + const char* path = capture.symbols[i].filepath; + if (!path) continue; + std::string fpath(path); + size_t pos = fpath.find("runFeatureTests#offset="); + if (pos != std::string::npos && fpath.find("&size=", pos) != std::string::npos) + bCaptured_itself = true; + } + EXPECT_EQ(bCaptured_itself, true); + + TearDownRocprofiler(); +} + + +TEST_F(CodeobjTest, WhenRunningProfilerWithMultipleCaptureAndCopy) { + int result = ROCPROFILER_STATUS_ERROR; + + SetupRocprofiler(); + + rocprofiler_record_id_t id2; + rocprofiler_record_id_t id3; + + result = rocprofiler_codeobj_capture_create(&id2, ROCPROFILER_CAPTURE_COPY_FILE_AND_MEMORY, 0); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + result = rocprofiler_codeobj_capture_start(id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + result = rocprofiler_codeobj_capture_start(id2); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + result = rocprofiler_codeobj_capture_create(&id3, ROCPROFILER_CAPTURE_COPY_MEMORY, 0); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + result = rocprofiler_codeobj_capture_start(id3); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // Launch a kernel + LaunchVectorAddKernel(); + + result = rocprofiler_codeobj_capture_stop(id2); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + result = rocprofiler_codeobj_capture_stop(id3); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + result = rocprofiler_codeobj_capture_free(id3); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + rocprofiler_codeobj_symbols_t capture; + rocprofiler_codeobj_capture_get(id2, &capture); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + EXPECT_GE(capture.count, 1); + + for (int i=0; i<(int)capture.count; i++) { + EXPECT_NE(capture.symbols[i].base_address, 0); + EXPECT_NE(capture.symbols[i].clock_start.value, 0); + EXPECT_NE(capture.symbols[i].data, nullptr); + EXPECT_NE(capture.symbols[i].size, 0); + } + + result = rocprofiler_codeobj_capture_stop(id); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + result = rocprofiler_codeobj_capture_free(id2); + EXPECT_EQ(ROCPROFILER_STATUS_SUCCESS, result); + + // Expect to return error when running get() after free() + result = rocprofiler_codeobj_capture_get(id3, &capture); + EXPECT_NE(ROCPROFILER_STATUS_SUCCESS, result); + + TearDownRocprofiler(); +} + /* * ################################################### * ############ Plugin tests ################ diff --git a/tests-v2/unittests/CMakeLists.txt b/tests-v2/unittests/CMakeLists.txt index e907473fdd..432a4988c1 100644 --- a/tests-v2/unittests/CMakeLists.txt +++ b/tests-v2/unittests/CMakeLists.txt @@ -5,6 +5,8 @@ find_library(PCIACCESS_LIBRARIES pciaccess REQUIRED) enable_testing() find_package(GTest REQUIRED) +find_library(GDB rocm-dbgapi PATHS ${ROCM_PATH} REQUIRED) + # Getting Source files for ROCProfiler, Hardware, HSA, Memory, Session, Counters, Utils set(CORE_MEMORY_DIR ${PROJECT_SOURCE_DIR}/src/core/memory) file(GLOB CORE_MEMORY_SRC_FILES ${CORE_MEMORY_DIR}/*.cpp) @@ -35,14 +37,15 @@ 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_ATT_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/session/att/*.cpp) file(GLOB ROCPROFILER_SRC_CLASS_FILES ${CMAKE_CURRENT_SOURCE_DIR}/rocprofiler_singleton.cpp) +file(GLOB ROCPROFILER_ISA_SRC_FILES ${PROJECT_SOURCE_DIR}/src/core/isa_capture/*.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} ${ROCPROFILER_ATT_SRC_FILES}) +set(ROCPROFILER_SRC_FILES ${ROCPROFILER_SRC_API_FILES} ${ROCPROFILER_ATT_SRC_FILES} + ${ROCPROFILER_ISA_SRC_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) @@ -98,7 +101,7 @@ target_compile_definitions( target_link_libraries( runUnitTests PRIVATE rocprofiler_tool ${AQLPROFILE_LIB} hsa-runtime64::hsa-runtime64 - GTest::gtest GTest::gtest_main stdc++fs ${PCIACCESS_LIBRARIES}) + GTest::gtest GTest::gtest_main stdc++fs ${PCIACCESS_LIBRARIES} ${GDB} dw elf c dl) add_dependencies(tests runUnitTests) install(TARGETS runUnitTests