Remove -Wno-missing-field-initializers from build flags (#810)
* Remove -Wno-missing-field-initializers
- Compiler errors if missing field initializers
* Update lib/rocprofiler-sdk/counters/evaluate_ast.cpp
- copy over dispatch ID in perform_reduction/evaluate
[ROCm/rocprofiler-sdk commit: 48273d6a65]
This commit is contained in:
Jonathan R. Madsen
committed by
GitHub
GitHub
parent
843398eb3c
commit
399f2cc660
@@ -63,9 +63,8 @@ target_link_libraries(rocprofiler-build-flags INTERFACE rocprofiler::rocprofiler
|
||||
# ----------------------------------------------------------------------------------------#
|
||||
# set the compiler flags
|
||||
#
|
||||
rocprofiler_target_compile_options(
|
||||
rocprofiler-build-flags INTERFACE "-W" "-Wall" "-Wno-unknown-pragmas"
|
||||
"-Wno-missing-field-initializers")
|
||||
rocprofiler_target_compile_options(rocprofiler-build-flags
|
||||
INTERFACE "-W" "-Wall" "-Wno-unknown-pragmas")
|
||||
|
||||
# ----------------------------------------------------------------------------------------#
|
||||
# extra flags for debug information in debug or optimized binaries
|
||||
|
||||
@@ -39,7 +39,7 @@ hsa_ven_amd_aqlprofile_id_query_t
|
||||
get_query_info(rocprofiler_agent_id_t agent, const counters::Metric& metric)
|
||||
{
|
||||
auto aql_agent = *CHECK_NOTNULL(rocprofiler::agent::get_aql_agent(agent));
|
||||
aqlprofile_pmc_profile_t profile = {.agent = aql_agent};
|
||||
aqlprofile_pmc_profile_t profile = {.agent = aql_agent, .events = nullptr, .event_count = 0};
|
||||
hsa_ven_amd_aqlprofile_id_query_t query = {metric.block().c_str(), 0, 0};
|
||||
if(aqlprofile_get_pmc_info(&profile, AQLPROFILE_INFO_BLOCK_ID, &query) != HSA_STATUS_SUCCESS)
|
||||
{
|
||||
@@ -113,4 +113,4 @@ get_dim_info(rocprofiler_agent_id_t agent,
|
||||
return ROCPROFILER_STATUS_SUCCESS;
|
||||
}
|
||||
} // namespace aql
|
||||
} // namespace rocprofiler
|
||||
} // namespace rocprofiler
|
||||
|
||||
@@ -103,7 +103,14 @@ findDeviceMetrics(const rocprofiler_agent_t& agent, const std::unordered_set<std
|
||||
hsa_ven_amd_aqlprofile_id_query_t
|
||||
v1_get_query_info(hsa_agent_t agent, const counters::Metric& metric)
|
||||
{
|
||||
hsa_ven_amd_aqlprofile_profile_t profile = {.agent = agent};
|
||||
hsa_ven_amd_aqlprofile_profile_t profile = {.agent = agent,
|
||||
.type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC,
|
||||
.events = nullptr,
|
||||
.event_count = 0,
|
||||
.parameters = nullptr,
|
||||
.parameter_count = 0,
|
||||
.output_buffer = {nullptr, 0},
|
||||
.command_buffer = {nullptr, 0}};
|
||||
hsa_ven_amd_aqlprofile_id_query_t query = {metric.block().c_str(), 0, 0};
|
||||
if(hsa_ven_amd_aqlprofile_get_info(&profile, HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_ID, &query) !=
|
||||
HSA_STATUS_SUCCESS)
|
||||
@@ -120,7 +127,11 @@ v1_get_block_counters(hsa_agent_t agent, const hsa_ven_amd_aqlprofile_event_t& e
|
||||
hsa_ven_amd_aqlprofile_profile_t query = {.agent = agent,
|
||||
.type = HSA_VEN_AMD_AQLPROFILE_EVENT_TYPE_PMC,
|
||||
.events = &event,
|
||||
.event_count = 1};
|
||||
.event_count = 1,
|
||||
.parameters = nullptr,
|
||||
.parameter_count = 0,
|
||||
.output_buffer = {nullptr, 0},
|
||||
.command_buffer = {nullptr, 0}};
|
||||
uint32_t max_block_counters = 0;
|
||||
if(hsa_ven_amd_aqlprofile_get_info(&query,
|
||||
HSA_VEN_AMD_AQLPROFILE_INFO_BLOCK_COUNTERS,
|
||||
|
||||
@@ -58,7 +58,7 @@ get_reduce_op_type_from_string(const std::string& op)
|
||||
std::vector<rocprofiler_record_counter_t>*
|
||||
perform_reduction(ReduceOperation reduce_op, std::vector<rocprofiler_record_counter_t>* input_array)
|
||||
{
|
||||
rocprofiler_record_counter_t result{.id = 0, .counter_value = 0};
|
||||
rocprofiler_record_counter_t result{.id = 0, .counter_value = 0, .dispatch_id = 0};
|
||||
if(input_array->empty()) return input_array;
|
||||
switch(reduce_op)
|
||||
{
|
||||
@@ -81,26 +81,30 @@ perform_reduction(ReduceOperation reduce_op, std::vector<rocprofiler_record_coun
|
||||
}
|
||||
case REDUCE_SUM:
|
||||
{
|
||||
result = std::accumulate(input_array->begin(),
|
||||
input_array->end(),
|
||||
rocprofiler_record_counter_t{.id = 0, .counter_value = 0},
|
||||
[](auto& a, auto& b) {
|
||||
return rocprofiler_record_counter_t{
|
||||
.id = a.id,
|
||||
.counter_value = a.counter_value + b.counter_value};
|
||||
});
|
||||
result = std::accumulate(
|
||||
input_array->begin(),
|
||||
input_array->end(),
|
||||
rocprofiler_record_counter_t{.id = 0, .counter_value = 0, .dispatch_id = 0},
|
||||
[](auto& a, auto& b) {
|
||||
return rocprofiler_record_counter_t{
|
||||
.id = a.id,
|
||||
.counter_value = a.counter_value + b.counter_value,
|
||||
.dispatch_id = a.dispatch_id};
|
||||
});
|
||||
break;
|
||||
}
|
||||
case REDUCE_AVG:
|
||||
{
|
||||
result = std::accumulate(input_array->begin(),
|
||||
input_array->end(),
|
||||
rocprofiler_record_counter_t{.id = 0, .counter_value = 0},
|
||||
[](auto& a, auto& b) {
|
||||
return rocprofiler_record_counter_t{
|
||||
.id = a.id,
|
||||
.counter_value = a.counter_value + b.counter_value};
|
||||
});
|
||||
result = std::accumulate(
|
||||
input_array->begin(),
|
||||
input_array->end(),
|
||||
rocprofiler_record_counter_t{.id = 0, .counter_value = 0, .dispatch_id = 0},
|
||||
[](auto& a, auto& b) {
|
||||
return rocprofiler_record_counter_t{
|
||||
.id = a.id,
|
||||
.counter_value = a.counter_value + b.counter_value,
|
||||
.dispatch_id = a.dispatch_id};
|
||||
});
|
||||
result.counter_value /= input_array->size();
|
||||
break;
|
||||
}
|
||||
@@ -214,8 +218,9 @@ EvaluateAST::EvaluateAST(rocprofiler_counter_id_t out_id,
|
||||
if(_type == NodeType::NUMBER_NODE)
|
||||
{
|
||||
_raw_value = std::get<int64_t>(ast.value);
|
||||
_static_value.push_back(
|
||||
{.id = 0, .counter_value = static_cast<double>(std::get<int64_t>(ast.value))});
|
||||
_static_value.push_back({.id = 0,
|
||||
.counter_value = static_cast<double>(std::get<int64_t>(ast.value)),
|
||||
.dispatch_id = 0});
|
||||
}
|
||||
|
||||
for(const auto& nextAst : ast.counter_set)
|
||||
@@ -590,24 +595,30 @@ EvaluateAST::evaluate(
|
||||
case ADDITION_NODE:
|
||||
return perform_op([](auto& a, auto& b) {
|
||||
return rocprofiler_record_counter_t{
|
||||
.id = a.id, .counter_value = a.counter_value + b.counter_value};
|
||||
.id = a.id,
|
||||
.counter_value = a.counter_value + b.counter_value,
|
||||
.dispatch_id = a.dispatch_id};
|
||||
});
|
||||
case SUBTRACTION_NODE:
|
||||
return perform_op([](auto& a, auto& b) {
|
||||
return rocprofiler_record_counter_t{
|
||||
.id = a.id, .counter_value = a.counter_value - b.counter_value};
|
||||
.id = a.id,
|
||||
.counter_value = a.counter_value - b.counter_value,
|
||||
.dispatch_id = a.dispatch_id};
|
||||
});
|
||||
case MULTIPLY_NODE:
|
||||
return perform_op([](auto& a, auto& b) {
|
||||
return rocprofiler_record_counter_t{
|
||||
.id = a.id, .counter_value = a.counter_value * b.counter_value};
|
||||
.id = a.id,
|
||||
.counter_value = a.counter_value * b.counter_value,
|
||||
.dispatch_id = a.dispatch_id};
|
||||
});
|
||||
case DIVIDE_NODE:
|
||||
return perform_op([](auto& a, auto& b) {
|
||||
return rocprofiler_record_counter_t{
|
||||
.id = a.id,
|
||||
.counter_value =
|
||||
(b.counter_value == 0 ? 0 : a.counter_value / b.counter_value)};
|
||||
.id = a.id,
|
||||
.counter_value = (b.counter_value == 0 ? 0 : a.counter_value / b.counter_value),
|
||||
.dispatch_id = a.dispatch_id};
|
||||
});
|
||||
case REFERENCE_NODE:
|
||||
{
|
||||
|
||||
+28
-22
@@ -646,16 +646,17 @@ TEST(evaluate_ast, counter_reduction_sum)
|
||||
sum_vec(base_counter_data["KRUEGER"])),
|
||||
2},
|
||||
{"KRAMER",
|
||||
plus_vec(times_vec(std::vector<rocprofiler_record_counter_t>{{.id = 0,
|
||||
.counter_value = 5.0}},
|
||||
sum_vec(base_counter_data["VOORHEES"])),
|
||||
plus_vec(times_vec(
|
||||
std::vector<rocprofiler_record_counter_t>{
|
||||
{.id = 0, .counter_value = 5.0, .dispatch_id = 0}},
|
||||
sum_vec(base_counter_data["VOORHEES"])),
|
||||
sum_vec(base_counter_data["KRUEGER"])),
|
||||
2},
|
||||
{"GHOSTFACE",
|
||||
plus_vec(sum_vec(base_counter_data["VOORHEES"]),
|
||||
divide_vec(sum_vec(base_counter_data["KRUEGER"]),
|
||||
std::vector<rocprofiler_record_counter_t>{
|
||||
{.id = 0, .counter_value = 5.0}})),
|
||||
{.id = 0, .counter_value = 5.0, .dispatch_id = 0}})),
|
||||
2},
|
||||
};
|
||||
|
||||
@@ -726,16 +727,17 @@ TEST(evaluate_ast, counter_reduction_min)
|
||||
min_vec(base_counter_data["KRUEGER"])),
|
||||
2},
|
||||
{"KRAMER",
|
||||
plus_vec(times_vec(std::vector<rocprofiler_record_counter_t>{{.id = 0,
|
||||
.counter_value = 5.0}},
|
||||
min_vec(base_counter_data["VOORHEES"])),
|
||||
plus_vec(times_vec(
|
||||
std::vector<rocprofiler_record_counter_t>{
|
||||
{.id = 0, .counter_value = 5.0, .dispatch_id = 0}},
|
||||
min_vec(base_counter_data["VOORHEES"])),
|
||||
min_vec(base_counter_data["KRUEGER"])),
|
||||
2},
|
||||
{"GHOSTFACE",
|
||||
plus_vec(min_vec(base_counter_data["VOORHEES"]),
|
||||
divide_vec(min_vec(base_counter_data["KRUEGER"]),
|
||||
std::vector<rocprofiler_record_counter_t>{
|
||||
{.id = 0, .counter_value = 5.0}})),
|
||||
{.id = 0, .counter_value = 5.0, .dispatch_id = 0}})),
|
||||
2},
|
||||
};
|
||||
|
||||
@@ -806,16 +808,17 @@ TEST(evaluate_ast, counter_reduction_max)
|
||||
max_vec(base_counter_data["KRUEGER"])),
|
||||
2},
|
||||
{"KRAMER",
|
||||
plus_vec(times_vec(std::vector<rocprofiler_record_counter_t>{{.id = 0,
|
||||
.counter_value = 5.0}},
|
||||
max_vec(base_counter_data["VOORHEES"])),
|
||||
plus_vec(times_vec(
|
||||
std::vector<rocprofiler_record_counter_t>{
|
||||
{.id = 0, .counter_value = 5.0, .dispatch_id = 0}},
|
||||
max_vec(base_counter_data["VOORHEES"])),
|
||||
max_vec(base_counter_data["KRUEGER"])),
|
||||
2},
|
||||
{"GHOSTFACE",
|
||||
plus_vec(max_vec(base_counter_data["VOORHEES"]),
|
||||
divide_vec(max_vec(base_counter_data["KRUEGER"]),
|
||||
std::vector<rocprofiler_record_counter_t>{
|
||||
{.id = 0, .counter_value = 5.0}})),
|
||||
{.id = 0, .counter_value = 5.0, .dispatch_id = 0}})),
|
||||
2},
|
||||
};
|
||||
|
||||
@@ -888,16 +891,17 @@ TEST(evaluate_ast, counter_reduction_avg)
|
||||
avg_vec(base_counter_data["KRUEGER"])),
|
||||
2},
|
||||
{"KRAMER",
|
||||
plus_vec(times_vec(std::vector<rocprofiler_record_counter_t>{{.id = 0,
|
||||
.counter_value = 5.0}},
|
||||
avg_vec(base_counter_data["VOORHEES"])),
|
||||
plus_vec(times_vec(
|
||||
std::vector<rocprofiler_record_counter_t>{
|
||||
{.id = 0, .counter_value = 5.0, .dispatch_id = 0}},
|
||||
avg_vec(base_counter_data["VOORHEES"])),
|
||||
avg_vec(base_counter_data["KRUEGER"])),
|
||||
2},
|
||||
{"GHOSTFACE",
|
||||
plus_vec(avg_vec(base_counter_data["VOORHEES"]),
|
||||
divide_vec(avg_vec(base_counter_data["KRUEGER"]),
|
||||
std::vector<rocprofiler_record_counter_t>{
|
||||
{.id = 0, .counter_value = 5.0}})),
|
||||
{.id = 0, .counter_value = 5.0, .dispatch_id = 0}})),
|
||||
2},
|
||||
};
|
||||
|
||||
@@ -956,18 +960,20 @@ TEST(evaluate_ast, evaluate_mixed_counters)
|
||||
std::vector<std::tuple<std::string, std::vector<rocprofiler_record_counter_t>, int64_t>>
|
||||
derived_counters = {
|
||||
{"BATES",
|
||||
times_vec(std::vector<rocprofiler_record_counter_t>{{.id = 0, .counter_value = 32}},
|
||||
sum_vec(base_counter_data["VOORHEES"])),
|
||||
times_vec(
|
||||
std::vector<rocprofiler_record_counter_t>{
|
||||
{.id = 0, .counter_value = 32, .dispatch_id = 0}},
|
||||
sum_vec(base_counter_data["VOORHEES"])),
|
||||
2},
|
||||
{"KRAMER",
|
||||
times_vec(
|
||||
sum_vec(base_counter_data["KRUEGER"]),
|
||||
std::vector<rocprofiler_record_counter_t>{{.id = 0, .counter_value = 8.0 / 5.0}}),
|
||||
times_vec(sum_vec(base_counter_data["KRUEGER"]),
|
||||
std::vector<rocprofiler_record_counter_t>{
|
||||
{.id = 0, .counter_value = 8.0 / 5.0, .dispatch_id = 0}}),
|
||||
3},
|
||||
{"TORRANCE",
|
||||
times_vec(sum_vec(base_counter_data["KRUEGER"]),
|
||||
std::vector<rocprofiler_record_counter_t>{
|
||||
{.id = 0, .counter_value = 104.0 / (156.0 * 8.0)}}),
|
||||
{.id = 0, .counter_value = 104.0 / (156.0 * 8.0), .dispatch_id = 0}}),
|
||||
4},
|
||||
};
|
||||
|
||||
|
||||
@@ -301,12 +301,13 @@ WriteInterceptor(const void* packets,
|
||||
.dispatch_id = dispatch_id,
|
||||
.private_segment_size = kernel_pkt.kernel_dispatch.private_segment_size,
|
||||
.group_segment_size = kernel_pkt.kernel_dispatch.group_segment_size,
|
||||
.workgroup_size = rocprofiler_dim3_t{kernel_pkt.kernel_dispatch.workgroup_size_x,
|
||||
.workgroup_size = rocprofiler_dim3_t{kernel_pkt.kernel_dispatch.workgroup_size_x,
|
||||
kernel_pkt.kernel_dispatch.workgroup_size_y,
|
||||
kernel_pkt.kernel_dispatch.workgroup_size_z},
|
||||
.grid_size = rocprofiler_dim3_t{kernel_pkt.kernel_dispatch.grid_size_x,
|
||||
.grid_size = rocprofiler_dim3_t{kernel_pkt.kernel_dispatch.grid_size_x,
|
||||
kernel_pkt.kernel_dispatch.grid_size_y,
|
||||
kernel_pkt.kernel_dispatch.grid_size_z}}};
|
||||
kernel_pkt.kernel_dispatch.grid_size_z},
|
||||
.reserved_padding = {0}}};
|
||||
|
||||
{
|
||||
auto tracer_data = callback_record;
|
||||
|
||||
@@ -84,7 +84,64 @@ destroy_queue(hsa_queue_t* hsa_queue)
|
||||
|
||||
constexpr rocprofiler_agent_t default_agent =
|
||||
rocprofiler_agent_t{.size = sizeof(rocprofiler_agent_t),
|
||||
.id = rocprofiler_agent_id_t{std::numeric_limits<uint64_t>::max()}};
|
||||
.id = rocprofiler_agent_id_t{std::numeric_limits<uint64_t>::max()},
|
||||
.type = ROCPROFILER_AGENT_TYPE_NONE,
|
||||
.cpu_cores_count = 0,
|
||||
.simd_count = 0,
|
||||
.mem_banks_count = 0,
|
||||
.caches_count = 0,
|
||||
.io_links_count = 0,
|
||||
.cpu_core_id_base = 0,
|
||||
.simd_id_base = 0,
|
||||
.max_waves_per_simd = 0,
|
||||
.lds_size_in_kb = 0,
|
||||
.gds_size_in_kb = 0,
|
||||
.num_gws = 0,
|
||||
.wave_front_size = 0,
|
||||
.num_xcc = 0,
|
||||
.cu_count = 0,
|
||||
.array_count = 0,
|
||||
.num_shader_banks = 0,
|
||||
.simd_arrays_per_engine = 0,
|
||||
.cu_per_simd_array = 0,
|
||||
.simd_per_cu = 0,
|
||||
.max_slots_scratch_cu = 0,
|
||||
.gfx_target_version = 0,
|
||||
.vendor_id = 0,
|
||||
.device_id = 0,
|
||||
.location_id = 0,
|
||||
.domain = 0,
|
||||
.drm_render_minor = 0,
|
||||
.num_sdma_engines = 0,
|
||||
.num_sdma_xgmi_engines = 0,
|
||||
.num_sdma_queues_per_engine = 0,
|
||||
.num_cp_queues = 0,
|
||||
.max_engine_clk_ccompute = 0,
|
||||
.max_engine_clk_fcompute = 0,
|
||||
.sdma_fw_version = {},
|
||||
.fw_version = {},
|
||||
.capability = {},
|
||||
.cu_per_engine = 0,
|
||||
.max_waves_per_cu = 0,
|
||||
.family_id = 0,
|
||||
.workgroup_max_size = 0,
|
||||
.grid_max_size = 0,
|
||||
.local_mem_size = 0,
|
||||
.hive_id = 0,
|
||||
.gpu_id = 0,
|
||||
.workgroup_max_dim = {0, 0, 0},
|
||||
.grid_max_dim = {0, 0, 0},
|
||||
.mem_banks = nullptr,
|
||||
.caches = nullptr,
|
||||
.io_links = nullptr,
|
||||
.name = nullptr,
|
||||
.vendor_name = nullptr,
|
||||
.product_name = nullptr,
|
||||
.model_name = nullptr,
|
||||
.num_pc_sampling_configs = 0,
|
||||
.pc_sampling_configs = nullptr,
|
||||
.node_id = 0,
|
||||
.logical_node_id = 0};
|
||||
} // namespace
|
||||
|
||||
void
|
||||
|
||||
+2
-3
@@ -848,9 +848,8 @@ struct poll_kfd_t
|
||||
|
||||
// Create fd for notifying thread when we want to wake it up, and an eventfd for any events
|
||||
// to this thread
|
||||
file_handles.emplace_back(pollfd{
|
||||
.fd = eventfd(0, DEFAULT_FLAGS),
|
||||
});
|
||||
file_handles.emplace_back(
|
||||
pollfd{.fd = eventfd(0, DEFAULT_FLAGS), .events = 0, .revents = 0});
|
||||
|
||||
fd_t thread_pipes[2]{};
|
||||
|
||||
|
||||
Reference in New Issue
Block a user