AI/rocm-systems
2
0
派生 0
代码 工单 合并请求 工作流 软件包 项目 发布 百科 活动

Block list extending

浏览代码 
Change-Id: Id17efde25fce287296e80f2b37c77b15aa59b561


[ROCm/ROCR-Runtime commit: c533229bc1]
这个提交包含在：
Evgeny
2017-06-22 18:41:05 -05:00
提交者 Evgeny Shcherbakov
父节点 9725267667
当前提交 7892cc861c
共修改 17 个文件，包含 453 行新增和 242 行删除
下载 Patch 文件 下载 Diff 文件 展开所有文件 折叠所有文件
projects/rocr-runtime/runtime/hsa-ext-aql-profile/inc/hsa_ext_amd_aql_profile.h
+5
查看文件
@@ -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;
projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.cpp
+230 -182
查看文件
@@ -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<hsa_ext_amd_aql_profile_info_data_t*>(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<hsa_ext_amd_aql_profile_info_data_t*>(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;
projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile.h
-1
查看文件
@@ -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;
projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/aql_profile_exception.h
+24
查看文件
@@ -0,0 +1,24 @@
#ifndef _AQL_PROFILE_EXCEPTION_H_
#define _AQL_PROFILE_EXCEPTION_H_
#include <string>
#include <sstream>
namespace aql_profile {
template <typename T> 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_
projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/gfx8_factory.cpp
+6 -3
查看文件
@@ -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);
projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/gfx9_factory.cpp
+31 -13
查看文件
@@ -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);
projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/logger.h
+98
查看文件
@@ -0,0 +1,98 @@
#ifndef _LOGGER_H_
#define _LOGGER_H_
#include <time.h>
#include <stdio.h>
#include <unistd.h>
#include <sys/types.h>
#include <sys/syscall.h>
#include <sys/file.h>
#include <stdarg.h>
#include <stdlib.h>
#include <string>
#include <iostream>
#include <sstream>
#include <fstream>
#include <exception>
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 <typename T> 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_
projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/aqlprofile/pm4_factory.h
+17 -5
查看文件
@@ -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<uint32_t>(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<uint32_t>(std::string("Bad block id table, block_id"),
block_id);
} else
throw aql_profile_exception<uint32_t>(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>(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();
projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/perfcounter/gfx9_block_info.cpp
-2
查看文件
@@ -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,
projects/rocr-runtime/runtime/hsa-ext-aql-profile/src/perfcounter/gfx9_block_info.h
+4 -4
查看文件
@@ -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,
projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test.cpp
+5 -5
查看文件
@@ -25,7 +25,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#include <assert.h>
#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;
}
projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_aql.h
+3 -12
查看文件
@@ -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_
projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_assert.h
+13
查看文件
@@ -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_
projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_hsa.cpp
+4 -4
查看文件
@@ -25,7 +25,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#include <assert.h>
#include "test_assert.h"
#include <atomic>
//#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
projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pgen_pmc.h
+5 -4
查看文件
@@ -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);
}
projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pgen_sqtt.h
+6 -5
查看文件
@@ -32,6 +32,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE.
#include <iomanip>
#include <fstream>
#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);
}
projects/rocr-runtime/runtime/hsa-ext-aql-profile/test/ctrl/test_pmgr.cpp
+2 -2
查看文件
@@ -26,7 +26,7 @@ OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#include <atomic>
#include <assert.h>
#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);
}