From 21eaf44a4ba439565691c9559ffd6f2c027f33c7 Mon Sep 17 00:00:00 2001 From: Sushma Vaddireddy Date: Fri, 24 Apr 2026 01:01:53 +0000 Subject: [PATCH] Add support for gfx1250 --- projects/rocprofiler-sdk/CHANGELOG.md | 20 + projects/rocprofiler-sdk/VERSION | 2 +- .../Modules/rocprofiler-sdk-utilities.cmake | 3 +- .../cxx/codeobj/disassembly.hpp | 62 +- .../cxx/serialization/save.hpp | 6 + .../include/rocprofiler-sdk/pc_sampling.h | 34 +- .../source/lib/att-tool/waitcnt/gfx12.cpp | 45 +- .../tests/att_decoder_waitcnt_test.cpp | 14 +- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 7 +- .../rocprofiler-sdk/counters/evaluate_ast.cpp | 107 ++-- .../rocprofiler-sdk/counters/id_decode.hpp | 100 ++- .../counters/tests/dimension.cpp | 33 +- .../counters/tests/evaluate_ast_test.cpp | 14 +- .../source/lib/rocprofiler-sdk/hsa/abi.cpp | 3 + .../lib/rocprofiler-sdk/hsa/details/fmt.hpp | 69 +- .../source/lib/rocprofiler-sdk/hsa/queue.cpp | 130 +++- .../hsa/rocprofiler_packet.hpp | 15 +- .../pc_sampling/hsa_adapter.cpp | 2 +- .../pc_sampling/ioctl/ioctl_adapter.cpp | 8 + .../pc_sampling/parser/CMakeLists.txt | 3 +- .../pc_sampling/parser/correlation.hpp | 9 +- .../pc_sampling/parser/gfx1250.hpp | 30 + .../parser/pc_record_interface.cpp | 9 +- .../pc_sampling/parser/tests/CMakeLists.txt | 25 + .../pc_sampling/parser/tests/gfx1250test.cpp | 593 ++++++++++++++++++ .../pc_sampling/parser/tests/gfxtest.hpp | 36 ++ .../pc_sampling/parser/translation.hpp | 120 +++- .../lib/rocprofiler-sdk/pc_sampling/utils.cpp | 6 + .../lib/rocprofiler-sdk/pc_sampling/utils.hpp | 8 +- .../source/share/rocprofiler-sdk/config.yaml | 279 ++++++++ .../exec_mask_manipulation.cpp | 21 +- .../tests/pc_sampling/main.cpp | 60 +- .../rocprofiler-sdk/tests/pc_sampling/pcs.cpp | 9 + .../pc_sampling/exec_mask_manipulation/csv.py | 4 +- .../transpose_multiple_agents/csv.py | 6 +- .../exec-mask-manipulation/CMakeLists.txt | 6 + .../exec-mask-manipulation/conftest.py | 20 + .../exec-mask-manipulation/validate.py | 22 +- .../transpose-multiple-agents/validate.py | 22 +- 39 files changed, 1759 insertions(+), 203 deletions(-) create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/gfx1250.hpp create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/gfx1250test.cpp diff --git a/projects/rocprofiler-sdk/CHANGELOG.md b/projects/rocprofiler-sdk/CHANGELOG.md index dd00fe93f305..b22b30be51ed 100644 --- a/projects/rocprofiler-sdk/CHANGELOG.md +++ b/projects/rocprofiler-sdk/CHANGELOG.md @@ -328,3 +328,23 @@ Full documentation for ROCprofiler-SDK is available at [rocm.docs.amd.com/projec ### Removed - Counter collection support for plain text (`.txt`) input files has been deprecated due to lack of schema validation and input sanitization. Only structured file formats (JSON and YAML) with schema validation are supported. + +## ROCprofiler-SDK 1.3.0 + +### Added + +- gfx1250 PMC support, including counter definitions and mappings for CHA, CHC, GLARBC, GL1A, GL1C, GRBMA, and GRBMH new blocks. +- gfx1250 ATT support with per-XCC trace configuration, buffer programming, token-mask updates, and trace readback for multi-XCC devices. + - Extended ATT waitcnt decode for async and tensor-related instructions, adding support for `s_wait_asynccnt` and `s_wait_tensorcnt`. +- gfx1250 stochastic PC sampling support: + - gfx1250-specific parser with sample routing. + - New stochastic sample fields: `sampling_lock_error`, `async_cnt`, `tensor_cnt`, `xnack_cnt`, and chiplet-aware metadata. + - Memory dependency counters reveal per-wave memory latency. + - Updated public PC sampling API structs and serialization for the new fields. + - Kernel/IOCTL enablement and dedicated gfx1250 parser tests. + +### Changed + +- Counter dimension encoding changed from fixed-width to variable-width allocation per dimension type. +- Dimension selection and reduction logic now uses explicit dimension masks and single-index selection. +- HSA queue interception extended to handle AMD extended kernel dispatch packets. diff --git a/projects/rocprofiler-sdk/VERSION b/projects/rocprofiler-sdk/VERSION index 0495c4a88cae..f0bb29e76388 100644 --- a/projects/rocprofiler-sdk/VERSION +++ b/projects/rocprofiler-sdk/VERSION @@ -1 +1 @@ -1.2.3 +1.3.0 diff --git a/projects/rocprofiler-sdk/cmake/Modules/rocprofiler-sdk-utilities.cmake b/projects/rocprofiler-sdk/cmake/Modules/rocprofiler-sdk-utilities.cmake index 69573ad447ae..1e1fc0a4c500 100644 --- a/projects/rocprofiler-sdk/cmake/Modules/rocprofiler-sdk-utilities.cmake +++ b/projects/rocprofiler-sdk/cmake/Modules/rocprofiler-sdk-utilities.cmake @@ -91,7 +91,8 @@ function(rocprofiler_sdk_pc_sampling_stochastic_disabled _VAR) list(GET rocprofiler-sdk-tests-gfx-info 0 pc-sampling-gpu-0-gfx-info) if("${pc-sampling-gpu-0-gfx-info}" MATCHES "^gfx94[0-9]$" - OR "${pc-sampling-gpu-0-gfx-info}" MATCHES "^gfx95[0-9]$") + OR "${pc-sampling-gpu-0-gfx-info}" MATCHES "^gfx95[0-9]$" + OR "${pc-sampling-gpu-0-gfx-info}" MATCHES "^gfx1250$") # PC sampling is enabled on this architecture. set(${_VAR} FALSE diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/codeobj/disassembly.hpp b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/codeobj/disassembly.hpp index c054143e62af..4b277bffb368 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/codeobj/disassembly.hpp +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/codeobj/disassembly.hpp @@ -71,6 +71,44 @@ namespace codeobj { namespace disassembly { +struct FallbackOp +{ + const char* str{}; + size_t size{}; +}; + +inline FallbackOp +get_fallback_op_gfx1250(uint32_t header) +{ + if((header >> 7) == 0) return {"v_vop_generic ", 4}; + + switch(header) + { + case 0b11001100: return {"v_vop3p_generic ", 8}; + case 0b11001111: return {"v_vopd3_generic ", 12}; + default: break; + } + + header >>= 2; + + switch(header) + { + case 0b110001: return {"buffer_load_generic ", 12}; + case 0b111011: return {"global_load_generic ", 12}; + case 0b110100: return {"image_load_generic ", 12}; + case 0b110010: return {"v_vopd_generic ", 8}; + case 0b110101: return {"v_vopsd_generic ", 8}; + case 0b110110: return {"ds_generic ", 8}; + case 0b111101: return {"s_load_generic ", 8}; + default: break; + } + + header >>= 4; + if(header == 0b10) return {"s_sop ", 4}; + + return {nullptr, 0}; +} + class CodeObjectBinary { public: @@ -206,6 +244,8 @@ class DisassemblyInstance &DisassemblyInstance::inst_callback, [](uint64_t, void*) {}, &info)); + + if(input_isa.find("gfx1250") != std::string::npos) gfxip = 1250; } ~DisassemblyInstance() { @@ -218,8 +258,25 @@ class DisassemblyInstance uint64_t size_read; uint64_t addr_in_buffer = reinterpret_cast(buffer.data()) + faddr; - THROW_COMGR( - amd_comgr_disassemble_instruction(info, addr_in_buffer, (void*) this, &size_read)); + auto _status = + amd_comgr_disassemble_instruction(info, addr_in_buffer, (void*) this, &size_read); + if(_status != AMD_COMGR_STATUS_SUCCESS) + { + if(faddr + 4 > buffer.size()) THROW_COMGR(_status); + + uint32_t read = 0; + std::memcpy(&read, buffer.data() + faddr, 4); + + FallbackOp fallback{}; + + if(gfxip == 1250) fallback = get_fallback_op_gfx1250(read >> 24); + + if(fallback.str == nullptr || fallback.size == 0) THROW_COMGR(_status); + + this->last_instruction = fallback.str; + size_read = fallback.size; + } + return {std::move(this->last_instruction), size_read}; } @@ -295,6 +352,7 @@ class DisassemblyInstance return offset; } + int gfxip{}; std::vector buffer{}; std::string last_instruction{}; amd_comgr_disassembly_info_t info{}; diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization/save.hpp b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization/save.hpp index 2410e69b0387..f60a72a2cc8a 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization/save.hpp +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization/save.hpp @@ -1058,6 +1058,9 @@ save(ArchiveT& ar, rocprofiler_pc_sampling_snapshot_v0_t data) ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_exp", arb_state_stall_exp); ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_misc", arb_state_stall_misc); ROCP_SDK_SAVE_DATA_BITFIELD("arb_state_stall_brmsg", arb_state_stall_brmsg); + + // sampling lock error + ROCP_SDK_SAVE_DATA_BITFIELD("lck_err", sampling_lock_error); } template @@ -1070,6 +1073,9 @@ save(ArchiveT& ar, rocprofiler_pc_sampling_memory_counters_t data) ROCP_SDK_SAVE_DATA_BITFIELD("sample_cnt", sample_cnt); ROCP_SDK_SAVE_DATA_BITFIELD("ds_cnt", ds_cnt); ROCP_SDK_SAVE_DATA_BITFIELD("km_cnt", km_cnt); + ROCP_SDK_SAVE_DATA_BITFIELD("async_cnt", async_cnt); + ROCP_SDK_SAVE_DATA_BITFIELD("tensor_cnt", tensor_cnt); + ROCP_SDK_SAVE_DATA_BITFIELD("xnack_cnt", xnack_cnt); } template diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/pc_sampling.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/pc_sampling.h index a28dfb2b975a..ce5647cd19a8 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/pc_sampling.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/pc_sampling.h @@ -388,9 +388,9 @@ typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_snapshot_v0_ uint32_t arb_state_stall_brmsg : 1; ///< branch/message instruction was stalled uint32_t arb_state_state_reserved : 1; ///< reserved for the future use // We have two reserved bits - uint32_t dual_issue_valu : 1; - uint32_t reserved1 : 1; ///< reserved for the future use - uint32_t reserved2 : 3; ///< reserved for the future use + uint32_t dual_issue_valu : 1; + uint32_t sampling_lock_error : 1; + uint32_t reserved2 : 3; ///< reserved for the future use /// @var reason_not_issued /// @brief The reason for not issuing an instruction. The field takes one of the value defined @@ -399,6 +399,10 @@ typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_snapshot_v0_ /// @brief VALU instruction was stalled when a sample was generated /// @var dual_issue_valu /// @brief Two VALU instructions were issued for coexecution (MI3xx specific) + /// @var sampling_lock_error + /// @brief At least one wave was locked out from taking a sample, + /// due to the latency instroduced by current sample read. + /// Too many samples with this bit on indicates that the sampling frequency is too high } rocprofiler_pc_sampling_snapshot_v0_t; /** @@ -406,12 +410,17 @@ typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_snapshot_v0_ */ typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_memory_counters_t { - uint32_t load_cnt : 6; - uint32_t store_cnt : 6; - uint32_t bvh_cnt : 3; - uint32_t sample_cnt : 6; - uint32_t ds_cnt : 6; - uint32_t km_cnt : 5; + uint64_t load_cnt : 6; + uint64_t store_cnt : 6; + uint64_t bvh_cnt : 3; + uint64_t sample_cnt : 6; + uint64_t ds_cnt : 6; + uint64_t km_cnt : 5; + // The following counters are relevant for gfx1250 only. + uint64_t async_cnt : 6; + uint64_t tensor_cnt : 6; + uint64_t xnack_cnt : 6; + uint64_t reserved : 14; ///< reserved for the future use /// @var load_cnt /// @brief Counts the number of VMEM load instructions issued but not yet completed. @@ -426,6 +435,13 @@ typedef struct ROCPROFILER_SDK_EXPERIMENTAL rocprofiler_pc_sampling_memory_count /// @var km_cnt /// @brief Counts the number of scalar memory reads and memory instructions issued but not yet /// completed. + /// @var async_cnt + /// @brief Counts the number of async instructions issued but not yet completed. + /// @var tensor_cnt + /// @brief Counts the number of tensor instructions issued but not yet completed. + /// @var xnack_cnt + /// @brief Counts the number of oustanding memory instruction not yet reported + /// XNACK ackwlodgement. } rocprofiler_pc_sampling_memory_counters_t; /** diff --git a/projects/rocprofiler-sdk/source/lib/att-tool/waitcnt/gfx12.cpp b/projects/rocprofiler-sdk/source/lib/att-tool/waitcnt/gfx12.cpp index 0474df55b995..18d16c5bfd4d 100644 --- a/projects/rocprofiler-sdk/source/lib/att-tool/waitcnt/gfx12.cpp +++ b/projects/rocprofiler-sdk/source/lib/att-tool/waitcnt/gfx12.cpp @@ -45,6 +45,8 @@ enum MemoryInstType TYPE_FLAT_LOAD, TYPE_FLAT_STOR, TYPE_BVH, + TYPE_TENSOR, + TYPE_ASYNC, TYPE_WAITCNT }; @@ -65,6 +67,8 @@ union MemoryInst int stcnt : 1; int sampl : 1; int idle : 1; + int tensr : 1; + int async : 1; }; int raw = 0; }; @@ -82,14 +86,14 @@ classify(const std::string& inst) { if(inst.find("s_wait_alu") != npos) return MemoryInstType::TYPE_NOT_MEM; + MemoryInst type = MemoryInstType::TYPE_WAITCNT; + if(inst.find("s_wait_idle") != npos) { - MemoryInst type = MemoryInstType::TYPE_WAITCNT; - type.idle = true; + type.idle = true; return type; } - MemoryInst type = MemoryInstType::TYPE_WAITCNT; if(inst.find("dscnt") != npos) type.dscnt = true; if(inst.find("bvhcnt") != npos) type.bvhcn = true; if(inst.find("expcnt") != npos) type.expcn = true; @@ -97,6 +101,8 @@ classify(const std::string& inst) if(inst.find("loadcnt") != npos) type.ldcnt = true; if(inst.find("storecnt") != npos) type.stcnt = true; if(inst.find("samplecnt") != npos) type.sampl = true; + if(inst.find("asynccnt") != npos) type.async = true; + if(inst.find("tensorcnt") != npos) type.tensr = true; return type; } @@ -130,6 +136,17 @@ classify(const std::string& inst) else return MemoryInstType::TYPE_LDS; } + else if(inst.find("tensor_") <= 1) + { + return MemoryInstType::TYPE_TENSOR; + } + else if(inst.find("_async_") != npos) + { + if(inst.find("global_load_async_") == 0 || inst.find("global_store_async_") == 0) + { + return MemoryInstType::TYPE_ASYNC; + } + } bool bStore = inst.find("store") != npos || inst.find("global_wb") == 0; @@ -155,6 +172,8 @@ WaitcntList::gfx12_construct(const wave_t& wave, isa_map_t& isa_map) MemoryCounter kmcnt("kmcnt"); MemoryCounter expcnt("expcnt"); MemoryCounter bvhcnt("bvhcnt"); + MemoryCounter async("asynccnt"); + MemoryCounter tensor("tensorcnt"); std::vector flat_load{}; std::vector flat_stor{}; @@ -231,6 +250,14 @@ WaitcntList::gfx12_construct(const wave_t& wave, isa_map_t& isa_map) { bvhcnt.list.push_back(line_number); } + else if(type.inst == MemoryInstType::TYPE_TENSOR) + { + tensor.list.push_back(line_number); + } + else if(type.inst == MemoryInstType::TYPE_ASYNC) + { + async.list.push_back(line_number); + } else if(type.inst == MemoryInstType::TYPE_WAITCNT) { if(type.bvhcn) @@ -263,6 +290,16 @@ WaitcntList::gfx12_construct(const wave_t& wave, isa_map_t& isa_map) if(auto joined = loadcnt.handle_mem_op(inst_str, flat_load)) mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); } + if(type.tensr) + { + if(auto joined = tensor.handle_mem_op(inst_str, empty_list)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } + if(type.async) + { + if(auto joined = async.handle_mem_op(inst_str, empty_list)) + mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(*joined)}); + } if(type.dscnt) { if(auto joined = dscnt.handle_mem_op(inst_str, flat_load)) @@ -287,6 +324,8 @@ WaitcntList::gfx12_construct(const wave_t& wave, isa_map_t& isa_map) kmcnt.clearTo(all); expcnt.clearTo(all); bvhcnt.clearTo(all); + tensor.clearTo(all); + async.clearTo(all); mem_unroll.emplace_back(LineWaitcnt{line_number, std::move(all)}); } diff --git a/projects/rocprofiler-sdk/source/lib/att-tool/waitcnt/tests/att_decoder_waitcnt_test.cpp b/projects/rocprofiler-sdk/source/lib/att-tool/waitcnt/tests/att_decoder_waitcnt_test.cpp index 2ea0378957dc..62ecbe7518d6 100644 --- a/projects/rocprofiler-sdk/source/lib/att-tool/waitcnt/tests/att_decoder_waitcnt_test.cpp +++ b/projects/rocprofiler-sdk/source/lib/att-tool/waitcnt/tests/att_decoder_waitcnt_test.cpp @@ -292,6 +292,16 @@ TEST(att_decoder_waitcnt_test, gfx12) append_isa(42, "s_wait_idle"); append_isa(43, "invalid"); + append_isa(44, "global_load_async_"); + append_isa(45, "global_store_async_"); + + append_isa(46, "tensor_load"); + append_isa(47, "tensor_store"); + append_isa(48, "tensor_save"); + + append_isa(49, "s_wait_tensorcnt 0"); + append_isa(50, "s_wait_asynccnt 0"); + std::vector insts{}; for(size_t i = 0; i < isa_map.size(); i++) { @@ -321,7 +331,7 @@ TEST(att_decoder_waitcnt_test, gfx12) ASSERT_EQ(dependencies.at(dep).size(), set.size()); }; - ASSERT_EQ(dependencies.size(), 12); + ASSERT_EQ(dependencies.size(), 14); set_equal(6, {2, 3}); set_equal(7, {4}); set_equal(8, {5}); @@ -334,6 +344,8 @@ TEST(att_decoder_waitcnt_test, gfx12) set_equal(38, {32, 33, 34, 35}); set_equal(39, {36}); set_equal(42, {40, 41}); + set_equal(49, {46, 47, 48}); + set_equal(50, {44, 45}); } TEST(att_decoder_waitcnt_test, fail_conditions) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp index 319b69ec987e..1c94a0b4f26e 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -1386,7 +1386,12 @@ construct_counter_collection_profile(rocprofiler_agent_id_t agent_id, // search the gpu agent counter info for a counter with a matching name bool counter_found = false; - for(const auto& citr : gpu_agents_counter_info.at(agent_id)) + auto counter_vec = gpu_agents_counter_info.find(agent_id); + ROCP_FATAL_IF(counter_vec == gpu_agents_counter_info.end()) + << "No counter information found for agent " << agent_v->node_id << " (gpu-" + << agent_v->gpu_index << ", " << agent_v->name << "). Unable to find counter: " << itr; + + for(const auto& citr : counter_vec->second) { if(name_v == std::string_view{citr.name}) { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp index 3d9cf41a5ab3..b48f4407ee2a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/evaluate_ast.cpp @@ -31,6 +31,8 @@ #include "lib/rocprofiler-sdk/counters/parser/raw_ast.hpp" #include "lib/rocprofiler-sdk/counters/parser/reader.hpp" +#include + #include #include @@ -140,16 +142,13 @@ perform_reduction( } std::unordered_map> rec_groups; - size_t bit_length = DIM_BIT_LENGTH / ROCPROFILER_DIMENSION_LAST; for(auto& rec : *input_array) { for(auto dim : _reduce_dimension_set) { - int64_t mask_dim = (MAX_64 >> (64 - bit_length)) << ((dim - 1) * bit_length); - - rec.id = rec.id | mask_dim; - rec.id = rec.id ^ mask_dim; + uint64_t mask = get_dim_mask(dim); + rec.id = (rec.id | mask) ^ mask; } rec_groups[rec.id].push_back(rec); } @@ -173,42 +172,57 @@ perform_reduction( return input_array; } -int64_t -get_int_encoded_dimensions_from_string(const std::string& rangeStr) +/** + * @brief Parse dimension selection string to get a single selected index. + * + * This function parses a string containing a single dimension index value. + * Only single index selection is supported (e.g., "5"). Range-based selection + * (e.g., "1:4") and multiple values (e.g., "1,2,3") are not supported. + * + * @param indexStr The string containing the dimension index + * @param dim The dimension type for validation + * @return size_t The selected dimension index value + */ +size_t +parse_dimension_selection(const std::string& indexStr, + rocprofiler_profile_counter_instance_types dim) { - int64_t result = 0; - std::istringstream iss(rangeStr); - std::string token; - size_t bit_length = DIM_BIT_LENGTH / ROCPROFILER_DIMENSION_LAST; + auto token = sdk::parse::strip(std::string{indexStr}, " \t\n\r"); - while(std::getline(iss, token, ',')) + if(token.empty()) { - token.erase(std::remove_if(token.begin(), token.end(), ::isspace), token.end()); + throw std::runtime_error( + fmt::format("Empty dimension selection for dimension type {}.", static_cast(dim))); + } - size_t dash_pos = token.find(':'); - if(dash_pos != std::string::npos) - { - throw std::runtime_error( - fmt::format("Range based selection not supported by Dimension API. only select " - "single value for each dimension.")); - int start = std::stoi(token.substr(0, dash_pos)); - int end = std::stoi(token.substr(dash_pos + 1)); - result |= (1LL << std::min(64, end + 1)) - (1LL << std::max(start, 0)); - } - else - { - int num = std::stoi(token); - if(num < (1 << bit_length)) - { - result |= (1LL << num); - } - else - { - throw std::runtime_error(fmt::format("Dimension value exceeds max allowed.")); - } - } + if(token.find(':') != std::string::npos || token.find(',') != std::string::npos) + { + throw std::runtime_error( + fmt::format("Range or multi-value selection not supported by Dimension API. " + "Only select a single index for each dimension.")); } - return result; + + auto bit_length = get_dim_bit_length(dim); + auto max_value = (size_t{1} << bit_length) - 1; + + int num = sdk::parse::from_string(token); + if(num < 0) + { + throw std::runtime_error(fmt::format( + "Dimension value {} is negative for dimension type {}.", num, static_cast(dim))); + } + + auto unum = static_cast(num); + if(unum > max_value) + { + throw std::runtime_error( + fmt::format("Dimension value {} exceeds max allowed {} for dimension type {}.", + num, + max_value, + static_cast(dim))); + } + + return unum; } std::vector* @@ -216,23 +230,24 @@ perform_selection(std::map* input_array) { if(input_array->empty()) return input_array; + for(auto& dim_pair : dimension_map) { - int64_t encoded_dim_values = get_int_encoded_dimensions_from_string(dim_pair.second); - size_t bit_length = DIM_BIT_LENGTH / ROCPROFILER_DIMENSION_LAST; - int64_t mask = (MAX_64 >> (64 - bit_length)) << ((dim_pair.first - 1) * bit_length); + size_t selected_index = parse_dimension_selection(dim_pair.second, dim_pair.first); + uint64_t mask = get_dim_mask(dim_pair.first); input_array->erase(std::remove_if(input_array->begin(), input_array->end(), [&](rocprofiler_counter_record_t& rec) { - bool should_remove = - (encoded_dim_values & - (1 << rocprofiler::counters::rec_to_dim_pos( - rec.id, dim_pair.first))) == 0; + size_t dim_val = + rec_to_dim_pos(rec.id, dim_pair.first); + + bool should_remove = (dim_val != selected_index); + if(!should_remove) { - rec.id = rec.id | mask; - rec.id = rec.id ^ mask; + // Clear the dimension bits in the record + rec.id = (rec.id | mask) ^ mask; } return should_remove; }), @@ -580,7 +595,7 @@ EvaluateAST::get_required_counters(const std::unordered_mapget_required_counters(asts, counters); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.hpp index 294532863986..e1e94b964661 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/id_decode.hpp @@ -51,6 +51,69 @@ enum rocprofiler_profile_counter_instance_types ROCPROFILER_DIMENSION_LAST }; +// Variable bit allocation for each dimension type +// Layout (48 bits total for dimensions): +// XCC: 6 bits (supports up to 64) +// AID: 6 bits (supports up to 64) +// SHADER_ENGINE: 6 bits (supports up to 64) +// AGENT: 6 bits (supports up to 64) +// SHADER_ARRAY: 6 bits (supports up to 64) +// WGP: 6 bits (supports up to 64) +// INSTANCE: 10 bits (supports up to 1024 instances) +// Total used: 46 bits, 2 bits reserved +constexpr uint8_t DIM_BIT_LENGTHS[ROCPROFILER_DIMENSION_LAST] = { + 0, // DIMENSION_NONE + 6, // DIMENSION_XCC + 6, // DIMENSION_AID + 6, // DIMENSION_SHADER_ENGINE + 6, // DIMENSION_AGENT + 6, // DIMENSION_SHADER_ARRAY + 6, // DIMENSION_WGP + 10, // DIMENSION_INSTANCE +}; + +// Precomputed bit offsets for each dimension +// Offset is the starting bit position within the 48-bit dimension field +constexpr uint8_t DIM_BIT_OFFSETS[ROCPROFILER_DIMENSION_LAST] = { + 0, // DIMENSION_NONE + 0, // DIMENSION_XCC: bits 0-5 + 6, // DIMENSION_AID: bits 6-11 + 12, // DIMENSION_SHADER_ENGINE: bits 12-17 + 18, // DIMENSION_AGENT: bits 18-23 + 24, // DIMENSION_SHADER_ARRAY: bits 24-29 + 30, // DIMENSION_WGP: bits 30-35 + 36, // DIMENSION_INSTANCE: bits 36-45 +}; + +// Compile-time validation that dimension bit allocation does not exceed 48 bits +static_assert(DIM_BIT_OFFSETS[ROCPROFILER_DIMENSION_INSTANCE] + + DIM_BIT_LENGTHS[ROCPROFILER_DIMENSION_INSTANCE] <= + DIM_BIT_LENGTH, + "Dimension bit allocation exceeds 48-bit limit"); + +// Helper function to get bit length for a dimension +constexpr uint64_t +get_dim_bit_length(rocprofiler_profile_counter_instance_types dim) +{ + return (dim < ROCPROFILER_DIMENSION_LAST) ? DIM_BIT_LENGTHS[dim] : 0; +} + +// Helper function to get bit offset for a dimension +constexpr uint64_t +get_dim_bit_offset(rocprofiler_profile_counter_instance_types dim) +{ + return (dim < ROCPROFILER_DIMENSION_LAST) ? DIM_BIT_OFFSETS[dim] : 0; +} + +// Helper function to get the bitmask for a dimension (shifted to correct position) +constexpr uint64_t +get_dim_mask(rocprofiler_profile_counter_instance_types dim) +{ + uint64_t bit_length = get_dim_bit_length(dim); + uint64_t bit_offset = get_dim_bit_offset(dim); + return ((1ULL << bit_length) - 1) << bit_offset; +} + using DimensionMap = std::unordered_map; @@ -103,27 +166,25 @@ rocprofiler::counters::set_dim_in_rec(rocprofiler_counter_instance_id_t& rocprofiler_profile_counter_instance_types dim, size_t value) { - uint64_t bit_length = DIM_BIT_LENGTH / ROCPROFILER_DIMENSION_LAST; - if(dim == ROCPROFILER_DIMENSION_NONE) { // Set all 48 bits of dimension - id = (id & ~(MAX_64 >> COUNTER_BIT_LENGTH)) | value; - bit_length = DIM_BIT_LENGTH; - } - else - { - uint64_t mask = (MAX_64 >> (BITS_IN_UINT64 - bit_length)) << ((dim - 1) * bit_length); - // Reset bits to 0 for dimension. Does so by getting the bit length as F's then - // shifiting that into the position of dim. Not's that value and then and's it - // with id. - id = (id & ~(mask)); - // Set the value for the dimenstion - id = id | (value << ((dim - 1) * bit_length)); + id = (id & ~(MAX_64 >> COUNTER_BIT_LENGTH)) | value; + CHECK(value <= (MAX_64 >> COUNTER_BIT_LENGTH)) << "Dimension value exceeds max allowed"; + return; } - CHECK(value <= (MAX_64 >> (BITS_IN_UINT64 - bit_length))) - << "Dimension value exceeds max allowed"; + uint64_t bit_length = get_dim_bit_length(dim); + uint64_t bit_offset = get_dim_bit_offset(dim); + uint64_t mask = get_dim_mask(dim); + + CHECK(bit_length > 0) << "Invalid dimension type"; + CHECK(value <= ((1ULL << bit_length) - 1)) + << "Dimension value " << value << " exceeds max allowed for dimension " + << static_cast(dim) << " (max: " << ((1ULL << bit_length) - 1) << ")"; + + // Reset bits to 0 for dimension, then set the value + id = (id & ~mask) | (value << bit_offset); } void @@ -152,9 +213,10 @@ rocprofiler::counters::rec_to_dim_pos(rocprofiler_counter_instance_id_t return id & (MAX_64 >> COUNTER_BIT_LENGTH); } - size_t bit_length = DIM_BIT_LENGTH / ROCPROFILER_DIMENSION_LAST; - id = id & ((MAX_64 >> (BITS_IN_UINT64 - bit_length)) << ((dim - 1) * bit_length)); - return id >> ((dim - 1) * bit_length); + uint64_t bit_offset = get_dim_bit_offset(dim); + uint64_t mask = get_dim_mask(dim); + + return (id & mask) >> bit_offset; } // Counter ID encoding/decoding implementations diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp index 61fc8644d36b..71538d352b5e 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/dimension.cpp @@ -82,10 +82,8 @@ check_counter_id(rocprofiler_counter_instance_id_t id, uint64_t expected_base_me TEST(dimension, set_get) { using namespace rocprofiler::counters; - int64_t max_counter_val = (std::numeric_limits::max() >> - (64 - (DIM_BIT_LENGTH / ROCPROFILER_DIMENSION_LAST))); - rocprofiler_counter_instance_id_t test_id = 0; - rocprofiler_counter_id_t test_counter{.handle = 123}; + auto test_id = rocprofiler_counter_instance_id_t{0}; + auto test_counter = rocprofiler_counter_id_t{.handle = 123}; set_counter_in_rec(test_id, test_counter); // 0x007B000000000000 = decimal counter id 123 << DIM_BIT_LENGTH @@ -104,29 +102,27 @@ TEST(dimension, set_get) for(size_t i = 1; i < static_cast(ROCPROFILER_DIMENSION_LAST); i++) { auto dim = static_cast(i); + set_dim_in_rec(test_id, dim, i); check_dim_pos(test_id, dim, i); set_dim_in_rec(test_id, dim, i * multi_factor); for(size_t j = 1; j < static_cast(ROCPROFILER_DIMENSION_LAST); j++) { if(i == j) continue; - set_dim_in_rec(test_id, - static_cast(j), - max_counter_val); - check_dim_pos(test_id, - static_cast(j), - max_counter_val); + auto other_dim = static_cast(j); + uint64_t other_max_val = (1ULL << get_dim_bit_length(other_dim)) - 1; + set_dim_in_rec(test_id, other_dim, other_max_val); + check_dim_pos(test_id, other_dim, other_max_val); check_dim_pos(test_id, dim, i * multi_factor); } for(size_t j = static_cast(ROCPROFILER_DIMENSION_LAST - 1); j > 0; j--) { if(i == j) continue; - set_dim_in_rec(test_id, - static_cast(j), - max_counter_val); - check_dim_pos( - test_id, (rocprofiler_profile_counter_instance_types) j, max_counter_val); + auto other_dim = static_cast(j); + uint64_t other_max_val = (1ULL << get_dim_bit_length(other_dim)) - 1; + set_dim_in_rec(test_id, other_dim, other_max_val); + check_dim_pos(test_id, other_dim, other_max_val); check_dim_pos(test_id, dim, i * multi_factor); } @@ -139,12 +135,15 @@ TEST(dimension, set_get) for(size_t i = static_cast(ROCPROFILER_DIMENSION_LAST - 1); i > 0; i--) { - auto dim = static_cast(i); + auto dim = static_cast(i); + uint64_t dim_max_val = (1ULL << get_dim_bit_length(dim)) - 1; + set_dim_in_rec(test_id, dim, i * 5); check_dim_pos(test_id, dim, i * 5); set_dim_in_rec(test_id, dim, i * 3); check_dim_pos(test_id, dim, i * 3); - for(size_t j = 1; j < 64; j++) + // Test all values up to the max for this dimension + for(size_t j = 1; j <= dim_max_val; j++) { test_id = 0; set_dim_in_rec(test_id, dim, j); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp index 05ec89af428b..9410c341a5d0 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/evaluate_ast_test.cpp @@ -1312,12 +1312,10 @@ TEST(evatuate_ast, evaluate_select) } for(auto& dim_pair : dims) { - size_t bit_length = DIM_BIT_LENGTH / ROCPROFILER_DIMENSION_LAST; - int64_t mask = (MAX_64 >> (64 - bit_length)) << ((dim_pair.first - 1) * bit_length); + uint64_t mask = get_dim_mask(dim_pair.first); for(auto& rec : a) { - rec.id = rec.id | mask; - rec.id = rec.id ^ mask; + rec.id = (rec.id | mask) ^ mask; } } return a; @@ -1421,8 +1419,6 @@ TEST(evaluate_ast, counter_reduction_dimension) { using namespace rocprofiler::counters; - size_t bit_length = DIM_BIT_LENGTH / ROCPROFILER_DIMENSION_LAST; - auto get_base_rec_id = [](uint64_t counter_id) { rocprofiler_counter_instance_id_t base_id = 0; set_counter_in_rec(base_id, {.handle = counter_id}); @@ -1435,10 +1431,8 @@ TEST(evaluate_ast, counter_reduction_dimension) std::vector result; for(auto rec : a) { - int64_t mask_dim = (MAX_64 >> (64 - bit_length)) << (bit_length * 0); - - rec.id = rec.id | mask_dim; - rec.id = rec.id ^ mask_dim; + uint64_t mask = get_dim_mask(ROCPROFILER_DIMENSION_XCC); + rec.id = (rec.id | mask) ^ mask; if(groups_dim.find(rec.id) == groups_dim.end()) { groups_dim[rec.id] = rec; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/abi.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/abi.cpp index 6faef01d575a..b04be51dbaf8 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/abi.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/abi.cpp @@ -322,6 +322,9 @@ ROCP_SDK_ENFORCE_ABI(::AmdExtTable, hsa_amd_agent_preload_fn, 80); #if HSA_AMD_EXT_API_TABLE_STEP_VERSION >= 0x0C ROCP_SDK_ENFORCE_ABI(::AmdExtTable, hsa_amd_svm_discard_batch_async_fn, 81); #endif +#if HSA_AMD_EXT_API_TABLE_STEP_VERSION >= 0x0D +ROCP_SDK_ENFORCE_ABI(::AmdExtTable, hsa_amd_signal_get_event_id_fn, 82); +#endif ROCP_SDK_ENFORCE_ABI(::ImageExtTable, hsa_ext_image_get_capability_fn, 1); ROCP_SDK_ENFORCE_ABI(::ImageExtTable, hsa_ext_image_data_get_info_fn, 2); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/details/fmt.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/details/fmt.hpp index 8c48c767f987..8ccc45566d01 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/details/fmt.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/details/fmt.hpp @@ -83,6 +83,69 @@ struct formatter } }; +#if HSA_AMD_EXT_API_TABLE_STEP_VERSION >= 0x0D +template <> +struct formatter +{ + template + constexpr auto parse(ParseContext& ctx) + { + return ctx.begin(); + } + + template + auto format(hsa_amd_ext_perf_hint_t const& hint, Ctx& ctx) const + { + return fmt::format_to(ctx.out(), + "[EXT_PERF_HINT, group_mem_carveout={}, reverse_dispatch_order={}, " + "hint_val={}]", + hint.group_mem_carveout, + hint.reverse_dispatch_order, + hint.hint_val); + } +}; + +template <> +struct formatter +{ + template + constexpr auto parse(ParseContext& ctx) + { + return ctx.begin(); + } + + template + auto format(hsa_amd_ext_kernel_dispatch_packet_t const& pkt, Ctx& ctx) const + { + return fmt::format_to(ctx.out(), + "[EXT_KERNEL_DISPATCH, header={}, amd_format={}, setup={}, " + "workgroup_size=[{}, {}, {}], cluster_count=[{}, {}, {}], " + "cluster_size=[{}, {}, {}], perf_hint={}, private_size={}, " + "group_size={}, kernel_object={:x}, kern_arg={}, dep_signal={}, " + "completion_signal={}]", + pkt.header, + pkt.amd_format, + pkt.setup, + pkt.workgroup_size_x, + pkt.workgroup_size_y, + pkt.workgroup_size_z, + pkt.cluster_count_x, + pkt.cluster_count_y, + pkt.cluster_count_z, + pkt.cluster_size_x, + pkt.cluster_size_y, + pkt.cluster_size_z, + pkt.perf_hint, + pkt.private_segment_size, + pkt.group_segment_size, + pkt.kernel_object, + pkt.kernarg_address, + pkt.dep_signal.handle, + pkt.completion_signal.handle); + } +}; +#endif + template <> struct formatter { @@ -157,7 +220,11 @@ struct formatter switch(t) { case 0: - // PM4 packet +#if HSA_AMD_EXT_API_TABLE_STEP_VERSION >= 0x0D + // Vendor specific packet - check AMD format + if(pkt.ext_kernel_dispatch.amd_format == HSA_AMD_PACKET_TYPE_EXT_KERNEL_DISPATCH) + return fmt::format_to(ctx.out(), "{}", pkt.ext_kernel_dispatch); +#endif return fmt::format_to(ctx.out(), "{}", pkt.ext_amd_aql_pm4); case 2: // Kernel dispatch diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp index e63684d91cec..5433d51c5898 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -68,6 +68,11 @@ static_assert(offsetof(hsa_ext_amd_aql_pm4_packet_t, completion_signal) == static_assert(offsetof(hsa_ext_amd_aql_pm4_packet_t, completion_signal) == offsetof(hsa_barrier_or_packet_t, completion_signal), "unexpected ABI incompatibility"); +#if HSA_AMD_EXT_API_TABLE_STEP_VERSION >= 0x0D +static_assert(offsetof(hsa_ext_amd_aql_pm4_packet_t, completion_signal) == + offsetof(hsa_amd_ext_kernel_dispatch_packet_t, completion_signal), + "unexpected ABI incompatibility"); +#endif namespace rocprofiler { @@ -350,6 +355,16 @@ WriteInterceptor(const void* packets, { ++num_dispatch_packets; } +#if HSA_AMD_EXT_API_TABLE_STEP_VERSION >= 0x0D + else if(packet_type == HSA_PACKET_TYPE_VENDOR_SPECIFIC) + { + const auto& ext_packet = packets_arr[i].ext_kernel_dispatch; + if(ext_packet.amd_format == HSA_AMD_PACKET_TYPE_EXT_KERNEL_DISPATCH) + { + ++num_dispatch_packets; + } + } +#endif } if(num_dispatch_packets == 0) @@ -438,7 +453,21 @@ WriteInterceptor(const void* packets, bit_extract(original_packet.header, HSA_PACKET_HEADER_TYPE, HSA_PACKET_HEADER_TYPE + HSA_PACKET_HEADER_WIDTH_TYPE - 1); - if(packet_type != HSA_PACKET_TYPE_KERNEL_DISPATCH) + bool is_kernel_dispatch = (packet_type == HSA_PACKET_TYPE_KERNEL_DISPATCH); + bool is_ext_kernel_dispatch = false; + +#if HSA_AMD_EXT_API_TABLE_STEP_VERSION >= 0x0D + if(packet_type == HSA_PACKET_TYPE_VENDOR_SPECIFIC) + { + const auto& ext_packet = _packets[i].ext_kernel_dispatch; + if(ext_packet.amd_format == HSA_AMD_PACKET_TYPE_EXT_KERNEL_DISPATCH) + { + is_ext_kernel_dispatch = true; + } + } +#endif + + if(!is_kernel_dispatch && !is_ext_kernel_dispatch) { transformed_packets.emplace_back(_packets[i]); continue; @@ -461,9 +490,55 @@ WriteInterceptor(const void* packets, ROCPROFILER_KERNEL_DISPATCH_ENQUEUE, internal_corr_id); - const auto original_completion_signal = original_packet.completion_signal; + // Lambda to extract packet info regardless of packet type + auto extract_packet_info = [](const rocprofiler_packet& pkt, bool is_ext) { + struct packet_info + { + hsa_signal_t completion_signal; + uint64_t kernel_object; + uint32_t private_segment_size; + uint32_t group_segment_size; + rocprofiler_dim3_t workgroup_size; + rocprofiler_dim3_t grid_size; + }; + +#if HSA_AMD_EXT_API_TABLE_STEP_VERSION >= 0x0D + if(is_ext) + { + const auto& e = pkt.ext_kernel_dispatch; + return packet_info{e.completion_signal, + e.kernel_object, + e.private_segment_size, + e.group_segment_size, + {e.workgroup_size_x, e.workgroup_size_y, e.workgroup_size_z}, + {static_cast(e.cluster_count_x) * + static_cast(e.cluster_size_x) * + static_cast(e.workgroup_size_x), + static_cast(e.cluster_count_y) * + static_cast(e.cluster_size_y) * + static_cast(e.workgroup_size_y), + static_cast(e.cluster_count_z) * + static_cast(e.cluster_size_z) * + static_cast(e.workgroup_size_z)}}; + } +#else + (void) is_ext; +#endif + { + const auto& s = pkt.kernel_dispatch; + return packet_info{s.completion_signal, + s.kernel_object, + s.private_segment_size, + s.group_segment_size, + {s.workgroup_size_x, s.workgroup_size_y, s.workgroup_size_z}, + {s.grid_size_x, s.grid_size_y, s.grid_size_z}}; + } + }; + + const auto pkt_info = extract_packet_info(_packets[i], is_ext_kernel_dispatch); + const auto original_completion_signal = pkt_info.completion_signal; const bool existing_completion_signal = (original_completion_signal.handle != 0); - const uint64_t kernel_id = code_object::get_kernel_id(original_packet.kernel_object); + const uint64_t kernel_id = code_object::get_kernel_id(pkt_info.kernel_object); // Copy kernel pkt, copy is to allow for signal to be modified _packet_data.kernel_packet = _packets[i]; @@ -473,8 +548,14 @@ WriteInterceptor(const void* packets, // create our own signal that we can get a callback on. if there is an original // completion signal we will create a barrier packet, assign the original completion // signal that that barrier packet, and add it right after the kernel packet - _packet_data.pooled_signal = - queue.create_signal(0, &kernel_packet.kernel_dispatch.completion_signal, true); +#if HSA_AMD_EXT_API_TABLE_STEP_VERSION >= 0x0D + if(is_ext_kernel_dispatch) + _packet_data.pooled_signal = queue.create_signal( + 0, &kernel_packet.ext_kernel_dispatch.completion_signal, true); + else +#endif + _packet_data.pooled_signal = + queue.create_signal(0, &kernel_packet.kernel_dispatch.completion_signal, true); // computes the "size" based on the offset of reserved_padding field constexpr auto kernel_dispatch_info_rt_size = @@ -483,27 +564,22 @@ WriteInterceptor(const void* packets, static_assert(kernel_dispatch_info_rt_size < sizeof(rocprofiler_kernel_dispatch_info_t), "failed to compute size field based on offset of reserved_padding field"); - auto dispatch_id = ++sequence_counter; - _packet_data.callback_record = callback_record_t{ - sizeof(callback_record_t), - rocprofiler_timestamp_t{0}, - rocprofiler_timestamp_t{0}, - rocprofiler_kernel_dispatch_info_t{ - .size = kernel_dispatch_info_rt_size, - .agent_id = queue.get_agent().get_rocp_agent()->id, - .queue_id = queue.get_id(), - .kernel_id = kernel_id, - .dispatch_id = dispatch_id, - .private_segment_size = kernel_packet.kernel_dispatch.private_segment_size, - .group_segment_size = kernel_packet.kernel_dispatch.group_segment_size, - .workgroup_size = - rocprofiler_dim3_t{kernel_packet.kernel_dispatch.workgroup_size_x, - kernel_packet.kernel_dispatch.workgroup_size_y, - kernel_packet.kernel_dispatch.workgroup_size_z}, - .grid_size = rocprofiler_dim3_t{kernel_packet.kernel_dispatch.grid_size_x, - kernel_packet.kernel_dispatch.grid_size_y, - kernel_packet.kernel_dispatch.grid_size_z}, - .reserved_padding = {0}}}; + auto dispatch_id = ++sequence_counter; + _packet_data.callback_record = + callback_record_t{sizeof(callback_record_t), + rocprofiler_timestamp_t{0}, + rocprofiler_timestamp_t{0}, + rocprofiler_kernel_dispatch_info_t{ + .size = kernel_dispatch_info_rt_size, + .agent_id = queue.get_agent().get_rocp_agent()->id, + .queue_id = queue.get_id(), + .kernel_id = kernel_id, + .dispatch_id = dispatch_id, + .private_segment_size = pkt_info.private_segment_size, + .group_segment_size = pkt_info.group_segment_size, + .workgroup_size = pkt_info.workgroup_size, + .grid_size = pkt_info.grid_size, + .reserved_padding = {0}}}; { auto tracer_data = _packet_data.callback_record; @@ -623,7 +699,7 @@ WriteInterceptor(const void* packets, get_core_table()->hsa_signal_store_screlease_fn(completion_signal, 0); } - ROCP_FATAL_IF(packet_type != HSA_PACKET_TYPE_KERNEL_DISPATCH) + ROCP_FATAL_IF(!(is_kernel_dispatch || is_ext_kernel_dispatch)) << "get_kernel_id below might need to be updated"; { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp index fd42f9bcaa74..d3e4701552aa 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp @@ -43,9 +43,12 @@ union rocprofiler_packet { hsa_ext_amd_aql_pm4_packet_t ext_amd_aql_pm4; hsa_kernel_dispatch_packet_t kernel_dispatch; - hsa_barrier_and_packet_t barrier_and; - hsa_barrier_or_packet_t barrier_or; - amd_aql_intercept_marker_t marker; +#if HSA_AMD_EXT_API_TABLE_STEP_VERSION >= 0x0D + hsa_amd_ext_kernel_dispatch_packet_t ext_kernel_dispatch; +#endif + hsa_barrier_and_packet_t barrier_and; + hsa_barrier_or_packet_t barrier_or; + amd_aql_intercept_marker_t marker; rocprofiler_packet() : ext_amd_aql_pm4{null_amd_aql_pm4_packet} @@ -71,6 +74,12 @@ union rocprofiler_packet : marker{val} {} +#if HSA_AMD_EXT_API_TABLE_STEP_VERSION >= 0x0D + rocprofiler_packet(hsa_amd_ext_kernel_dispatch_packet_t val) + : ext_kernel_dispatch{val} + {} +#endif + ~rocprofiler_packet() = default; rocprofiler_packet(const rocprofiler_packet&) = default; rocprofiler_packet(rocprofiler_packet&&) noexcept = default; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp index 77c0ab9351d6..0a7840d7d46b 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/hsa_adapter.cpp @@ -320,7 +320,7 @@ pc_sampling_service_finish_configuration(context::pc_sampling_service* service) pc_sampling::utils::get_matching_hsa_pcs_units(agent_session->unit), agent_session->interval, pc_sampling::utils::get_hsa_pcs_latency(), - pc_sampling::utils::get_hsa_pcs_buffer_size(), + pc_sampling::utils::get_hsa_pcs_buffer_size(agent_session->agent->gfx_target_version), data_ready_callback, agent_session.get(), &agent_session->hsa_pc_sampling); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.cpp index dc51bb4d90d6..770251df9d98 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/ioctl/ioctl_adapter.cpp @@ -318,6 +318,14 @@ is_pc_sampling_method_supported(rocprofiler_pc_sampling_method_t method, else return ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_KERNEL; } + else if(agent_name.find("gfx1250") == 0) + { + // 1.7 version enables stochastic PC sampling on gfx1250 + if(pcs_ioctl_version >= PC_SAMPLING_IOCTL_COMPUTE_VERSION(1, 7)) + return ROCPROFILER_STATUS_SUCCESS; + else + return ROCPROFILER_STATUS_ERROR_INCOMPATIBLE_KERNEL; + } } else { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/CMakeLists.txt index 431069168b75..b7f16a313ee6 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/CMakeLists.txt @@ -8,7 +8,8 @@ set(ROCPROFILER_LIB_PC_SAMPLING_PARSER_HEADERS rocr.h translation.hpp gfx950.hpp - gfx12.hpp) + gfx12.hpp + gfx1250.hpp) target_sources( rocprofiler-sdk-object-library PRIVATE ${ROCPROFILER_LIB_PC_SAMPLING_PARSER_SOURCES} diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/correlation.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/correlation.hpp index c8073c69989c..fab1c2696ea7 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/correlation.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/correlation.hpp @@ -388,7 +388,14 @@ pcsample_status_t inline parse_buffer(generic_sample_t* buffer, } else if(gfxip_major == 12) { - parseSample_func = _parse_buffer; + if(gfxip_minor == 5) + { + parseSample_func = _parse_buffer; + } + else + { + parseSample_func = _parse_buffer; + } } else { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/gfx1250.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/gfx1250.hpp new file mode 100644 index 000000000000..350e5b4c929a --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/gfx1250.hpp @@ -0,0 +1,30 @@ +// MIT License +// +// Copyright (c) 2023-2026 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +// FIXME: check what can be inherited +class GFX1250 : public GFX12 +{ +public: + static constexpr uint32_t gfx_ip_minor = 5; +}; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/pc_record_interface.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/pc_record_interface.cpp index 8d418df50e72..753369c4b536 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/pc_record_interface.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/pc_record_interface.cpp @@ -109,7 +109,14 @@ PCSamplingParserContext::parse(const upcoming_samples_t& upcoming, } else if(gfxip_major == 12) { - parseSample_func = _get_parse_func_for_method(pcs_method); + if(gfxip_minor == 5) + { + parseSample_func = _get_parse_func_for_method(pcs_method); + } + else + { + parseSample_func = _get_parse_func_for_method(pcs_method); + } } else { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/CMakeLists.txt index 46feb0cf517e..58b0f492c9ca 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/CMakeLists.txt @@ -21,6 +21,9 @@ set(ROCPROFILER_LIB_PC_SAMPLING_PARSER_MULTIGPU_TEST_SOURCES set(ROCPROFILER_LIB_PC_SAMPLING_PARSER_GFX12_TEST_SOURCES ${ROCPROFILER_LIB_PC_SAMPLING_PARSER_TEST_SOURCES} gfxtest.hpp gfx12test.cpp) +set(ROCPROFILER_LIB_PC_SAMPLING_PARSER_GFX1250_TEST_SOURCES + ${ROCPROFILER_LIB_PC_SAMPLING_PARSER_TEST_SOURCES} gfxtest.hpp gfx1250test.cpp) + add_executable(pcs_gfx9_test) target_sources(pcs_gfx9_test @@ -117,3 +120,25 @@ rocprofiler_add_unit_test( ENVIRONMENT "LD_LIBRARY_PATH=${CMAKE_INSTALL_PREFIX}/lib:${CMAKE_INSTALL_PREFIX}/llvm/lib:/opt/rocm/lib:/opt/rocm/llvm/lib:$ENV{LD_LIBRARY_PATH}" FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") + +add_executable(pcs_gfx1250_test) + +target_sources(pcs_gfx1250_test + PRIVATE ${ROCPROFILER_LIB_PC_SAMPLING_PARSER_GFX1250_TEST_SOURCES}) +target_include_directories(pcs_gfx1250_test PRIVATE ${PCTEST_INCLUDE_DIR}) + +target_link_libraries( + pcs_gfx1250_test + PRIVATE rocprofiler-sdk::rocprofiler-sdk-common-library + rocprofiler-sdk::rocprofiler-sdk-static-library GTest::gtest + GTest::gtest_main) + +rocprofiler_add_unit_test( + TARGET pcs_gfx1250_test + SOURCES ${ROCPROFILER_LIB_PC_SAMPLING_PARSER_GFX1250_TEST_SOURCES} + LABELS "unittests" + TIMEOUT 45 + DISABLED TRUE + ENVIRONMENT + "LD_LIBRARY_PATH=${CMAKE_INSTALL_PREFIX}/lib:${CMAKE_INSTALL_PREFIX}/llvm/lib:/opt/rocm/lib:/opt/rocm/llvm/lib:$ENV{LD_LIBRARY_PATH}" + FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/gfx1250test.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/gfx1250test.cpp new file mode 100644 index 000000000000..4d6d19ee665b --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/gfx1250test.cpp @@ -0,0 +1,593 @@ +// MIT License +// +// Copyright (c) 2023-2026 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#ifdef NDEBUG +# undef NDEBUG +#endif + +#include "lib/rocprofiler-sdk/pc_sampling/parser/pc_record_interface.hpp" +#include "lib/rocprofiler-sdk/pc_sampling/parser/tests/gfxtest.hpp" +#include "lib/rocprofiler-sdk/pc_sampling/parser/tests/mocks.hpp" + +#include + +#include +#include + +#define GFXIP_MAJOR 12 + +#define RECORD_INST_TYPE(x) \ + { \ + PcSamplingRecordT sample{}; \ + sample.inst_type = ROCPROFILER_PC_SAMPLING_INSTRUCTION##_##x; \ + this->snapshots.push_back(sample); \ + } + +#define GENERATE_RECORDS_INST_TYPE() \ + RECORD_INST_TYPE(TYPE_VALU); \ + RECORD_INST_TYPE(TYPE_SCALAR); \ + RECORD_INST_TYPE(TYPE_TEX); \ + RECORD_INST_TYPE(TYPE_LDS); \ + RECORD_INST_TYPE(TYPE_LDS_DIRECT); \ + RECORD_INST_TYPE(TYPE_EXPORT); \ + RECORD_INST_TYPE(TYPE_MESSAGE); \ + RECORD_INST_TYPE(TYPE_BARRIER); \ + RECORD_INST_TYPE(TYPE_BRANCH_NOT_TAKEN); \ + RECORD_INST_TYPE(TYPE_BRANCH_TAKEN); \ + RECORD_INST_TYPE(TYPE_JUMP); \ + RECORD_INST_TYPE(TYPE_OTHER); \ + RECORD_INST_TYPE(TYPE_NO_INST); \ + RECORD_INST_TYPE(TYPE_DUAL_VALU); \ + RECORD_INST_TYPE(TYPE_FLAT); \ + RECORD_INST_TYPE(TYPE_MATRIX); + +#define RECORD_NOT_ISSUED_REASON(x) \ + { \ + PcSamplingRecordT sample{}; \ + sample.snapshot.reason_not_issued = ROCPROFILER_PC_SAMPLING_INSTRUCTION_NOT_ISSUED##_##x; \ + this->snapshots.push_back(sample); \ + } + +#define GENERATE_RECORDS_NOT_ISSUED_REASON(x) \ + RECORD_NOT_ISSUED_REASON(REASON_NO_INSTRUCTION_AVAILABLE); \ + RECORD_NOT_ISSUED_REASON(REASON_ALU_DEPENDENCY); \ + RECORD_NOT_ISSUED_REASON(REASON_WAITCNT); \ + RECORD_NOT_ISSUED_REASON(REASON_ARBITER_NOT_WIN); \ + RECORD_NOT_ISSUED_REASON(REASON_SLEEP_WAIT); \ + RECORD_NOT_ISSUED_REASON(REASON_BARRIER_WAIT); \ + RECORD_NOT_ISSUED_REASON(REASON_OTHER_WAIT); \ + RECORD_NOT_ISSUED_REASON(REASON_INTERNAL_INSTRUCTION); + +#define RECORD_ARBSTATE_ISSUE_STALL(x, y) \ + { \ + PcSamplingRecordT sample{}; \ + sample.snapshot.arb_state##_##x = 1; \ + sample.snapshot.arb_state##_##y = 1; \ + this->snapshots.push_back(sample); \ + } + +// Respecting the order of elements in GFX12:arb_state that match the order of arb_state bits +// in perf_snapshot_data register +#define RECORD_ARBSTATE_ISSUE(x) \ + RECORD_ARBSTATE_ISSUE_STALL(x, stall_brmsg); \ + RECORD_ARBSTATE_ISSUE_STALL(x, stall_exp); \ + RECORD_ARBSTATE_ISSUE_STALL(x, stall_lds_direct); \ + RECORD_ARBSTATE_ISSUE_STALL(x, stall_lds); \ + RECORD_ARBSTATE_ISSUE_STALL(x, stall_vmem_tex); \ + RECORD_ARBSTATE_ISSUE_STALL(x, stall_scalar); \ + RECORD_ARBSTATE_ISSUE_STALL(x, stall_valu); + +// Respecting the order of elements in GFX12:arb_state that match the order of arb_state bits +// in perf_snapshot_data register +#define GENERATE_RECORDS_ARBSTATE_ISSUE() \ + RECORD_ARBSTATE_ISSUE(issue_brmsg); \ + RECORD_ARBSTATE_ISSUE(issue_exp); \ + RECORD_ARBSTATE_ISSUE(issue_lds_direct); \ + RECORD_ARBSTATE_ISSUE(issue_lds); \ + RECORD_ARBSTATE_ISSUE(issue_vmem_tex); \ + RECORD_ARBSTATE_ISSUE(issue_scalar); \ + RECORD_ARBSTATE_ISSUE(issue_valu); + +#define NON_GFX12_ARBSTATE_IS_ZERO(x, y) \ + EXPECT_EQ(x.snapshot.arb_state_issue_misc, 0); \ + EXPECT_EQ(y.snapshot.arb_state_issue_misc, 0); \ + EXPECT_EQ(x.snapshot.arb_state_issue_matrix, 0); \ + EXPECT_EQ(y.snapshot.arb_state_issue_matrix, 0); \ + EXPECT_EQ(x.snapshot.arb_state_issue_flat, 0); \ + EXPECT_EQ(y.snapshot.arb_state_issue_flat, 0); \ + \ + EXPECT_EQ(x.snapshot.arb_state_stall_misc, 0); \ + EXPECT_EQ(y.snapshot.arb_state_stall_misc, 0); \ + EXPECT_EQ(x.snapshot.arb_state_stall_matrix, 0); \ + EXPECT_EQ(y.snapshot.arb_state_stall_matrix, 0); \ + EXPECT_EQ(x.snapshot.arb_state_stall_flat, 0); \ + EXPECT_EQ(y.snapshot.arb_state_stall_flat, 0); + +#define MATCH_ARBSTATE(x, y) \ + EXPECT_EQ(x.snapshot.arb_state_issue_valu, y.snapshot.arb_state_issue_valu); \ + EXPECT_EQ(x.snapshot.arb_state_issue_lds, y.snapshot.arb_state_issue_lds); \ + EXPECT_EQ(x.snapshot.arb_state_issue_lds_direct, y.snapshot.arb_state_issue_lds_direct); \ + EXPECT_EQ(x.snapshot.arb_state_issue_scalar, y.snapshot.arb_state_issue_scalar); \ + EXPECT_EQ(x.snapshot.arb_state_issue_vmem_tex, y.snapshot.arb_state_issue_vmem_tex); \ + EXPECT_EQ(x.snapshot.arb_state_issue_exp, y.snapshot.arb_state_issue_exp); \ + EXPECT_EQ(x.snapshot.arb_state_issue_brmsg, y.snapshot.arb_state_issue_brmsg); \ + \ + EXPECT_EQ(x.snapshot.arb_state_stall_valu, y.snapshot.arb_state_stall_valu); \ + EXPECT_EQ(x.snapshot.arb_state_stall_lds, y.snapshot.arb_state_stall_lds); \ + EXPECT_EQ(x.snapshot.arb_state_stall_lds_direct, y.snapshot.arb_state_stall_lds_direct); \ + EXPECT_EQ(x.snapshot.arb_state_stall_scalar, y.snapshot.arb_state_stall_scalar); \ + EXPECT_EQ(x.snapshot.arb_state_stall_vmem_tex, y.snapshot.arb_state_stall_vmem_tex); \ + EXPECT_EQ(x.snapshot.arb_state_stall_exp, y.snapshot.arb_state_stall_exp); \ + EXPECT_EQ(x.snapshot.arb_state_stall_brmsg, y.snapshot.arb_state_stall_brmsg); \ + \ + NON_GFX12_ARBSTATE_IS_ZERO(x, y) + +// Design decision +// Due to some potential future changes in the subfamily, +// I've decided to keep this as a separate test and not spend much time +// in merging the the GFX12. In the future, when we'll have all details on deck, +// we might find the right way to inherit common behavior in GFX12, +// and separate only specifics to GFX1250. + +template +class InstTypeTestGFX1250 : public InstTypeTest +{ +public: + void generate_records_inst_type() override { GENERATE_RECORDS_INST_TYPE(); } +}; + +template +class StallReasonTestGFX1250 : public StallReasonTest +{ +public: + void generate_records_not_issued_reason() override { GENERATE_RECORDS_NOT_ISSUED_REASON(); } +}; + +template +class ArbStateTestGFX1250 : public ArbStateTest +{ +public: + void generate_records_arbstate_issue() override { GENERATE_RECORDS_ARBSTATE_ISSUE(); } + + void match_arbstate(PcSamplingRecordT& x, PcSamplingRecordT& y) override + { + MATCH_ARBSTATE(x, y); + } +}; + +// TODO: check if it's possible to reuse the GFX12 test +template +class WaveIssueAndErrorTestGFX1250 +: public WaveIssueAndErrorTest +{ + // Encodes bits from the perf_snapshot_data register + union perf_snapshot_data_t + { + struct + { + uint32_t valid : 1; + uint32_t issued : 1; + uint32_t reserved : 30; + }; + uint32_t raw; + }; + + // specific + void FillBuffers() override + { + this->buffer->genUpcomingSamples(4); + for(int valid = 0; valid <= 1; valid++) + for(int issued = 0; issued <= 1; issued++) + genPCSample(valid, issued); + } + + // Could be reused with assumption that the num_combinations will be overridden + void CheckBuffers() override + { + // TODO: fix this + const int num_combinations = 4; + auto parsed = this->buffer->get_parsed_buffer(GFX1250::gfx_ip_major, GFX1250::gfx_ip_minor); + EXPECT_EQ(parsed.size(), 1); + EXPECT_EQ(parsed[0].size(), num_combinations); + EXPECT_EQ(this->compare.size(), num_combinations); + + for(size_t i = 0; i < num_combinations; i++) + { + if(this->compare[i].valid) + { + EXPECT_EQ(this->compare[i].valid_record.wave_issued, parsed[0][i].wave_issued); + // dual_issue_valu doesn't exist on GFX12, so we expect it to be 0 always + EXPECT_EQ(0, parsed[0][i].snapshot.dual_issue_valu); + } + else + { + // Internally (inside the parser) invalid samples are represented with + // PcSamplingRecordT of size 0. Eventually, those records are replaced with the + // PcSamplingRecordInvalidT prior to putting inside the SDK buffer. + EXPECT_EQ(parsed[0][i].size, 0); + } + } + } + + void genPCSample(bool valid, bool issued) + { + typename WaveIssueAndErrorTest:: + pc_sampling_test_record_t record{}; + record.valid = valid; + if(record.valid) + { + // Fill in the data for the valid record. + auto& sample = record.valid_record; + + // TODO: Since code objects are not mocked, use pc.code_object_offset + // as the absolute physical address of the mocked PC. + sample.pc.code_object_offset = this->dispatch->unique_id; + + sample.correlation_id.internal = this->dispatch->getMockId().raw; + + sample.wave_issued = issued; + + EXPECT_NE(this->dispatch.get(), nullptr); + } + + this->compare.push_back(record); + + perf_snapshot_data_t perf_snapshot_data{}; + perf_snapshot_data.valid = valid; + perf_snapshot_data.issued = issued; + + perf_sample_snapshot_v1 pss; + pss.perf_snapshot_data = perf_snapshot_data.raw; + pss.correlation_id = this->dispatch->getMockId().raw; + this->dispatch->submit(std::move(pss)); + }; +}; + +// TODO: Think about how to merge the HWIDTest for GFX12 and GFX1250 +template +class HwIdTestGFX1250 : public WaveSnapTest +{ + // The combined hw_id1 and hw_id2 encoded by ROCr's 2nd level trap handler + union gfx12_hw_id_t + { + uint32_t raw; + struct + { + uint32_t wave_id : 5; ///< wave_id[4:0] + uint32_t queue_id : 4; ///< queue_id[8:5] + uint32_t reserved0 : 1; ///< reserved [9] + uint32_t cu_or_wgp_id : 4; ///< wgp_id[13:10] + uint32_t simd_id : 2; ///< simd_id[15:14] + uint32_t shader_array_id : 1; ///< sa_id[16] + uint32_t microengine_id : 1; ///< me_id[17] + uint32_t shader_engine_id : 2; ///< se_id[19:18] + uint32_t pipe_id : 2; ///< pipe_id[21:20] + uint32_t reserved1 : 1; ///< reserved [22] + uint32_t workgroup_id : 5; ///< wg_id[27:23] + uint32_t vm_id : 4; ///< vm_id[31:28] + }; + }; + + void FillBuffers() override + { + gfx12_hw_id_t hw_id_val0{}; + hw_id_val0.wave_id = 0; + hw_id_val0.simd_id = 0; + hw_id_val0.cu_or_wgp_id = 0; + hw_id_val0.shader_array_id = 0; + hw_id_val0.shader_engine_id = 0; + hw_id_val0.queue_id = 0; + hw_id_val0.pipe_id = 0; + hw_id_val0.microengine_id = 0; + hw_id_val0.workgroup_id = 0; + hw_id_val0.vm_id = 0; + + gfx12_hw_id_t hw_id_val1{}; + hw_id_val1.wave_id = 15; + hw_id_val1.simd_id = 3; + hw_id_val1.cu_or_wgp_id = 15; + hw_id_val1.shader_array_id = 1; + hw_id_val1.shader_engine_id = 2; + hw_id_val1.queue_id = 7; + hw_id_val1.pipe_id = 3; + hw_id_val1.microengine_id = 1; + hw_id_val1.workgroup_id = 15; + hw_id_val1.vm_id = 15; + + gfx12_hw_id_t hw_id_val2{}; + hw_id_val2.wave_id = 7; + hw_id_val2.simd_id = 2; + hw_id_val2.cu_or_wgp_id = 6; + hw_id_val2.shader_array_id = 0; + hw_id_val2.shader_engine_id = 3; + hw_id_val2.queue_id = 3; + hw_id_val2.pipe_id = 2; + hw_id_val2.microengine_id = 1; + hw_id_val2.workgroup_id = 8; + hw_id_val2.vm_id = 9; + + this->buffer->genUpcomingSamples(3); + genPCSample(hw_id_val0); + genPCSample(hw_id_val1); + genPCSample(hw_id_val2); + } + + void CheckBuffers() override + { + auto parsed = this->buffer->get_parsed_buffer(GFXIP_MAJOR); // GFXIP==12 + EXPECT_EQ(parsed.size(), 1); + EXPECT_EQ(parsed[0].size(), 3); + EXPECT_EQ(compare.size(), 3); + + for(size_t i = 0; i < 3; i++) + { + // Comparing individual fields + EXPECT_EQ(compare[i].hw_id.wave_id, parsed[0][i].hw_id.wave_id); + EXPECT_EQ(compare[i].hw_id.simd_id, parsed[0][i].hw_id.simd_id); + EXPECT_EQ(compare[i].hw_id.pipe_id, parsed[0][i].hw_id.pipe_id); + EXPECT_EQ(compare[i].hw_id.cu_or_wgp_id, parsed[0][i].hw_id.cu_or_wgp_id); + EXPECT_EQ(compare[i].hw_id.shader_array_id, parsed[0][i].hw_id.shader_array_id); + EXPECT_EQ(compare[i].hw_id.shader_engine_id, parsed[0][i].hw_id.shader_engine_id); + EXPECT_EQ(compare[i].hw_id.workgroup_id, parsed[0][i].hw_id.workgroup_id); + EXPECT_EQ(compare[i].hw_id.vm_id, parsed[0][i].hw_id.vm_id); + EXPECT_EQ(compare[i].hw_id.queue_id, parsed[0][i].hw_id.queue_id); + EXPECT_EQ(compare[i].hw_id.microengine_id, parsed[0][i].hw_id.microengine_id); + } + } + + void genPCSample(gfx12_hw_id_t hw_id) + { + // ROCr doesn't deliver the info store in hw_id2 + PcSamplingRecordT sample; + ::memset(&sample, 0, sizeof(sample)); + + // Unpacking individual fields + // NOTE: chiplet is tested in a WaveOtherFieldsTest test, because it's not + // transferred via hw_id, but chiplet_and_wave_id field. + sample.hw_id.wave_id = hw_id.wave_id; + sample.hw_id.simd_id = hw_id.simd_id; + sample.hw_id.cu_or_wgp_id = hw_id.cu_or_wgp_id; + sample.hw_id.shader_array_id = hw_id.shader_array_id; + sample.hw_id.shader_engine_id = hw_id.shader_engine_id; + sample.hw_id.pipe_id = hw_id.pipe_id; + sample.hw_id.workgroup_id = hw_id.workgroup_id; + sample.hw_id.vm_id = hw_id.vm_id; + sample.hw_id.queue_id = hw_id.queue_id; + sample.hw_id.microengine_id = hw_id.microengine_id; + + compare.push_back(sample); + + perf_sample_snapshot_v1 snap; + ::memset(&snap, 0, sizeof(snap)); + + // raw register value + snap.hw_id = hw_id.raw; + snap.correlation_id = this->dispatch->getMockId().raw; + snap.perf_snapshot_data |= 0x1; // sample is valid + + EXPECT_NE(this->dispatch.get(), nullptr); + this->dispatch->submit(snap); + }; + + std::vector compare; +}; + +// TODO: this needs some replumbing based on the new counter definition format +template +class MemoryCountersTest : public WaveSnapTest +{ + union perf_snapshot_data + { + uint32_t raw; + struct + { + uint32_t reserved : 20; ///< bits 19:0 + uint32_t async_cnt : 6; ///< bits 25:20 + uint32_t tensor_cnt : 6; ///< bits 31:26 + }; + }; + + union perf_snapshot_data1 + { + uint32_t raw; + struct + { + uint32_t reserved : 26; ///< bits 25:0 + uint32_t xnack_cnt : 6; ///< bits 31:26 + }; + }; + + union perf_snapshot_data2 + { + uint32_t raw; + struct + { + uint32_t load_cnt : 6; ///< bits 5:0 + uint32_t store_cnt : 6; ///< bits 11:6 + uint32_t bvh_cnt : 3; ///< bits 14:12 + uint32_t sample_cnt : 6; ///< bits 20:15 + uint32_t ds_cnt : 6; ///< bits 26:21 + uint32_t km_cnt : 5; ///< bits 31:27 + }; + }; + + void FillBuffers() override + { + this->buffer->genUpcomingSamples(4); + genPCSample(0, 0, 0, 0, 0, 0, 0, 0, 0); // All zeros + genPCSample(1, 2, 3, 4, 5, 6, 7, 8, 9); // Counting + genPCSample(3, 5, 7, 11, 13, 17, 19, 23, 29); // Some prime numbers + genPCSample(29, 23, 19, 17, 13, 11, 7, 5, 7); // Some reversed primes + } + + void CheckBuffers() override + { + // Test appliclable only to stochastic sampling records + bool is_stoch_sampling_record = + std::is_same::value; + + EXPECT_EQ(is_stoch_sampling_record, true); + + auto parsed = this->buffer->get_parsed_buffer(GFX1250::gfx_ip_major, GFX1250::gfx_ip_minor); + EXPECT_EQ(parsed.size(), 1); + EXPECT_EQ(parsed[0].size(), 4); + EXPECT_EQ(compare.size(), 4); + + for(size_t i = 0; i < 4; i++) + { + EXPECT_EQ(1, parsed[0][i].flags.has_memory_counter); + EXPECT_EQ(compare[i].memory_counters.load_cnt, parsed[0][i].memory_counters.load_cnt); + EXPECT_EQ(compare[i].memory_counters.store_cnt, parsed[0][i].memory_counters.store_cnt); + EXPECT_EQ(compare[i].memory_counters.bvh_cnt, parsed[0][i].memory_counters.bvh_cnt); + EXPECT_EQ(compare[i].memory_counters.sample_cnt, + parsed[0][i].memory_counters.sample_cnt); + EXPECT_EQ(compare[i].memory_counters.ds_cnt, parsed[0][i].memory_counters.ds_cnt); + EXPECT_EQ(compare[i].memory_counters.km_cnt, parsed[0][i].memory_counters.km_cnt); + EXPECT_EQ(compare[i].memory_counters.async_cnt, parsed[0][i].memory_counters.async_cnt); + EXPECT_EQ(compare[i].memory_counters.tensor_cnt, + parsed[0][i].memory_counters.tensor_cnt); + EXPECT_EQ(compare[i].memory_counters.xnack_cnt, parsed[0][i].memory_counters.xnack_cnt); + } + } + + void genPCSample(int load_cnt, + int store_cnt, + int bvh_cnt, + int sample_cnt, + int ds_cnt, + int km_cnt, + int async_cnt, + int tensor_cnt, + int xnack_cnt) + { + PcSamplingRecordT sample; + ::memset(&sample, 0, sizeof(sample)); + + sample.flags.has_memory_counter = 1; + sample.memory_counters.load_cnt = load_cnt; + sample.memory_counters.store_cnt = store_cnt; + sample.memory_counters.bvh_cnt = bvh_cnt; + sample.memory_counters.sample_cnt = sample_cnt; + sample.memory_counters.ds_cnt = ds_cnt; + sample.memory_counters.km_cnt = km_cnt; + sample.memory_counters.async_cnt = async_cnt; + sample.memory_counters.tensor_cnt = tensor_cnt; + sample.memory_counters.xnack_cnt = xnack_cnt; + + compare.push_back(sample); + + perf_sample_snapshot_v1 snap; + ::memset(&snap, 0, sizeof(snap)); + + perf_snapshot_data data{}; + data.async_cnt = async_cnt; + data.tensor_cnt = tensor_cnt; + + perf_snapshot_data1 data1{}; + data1.xnack_cnt = xnack_cnt; + + perf_snapshot_data2 data2{}; + data2.load_cnt = load_cnt; + data2.store_cnt = store_cnt; + data2.bvh_cnt = bvh_cnt; + data2.sample_cnt = sample_cnt; + data2.ds_cnt = ds_cnt; + data2.km_cnt = km_cnt; + + snap.perf_snapshot_data = data.raw; + snap.perf_snapshot_data1 = data1.raw; + snap.perf_snapshot_data2 = data2.raw; + snap.correlation_id = this->dispatch->getMockId().raw; + + // to ensure all stochastic samples are generated properly, + // marked them as valid + snap.perf_snapshot_data |= 0x1; // set the bit indicating the sample is valid + + EXPECT_NE(this->dispatch.get(), nullptr); + this->dispatch->submit(snap); + }; + + std::vector compare; +}; + +// TODO: check if it's possible to reuse the GFX12 test +template +class SamplingLockErrorTestGFX1250 : public SamplingLockErrorTest +{ + // Encodes bits from the perf_snapshot_data register with the focus + // on relevant bits for this test + union perf_snapshot_data_t + { + struct + { + uint32_t valid : 1; // bit 0 indicates if the sample is valid + uint32_t reserved0 : 13; + uint32_t sampling_lock_error : 1; ///< bit 14 indicates sampling lock error + uint32_t reserved1 : 17; + }; + uint32_t raw; + }; + + void genPCSample(bool sampling_lock_error) + { + PcSamplingRecordT record{}; + record.snapshot.sampling_lock_error = sampling_lock_error; + this->snapshots.push_back(record); + + perf_snapshot_data_t perf_snapshot_data{}; + // sample is always valid + perf_snapshot_data.valid = 1; + perf_snapshot_data.sampling_lock_error = sampling_lock_error; + + perf_sample_snapshot_v1 pss{}; + pss.perf_snapshot_data = perf_snapshot_data.raw; + pss.correlation_id = this->dispatch->getMockId().raw; + EXPECT_NE(this->dispatch.get(), nullptr); + this->dispatch->submit(std::move(pss)); + }; +}; + +TEST(pcs_parser, gfx1250_test) +{ + // Tests specific to stochastic sampling only + WaveCntTest{}.Test(); + InstTypeTestGFX1250{}.Test(); + StallReasonTestGFX1250{}.Test(); + ArbStateTestGFX1250{}.Test(); + WaveIssueAndErrorTestGFX1250{} + .Test(); + + // Tests common for both host trap and stochastic sampling. + HwIdTestGFX1250{}.Test(); + HwIdTestGFX1250{}.Test(); + WaveOtherFieldsTest{}.Test(); + WaveOtherFieldsTest{}.Test(); + MidMacroPCCorrection{}.Test(); + MidMacroPCCorrection{}.Test(); + + // test specific to GFX1250 + MemoryCountersTest{}.Test(); + SamplingLockErrorTestGFX1250{}.Test(); + + std::cout << "GFX1250 Test Done." << std::endl; +} diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/gfxtest.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/gfxtest.hpp index a19adbfa86a3..19e80d152286 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/gfxtest.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/tests/gfxtest.hpp @@ -73,6 +73,13 @@ class WaveSnapTest snap.perf_snapshot_data1 = wave_cnt; snap.perf_snapshot_data1 |= (arb_issue << 6) | (arb_stall << 14); } + else if constexpr(std::is_same_v) + { + snap.perf_snapshot_data = (inst_type << 2) | (reason << 6); + snap.perf_snapshot_data |= 0x1; // sample is valid + snap.perf_snapshot_data1 = wave_cnt; + snap.perf_snapshot_data1 |= (arb_issue << 9) | (arb_stall << 17); + } EXPECT_NE(dispatch.get(), nullptr); dispatch->submit(packet_union_t{.snap = snap}); @@ -380,3 +387,32 @@ class MidMacroPCCorrection : public WaveSnapTest ///< testing data std::vector compare; }; + +template +class SamplingLockErrorTest : public WaveSnapTest +{ + void FillBuffers() override + { + this->buffer->genUpcomingSamples(2); + // sampling_lock_error = 0; + this->genPCSample(0); + // sapmling lock error = 1 + this->genPCSample(1); + } + + // Could be reused with assumption that the num_combinations will be overridden + void CheckBuffers() override + { + auto parsed = this->buffer->get_parsed_buffer(GFX1250::gfx_ip_major, GFX1250::gfx_ip_minor); + EXPECT_EQ(parsed.size(), 1); + EXPECT_EQ(parsed[0].size(), 2); + EXPECT_EQ(this->snapshots.size(), 2); + + EXPECT_EQ(this->snapshots[0].snapshot.sampling_lock_error, + parsed[0][0].snapshot.sampling_lock_error); + EXPECT_EQ(this->snapshots[1].snapshot.sampling_lock_error, + parsed[0][1].snapshot.sampling_lock_error); + } + + virtual void genPCSample(bool sampling_lock_error) = 0; +}; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/translation.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/translation.hpp index 79d110dbf9fd..4183a3b9701c 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/translation.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/parser/translation.hpp @@ -24,6 +24,7 @@ #include "lib/rocprofiler-sdk/pc_sampling/parser/gfx11.hpp" #include "lib/rocprofiler-sdk/pc_sampling/parser/gfx12.hpp" +#include "lib/rocprofiler-sdk/pc_sampling/parser/gfx1250.hpp" #include "lib/rocprofiler-sdk/pc_sampling/parser/gfx9.hpp" #include "lib/rocprofiler-sdk/pc_sampling/parser/gfx950.hpp" #include "lib/rocprofiler-sdk/pc_sampling/parser/parser_types.hpp" @@ -342,7 +343,7 @@ copySample(const void* sa // Extracting data from the perf_snapshot_data register auto perf_snapshot_data = sample_.perf_snapshot_data; - // The sample is valid if perf_snapshot_data.valid == 1 + // The sample is valid if perf_snapshot_data.valid == 1 auto valid = static_cast(EXTRACT_BITS(perf_snapshot_data, 0, 0)); if(!valid) { @@ -403,6 +404,123 @@ copySample(const void* sa return ret; } +/** + * @brief Host trap V0 sample for GFX1250 + */ +template <> +inline rocprofiler_pc_sampling_record_host_trap_v0_t +copySample(const void* sample) +{ + // Host-trap samples are same for GFX12 and GFX1250 + auto ret = copySample(sample); + // Only difference are chiplets that exist on GFX1250 + copyChipletId(ret, *static_cast(sample)); + return ret; +} + +/** + * @brief Stochastic sample for GFX1250 + */ +template <> +inline rocprofiler_pc_sampling_record_stochastic_v0_t +copySample(const void* sample) +{ + // Differences compared to the GFX12: + // - SAMPLING_LOCK_ERR introduced in the GFX1250 and means: + // A wave tried taking a snapshot but was unable to as the previous snapshot had not yet been + // read. + // - ARB_STATE_* begins at different offset + // - XCNT, ASYNC_CNT, TENSOR_CNT introduced in GFX1250. + // - GFX1250 uses chiplets + // Due to all differences and potential adaptation to the future subfamily members, + // we decided to have a separate implementation of copySample for GFX1250 stochastic. + + const auto& sample_ = *static_cast(sample); + + // Extracting data from the perf_snapshot_data register + auto perf_snapshot_data = sample_.perf_snapshot_data; + // The sample is valid if perf_snapshot_data.valid == 1 + auto valid = static_cast(EXTRACT_BITS(perf_snapshot_data, 0, 0)); + if(!valid) + { + // To reduce refactoring of the PC sampling parser, we agreed to internally represent + // invalid samples with `rocprofiler_pc_sampling_record_stochastic_v0_t` with size 0. + // Eventually, those records are replaced with rocprofiler_pc_sampling_record_invalid_t + // and placed into the SDK buffer consumed by the end tool. + rocprofiler_pc_sampling_record_stochastic_v0_t invalid{}; + invalid.size = 0; + // No need to further process invalid samples + return invalid; + } + + auto ret = copySampleHeader(sample_); + // GFX1250 has chiplets + copyChipletId(ret, sample_); + // ROCr uses same hw_id struct for both GFX12 and GFX1250. + copyHwId(ret.hw_id, sample_.hw_id); + + // wave issued an instruction + ret.wave_issued = EXTRACT_BITS(perf_snapshot_data, 1, 1); + // TODO: we need to handle new instruction type I think + // type of issued instruction, valid only if `ret.wave_issued` is true. + ret.inst_type = translate_inst(EXTRACT_BITS(perf_snapshot_data, 5, 2)); + // reason for not issuing an instruction, valid only if `ret.wave_issued` is false + ret.snapshot.reason_not_issued = + translate_reason(EXTRACT_BITS(perf_snapshot_data, 8, 6)); + + // Sampling lock error indicates that reading this sample took to long + // so at least one wave was locked out from taking a sample. + ret.snapshot.sampling_lock_error = EXTRACT_BITS(perf_snapshot_data, 14, 14); + + // arbiter state information + auto perf_snapshot_data1 = sample_.perf_snapshot_data1; + uint16_t arb_state = EXTRACT_BITS(perf_snapshot_data1, 24, 9); + + ret.snapshot.arb_state_issue_brmsg = EXTRACT_BITS(arb_state, 0, 0); + ret.snapshot.arb_state_issue_exp = EXTRACT_BITS(arb_state, 1, 1); + ret.snapshot.arb_state_issue_lds_direct = EXTRACT_BITS(arb_state, 2, 2); + ret.snapshot.arb_state_issue_lds = EXTRACT_BITS(arb_state, 3, 3); + ret.snapshot.arb_state_issue_vmem_tex = EXTRACT_BITS(arb_state, 4, 4); + ret.snapshot.arb_state_issue_scalar = EXTRACT_BITS(arb_state, 5, 5); + ret.snapshot.arb_state_issue_valu = EXTRACT_BITS(arb_state, 6, 6); + + ret.snapshot.arb_state_stall_brmsg = EXTRACT_BITS(arb_state, 8, 8); + ret.snapshot.arb_state_stall_exp = EXTRACT_BITS(arb_state, 9, 9); + ret.snapshot.arb_state_stall_lds_direct = EXTRACT_BITS(arb_state, 10, 10); + ret.snapshot.arb_state_stall_lds = EXTRACT_BITS(arb_state, 11, 11); + ret.snapshot.arb_state_stall_vmem_tex = EXTRACT_BITS(arb_state, 12, 12); + ret.snapshot.arb_state_stall_scalar = EXTRACT_BITS(arb_state, 13, 13); + ret.snapshot.arb_state_stall_valu = EXTRACT_BITS(arb_state, 14, 14); + + ret.wave_count = EXTRACT_BITS(perf_snapshot_data1, 5, 0); + + // Memory counters exist on GFX1250. + ret.flags.has_memory_counter = true; + // Extracting memory counters from the perf_snapshot_data2 register + auto perf_snapshot_data2 = sample_.perf_snapshot_data2; + ret.memory_counters.load_cnt = EXTRACT_BITS(perf_snapshot_data2, 5, 0); + ret.memory_counters.store_cnt = EXTRACT_BITS(perf_snapshot_data2, 11, 6); + ret.memory_counters.bvh_cnt = EXTRACT_BITS(perf_snapshot_data2, 14, 12); + ret.memory_counters.sample_cnt = EXTRACT_BITS(perf_snapshot_data2, 20, 15); + ret.memory_counters.ds_cnt = EXTRACT_BITS(perf_snapshot_data2, 26, 21); + ret.memory_counters.km_cnt = EXTRACT_BITS(perf_snapshot_data2, 31, 27); + // counters available in perf_snapshot_data + ret.memory_counters.async_cnt = EXTRACT_BITS(perf_snapshot_data, 25, 20); + ret.memory_counters.tensor_cnt = EXTRACT_BITS(perf_snapshot_data, 31, 26); + // counters available in perf_snapshot_data1 + ret.memory_counters.xnack_cnt = EXTRACT_BITS(perf_snapshot_data1, 31, 26); + + // ensure that the wave_id of snapshot_data matches the hw_id.wave_id + auto sampled_wave_id = EXTRACT_BITS(perf_snapshot_data, 13, 9); + if(sampled_wave_id != ret.hw_id.wave_id) + { + ROCP_FATAL << "sampled_wave_id: " << sampled_wave_id + << " mismatches the hw_id.wave_id: " << ret.hw_id.wave_id; + } + + return ret; +} + /** * @brief The default implementation assumes no correction is needed. */ diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/utils.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/utils.cpp index 407729d7eb3c..5978e24ae2d7 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/utils.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/utils.cpp @@ -69,6 +69,12 @@ get_matching_hsa_pcs_units(rocprofiler_pc_sampling_unit_t unit) ROCP_FATAL << "Illegal pc sampling unit " << unit; } + +size_t get_hsa_pcs_buffer_size(uint32_t /*gfx_target_version*/) +{ + // Use bigger buffer to reduce number of dropped samples on actual silicon + return 64 * 1024 * sizeof(perf_sample_hosttrap_v1_t); // 4MB +} } // namespace utils } // namespace pc_sampling } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/utils.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/utils.hpp index 8b1ff929345d..7e411b65ba26 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/utils.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/pc_sampling/utils.hpp @@ -54,12 +54,8 @@ get_hsa_pcs_latency() return 1000; } -constexpr size_t -get_hsa_pcs_buffer_size() -{ - // TODO: Find the minimum size of all buffers and use that. - return 64 * 1024 * sizeof(perf_sample_hosttrap_v1_t); // 4MB -} +size_t +get_hsa_pcs_buffer_size(uint32_t gfx_target_version); } // namespace utils } // namespace pc_sampling } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk/config.yaml b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk/config.yaml index 8ac734cfcf63..7b3864e80bcc 100644 --- a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk/config.yaml +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk/config.yaml @@ -37,6 +37,62 @@ rocprofiler-sdk: - architectures: - gfx90a expression: reduce(SQ_THREAD_CYCLES_VALU,sum)/reduce(SQ_ACTIVE_INST_VALU,sum) + - name: CHA_BUSY + description: Count busy cyles. Not windowable. + properties: [] + definitions: + - architectures: + - gfx1250 + block: CHA + event: 0 + - name: CHA_CYCLE + description: Number of cycles. Not windowable. + properties: [] + definitions: + - architectures: + - gfx1250 + block: CHA + event: 15 + - name: CHC_REQ_READ + description: Sum of REQ_READ_32B, REQ_READ_64B, REQ_READ_128B, REQ_READ_256B + properties: [] + definitions: + - architectures: + - gfx1250 + block: CHC + event: 11 + - name: CHC_REQ_READ_32B + description: Number of read requests with logical size 32B + properties: [] + definitions: + - architectures: + - gfx1250 + block: CHC + event: 12 + - name: CHC_REQ_READ_64B + description: Number of read requests with logical size 64B + properties: [] + definitions: + - architectures: + - gfx1250 + block: CHC + event: 13 + - name: CHC_REQ_READ_128B + description: Number of read requests with logical size 128B + properties: [] + definitions: + - architectures: + - gfx1250 + block: CHC + event: 14 + - name: CHC_REQ_READ_256B + description: Number of read requests with logical size 256B + properties: [] + definitions: + - architectures: + - gfx1250 + block: CHC + event: 15 - name: CPC_CPC_STAT_BUSY description: CPC Busy. properties: [] @@ -48,6 +104,7 @@ rocprofiler-sdk: - gfx941 - gfx942 - gfx950 + - gfx1250 block: CPC event: 25 - name: CPC_CPC_STAT_IDLE @@ -61,6 +118,7 @@ rocprofiler-sdk: - gfx941 - gfx942 - gfx950 + - gfx1250 block: CPC event: 26 - name: CPC_CPC_STAT_STALL @@ -74,6 +132,7 @@ rocprofiler-sdk: - gfx941 - gfx942 - gfx950 + - gfx1250 block: CPC event: 27 - name: CPC_CPC_TCIU_BUSY @@ -87,6 +146,7 @@ rocprofiler-sdk: - gfx941 - gfx942 - gfx950 + - gfx1250 block: CPC event: 28 - name: CPC_CPC_TCIU_IDLE @@ -100,6 +160,7 @@ rocprofiler-sdk: - gfx941 - gfx942 - gfx950 + - gfx1250 block: CPC event: 29 - name: CPC_CPC_UTCL2IU_BUSY @@ -113,6 +174,7 @@ rocprofiler-sdk: - gfx941 - gfx942 - gfx950 + - gfx1250 block: CPC event: 30 - name: CPC_CPC_UTCL2IU_IDLE @@ -126,6 +188,7 @@ rocprofiler-sdk: - gfx941 - gfx942 - gfx950 + - gfx1250 block: CPC event: 31 - name: CPC_CPC_UTCL2IU_STALL @@ -152,6 +215,7 @@ rocprofiler-sdk: - gfx941 - gfx942 - gfx950 + - gfx1250 block: CPC event: 13 - name: CPC_ME1_DC0_SPI_BUSY @@ -165,6 +229,7 @@ rocprofiler-sdk: - gfx941 - gfx942 - gfx950 + - gfx1250 block: CPC event: 33 - name: CPC_UTCL1_STALL_ON_TRANSLATION @@ -344,6 +409,10 @@ rocprofiler-sdk: - gfx950 block: CPF event: 23 + - architectures: + - gfx1250 + block: CPF + event: 24 - name: CPF_CPF_STAT_IDLE description: CPF Idle. properties: [] @@ -722,6 +791,56 @@ rocprofiler-sdk: - gfx1031 - gfx1032 expression: 100*reduce(GRBM_GDS_BUSY,max)/reduce(GRBM_GUI_ACTIVE,max) + - name: GLARBC_CYCLE + description: Number of cycles. Not windowable. + properties: [] + definitions: + - architectures: + - gfx1250 + block: GLARBC + event: 0 + - name: GLARBC_BUSY + description: Number of cycles with a request pending. Not Windowed. + properties: [] + definitions: + - architectures: + - gfx1250 + block: GLARBC + event: 1 + - name: GL1A_BUSY + description: Number of cycles. Not windowable. + properties: [] + definitions: + - architectures: + - gfx1250 + block: GL1A + event: 0 + - name: GL1A_CYCLE + description: Number of cycles. Not windowable. + properties: [] + definitions: + - architectures: + - gfx1250 + block: GL1A + event: 77 + - name: GL1C_GL2_REQ_READ_LEVEL + description: Number of cycles in-flight for reads, atomics with return, NOP_RTN0. Average GL2 + read latency = GL1C_PERF_SEL_GL2_REQ_READ_LEVEL/SUM_OF_REQ_COUNTERS + properties: [] + definitions: + - architectures: + - gfx1250 + block: GL1C + event: 10 + - name: GL1C_GL2_REQ_WRITE_LEVEL + description: Number of cycles in-flight for writes, atomics without return, NOP_ACK. Average + GL2 read latency = GL1C_PERF_SEL_GL2_REQ_WRITE_LEVEL/SUM_OF_REQ_COUNTERS + properties: [] + definitions: + - architectures: + - gfx1250 + block: GL1C + event: 19 - name: GL2C_EA_RDREQ description: Number of GL2C/EA read requests (either 32-byte or 64-byte or 128-byte) for all clients. @@ -1321,6 +1440,7 @@ rocprofiler-sdk: - gfx12 - gfx1200 - gfx1201 + - gfx1250 - gfx9 - gfx900 - gfx906 @@ -1343,6 +1463,7 @@ rocprofiler-sdk: - gfx941 - gfx942 - gfx950 + - gfx1250 block: GRBM event: 30 - name: GRBM_CPF_BUSY @@ -1439,6 +1560,7 @@ rocprofiler-sdk: - gfx12 - gfx1200 - gfx1201 + - gfx1250 - gfx9 - gfx900 - gfx906 @@ -1512,6 +1634,38 @@ rocprofiler-sdk: - gfx950 block: GRBM event: 34 + - name: GRBMA_COUNT + description: Tie High - Count Number of Clocks + properties: [] + definitions: + - architectures: + - gfx1250 + block: GRBMA + event: 0 + - name: GRBMA_SPI_BUSY + description: Any of the Shader Pipe Interpolators (SPI OR SQG) are busy. + properties: [] + definitions: + - architectures: + - gfx1250 + block: GRBMA + event: 9 + - name: GRBMH_COUNT + description: Tie High - Count Number of Clocks + properties: [] + definitions: + - architectures: + - gfx1250 + block: GRBMH + event: 0 + - name: GRBMH_SPI_BUSY + description: Any of the Shader Pipe Interpolators (SPI OR SQG) are busy. + properties: [] + definitions: + - architectures: + - gfx1250 + block: GRBMH + event: 9 - name: GpuUtil description: 'Unit: percent' properties: [] @@ -2132,6 +2286,10 @@ rocprofiler-sdk: - gfx950 block: SPI event: 79 + - architectures: + - gfx1250 + block: SPI + event: 144 - name: SPI_RA_REQ_NO_ALLOC_CSN description: Arb cycles with CSn req and no CSn alloc. Source is RA0 properties: [] @@ -2972,6 +3130,10 @@ rocprofiler-sdk: - gfx1201 block: SQ event: 302 + - architectures: + - gfx1250 + block: SQ + event: 642 - name: SQC_ICACHE_MISSES description: Number of cache misses, includes uncached requests. {per-Bank, nondeterministic, C2} @@ -2983,6 +3145,10 @@ rocprofiler-sdk: - gfx1201 block: SQ event: 303 + - architectures: + - gfx1250 + block: SQ + event: 643 - name: SQC_ICACHE_INPUT_VALID_READYB description: ' Input stalled by SQC (per-SQ, nondeterministic, unwindowed)' properties: [] @@ -3022,6 +3188,10 @@ rocprofiler-sdk: - gfx950 block: SQ event: 273 + - architectures: + - gfx1250 + block: SQ + event: 644 - name: SQC_ICACHE_REQ description: Number of requests. (per-SQ, per-Bank) properties: [] @@ -3041,6 +3211,10 @@ rocprofiler-sdk: - gfx1201 block: SQ event: 301 + - architectures: + - gfx1250 + block: SQ + event: 641 - name: SQC_LDS_BANK_CONFLICT description: Number of cycles LDS is stalled by bank conflicts. (emulated, C1) properties: [] @@ -3492,6 +3666,7 @@ rocprofiler-sdk: - gfx12 - gfx1200 - gfx1201 + - gfx1250 - gfx90a - gfx908 - gfx940 @@ -3511,6 +3686,7 @@ rocprofiler-sdk: - gfx941 - gfx942 - gfx950 + - gfx1250 block: SQ event: 2 - name: SQ_IFETCH @@ -3897,6 +4073,7 @@ rocprofiler-sdk: - gfx940 - gfx941 - gfx942 + - gfx1250 block: SQ event: 70 - architectures: @@ -4991,6 +5168,10 @@ rocprofiler-sdk: - gfx950 block: SQ event: 14 + - architectures: + - gfx1250 + block: SQ + event: 7 - name: SQ_LDS_ADDR_CONFLICT description: Number of cycles LDS (local data store) is stalled by address conflicts. This value is returned on a per-SE (aggregate of values in SIMDs in the SE) basis. @@ -5422,6 +5603,7 @@ rocprofiler-sdk: - gfx12 - gfx1200 - gfx1201 + - gfx1250 - gfx9 - gfx900 - gfx906 @@ -5667,6 +5849,10 @@ rocprofiler-sdk: - gfx950 block: SQ event: 95 + - architectures: + - gfx1250 + block: SQ + event: 32 - name: SQ_INSTS_VALU_FLOPS_FP16 description: Counts FLOPS per instruction on float 16 excluding MFMA/SMFMA. properties: [] @@ -5741,6 +5927,99 @@ rocprofiler-sdk: - gfx950 block: SQ event: 74 + - name: SQ_INST_CYCLES_VALU_WMMA + description: Number of cycles of work sent for wmma. {emulated} + properties: [] + definitions: + - architectures: + - gfx1250 + block: SQ + event: 359 + - name: SQ_INSTS_VEC32_VALU_WMMA + description: Number of valu WMMA instructions issued per vec32. {emulated} + properties: [] + definitions: + - architectures: + - gfx1250 + block: SQ + event: 237 + - name: SQ_VALU_WMMA_FLOP_FP32 + description: SGEMM flop, only report 1 per issue cycle, it stands for 128 Ops per SIMD32. + {emulated, windowed, deterministic} + properties: [] + definitions: + - architectures: + - gfx1250 + block: SQ + event: 288 + - name: SQ_VALU_WMMA_FLOP_FP16 + description: HPGEMM flop, only report 1 per issue cycle, it stands for 2048 Ops per SIMD32, + sparse will report 2 per issue cycle.. {emulated, windowed, deterministic} + properties: [] + definitions: + - architectures: + - gfx1250 + block: SQ + event: 290 + - name: SQ_VALU_WMMA_FLOP_FP8 + description: FP8 GEMM flop, only report 1 per issue cycle, it stands for 4096 Ops per SIMD32, + sparse will report 2 per issue cycle. Count as FP8 is the largest input data type (other + input is FP4, FP6 or FP8).. {emulated, windowed, deterministic} + properties: [] + definitions: + - architectures: + - gfx1250 + block: SQ + event: 292 + - name: SQ_VALU_WMMA_FLOP_FP6 + description: FP6 GEMM flop, only report 1 per issue cycle, it stands for 8192 Ops per SIMD32, + sparse will report 2 per issue cycle (but there is no sparse currently). Count as FP6 is the + largest input data type (other input is FP4 or FP6).. {emulated, windowed, deterministic} + properties: [] + definitions: + - architectures: + - gfx1250 + block: SQ + event: 293 + - name: SQ_VALU_WMMA_FLOP_FP4 + description: FP4 GEMM flop, only report 1 per issue cycle, it stands for 8192 Ops per SIMD32, + sparse will report 2 per issue cycle (but there is no sparse currently). Count as FP4 if + both sources are FP4.. {emulated, windowed, deterministic} + properties: [] + definitions: + - architectures: + - gfx1250 + block: SQ + event: 294 + - name: SQ_VALU_WMMA_FLOP_F6F4 + description: F6 or F4 GEMM flop, only report 1 per issue cycle, it stands for 4096 Ops per + SIMD32, sparse will report 2 per issue cycle. Mixed precision WMMA counts as FLOP_F6F4 if + input operands are only F6 and F4 (not F8). If any are F8, it counts as FLOP_F8. {emulated, + windowed, deterministic} + properties: [] + definitions: + - architectures: + - gfx1250 + block: SQ + event: 297 + - name: SQ_VALU_WMMA_FLOP_BF16 + description: BF16GEMM flop, only report 1 per issue cycle, it stands for 2048 Ops per SIMD32, + sparse will report 2 per issue cycle.. {emulated, windowed, deterministic} + properties: [] + definitions: + - architectures: + - gfx1250 + block: SQ + event: 291 + - name: SQ_VALU_WMMA_FLOP_I8 + description: I8GEMM flop, only report 1 per issue cycle, it stands for 4096 Ops per SIMD32, + sparse will report 2 per issue cycle.. {emulated, windowed, deterministic} + properties: [] + definitions: + - architectures: + - gfx1250 + block: SQ + event: 295 - name: SQ_INSTS_LDS_LOAD description: Number of LDS load instructions issued . (per-simd, emulated) properties: [] diff --git a/projects/rocprofiler-sdk/tests/bin/pc-sampling/exec-mask-manipulation/exec_mask_manipulation.cpp b/projects/rocprofiler-sdk/tests/bin/pc-sampling/exec-mask-manipulation/exec_mask_manipulation.cpp index 5b0effb6bb46..2f37a332b1f4 100644 --- a/projects/rocprofiler-sdk/tests/bin/pc-sampling/exec-mask-manipulation/exec_mask_manipulation.cpp +++ b/projects/rocprofiler-sdk/tests/bin/pc-sampling/exec-mask-manipulation/exec_mask_manipulation.cpp @@ -54,11 +54,11 @@ check_hip_error(void); // ====================================================== __global__ void -kernel1(const int c) +kernel1(const int c, const int iter_num) { int a = 0; #pragma nounroll - for(int i = 0; i < ITER_NUM; i++) + for(int i = 0; i < iter_num; i++) { asm volatile("v_mov_b32 %0 %1\n" : "=v"(a) : "s"(c)); asm volatile("v_mov_b32 %0 %1\n" : "=v"(a) : "s"(c)); @@ -164,11 +164,11 @@ kernel1(const int c) } __global__ void -kernel2(const int c) +kernel2(const int c, const int iter_num) { int a = 0; #pragma nounroll - for(int i = 0; i < ITER_NUM; i++) + for(int i = 0; i < iter_num; i++) { asm volatile("s_mov_b32 %0 %1\n" : "=s"(a) : "s"(c)); asm volatile("s_mov_b32 %0 %1\n" : "=s"(a) : "s"(c)); @@ -274,14 +274,14 @@ kernel2(const int c) } __global__ void -kernel3(const float c) +kernel3(const float c, const int iter_num) { double a = threadIdx.x; float i = 0; float d = threadIdx.x; float e = 0; int tid_even = threadIdx.x % 2; - for(int j = 0; j < ITER_NUM; j++) + for(int j = 0; j < iter_num; j++) { if(tid_even == 0) { @@ -501,21 +501,22 @@ run_kernel() HIP_API_CALL(hipDeviceGetAttribute(&wave_size, hipDeviceAttributeWarpSize, 0)); // Get device properties to retrieve GFXIP version - uint32_t num_blocks = BLOCK_SIZE; + size_t num_blocks = BLOCK_SIZE; + size_t num_iters = ITER_NUM; for(int i = 1; i <= wave_size; i++) { if(i % 2 == 1) - kernel1<<>>(i); + kernel1<<>>(i, num_iters); else - kernel2<<>>(i); + kernel2<<>>(i, num_iters); check_hip_error(); HIP_API_CALL(hipDeviceSynchronize()); } float arg = 0; - kernel3<<>>(arg); + kernel3<<>>(arg, num_iters); check_hip_error(); HIP_API_CALL(hipDeviceSynchronize()); } diff --git a/projects/rocprofiler-sdk/tests/pc_sampling/main.cpp b/projects/rocprofiler-sdk/tests/pc_sampling/main.cpp index b012d2b8eb48..2fde37cc9b5d 100644 --- a/projects/rocprofiler-sdk/tests/pc_sampling/main.cpp +++ b/projects/rocprofiler-sdk/tests/pc_sampling/main.cpp @@ -26,6 +26,7 @@ #include #include #include +#include namespace { @@ -61,19 +62,18 @@ check_hip_error(void); } // namespace __global__ void -matrix_multiply(float* A, float* B, float* Out, int /*m*/, int n, int k) +matrix_multiply(float* A, float* B, float* Out, int m, int n, int k) { int gid_x = blockDim.x * blockIdx.x + threadIdx.x; int gid_y = blockDim.y * blockIdx.y + threadIdx.y; - if(gid_x < N && gid_y < M) + if(gid_x < n && gid_y < m) { float sum = 0; for(int i = 0; i < k; ++i) { sum += A[gid_y * k + i] * B[i * n + gid_x]; } - Out[gid_y * n + gid_x] = sum; } } @@ -138,45 +138,37 @@ matrix_multiply_tile(float* A, float* B, float* Out, int m, int n, int k) #endif void -run_hip_app() +run_hip_app(const std::string& /*arch_name*/) { - std::vector A(M * K); - std::vector B(K * N); - std::vector Out(M * N); + size_t m = M, n = N, k = K; - // Randomly initialize the matrices - for(int i = 0; i < M * K; ++i) - { - A[i] = (float) rand() / (float) RAND_MAX; - } + std::vector A(m * k); + std::vector B(k * n); + std::vector Out(m * n); - for(int i = 0; i < K * N; ++i) - { - B[i] = (float) rand() / (float) RAND_MAX; - } + for(size_t i = 0; i < m * k; ++i) + A[i] = static_cast(rand()) / static_cast(RAND_MAX); + for(size_t i = 0; i < k * n; ++i) + B[i] = static_cast(rand()) / static_cast(RAND_MAX); - // Allocate GPU Memory float *d_A, *d_B, *d_Out; - HIP_API_CALL(hipMalloc(&d_A, sizeof(float) * M * K)); - HIP_API_CALL(hipMalloc(&d_B, sizeof(float) * K * N)); - HIP_API_CALL(hipMalloc(&d_Out, sizeof(float) * M * N)); + HIP_API_CALL(hipMalloc(&d_A, sizeof(float) * m * k)); + HIP_API_CALL(hipMalloc(&d_B, sizeof(float) * k * n)); + HIP_API_CALL(hipMalloc(&d_Out, sizeof(float) * m * n)); - // Copy data to GPU - HIP_API_CALL(hipMemcpy(d_A, A.data(), sizeof(float) * M * K, hipMemcpyHostToDevice)); - HIP_API_CALL(hipMemcpy(d_B, B.data(), sizeof(float) * K * N, hipMemcpyHostToDevice)); - - // Run the kernel dim3 block_size(BLOCK_SIZE_X, BLOCK_SIZE_Y); - dim3 grid_size((M + block_size.x - 1) / block_size.x, (N + block_size.y - 1) / block_size.y); - matrix_multiply<<>>(d_A, d_B, d_Out, M, N, K); + dim3 grid_size((m + block_size.x - 1) / block_size.x, (n + block_size.y - 1) / block_size.y); + + matrix_multiply<<>>( + d_A, d_B, d_Out, static_cast(m), static_cast(n), static_cast(k)); check_hip_error(); - matrix_multiply_tile<<>>(d_A, d_B, d_Out, M, N, K); + + matrix_multiply_tile<<>>( + d_A, d_B, d_Out, static_cast(m), static_cast(n), static_cast(k)); check_hip_error(); - // Copy data back to CPU - HIP_API_CALL(hipMemcpy(Out.data(), d_Out, sizeof(float) * M * N, hipMemcpyDeviceToHost)); + HIP_API_CALL(hipMemcpy(Out.data(), d_Out, sizeof(float) * m * n, hipMemcpyDeviceToHost)); - // Free GPU Memory HIP_API_CALL(hipFree(d_A)); HIP_API_CALL(hipFree(d_B)); HIP_API_CALL(hipFree(d_Out)); @@ -199,10 +191,14 @@ main(int /*argc*/, char** /*argv*/) assert(status == hipSuccess); assert(deviceId == currDeviceId); + hipDeviceProp_t device_props{}; + HIP_API_CALL(hipGetDeviceProperties(&device_props, deviceId)); + std::string arch_name = device_props.gcnArchName; + for(int i = 0; i < 1; i++) { std::cout << "<<< MatMul starts" << std::endl; - run_hip_app(); + run_hip_app(arch_name); std::cout << ">>> MatMul ends" << std::endl; } diff --git a/projects/rocprofiler-sdk/tests/pc_sampling/pcs.cpp b/projects/rocprofiler-sdk/tests/pc_sampling/pcs.cpp index 2570dc2907b1..df10edc280c7 100644 --- a/projects/rocprofiler-sdk/tests/pc_sampling/pcs.cpp +++ b/projects/rocprofiler-sdk/tests/pc_sampling/pcs.cpp @@ -326,6 +326,15 @@ print_sample_common_fields(std::ostream& os, const PcSamplingRecordT* pc_sample) << "wave_in_group: " << std::setw(2) << static_cast(pc_sample->wave_in_group) << ", " << "chiplet: " << std::setw(2) << static_cast(pc_sample->hw_id.chiplet) << ", " + << "shader engine: " << std::setw(2) + << static_cast(pc_sample->hw_id.shader_engine_id) << ", " + << "shader array: " << std::setw(2) + << static_cast(pc_sample->hw_id.shader_array_id) << ", " + << "CU (WGP): " << std::setw(2) << static_cast(pc_sample->hw_id.cu_or_wgp_id) + << ", " + << "wave slot: " << std::setw(2) << static_cast(pc_sample->hw_id.wave_id) + << ", " + << "simd: " << std::setw(2) << static_cast(pc_sample->hw_id.simd_id) << ", " << "dispatch_id: " << std::setw(7) << pc_sample->dispatch_id << "," << "correlation: {internal=" << std::setw(7) << pc_sample->correlation_id.internal << ", " << "external=" << std::setw(5) << pc_sample->correlation_id.external.value << "}, "; diff --git a/projects/rocprofiler-sdk/tests/pytest-packages/pc_sampling/exec_mask_manipulation/csv.py b/projects/rocprofiler-sdk/tests/pytest-packages/pc_sampling/exec_mask_manipulation/csv.py index 0d8470009f0e..2f9421bcfc0c 100644 --- a/projects/rocprofiler-sdk/tests/pytest-packages/pc_sampling/exec_mask_manipulation/csv.py +++ b/projects/rocprofiler-sdk/tests/pytest-packages/pc_sampling/exec_mask_manipulation/csv.py @@ -27,10 +27,10 @@ import pandas as pd -def stochastic_assert(df, df_condition_selection, max_failing_samples=20): +def stochastic_assert(df, df_condition_selection, max_failing_samples=30): # TODO: When asserting certain conditions related to exec_masks for all samples, # we observe some failures. - # This usually happens because some small number of samples (e.g., 1-20 out of 100k) + # This usually happens because some small number of samples (e.g., 1-30 out of 100k) # do not satisfy the condition. This is either a regression in the ROCr 2nd level trap # handler (as sometimes execution mask or correlation ID mismatches), or # just stochastic nature of the sampling (meaning our checks are too strict). diff --git a/projects/rocprofiler-sdk/tests/pytest-packages/pc_sampling/transpose_multiple_agents/csv.py b/projects/rocprofiler-sdk/tests/pytest-packages/pc_sampling/transpose_multiple_agents/csv.py index af7af4283335..7ceac9fabba8 100644 --- a/projects/rocprofiler-sdk/tests/pytest-packages/pc_sampling/transpose_multiple_agents/csv.py +++ b/projects/rocprofiler-sdk/tests/pytest-packages/pc_sampling/transpose_multiple_agents/csv.py @@ -35,8 +35,8 @@ def validate_all_agents_are_sampled( input_kernel_trace_csv: pd.DataFrame, input_agent_info_csv: pd.DataFrame, ): - transpose_kernel_source_line_start = 137 - transpose_kernel_source_line_end = 145 + transpose_kernel_source_line_start = 181 + transpose_kernel_source_line_end = 189 gfx9_gfx12_agents_df = input_agent_info_csv[ input_agent_info_csv["Name"].apply( @@ -78,7 +78,7 @@ def validate_all_agents_are_sampled( # extract decoded samples that are mapped to the transpose.cpp file transpose_samples_df = samples_df[ samples_df["Instruction_Comment"].apply( - lambda comment: "transpose-all-agents.cpp" in comment + lambda comment: "transpose.cpp" in comment ) ].copy() # determine the line number for each sample diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/CMakeLists.txt index a5f85c983854..c12170faa1eb 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/CMakeLists.txt @@ -64,6 +64,8 @@ rocprofiler_add_integration_validate_test( ${CMAKE_CURRENT_BINARY_DIR}/pc_sampling_cmd_input/out_pc_sampling_stochastic.csv --input-json ${CMAKE_CURRENT_BINARY_DIR}/pc_sampling_cmd_input/out_results.json + --input-agent-info-csv + ${CMAKE_CURRENT_BINARY_DIR}/pc_sampling_cmd_input/out_agent_info.csv --all-sampled False TIMEOUT 60 @@ -99,6 +101,8 @@ rocprofiler_add_integration_validate_test( ${CMAKE_CURRENT_BINARY_DIR}/pc_sampling_json_input/out_pc_sampling_stochastic.csv --input-json ${CMAKE_CURRENT_BINARY_DIR}/pc_sampling_json_input/out_results.json + --input-agent-info-csv + ${CMAKE_CURRENT_BINARY_DIR}/pc_sampling_cmd_input/out_agent_info.csv --all-sampled False TIMEOUT 60 @@ -133,6 +137,8 @@ rocprofiler_add_integration_validate_test( ${CMAKE_CURRENT_BINARY_DIR}/pc_sampling_yaml_input/out_pc_sampling_stochastic.csv --input-json ${CMAKE_CURRENT_BINARY_DIR}/pc_sampling_yaml_input/out_results.json + --input-agent-info-csv + ${CMAKE_CURRENT_BINARY_DIR}/pc_sampling_cmd_input/out_agent_info.csv --all-sampled False TIMEOUT 60 diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/conftest.py index 0d7d59f5a2d2..ec9932d80dd6 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/conftest.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/conftest.py @@ -50,6 +50,12 @@ def pytest_addoption(parser): help="All SW and HW units must be sampled.", ) + parser.addoption( + "--input-agent-info-csv", + action="store", + help="Path to CSV file containing agents information.", + ) + @pytest.fixture def input_csv(request): @@ -96,3 +102,17 @@ def input_json(request): def all_sampled(request): _all_sampled_str = request.config.getoption("--all-sampled") return _all_sampled_str == "True" + + +@pytest.fixture +def input_agent_info_csv(request): + filename = request.config.getoption("--input-agent-info-csv") + with open(filename, "r") as inp: + return pd.read_csv( + inp, + na_filter=False, # parse empty fields as "" + keep_default_na=False, # parse empty fields as "" + dtype={ + "Name": str, + }, + ) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/validate.py index 629b461f3124..fd55f99dd319 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/exec-mask-manipulation/validate.py @@ -62,7 +62,12 @@ def test_validate_pc_sampling_exec_mask_manipulation_json( # ======================== Validating fields specific for stochastic sampling -def test_validate_pc_sampling_stochastic_specific_csv(input_csv: pd.DataFrame): +def test_validate_pc_sampling_stochastic_specific_csv( + input_csv: pd.DataFrame, input_agent_info_csv: pd.DataFrame +): + if tmp_disable_for_gfx12(input_agent_info_csv): + pytest.skip("Stochastic sampling specific checks are not implemented for GFX12") + from rocprofiler_sdk.pc_sampling.stochastic.csv.gfx9 import ( validate_stochastic_samples_csv, ) @@ -70,7 +75,12 @@ def test_validate_pc_sampling_stochastic_specific_csv(input_csv: pd.DataFrame): validate_stochastic_samples_csv(input_csv) -def test_validate_pc_sampling_stochastic_specific_json(input_json): +def test_validate_pc_sampling_stochastic_specific_json( + input_json, input_agent_info_csv: pd.DataFrame +): + if tmp_disable_for_gfx12(input_agent_info_csv): + pytest.skip("Stochastic sampling specific checks are not implemented for GFX12") + from rocprofiler_sdk.pc_sampling.stochastic.json.gfx9 import ( validate_stochastic_samples_json, ) @@ -78,6 +88,14 @@ def test_validate_pc_sampling_stochastic_specific_json(input_json): validate_stochastic_samples_json(input_json["rocprofiler-sdk-tool"]) +def tmp_disable_for_gfx12(input_agent_info_csv: pd.DataFrame): + """ + If any of the agents are from GFX12 family, we temporarily disable + stochastic sampling specific checks, because they're not fully implemented. + """ + return input_agent_info_csv["Name"].str.contains("gfx12").any() + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/transpose-multiple-agents/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/transpose-multiple-agents/validate.py index 3c8e2cecac92..7e1d25e33ba3 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/transpose-multiple-agents/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/pc-sampling/stochastic/transpose-multiple-agents/validate.py @@ -47,7 +47,12 @@ def test_multi_agent_support( # =================== validation specific to stochastic sampling -def test_validate_pc_sampling_stochastic_specific_csv(input_samples_csv: pd.DataFrame): +def test_validate_pc_sampling_stochastic_specific_csv( + input_samples_csv: pd.DataFrame, input_agent_info_csv: pd.DataFrame +): + if tmp_disable_for_gfx12(input_agent_info_csv): + pytest.skip("Stochastic sampling specific checks are not implemented for GFX12") + from rocprofiler_sdk.pc_sampling.stochastic.csv.gfx9 import ( validate_stochastic_samples_csv, ) @@ -55,7 +60,12 @@ def test_validate_pc_sampling_stochastic_specific_csv(input_samples_csv: pd.Data validate_stochastic_samples_csv(input_samples_csv) -def test_validate_pc_sampling_stochastic_specific_json(input_samples_json): +def test_validate_pc_sampling_stochastic_specific_json( + input_samples_json, input_agent_info_csv: pd.DataFrame +): + if tmp_disable_for_gfx12(input_agent_info_csv): + pytest.skip("Stochastic sampling specific checks are not implemented for GFX12") + from rocprofiler_sdk.pc_sampling.stochastic.json.gfx9 import ( validate_stochastic_samples_json, ) @@ -63,6 +73,14 @@ def test_validate_pc_sampling_stochastic_specific_json(input_samples_json): validate_stochastic_samples_json(input_samples_json["rocprofiler-sdk-tool"]) +def tmp_disable_for_gfx12(input_agent_info_csv: pd.DataFrame): + """ + If any of the agents are from GFX12 family, we temporarily disable + stochastic sampling specific checks, because they're not fully implemented. + """ + return input_agent_info_csv["Name"].str.contains("gfx12").any() + + if __name__ == "__main__": exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) sys.exit(exit_code)