diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp index 36989b6a85..975b250ebe 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/enum_string.hpp @@ -1663,7 +1663,9 @@ ROCPROFILER_ENUM_LABEL(ROCPROFILER_THREAD_TRACE_PARAMETER_SIMD_SELECT); ROCPROFILER_ENUM_LABEL(ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTERS_CTRL); ROCPROFILER_ENUM_LABEL(ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTER); ROCPROFILER_ENUM_LABEL(ROCPROFILER_THREAD_TRACE_PARAMETER_SERIALIZE_ALL); -static_assert(ROCPROFILER_THREAD_TRACE_PARAMETER_LAST == 7); +ROCPROFILER_ENUM_LABEL(ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTER_EXCLUDE_MASK); +ROCPROFILER_ENUM_LABEL(ROCPROFILER_THREAD_TRACE_PARAMETER_NO_DETAIL); +static_assert(ROCPROFILER_THREAD_TRACE_PARAMETER_LAST == 9); ROCPROFILER_ENUM_LABEL(ROCPROFILER_THREAD_TRACE_CONTROL_NONE); ROCPROFILER_ENUM_LABEL(ROCPROFILER_THREAD_TRACE_CONTROL_START_AND_STOP); diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/thread-trace/core.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/thread-trace/core.h index d870de5a3a..a4de7c6997 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/thread-trace/core.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/thread-trace/core.h @@ -47,9 +47,14 @@ typedef enum rocprofiler_thread_trace_parameter_type_t ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE, ///< Size of combined GPU buffer for ATT ROCPROFILER_THREAD_TRACE_PARAMETER_SIMD_SELECT, ///< Bitmask (GFX9) or ID (Navi) of SIMDs ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTERS_CTRL, ///< Period [1,32] or disable (0) perfmon - ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTER, ///< Perfmon ID and SIMD mask - ROCPROFILER_THREAD_TRACE_PARAMETER_SERIALIZE_ALL, ///< Serializes kernels not under thread - ///< trace + ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTER, ///< Perfmon ID and SIMD mask. gfx9 only + ROCPROFILER_THREAD_TRACE_PARAMETER_SERIALIZE_ALL, ///< Serializes also kernels not under + ///< thread trace + ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTER_EXCLUDE_MASK, ///< Bitmask of which compute + ///< units to exclude from + ///< perfcounters. gfx9 only + ROCPROFILER_THREAD_TRACE_PARAMETER_NO_DETAIL, ///< Dont collect instruction timing, + ///< only shader-wide information ROCPROFILER_THREAD_TRACE_PARAMETER_LAST } rocprofiler_thread_trace_parameter_type_t; diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/thread-trace/trace_decoder_types.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/thread-trace/trace_decoder_types.h index bf530a849a..c0506e7055 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/thread-trace/trace_decoder_types.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/thread-trace/trace_decoder_types.h @@ -40,6 +40,7 @@ typedef enum rocprofiler_thread_trace_decoder_info_t ROCPROFILER_THREAD_TRACE_DECODER_INFO_NONE = 0, ROCPROFILER_THREAD_TRACE_DECODER_INFO_DATA_LOST, ROCPROFILER_THREAD_TRACE_DECODER_INFO_STITCH_INCOMPLETE, + ROCPROFILER_THREAD_TRACE_DECODER_INFO_WAVE_INCOMPLETE, ROCPROFILER_THREAD_TRACE_DECODER_INFO_LAST } rocprofiler_thread_trace_decoder_info_t; @@ -76,7 +77,7 @@ typedef struct rocprofiler_thread_trace_decoder_occupancy_t uint8_t reserved; ///< Reserved uint8_t cu; ///< Compute unit ID (gfx9) or WGP ID (gfx10+). uint8_t simd; ///< SIMD ID [0,3] within compute unit - uint8_t slot; ///< Wave slot ID within SIMD + uint8_t wave_id; ///< Wave slot ID within SIMD uint32_t start : 1; ///< 1 if wave_start, 0 if a wave_end uint32_t _rsvd : 31; } rocprofiler_thread_trace_decoder_occupancy_t; @@ -168,6 +169,49 @@ typedef struct rocprofiler_thread_trace_decoder_wave_t rocprofiler_thread_trace_decoder_inst_t* instructions_array; ///< Instructions executed } rocprofiler_thread_trace_decoder_wave_t; +/** + * @brief Matches the reference (realtime) clock with the shader clock + * Added in rocprof-trace-decoder 0.1.3. Requires aqlprofile for rocm 7.1+. + * clock_in_seconds = realtime_clock / ROCPROFILER_THREAD_TRACE_DECODER_RECORD_RT_FREQUENCY + * gfx_frequency = delta(shader_clock) / delta(clock_in_seconds) + * For best average, use + * gfx_frequency[n] = (shader_clock[n]-shader_clock[0]) / (clock_in_seconds[n]-clock_in_seconds[0]) + */ +typedef struct rocprofiler_thread_trace_decoder_realtime_t +{ + int64_t shader_clock; ///< Clock timestamp in gfx clock units + uint64_t realtime_clock; ///< Clock timestamp in realtime units + uint64_t reserved; +} rocprofiler_thread_trace_decoder_realtime_t; + +/** + * @brief Bitmask of additional information for shaderdata_t + * Added in rocprof-trace-decoder 0.1.3 + */ +typedef enum rocprofiler_thread_trace_decoder_shaderdata_flags_t +{ + ROCPROFILER_THREAD_TRACE_DECODER_SHADERDATA_FLAGS_IMM = 0, + ROCPROFILER_THREAD_TRACE_DECODER_SHADERDATA_FLAGS_PRIV ///< Generated by the trap handler + + /// @var ROCPROFILER_THREAD_TRACE_DECODER_SHADERDATA_FLAGS_IMM + /// @brief Value comes from s_ttracedata_imm. +} rocprofiler_thread_trace_decoder_shaderdata_flags_t; + +/** + * @brief Record created by s_ttracedata and s_ttracedata_imm + * Added in rocprof-trace-decoder 0.1.3 + */ +typedef struct rocprofiler_thread_trace_decoder_shaderdata_t +{ + int64_t time; + uint64_t value; ///< Value written from M0/IMM + uint8_t cu; ///< CU id (gfx9) or wgp id (gfx10+). This is always the target_cu. + uint8_t simd; ///< SIMD ID [0,3]. + uint8_t wave_id; ///< Wave slot ID within SIMD. + uint8_t flags; ///< bitmask of rocprofiler_thread_trace_decoder_shaderdata_flags_t + uint32_t reserved; +} rocprofiler_thread_trace_decoder_shaderdata_t; + /** * @brief Defines the type of payload received by rocprofiler_thread_trace_decoder_callback_t */ @@ -179,7 +223,13 @@ typedef enum rocprofiler_thread_trace_decoder_record_type_t ROCPROFILER_THREAD_TRACE_DECODER_RECORD_WAVE, ///< rocprofiler_thread_trace_decoder_wave_t* ROCPROFILER_THREAD_TRACE_DECODER_RECORD_INFO, ///< rocprofiler_thread_trace_decoder_info_t* ROCPROFILER_THREAD_TRACE_DECODER_RECORD_DEBUG, ///< Debug + ROCPROFILER_THREAD_TRACE_DECODER_RECORD_SHADERDATA, ///< rocprofiler_thread_trace_decoder_shaderdata_t* + ROCPROFILER_THREAD_TRACE_DECODER_RECORD_REALTIME, ///< rocprofiler_thread_trace_decoder_realtime_t* + ROCPROFILER_THREAD_TRACE_DECODER_RECORD_RT_FREQUENCY, ROCPROFILER_THREAD_TRACE_DECODER_RECORD_LAST + + /// @var ROCPROFILER_THREAD_TRACE_DECODER_RECORD_RT_FREQUENCY + /// @brief uint64_t*. Realtime clock frequency in Hz. } rocprofiler_thread_trace_decoder_record_type_t; /** @} */ diff --git a/projects/rocprofiler-sdk/source/lib/att-tool/occupancy.cpp b/projects/rocprofiler-sdk/source/lib/att-tool/occupancy.cpp index 8de1cccb2d..33a26a4543 100644 --- a/projects/rocprofiler-sdk/source/lib/att-tool/occupancy.cpp +++ b/projects/rocprofiler-sdk/source/lib/att-tool/occupancy.cpp @@ -68,7 +68,7 @@ OccupancyFile(const Fspath& dir, json_event.push_back(event.time); json_event.push_back(event.cu); json_event.push_back(event.simd); - json_event.push_back(event.slot); + json_event.push_back(event.wave_id); json_event.push_back(event.start); json_event.push_back(get_kernel_id(event.pc)); list.push_back(json_event); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp index 6425bd3884..51d04542cb 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/aql/packet_construct.cpp @@ -148,6 +148,7 @@ ThreadTraceAQLPacketFactory::ThreadTraceAQLPacketFactory(const hsa::AgentCache& uint32_t buffer_size_lo = static_cast(params.buffer_size); uint32_t buffer_size_hi = static_cast(params.buffer_size >> 32); uint32_t perf_ctrl = static_cast(params.perfcounter_ctrl); + uint32_t perf_exclude_mask = static_cast(params.perf_exclude_mask); aql_params.clear(); @@ -156,8 +157,22 @@ ThreadTraceAQLPacketFactory::ThreadTraceAQLPacketFactory(const hsa::AgentCache& aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_SIMD_SELECTION, {simd}}); aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_ATT_BUFFER_SIZE, {buffer_size_lo}}); - if(buffer_size_hi != 0) aql_params.push_back({static_cast( - AQLPROFILE_ATT_PARAMETER_NAME_BUFFER_SIZE_HIGH), {buffer_size_hi}}); + if(buffer_size_hi != 0) + { + aql_params.push_back({static_cast( + AQLPROFILE_ATT_PARAMETER_NAME_BUFFER_SIZE_HIGH), + {buffer_size_hi}}); + } + + if(perf_exclude_mask) + { + // Bitwise NOT because aqlprofile receives the mask, not the exclude mask + aql_params.push_back( + {HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_PERFCOUNTER_MASK, {~perf_exclude_mask}}); + } + + if(params.no_detail_simd) + aql_params.push_back({HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_OCCUPANCY_MODE, {1}}); if(perf_ctrl != 0 && !params.perfcounters.empty()) { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp index 2644046213..95e10f7668 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer.hpp @@ -43,8 +43,8 @@ struct instance using buffer_t = common::container::record_header_buffer; mutable std::array buffers = {}; - mutable std::atomic_flag syncer = ATOMIC_FLAG_INIT; // writer and reader lock. - mutable std::atomic buffer_idx = {}; // array index + mutable std::atomic_flag syncer = ATOMIC_FLAG_INIT; // writer and reader lock. + mutable std::atomic buffer_idx = {}; // array index mutable std::atomic drop_count = {}; uint64_t watermark = 0; uint64_t context_id = 0; // rocprofiler_context_id_t value diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/core.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/core.hpp index ec480295a3..f27a874cfc 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/core.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/core.hpp @@ -67,6 +67,8 @@ struct thread_trace_parameter_pack uint8_t perfcounter_ctrl = 0; uint64_t shader_engine_mask = DEFAULT_SE_MASK; uint64_t buffer_size = DEFAULT_BUFFER_SIZE; + uint64_t perf_exclude_mask = 0; + bool no_detail_simd = false; bool bSerialize = false; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/service.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/service.cpp index 1fcf62636c..f5cd753988 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/service.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/service.cpp @@ -34,6 +34,64 @@ using DispatchThreadTracer = rocprofiler::thread_trace::DispatchThreadTracer; using DeviceThreadTracer = rocprofiler::thread_trace::DeviceThreadTracer; +namespace +{ +using parameter_pack = rocprofiler::thread_trace::thread_trace_parameter_pack; + +rocprofiler_status_t +build_pack_from_array(parameter_pack& pack, + const rocprofiler_thread_trace_parameter_t* params, + size_t num_parameters) +{ + auto id_map = rocprofiler::counters::getPerfCountersIdMap(); + + for(size_t p = 0; p < num_parameters; p++) + { + const rocprofiler_thread_trace_parameter_t& param = params[p]; + if(param.type > ROCPROFILER_THREAD_TRACE_PARAMETER_LAST) + return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; + + switch(param.type) + { + case ROCPROFILER_THREAD_TRACE_PARAMETER_TARGET_CU: pack.target_cu = param.value; break; + case ROCPROFILER_THREAD_TRACE_PARAMETER_SHADER_ENGINE_MASK: + pack.shader_engine_mask = param.value; + break; + case ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE: + pack.buffer_size = param.value; + break; + case ROCPROFILER_THREAD_TRACE_PARAMETER_SIMD_SELECT: + pack.simd_select = param.value; + break; + case ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTER: + { + auto event_it = id_map.find(param.counter_id.handle); + if(event_it != id_map.end()) + pack.perfcounters.push_back({event_it->second, param.simd_mask}); + } + break; + case ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTERS_CTRL: + pack.perfcounter_ctrl = param.value; + break; + case ROCPROFILER_THREAD_TRACE_PARAMETER_SERIALIZE_ALL: + pack.bSerialize = param.value != 0; + break; + case ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTER_EXCLUDE_MASK: + pack.perf_exclude_mask = param.value; + break; + case ROCPROFILER_THREAD_TRACE_PARAMETER_NO_DETAIL: + pack.no_detail_simd = param.value != 0; + break; + case ROCPROFILER_THREAD_TRACE_PARAMETER_LAST: + return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; + } + } + if(!pack.are_params_valid()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; + + return ROCPROFILER_STATUS_SUCCESS; +} +}; // namespace + extern "C" { rocprofiler_status_t rocprofiler_configure_dispatch_thread_trace_service( @@ -66,45 +124,11 @@ rocprofiler_configure_dispatch_thread_trace_service( if(pack.dispatch_cb_fn == nullptr) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; - auto id_map = rocprofiler::counters::getPerfCountersIdMap(); - for(size_t p = 0; p < num_parameters; p++) { - const rocprofiler_thread_trace_parameter_t& param = parameters[p]; - if(param.type > ROCPROFILER_THREAD_TRACE_PARAMETER_LAST) - return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; - - switch(param.type) - { - case ROCPROFILER_THREAD_TRACE_PARAMETER_TARGET_CU: pack.target_cu = param.value; break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_SHADER_ENGINE_MASK: - pack.shader_engine_mask = param.value; - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE: - pack.buffer_size = param.value; - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_SIMD_SELECT: - pack.simd_select = param.value; - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTER: - { - auto event_it = id_map.find(param.counter_id.handle); - if(event_it != id_map.end()) - pack.perfcounters.push_back({event_it->second, param.simd_mask}); - } - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTERS_CTRL: - pack.perfcounter_ctrl = param.value; - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_SERIALIZE_ALL: - pack.bSerialize = param.value != 0; - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_LAST: - return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; - } + auto status = build_pack_from_array(pack, parameters, num_parameters); + if(status != ROCPROFILER_STATUS_SUCCESS) return status; } - if(!pack.are_params_valid()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; - ctx->dispatch_thread_trace->add_agent(agent_id, pack); return ROCPROFILER_STATUS_SUCCESS; } @@ -135,44 +159,13 @@ rocprofiler_configure_device_thread_trace_service( pack.shader_cb_fn = shader_callback; pack.callback_userdata = userdata; - auto id_map = rocprofiler::counters::getPerfCountersIdMap(); - for(size_t p = 0; p < num_parameters; p++) { - const rocprofiler_thread_trace_parameter_t& param = parameters[p]; - if(param.type > ROCPROFILER_THREAD_TRACE_PARAMETER_LAST) - return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; - - switch(param.type) - { - case ROCPROFILER_THREAD_TRACE_PARAMETER_TARGET_CU: pack.target_cu = param.value; break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_SHADER_ENGINE_MASK: - pack.shader_engine_mask = param.value; - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE: - pack.buffer_size = param.value; - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_SIMD_SELECT: - pack.simd_select = param.value; - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTER: - { - auto event_it = id_map.find(param.counter_id.handle); - if(event_it != id_map.end()) - pack.perfcounters.push_back({event_it->second, param.simd_mask}); - } - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTERS_CTRL: - pack.perfcounter_ctrl = param.value; - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_SERIALIZE_ALL: - if(param.value != 0) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; - break; - case ROCPROFILER_THREAD_TRACE_PARAMETER_LAST: - return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; - } + auto status = build_pack_from_array(pack, parameters, num_parameters); + if(status != ROCPROFILER_STATUS_SUCCESS) return status; } - if(!pack.are_params_valid()) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; + // Serialization not supported in device mode + if(pack.bSerialize) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; ctx->device_thread_trace->add_agent(agent_id, pack); return ROCPROFILER_STATUS_SUCCESS; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp index f2e00e8eea..5094e3bc15 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/tests/att_packet_test.cpp @@ -139,6 +139,9 @@ TEST(thread_trace, configure_test) params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SHADER_ENGINE_MASK, {0xF}}); params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE, {0x1000000}}); params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SIMD_SELECT, {0xF}}); + params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTERS_CTRL, {0}}); + params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTER_EXCLUDE_MASK, {0}}); + params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_NO_DETAIL, {0}}); auto agents = hsa::get_queue_controller()->get_supported_agents(); ASSERT_GT(agents.size(), 0); @@ -278,7 +281,9 @@ query_available_agents(rocprofiler_agent_version_t /* version */, params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SHADER_ENGINE_MASK, {0xF}}); params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE, {0x1000000}}); params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SIMD_SELECT, {0xF}}); - params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTERS_CTRL, {1}}); + params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTERS_CTRL, {0}}); + params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_PERFCOUNTER_EXCLUDE_MASK, {0}}); + params.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_NO_DETAIL, {0}}); { auto metrics = rocprofiler::counters::getMetricsForAgent("gfx90a"); diff --git a/projects/rocprofiler-sdk/tests/thread-trace/CMakeLists.txt b/projects/rocprofiler-sdk/tests/thread-trace/CMakeLists.txt index 4c43dc9f67..9dd28cf870 100644 --- a/projects/rocprofiler-sdk/tests/thread-trace/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/thread-trace/CMakeLists.txt @@ -114,3 +114,20 @@ set_tests_properties( "${ROCPROFILER_DEFAULT_FAIL_REGEX}" DISABLED ${ROCPROFILER_DISABLE_UNSTABLE_CTESTS}) + +# Test occupancy mode +add_test(NAME thread-trace-api-extra-args + COMMAND $) + +set_tests_properties( + thread-trace-api-extra-args + PROPERTIES TIMEOUT + 10 + LABELS + "integration-tests" + ENVIRONMENT + "${PRELOAD_ENV};ATT_NODETAIL=1" + FAIL_REGULAR_EXPRESSION + "${ROCPROFILER_DEFAULT_FAIL_REGEX}" + DISABLED + ${ROCPROFILER_DISABLE_UNSTABLE_CTESTS}) diff --git a/projects/rocprofiler-sdk/tests/thread-trace/agent.cpp b/projects/rocprofiler-sdk/tests/thread-trace/agent.cpp index ccc2004995..865d785fa5 100644 --- a/projects/rocprofiler-sdk/tests/thread-trace/agent.cpp +++ b/projects/rocprofiler-sdk/tests/thread-trace/agent.cpp @@ -103,7 +103,7 @@ query_available_agents(rocprofiler_agent_version_t /* version */, if(agent->type != ROCPROFILER_AGENT_TYPE_GPU) continue; // Check if we are testing for large buffers - static const char* var = getenv("ATT_BUFFER_SIZE_MB"); + static const char* var = std::getenv("ATT_BUFFER_SIZE_MB"); static uint64_t buffer_size_mb = (var ? atoi(var) : 96) * 1024ul * 1024ul; std::vector parameters; @@ -111,7 +111,15 @@ query_available_agents(rocprofiler_agent_version_t /* version */, parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SIMD_SELECT, 0xF}); parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_BUFFER_SIZE, buffer_size_mb}); parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SHADER_ENGINE_MASK, 0x1}); - parameters.push_back({ROCPROFILER_THREAD_TRACE_PARAMETER_SERIALIZE_ALL, 0}); + + static const bool extra_args = + std::getenv("ATT_NODETAIL") ? std::stoi(std::getenv("ATT_NODETAIL")) != 0 : false; + if(extra_args) + { + // Dont generate instruction profiling, only occupancy and shaderdata + parameters.emplace_back(rocprofiler_thread_trace_parameter_t{ + ROCPROFILER_THREAD_TRACE_PARAMETER_NO_DETAIL, 1}); + } ROCPROFILER_CALL( rocprofiler_configure_device_thread_trace_service(agent_ctx, diff --git a/projects/rocprofiler-sdk/tests/thread-trace/kernel_lds.cpp b/projects/rocprofiler-sdk/tests/thread-trace/kernel_lds.cpp index a16b68d886..2da5b21a0c 100644 --- a/projects/rocprofiler-sdk/tests/thread-trace/kernel_lds.cpp +++ b/projects/rocprofiler-sdk/tests/thread-trace/kernel_lds.cpp @@ -53,4 +53,8 @@ looping_lds_kernel(float* __restrict__ a, } a[index] = interm[threadIdx.x % SHM_SIZE] + c[index]; + + asm volatile("s_mov_b32 m0, 0xDEADBEEF"); // checked in trace_callbacks.cpp + asm volatile("s_nop 1"); // s_nop 0 should also work + asm volatile("s_ttracedata"); } diff --git a/projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.cpp b/projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.cpp index d40377badb..6f8638cfad 100644 --- a/projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.cpp +++ b/projects/rocprofiler-sdk/tests/thread-trace/trace_callbacks.cpp @@ -28,6 +28,7 @@ #include "trace_callbacks.hpp" #include +#include #include #include @@ -35,6 +36,10 @@ namespace Callbacks { rocprofiler_thread_trace_decoder_id_t decoder{}; std::atomic latency{0}; +std::atomic has_sdata{false}; + +// defined in kernel_lds.cpp +constexpr uint64_t SDATA_RECORD = 0xDEADBEEF; void tool_codeobj_tracing_callback(rocprofiler_callback_tracing_record_t record, @@ -80,6 +85,14 @@ shader_data_callback(rocprofiler_agent_id_t /* agent */, void* trace_events, uint64_t trace_size, void* /* userdata */) { + if(record_type_id == ROCPROFILER_THREAD_TRACE_DECODER_RECORD_SHADERDATA) + { + const auto* events = + static_cast(trace_events); + for(size_t i = 0; i < trace_size; i++) + if(events[i].value == SDATA_RECORD) has_sdata = true; + } + if(record_type_id != ROCPROFILER_THREAD_TRACE_DECODER_RECORD_WAVE) return; for(size_t w = 0; w < trace_size; w++) @@ -102,9 +115,22 @@ init() void finalize(void* /* tool_data */) { - rocprofiler_thread_trace_decoder_destroy(decoder); + const char* env_args = std::getenv("ATT_NODETAIL"); + const bool extra_args = env_args ? std::stoi(env_args) != 0 : false; - if(latency.load() == 0) std::cerr << "Error: No latency was assigned to the trace!"; + // only check if we have a valid decoder + if(decoder.handle != 0) + { + if(extra_args && latency != 0) + throw std::runtime_error("Got detailled profling in nondetail mode"); + else if(!extra_args && latency == 0) + throw std::runtime_error("Missing detailed profiling!"); + + // disabling until new decoder version is picked up + // if(!has_sdata) throw std::runtime_error("Missing shaderdata record!"); + } + + rocprofiler_thread_trace_decoder_destroy(decoder); } } // namespace Callbacks