From 7892cc861c573d6afa4be65b6cca1b1146663f73 Mon Sep 17 00:00:00 2001 From: Evgeny Date: Thu, 22 Jun 2017 18:41:05 -0500 Subject: [PATCH] Block list extending Change-Id: Id17efde25fce287296e80f2b37c77b15aa59b561 [ROCm/ROCR-Runtime commit: c533229bc170003533774b64a905f42d3b31843c] --- .../inc/hsa_ext_amd_aql_profile.h | 5 + .../src/aqlprofile/aql_profile.cpp | 412 ++++++++++-------- .../src/aqlprofile/aql_profile.h | 1 - .../src/aqlprofile/aql_profile_exception.h | 24 + .../src/aqlprofile/gfx8_factory.cpp | 9 +- .../src/aqlprofile/gfx9_factory.cpp | 44 +- .../src/aqlprofile/logger.h | 98 +++++ .../src/aqlprofile/pm4_factory.h | 22 +- .../src/perfcounter/gfx9_block_info.cpp | 2 - .../src/perfcounter/gfx9_block_info.h | 8 +- .../hsa-ext-aql-profile/test/ctrl/test.cpp | 10 +- .../hsa-ext-aql-profile/test/ctrl/test_aql.h | 15 +- .../test/ctrl/test_assert.h | 13 + .../test/ctrl/test_hsa.cpp | 8 +- .../test/ctrl/test_pgen_pmc.h | 9 +- .../test/ctrl/test_pgen_sqtt.h | 11 +- .../test/ctrl/test_pmgr.cpp | 4 +- 17 files changed, 453 insertions(+), 242 deletions(-) create mode 100644 projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile_exception.h create mode 100644 projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/logger.h create mode 100644 projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_assert.h diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/inc/hsa_ext_amd_aql_profile.h b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/inc/hsa_ext_amd_aql_profile.h index 7d69afdc87..abeb2d88db 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/inc/hsa_ext_amd_aql_profile.h +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/inc/hsa_ext_amd_aql_profile.h @@ -91,9 +91,11 @@ typedef enum { HSA_EXT_AQL_PROFILE_BLOCK_PASC, HSA_EXT_AQL_PROFILE_BLOCK_SPI, HSA_EXT_AQL_PROFILE_BLOCK_SQ, + HSA_EXT_AQL_PROFILE_BLOCK_SQES, HSA_EXT_AQL_PROFILE_BLOCK_SQGS, HSA_EXT_AQL_PROFILE_BLOCK_SQVS, HSA_EXT_AQL_PROFILE_BLOCK_SQPS, + HSA_EXT_AQL_PROFILE_BLOCK_SQLS, HSA_EXT_AQL_PROFILE_BLOCK_SQHS, HSA_EXT_AQL_PROFILE_BLOCK_SQCS, HSA_EXT_AQL_PROFILE_BLOCK_SX, @@ -106,8 +108,11 @@ typedef enum { HSA_EXT_AQL_PROFILE_BLOCK_VGT, HSA_EXT_AQL_PROFILE_BLOCK_IA, HSA_EXT_AQL_PROFILE_BLOCK_MC, + HSA_EXT_AQL_PROFILE_BLOCK_SRBM, HSA_EXT_AQL_PROFILE_BLOCK_TCS, HSA_EXT_AQL_PROFILE_BLOCK_WD, + HSA_EXT_AQL_PROFILE_BLOCK_CPG, + HSA_EXT_AQL_PROFILE_BLOCK_CPC, HSA_EXT_AQL_PROFILE_BLOCKS_NUMBER } hsa_ext_amd_aql_profile_block_name_t; diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.cpp b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.cpp index 62671139fc..26eadca6b5 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.cpp +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.cpp @@ -6,6 +6,7 @@ #include "perf_counter.h" // perfcounter #include "thread_trace.h" // threadtrace #include "gpu_block_info.h" +#include "logger.h" #define PUBLIC_API __attribute__((visibility("default"))) @@ -143,92 +144,101 @@ PUBLIC_API hsa_status_t hsa_ext_amd_aql_profile_validate_event( // Method to populate the provided AQL packet with profiling start commands PUBLIC_API hsa_status_t hsa_ext_amd_aql_profile_start( const hsa_ext_amd_aql_profile_profile_t* profile, aql_profile::packet_t* aql_start_packet) { - aql_profile::Pm4Factory* pm4_factory = aql_profile::Pm4Factory::Create(profile); - if (pm4_factory == NULL) return HSA_STATUS_ERROR; + aql_profile::Logger logger; - pm4_profile::CommandWriter* cmdWriter = pm4_factory->getCommandWriter(); - if (cmdWriter == NULL) return HSA_STATUS_ERROR; + try { + aql_profile::Pm4Factory* pm4_factory = aql_profile::Pm4Factory::Create(profile); + if (pm4_factory == NULL) return HSA_STATUS_ERROR; - pm4_profile::DefaultCmdBuf commands; - aql_profile::CommandBufferMgr cmdBufMgr(profile); - if (cmdBufMgr.getSize() == 0) return HSA_STATUS_ERROR; + pm4_profile::CommandWriter* cmdWriter = pm4_factory->getCommandWriter(); + if (cmdWriter == NULL) return HSA_STATUS_ERROR; - if (profile->type == HSA_EXT_AQL_PROFILE_EVENT_PMC) { - pm4_profile::Pmu* pmcMgr = pm4_factory->getPmcMgr(); - if (pmcMgr == NULL) return HSA_STATUS_ERROR; + pm4_profile::DefaultCmdBuf commands; + aql_profile::CommandBufferMgr cmdBufMgr(profile); + if (cmdBufMgr.getSize() == 0) return HSA_STATUS_ERROR; - pm4_profile::CountersMap countersMap; - for (const hsa_ext_amd_aql_profile_event_t* p = profile->events; - p < profile->events + profile->event_count; ++p) { - countersMap[pm4_factory->getBlockId(p)].push_back(p->counter_id); - } + if (profile->type == HSA_EXT_AQL_PROFILE_EVENT_PMC) { + pm4_profile::Pmu* pmcMgr = pm4_factory->getPmcMgr(); + if (pmcMgr == NULL) return HSA_STATUS_ERROR; - // Generate start commands - pmcMgr->begin(&commands, cmdWriter, countersMap); - cmdBufMgr.setPreSize(commands.Size()); + pm4_profile::CountersMap countersMap; + for (const hsa_ext_amd_aql_profile_event_t* p = profile->events; + p < profile->events + profile->event_count; ++p) { + countersMap[pm4_factory->getBlockId(p)].push_back(p->counter_id); + } - // Generate stop commands - const uint32_t data_size = - pmcMgr->end(&commands, cmdWriter, countersMap, profile->output_buffer.ptr); - if (data_size == 0) return HSA_STATUS_ERROR; - assert(data_size <= profile->output_buffer.size); - if (data_size > profile->output_buffer.size) return HSA_STATUS_ERROR; - } else if (profile->type == HSA_EXT_AQL_PROFILE_EVENT_SQTT) { - pm4_profile::ThreadTrace* sqttMgr = pm4_factory->getSqttMgr(); - if (sqttMgr == NULL) return HSA_STATUS_ERROR; + // Generate start commands + pmcMgr->begin(&commands, cmdWriter, countersMap); + cmdBufMgr.setPreSize(commands.Size()); - pm4_profile::ThreadTraceConfig sqtt_config; - sqttMgr->InitThreadTraceConfig(&sqtt_config); - if (profile->parameters) { - for (const hsa_ext_amd_aql_profile_parameters_t* p = profile->parameters; - p < (profile->parameters + profile->parameter_count); ++p) { - switch (p->parameter_name) { - case HSA_EXT_AQL_PROFILE_PARAM_COMPUTE_UNIT_TARGET: - sqtt_config.threadTraceTargetCu = p->value; - break; - case HSA_EXT_AQL_PROFILE_PARAM_VM_ID_MASK: - sqtt_config.threadTraceVmIdMask = p->value; - break; - case HSA_EXT_AQL_PROFILE_PARAM_MASK: - sqtt_config.threadTraceMask = p->value; - break; - case HSA_EXT_AQL_PROFILE_PARAM_TOKEN_MASK: - sqtt_config.threadTraceTokenMask = p->value; - break; - case HSA_EXT_AQL_PROFILE_PARAM_TOKEN_MASK2: - sqtt_config.threadTraceTokenMask2 = p->value; - break; - default: - return HSA_STATUS_ERROR; + // Generate stop commands + const uint32_t data_size = + pmcMgr->end(&commands, cmdWriter, countersMap, profile->output_buffer.ptr); + if (data_size == 0) return HSA_STATUS_ERROR; + assert(data_size <= profile->output_buffer.size); + if (data_size > profile->output_buffer.size) return HSA_STATUS_ERROR; + + } else if (profile->type == HSA_EXT_AQL_PROFILE_EVENT_SQTT) { + pm4_profile::ThreadTrace* sqttMgr = pm4_factory->getSqttMgr(); + if (sqttMgr == NULL) return HSA_STATUS_ERROR; + + pm4_profile::ThreadTraceConfig sqtt_config; + sqttMgr->InitThreadTraceConfig(&sqtt_config); + if (profile->parameters) { + for (const hsa_ext_amd_aql_profile_parameters_t* p = profile->parameters; + p < (profile->parameters + profile->parameter_count); ++p) { + switch (p->parameter_name) { + case HSA_EXT_AQL_PROFILE_PARAM_COMPUTE_UNIT_TARGET: + sqtt_config.threadTraceTargetCu = p->value; + break; + case HSA_EXT_AQL_PROFILE_PARAM_VM_ID_MASK: + sqtt_config.threadTraceVmIdMask = p->value; + break; + case HSA_EXT_AQL_PROFILE_PARAM_MASK: + sqtt_config.threadTraceMask = p->value; + break; + case HSA_EXT_AQL_PROFILE_PARAM_TOKEN_MASK: + sqtt_config.threadTraceTokenMask = p->value; + break; + case HSA_EXT_AQL_PROFILE_PARAM_TOKEN_MASK2: + sqtt_config.threadTraceTokenMask2 = p->value; + break; + default: + ERR_LOGGING(logger) << "Bad SQTT parameter name (" << p->parameter_name << ")"; + return HSA_STATUS_ERROR; + } } } - } - sqttMgr->Init(&sqtt_config); + sqttMgr->Init(&sqtt_config); - sqttMgr->setSqttDataBuff((uint8_t*)profile->output_buffer.ptr, profile->output_buffer.size); + sqttMgr->setSqttDataBuff((uint8_t*)profile->output_buffer.ptr, profile->output_buffer.size); - const uint32_t status_size = sqttMgr->StatusSizeInfo(); - void* status_ptr = cmdBufMgr.setPostfix(status_size); - if (status_ptr == NULL) return HSA_STATUS_ERROR; - // Control buffer registering - sqttMgr->setSqttCtrlBuff((uint32_t*)status_ptr); + const uint32_t status_size = sqttMgr->StatusSizeInfo(); + void* status_ptr = cmdBufMgr.setPostfix(status_size); + if (status_ptr == NULL) return HSA_STATUS_ERROR; + // Control buffer registering + sqttMgr->setSqttCtrlBuff((uint32_t*)status_ptr); - // Generate start commands - sqttMgr->BeginSession(&commands, cmdWriter); - cmdBufMgr.setPreSize(commands.Size()); - // Generate stop commands - sqttMgr->StopSession(&commands, cmdWriter); - } else + // Generate start commands + sqttMgr->BeginSession(&commands, cmdWriter); + cmdBufMgr.setPreSize(commands.Size()); + // Generate stop commands + sqttMgr->StopSession(&commands, cmdWriter); + } else + return HSA_STATUS_ERROR; + + if (!cmdBufMgr.checkTotalSize(commands.Size())) return HSA_STATUS_ERROR; + + const aql_profile::descriptor_t pre_descr = cmdBufMgr.getPreDescr(); + const aql_profile::descriptor_t post_descr = cmdBufMgr.getPostDescr(); + memcpy(pre_descr.ptr, commands.Base(), pre_descr.size); + memcpy(post_descr.ptr, commands.Base() + pre_descr.size, post_descr.size); + // Populate start aql packet + aql_profile::populateAql(pre_descr.ptr, pre_descr.size, cmdWriter, aql_start_packet); + } catch (std::exception& e) { + ERR_LOGGING(logger) << e.what(); return HSA_STATUS_ERROR; - - if (!cmdBufMgr.checkTotalSize(commands.Size())) return HSA_STATUS_ERROR; - - const aql_profile::descriptor_t pre_descr = cmdBufMgr.getPreDescr(); - const aql_profile::descriptor_t post_descr = cmdBufMgr.getPostDescr(); - memcpy(pre_descr.ptr, commands.Base(), pre_descr.size); - memcpy(post_descr.ptr, commands.Base() + pre_descr.size, post_descr.size); - // Populate start aql packet - aql_profile::populateAql(pre_descr.ptr, pre_descr.size, cmdWriter, aql_start_packet); + } return HSA_STATUS_SUCCESS; } @@ -236,18 +246,25 @@ PUBLIC_API hsa_status_t hsa_ext_amd_aql_profile_start( // Method to populate the provided AQL packet with profiling stop commands PUBLIC_API hsa_status_t hsa_ext_amd_aql_profile_stop( const hsa_ext_amd_aql_profile_profile_t* profile, aql_profile::packet_t* aql_stop_packet) { - aql_profile::Pm4Factory* pm4_factory = aql_profile::Pm4Factory::Create(profile); - if (pm4_factory == NULL) return HSA_STATUS_ERROR; + aql_profile::Logger logger; - pm4_profile::CommandWriter* cmdWriter = pm4_factory->getCommandWriter(); - if (cmdWriter == NULL) return HSA_STATUS_ERROR; + try { + aql_profile::Pm4Factory* pm4_factory = aql_profile::Pm4Factory::Create(profile); + if (pm4_factory == NULL) return HSA_STATUS_ERROR; - aql_profile::CommandBufferMgr cmdBufMgr(profile); - if (cmdBufMgr.getSize() == 0) return HSA_STATUS_ERROR; + pm4_profile::CommandWriter* cmdWriter = pm4_factory->getCommandWriter(); + if (cmdWriter == NULL) return HSA_STATUS_ERROR; - const aql_profile::descriptor_t post_descr = cmdBufMgr.getPostDescr(); - // Populate stop aql packet - aql_profile::populateAql(post_descr.ptr, post_descr.size, cmdWriter, aql_stop_packet); + aql_profile::CommandBufferMgr cmdBufMgr(profile); + if (cmdBufMgr.getSize() == 0) return HSA_STATUS_ERROR; + + const aql_profile::descriptor_t post_descr = cmdBufMgr.getPostDescr(); + // Populate stop aql packet + aql_profile::populateAql(post_descr.ptr, post_descr.size, cmdWriter, aql_stop_packet); + } catch (std::exception& e) { + ERR_LOGGING(logger) << e.what(); + return HSA_STATUS_ERROR; + } return HSA_STATUS_SUCCESS; } @@ -255,13 +272,21 @@ PUBLIC_API hsa_status_t hsa_ext_amd_aql_profile_stop( // Legacy devices, converting of the profiling AQL packet to PM4 packet blob PUBLIC_API hsa_status_t hsa_ext_amd_aql_profile_legacy_get_pm4(const aql_profile::packet_t* aql_packet, void* data) { - // Populate GFX8 pm4 packet blob - // Adding HSA barrier acquire packet - data = aql_profile::legacyAqlAcquire(aql_packet, data); - // Adding PM4 command packet - data = aql_profile::legacyPm4(aql_packet, data); - // Adding HSA barrier release packet - data = aql_profile::legacyAqlRelease(aql_packet, data); + aql_profile::Logger logger; + + try { + // Populate GFX8 pm4 packet blob + // Adding HSA barrier acquire packet + data = aql_profile::legacyAqlAcquire(aql_packet, data); + // Adding PM4 command packet + data = aql_profile::legacyPm4(aql_packet, data); + // Adding HSA barrier release packet + data = aql_profile::legacyAqlRelease(aql_packet, data); + } catch (std::exception& e) { + ERR_LOGGING(logger) << e.what(); + return HSA_STATUS_ERROR; + } + return HSA_STATUS_SUCCESS; } @@ -270,25 +295,32 @@ PUBLIC_API hsa_status_t hsa_ext_amd_aql_profile_get_info(const hsa_ext_amd_aql_profile_profile_t* profile, hsa_ext_amd_aql_profile_info_type_t attribute, void* value) { hsa_status_t status = HSA_STATUS_SUCCESS; + aql_profile::Logger logger; - switch (attribute) { - case HSA_EXT_AQL_PROFILE_INFO_COMMAND_BUFFER_SIZE: - *(uint32_t*)value = 0x1000; // a current approximation as 4K is big enaugh - break; - case HSA_EXT_AQL_PROFILE_INFO_PMC_DATA_SIZE: - *(uint32_t*)value = 0x1000; // a current approximation as 4K is big enaugh - break; - case HSA_EXT_AQL_PROFILE_INFO_PMC_DATA: - reinterpret_cast(value)->pmc_data.result = 0; - status = hsa_ext_amd_aql_profile_iterate_data(profile, aql_profile::default_pmcdata_callback, - value); - break; - case HSA_EXT_AQL_PROFILE_INFO_SQTT_DATA: - status = hsa_ext_amd_aql_profile_iterate_data(profile, aql_profile::default_sqttdata_callback, - value); - break; - default: - status = HSA_STATUS_ERROR_INVALID_ARGUMENT; + try { + switch (attribute) { + case HSA_EXT_AQL_PROFILE_INFO_COMMAND_BUFFER_SIZE: + *(uint32_t*)value = 0x1000; // a current approximation as 4K is big enaugh + break; + case HSA_EXT_AQL_PROFILE_INFO_PMC_DATA_SIZE: + *(uint32_t*)value = 0x1000; // a current approximation as 4K is big enaugh + break; + case HSA_EXT_AQL_PROFILE_INFO_PMC_DATA: + reinterpret_cast(value)->pmc_data.result = 0; + status = hsa_ext_amd_aql_profile_iterate_data(profile, + aql_profile::default_pmcdata_callback, value); + break; + case HSA_EXT_AQL_PROFILE_INFO_SQTT_DATA: + status = hsa_ext_amd_aql_profile_iterate_data( + profile, aql_profile::default_sqttdata_callback, value); + break; + default: + status = HSA_STATUS_ERROR_INVALID_ARGUMENT; + ERR_LOGGING(logger) << "Invalid attribute (" << attribute << ")"; + } + } catch (std::exception& e) { + ERR_LOGGING(logger) << e.what(); + return HSA_STATUS_ERROR; } return status; @@ -299,93 +331,109 @@ PUBLIC_API hsa_status_t hsa_ext_amd_aql_profile_iterate_data(const hsa_ext_amd_aql_profile_profile_t* profile, hsa_ext_amd_aql_profile_data_callback_t callback, void* data) { hsa_status_t status = HSA_STATUS_SUCCESS; - aql_profile::Pm4Factory* pm4_factory = aql_profile::Pm4Factory::Create(profile); - if (pm4_factory == NULL) return HSA_STATUS_ERROR; + aql_profile::Logger logger; - if (profile->type == HSA_EXT_AQL_PROFILE_EVENT_PMC) { - uint32_t info_size = 0; - void* info_data; - uint64_t* samples = (uint64_t*)profile->output_buffer.ptr; - const uint32_t sample_count = profile->output_buffer.size / sizeof(uint64_t); - uint32_t sample_index = 0; + try { + aql_profile::Pm4Factory* pm4_factory = aql_profile::Pm4Factory::Create(profile); + if (pm4_factory == NULL) return HSA_STATUS_ERROR; - pm4_profile::Pmu* pmcMgr = pm4_factory->getPmcMgr(); - if (pmcMgr == NULL) return HSA_STATUS_ERROR; + if (profile->type == HSA_EXT_AQL_PROFILE_EVENT_PMC) { + uint32_t info_size = 0; + void* info_data; + uint64_t* samples = (uint64_t*)profile->output_buffer.ptr; + const uint32_t sample_count = profile->output_buffer.size / sizeof(uint64_t); + uint32_t sample_index = 0; - for (const hsa_ext_amd_aql_profile_event_t* p = profile->events; - p < (profile->events + profile->event_count); ++p) { - const pm4_profile::GpuBlockInfo* block_info = pm4_factory->getBlockInfo(p); - if (block_info == NULL) return HSA_STATUS_ERROR; - const pm4_profile::CntlMethod method = pm4_factory->getBlockInfo(p)->method; - // A perfcounter data sample per ShaderEngine - const uint32_t block_samples_count = (method == pm4_profile::CntlMethodBySe || - method == pm4_profile::CntlMethodBySeAndInstance) - ? pmcMgr->getNumSe() - : 1; - for (uint32_t i = 0; i < block_samples_count; ++i) { - assert(sample_index < sample_count); - if (sample_index >= sample_count) return HSA_STATUS_ERROR; + pm4_profile::Pmu* pmcMgr = pm4_factory->getPmcMgr(); + if (pmcMgr == NULL) return HSA_STATUS_ERROR; + + for (const hsa_ext_amd_aql_profile_event_t* p = profile->events; + p < (profile->events + profile->event_count); ++p) { + const pm4_profile::GpuBlockInfo* block_info = pm4_factory->getBlockInfo(p); + if (block_info == NULL) return HSA_STATUS_ERROR; + const pm4_profile::CntlMethod method = pm4_factory->getBlockInfo(p)->method; + // A perfcounter data sample per ShaderEngine + const uint32_t block_samples_count = (method == pm4_profile::CntlMethodBySe || + method == pm4_profile::CntlMethodBySeAndInstance) + ? pmcMgr->getNumSe() + : 1; + for (uint32_t i = 0; i < block_samples_count; ++i) { + assert(sample_index < sample_count); + if (sample_index >= sample_count) return HSA_STATUS_ERROR; + + hsa_ext_amd_aql_profile_info_data_t sample_info; + sample_info.sample_id = i; + sample_info.pmc_data.event = *p; + sample_info.pmc_data.result = samples[sample_index]; + status = callback(HSA_EXT_AQL_PROFILE_INFO_PMC_DATA, &sample_info, data); + if (status == HSA_STATUS_INFO_BREAK) { + status = HSA_STATUS_SUCCESS; + break; + } + if (status != HSA_STATUS_SUCCESS) { + ERR_LOGGING(logger) << "PMC data callback error, sample_id(" << i << ") status(" + << status << ")"; + break; + } + ++sample_index; + } + } + } else if (profile->type == HSA_EXT_AQL_PROFILE_EVENT_SQTT) { + pm4_profile::ThreadTrace* sqttMgr = pm4_factory->getSqttMgr(); + if (sqttMgr == NULL) return HSA_STATUS_ERROR; + + aql_profile::CommandBufferMgr cmdBufMgr(profile); + if (cmdBufMgr.getSize() == 0) return HSA_STATUS_ERROR; + + const uint32_t status_size = sqttMgr->StatusSizeInfo(); + // Control buffer was allocated as the CmdBuffer postfix partition + void* status_ptr = cmdBufMgr.setPostfix(status_size); + if (status_ptr == NULL) return HSA_STATUS_ERROR; + // Control buffer registering + sqttMgr->setSqttCtrlBuff((uint32_t*)status_ptr); + // Validate SQTT status and normalize WRPTR + if (sqttMgr->Validate() == false) return HSA_STATUS_ERROR; + + const uint32_t se_number = sqttMgr->getNumSe(); + // Casting status pointer to SQTT control per ShaderEngine array + aql_profile::sqtt_ctrl_t* sqtt_ctrl = (aql_profile::sqtt_ctrl_t*)status_ptr; + assert(status_size == sizeof(aql_profile::sqtt_ctrl_t) * se_number); + if (status_size != sizeof(aql_profile::sqtt_ctrl_t) * se_number) { + return HSA_STATUS_ERROR; + } + // SQTT output buffer and capacity per ShaderEngine + void* sample_ptr = profile->output_buffer.ptr; + const uint32_t sample_capacity = profile->output_buffer.size / se_number; + // The samples sizes are returned in the control buffer + for (int i = 0; i < se_number; ++i) { + // WPTR specifies the index in thread trace buffer where next token will be + // written by hardware. The index is incremented by size of 32 bytes. + uint32_t sample_size = sqtt_ctrl[i].writePtr * TT_WRITE_PTR_BLK; hsa_ext_amd_aql_profile_info_data_t sample_info; sample_info.sample_id = i; - sample_info.pmc_data.event = *p; - sample_info.pmc_data.result = samples[sample_index]; - status = callback(HSA_EXT_AQL_PROFILE_INFO_PMC_DATA, &sample_info, data); + sample_info.sqtt_data.ptr = sample_ptr; + sample_info.sqtt_data.size = sample_size; + status = callback(HSA_EXT_AQL_PROFILE_INFO_SQTT_DATA, &sample_info, data); if (status == HSA_STATUS_INFO_BREAK) { status = HSA_STATUS_SUCCESS; break; } - if (status != HSA_STATUS_SUCCESS) break; - ++sample_index; + if (status != HSA_STATUS_SUCCESS) { + ERR_LOGGING(logger) << "SQTT data callback error, sample_id(" << i << ") status(" + << status << ")"; + break; + } + + sample_ptr += sample_capacity; } + } else { + ERR_LOGGING(logger) << "Bad profile type (" << profile->type << ")"; + status = HSA_STATUS_ERROR; } - } else if (profile->type == HSA_EXT_AQL_PROFILE_EVENT_SQTT) { - pm4_profile::ThreadTrace* sqttMgr = pm4_factory->getSqttMgr(); - if (sqttMgr == NULL) return HSA_STATUS_ERROR; - - aql_profile::CommandBufferMgr cmdBufMgr(profile); - if (cmdBufMgr.getSize() == 0) return HSA_STATUS_ERROR; - - const uint32_t status_size = sqttMgr->StatusSizeInfo(); - // Control buffer was allocated as the CmdBuffer postfix partition - void* status_ptr = cmdBufMgr.setPostfix(status_size); - if (status_ptr == NULL) return HSA_STATUS_ERROR; - // Control buffer registering - sqttMgr->setSqttCtrlBuff((uint32_t*)status_ptr); - // Validate SQTT status and normalize WRPTR - if (sqttMgr->Validate() == false) return HSA_STATUS_ERROR; - - const uint32_t se_number = sqttMgr->getNumSe(); - // Casting status pointer to SQTT control per ShaderEngine array - aql_profile::sqtt_ctrl_t* sqtt_ctrl = (aql_profile::sqtt_ctrl_t*)status_ptr; - assert(status_size == sizeof(aql_profile::sqtt_ctrl_t) * se_number); - if (status_size != sizeof(aql_profile::sqtt_ctrl_t) * se_number) { - return HSA_STATUS_ERROR; - } - // SQTT output buffer and capacity per ShaderEngine - void* sample_ptr = profile->output_buffer.ptr; - const uint32_t sample_capacity = profile->output_buffer.size / se_number; - // The samples sizes are returned in the control buffer - for (int i = 0; i < se_number; ++i) { - // WPTR specifies the index in thread trace buffer where next token will be - // written by hardware. The index is incremented by size of 32 bytes. - uint32_t sample_size = sqtt_ctrl[i].writePtr * TT_WRITE_PTR_BLK; - - hsa_ext_amd_aql_profile_info_data_t sample_info; - sample_info.sample_id = i; - sample_info.sqtt_data.ptr = sample_ptr; - sample_info.sqtt_data.size = sample_size; - status = callback(HSA_EXT_AQL_PROFILE_INFO_SQTT_DATA, &sample_info, data); - if (status == HSA_STATUS_INFO_BREAK) { - status = HSA_STATUS_SUCCESS; - break; - } - if (status != HSA_STATUS_SUCCESS) break; - - sample_ptr += sample_capacity; - } - } else { - status = HSA_STATUS_ERROR; + } catch (std::exception& e) { + ERR_LOGGING(logger) << e.what(); + return HSA_STATUS_ERROR; } return status; diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.h b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.h index c8b12622d2..2764a0e702 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.h +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.h @@ -8,7 +8,6 @@ class CommandWriter; } namespace aql_profile { - typedef hsa_ext_amd_aql_profile_descriptor_t descriptor_t; typedef hsa_ext_amd_aql_profile_profile_t profile_t; typedef hsa_ext_amd_aql_profile_info_type_t info_type_t; diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile_exception.h b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile_exception.h new file mode 100644 index 0000000000..bfba9dd0df --- /dev/null +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile_exception.h @@ -0,0 +1,24 @@ +#ifndef _AQL_PROFILE_EXCEPTION_H_ +#define _AQL_PROFILE_EXCEPTION_H_ + +#include +#include + +namespace aql_profile { + +template class aql_profile_exception : public std::exception { + public: + aql_profile_exception(const std::string& m, const T& v) : msg(m), val(v) {} + virtual const char* what() const throw() { + std::ostringstream oss; + oss << msg << "(" << val << ")"; + return strdup(oss.str().c_str()); + } + + private: + std::string msg; + T val; +}; +} + +#endif // _AQL_PROFILE_EXCEPTION_H_ diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/gfx8_factory.cpp b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/gfx8_factory.cpp index e9a30fb660..a0ecbc4107 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/gfx8_factory.cpp +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/gfx8_factory.cpp @@ -16,15 +16,18 @@ uint32_t Gfx8Factory::block_id_table[HSA_EXT_AQL_PROFILE_BLOCKS_NUMBER] = { pm4_profile::kHsaViCounterBlockIdDb0, pm4_profile::kHsaViCounterBlockIdGrbm, pm4_profile::kHsaViCounterBlockIdGrbmSe, pm4_profile::kHsaViCounterBlockIdPaSu, pm4_profile::kHsaViCounterBlockIdPaSc, pm4_profile::kHsaViCounterBlockIdSpi, - pm4_profile::kHsaViCounterBlockIdSq, pm4_profile::kHsaViCounterBlockIdSqGs, - pm4_profile::kHsaViCounterBlockIdSqVs, pm4_profile::kHsaViCounterBlockIdSqPs, + pm4_profile::kHsaViCounterBlockIdSq, pm4_profile::kHsaViCounterBlockIdSqEs, + pm4_profile::kHsaViCounterBlockIdSqGs, pm4_profile::kHsaViCounterBlockIdSqVs, + pm4_profile::kHsaViCounterBlockIdSqPs, pm4_profile::kHsaViCounterBlockIdSqLs, pm4_profile::kHsaViCounterBlockIdSqHs, pm4_profile::kHsaViCounterBlockIdSqCs, pm4_profile::kHsaViCounterBlockIdSx, pm4_profile::kHsaViCounterBlockIdTa0, pm4_profile::kHsaViCounterBlockIdTca0, pm4_profile::kHsaViCounterBlockIdTcc0, pm4_profile::kHsaViCounterBlockIdTd0, pm4_profile::kHsaViCounterBlockIdTcp0, pm4_profile::kHsaViCounterBlockIdGds, pm4_profile::kHsaViCounterBlockIdVgt, pm4_profile::kHsaViCounterBlockIdIa, pm4_profile::kHsaViCounterBlockIdMc, - pm4_profile::kHsaViCounterBlockIdTcs, pm4_profile::kHsaViCounterBlockIdWd}; + pm4_profile::kHsaViCounterBlockIdSrbm, pm4_profile::kHsaViCounterBlockIdTcs, + pm4_profile::kHsaViCounterBlockIdWd, pm4_profile::kHsaViCounterBlockIdCpg, + pm4_profile::kHsaViCounterBlockIdCpc}; pm4_profile::CommandWriter* Gfx8Factory::getCommandWriter() { return new pm4_profile::gfx8::Gfx8CmdWriter(false, true); diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/gfx9_factory.cpp b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/gfx9_factory.cpp index 577ca8ed08..7a293a5e90 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/gfx9_factory.cpp +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/gfx9_factory.cpp @@ -12,19 +12,37 @@ namespace aql_profile { // GFX9 block ID mapping table uint32_t Gfx9Factory::block_id_table[HSA_EXT_AQL_PROFILE_BLOCKS_NUMBER] = { - pm4_profile::kHsaAiCounterBlockIdCb0, pm4_profile::kHsaAiCounterBlockIdCpf, - pm4_profile::kHsaAiCounterBlockIdDb0, pm4_profile::kHsaAiCounterBlockIdGrbm, - pm4_profile::kHsaAiCounterBlockIdGrbmSe, pm4_profile::kHsaAiCounterBlockIdPaSu, - pm4_profile::kHsaAiCounterBlockIdPaSc, pm4_profile::kHsaAiCounterBlockIdSpi, - pm4_profile::kHsaAiCounterBlockIdSq, pm4_profile::kHsaAiCounterBlockIdSqGs, - pm4_profile::kHsaAiCounterBlockIdSqVs, pm4_profile::kHsaAiCounterBlockIdSqPs, - pm4_profile::kHsaAiCounterBlockIdSqHs, pm4_profile::kHsaAiCounterBlockIdSqCs, - pm4_profile::kHsaAiCounterBlockIdSx, pm4_profile::kHsaAiCounterBlockIdTa0, - pm4_profile::kHsaAiCounterBlockIdTca0, pm4_profile::kHsaAiCounterBlockIdTcc0, - pm4_profile::kHsaAiCounterBlockIdTd0, pm4_profile::kHsaAiCounterBlockIdTcp0, - pm4_profile::kHsaAiCounterBlockIdGds, pm4_profile::kHsaAiCounterBlockIdVgt, - pm4_profile::kHsaAiCounterBlockIdIa, pm4_profile::kHsaAiCounterBlockIdMc, - pm4_profile::kHsaAiCounterBlockIdTcs, pm4_profile::kHsaAiCounterBlockIdWd}; + pm4_profile::kHsaAiCounterBlockIdCb0, + kBadBlockId /*CPF*/, + pm4_profile::kHsaAiCounterBlockIdDb0, + pm4_profile::kHsaAiCounterBlockIdGrbm, + pm4_profile::kHsaAiCounterBlockIdGrbmSe, + pm4_profile::kHsaAiCounterBlockIdPaSu, + pm4_profile::kHsaAiCounterBlockIdPaSc, + pm4_profile::kHsaAiCounterBlockIdSpi, + pm4_profile::kHsaAiCounterBlockIdSq, + kBadBlockId /*GFX8:SQES*/, + pm4_profile::kHsaAiCounterBlockIdSqGs, + pm4_profile::kHsaAiCounterBlockIdSqVs, + pm4_profile::kHsaAiCounterBlockIdSqPs, + kBadBlockId /*GFX8:SQLS*/, + pm4_profile::kHsaAiCounterBlockIdSqHs, + pm4_profile::kHsaAiCounterBlockIdSqCs, + pm4_profile::kHsaAiCounterBlockIdSx, + pm4_profile::kHsaAiCounterBlockIdTa0, + pm4_profile::kHsaAiCounterBlockIdTca0, + pm4_profile::kHsaAiCounterBlockIdTcc0, + pm4_profile::kHsaAiCounterBlockIdTd0, + pm4_profile::kHsaAiCounterBlockIdTcp0, + pm4_profile::kHsaAiCounterBlockIdGds, + pm4_profile::kHsaAiCounterBlockIdVgt, + pm4_profile::kHsaAiCounterBlockIdIa, + pm4_profile::kHsaAiCounterBlockIdMc, + kBadBlockId /*SRBM*/, + pm4_profile::kHsaAiCounterBlockIdTcs, + pm4_profile::kHsaAiCounterBlockIdWd, + kBadBlockId /*CPG*/, + kBadBlockId /*CPC*/}; pm4_profile::CommandWriter* Gfx9Factory::getCommandWriter() { return new pm4_profile::gfx9::Gfx9CmdWriter(false, true); diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/logger.h b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/logger.h new file mode 100644 index 0000000000..0034c04b2c --- /dev/null +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/logger.h @@ -0,0 +1,98 @@ +#ifndef _LOGGER_H_ +#define _LOGGER_H_ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace aql_profile { + +class Logger { + public: + void msg(const std::string& m) { log(m); } + + void prn(const char* fmt, ...) { + const size_t formatted_size = 256; + char formatted_string[formatted_size]; + va_list argptr; + va_start(argptr, fmt); + vsnprintf(formatted_string, formatted_size, fmt, argptr); + va_end(argptr); + msg(formatted_string); + } + + template Logger& operator<<(const T& m) { + std::ostringstream oss; + oss << m; + if (!streaming) + log(oss.str()); + else + put(oss.str()); + streaming = true; + return *this; + } + + typedef void (*manip_t)(Logger&); + Logger& operator<<(manip_t f) { + f(*this); + return *this; + } + + static void endl(Logger& logger) { logger.streaming = false; } + + Logger() : file(NULL), dirty(false), streaming(false) { + const char* path = getenv("HSA_EXT_AQL_PROFILE_LOG"); + if (path != NULL) { + file = fopen("/tmp/aql_profile_log.txt", "a"); + } + } + ~Logger() { + if (file != NULL) { + if (dirty) put("\n"); + fclose(file); + } + } + + private: + void put(const std::string& m) { + if (file != NULL) { + dirty = true; + flock(fileno(file), LOCK_EX); + fprintf(file, "%s", m.c_str()); + fflush(file); + flock(fileno(file), LOCK_UN); + } + } + + void log(const std::string& m) { + const time_t rawtime = time(NULL); + const tm* tm_info = localtime(&rawtime); + char tm_str[26]; + strftime(tm_str, 26, "%Y-%m-%d %H:%M:%S", tm_info); + std::ostringstream oss; + oss << "\n<" << tm_str << std::dec << " pid" << syscall(__NR_getpid) << " tid" + << syscall(__NR_gettid) << "> " << m; + put(oss.str()); + } + + FILE* file; + bool dirty; + bool streaming; +}; + +} // aql_profile + +#define ERR_LOGGING(logger) \ + (logger << aql_profile::Logger::endl << "Error: " << __FUNCTION__ << "(): ") + +#endif // _LOGGER_H_ diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/pm4_factory.h b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/pm4_factory.h index fbe41d23bf..9402109ae3 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/pm4_factory.h +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/pm4_factory.h @@ -8,6 +8,7 @@ #include "aql_profile.h" #include "gpu_block_info.h" +#include "aql_profile_exception.h" namespace pm4_profile { class CommandWriter; @@ -23,23 +24,31 @@ namespace aql_profile { class Pm4Factory { public: + enum { kBadBlockId = UINT_MAX }; + static Pm4Factory* Create(const hsa_ext_amd_aql_profile_profile_t* profile); virtual pm4_profile::CommandWriter* getCommandWriter() = 0; virtual pm4_profile::Pmu* getPmcMgr() = 0; virtual pm4_profile::ThreadTrace* getSqttMgr() = 0; uint32_t getBlockId(const event_t* event) { - const hsa_ext_amd_aql_profile_block_name_t block_name = event->block_name; + const hsa_ext_amd_aql_profile_block_name_t& block_name = event->block_name; + if (block_name >= tables.get_block_id_count()) + throw aql_profile_exception(std::string("Invalid block name, block_name"), + block_name); return (block_name < tables.get_block_id_count()) ? tables.get_block_id_ptr()[block_name] + event->block_index - : UINT_MAX; + : kBadBlockId; } const pm4_profile::GpuBlockInfo* getBlockInfo(const uint32_t& block_id) { const pm4_profile::GpuBlockInfo* info = NULL; if (block_id < tables.get_block_info_count()) { info = tables.get_block_info_ptr() + block_id; - if (info->counterGroupId != block_id) info = NULL; - } + if (info->counterGroupId != block_id) + throw aql_profile_exception(std::string("Bad block id table, block_id"), + block_id); + } else + throw aql_profile_exception(std::string("Invalid block id, block_id"), block_id); return info; } const pm4_profile::GpuBlockInfo* getBlockInfo(const event_t* event) { @@ -107,7 +116,10 @@ inline Pm4Factory* Pm4Factory::Create(const hsa_ext_amd_aql_profile_profile_t* p char agent_name[64]; hsa_agent_get_info(profile->agent, HSA_AGENT_INFO_NAME, agent_name); - if (strncmp(agent_name, "gfx8", 4) == 0) { + if (strncmp(agent_name, "gfx801", 6) == 0) { + throw aql_profile_exception(std::string("GFX8 Carrizo is not supported "), + agent_name); + } else if (strncmp(agent_name, "gfx8", 4) == 0) { instance = new Gfx8Factory(); } else if (strncmp(agent_name, "gfx9", 4) == 0) { instance = new Gfx9Factory(); diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/perfcounter/gfx9_block_info.cpp b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/perfcounter/gfx9_block_info.cpp index 670319d85e..0843b852d6 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/perfcounter/gfx9_block_info.cpp +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/perfcounter/gfx9_block_info.cpp @@ -24,8 +24,6 @@ GpuBlockInfo Gfx9HwBlocks[] = { {"AI_CPF", kHsaAiCounterBlockIdCpf, AI_MAX_NUM_SHADER_ENGINES, 2, 1, CntlMethodNone, 19, AI_COUNTER_NUM_PER_CPF, 0, 0, true, 0, 0, false, 0, 0}, */ - {"AI_CB3", kHsaAiCounterBlockIdCpf, AI_MAX_NUM_SHADER_ENGINES, 2, AI_NUM_CB, - CntlMethodBySeAndInstance, 395, AI_COUNTER_NUM_PER_CB, 0, 0, true, 0, 0, false, 0, 0}, // Counter block DB {"AI_DB0", kHsaAiCounterBlockIdDb0, AI_MAX_NUM_SHADER_ENGINES, 2, AI_NUM_DB, diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/perfcounter/gfx9_block_info.h b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/perfcounter/gfx9_block_info.h index 79cb0fb050..c5ef546d53 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/perfcounter/gfx9_block_info.h +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/perfcounter/gfx9_block_info.h @@ -84,8 +84,8 @@ typedef enum HsaAiCounterBlockId { kHsaAiCounterBlockIdCb2, kHsaAiCounterBlockIdCb3, - // Temp commented for Vega10 - kHsaAiCounterBlockIdCpf, + // Temp commented out for Vega10 + // kHsaAiCounterBlockIdCpf, kHsaAiCounterBlockIdDb0, kHsaAiCounterBlockIdDb1, @@ -192,8 +192,8 @@ typedef enum HsaAiCounterBlockId { // Temp commented out for Vega10 // kHsaAiCounterBlockIdCpg, - // Temp commented for Vega10 - kHsaAiCounterBlockIdCpc, + // Temp commented out for Vega10 + // kHsaAiCounterBlockIdCpc, // Counters retrieved by KFD kHsaAiCounterBlockIdIommuV2, diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test.cpp b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test.cpp index 5da9952fd1..da865e1d4d 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test.cpp +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test.cpp @@ -25,7 +25,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE. *******************************************************************************/ -#include +#include "test_assert.h" #include "simple_convolution.h" #include "test_hsa.h" #include "test_pgen_pmc.h" @@ -48,14 +48,14 @@ int main(int argc, char* argv[]) { test_aql = new TestPGenPMC(test_aql); else if (sqtt_enable) test_aql = new TestPGenSQTT(test_aql); - assert(test_aql != NULL); + test_assert(test_aql != NULL); if (test_aql == NULL) return 1; // Initialization of Hsa Runtime ret_val = test_aql->initialize(argc, argv); if (ret_val == false) { std::cout << "Error in the test initialization" << std::endl; - assert(ret_val); + test_assert(ret_val); return 1; } @@ -63,7 +63,7 @@ int main(int argc, char* argv[]) { ret_val = test_aql->setup(); if (ret_val == false) { std::cout << "Error in creating hsa resources" << std::endl; - assert(ret_val); + test_assert(ret_val); return 1; } @@ -71,7 +71,7 @@ int main(int argc, char* argv[]) { ret_val = test_aql->run(); if (ret_val == false) { std::cout << "Error in running the test kernel" << std::endl; - assert(ret_val); + test_assert(ret_val); return 1; } diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_aql.h b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_aql.h index 3504cf0d41..bee82a5867 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_aql.h +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_aql.h @@ -25,22 +25,13 @@ OF THE POSSIBILITY OF SUCH DAMAGE. *******************************************************************************/ -#ifndef _TESTAQL_H_ -#define _TESTAQL_H_ +#ifndef _TEST_AQL_H_ +#define _TEST_AQL_H_ #include "hsa.h" #include "hsa_rsrc_factory.h" #include "hsa_ext_amd_aql_profile.h" -#define test_assert(cond) \ - { \ - if (cond) { \ - std::cout << "ASSERT FAILED: " << #cond << " : " << __FILE__ << "(" << __LINE__ << ")" \ - << std::endl; \ - abort(); \ - } \ - } - // Test AQL interface class TestAql { TestAql* const test_aql; @@ -84,4 +75,4 @@ class TestAql { virtual bool cleanup() { return (test_aql) ? test_aql->cleanup() : true; } }; -#endif // _TESTAQL_H_ +#endif // _TEST_AQL_H_ diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_assert.h b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_assert.h new file mode 100644 index 0000000000..96292a68ed --- /dev/null +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_assert.h @@ -0,0 +1,13 @@ +#ifndef _TEST_ASSERT_H_ +#define _TEST_ASSERT_H_ + +#define test_assert(cond) \ + { \ + if (!(cond)) { \ + std::cout << "ASSERT FAILED(" << #cond << ") at \"" << __FILE__ << "\" line " << __LINE__ \ + << std::endl; \ + abort(); \ + } \ + } + +#endif // _TEST_ASSERT_H_ diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_hsa.cpp b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_hsa.cpp index 5b0427fa0d..69bfca56ed 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_hsa.cpp +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_hsa.cpp @@ -25,7 +25,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE. *******************************************************************************/ -#include +#include "test_assert.h" #include //#include "os.h" @@ -72,7 +72,7 @@ bool TestHSA::initialize(int arg_cnt, char** arg_list) { } else if (agentName.compare(0, 4, "gfx9") == 0) { brig_path_obj_.append("gfx9"); } else { - assert(false); + test_assert(false); return false; } brig_path_obj_.append("_" + name_ + ".hsaco"); @@ -92,7 +92,7 @@ bool TestHSA::setup() { void* ptr = (des.local) ? hsa_rsrc_->AllocateLocalMemory(agent_info_, des.size) : hsa_rsrc_->AllocateSysMemory(agent_info_, des.size); des.ptr = ptr; - assert(ptr != NULL); + test_assert(ptr != NULL); if (ptr == NULL) return false; } test_->init(); @@ -137,7 +137,7 @@ bool TestHSA::run() { size_t size_info = 0; hsa_executable_symbol_get_info( kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &size_info); - assert(kernarg_segment_size == size_info); + test_assert(kernarg_segment_size == size_info); if (kernarg_segment_size != size_info) return false; // Retrieve handle of the code block diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pgen_pmc.h b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pgen_pmc.h index b3f5239c8b..d83591f104 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pgen_pmc.h +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pgen_pmc.h @@ -28,6 +28,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE. #ifndef _TEST_PGEN_PMC_H_ #define _TEST_PGEN_PMC_H_ +#include "test_assert.h" #include "test_pgen.h" hsa_status_t TestPGenPMC_Callback(hsa_ext_amd_aql_profile_info_type_t info_type, @@ -104,12 +105,12 @@ class TestPGenPMC : public TestPGen { command_buffer_alignment = buffer_alignment; status = hsa_ext_amd_aql_profile_get_info( &profile, HSA_EXT_AQL_PROFILE_INFO_COMMAND_BUFFER_SIZE, &command_buffer_size); - assert(status == HSA_STATUS_SUCCESS); + test_assert(status == HSA_STATUS_SUCCESS); output_buffer_alignment = buffer_alignment; status = hsa_ext_amd_aql_profile_get_info(&profile, HSA_EXT_AQL_PROFILE_INFO_PMC_DATA_SIZE, &output_buffer_size); - assert(status == HSA_STATUS_SUCCESS); + test_assert(status == HSA_STATUS_SUCCESS); // Application is allocating the command buffer // Allocate(command_buffer_alignment, command_buffer_size, @@ -128,12 +129,12 @@ class TestPGenPMC : public TestPGen { // Populating the AQL start packet status = hsa_ext_amd_aql_profile_start(&profile, PrePacket()); - assert(status == HSA_STATUS_SUCCESS); + test_assert(status == HSA_STATUS_SUCCESS); if (status != HSA_STATUS_SUCCESS) return false; // Populating the AQL stop packet status = hsa_ext_amd_aql_profile_stop(&profile, PostPacket()); - assert(status == HSA_STATUS_SUCCESS); + test_assert(status == HSA_STATUS_SUCCESS); return (status == HSA_STATUS_SUCCESS); } diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pgen_sqtt.h b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pgen_sqtt.h index c4728258c7..6cdf0ce96e 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pgen_sqtt.h +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pgen_sqtt.h @@ -32,6 +32,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE. #include #include +#include "test_assert.h" #include "test_pgen.h" hsa_status_t TestPGenSQTT_Callback(hsa_ext_amd_aql_profile_info_type_t info_type, @@ -65,11 +66,11 @@ class TestPGenSQTT : public TestPGen { << ") size(" << dec << it->sqtt_data.size << ")" << std::endl; void* sys_buf = getRsrcFactory()->AllocateSysMemory(getAgentInfo(), it->sqtt_data.size); - assert(sys_buf != NULL); + test_assert(sys_buf != NULL); if (sys_buf == NULL) return HSA_STATUS_ERROR; hsa_status_t status = hsa_memory_copy(sys_buf, it->sqtt_data.ptr, it->sqtt_data.size); - assert(status == HSA_STATUS_SUCCESS); + test_assert(status == HSA_STATUS_SUCCESS); if (status != HSA_STATUS_SUCCESS) return status; std::string file_name; @@ -125,7 +126,7 @@ class TestPGenSQTT : public TestPGen { command_buffer_alignment = buffer_alignment; status = hsa_ext_amd_aql_profile_get_info( &profile, HSA_EXT_AQL_PROFILE_INFO_COMMAND_BUFFER_SIZE, &command_buffer_size); - assert(status == HSA_STATUS_SUCCESS); + test_assert(status == HSA_STATUS_SUCCESS); output_buffer_alignment = buffer_alignment; output_buffer_size = buffer_size; @@ -146,12 +147,12 @@ class TestPGenSQTT : public TestPGen { // Populating the AQL start packet status = hsa_ext_amd_aql_profile_start(&profile, PrePacket()); - assert(status == HSA_STATUS_SUCCESS); + test_assert(status == HSA_STATUS_SUCCESS); if (status != HSA_STATUS_SUCCESS) return false; // Populating the AQL stop packet status = hsa_ext_amd_aql_profile_stop(&profile, PostPacket()); - assert(status == HSA_STATUS_SUCCESS); + test_assert(status == HSA_STATUS_SUCCESS); return (status == HSA_STATUS_SUCCESS); } diff --git a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.cpp b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.cpp index 54f311a30d..cef2b8cc5c 100644 --- a/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.cpp +++ b/projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.cpp @@ -26,7 +26,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE. *******************************************************************************/ #include -#include +#include "test_assert.h" #include "test_pmgr.h" @@ -115,7 +115,7 @@ bool TestPMgr::run() { bool TestPMgr::initialize(int argc, char** argv) { TestAql::initialize(argc, argv); hsa_status_t status = hsa_signal_create(1, 0, NULL, &postSignal); - assert(status == HSA_STATUS_SUCCESS); + test_assert(status == HSA_STATUS_SUCCESS); return (status == HSA_STATUS_SUCCESS); }