AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

dispatch data th-id, queue-id, kernel-object

Browse Source 
Change-Id: Ic5593603c6587cad17be33618d011a9aefc648ab
This commit is contained in:
Evgeny
2018-10-16 19:33:48 -05:00
parent 13fe0f9f69
commit c05bded17c
11 changed files with 523 additions and 23 deletions
Download Patch File Download Diff File Expand all files Collapse all files
test/ctrl/test_hsa.cpp
+12 -10
View File
@@ -50,15 +50,6 @@ HsaRsrcFactory* TestHsa::HsaInstantiate(const uint32_t agent_ind) {
return NULL;
}
std::clog << "> Using agent[" << agent_ind << "] : " << agent_info_->name << std::endl;
// Create an instance of Aql Queue
if (hsa_queue_ == NULL) {
uint32_t num_pkts = 128;
if (hsa_rsrc_->CreateQueue(agent_info_, num_pkts, &hsa_queue_) == false) {
hsa_queue_ = NULL;
TEST_ASSERT(false);
}
}
}
return hsa_rsrc_;
}
@@ -74,6 +65,15 @@ void TestHsa::HsaShutdown() {
bool TestHsa::Initialize(int /*arg_cnt*/, char** /*arg_list*/) {
std::clog << "TestHsa::Initialize :" << std::endl;
// Create an instance of Aql Queue
if (hsa_queue_ == NULL) {
uint32_t num_pkts = 128;
if (hsa_rsrc_->CreateQueue(agent_info_, num_pkts, &hsa_queue_) == false) {
hsa_queue_ = NULL;
TEST_ASSERT(false);
}
}
// Instantiate a Timer object
setup_timer_idx_ = hsa_timer_.CreateTimer();
dispatch_timer_idx_ = hsa_timer_.CreateTimer();
@@ -222,7 +222,7 @@ bool TestHsa::Run() {
// Submit AQL packet to the queue
const uint64_t que_idx = hsa_rsrc_->Submit(hsa_queue_, &aql);
std::clog << "> Waiting on kernel dispatch signal, que_idx=" << que_idx << std::endl;
std::clog << "> Waiting on kernel dispatch signal, que_idx=" << que_idx << std::endl << std::flush;
// Wait on the dispatch signal until the kernel is finished.
// Update wait condition to HSA_WAIT_STATE_ACTIVE for Polling
@@ -283,5 +283,7 @@ void TestHsa::PrintTime() {
bool TestHsa::Cleanup() {
hsa_executable_destroy(hsa_exec_);
hsa_signal_destroy(hsa_signal_);
hsa_queue_destroy(hsa_queue_);
hsa_queue_ = NULL;
return true;
}