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

ROC profiler prototype sources importing

Browse Source
This commit is contained in:
Evgeny
2017-11-09 17:26:19 -06:00
parent 54156e1953
commit 85278f08a0
63 changed files with 7598 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
test/ctrl/run_kernel.h
+85
View File
@@ -0,0 +1,85 @@
/******************************************************************************
Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
Redistribution and use in source and binary forms, with or without modification,
are permitted provided that the following conditions are met:
Redistributions of source code must retain the above copyright notice, this list
of conditions and the following disclaimer.
Redistributions in binary form must reproduce the above copyright notice, this
list of conditions and the following disclaimer in the documentation and/or
other materials provided with the distribution.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#ifndef TEST_CTRL_RUN_KERNEL_H_
#define TEST_CTRL_RUN_KERNEL_H_
#include "ctrl/test_hsa.h"
#include "util/test_assert.h"
template <class Kernel, class Test> bool RunKernel(int argc, char* argv[]) {
bool ret_val = false;
// Create test kernel object
Kernel test_kernel;
TestAql* test_aql = new TestHsa(&test_kernel);
test_aql = new Test(test_aql);
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::cerr << "Error in the test initialization" << std::endl;
// TEST_ASSERT(ret_val);
return false;
}
// Setup Hsa resources needed for execution
ret_val = test_aql->Setup();
if (ret_val == false) {
std::cerr << "Error in creating hsa resources" << std::endl;
TEST_ASSERT(ret_val);
return false;
}
// Run test kernel
ret_val = test_aql->Run();
if (ret_val == false) {
std::cerr << "Error in running the test kernel" << std::endl;
TEST_ASSERT(ret_val);
return false;
}
// Verify the results of the execution
ret_val = test_aql->VerifyResults();
if (ret_val) {
std::clog << "Test : Passed" << std::endl;
} else {
std::clog << "Test : Failed" << std::endl;
}
// Print time taken by sample
test_aql->PrintTime();
test_aql->Cleanup();
delete test_aql;
return ret_val;
}
#endif // TEST_CTRL_RUN_KERNEL_H_
test/ctrl/standalone_test.cpp
+126
View File
@@ -0,0 +1,126 @@
#include <hsa.h>
#include <string.h>
#include <iostream>
#include "ctrl/run_kernel.h"
#include "ctrl/test_aql.h"
#include "ctrl/test_hsa.h"
#include "inc/rocprofiler.h"
#include "simple_convolution/simple_convolution.h"
#include "util/test_assert.h"
int main(int argc, char** argv) {
bool ret_val = false;
// HSA status
hsa_status_t status = HSA_STATUS_ERROR;
// Profiling context
rocprofiler_t* context = NULL;
// Profiling properties
rocprofiler_properties_t properties;
// Number of context invocation
uint32_t invocation = 0;
#if 0
// Profiling info objects
const unsigned info_count = 1;
rocprofiler_info_t info[info_count];
// PMC events
memset(info, 0, sizeof(info));
info[0].type = ROCPROFILER_TYPE_METRIC;
info[0].name = "SQ_WAVES";
#else
// Profiling info objects
const unsigned info_count = 3;
rocprofiler_info_t info[info_count];
// PMC events
memset(info, 0, sizeof(info));
info[0].type = ROCPROFILER_TYPE_METRIC;
info[0].name = "SQ_WAVES";
info[1].type = ROCPROFILER_TYPE_METRIC;
info[1].name = "SQ_ITEMS";
// Tracing parameters
const unsigned parameter_count = 2;
rocprofiler_parameter_t parameters[parameter_count];
info[2].name = "THREAD_TRACE";
info[2].type = ROCPROFILER_TYPE_TRACE;
info[2].parameters = parameters;
info[2].parameter_count = parameter_count;
parameters[0].parameter_name = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_MASK;
parameters[0].value = 0;
parameters[1].parameter_name = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK;
parameters[1].value = 0;
#endif
// Creating profiling context
properties = {};
properties.queue_depth = 128;
status = rocprofiler_open(TestHsa::HsaAgentId(), info, info_count, &context, ROCPROFILER_MODE_STANDALONE|ROCPROFILER_MODE_OWNQUEUE, &properties);
TEST_STATUS(status == HSA_STATUS_SUCCESS);
TestHsa::SetQueue(properties.queue);
// Adding dispatch observer
status = rocprofiler_dispatch_observer(rocprofiler_dispatch_callback, context);
TEST_STATUS(status == HSA_STATUS_SUCCESS);
// Querying the number of context invocation
status = rocprofiler_invocation(context, &invocation);
TEST_STATUS(status == HSA_STATUS_SUCCESS);
// Dispatching profiled kernel n-times to collect all counter groups data
unsigned n = 0;
while(1) {
std::cout << "> " << n << "/" << invocation << std::endl;
#if 0
status = rocprofiler_start(context);
TEST_STATUS(status == HSA_STATUS_SUCCESS);
ret_val = RunKernel<SimpleConvolution, TestAql>(argc, argv);
status = rocprofiler_stop(context);
TEST_STATUS(status == HSA_STATUS_SUCCESS);
#else
ret_val = RunKernel<SimpleConvolution, TestAql>(argc, argv);
#endif
status = rocprofiler_sample(context);
TEST_STATUS(status == HSA_STATUS_SUCCESS);
for (rocprofiler_info_t* p = info; p < info + info_count; ++p) {
std::cout << (p - info) << ": " << p->name;
switch (p->data.kind) {
case ROCPROFILER_INT64:
std::cout << std::dec << " result64 (" << p->data.result64 << ")" << std::endl;
break;
case ROCPROFILER_BYTES: {
const char* ptr = reinterpret_cast<const char*>(p->data.result_bytes.ptr);
uint64_t size = 0;
for (unsigned i = 0; i < p->data.result_bytes.instance_count; ++i) {
size = *reinterpret_cast<const uint64_t*>(ptr);
const char* data = ptr + sizeof(size);
std::cout << std::endl;
std::cout << std::hex << " data (" << (void*)data << ")" << std::endl;
std::cout << std::dec << " size (" << size << ")" << std::endl;
ptr = data + size;
}
break;
}
default:
std::cout << "result kind (" << p->data.kind << ")" << std::endl;
TEST_ASSERT(false);
}
}
++n;
if (n < invocation) {
status = rocprofiler_next(context);
TEST_STATUS(status == HSA_STATUS_SUCCESS);
continue;
}
break;
}
// Finishing cleanup
// Deleting profiling context will delete all allocated resources
status = rocprofiler_close(context);
TEST_STATUS(status == HSA_STATUS_SUCCESS);
return (ret_val) ? 0 : 1;
}
test/ctrl/test.cpp
+14
View File
@@ -0,0 +1,14 @@
#include <hsa.h>
#include <string.h>
#include <iostream>
#include "ctrl/run_kernel.h"
#include "ctrl/test_aql.h"
#include "simple_convolution/simple_convolution.h"
int main(int argc, char** argv) {
TestHsa::HsaInstantiate();
for (int i = 0; i < 3; ++i) RunKernel<SimpleConvolution, TestAql>(argc, argv);
TestHsa::HsaShutdown();
return 0;
}
test/ctrl/test_aql.h
+80
View File
@@ -0,0 +1,80 @@
/******************************************************************************
Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
Redistribution and use in source and binary forms, with or without modification,
are permitted provided that the following conditions are met:
Redistributions of source code must retain the above copyright notice, this list
of conditions and the following disclaimer.
Redistributions in binary form must reproduce the above copyright notice, this
list of conditions and the following disclaimer in the documentation and/or
other materials provided with the distribution.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#ifndef TEST_CTRL_TEST_AQL_H_
#define TEST_CTRL_TEST_AQL_H_
#include <hsa.h>
#include <hsa_ven_amd_aqlprofile.h>
#include "util/hsa_rsrc_factory.h"
// Test AQL interface
class TestAql {
public:
explicit TestAql(TestAql* t = 0) : test_(t) {}
virtual ~TestAql() { if (test_) delete test_; }
TestAql* Test() { return test_; }
virtual AgentInfo* GetAgentInfo() { return (test_) ? test_->GetAgentInfo() : 0; }
virtual hsa_queue_t* GetQueue() { return (test_) ? test_->GetQueue() : 0; }
virtual HsaRsrcFactory* GetRsrcFactory() { return (test_) ? test_->GetRsrcFactory() : 0; }
// Initialize application environment including setting
// up of various configuration parameters based on
// command line arguments
// @return bool true on success and false on failure
virtual bool Initialize(int argc, char** argv) {
return (test_) ? test_->Initialize(argc, argv) : true;
}
// Setup application parameters for exectuion
// @return bool true on success and false on failure
virtual bool Setup() { return (test_) ? test_->Setup() : true; }
// Run the kernel
// @return bool true on success and false on failure
virtual bool Run() { return (test_) ? test_->Run() : true; }
// Verify results
// @return bool true on success and false on failure
virtual bool VerifyResults() { return (test_) ? test_->VerifyResults() : true; }
// Print to console the time taken to execute kernel
virtual void PrintTime() {
if (test_) test_->PrintTime();
}
// Release resources e.g. memory allocations
// @return bool true on success and false on failure
virtual bool Cleanup() { return (test_) ? test_->Cleanup() : true; }
private:
TestAql* const test_;
};
#endif // TEST_CTRL_TEST_AQL_H_
test/ctrl/test_hsa.cpp
+252
View File
@@ -0,0 +1,252 @@
/******************************************************************************
Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
Redistribution and use in source and binary forms, with or without modification,
are permitted provided that the following conditions are met:
Redistributions of source code must retain the above copyright notice, this list
of conditions and the following disclaimer.
Redistributions in binary form must reproduce the above copyright notice, this
list of conditions and the following disclaimer in the documentation and/or
other materials provided with the distribution.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#include "ctrl/test_hsa.h"
#include <atomic>
#include "util/test_assert.h"
#include "util/helper_funcs.h"
#include "util/hsa_rsrc_factory.h"
HsaRsrcFactory* TestHsa::hsa_rsrc_ = NULL;
AgentInfo* TestHsa::agent_info_ = NULL;
hsa_queue_t* TestHsa::hsa_queue_ = NULL;
uint32_t TestHsa::agent_id_ = 0;
HsaRsrcFactory* TestHsa::HsaInstantiate(const uint32_t agent_ind) {
// Instantiate an instance of Hsa Resources Factory
if (hsa_rsrc_ == NULL) {
agent_id_ = agent_ind;
hsa_rsrc_ = HsaRsrcFactory::Create();
// Print properties of the agents
hsa_rsrc_->PrintGpuAgents("> GPU agents");
// Create an instance of Gpu agent
if (!hsa_rsrc_->GetGpuAgentInfo(agent_ind, &agent_info_)) {
agent_info_ = NULL;
std::cerr << "> error: agent[" << agent_ind << "] is not found" << std::endl;
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;
}
}
}
return hsa_rsrc_;
}
void TestHsa::HsaShutdown() { if (hsa_rsrc_) hsa_rsrc_->Destroy(); }
bool TestHsa::Initialize(int arg_cnt, char** arg_list) {
std::clog << "TestHsa::Initialize :" << std::endl;
// Instantiate a Timer object
setup_timer_idx_ = hsa_timer_.CreateTimer();
dispatch_timer_idx_ = hsa_timer_.CreateTimer();
hsa_rsrc_ = HsaInstantiate(agent_id_);
if (hsa_rsrc_ == NULL) {
TEST_ASSERT(false);
return false;
}
// Obtain handle of signal
hsa_rsrc_->CreateSignal(1, &hsa_signal_);
// Obtain the code object file name
std::string agentName(agent_info_->name);
if (agentName.compare(0, 4, "gfx8") == 0) {
brig_path_obj_.append("gfx8");
} else if (agentName.compare(0, 4, "gfx9") == 0) {
brig_path_obj_.append("gfx9");
} else {
TEST_ASSERT(false);
return false;
}
brig_path_obj_.append("_" + name_ + ".hsaco");
return true;
}
bool TestHsa::Setup() {
std::clog << "TestHsa::setup :" << std::endl;
// Start the timer object
hsa_timer_.StartTimer(setup_timer_idx_);
mem_map_t& mem_map = test_->GetMemMap();
for (mem_it_t it = mem_map.begin(); it != mem_map.end(); ++it) {
mem_descr_t& des = it->second;
void* ptr = (des.local) ? hsa_rsrc_->AllocateLocalMemory(agent_info_, des.size)
: hsa_rsrc_->AllocateSysMemory(agent_info_, des.size);
des.ptr = ptr;
TEST_ASSERT(ptr != NULL);
if (ptr == NULL) return false;
}
test_->Init();
// Load and Finalize Kernel Code Descriptor
char* brig_path = (char*)brig_path_obj_.c_str();
const bool ret_val =
hsa_rsrc_->LoadAndFinalize(agent_info_, brig_path, strdup(name_.c_str()), &kernel_code_desc_);
if (ret_val == false) {
std::cerr << "Error in loading and finalizing Kernel" << std::endl;
return ret_val;
}
// Stop the timer object
hsa_timer_.StopTimer(setup_timer_idx_);
setup_time_taken_ = hsa_timer_.ReadTimer(setup_timer_idx_);
total_time_taken_ = setup_time_taken_;
return true;
}
bool TestHsa::Run() {
std::clog << "TestHsa::run :" << std::endl;
const uint32_t work_group_size = 64;
const uint32_t work_grid_size = test_->GetGridSize();
uint32_t group_segment_size = 0;
uint32_t private_segment_size = 0;
const size_t kernarg_segment_size = test_->GetKernargSize();
uint64_t code_handle = 0;
// Retrieve the amount of group memory needed
hsa_executable_symbol_get_info(
kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &group_segment_size);
// Retrieve the amount of private memory needed
hsa_executable_symbol_get_info(kernel_code_desc_,
HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
&private_segment_size);
// Check the kernel args size
size_t size_info = 0;
hsa_executable_symbol_get_info(
kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_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
hsa_executable_symbol_get_info(kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
&code_handle);
// Initialize the dispatch packet.
hsa_kernel_dispatch_packet_t aql;
memset(&aql, 0, sizeof(aql));
// Set the packet's type, barrier bit, acquire and release fences
aql.header = HSA_PACKET_TYPE_KERNEL_DISPATCH;
aql.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCACQUIRE_FENCE_SCOPE;
aql.header |= HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_SCRELEASE_FENCE_SCOPE;
// Populate Aql packet with default values
aql.setup = 1;
aql.grid_size_x = work_grid_size;
aql.grid_size_y = 1;
aql.grid_size_z = 1;
aql.workgroup_size_x = work_group_size;
aql.workgroup_size_y = 1;
aql.workgroup_size_z = 1;
// Bind the kernel code descriptor and arguments
aql.kernel_object = code_handle;
aql.kernarg_address = test_->GetKernargPtr();
aql.group_segment_size = group_segment_size;
aql.private_segment_size = private_segment_size;
// Initialize Aql packet with handle of signal
aql.completion_signal = hsa_signal_;
// Compute the write index of queue and copy Aql packet into it
const uint64_t que_idx = hsa_queue_load_write_index_relaxed(hsa_queue_);
const uint32_t mask = hsa_queue_->size - 1;
std::clog << "> Executing kernel: \"" << name_ << "\"" << std::endl;
// Start the timer object
hsa_timer_.StartTimer(dispatch_timer_idx_);
// Disable packet so that submission to HW is complete
const auto header = aql.header;
aql.header = HSA_PACKET_TYPE_INVALID << HSA_PACKET_HEADER_TYPE;
// Copy Aql packet into queue buffer
((hsa_kernel_dispatch_packet_t*)(hsa_queue_->base_address))[que_idx & mask] = aql;
// After AQL packet is fully copied into queue buffer
// update packet header from invalid state to valid state
std::atomic_thread_fence(std::memory_order_release);
((hsa_kernel_dispatch_packet_t*)(hsa_queue_->base_address))[que_idx & mask].header = header;
// Increment the write index and ring the doorbell to dispatch the kernel.
hsa_queue_store_write_index_relaxed(hsa_queue_, (que_idx + 1));
hsa_signal_store_relaxed(hsa_queue_->doorbell_signal, que_idx);
std::clog << "> Waiting on kernel dispatch signal, que_idx=" << que_idx << std::endl;
// Wait on the dispatch signal until the kernel is finished.
// Update wait condition to HSA_WAIT_STATE_ACTIVE for Polling
hsa_signal_wait_acquire(hsa_signal_, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)-1,
HSA_WAIT_STATE_BLOCKED);
// Stop the timer object
hsa_timer_.StopTimer(dispatch_timer_idx_);
dispatch_time_taken_ = hsa_timer_.ReadTimer(dispatch_timer_idx_);
total_time_taken_ += dispatch_time_taken_;
// Copy kernel buffers from local memory into system memory
hsa_rsrc_->TransferData(test_->GetOutputPtr(), test_->GetLocalPtr(), test_->GetOutputSize(),
false);
test_->PrintOutput();
return true;
}
bool TestHsa::VerifyResults() {
// Compare the results and see if they match
const void* const refout_ptr = test_->GetRefoutPtr();
const int32_t cmp_val =
(refout_ptr != NULL) ? memcmp(test_->GetOutputPtr(), refout_ptr, test_->GetOutputSize()) : 0;
return (cmp_val == 0);
}
void TestHsa::PrintTime() {
std::clog << "Time taken for Setup by " << this->name_ << " : " << this->setup_time_taken_
<< std::endl;
std::clog << "Time taken for Dispatch by " << this->name_ << " : " << this->dispatch_time_taken_
<< std::endl;
std::clog << "Time taken in Total by " << this->name_ << " : " << this->total_time_taken_
<< std::endl;
}
bool TestHsa::Cleanup() { return true; }
test/ctrl/test_hsa.h
+125
View File
@@ -0,0 +1,125 @@
/******************************************************************************
Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
Redistribution and use in source and binary forms, with or without modification,
are permitted provided that the following conditions are met:
Redistributions of source code must retain the above copyright notice, this list
of conditions and the following disclaimer.
Redistributions in binary form must reproduce the above copyright notice, this
list of conditions and the following disclaimer in the documentation and/or
other materials provided with the distribution.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#ifndef TEST_CTRL_TEST_HSA_H_
#define TEST_CTRL_TEST_HSA_H_
#include "ctrl/test_aql.h"
#include "ctrl/test_kernel.h"
#include "util/hsa_rsrc_factory.h"
#include "util/perf_timer.h"
// Class implements HSA test
class TestHsa : public TestAql {
public:
// Instantiate HSA resources
static HsaRsrcFactory* HsaInstantiate(const uint32_t agent_ind = agent_id_);
static void HsaShutdown();
static void SetQueue(hsa_queue_t* queue) { hsa_queue_ = queue; }
static uint32_t HsaAgentId() { return agent_id_; }
// Constructor
explicit TestHsa(TestKernel* test) : test_(test), name_(test->Name()) {
total_time_taken_ = 0;
setup_time_taken_ = 0;
dispatch_time_taken_ = 0;
}
// Get methods for Agent Info, HAS queue, HSA Resourcse Manager
AgentInfo* GetAgentInfo() { return agent_info_; }
hsa_queue_t* GetQueue() { return hsa_queue_; }
HsaRsrcFactory* GetRsrcFactory() { return hsa_rsrc_; }
// Initialize application environment including setting
// up of various configuration parameters based on
// command line arguments
// @return bool true on success and false on failure
bool Initialize(int argc, char** argv);
// Setup application parameters for exectuion
// @return bool true on success and false on failure
bool Setup();
// Run the BinarySearch kernel
// @return bool true on success and false on failure
bool Run();
// Verify against reference implementation
// @return bool true on success and false on failure
bool VerifyResults();
// Print to console the time taken to execute kernel
void PrintTime();
// Release resources e.g. memory allocations
// @return bool true on success and false on failure
bool Cleanup();
private:
typedef TestKernel::mem_descr_t mem_descr_t;
typedef TestKernel::mem_map_t mem_map_t;
typedef TestKernel::mem_it_t mem_it_t;
// Test object
TestKernel* test_;
// Path of Brig file
std::string brig_path_obj_;
// Used to track time taken to run the sample
double total_time_taken_;
double setup_time_taken_;
double dispatch_time_taken_;
// Handle of signal
hsa_signal_t hsa_signal_;
// Handle of Kernel Code Descriptor
hsa_executable_symbol_t kernel_code_desc_;
// Instance of timer object
uint32_t setup_timer_idx_;
uint32_t dispatch_timer_idx_;
PerfTimer hsa_timer_;
// Instance of Hsa Resources Factory
static HsaRsrcFactory* hsa_rsrc_;
// GPU id
static uint32_t agent_id_;
// Handle to an Hsa Gpu Agent
static AgentInfo* agent_info_;
// Handle to an Hsa Queue
static hsa_queue_t* hsa_queue_;
// Test kernel name
std::string name_;
};
#endif // TEST_CTRL_TEST_HSA_H_
test/ctrl/test_kernel.h
+107
View File
@@ -0,0 +1,107 @@
/******************************************************************************
Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
Redistribution and use in source and binary forms, with or without modification,
are permitted provided that the following conditions are met:
Redistributions of source code must retain the above copyright notice, this list
of conditions and the following disclaimer.
Redistributions in binary form must reproduce the above copyright notice, this
list of conditions and the following disclaimer in the documentation and/or
other materials provided with the distribution.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#ifndef TEST_CTRL_TEST_KERNEL_H_
#define TEST_CTRL_TEST_KERNEL_H_
#include <stdint.h>
#include <map>
// Class implements kernel test
class TestKernel {
public:
// Memory descriptors IDs
enum { INPUT_DES_ID, OUTPUT_DES_ID, LOCAL_DES_ID, MASK_DES_ID, KERNARG_DES_ID, REFOUT_DES_ID };
// Memory descriptors vector declaration
struct mem_descr_t {
void* ptr;
uint32_t size;
bool local;
};
// Memory map declaration
typedef std::map<uint32_t, mem_descr_t> mem_map_t;
typedef mem_map_t::iterator mem_it_t;
typedef mem_map_t::const_iterator mem_const_it_t;
virtual ~TestKernel() {}
// Initialize method
virtual void Init() = 0;
// Return kernel memory map
mem_map_t& GetMemMap() { return mem_map_; }
// Return NULL descriptor
static mem_descr_t NullDescriptor() { return {NULL, 0, 0}; }
// Methods to get the kernel attributes
void* GetKernargPtr() const { return GetDescr(KERNARG_DES_ID).ptr; }
uint32_t GetKernargSize() const { return GetDescr(KERNARG_DES_ID).size; }
void* GetOutputPtr() const { return GetDescr(OUTPUT_DES_ID).ptr; }
uint32_t GetOutputSize() const { return GetDescr(OUTPUT_DES_ID).size; }
void* GetLocalPtr() const { return GetDescr(LOCAL_DES_ID).ptr; }
void* GetRefoutPtr() const { return GetDescr(REFOUT_DES_ID).ptr; }
virtual uint32_t GetGridSize() const = 0;
// Print output
virtual void PrintOutput() const = 0;
// Return name
virtual std::string Name() const = 0;
protected:
// Set system memory descriptor
bool SetSysDescr(const uint32_t& id, const uint32_t& size) {
return SetMemDescr(id, size, false);
}
// Set local memory descriptor
bool SetLocalDescr(const uint32_t& id, const uint32_t& size) {
return SetMemDescr(id, size, true);
}
// Get memory descriptor
mem_descr_t GetDescr(const uint32_t& id) const {
mem_const_it_t it = mem_map_.find(id);
return (it != mem_map_.end()) ? it->second : NullDescriptor();
}
private:
// Set memory descriptor
bool SetMemDescr(const uint32_t& id, const uint32_t& size, const bool& local) {
const mem_descr_t des = {NULL, size, local};
auto ret = mem_map_.insert(mem_map_t::value_type(id, des));
return ret.second;
}
// Kernel memory map object
mem_map_t mem_map_;
};
#endif // TEST_CTRL_TEST_KERNEL_H_
test/ctrl/test_pgen.h
+45
View File
@@ -0,0 +1,45 @@
/******************************************************************************
Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
Redistribution and use in source and binary forms, with or without modification,
are permitted provided that the following conditions are met:
Redistributions of source code must retain the above copyright notice, this list
of conditions and the following disclaimer.
Redistributions in binary form must reproduce the above copyright notice, this
list of conditions and the following disclaimer in the documentation and/or
other materials provided with the distribution.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#ifndef TEST_CTRL_TEST_PGEN_H_
#define TEST_CTRL_TEST_PGEN_H_
#include "ctrl/test_pmgr.h"
// SimpleConvolution: Class implements OpenCL SimpleConvolution sample
class TestPGen : public TestPMgr {
protected:
typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
packet_t* PrePacket() { return reinterpret_cast<packet_t*>(&pre_packet_); }
packet_t* PostPacket() { return reinterpret_cast<packet_t*>(&post_packet_); }
public:
explicit TestPGen(TestAql* t) : TestPMgr(t) {}
};
#endif // TEST_CTRL_TEST_PGEN_H_
test/ctrl/test_pgen_rocp.h
+78
View File
@@ -0,0 +1,78 @@
/******************************************************************************
Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
Redistribution and use in source and binary forms, with or without modification,
are permitted provided that the following conditions are met:
Redistributions of source code must retain the above copyright notice, this list
of conditions and the following disclaimer.
Redistributions in binary form must reproduce the above copyright notice, this
list of conditions and the following disclaimer in the documentation and/or
other materials provided with the distribution.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#ifndef TEST_CTRL_TEST_PGEN_ROCP_H_
#define TEST_CTRL_TEST_PGEN_ROCP_H_
#include <list>
#include <vector>
#include "ctrl/test_pgen.h"
#include "util/test_assert.h"
hsa_status_t TestPGenRocpCallback(hsa_ven_amd_aqlprofile_info_type_t info_type,
hsa_ven_amd_aqlprofile_info_data_t* info_data,
void* callback_data) {
hsa_status_t status = HSA_STATUS_SUCCESS;
typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> passed_data_t;
reinterpret_cast<passed_data_t*>(callback_data)->push_back(*info_data);
return status;
}
// Class implements PMC profiling
class TestPGenRocp : public TestPGen {
public:
explicit TestPGenRocp(TestAql* t) : TestPGen(t) { std::clog << "Test: PGen ROCP" << std::endl; }
bool Initialize(int /*arg_cnt*/, char** /*arg_list*/) {
status = rocprofiler_on_dispatch(&profile_, PrePacket(), PostPacket());
TEST_STATUS(status != HSA_STATUS_SUCCESS);
return (status == HSA_STATUS_SUCCESS);
}
private:
bool BuildPackets() { return true; }
bool DumpData() {
std::clog << "TestPGenRocp::DumpData :" << std::endl;
typedef std::vector<hsa_ven_amd_aqlprofile_info_data_t> callback_data_t;
callback_data_t data;
api_.hsa_ven_amd_aqlprofile_iterate_data(&profile_, TestPGenRocpCallback, &data);
for (callback_data_t::iterator it = data.begin(); it != data.end(); ++it) {
std::cout << std::dec << "event(block(" << it->pmc_data.event.block_name << "_"
<< it->pmc_data.event.block_index << "), id(" << it->pmc_data.event.counter_id
<< ")), sample(" << it->sample_id << "), result(" << it->pmc_data.result << ")"
<< std::endl;
}
return true;
}
};
#endif // TEST_CTRL_TEST_PGEN_ROCP_H_
test/ctrl/test_pmgr.cpp
+144
View File
@@ -0,0 +1,144 @@
/******************************************************************************
Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
Redistribution and use in source and binary forms, with or without modification,
are permitted provided that the following conditions are met:
Redistributions of source code must retain the above copyright notice, this list
of conditions and the following disclaimer.
Redistributions in binary form must reproduce the above copyright notice, this
list of conditions and the following disclaimer in the documentation and/or
other materials provided with the distribution.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#include "ctrl/test_pmgr.h"
#include <atomic>
#include "ctrl/test_assert.h"
bool TestPMgr::AddPacketGfx9(const packet_t* packet) {
packet_t aql_packet = *packet;
// Compute the write index of queue and copy Aql packet into it
uint64_t que_idx = hsa_queue_load_write_index_relaxed(GetQueue());
const uint32_t mask = GetQueue()->size - 1;
packet_t* slot = (reinterpret_cast<packet_t*>(GetQueue()->base_address)) + (que_idx & mask);
// Disable packet so that submission to HW is complete
const auto header = HSA_PACKET_TYPE_VENDOR_SPECIFIC << HSA_PACKET_HEADER_TYPE;
aql_packet.header &= (~((1ul << HSA_PACKET_HEADER_WIDTH_TYPE) - 1)) << HSA_PACKET_HEADER_TYPE;
aql_packet.header |= HSA_PACKET_TYPE_INVALID << HSA_PACKET_HEADER_TYPE;
// Copy Aql packet into queue buffer
*slot = aql_packet;
// After AQL packet is fully copied into queue buffer
// update packet header from invalid state to valid state
auto header_atomic_ptr =
reinterpret_cast<std::atomic<uint16_t>*>(&slot->header);
header_atomic_ptr->store(header, std::memory_order_release);
// Increment the write index and ring the doorbell to dispatch the kernel.
hsa_queue_store_write_index_relaxed(GetQueue(), (que_idx + 1));
hsa_signal_store_relaxed(GetQueue()->doorbell_signal, que_idx);
return true;
}
bool TestPMgr::AddPacketGfx8(const packet_t* packet) {
// Create legacy devices PM4 data
const hsa_ext_amd_aql_pm4_packet_t* aql_packet = (const hsa_ext_amd_aql_pm4_packet_t*)packet;
slot_pm4_t data;
api_.hsa_ven_amd_aqlprofile_legacy_get_pm4(aql_packet, reinterpret_cast<void*>(data.words));
// Compute the write index of queue and copy Aql packet into it
uint64_t que_idx = hsa_queue_load_write_index_relaxed(GetQueue());
const uint32_t mask = GetQueue()->size - 1;
// Copy Aql/Pm4 blob into queue buffer
packet_t* ptr = (reinterpret_cast<packet_t*>(GetQueue()->base_address)) + (que_idx & mask);
slot_pm4_t* slot = reinterpret_cast<slot_pm4_t*>(ptr);
for (unsigned i = 1; i < SLOT_PM4_SIZE_DW; ++i) {
slot->words[i] = data.words[i];
}
// To maintain global order to ensure the prior copy of the packet contents is made visible
// before the header is updated.
// With in-order CP it will wait until the first packet in the blob will be valid
std::atomic<uint32_t>* header_atomic_ptr =
reinterpret_cast<std::atomic<uint32_t>*>(&slot->words[0]);
header_atomic_ptr->store(data.words[0], std::memory_order_release);
// Increment the write index and ring the doorbell to dispatch the kernel.
que_idx += SLOT_PM4_SIZE_AQLP - 1;
hsa_queue_store_write_index_relaxed(GetQueue(), (que_idx + 1));
hsa_signal_store_relaxed(GetQueue()->doorbell_signal, que_idx);
return true;
}
bool TestPMgr::AddPacket(const packet_t* packet) {
const char* agent_name = GetAgentInfo()->name;
return (strncmp(agent_name, "gfx8", 4) == 0) ? AddPacketGfx8(packet) : AddPacketGfx9(packet);
}
bool TestPMgr::Run() {
// Build Aql Pkts
const bool active = BuildPackets();
if (active) {
// Submit Pre-Dispatch Aql packet
AddPacket(&pre_packet_);
}
Test()->Run();
if (active) {
// Set post packet completion signal
post_packet_.completion_signal = post_signal_;
// Submit Post-Dispatch Aql packet
AddPacket(&post_packet_);
// Wait for Post-Dispatch packet to complete
hsa_signal_wait_acquire(post_signal_, HSA_SIGNAL_CONDITION_LT, 1, (uint64_t)-1,
HSA_WAIT_STATE_BLOCKED);
// Dumping profiling data
DumpData();
}
return true;
}
bool TestPMgr::Initialize(int argc, char** argv) {
TestAql::Initialize(argc, argv);
hsa_status_t status = HSA_STATUS_ERROR;
status = hsa_signal_create(1, 0, NULL, &post_signal_);
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
status = hsa_system_get_extension_table(HSA_EXTENSION_AMD_AQLPROFILE, 1, 0, &api_);
TEST_ASSERT(status == HSA_STATUS_SUCCESS);
return true;
}
TestPMgr::TestPMgr(TestAql* t) : TestAql(t), api_({0}) {
memset(&pre_packet_, 0, sizeof(pre_packet_));
memset(&post_packet_, 0, sizeof(post_packet_));
dummy_signal_.handle = 0;
post_signal_ = dummy_signal_;
memset(&api_, 0, sizeof(api_));
}
test/ctrl/test_pmgr.h
+70
View File
@@ -0,0 +1,70 @@
/******************************************************************************
Copyright ©2013 Advanced Micro Devices, Inc. All rights reserved.
Redistribution and use in source and binary forms, with or without modification,
are permitted provided that the following conditions are met:
Redistributions of source code must retain the above copyright notice, this list
of conditions and the following disclaimer.
Redistributions in binary form must reproduce the above copyright notice, this
list of conditions and the following disclaimer in the documentation and/or
other materials provided with the distribution.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED.
IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT,
INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE,
DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF
LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE
OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED
OF THE POSSIBILITY OF SUCH DAMAGE.
*******************************************************************************/
#ifndef TEST_CTRL_TEST_PMGR_H_
#define TEST_CTRL_TEST_PMGR_H_
#include <hsa.h>
#include <hsa_ven_amd_aqlprofile.h>
#include <atomic>
#include "ctrl/test_aql.h"
// Class implements profiling manager
class TestPMgr : public TestAql {
public:
typedef hsa_ext_amd_aql_pm4_packet_t packet_t;
explicit TestPMgr(TestAql* t);
bool Run();
protected:
packet_t pre_packet_;
packet_t post_packet_;
hsa_signal_t dummy_signal_;
hsa_signal_t post_signal_;
hsa_ven_amd_aqlprofile_1_00_pfn_t api_;
virtual bool BuildPackets() { return false; }
virtual bool DumpData() { return false; }
virtual bool Initialize(int argc, char** argv);
private:
enum {
SLOT_PM4_SIZE_DW = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(uint32_t),
SLOT_PM4_SIZE_AQLP = HSA_VEN_AMD_AQLPROFILE_LEGACY_PM4_PACKET_SIZE / sizeof(packet_t)
};
struct slot_pm4_t {
uint32_t words[SLOT_PM4_SIZE_DW];
};
bool AddPacket(const packet_t* packet);
bool AddPacketGfx8(const packet_t* packet);
bool AddPacketGfx9(const packet_t* packet);
};
#endif // TEST_CTRL_TEST_PMGR_H_
test/ctrl/thr_tool.cpp
+297
View File
@@ -0,0 +1,297 @@
#include <hsa.h>
#include <pthread.h>
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <map>
#include <vector>
#include "inc/rocprofiler.h"
#include "util/xml.h"
#define PUBLIC_API __attribute__((visibility("default")))
#define CONSTRUCTOR_API __attribute__((constructor))
#define DESTRUCTOR_API __attribute__((destructor))
// Tool thread
pthread_t thread;
pthread_attr_t thr_attr;
bool thr_stop = false;
struct dispatch_data_t {
rocprofiler_info_t* info;
unsigned info_count;
unsigned group_index;
};
struct context_entry_t {
rocprofiler_group_t* group;
rocprofiler_info_t* info;
unsigned info_count;
rocprofiler_callback_data_t data;
};
pthread_mutex_t mutex = PTHREAD_MUTEX_INITIALIZER;
unsigned context_array_size = 1;
context_entry_t* context_array = NULL;
unsigned context_array_index = 0;
const char* file_name;
FILE* file_handle = NULL;
void check_status(hsa_status_t status) {
if (status != HSA_STATUS_SUCCESS) {
const char* error_string = NULL;
rocprofiler_error_string(&error_string);
fprintf(stderr, "ERROR: %s\n", error_string);
exit(1);
}
}
unsigned align_size(unsigned size, unsigned alignment) { return ((size + alignment - 1) & ~(alignment - 1)); }
void print_info(FILE* file, const rocprofiler_info_t* info, const unsigned info_count, const char* str) {
if (str) fprintf(file, "%s:\n", str);
for (unsigned i= 0; i < info_count; ++i) {
const rocprofiler_info_t* p = &info[i];
fprintf(file, " %s ", p->name);
switch (p->data.kind) {
case ROCPROFILER_INT64:
fprintf(file, "(%lu)\n", p->data.result64);
break;
case ROCPROFILER_BYTES: {
fprintf(file, "(\n");
const char* ptr = reinterpret_cast<const char*>(p->data.result_bytes.ptr);
uint64_t size = 0;
for (unsigned i = 0; i < p->data.result_bytes.instance_count; ++i) {
size = *reinterpret_cast<const uint64_t*>(ptr);
const char* data = ptr + sizeof(size);
fprintf(file, " data (%p), size (%lu)\n", data, size);
size = align_size(size, sizeof(uint64_t));
ptr = data + size;
}
fprintf(file, " )\n");
break;
}
default:
std::cout << "Bad result kind (" << p->data.kind << ")" << std::endl;
}
}
}
void print_group(FILE* file, const rocprofiler_group_t* group, const char* str) {
if (str) fprintf(file, "%s:\n", str);
for (unsigned i= 0; i < group->info_count; ++i) {
print_info(file, group->info[i], 1, NULL);
}
}
void store_context(context_entry_t context_entry) {
if(pthread_mutex_lock(&mutex) != 0) {
perror("pthread_mutex_lock");
exit(1);
}
if ((context_array == NULL) || (context_array_index >= context_array_size)) {
context_array_size *= 2;
context_array = reinterpret_cast<context_entry_t*>(realloc(context_array, context_array_size * sizeof(context_entry_t)));
}
context_array_index += 1;
context_array[context_array_index - 1] = context_entry;
if(pthread_mutex_unlock(&mutex) != 0) {
perror("pthread_mutex_unlock");
exit(1);
}
}
void dump_context(FILE *file, unsigned index) {
hsa_status_t status = HSA_STATUS_ERROR;
if (pthread_mutex_lock(&mutex) != 0) {
perror("pthread_mutex_lock");
exit(1);
}
context_entry_t* entry = &context_array[index];
rocprofiler_group_t* group = entry->group;
const rocprofiler_info_t* info = entry->info;
const unsigned info_count = entry->info_count;
fprintf(file, "Dispatch[%u], kernel_object(0x%lx):\n", index, entry->data.kernel_object);
if (pthread_mutex_unlock(&mutex) != 0) {
perror("pthread_mutex_unlock");
exit(1);
}
status = rocprofiler_get_group_data(group);
check_status(status);
//print_group(file, group, "Group[0] data");
status = rocprofiler_get_metrics_data(group->context);
check_status(status);
print_info(file, info, info_count, NULL);
// Finishing cleanup
// Deleting profiling context will delete all allocated resources
rocprofiler_close(group->context);
}
// Provided standard profiling callback
hsa_status_t dispatch_callback(
const rocprofiler_callback_data_t* callback_data,
void* user_data,
rocprofiler_group_t** group) {
hsa_status_t status = HSA_STATUS_ERROR;
// Passed tool data
dispatch_data_t* tool_data = reinterpret_cast<dispatch_data_t*>(user_data);
// Profiling context
rocprofiler_t* context = NULL;
// Open profiling context
status = rocprofiler_open(0, tool_data->info, tool_data->info_count, &context, 0, NULL);
check_status(status);
rocprofiler_group_t* groups = NULL;
uint32_t group_count = 0;
status = rocprofiler_get_groups(context, &groups, &group_count);
check_status(status);
assert(group_count == 1);
*group = &groups[0];
store_context({*group, tool_data->info, tool_data->info_count, *callback_data});
return status;
}
void* dumping_data(void*) {
unsigned index = 0;
do {
while (index < context_array_index) {
dump_context(file_handle, index);
++index;
}
} while (!thr_stop);
return NULL;
}
CONSTRUCTOR_API void constructor() {
std::map<std::string, hsa_ven_amd_aqlprofile_parameter_name_t> parameters_dict;
parameters_dict["HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_COMPUTE_UNIT_TARGET"] = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_COMPUTE_UNIT_TARGET;
parameters_dict["HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_VM_ID_MASK"] = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_VM_ID_MASK;
parameters_dict["HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_MASK"] = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_MASK;
parameters_dict["HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK"] = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK;
parameters_dict["HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK2"] = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK2;
#ifdef TOOL_THREAD
int err = pthread_attr_init(&thr_attr);
if (err) { errno = err; perror("pthread_attr_init"); exit(1); }
err = pthread_create(&thread, &thr_attr, dumping_data, NULL);
if (err) { errno = err; perror("pthread_create"); exit(1); }
#endif
// Set output file
file_name = getenv("ROCP_OUTPUT");
if (file_name != NULL) {
file_handle = fopen(file_name, "w");
if (file_handle == NULL) {
perror("fopen");
exit(1);
}
} else file_handle = stdout;
// Getting input
const char* xml_name = getenv("ROCP_INPUT");
if (xml_name == NULL) {
fprintf(stderr, "ROCProfiler: input is not specified, ROCP_INPUT env");
exit(1);
}
printf("ROCProfiler: input from \"%s\"\n", xml_name);
xml::Xml* xml = new xml::Xml(xml_name);
// Getting metrics
auto metrics_list = xml->GetNodes("top.metric");
std::vector<std::string> metrics_vec;
for (auto* entry : metrics_list) {
const std::string entry_str = entry->opts["name"];
size_t pos1 = 0;
while(pos1 < entry_str.length()) {
const size_t pos2 = entry_str.find(",", pos1);
const std::string metric_name = entry_str.substr(pos1, pos2 - pos1);
metrics_vec.push_back(metric_name);
if (pos2 == std::string::npos) break;
pos1 = pos2 + 1;
}
}
// Getting traces
auto traces_list = xml->GetNodes("top.trace");
const unsigned info_count = metrics_vec.size() + traces_list.size();
rocprofiler_info_t* info= new rocprofiler_info_t[info_count];
memset(info, 0, info_count * sizeof(rocprofiler_info_t));
printf(" %d metrics\n", (int) metrics_vec.size());
for (unsigned i = 0; i < metrics_vec.size(); ++i) {
const std::string& name = metrics_vec[i];
printf("%s%s", (i == 0) ? " " : ", ", name.c_str());
info[i] = {};
info[i].type = ROCPROFILER_TYPE_METRIC;
info[i].name = strdup(name.c_str());
}
if (metrics_vec.size()) printf("\n");
printf(" %d traces\n", (int) traces_list.size());
unsigned index = metrics_vec.size();
for (auto* entry : traces_list) {
auto params_list = xml->GetNodes("top.trace.parameters");
if (params_list.size() != 1) {
fprintf(stderr, "ROCProfiler: Single input 'parameters' section is supported\n");
exit(1);
}
const std::string& name = entry->opts["name"];
printf(" %s (\n", name.c_str());
info[index] = {};
info[index].type = ROCPROFILER_TYPE_TRACE;
info[index].name = strdup(name.c_str());
for (auto* params : params_list) {
const unsigned parameter_count = params->opts.size();
rocprofiler_parameter_t *parameters = new rocprofiler_parameter_t[parameter_count];
unsigned p_index = 0;
for (auto& v : params->opts) {
const std::string parameter_name = v.first;
if (parameters_dict.find(parameter_name) == parameters_dict.end()) {
fprintf(stderr, "ROCProfiler: unknown trace parameter %s\n", parameter_name.c_str());
exit(1);
}
const uint32_t value = strtol(v.second.c_str(), NULL, 0);
printf(" %s = 0x%x\n", parameter_name.c_str(), value);
parameters[p_index] = {};
parameters[p_index].parameter_name = parameters_dict[parameter_name];
parameters[p_index].value = value;
++p_index;
}
info[index].parameters = parameters;
info[index].parameter_count = parameter_count;
}
printf(" )\n");
++index;
}
if (info_count) {
// Adding dispatch observer
dispatch_data_t* dispatch_data = new dispatch_data_t{};
dispatch_data->info = info;
dispatch_data->info_count = info_count;
dispatch_data->group_index = 0;
rocprofiler_dispatch_observer(dispatch_callback, dispatch_data);
}
}
DESTRUCTOR_API void destructor() {
printf("\nROCPRofiler: %u contexts collected", context_array_index);
thr_stop = true;
#ifdef TOOL_THREAD
pthread_join(thread, NULL);
#else
dumping_data(NULL);
#endif
}
test/ctrl/tool.cpp
+313
View File
@@ -0,0 +1,313 @@
#include <hsa.h>
#include <pthread.h>
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
#include <map>
#include <vector>
#include "inc/rocprofiler.h"
#include "util/xml.h"
#define PUBLIC_API __attribute__((visibility("default")))
#define CONSTRUCTOR_API __attribute__((constructor))
#define DESTRUCTOR_API __attribute__((destructor))
struct dispatch_data_t {
rocprofiler_info_t* info;
unsigned info_count;
unsigned group_index;
};
struct context_entry_t {
rocprofiler_group_t* group;
rocprofiler_info_t* info;
unsigned info_count;
rocprofiler_callback_data_t data;
};
pthread_mutex_t mutex = PTHREAD_MUTEX_INITIALIZER;
unsigned context_array_size = 1;
context_entry_t* context_array = NULL;
unsigned context_array_index = 0;
unsigned dump_index = 0;
const char* file_name = NULL;
FILE* file_handle = NULL;
void check_status(hsa_status_t status) {
if (status != HSA_STATUS_SUCCESS) {
const char* error_string = NULL;
rocprofiler_error_string(&error_string);
fprintf(stderr, "ERROR: %s\n", error_string);
exit(1);
}
}
hsa_status_t trace_data_cb(
hsa_ven_amd_aqlprofile_info_type_t info_type,
hsa_ven_amd_aqlprofile_info_data_t* info_data,
void* data)
{
hsa_status_t status = HSA_STATUS_SUCCESS;
if (info_type == HSA_VEN_AMD_AQLPROFILE_INFO_SQTT_DATA) {
printf(" data ptr (%p), size(%u)\n", info_data->sqtt_data.ptr, info_data->sqtt_data.size);
} else status = HSA_STATUS_ERROR;
return status;
}
unsigned align_size(unsigned size, unsigned alignment) { return ((size + alignment - 1) & ~(alignment - 1)); }
void print_info(FILE* file, const rocprofiler_info_t* info, const unsigned info_count, rocprofiler_t* context, const char* str) {
if (str) fprintf(file, "%s:\n", str);
for (unsigned i= 0; i < info_count; ++i) {
const rocprofiler_info_t* p = &info[i];
fprintf(file, " %s ", p->name);
switch (p->data.kind) {
case ROCPROFILER_INT64:
fprintf(file, "(%lu)\n", p->data.result_int64);
break;
case ROCPROFILER_BYTES: {
fprintf(file, "(\n");
if (p->data.result_bytes.copy) {
fprintf(file, " system memory copy\n");
const char* ptr = reinterpret_cast<const char*>(p->data.result_bytes.ptr);
uint64_t size = 0;
for (unsigned i = 0; i < p->data.result_bytes.instance_count; ++i) {
size = *reinterpret_cast<const uint64_t*>(ptr);
const char* data = ptr + sizeof(size);
fprintf(file, " data (%p), size (%lu)\n", data, size);
size = align_size(size, sizeof(uint64_t));
ptr = data + size;
}
} else {
fprintf(file, " local memory buffer\n");
rocprofiler_iterate_trace_data(context, trace_data_cb, NULL);
}
fprintf(file, " )\n");
break;
}
default:
std::cout << "Bad result kind (" << p->data.kind << ")" << std::endl;
}
}
}
void print_group(FILE* file, const rocprofiler_group_t* group, const char* str) {
if (str) fprintf(file, "%s:\n", str);
for (unsigned i= 0; i < group->info_count; ++i) {
print_info(file, group->info[i], 1, group->context, NULL);
}
}
void store_entry(const context_entry_t& context_entry) {
if(pthread_mutex_lock(&mutex) != 0) {
perror("pthread_mutex_lock");
exit(1);
}
if ((context_array == NULL) || (context_array_index >= context_array_size)) {
context_array_size *= 2;
context_array = reinterpret_cast<context_entry_t*>(realloc(context_array, context_array_size * sizeof(context_entry_t)));
}
context_array[context_array_index] = context_entry;
context_array_index += 1;
if (pthread_mutex_unlock(&mutex) != 0) {
perror("pthread_mutex_unlock");
exit(1);
}
}
void dump_context(FILE *file, context_entry_t* entry, unsigned index) {
hsa_status_t status = HSA_STATUS_ERROR;
rocprofiler_group_t* group = entry->group;
const rocprofiler_info_t* info = entry->info;
const unsigned info_count = entry->info_count;
fprintf(file, "Dispatch[%u], kernel_object(0x%lx):\n", index, entry->data.kernel_object);
status = rocprofiler_get_group_data(group);
check_status(status);
//print_group(file, group, "Group[0] data");
status = rocprofiler_get_metrics_data(group->context);
check_status(status);
print_info(file, info, info_count, group->context, NULL);
// Finishing cleanup
// Deleting profiling context will delete all allocated resources
rocprofiler_close(group->context);
dump_index = index;
}
void dumping_data() {
if (pthread_mutex_lock(&mutex) != 0) {
perror("pthread_mutex_lock");
exit(1);
}
for (unsigned index = 0; index < context_array_index; ++index) {
dump_context(file_handle, &context_array[index], index);
}
if (pthread_mutex_unlock(&mutex) != 0) {
perror("pthread_mutex_unlock");
exit(1);
}
}
// profiling callback
hsa_status_t dispatch_callback(
const rocprofiler_callback_data_t* callback_data,
void* user_data,
rocprofiler_group_t** group) {
hsa_status_t status = HSA_STATUS_ERROR;
// Passed tool data
dispatch_data_t* tool_data = reinterpret_cast<dispatch_data_t*>(user_data);
// Profiling context
rocprofiler_t* context = NULL;
// context properties
rocprofiler_properties_t properties{};
// Open profiling context
status = rocprofiler_open(0, tool_data->info, tool_data->info_count, &context, 0, &properties);
check_status(status);
rocprofiler_group_t* groups = NULL;
uint32_t group_count = 0;
status = rocprofiler_get_groups(context, &groups, &group_count);
check_status(status);
assert(group_count == 1);
*group = &groups[0];
context_entry_t entry;
entry.group = *group;
entry.info = tool_data->info;
entry.info_count = tool_data->info_count;
entry.data = *callback_data;
store_entry(entry);
return status;
}
CONSTRUCTOR_API void constructor() {
std::map<std::string, hsa_ven_amd_aqlprofile_parameter_name_t> parameters_dict;
parameters_dict["HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_COMPUTE_UNIT_TARGET"] = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_COMPUTE_UNIT_TARGET;
parameters_dict["HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_VM_ID_MASK"] = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_VM_ID_MASK;
parameters_dict["HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_MASK"] = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_MASK;
parameters_dict["HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK"] = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK;
parameters_dict["HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK2"] = HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK2;
// Set output file
file_name = getenv("ROCP_OUTPUT");
if (file_name != NULL) {
file_handle = fopen(file_name, "w");
if (file_handle == NULL) {
perror("fopen");
exit(1);
}
} else file_handle = stdout;
// Getting input
const char* xml_name = getenv("ROCP_INPUT");
if (xml_name == NULL) {
fprintf(stderr, "ROCProfiler: input is not specified, ROCP_INPUT env");
exit(1);
}
printf("ROCProfiler: input from \"%s\"\n", xml_name);
xml::Xml* xml = new xml::Xml(xml_name);
// Getting metrics
auto metrics_list = xml->GetNodes("top.metric");
std::vector<std::string> metrics_vec;
for (auto* entry : metrics_list) {
const std::string entry_str = entry->opts["name"];
size_t pos1 = 0;
while(pos1 < entry_str.length()) {
const size_t pos2 = entry_str.find(",", pos1);
const std::string metric_name = entry_str.substr(pos1, pos2 - pos1);
metrics_vec.push_back(metric_name);
if (pos2 == std::string::npos) break;
pos1 = pos2 + 1;
}
}
// Getting traces
auto traces_list = xml->GetNodes("top.trace");
const unsigned info_count = metrics_vec.size() + traces_list.size();
rocprofiler_info_t* info= new rocprofiler_info_t[info_count];
memset(info, 0, info_count * sizeof(rocprofiler_info_t));
printf(" %d metrics\n", (int) metrics_vec.size());
for (unsigned i = 0; i < metrics_vec.size(); ++i) {
const std::string& name = metrics_vec[i];
printf("%s%s", (i == 0) ? " " : ", ", name.c_str());
info[i] = {};
info[i].type = ROCPROFILER_TYPE_METRIC;
info[i].name = strdup(name.c_str());
}
if (metrics_vec.size()) printf("\n");
printf(" %d traces\n", (int) traces_list.size());
unsigned index = metrics_vec.size();
for (auto* entry : traces_list) {
auto params_list = xml->GetNodes("top.trace.parameters");
if (params_list.size() != 1) {
fprintf(stderr, "ROCProfiler: Single input 'parameters' section is supported\n");
exit(1);
}
const std::string& name = entry->opts["name"];
const bool to_copy_data = (entry->opts["copy"] == "true");
printf(" %s (\n", name.c_str());
info[index] = {};
info[index].type = ROCPROFILER_TYPE_TRACE;
info[index].name = strdup(name.c_str());
info[index].data.result_bytes.copy = to_copy_data;
for (auto* params : params_list) {
const unsigned parameter_count = params->opts.size();
rocprofiler_parameter_t *parameters = new rocprofiler_parameter_t[parameter_count];
unsigned p_index = 0;
for (auto& v : params->opts) {
const std::string parameter_name = v.first;
if (parameters_dict.find(parameter_name) == parameters_dict.end()) {
fprintf(stderr, "ROCProfiler: unknown trace parameter %s\n", parameter_name.c_str());
exit(1);
}
const uint32_t value = strtol(v.second.c_str(), NULL, 0);
printf(" %s = 0x%x\n", parameter_name.c_str(), value);
parameters[p_index] = {};
parameters[p_index].parameter_name = parameters_dict[parameter_name];
parameters[p_index].value = value;
++p_index;
}
info[index].parameters = parameters;
info[index].parameter_count = parameter_count;
}
printf(" )\n");
++index;
}
if (info_count) {
// Adding dispatch observer
dispatch_data_t* dispatch_data = new dispatch_data_t{};
dispatch_data->info = info;
dispatch_data->info_count = info_count;
dispatch_data->group_index = 0;
rocprofiler_set_dispatch_observer(dispatch_callback, dispatch_data);
}
}
DESTRUCTOR_API void destructor() {
printf("\nROCPRofiler: %u contexts collected", context_array_index);
if (file_name == NULL) {
printf("\n");
} else {
printf(", dumping to %s\n", file_name);
}
dumping_data();
}