Adding new trace decoder record types and new ATT parameters (#195)
* Adding new trace decoder record types and new ATT parameters * Add compatiblity with decoder 0.1.2 * Added RT * Format * Add logging to sdata values * Review comment * Review comments * Update projects/rocprofiler-sdk/source/include/rocprofiler-sdk/experimental/thread-trace/trace_decoder_types.h
このコミットが含まれているのは:
Giovanni Lenzi Baraldi
GitHub
@@ -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);
|
||||
|
||||
+8
-3
@@ -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;
|
||||
|
||||
|
||||
+51
-1
@@ -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;
|
||||
|
||||
/** @} */
|
||||
|
||||
@@ -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);
|
||||
|
||||
@@ -148,6 +148,7 @@ ThreadTraceAQLPacketFactory::ThreadTraceAQLPacketFactory(const hsa::AgentCache&
|
||||
uint32_t buffer_size_lo = static_cast<uint32_t>(params.buffer_size);
|
||||
uint32_t buffer_size_hi = static_cast<uint32_t>(params.buffer_size >> 32);
|
||||
uint32_t perf_ctrl = static_cast<uint32_t>(params.perfcounter_ctrl);
|
||||
uint32_t perf_exclude_mask = static_cast<uint32_t>(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<hsa_ven_amd_aqlprofile_parameter_name_t>(
|
||||
AQLPROFILE_ATT_PARAMETER_NAME_BUFFER_SIZE_HIGH), {buffer_size_hi}});
|
||||
if(buffer_size_hi != 0)
|
||||
{
|
||||
aql_params.push_back({static_cast<hsa_ven_amd_aqlprofile_parameter_name_t>(
|
||||
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())
|
||||
{
|
||||
|
||||
@@ -43,8 +43,8 @@ struct instance
|
||||
using buffer_t = common::container::record_header_buffer;
|
||||
|
||||
mutable std::array<buffer_t, 2> buffers = {};
|
||||
mutable std::atomic_flag syncer = ATOMIC_FLAG_INIT; // writer and reader lock.
|
||||
mutable std::atomic<uint32_t> buffer_idx = {}; // array index
|
||||
mutable std::atomic_flag syncer = ATOMIC_FLAG_INIT; // writer and reader lock.
|
||||
mutable std::atomic<uint32_t> buffer_idx = {}; // array index
|
||||
mutable std::atomic<uint64_t> drop_count = {};
|
||||
uint64_t watermark = 0;
|
||||
uint64_t context_id = 0; // rocprofiler_context_id_t value
|
||||
|
||||
@@ -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;
|
||||
|
||||
|
||||
@@ -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;
|
||||
|
||||
+6
-1
@@ -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");
|
||||
|
||||
@@ -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 $<TARGET_FILE:thread-trace-api-agent-test>)
|
||||
|
||||
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})
|
||||
|
||||
@@ -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<rocprofiler_thread_trace_parameter_t> 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,
|
||||
|
||||
@@ -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");
|
||||
}
|
||||
|
||||
@@ -28,6 +28,7 @@
|
||||
#include "trace_callbacks.hpp"
|
||||
|
||||
#include <unistd.h>
|
||||
#include <atomic>
|
||||
#include <cassert>
|
||||
#include <fstream>
|
||||
|
||||
@@ -35,6 +36,10 @@ namespace Callbacks
|
||||
{
|
||||
rocprofiler_thread_trace_decoder_id_t decoder{};
|
||||
std::atomic<size_t> latency{0};
|
||||
std::atomic<bool> 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<rocprofiler_thread_trace_decoder_shaderdata_t*>(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
|
||||
|
||||
新しいイシューから参照
ユーザーをブロックする