ROC profiler prototype sources importing
[ROCm/rocprofiler commit: 85278f08a0]
This commit is contained in:
Evgeny
@@ -0,0 +1,41 @@
|
||||
cmake_minimum_required ( VERSION 3.5.0 )
|
||||
set ( CMAKE_VERBOSE_MAKEFILE TRUE CACHE BOOL "Verbose Output" FORCE )
|
||||
|
||||
set ( EXE_NAME "ctrl" )
|
||||
|
||||
if ( NOT DEFINED TEST_DIR )
|
||||
set ( TEST_DIR ${CMAKE_CURRENT_SOURCE_DIR} )
|
||||
project ( ${EXE_NAME} )
|
||||
## Set build environment
|
||||
include ( env )
|
||||
set ( ROCPROFILER_TARGET "rocprofiler64" )
|
||||
endif ()
|
||||
|
||||
## Util sources
|
||||
file( GLOB UTIL_SRC "${TEST_DIR}/util/*.cpp" )
|
||||
|
||||
## Test control sources
|
||||
set ( CTRL_SRC
|
||||
${TEST_DIR}/ctrl/test.cpp
|
||||
${TEST_DIR}/ctrl/test_hsa.cpp
|
||||
)
|
||||
|
||||
## Test kernels sources
|
||||
set ( TEST_NAME simple_convolution )
|
||||
set ( KERN_SRC ${TEST_DIR}/${TEST_NAME}/${TEST_NAME}.cpp )
|
||||
execute_process ( COMMAND sh -xc "cp ${TEST_DIR}/${TEST_NAME}/*.hsaco ${PROJECT_BINARY_DIR}" )
|
||||
|
||||
## Building test executable
|
||||
add_executable ( ${EXE_NAME} ${KERN_SRC} ${CTRL_SRC} ${UTIL_SRC} )
|
||||
target_include_directories ( ${EXE_NAME} PRIVATE ${TEST_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH} )
|
||||
target_link_libraries( ${EXE_NAME} ${ROCPROFILER_TARGET} ${HSA_RUNTIME_LIB} c stdc++ dl pthread rt atomic )
|
||||
execute_process ( COMMAND sh -xc "cp ${TEST_DIR}/run.sh ${PROJECT_BINARY_DIR}" )
|
||||
execute_process ( COMMAND sh -xc "cp ${TEST_DIR}/*.xml ${PROJECT_BINARY_DIR}" )
|
||||
|
||||
## Build test library
|
||||
set ( TEST_LIB "tool" )
|
||||
set ( TEST_LIB_SRC ${TEST_DIR}/ctrl/tool.cpp )
|
||||
add_library ( ${TEST_LIB} SHARED ${TEST_LIB_SRC} )
|
||||
target_include_directories ( ${TEST_LIB} PRIVATE ${TEST_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH} )
|
||||
target_link_libraries( ${TEST_LIB} ${ROCPROFILER_TARGET} ${HSA_RUNTIME_LIB} c stdc++ dl pthread rt atomic )
|
||||
|
||||
@@ -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_
|
||||
@@ -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;
|
||||
}
|
||||
@@ -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;
|
||||
}
|
||||
@@ -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_
|
||||
@@ -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; }
|
||||
@@ -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_
|
||||
@@ -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_
|
||||
@@ -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_
|
||||
@@ -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_
|
||||
@@ -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_));
|
||||
}
|
||||
@@ -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_
|
||||
@@ -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
|
||||
}
|
||||
@@ -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();
|
||||
}
|
||||
@@ -0,0 +1,7 @@
|
||||
<metric name=CPC_ME1_STALL_WAIT_ON_RCIU_READ,SQ_WAVES,SQ_WAVE_READY,SQ_CYCLES,SQ_ITEMS,WAVE_STALLS_RATE ></metric>
|
||||
<trace name=SQTT copy=true >
|
||||
<parameters
|
||||
HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_MASK=0xf
|
||||
HSA_VEN_AMD_AQLPROFILE_PARAMETER_NAME_TOKEN_MASK=0xf
|
||||
></parameters>
|
||||
</trace>
|
||||
@@ -0,0 +1,15 @@
|
||||
<gfx8>
|
||||
<metric name=SQ_CYCLES block=SQ event=2 ></metric>
|
||||
<metric name=SQ_WAVES block=SQ event=4 ></metric>
|
||||
<metric name=SQ_ITEMS block=SQ event=14 ></metric>
|
||||
<metric name=SQ_WAVE_READY block=SQ event=47 ></metric>
|
||||
<metric name=TCC_CYCLE block=TCC event=1 ></metric>
|
||||
<metric name=TCC_REQ block=TCC event=3 ></metric>
|
||||
<metric name=TCC_WRITEBACK block=TCC event=22 ></metric>
|
||||
<metric name=CPC_ALWAYS_COUNT block=CPC event=0 ></metric>
|
||||
<metric name=CPC_ME1_STALL_WAIT_ON_RCIU_READ block=CPC event=8 ></metric>
|
||||
</gfx8>
|
||||
|
||||
<global>
|
||||
<metric name=WAVE_STALLS_RATE expr=CPC_ME1_STALL_WAIT_ON_RCIU_READ*(SQ_WAVES+SQ_WAVE_READY)*100/(SQ_CYCLES/SQ_ITEMS) ></metric>
|
||||
</global>
|
||||
Executable
+27
@@ -0,0 +1,27 @@
|
||||
#!/bin/sh
|
||||
|
||||
tbin=./test/ctrl
|
||||
|
||||
#export HSA_LIB=/home/evgeny/pkg/compute-psdb-16453/lib
|
||||
export HSA_LIB=/home/evgeny/git/compute/out/ubuntu-16.04/16.04/lib
|
||||
export OCL_LIB=/home/evgeny/pkg/opencl_modified/opencl_x86_64/lib
|
||||
#export OCL_LIB=/home/evgeny/Perforce/eshcherb_opencl/drivers/opencl/dist/linux/debug/lib/x86_64
|
||||
|
||||
export LD_LIBRARY_PATH=$PWD:$HSA_LIB:$OCL_LIB
|
||||
export ROCPROFILER_LOG=1
|
||||
|
||||
export HSA_TOOLS_LIB=librocprofiler64.so
|
||||
export ROCP_TOOL_LIB=test/libtool.so
|
||||
export ROCP_HSA_INTERCEPT=1
|
||||
export ROCP_METRICS=metrics.xml
|
||||
export ROCP_INPUT=input.xml
|
||||
unset ROCP_PROXY_QUEUE
|
||||
|
||||
echo "Run simple profiling test"
|
||||
if [ -n "$1" ] ; then
|
||||
eval "$*"
|
||||
else
|
||||
eval $tbin
|
||||
fi
|
||||
|
||||
exit 0
|
||||
Binary file not shown.
Binary file not shown.
@@ -0,0 +1,81 @@
|
||||
/******************************************************************************
|
||||
|
||||
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.
|
||||
|
||||
********************************************************************************/
|
||||
|
||||
/**
|
||||
* SimpleConvolution is where each pixel of the output image
|
||||
* is the weighted sum of the neighborhood pixels of the input image
|
||||
* The neighborhood is defined by the dimensions of the mask and
|
||||
* weight of each neighbor is defined by the mask itself.
|
||||
* @param output Output matrix after performing convolution
|
||||
* @param input Input matrix on which convolution is to be performed
|
||||
* @param mask mask matrix using which convolution was to be performed
|
||||
* @param inputDimensions dimensions of the input matrix
|
||||
* @param maskDimensions dimensions of the mask matrix
|
||||
*/
|
||||
__kernel void SimpleConvolution(__global uint * output,
|
||||
__global uint * input,
|
||||
__global float * mask,
|
||||
const uint2 inputDimensions,
|
||||
const uint2 maskDimensions) {
|
||||
|
||||
uint tid = get_global_id(0);
|
||||
|
||||
uint width = inputDimensions.x;
|
||||
uint height = inputDimensions.y;
|
||||
|
||||
uint x = tid%width;
|
||||
uint y = tid/width;
|
||||
|
||||
uint maskWidth = maskDimensions.x;
|
||||
uint maskHeight = maskDimensions.y;
|
||||
|
||||
uint vstep = (maskWidth -1)/2;
|
||||
uint hstep = (maskHeight -1)/2;
|
||||
|
||||
// find the left, right, top and bottom indices such that
|
||||
// the indices do not go beyond image boundaires
|
||||
uint left = (x < vstep) ? 0 : (x - vstep);
|
||||
uint right = ((x + vstep) >= width) ? width - 1 : (x + vstep);
|
||||
uint top = (y < hstep) ? 0 : (y - hstep);
|
||||
uint bottom = ((y + hstep) >= height)? height - 1: (y + hstep);
|
||||
|
||||
// initializing wighted sum value
|
||||
float sumFX = 0;
|
||||
|
||||
for(uint i = left; i <= right; ++i) {
|
||||
for(uint j = top ; j <= bottom; ++j) {
|
||||
// performing wighted sum within the mask boundaries
|
||||
uint maskIndex = (j - (y - hstep)) * maskWidth + (i - (x - vstep));
|
||||
uint index = j * width + i;
|
||||
sumFX += ((float)input[index] * mask[maskIndex]);
|
||||
}
|
||||
}
|
||||
|
||||
// To round to the nearest integer
|
||||
sumFX += 0.5f;
|
||||
output[tid] = (uint)sumFX;
|
||||
}
|
||||
@@ -0,0 +1,390 @@
|
||||
/******************************************************************************
|
||||
|
||||
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 "simple_convolution/simple_convolution.h"
|
||||
|
||||
#include <string.h>
|
||||
#include <iostream>
|
||||
|
||||
#include "util/helper_funcs.h"
|
||||
#include "util/test_assert.h"
|
||||
|
||||
const uint32_t SimpleConvolution::input_data_[]{
|
||||
15, 201, 51, 89, 92, 34, 96, 66, 11, 225, 161, 96, 81, 211, 108, 124, 202, 244, 182,
|
||||
90, 215, 92, 98, 20, 44, 225, 55, 247, 202, 0, 45, 218, 202, 97, 51, 39, 131, 147,
|
||||
105, 143, 116, 11, 239, 198, 222, 92, 67, 169, 81, 250, 3, 40, 86, 101, 60, 131, 70,
|
||||
116, 123, 17, 117, 168, 236, 64, 10, 31, 103, 142, 179, 209, 29, 40, 220, 13, 239, 187,
|
||||
105, 50, 100, 186, 44, 104, 227, 131, 205, 32, 6, 20, 149, 130, 38, 10, 43, 18, 75,
|
||||
53, 50, 178, 195, 230, 132, 225, 14, 96, 238, 253, 27, 88, 48, 128, 18, 92, 232, 246,
|
||||
224, 182, 23, 231, 203, 172, 105, 241, 183, 148, 4, 2, 202, 55, 181, 142, 29, 57, 111,
|
||||
43, 153, 93, 41, 181, 181, 89, 54, 200, 182, 31, 190, 150, 213, 213, 126, 160, 130, 232,
|
||||
146, 57, 125, 151, 59, 71, 206, 240, 213, 236, 42, 68, 24, 195, 162, 65, 121, 87, 155,
|
||||
175, 31, 81, 207, 222, 232, 164, 180, 102, 69, 55, 79, 216, 112, 204, 112, 171, 19, 63,
|
||||
156, 233, 43, 198, 46, 67, 138, 208, 132, 4, 39, 32, 180, 71, 113, 131, 38, 90, 40,
|
||||
219, 193, 109, 18, 16, 70, 131, 220, 182, 46, 240, 245, 203, 217, 32, 146, 7, 100, 28,
|
||||
216, 233, 32, 255, 9, 213, 71, 123, 88, 110, 213, 128, 74, 150, 238, 93, 166, 52, 224,
|
||||
131, 234, 15, 115, 224, 218, 76, 1, 108, 84, 101, 137, 44, 79, 170, 44, 88, 127, 116,
|
||||
211, 216, 226, 168, 88, 45, 63, 70, 138, 230, 123, 107, 105, 101, 122, 220, 70, 84, 41,
|
||||
71, 193, 125, 173, 75, 169, 252, 245, 213, 84, 117, 73, 40, 77, 44, 209, 166, 90, 16,
|
||||
237, 229, 246, 104, 80, 95, 206, 202, 60, 20, 31, 101, 92, 225, 226, 9, 44, 140, 5,
|
||||
34, 97, 89, 151, 171, 129, 229, 216, 82, 139, 51, 99, 120, 24, 89, 225, 104, 185, 175,
|
||||
50, 246, 196, 82, 91, 32, 51, 62, 42, 96, 202, 47, 130, 44, 137, 26, 215, 10, 255,
|
||||
176, 93, 138, 227, 193, 3, 251, 27, 229, 100, 212, 149, 151, 202, 89, 233, 38, 122, 29,
|
||||
100, 164, 125, 46, 212, 0, 90, 93, 26, 50, 103, 25, 226, 197, 164, 198, 135, 168, 194,
|
||||
162, 141, 38, 119, 34, 190, 66, 124, 167, 104, 247, 197, 204, 156, 67, 251, 112, 67, 85,
|
||||
205, 93, 135, 53, 119, 106, 251, 28, 49, 130, 196, 243, 36, 82, 26, 155, 117, 216, 221,
|
||||
241, 128, 70, 233, 70, 18, 133, 137, 14, 245, 204, 99, 195, 42, 235, 248, 161, 86, 243,
|
||||
190, 135, 118, 130, 123, 154, 213, 150, 54, 74, 111, 20, 60, 240, 90, 37, 54, 109, 171,
|
||||
191, 123, 161, 140, 222, 100, 182, 202, 93, 88, 32, 80, 23, 168, 198, 153, 36, 97, 111,
|
||||
187, 151, 185, 43, 172, 245, 27, 6, 27, 82, 115, 199, 18, 239, 104, 158, 206, 205, 85,
|
||||
152, 42, 174, 185, 123, 197, 98, 65, 95, 135, 163, 206, 66, 59, 136, 109, 231, 125, 137,
|
||||
237, 153, 219, 97, 96, 237, 81, 201, 140, 31, 150, 226, 183, 192, 144, 113, 59, 86, 212,
|
||||
125, 182, 91, 33, 132, 158, 92, 12, 12, 68, 138, 149, 50, 36, 113, 147, 133, 95, 229,
|
||||
78, 235, 4, 228, 206, 188, 165, 95, 45, 225, 181, 1, 94, 107, 93, 128, 240, 251, 220,
|
||||
252, 7, 32, 135, 156, 83, 171, 14, 230, 48, 109, 203, 126, 89, 208, 99, 39, 140, 9,
|
||||
134, 185, 234, 60, 187, 73, 167, 24, 201, 152, 20, 166, 148, 27, 199, 28, 184, 26, 199,
|
||||
198, 0, 248, 52, 204, 119, 141, 157, 218, 181, 41, 227, 59, 227, 206, 119, 159, 23, 31,
|
||||
184, 224, 183, 204, 134, 76, 231, 77, 105, 160, 103, 48, 103, 104, 41, 155, 53, 160, 41,
|
||||
210, 123, 222, 252, 95, 26, 223, 45, 146, 126, 68, 177, 54, 37, 105, 3, 171, 182, 235,
|
||||
249, 31, 139, 97, 80, 243, 202, 121, 143, 0, 26, 184, 210, 149, 151, 207, 244, 177, 174,
|
||||
34, 67, 45, 102, 245, 100, 140, 95, 104, 55, 21, 83, 49, 53, 223, 147, 134, 210, 93,
|
||||
0, 97, 93, 26, 26, 48, 175, 178, 255, 164, 99, 174, 198, 167, 220, 45, 156, 64, 185,
|
||||
252, 168, 241, 18, 252, 35, 71, 219, 182, 205, 173, 19, 206, 15, 113, 232, 42, 161, 152,
|
||||
220, 160, 60, 64, 79, 3, 231, 43, 49, 132, 108, 235, 128, 21, 220, 146, 17, 255, 218,
|
||||
236, 182, 168, 154, 201, 118, 170, 58, 94, 212, 220, 246, 177, 125, 51, 241, 204, 55, 216,
|
||||
248, 104, 92, 100, 83, 221, 121, 48, 111, 138, 47, 73, 119, 230, 241, 17, 175, 103, 187,
|
||||
234, 198, 144, 199, 188, 65, 68, 240, 51, 17, 39, 11, 9, 143, 104, 109, 227, 70, 231,
|
||||
19, 181, 113, 66, 255, 233, 41, 241, 250, 217, 89, 182, 196, 31, 71, 139, 220, 137, 208,
|
||||
204, 188, 225, 243, 200, 234, 131, 48, 88, 102, 119, 63, 121, 44, 177, 188, 44, 154, 229,
|
||||
29, 149, 190, 118, 76, 130, 150, 147, 14, 114, 28, 222, 62, 217, 191, 50, 161, 170, 181,
|
||||
210, 2, 28, 73, 66, 149, 117, 243, 81, 162, 141, 55, 191, 35, 245, 54, 111, 120, 204,
|
||||
2, 134, 62, 31, 100, 125, 248, 36, 175, 153, 206, 101, 107, 209, 129, 181, 19, 22, 43,
|
||||
7, 104, 205, 149, 159, 140, 184, 149, 195, 39, 14, 143, 42, 148, 205, 73, 249, 74, 66,
|
||||
30, 250, 219, 237, 96, 71, 190, 225, 253, 210, 248, 40, 218, 96, 245, 111, 0, 130, 39,
|
||||
150, 69, 79, 165, 212, 122, 57, 162, 195, 51, 237, 6, 82, 231, 225, 63, 71, 41, 253,
|
||||
41, 38, 208, 33, 78, 170, 130, 68, 26, 131, 198, 66, 26, 12, 145, 191, 224, 11, 249,
|
||||
130, 207, 44, 112, 213, 126, 88, 183, 190, 160, 225, 187, 201, 8, 140, 235, 87, 55, 109,
|
||||
155, 81, 241, 98, 147, 11, 110, 37, 202, 79, 49, 195, 210, 0, 240, 66, 214, 110, 154,
|
||||
142, 44, 58, 111, 232, 4, 119, 117, 239, 207, 172, 93, 106, 254, 78, 205, 145, 89, 59,
|
||||
183, 35, 138, 232, 230, 92, 233, 214, 159, 191, 69, 58, 78, 114, 116, 189, 91, 121, 53,
|
||||
208, 104, 4, 125, 198, 111, 123, 20, 60, 13, 109, 120, 196, 145, 3, 172, 119, 95, 150,
|
||||
78, 255, 85, 147, 57, 163, 6, 174, 97, 97, 39, 151, 50, 144, 155, 175, 86, 11, 43,
|
||||
107, 71, 56, 216, 191, 253, 105, 194, 170, 225, 34, 64, 47, 34, 150, 195, 91, 58, 201,
|
||||
10, 155, 43, 49, 50, 93, 194, 206, 13, 25, 217, 56, 132, 33, 112, 92, 225, 109, 198,
|
||||
164, 23, 167, 199, 88, 215, 234, 238, 155, 69, 40, 100, 80, 196, 144, 129, 246, 237, 68,
|
||||
197, 250, 93, 159, 51, 225, 193, 163, 62, 163, 17, 4, 71, 41, 172, 15, 130, 132, 249,
|
||||
112, 31, 63, 152, 132, 143, 92, 20, 17, 83, 1, 86, 25, 252, 179, 185, 47, 149, 122,
|
||||
211, 211, 29, 229, 216, 101, 15, 133, 117, 145, 9, 111, 1, 40, 175, 154, 173, 62, 247,
|
||||
193, 80, 75, 194, 166, 100, 191, 90, 29, 239, 239, 152, 194, 195, 182, 168, 156, 27, 183,
|
||||
33, 145, 73, 43, 0, 75, 83, 175, 229, 0, 238, 221, 194, 63, 40, 133, 230, 140, 68,
|
||||
64, 170, 51, 48, 66, 246, 243, 248, 159, 144, 20, 87, 177, 165, 160, 220, 166, 235, 48,
|
||||
86, 209, 49, 68, 174, 243, 132, 214, 120, 106, 99, 189, 170, 13, 241, 219, 80, 232, 207,
|
||||
72, 135, 95, 92, 223, 16, 2, 127, 237, 169, 107, 29, 255, 61, 79, 68, 236, 67, 200,
|
||||
194, 188, 50, 38, 121, 221, 52, 107, 184, 132, 84, 136, 204, 219, 231, 41, 186, 248, 44,
|
||||
58, 229, 213, 166, 3, 212, 227, 82, 25, 207, 150, 225, 146, 82, 20, 185, 204, 242, 237,
|
||||
55, 170, 113, 139, 50, 62, 103, 26, 103, 34, 18, 148, 93, 247, 105, 3, 251, 62, 231,
|
||||
77, 87, 182, 227, 57, 73, 54, 77, 2, 2, 63, 239, 57, 234, 97, 197, 29, 159, 44,
|
||||
55, 7, 79, 74, 155, 172, 66, 5, 175, 61, 67, 150, 139, 155, 77, 111, 212, 151, 165,
|
||||
34, 153, 167, 98, 137, 225, 77, 234, 166, 107, 138, 211, 163, 145, 34, 237, 45, 206, 47,
|
||||
50, 126, 108, 117, 21, 248, 17, 98, 103, 230, 249, 12, 9, 147, 179, 107, 29, 149, 185,
|
||||
7, 59, 37, 146, 14, 200, 35, 49, 182, 80, 0, 230, 130, 126, 83, 248, 148, 75, 9,
|
||||
247, 178, 240, 240, 190, 249, 132, 114, 101, 161, 7, 30, 169, 67, 68, 59, 82, 12, 95,
|
||||
131, 195, 176, 131, 169, 51, 2, 252, 44, 150, 72, 54, 141, 250, 38, 126, 185, 31, 3,
|
||||
44, 132, 165, 52, 163, 78, 120, 231, 138, 202, 244, 234, 77, 183, 155, 209, 97, 207, 212,
|
||||
94, 251, 107, 166, 49, 249, 161, 88, 120, 91, 120, 123, 135, 253, 33, 188, 160, 112, 52,
|
||||
136, 250, 254, 125, 229, 76, 53, 128, 30, 150, 79, 243, 244, 75, 95, 155, 125, 88, 60,
|
||||
213, 209, 152, 78, 77, 32, 75, 110, 220, 236, 222, 17, 117, 217, 15, 242, 190, 92, 39,
|
||||
63, 123, 190, 143, 111, 178, 219, 206, 78, 88, 38, 138, 46, 247, 34, 124, 69, 66, 199,
|
||||
179, 31, 179, 145, 48, 41, 106, 64, 27, 41, 157, 67, 105, 24, 1, 249, 135, 179, 212,
|
||||
86, 1, 44, 124, 140, 91, 116, 175, 215, 185, 242, 159, 108, 17, 83, 254, 66, 124, 105,
|
||||
131, 151, 146, 32, 218, 252, 57, 219, 245, 193, 143, 201, 23, 145, 246, 148, 30, 82, 8,
|
||||
206, 41, 194, 192, 201, 47, 210, 28, 46, 20, 152, 151, 151, 48, 42, 184, 11, 38, 241,
|
||||
231, 28, 179, 119, 230, 202, 8, 220, 94, 39, 46, 103, 245, 88, 42, 181, 33, 90, 136,
|
||||
62, 136, 156, 214, 31, 52, 7, 74, 237, 19, 113, 223, 250, 141, 146, 113, 115, 92, 122,
|
||||
80, 187, 161, 126, 35, 150, 215, 78, 76, 249, 168, 212, 55, 48, 113, 14, 80, 166, 21,
|
||||
154, 147, 40, 12, 114, 35, 153, 5, 148, 12, 98, 15, 92, 29, 176, 219, 65, 71, 179,
|
||||
143, 147, 172, 56, 104, 227, 104, 218, 241, 185, 128, 7, 84, 20, 47, 96, 135, 82, 249,
|
||||
140, 231, 6, 238, 246, 99, 12, 167, 63, 77, 238, 242, 221, 130, 158, 21, 235, 129, 126,
|
||||
197, 114, 56, 69, 121, 140, 90, 169, 237, 225, 252, 231, 109, 228, 237, 91, 219, 81, 104,
|
||||
130, 144, 181, 113, 130, 147, 244, 32, 169, 223, 162, 39, 164, 21, 95, 234, 143, 236, 68,
|
||||
57, 217, 37, 53, 192, 147, 25, 174, 239, 245, 0, 87, 119, 144, 13, 232, 19, 160, 220,
|
||||
51, 73, 188, 214, 113, 96, 235, 209, 75, 122, 190, 144, 179, 151, 181, 233, 88, 73, 3,
|
||||
7, 56, 248, 7, 143, 112, 152, 156, 89, 171, 61, 53, 223, 135, 242, 181, 248, 83, 161,
|
||||
202, 158, 28, 136, 46, 208, 32, 228, 186, 121, 45, 189, 128, 102, 182, 136, 246, 38, 32,
|
||||
147, 127, 204, 208, 181, 171, 87, 167, 97, 80, 250, 2, 26, 153, 31, 163, 200, 239, 195,
|
||||
172, 169, 60, 218, 103, 188, 65, 30, 69, 55, 68, 102, 202, 196, 50, 154, 121, 221, 242,
|
||||
33, 63, 67, 28, 66, 93, 181, 97, 0, 126, 81, 196, 43, 251, 0, 5, 98, 189, 70,
|
||||
128, 3, 126, 197, 105, 72, 137, 155, 227, 3, 121, 214, 36, 184, 25, 65, 250, 118, 247,
|
||||
91, 119, 117, 173, 60, 160, 168, 60, 166, 10, 250, 237, 139, 253, 107, 80, 102, 180, 217,
|
||||
2, 151, 221, 123, 109, 1, 52, 134, 66, 46, 253, 57, 138, 117, 175, 55, 178, 79, 223,
|
||||
239, 245, 234, 233, 226, 117, 231, 78, 198, 78, 2, 159, 80, 154, 124, 204, 7, 126, 0,
|
||||
142, 193, 47, 140, 251, 185, 2, 170, 241, 180, 249, 208, 163, 239, 186, 141, 210, 48, 116,
|
||||
32, 246, 195, 34, 150, 19, 188, 19, 224, 196, 146, 224, 83, 83, 15, 224, 78, 201, 226,
|
||||
249, 186, 151, 243, 139, 58, 226, 70, 199, 181, 118, 60, 213, 109, 255, 248, 3, 19, 181,
|
||||
23, 243, 122, 169, 212, 205, 252, 228, 173, 75, 173, 144, 68, 104, 39, 55, 243, 98, 26,
|
||||
57, 41, 207, 175, 102, 165, 29, 102, 158, 32, 121, 83, 56, 109, 205, 225, 66, 155, 222,
|
||||
38, 73, 42, 212, 218, 110, 60, 1, 166, 48, 99, 193, 105, 141, 145, 25, 244, 54, 54,
|
||||
90, 213, 87, 212, 40, 143, 66, 246, 112, 132, 146, 79, 171, 220, 121, 128, 182, 232, 189,
|
||||
184, 143, 237, 27, 80, 86, 169, 226, 112, 158, 25, 166, 248, 238, 253, 204, 23, 141, 15,
|
||||
13, 254, 147, 160, 77, 63, 124, 199, 191, 50, 175, 124, 234, 62, 105, 6, 143, 192, 176,
|
||||
113, 48, 78, 139, 215, 71, 121, 213, 20, 144, 98, 35, 158, 96, 183, 62, 174, 246, 187,
|
||||
117, 182, 237, 37, 50, 216, 99, 156, 223, 243, 93, 143, 101, 142, 222, 240, 101, 37, 106,
|
||||
58, 57, 250, 157, 93, 153, 254, 20, 216, 172, 10, 147, 34, 192, 129, 71, 243, 90, 171,
|
||||
144, 57, 159, 238, 201, 4, 124, 167, 244, 225, 205, 95, 28, 7, 89, 185, 100, 243, 184,
|
||||
121, 203, 100, 131, 95, 135, 68, 224, 207, 56, 58, 122, 201, 115, 25, 183, 61, 30, 51,
|
||||
229, 18, 21, 178, 113, 49, 186, 203, 235, 31, 191, 163, 152, 138, 8, 28, 233, 143, 97,
|
||||
202, 95, 153, 4, 217, 98, 120, 243, 26, 182, 17, 77, 155, 36, 99, 78, 150, 149, 8,
|
||||
98, 128, 39, 33, 36, 192, 172, 45, 220, 149, 189, 61, 96, 28, 215, 100, 246, 58, 221,
|
||||
233, 84, 147, 251, 162, 47, 31, 5, 125, 181, 154, 134, 23, 27, 174, 57, 64, 110, 229,
|
||||
109, 75, 123, 43, 136, 219, 71, 95, 64, 61, 154, 29, 39, 238, 177, 34, 145, 225, 65,
|
||||
150, 94, 247, 49, 229, 15, 77, 147, 72, 141, 2, 45, 251, 77, 169, 38, 213, 132, 110,
|
||||
53, 196, 172, 207, 226, 212, 190, 148, 246, 79, 117, 56, 230, 212, 48, 23, 185, 63, 100,
|
||||
76, 136, 242, 78, 181, 237, 156, 95, 20, 113, 227, 131, 167, 168, 47, 119, 139, 3, 53,
|
||||
31, 250, 133, 149, 50, 107, 105, 99, 130, 34, 162, 231, 111, 42, 217, 190, 224, 199, 90,
|
||||
63, 220, 204, 35, 95, 115, 203, 143, 234, 86, 147, 32, 118, 141, 165, 11, 192, 16, 117,
|
||||
35, 147, 152, 198, 123, 7, 240, 84, 198, 209, 28, 33, 17, 248, 237, 52, 88, 97, 255,
|
||||
231, 76, 86, 122, 109, 204, 8, 18, 216, 201, 35, 77, 237, 183, 229, 179, 50, 237, 164,
|
||||
135, 179, 118, 164, 213, 135, 157, 195, 187, 245, 36, 187, 220, 113, 18, 87, 222, 222, 96,
|
||||
241, 183, 42, 21, 4, 23, 205, 233, 203, 0, 214, 112, 136, 138, 230, 44, 95, 110, 201,
|
||||
34, 41, 191, 71, 229, 155, 185, 247, 243, 151, 214, 84, 137, 141, 126, 159, 146, 149, 108,
|
||||
124, 97, 109, 82, 209, 245, 221, 183, 34, 60, 37, 236, 95, 79, 171, 167, 53, 71, 96,
|
||||
45, 58, 248, 3, 142, 129, 145, 12, 33, 36, 162, 142, 160, 3, 251, 243, 213, 240, 208,
|
||||
141, 19, 13, 178, 255, 109, 2, 170, 20, 55, 241, 116, 101, 44, 108, 105, 186, 238, 251,
|
||||
199, 15, 31, 106, 157, 191, 110, 152, 178, 67, 137, 131, 208, 156, 144, 131, 155, 253, 134,
|
||||
70, 18, 190, 55, 134, 35, 99, 243, 140, 30, 225, 135, 230, 240, 166, 81, 142, 102, 191,
|
||||
39, 25, 3, 177, 156, 211, 77, 45, 87, 233, 43, 221, 48, 61, 155, 103, 195, 191, 203,
|
||||
182, 75, 233, 152, 211, 208, 136, 121, 33, 23, 224, 224, 62, 249, 227, 239, 149, 183, 61,
|
||||
195, 15, 39, 238, 236, 87, 43, 136, 191, 239, 71, 138, 166, 147, 116, 62, 102, 68, 199,
|
||||
224, 101, 223, 193, 70, 29, 186, 42, 13, 80, 225, 75, 19, 241, 115, 1, 221, 202, 45,
|
||||
102, 137, 29, 174, 20, 195, 66, 136, 2, 168, 205, 201, 137, 50, 168, 74, 121, 198, 4,
|
||||
163, 212, 85, 133, 31, 105, 118, 146, 106, 84, 93, 152, 187, 231, 181, 105, 251, 121, 171,
|
||||
132, 123, 84, 81, 69, 221, 132, 238, 40, 253, 181, 45, 161, 137, 130, 39, 169, 235, 158,
|
||||
59, 86, 242, 153, 239, 173, 128, 165, 23, 123, 30, 195, 0, 154, 23, 81, 224, 245, 214,
|
||||
206, 30, 212, 131, 75, 117, 12, 206, 157, 181, 186, 59, 241, 17, 45, 138, 0, 219, 11,
|
||||
165, 243, 135, 196, 182, 135, 95, 205, 217, 63, 195, 175, 14, 225, 131, 145, 45, 249, 158,
|
||||
251, 150, 84, 182, 209, 70, 199, 255, 209, 199, 219, 220, 109, 206, 99, 50, 132, 234, 146,
|
||||
82, 195, 209, 22, 114, 223, 247, 246, 113, 37, 239, 16, 33, 134, 100, 215, 88, 170, 158,
|
||||
87, 123, 102, 50, 88, 211, 1, 187, 6, 134, 165, 152, 216, 105, 106, 239, 220, 74, 231,
|
||||
210, 187, 12, 194, 204, 45, 72, 49, 4, 160, 219, 162, 248, 87, 8, 43, 176, 220, 44,
|
||||
107, 227, 178, 17, 124, 139, 122, 230, 122, 87, 48, 97, 42, 236, 110, 236, 185, 155, 53,
|
||||
234, 159, 214, 198, 66, 206, 30, 75, 249, 206, 40, 38, 57, 11, 217, 74, 136, 100, 197,
|
||||
110, 223, 29, 159, 65, 71, 140, 175, 51, 69, 74, 105, 48, 234, 63, 246, 45, 13, 20,
|
||||
121, 7, 226, 161, 46, 28, 173, 7, 103, 53, 108, 45, 164, 76, 74, 68, 141, 145, 208,
|
||||
61, 197, 22, 136, 46, 70, 115, 110, 60, 161, 124, 81, 26, 132, 51, 188, 178, 79, 106,
|
||||
186, 183, 160, 39, 228, 68, 115, 46, 136, 1, 192, 89, 62, 133, 112, 198, 180, 182, 58,
|
||||
34, 243, 219, 158, 69, 245, 34, 120, 178, 213, 200, 28, 143, 128, 188, 182, 100, 1, 41,
|
||||
146, 137, 43, 82, 227, 105, 216, 83, 48, 140, 10, 106, 175, 254, 70, 77, 67, 59, 112,
|
||||
188, 237, 69, 133, 10, 212, 5, 198, 138, 105, 199, 180, 252, 81, 223, 79, 53, 73, 39,
|
||||
137, 121, 180, 148, 228, 99, 146, 42, 177, 214, 102, 33, 147, 84, 102, 25, 94, 59, 31,
|
||||
37, 197, 137, 237, 122, 133, 63, 90, 213, 116, 163, 253, 253, 29, 177, 145, 2, 21, 36,
|
||||
45, 198, 251, 147, 231, 143, 232, 78, 168, 71, 137, 199, 108, 79, 80, 90, 201, 214, 153,
|
||||
35, 172, 13, 199, 169, 11, 228, 91, 157, 231, 112, 193, 20, 54, 189, 167, 30, 77, 144,
|
||||
108, 245, 215, 246, 189, 68, 69, 14, 158, 14, 228, 55, 50, 145, 69, 249, 58, 80, 222,
|
||||
149, 237, 198, 5, 175, 218, 60, 109, 130, 91, 186, 18, 200, 175, 234, 190, 109, 46, 3,
|
||||
123, 204, 18, 96, 4, 68, 241, 73, 62, 44, 154, 29, 193, 136, 227, 199, 55, 189, 4,
|
||||
164, 64, 95, 95, 82, 39, 15, 60, 230, 124, 107, 233, 248, 55, 251, 89, 60, 63, 75,
|
||||
134, 126, 119, 32, 156, 57, 168, 127, 0, 224, 61, 5, 133, 125, 100, 228, 208, 140, 243,
|
||||
12, 114, 111, 119, 92, 104, 175, 87, 193, 236, 151, 13, 114, 21, 132, 146, 177, 189, 59,
|
||||
49, 190, 27, 110, 195, 160, 236, 40, 132, 188, 181, 120, 201, 40, 232, 65, 132, 80, 241,
|
||||
220, 18, 221, 115, 31, 79, 137, 164, 226, 58, 98, 29, 108, 32, 57, 219, 228, 218, 199,
|
||||
13, 95, 132, 195, 215, 77, 235, 191, 143, 112, 16, 128, 76, 35, 93, 191, 66, 173, 73,
|
||||
231, 143, 132, 73, 173, 240, 106, 231, 203, 78, 193, 147, 92, 33, 23, 31, 248, 100, 11,
|
||||
184, 243, 123, 201, 115, 200, 236, 209, 135, 47, 126, 209, 22, 14, 85, 95, 188, 69, 202,
|
||||
163, 17, 24, 101, 164, 117, 134, 187, 148, 127, 31, 159, 55, 19, 27, 1, 135, 227, 237,
|
||||
89, 107, 28, 216, 60, 51, 230, 145, 147, 163, 215, 93, 70, 232, 118, 172, 140, 235, 50,
|
||||
71, 128, 177, 103, 32, 233, 123, 60, 234, 2, 31, 216, 91, 139, 244, 52, 200, 40, 26,
|
||||
90, 188, 189, 49, 25, 4, 25, 144, 176, 166, 124, 227, 237, 252, 148, 85, 29, 125, 208,
|
||||
89, 104, 210, 121, 64, 46, 4, 53, 99, 204, 93, 125, 38, 25, 59, 88, 51, 64, 113,
|
||||
195, 241, 23, 64, 212, 5, 60, 104, 90, 90, 230, 42, 179, 78, 253, 44, 143, 44, 49,
|
||||
196, 143, 254, 34, 13, 36, 60, 73, 125, 112, 137, 239, 52, 122, 7, 116, 79, 12, 177,
|
||||
183, 103, 11, 158, 146, 190, 237, 143, 235, 124, 188, 28, 65, 76, 26, 100, 89, 63, 160,
|
||||
163, 188, 17, 44, 172, 69, 167, 179, 185, 246, 191, 107, 174, 38, 118, 76, 184, 53, 58,
|
||||
72, 32, 182, 5, 61, 248, 81, 88, 92, 170, 152, 253, 77, 84, 14, 122, 1, 83, 34,
|
||||
180, 13, 25, 115, 120, 199, 154, 238, 20, 83, 36, 79, 155, 68, 5, 160, 130, 254, 242,
|
||||
218, 90, 156, 114, 87, 234, 199, 101, 101, 200, 185, 135, 124, 198, 160, 240, 62, 104, 138,
|
||||
45, 125, 222, 81, 204, 122, 150, 210, 26, 24, 208, 12, 242, 42, 169, 101, 130, 148, 44,
|
||||
232, 249, 245, 161, 128, 113, 103, 33, 98, 166, 137, 236, 212, 7, 202, 38, 211, 69, 188,
|
||||
165, 95, 212, 118, 108, 199, 161, 22, 45, 35, 170, 90, 11, 163, 79, 173, 36, 193, 20,
|
||||
69, 35, 187, 207, 16, 144, 214, 219, 182, 170, 32, 114, 79, 128, 71, 198, 237, 15, 103,
|
||||
4, 60, 139, 175, 150, 151, 82, 230, 68, 119, 168, 89, 188, 204, 20, 140, 220, 165, 98,
|
||||
184, 91, 12, 217, 205, 92, 90, 20, 35, 71, 36, 138, 76, 96, 22, 251, 247, 173, 78,
|
||||
222, 241, 197, 134, 75, 130, 83, 96, 14, 47, 5, 113, 232, 96, 126, 193, 45, 218, 28,
|
||||
66, 253, 99, 103, 136, 176, 200, 158, 171, 191, 76, 249, 158, 62, 190, 37, 137, 65, 120,
|
||||
233, 80, 168, 238, 193, 145, 79, 63, 82, 125, 26, 111, 191, 24, 210, 39, 161, 131, 239,
|
||||
64, 46, 175, 140, 39, 77, 202, 230, 115, 84, 40, 235, 62, 120, 148, 45, 57, 37, 124,
|
||||
121, 120, 249, 148, 231, 185, 172, 186, 224, 77, 61, 207, 141, 107, 126, 26, 147, 204, 229,
|
||||
121, 63, 58, 161, 43, 120, 25, 191, 165, 83, 228, 34, 205, 92, 27, 97, 67, 213, 13,
|
||||
253, 182, 91, 59, 133, 233, 166, 4, 4, 57, 209, 233, 179, 16, 35, 85, 59, 155, 111,
|
||||
250, 65, 194, 223, 99, 144, 59, 127, 241, 127, 85, 255, 125, 11, 90, 184, 145, 68, 95,
|
||||
150, 72, 153, 103, 49, 76, 120, 85, 161, 179, 241, 16, 174, 51, 211, 142, 150, 99, 201,
|
||||
22, 85, 73, 108, 84, 199, 120, 175, 128, 9, 243, 223, 160, 59, 120, 8, 109, 197, 128,
|
||||
194, 103, 52, 180, 119, 227, 231, 75, 113, 126, 175, 59, 148, 4, 132, 1, 89, 75, 121,
|
||||
8, 204, 131, 251, 171, 36, 55, 36, 44, 165, 233, 172, 103, 80, 224, 28, 200, 195, 3,
|
||||
20, 53, 129, 195, 112, 22, 200, 244, 23, 34, 64, 145, 42, 12, 20, 38, 184, 56, 94,
|
||||
220, 101, 3, 198, 17, 107, 22, 242, 135, 222, 182, 138, 243, 235, 11, 182, 91, 34, 127,
|
||||
80, 58, 161, 145, 203, 204, 158, 224, 242, 86, 24, 81, 51, 126, 84, 249, 143, 191, 15,
|
||||
130, 70, 238, 57, 209, 225, 36, 221, 152, 128, 255, 24, 208, 57, 186, 97, 4, 134, 255,
|
||||
229, 121, 86, 254, 202, 137, 124, 31, 130, 12, 222, 146, 142, 37, 129, 199, 247, 98, 236,
|
||||
212, 251, 108, 211, 20, 60, 13, 206, 158, 18, 84};
|
||||
|
||||
SimpleConvolution::SimpleConvolution() {
|
||||
width_ = 64;
|
||||
height_ = 64;
|
||||
mask_width_ = 3;
|
||||
mask_height_ = mask_width_;
|
||||
randomize_seed_ = 0;
|
||||
|
||||
if (!IsPowerOf2(width_)) {
|
||||
width_ = RoundToPowerOf2(width_);
|
||||
}
|
||||
|
||||
if (!IsPowerOf2(height_)) {
|
||||
height_ = RoundToPowerOf2(height_);
|
||||
}
|
||||
|
||||
if (!(mask_width_ % 2)) {
|
||||
mask_width_++;
|
||||
}
|
||||
|
||||
if (!(mask_height_ % 2)) {
|
||||
mask_height_++;
|
||||
}
|
||||
|
||||
if (width_ * height_ < 256) {
|
||||
width_ = 64;
|
||||
height_ = 64;
|
||||
}
|
||||
|
||||
const uint32_t input_size_bytes = width_ * height_ * sizeof(uint32_t);
|
||||
const uint32_t mask_size_bytes = mask_width_ * mask_height_ * sizeof(float);
|
||||
|
||||
SetSysDescr(KERNARG_DES_ID, sizeof(kernel_args_t));
|
||||
SetSysDescr(INPUT_DES_ID, input_size_bytes);
|
||||
SetSysDescr(OUTPUT_DES_ID, input_size_bytes);
|
||||
SetLocalDescr(LOCAL_DES_ID, input_size_bytes);
|
||||
SetSysDescr(MASK_DES_ID, mask_size_bytes);
|
||||
SetSysDescr(REFOUT_DES_ID, input_size_bytes);
|
||||
|
||||
if (!randomize_seed_) TEST_ASSERT(sizeof(input_data_) <= input_size_bytes);
|
||||
}
|
||||
|
||||
void SimpleConvolution::Init() {
|
||||
std::clog << "SimpleConvolution::init :" << std::endl;
|
||||
|
||||
mem_descr_t input_des = GetDescr(INPUT_DES_ID);
|
||||
mem_descr_t local_des = GetDescr(LOCAL_DES_ID);
|
||||
mem_descr_t mask_des = GetDescr(MASK_DES_ID);
|
||||
mem_descr_t refout_des = GetDescr(REFOUT_DES_ID);
|
||||
mem_descr_t kernarg_des = GetDescr(KERNARG_DES_ID);
|
||||
|
||||
uint32_t* input = (uint32_t*)input_des.ptr;
|
||||
uint32_t* output_local = (uint32_t*)local_des.ptr;
|
||||
float* mask = (float*)mask_des.ptr;
|
||||
kernel_args_t* kernel_args = (kernel_args_t*)kernarg_des.ptr;
|
||||
|
||||
if (randomize_seed_) {
|
||||
// random initialisation of input
|
||||
FillRandom<uint32_t>(input, width_, height_, 0, 255, randomize_seed_);
|
||||
} else {
|
||||
// initialization with preset values
|
||||
memcpy(input, input_data_, width_ * height_ * sizeof(uint32_t));
|
||||
}
|
||||
|
||||
// Fill a blurr filter or some other filter of your choice
|
||||
const float val = 1.0f / (mask_width_ * 2.0f - 1.0f);
|
||||
for (uint32_t i = 0; i < (mask_width_ * mask_height_); i++) {
|
||||
mask[i] = 0;
|
||||
}
|
||||
for (uint32_t i = 0; i < mask_width_; i++) {
|
||||
uint32_t y = mask_height_ / 2;
|
||||
mask[y * mask_width_ + i] = val;
|
||||
}
|
||||
for (uint32_t i = 0; i < mask_height_; i++) {
|
||||
uint32_t x = mask_width_ / 2;
|
||||
mask[i * mask_width_ + x] = val;
|
||||
}
|
||||
|
||||
// Print the INPUT array.
|
||||
std::clog << std::dec;
|
||||
PrintArray<uint32_t>("> Input[0]", input, width_, 1);
|
||||
PrintArray<float>("> Mask", mask, mask_width_, mask_height_);
|
||||
|
||||
// Fill the kernel args
|
||||
kernel_args->arg1 = output_local;
|
||||
kernel_args->arg2 = input;
|
||||
kernel_args->arg3 = mask;
|
||||
kernel_args->arg4 = width_;
|
||||
kernel_args->arg41 = height_;
|
||||
kernel_args->arg5 = mask_width_;
|
||||
kernel_args->arg51 = mask_height_;
|
||||
|
||||
// Calculate the reference output
|
||||
memset(refout_des.ptr, 0, refout_des.size);
|
||||
ReferenceImplementation(reinterpret_cast<uint32_t*>(refout_des.ptr), input, mask, width_, height_,
|
||||
mask_width_, mask_height_);
|
||||
}
|
||||
|
||||
void SimpleConvolution::PrintOutput() const {
|
||||
PrintArray<uint32_t>("> Output[0]", reinterpret_cast<uint32_t*>(GetOutputPtr()), width_, 1);
|
||||
}
|
||||
|
||||
bool SimpleConvolution::ReferenceImplementation(uint32_t* output, const uint32_t* input,
|
||||
const float* mask, const uint32_t width,
|
||||
const uint32_t height, const uint32_t mask_width,
|
||||
const uint32_t mask_height) {
|
||||
const uint32_t vstep = (mask_width - 1) / 2;
|
||||
const uint32_t hstep = (mask_height - 1) / 2;
|
||||
|
||||
// for each pixel in the input
|
||||
for (uint32_t x = 0; x < width; x++) {
|
||||
for (uint32_t y = 0; y < height; y++) {
|
||||
// find the left, right, top and bottom indices such that
|
||||
// the indices do not go beyond image boundaires
|
||||
const uint32_t left = (x < vstep) ? 0 : (x - vstep);
|
||||
const uint32_t right = ((x + vstep) >= width) ? width - 1 : (x + vstep);
|
||||
const uint32_t top = (y < hstep) ? 0 : (y - hstep);
|
||||
const uint32_t bottom = ((y + hstep) >= height) ? height - 1 : (y + hstep);
|
||||
|
||||
// initializing wighted sum value
|
||||
float sum_fx = 0;
|
||||
for (uint32_t i = left; i <= right; ++i) {
|
||||
for (uint32_t j = top; j <= bottom; ++j) {
|
||||
// performing wighted sum within the mask boundaries
|
||||
uint32_t mask_idx = (j - (y - hstep)) * mask_width + (i - (x - vstep));
|
||||
uint32_t index = j * width + i;
|
||||
|
||||
// to round to the nearest integer
|
||||
sum_fx += ((float)input[index] * mask[mask_idx]);
|
||||
}
|
||||
}
|
||||
sum_fx += 0.5f;
|
||||
output[y * width + x] = uint32_t(sum_fx);
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
@@ -0,0 +1,96 @@
|
||||
/******************************************************************************
|
||||
|
||||
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_SIMPLE_CONVOLUTION_SIMPLE_CONVOLUTION_H_
|
||||
#define TEST_SIMPLE_CONVOLUTION_SIMPLE_CONVOLUTION_H_
|
||||
|
||||
#include <map>
|
||||
#include <vector>
|
||||
|
||||
#include "ctrl/test_kernel.h"
|
||||
|
||||
// Class implements SimpleConvolution kernel parameters
|
||||
class SimpleConvolution : public TestKernel {
|
||||
public:
|
||||
// Constructor
|
||||
SimpleConvolution();
|
||||
|
||||
// Initialize method
|
||||
void Init();
|
||||
|
||||
// Return compute grid size
|
||||
uint32_t GetGridSize() const { return width_ * height_; }
|
||||
|
||||
// Print output
|
||||
void PrintOutput() const;
|
||||
|
||||
// Return name
|
||||
std::string Name() const { return std::string("SimpleConvolution"); }
|
||||
|
||||
private:
|
||||
// Local kernel arguments declaration
|
||||
struct kernel_args_t {
|
||||
void* arg1;
|
||||
void* arg2;
|
||||
void* arg3;
|
||||
uint32_t arg4;
|
||||
uint32_t arg41;
|
||||
uint32_t arg5;
|
||||
uint32_t arg51;
|
||||
};
|
||||
|
||||
// Reference CPU implementation of Simple Convolution
|
||||
// @param output Output matrix after performing convolution
|
||||
// @param input Input matrix on which convolution is to be performed
|
||||
// @param mask mask matrix using which convolution was to be performed
|
||||
// @param input_dimensions dimensions of the input matrix
|
||||
// @param mask_dimensions dimensions of the mask matrix
|
||||
// @return bool true on success and false on failure
|
||||
bool ReferenceImplementation(uint32_t* output, const uint32_t* input, const float* mask,
|
||||
const uint32_t width, const uint32_t height,
|
||||
const uint32_t maskWidth, const uint32_t maskHeight);
|
||||
|
||||
// Width of the Input array
|
||||
uint32_t width_;
|
||||
|
||||
// Height of the Input array
|
||||
uint32_t height_;
|
||||
|
||||
// Mask dimensions
|
||||
uint32_t mask_width_;
|
||||
|
||||
// Mask dimensions
|
||||
uint32_t mask_height_;
|
||||
|
||||
// Randomize input data
|
||||
unsigned randomize_seed_;
|
||||
|
||||
// Input data
|
||||
static const uint32_t input_data_[];
|
||||
};
|
||||
|
||||
#endif // TEST_SIMPLE_CONVOLUTION_SIMPLE_CONVOLUTION_H_
|
||||
@@ -0,0 +1,154 @@
|
||||
module &m:1:0:$full:$large:$default;
|
||||
extension "amd:gcn";
|
||||
extension "IMAGE";
|
||||
|
||||
decl prog function &abort()();
|
||||
|
||||
prog kernel &__OpenCL_SimpleConvolution(kernarg_u64 %__global_offset_0,
|
||||
kernarg_u64 %output,
|
||||
kernarg_u64 %input,
|
||||
kernarg_u64 %mask,
|
||||
kernarg_u32 %inputDimensions[2],
|
||||
kernarg_u32 %maskDimensions[2]) {
|
||||
|
||||
pragma "AMD RTI", "ARGSTART:__OpenCL_SimpleConvolution";
|
||||
pragma "AMD RTI", "version:3:1:104";
|
||||
pragma "AMD RTI", "device:generic";
|
||||
pragma "AMD RTI", "uniqueid:1024";
|
||||
pragma "AMD RTI", "memory:private:0";
|
||||
pragma "AMD RTI", "memory:region:0";
|
||||
pragma "AMD RTI", "memory:local:0";
|
||||
pragma "AMD RTI", "value:__global_offset_0:u64:1:1:0";
|
||||
pragma "AMD RTI", "pointer:output:u32:1:1:96:uav:7:4:RW:0:0:0";
|
||||
pragma "AMD RTI", "pointer:input:u32:1:1:112:uav:7:4:RW:0:0:0";
|
||||
pragma "AMD RTI", "pointer:mask:float:1:1:128:uav:7:4:RW:0:0:0";
|
||||
pragma "AMD RTI", "value:inputDimensions:u32:2:1:144";
|
||||
pragma "AMD RTI", "constarg:4:inputDimensions";
|
||||
pragma "AMD RTI", "value:maskDimensions:u32:2:1:160";
|
||||
pragma "AMD RTI", "constarg:5:maskDimensions";
|
||||
pragma "AMD RTI", "function:1:0";
|
||||
pragma "AMD RTI", "memory:64bitABI";
|
||||
pragma "AMD RTI", "privateid:8";
|
||||
pragma "AMD RTI", "enqueue_kernel:0";
|
||||
pragma "AMD RTI", "kernel_index:0";
|
||||
pragma "AMD RTI", "reflection:0:size_t";
|
||||
pragma "AMD RTI", "reflection:1:uint*";
|
||||
pragma "AMD RTI", "reflection:2:uint*";
|
||||
pragma "AMD RTI", "reflection:3:float*";
|
||||
pragma "AMD RTI", "reflection:4:uint2";
|
||||
pragma "AMD RTI", "reflection:5:uint2";
|
||||
pragma "AMD RTI", "ARGEND:__OpenCL_SimpleConvolution";
|
||||
|
||||
@__OpenCL_SimpleConvolution_Entry:
|
||||
|
||||
// BB#0: // %entry
|
||||
|
||||
workitemabsid_u32 $s6, 0;
|
||||
cvt_u64_u32 $d0, $s6;
|
||||
ld_kernarg_align(8)_width(all)_u64 $d4, [%__global_offset_0];
|
||||
add_u64 $d0, $d0, $d4;
|
||||
cvt_u32_u64 $s5, $d0;
|
||||
ld_v2_kernarg_align(4)_width(all)_u32 ($s0, $s4), [%inputDimensions];
|
||||
ld_v2_kernarg_align(4)_width(all)_u32 ($s1, $s9), [%maskDimensions];
|
||||
rem_u32 $s7, $s5, $s0;
|
||||
add_u32 $s2, $s1, 4294967295;
|
||||
shr_u32 $s8, $s2, 1;
|
||||
add_u32 $s2, $s7, $s8;
|
||||
add_u32 $s3, $s0, 4294967295;
|
||||
cmp_ge_b1_u32 $c0, $s2, $s0;
|
||||
cmov_b32 $s2, $c0, $s3, $s2;
|
||||
sub_u32 $s3, $s7, $s8;
|
||||
cmp_lt_b1_u32 $c0, $s7, $s8;
|
||||
cmov_b32 $s3, $c0, 0, $s3;
|
||||
ld_kernarg_align(8)_width(all)_u64 $d1, [%output];
|
||||
cmp_le_b1_u32 $c0, $s3, $s2;
|
||||
cbr_b1 $c0, @BB0_2;
|
||||
|
||||
// BB#1:
|
||||
|
||||
mov_b32 $s6, 0;
|
||||
br @BB0_6;
|
||||
|
||||
// @BB0_2: // %for.cond32.preheader.lr.ph
|
||||
|
||||
@BB0_2:
|
||||
|
||||
div_u32 $s5, $s5, $s0;
|
||||
add_u32 $s9, $s9, 4294967295;
|
||||
shr_u32 $s9, $s9, 1;
|
||||
add_u32 $s10, $s5, $s9;
|
||||
add_u32 $s11, $s4, 4294967295;
|
||||
cmp_ge_b1_u32 $c0, $s10, $s4;
|
||||
cmov_b32 $s4, $c0, $s11, $s10;
|
||||
sub_u32 $s10, $s5, $s9;
|
||||
cmp_lt_b1_u32 $c0, $s5, $s9;
|
||||
cmov_b32 $s5, $c0, 0, $s10;
|
||||
ld_kernarg_align(8)_width(all)_u64 $d2, [%mask];
|
||||
ld_kernarg_align(8)_width(all)_u64 $d3, [%input];
|
||||
cvt_u64_u32 $d5, $s6;
|
||||
add_u64 $d4, $d4, $d5;
|
||||
cvt_u32_u64 $s6, $d4;
|
||||
div_u32 $s6, $s6, $s0;
|
||||
max_u32 $s10, $s9, $s6;
|
||||
sub_u32 $s12, $s10, $s6;
|
||||
max_u32 $s11, $s7, $s8;
|
||||
mov_b32 $s6, 0;
|
||||
mad_u32 $s12, $s1, $s12, $s11;
|
||||
sub_u32 $s7, $s12, $s7;
|
||||
sub_u32 $s9, $s10, $s9;
|
||||
mad_u32 $s9, $s0, $s9, $s11;
|
||||
sub_u32 $s8, $s9, $s8;
|
||||
|
||||
// @BB0_3: // %for.cond32.preheader
|
||||
|
||||
@BB0_3:
|
||||
|
||||
cmp_gt_b1_u32 $c0, $s5, $s4;
|
||||
mov_b32 $s9, $s7;
|
||||
mov_b32 $s10, $s8;
|
||||
mov_b32 $s11, $s5;
|
||||
cbr_b1 $c0, @BB0_5;
|
||||
|
||||
// @BB0_4: // %for.body35
|
||||
|
||||
@BB0_4:
|
||||
|
||||
cvt_u64_u32 $d4, $s9;
|
||||
shl_u64 $d4, $d4, 2;
|
||||
add_u64 $d4, $d2, $d4;
|
||||
ld_global_align(4)_f32 $s12, [$d4];
|
||||
cvt_u64_u32 $d4, $s10;
|
||||
shl_u64 $d4, $d4, 2;
|
||||
add_u64 $d4, $d3, $d4;
|
||||
ld_global_align(4)_u32 $s13, [$d4];
|
||||
cvt_f32_u32 $s13, $s13;
|
||||
mul_ftz_f32 $s12, $s13, $s12;
|
||||
add_u32 $s9, $s9, $s1;
|
||||
add_u32 $s10, $s10, $s0;
|
||||
add_u32 $s11, $s11, 1;
|
||||
add_ftz_f32 $s6, $s6, $s12;
|
||||
cmp_le_b1_u32 $c0, $s11, $s4;
|
||||
cbr_b1 $c0, @BB0_4;
|
||||
|
||||
// @BB0_5: // %for.inc48
|
||||
|
||||
@BB0_5:
|
||||
|
||||
add_u32 $s7, $s7, 1;
|
||||
add_u32 $s8, $s8, 1;
|
||||
add_u32 $s3, $s3, 1;
|
||||
cmp_le_b1_u32 $c0, $s3, $s2;
|
||||
cbr_b1 $c0, @BB0_3;
|
||||
|
||||
// @BB0_6: // %for.end50
|
||||
|
||||
@BB0_6:
|
||||
|
||||
and_b64 $d0, $d0, 4294967295;
|
||||
shl_u64 $d0, $d0, 2;
|
||||
add_u64 $d0, $d1, $d0;
|
||||
add_ftz_f32 $s0, $s6, 0F3f000000;
|
||||
cvt_ftz_u32_f32 $s0, $s0;
|
||||
st_global_align(4)_u32 $s0, [$d0];
|
||||
ret;
|
||||
};
|
||||
@@ -0,0 +1,87 @@
|
||||
/**********************************************************************
|
||||
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_UTIL_HELPER_FUNCS_H_
|
||||
#define TEST_UTIL_HELPER_FUNCS_H_
|
||||
|
||||
#include <time.h>
|
||||
#include <cmath>
|
||||
#include <iostream>
|
||||
#include <sstream>
|
||||
#include <string>
|
||||
|
||||
static inline void Error(std::string error_msg) {
|
||||
std::cerr << "Error: " << error_msg << std::endl;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
void PrintArray(const std::string header, const T* data, const int width, const int height) {
|
||||
std::clog << header << " :\n";
|
||||
for (int i = 0; i < height; i++) {
|
||||
std::clog << "> ";
|
||||
for (int j = 0; j < width; j++) {
|
||||
std::clog << data[i * width + j] << " ";
|
||||
}
|
||||
std::clog << "\n";
|
||||
}
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
bool FillRandom(T* array_ptr, const int width, const int height, const T range_min,
|
||||
const T range_max, unsigned int seed = 123) {
|
||||
if (!array_ptr) {
|
||||
Error("Cannot fill array. NULL pointer.");
|
||||
return false;
|
||||
}
|
||||
|
||||
if (!seed) seed = (unsigned int)time(NULL);
|
||||
|
||||
srand(seed);
|
||||
double range = double(range_max - range_min) + 1.0;
|
||||
|
||||
/* random initialisation of input */
|
||||
for (int i = 0; i < height; i++)
|
||||
for (int j = 0; j < width; j++) {
|
||||
int index = i * width + j;
|
||||
array_ptr[index] = range_min + T(range * rand() / (RAND_MAX + 1.0));
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
template <typename T> T RoundToPowerOf2(T val) {
|
||||
int bytes = sizeof(T);
|
||||
|
||||
val--;
|
||||
for (int i = 0; i < bytes; i++) val |= val >> (1 << i);
|
||||
val++;
|
||||
|
||||
return val;
|
||||
}
|
||||
|
||||
template <typename T> bool IsPowerOf2(T val) {
|
||||
long long long_val = val;
|
||||
return (((long_val & (-long_val)) - long_val == 0) && (long_val != 0));
|
||||
}
|
||||
|
||||
#endif // TEST_UTIL_HELPER_FUNCS_H_
|
||||
@@ -0,0 +1,372 @@
|
||||
/**********************************************************************
|
||||
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:
|
||||
|
||||
<95> Redistributions of source code must retain the above copyright notice, this list of
|
||||
conditions and the following disclaimer.
|
||||
<95> 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 "util/hsa_rsrc_factory.h"
|
||||
|
||||
#include <hsa.h>
|
||||
#include <hsa_ext_finalize.h>
|
||||
#include <stdint.h>
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#include <cassert>
|
||||
#include <fstream>
|
||||
#include <iostream>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
// Callback function to find and bind kernarg region of an agent
|
||||
static hsa_status_t FindMemRegionsCallback(hsa_region_t region, void* data) {
|
||||
hsa_region_global_flag_t flags;
|
||||
hsa_region_segment_t segment_id;
|
||||
|
||||
hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &segment_id);
|
||||
if (segment_id != HSA_REGION_SEGMENT_GLOBAL) {
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
AgentInfo* agent_info = (AgentInfo*)data;
|
||||
hsa_region_get_info(region, HSA_REGION_INFO_GLOBAL_FLAGS, &flags);
|
||||
if (flags & HSA_REGION_GLOBAL_FLAG_COARSE_GRAINED) {
|
||||
agent_info->coarse_region = region;
|
||||
}
|
||||
|
||||
if (flags & HSA_REGION_GLOBAL_FLAG_KERNARG) {
|
||||
agent_info->kernarg_region = region;
|
||||
}
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
// Callback function to get the number of agents
|
||||
static hsa_status_t GetHsaAgentsCallback(hsa_agent_t agent, void* data) {
|
||||
// Copy handle of agent and increment number of agents reported
|
||||
HsaRsrcFactory* rsrcFactory = reinterpret_cast<HsaRsrcFactory*>(data);
|
||||
|
||||
// Determine if device is a Gpu agent
|
||||
hsa_status_t status;
|
||||
hsa_device_type_t type;
|
||||
status = hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &type);
|
||||
CHECK_STATUS("Error Calling hsa_agent_get_info", status);
|
||||
if (type == HSA_DEVICE_TYPE_DSP) {
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
if (type == HSA_DEVICE_TYPE_CPU) {
|
||||
AgentInfo* agent_info = reinterpret_cast<AgentInfo*>(malloc(sizeof(AgentInfo)));
|
||||
agent_info->dev_id = agent;
|
||||
agent_info->dev_type = HSA_DEVICE_TYPE_CPU;
|
||||
rsrcFactory->AddAgentInfo(agent_info, false);
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
// Device is a Gpu agent, build an instance of AgentInfo
|
||||
AgentInfo* agent_info = reinterpret_cast<AgentInfo*>(malloc(sizeof(AgentInfo)));
|
||||
agent_info->dev_id = agent;
|
||||
agent_info->dev_type = HSA_DEVICE_TYPE_GPU;
|
||||
hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, agent_info->name);
|
||||
agent_info->max_wave_size = 0;
|
||||
hsa_agent_get_info(agent, HSA_AGENT_INFO_WAVEFRONT_SIZE, &agent_info->max_wave_size);
|
||||
agent_info->max_queue_size = 0;
|
||||
hsa_agent_get_info(agent, HSA_AGENT_INFO_QUEUE_MAX_SIZE, &agent_info->max_queue_size);
|
||||
agent_info->profile = hsa_profile_t(108);
|
||||
hsa_agent_get_info(agent, HSA_AGENT_INFO_PROFILE, &agent_info->profile);
|
||||
|
||||
// Initialize memory regions to zero
|
||||
agent_info->kernarg_region.handle = 0;
|
||||
agent_info->coarse_region.handle = 0;
|
||||
|
||||
// Find and Bind Memory regions of the Gpu agent
|
||||
hsa_agent_iterate_regions(agent, FindMemRegionsCallback, agent_info);
|
||||
|
||||
// Save the instance of AgentInfo
|
||||
rsrcFactory->AddAgentInfo(agent_info, true);
|
||||
|
||||
return HSA_STATUS_SUCCESS;
|
||||
}
|
||||
|
||||
// Constructor of the class
|
||||
HsaRsrcFactory::HsaRsrcFactory() {
|
||||
// Initialize the Hsa Runtime
|
||||
hsa_status_t status = hsa_init();
|
||||
CHECK_STATUS("Error in hsa_init", status);
|
||||
|
||||
// Discover the set of Gpu devices available on the platform
|
||||
status = hsa_iterate_agents(GetHsaAgentsCallback, this);
|
||||
CHECK_STATUS("Error Calling hsa_iterate_agents", status);
|
||||
}
|
||||
|
||||
// Destructor of the class
|
||||
HsaRsrcFactory::~HsaRsrcFactory() {
|
||||
hsa_status_t status = hsa_shut_down();
|
||||
CHECK_STATUS("Error in hsa_shut_down", status);
|
||||
}
|
||||
|
||||
// Get the count of Hsa Gpu Agents available on the platform
|
||||
//
|
||||
// @return uint32_t Number of Gpu agents on platform
|
||||
//
|
||||
uint32_t HsaRsrcFactory::GetCountOfGpuAgents() { return uint32_t(gpu_list_.size()); }
|
||||
|
||||
// Get the count of Hsa Cpu Agents available on the platform
|
||||
//
|
||||
// @return uint32_t Number of Cpu agents on platform
|
||||
//
|
||||
uint32_t HsaRsrcFactory::GetCountOfCpuAgents() { return uint32_t(cpu_list_.size()); }
|
||||
|
||||
// Get the AgentInfo handle of a Gpu device
|
||||
//
|
||||
// @param idx Gpu Agent at specified index
|
||||
//
|
||||
// @param agent_info Output parameter updated with AgentInfo
|
||||
//
|
||||
// @return bool true if successful, false otherwise
|
||||
//
|
||||
bool HsaRsrcFactory::GetGpuAgentInfo(uint32_t idx, AgentInfo** agent_info) {
|
||||
// Determine if request is valid
|
||||
uint32_t size = uint32_t(gpu_list_.size());
|
||||
if (idx >= size) {
|
||||
return false;
|
||||
}
|
||||
|
||||
// Copy AgentInfo from specified index
|
||||
*agent_info = gpu_list_[idx];
|
||||
return true;
|
||||
}
|
||||
|
||||
// Get the AgentInfo handle of a Cpu device
|
||||
//
|
||||
// @param idx Cpu Agent at specified index
|
||||
//
|
||||
// @param agent_info Output parameter updated with AgentInfo
|
||||
//
|
||||
// @return bool true if successful, false otherwise
|
||||
//
|
||||
bool HsaRsrcFactory::GetCpuAgentInfo(uint32_t idx, AgentInfo** agent_info) {
|
||||
// Determine if request is valid
|
||||
uint32_t size = uint32_t(cpu_list_.size());
|
||||
if (idx >= size) {
|
||||
return false;
|
||||
}
|
||||
|
||||
// Copy AgentInfo from specified index
|
||||
*agent_info = cpu_list_[idx];
|
||||
return true;
|
||||
}
|
||||
|
||||
// Create a Queue object and return its handle. The queue object is expected
|
||||
// to support user requested number of Aql dispatch packets.
|
||||
//
|
||||
// @param agent_info Gpu Agent on which to create a queue object
|
||||
//
|
||||
// @param num_Pkts Number of packets to be held by queue
|
||||
//
|
||||
// @param queue Output parameter updated with handle of queue object
|
||||
//
|
||||
// @return bool true if successful, false otherwise
|
||||
//
|
||||
bool HsaRsrcFactory::CreateQueue(AgentInfo* agent_info, uint32_t num_pkts, hsa_queue_t** queue) {
|
||||
hsa_status_t status;
|
||||
status = hsa_queue_create(agent_info->dev_id, num_pkts, HSA_QUEUE_TYPE_MULTI, NULL, NULL,
|
||||
UINT32_MAX, UINT32_MAX, queue);
|
||||
return (status == HSA_STATUS_SUCCESS);
|
||||
}
|
||||
|
||||
// Create a Signal object and return its handle.
|
||||
//
|
||||
// @param value Initial value of signal object
|
||||
//
|
||||
// @param signal Output parameter updated with handle of signal object
|
||||
//
|
||||
// @return bool true if successful, false otherwise
|
||||
//
|
||||
bool HsaRsrcFactory::CreateSignal(uint32_t value, hsa_signal_t* signal) {
|
||||
hsa_status_t status;
|
||||
status = hsa_signal_create(value, 0, NULL, signal);
|
||||
return (status == HSA_STATUS_SUCCESS);
|
||||
}
|
||||
|
||||
// Allocate memory for use by a kernel of specified size in specified
|
||||
// agent's memory region. Currently supports Global segment whose Kernarg
|
||||
// flag set.
|
||||
//
|
||||
// @param agent_info Agent from whose memory region to allocate
|
||||
//
|
||||
// @param size Size of memory in terms of bytes
|
||||
//
|
||||
// @return uint8_t* Pointer to buffer, null if allocation fails.
|
||||
//
|
||||
uint8_t* HsaRsrcFactory::AllocateLocalMemory(const AgentInfo* agent_info, size_t size) {
|
||||
hsa_status_t status;
|
||||
uint8_t* buffer = NULL;
|
||||
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
|
||||
|
||||
if (agent_info->coarse_region.handle != 0) {
|
||||
// Allocate in local memory if it is available
|
||||
status = hsa_memory_allocate(agent_info->coarse_region, size, (void**)&buffer);
|
||||
if (status == HSA_STATUS_SUCCESS) {
|
||||
status = hsa_memory_assign_agent(buffer, agent_info->dev_id, HSA_ACCESS_PERMISSION_RW);
|
||||
}
|
||||
} else {
|
||||
// Allocate in system memory if local memory is not available
|
||||
status = hsa_memory_allocate(agent_info->kernarg_region, size, (void**)&buffer);
|
||||
}
|
||||
|
||||
return (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
|
||||
}
|
||||
|
||||
// Allocate memory tp pass kernel parameters.
|
||||
//
|
||||
// @param agent_info Agent from whose memory region to allocate
|
||||
//
|
||||
// @param size Size of memory in terms of bytes
|
||||
//
|
||||
// @return uint8_t* Pointer to buffer, null if allocation fails.
|
||||
//
|
||||
uint8_t* HsaRsrcFactory::AllocateSysMemory(const AgentInfo* agent_info, size_t size) {
|
||||
hsa_status_t status;
|
||||
size = (size + MEM_PAGE_MASK) & ~MEM_PAGE_MASK;
|
||||
|
||||
uint8_t* buffer = NULL;
|
||||
status = hsa_memory_allocate(agent_info->kernarg_region, size, (void**)&buffer);
|
||||
return (status == HSA_STATUS_SUCCESS) ? buffer : NULL;
|
||||
}
|
||||
|
||||
// Transfer data method
|
||||
bool HsaRsrcFactory::TransferData(void* dest_buff, void* src_buff, uint32_t length,
|
||||
bool host_to_dev) {
|
||||
hsa_status_t status;
|
||||
status = hsa_memory_copy(dest_buff, src_buff, length);
|
||||
return (status == HSA_STATUS_SUCCESS);
|
||||
}
|
||||
|
||||
// Loads an Assembled Brig file and Finalizes it into Device Isa
|
||||
//
|
||||
// @param agent_info Gpu device for which to finalize
|
||||
//
|
||||
// @param brig_path File path of the Assembled Brig file
|
||||
//
|
||||
// @param kernel_name Name of the kernel to finalize
|
||||
//
|
||||
// @param code_desc Handle of finalized Code Descriptor that could
|
||||
// be used to submit for execution
|
||||
//
|
||||
// @return bool true if successful, false otherwise
|
||||
//
|
||||
bool HsaRsrcFactory::LoadAndFinalize(AgentInfo* agent_info, const char* brig_path,
|
||||
char* kernel_name, hsa_executable_symbol_t* code_desc) {
|
||||
// Finalize the Hsail object into code object
|
||||
hsa_status_t status;
|
||||
hsa_code_object_t code_object;
|
||||
|
||||
// Build the code object filename
|
||||
std::string filename(brig_path);
|
||||
std::clog << "Code object filename: " << filename << std::endl;
|
||||
|
||||
// Open the file containing code object
|
||||
std::ifstream codeStream(filename.c_str(), std::ios::binary | std::ios::ate);
|
||||
if (!codeStream) {
|
||||
std::cerr << "Error: failed to load " << filename << std::endl;
|
||||
assert(false);
|
||||
return false;
|
||||
}
|
||||
|
||||
// Allocate memory to read in code object from file
|
||||
size_t size = std::string::size_type(codeStream.tellg());
|
||||
char* codeBuff = (char*)AllocateSysMemory(agent_info, size);
|
||||
if (!codeBuff) {
|
||||
std::cerr << "Error: failed to allocate memory for code object." << std::endl;
|
||||
assert(false);
|
||||
return false;
|
||||
}
|
||||
|
||||
// Read the code object into allocated memory
|
||||
codeStream.seekg(0, std::ios::beg);
|
||||
std::copy(std::istreambuf_iterator<char>(codeStream), std::istreambuf_iterator<char>(), codeBuff);
|
||||
|
||||
// De-Serialize the code object that has been read into memory
|
||||
status = hsa_code_object_deserialize(codeBuff, size, NULL, &code_object);
|
||||
if (status != HSA_STATUS_SUCCESS) {
|
||||
std::cerr << "Failed to deserialize code object" << std::endl;
|
||||
return false;
|
||||
}
|
||||
|
||||
// Create executable.
|
||||
hsa_executable_t hsaExecutable;
|
||||
status =
|
||||
hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, "", &hsaExecutable);
|
||||
CHECK_STATUS("Error in creating executable object", status);
|
||||
|
||||
// Load code object.
|
||||
status = hsa_executable_load_code_object(hsaExecutable, agent_info->dev_id, code_object, "");
|
||||
CHECK_STATUS("Error in loading executable object", status);
|
||||
|
||||
// Freeze executable.
|
||||
status = hsa_executable_freeze(hsaExecutable, "");
|
||||
CHECK_STATUS("Error in freezing executable object", status);
|
||||
|
||||
// Get symbol handle.
|
||||
hsa_executable_symbol_t kernelSymbol;
|
||||
status = hsa_executable_get_symbol(hsaExecutable, NULL, kernel_name, agent_info->dev_id, 0,
|
||||
&kernelSymbol);
|
||||
CHECK_STATUS("Error in looking up kernel symbol", status);
|
||||
|
||||
// Update output parameter
|
||||
*code_desc = kernelSymbol;
|
||||
return true;
|
||||
}
|
||||
|
||||
// Add an instance of AgentInfo representing a Hsa Gpu agent
|
||||
void HsaRsrcFactory::AddAgentInfo(AgentInfo* agent_info, bool gpu) {
|
||||
// Add input to Gpu list
|
||||
if (gpu) {
|
||||
gpu_list_.push_back(agent_info);
|
||||
return;
|
||||
}
|
||||
|
||||
// Add input to Cpu list
|
||||
cpu_list_.push_back(agent_info);
|
||||
}
|
||||
|
||||
// Print the various fields of Hsa Gpu Agents
|
||||
bool HsaRsrcFactory::PrintGpuAgents(const std::string& header) {
|
||||
std::clog << header << " :" << std::endl;
|
||||
|
||||
AgentInfo* agent_info;
|
||||
int size = uint32_t(gpu_list_.size());
|
||||
for (int idx = 0; idx < size; idx++) {
|
||||
agent_info = gpu_list_[idx];
|
||||
|
||||
std::clog << "> agent[" << idx << "] :" << std::endl;
|
||||
std::clog << ">> Name : " << agent_info->name << std::endl;
|
||||
std::clog << ">> Max Wave Size : " << agent_info->max_wave_size << std::endl;
|
||||
std::clog << ">> Max Queue Size : " << agent_info->max_queue_size << std::endl;
|
||||
std::clog << ">> Kernarg Region Id : " << agent_info->coarse_region.handle << std::endl;
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
||||
HsaRsrcFactory* HsaRsrcFactory::instance_ = NULL;
|
||||
HsaRsrcFactory::mutex_t HsaRsrcFactory::mutex_;
|
||||
@@ -0,0 +1,234 @@
|
||||
/**********************************************************************
|
||||
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:
|
||||
|
||||
<95> Redistributions of source code must retain the above copyright notice, this list of
|
||||
conditions and the following disclaimer.
|
||||
<95> 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_UTIL_HSA_RSRC_FACTORY_H_
|
||||
#define TEST_UTIL_HSA_RSRC_FACTORY_H_
|
||||
|
||||
#include <hsa.h>
|
||||
#include <hsa_ext_finalize.h>
|
||||
#include <stdint.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
#include <string.h>
|
||||
|
||||
#include <iostream>
|
||||
#include <mutex>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#define HSA_ARGUMENT_ALIGN_BYTES 16
|
||||
#define HSA_QUEUE_ALIGN_BYTES 64
|
||||
#define HSA_PACKET_ALIGN_BYTES 64
|
||||
|
||||
#define CHECK_STATUS(msg, status) \
|
||||
if (status != HSA_STATUS_SUCCESS) { \
|
||||
const char* emsg = 0; \
|
||||
hsa_status_string(status, &emsg); \
|
||||
printf("%s: %s\n", msg, emsg ? emsg : "<unknown error>"); \
|
||||
exit(1); \
|
||||
}
|
||||
|
||||
static const unsigned MEM_PAGE_BYTES = 0x1000;
|
||||
static const unsigned MEM_PAGE_MASK = MEM_PAGE_BYTES - 1;
|
||||
|
||||
// Encapsulates information about a Hsa Agent such as its
|
||||
// handle, name, max queue size, max wavefront size, etc.
|
||||
typedef struct {
|
||||
// Handle of Agent
|
||||
hsa_agent_t dev_id;
|
||||
|
||||
// Agent type - Cpu = 0, Gpu = 1 or Dsp = 2
|
||||
uint32_t dev_type;
|
||||
|
||||
// Name of Agent whose length is less than 64
|
||||
char name[64];
|
||||
|
||||
// Max size of Wavefront size
|
||||
uint32_t max_wave_size;
|
||||
|
||||
// Max size of Queue buffer
|
||||
uint32_t max_queue_size;
|
||||
|
||||
// Hsail profile supported by agent
|
||||
hsa_profile_t profile;
|
||||
|
||||
// Memory region supporting kernel parameters
|
||||
hsa_region_t coarse_region;
|
||||
|
||||
// Memory region supporting kernel arguments
|
||||
hsa_region_t kernarg_region;
|
||||
|
||||
} AgentInfo;
|
||||
|
||||
class HsaRsrcFactory {
|
||||
public:
|
||||
typedef std::recursive_mutex mutex_t;
|
||||
|
||||
static HsaRsrcFactory* Create() {
|
||||
std::lock_guard<mutex_t> lck(mutex_);
|
||||
if (HsaRsrcFactory::instance_ == NULL) {
|
||||
HsaRsrcFactory::instance_ = new HsaRsrcFactory();
|
||||
}
|
||||
return instance_;
|
||||
}
|
||||
|
||||
static void Destroy() {
|
||||
std::lock_guard<mutex_t> lck(mutex_);
|
||||
if (instance_) delete instance_;
|
||||
instance_ = NULL;
|
||||
}
|
||||
|
||||
static HsaRsrcFactory& Instance() {
|
||||
hsa_status_t status = (instance_ != NULL) ? HSA_STATUS_SUCCESS : HSA_STATUS_ERROR;
|
||||
CHECK_STATUS("HsaRsrcFactory::Instance()", status);
|
||||
return *instance_;
|
||||
}
|
||||
|
||||
// Get the count of Hsa Gpu Agents available on the platform
|
||||
//
|
||||
// @return uint32_t Number of Gpu agents on platform
|
||||
//
|
||||
uint32_t GetCountOfGpuAgents();
|
||||
|
||||
// Get the count of Hsa Cpu Agents available on the platform
|
||||
//
|
||||
// @return uint32_t Number of Cpu agents on platform
|
||||
//
|
||||
uint32_t GetCountOfCpuAgents();
|
||||
|
||||
// Get the AgentInfo handle of a Gpu device
|
||||
//
|
||||
// @param idx Gpu Agent at specified index
|
||||
//
|
||||
// @param agent_info Output parameter updated with AgentInfo
|
||||
//
|
||||
// @return bool true if successful, false otherwise
|
||||
//
|
||||
bool GetGpuAgentInfo(uint32_t idx, AgentInfo** agent_info);
|
||||
|
||||
// Get the AgentInfo handle of a Cpu device
|
||||
//
|
||||
// @param idx Cpu Agent at specified index
|
||||
//
|
||||
// @param agent_info Output parameter updated with AgentInfo
|
||||
//
|
||||
// @return bool true if successful, false otherwise
|
||||
//
|
||||
bool GetCpuAgentInfo(uint32_t idx, AgentInfo** agent_info);
|
||||
|
||||
// Create a Queue object and return its handle. The queue object is expected
|
||||
// to support user requested number of Aql dispatch packets.
|
||||
//
|
||||
// @param agent_info Gpu Agent on which to create a queue object
|
||||
//
|
||||
// @param num_Pkts Number of packets to be held by queue
|
||||
//
|
||||
// @param queue Output parameter updated with handle of queue object
|
||||
//
|
||||
// @return bool true if successful, false otherwise
|
||||
//
|
||||
bool CreateQueue(AgentInfo* agent_info, uint32_t num_pkts, hsa_queue_t** queue);
|
||||
|
||||
// Create a Signal object and return its handle.
|
||||
//
|
||||
// @param value Initial value of signal object
|
||||
//
|
||||
// @param signal Output parameter updated with handle of signal object
|
||||
//
|
||||
// @return bool true if successful, false otherwise
|
||||
//
|
||||
bool CreateSignal(uint32_t value, hsa_signal_t* signal);
|
||||
|
||||
// Allocate memory for use by a kernel of specified size in specified
|
||||
// agent's memory region. Currently supports Global segment whose Kernarg
|
||||
// flag set.
|
||||
//
|
||||
// @param agent_info Agent from whose memory region to allocate
|
||||
//
|
||||
// @param size Size of memory in terms of bytes
|
||||
//
|
||||
// @return uint8_t* Pointer to buffer, null if allocation fails.
|
||||
//
|
||||
uint8_t* AllocateLocalMemory(const AgentInfo* agent_info, size_t size);
|
||||
|
||||
// Allocate memory tp pass kernel parameters.
|
||||
//
|
||||
// @param agent_info Agent from whose memory region to allocate
|
||||
//
|
||||
// @param size Size of memory in terms of bytes
|
||||
//
|
||||
// @return uint8_t* Pointer to buffer, null if allocation fails.
|
||||
//
|
||||
uint8_t* AllocateSysMemory(const AgentInfo* agent_info, size_t size);
|
||||
|
||||
// Transfer data method
|
||||
bool TransferData(void* dest_buff, void* src_buff, uint32_t length, bool host_to_dev);
|
||||
|
||||
// Loads an Assembled Brig file and Finalizes it into Device Isa
|
||||
//
|
||||
// @param agent_info Gpu device for which to finalize
|
||||
//
|
||||
// @param brig_path File path of the Assembled Brig file
|
||||
//
|
||||
// @param kernel_name Name of the kernel to finalize
|
||||
//
|
||||
// @param code_desc Handle of finalized Code Descriptor that could
|
||||
// be used to submit for execution
|
||||
//
|
||||
// @return bool true if successful, false otherwise
|
||||
//
|
||||
bool LoadAndFinalize(AgentInfo* agent_info, const char* brig_path, char* kernel_name,
|
||||
hsa_executable_symbol_t* code_desc);
|
||||
|
||||
// Add an instance of AgentInfo representing a Hsa Gpu agent
|
||||
void AddAgentInfo(AgentInfo* agent_info, bool gpu);
|
||||
|
||||
// Print the various fields of Hsa Gpu Agents
|
||||
bool PrintGpuAgents(const std::string& header);
|
||||
|
||||
private:
|
||||
// Constructor of the class. Will initialize the Hsa Runtime and
|
||||
// query the system topology to get the list of Cpu and Gpu devices
|
||||
HsaRsrcFactory();
|
||||
|
||||
// Destructor of the class
|
||||
~HsaRsrcFactory();
|
||||
|
||||
static HsaRsrcFactory* instance_;
|
||||
static mutex_t mutex_;
|
||||
|
||||
// Used to maintain a list of Hsa Queue handles
|
||||
std::vector<hsa_queue_t*> queue_list_;
|
||||
|
||||
// Used to maintain a list of Hsa Signal handles
|
||||
std::vector<hsa_signal_t*> signal_list_;
|
||||
|
||||
// Used to maintain a list of Hsa Gpu Agent Info
|
||||
std::vector<AgentInfo*> gpu_list_;
|
||||
|
||||
// Used to maintain a list of Hsa Cpu Agent Info
|
||||
std::vector<AgentInfo*> cpu_list_;
|
||||
};
|
||||
|
||||
#endif // TEST_UTIL_HSA_RSRC_FACTORY_H_
|
||||
@@ -0,0 +1,181 @@
|
||||
/**********************************************************************
|
||||
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:
|
||||
|
||||
<95> Redistributions of source code must retain the above copyright notice, this list of
|
||||
conditions and the following disclaimer.
|
||||
<95> 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 "util/perf_timer.h"
|
||||
|
||||
PerfTimer::PerfTimer() { freq_in_100mhz_ = MeasureTSCFreqHz(); }
|
||||
|
||||
PerfTimer::~PerfTimer() {
|
||||
while (!timers_.empty()) {
|
||||
Timer* temp = timers_.back();
|
||||
timers_.pop_back();
|
||||
delete temp;
|
||||
}
|
||||
}
|
||||
|
||||
// New cretaed timer instantance index will be returned
|
||||
int PerfTimer::CreateTimer() {
|
||||
Timer* newTimer = new Timer;
|
||||
newTimer->start = 0;
|
||||
newTimer->clocks = 0;
|
||||
|
||||
#ifdef _WIN32
|
||||
QueryPerformanceFrequency((LARGE_INTEGER*)&newTimer->freq);
|
||||
#else
|
||||
newTimer->freq = (long long)1.0E3;
|
||||
#endif
|
||||
|
||||
/* Push back the address of new Timer instance created */
|
||||
timers_.push_back(newTimer);
|
||||
return (int)(timers_.size() - 1);
|
||||
}
|
||||
|
||||
int PerfTimer::StartTimer(int index) {
|
||||
if (index >= (int)timers_.size()) {
|
||||
Error("Cannot reset timer. Invalid handle.");
|
||||
return FAILURE;
|
||||
}
|
||||
|
||||
#ifdef _WIN32
|
||||
// General Windows timing method
|
||||
#ifndef _AMD
|
||||
long long tmpStart;
|
||||
QueryPerformanceCounter((LARGE_INTEGER*)&(tmpStart));
|
||||
timers_[index]->start = (double)tmpStart;
|
||||
#else
|
||||
// AMD Windows timing method
|
||||
#endif
|
||||
#else
|
||||
// General Linux timing method
|
||||
#ifndef _AMD
|
||||
struct timeval s;
|
||||
gettimeofday(&s, 0);
|
||||
timers_[index]->start = s.tv_sec * 1.0E3 + ((double)(s.tv_usec / 1.0E3));
|
||||
#else
|
||||
// AMD timing method
|
||||
unsigned int unused;
|
||||
timers_[index]->start = __rdtscp(&unused);
|
||||
#endif
|
||||
#endif
|
||||
|
||||
return SUCCESS;
|
||||
}
|
||||
|
||||
|
||||
int PerfTimer::StopTimer(int index) {
|
||||
double n = 0;
|
||||
if (index >= (int)timers_.size()) {
|
||||
Error("Cannot reset timer. Invalid handle.");
|
||||
return FAILURE;
|
||||
}
|
||||
#ifdef _WIN32
|
||||
#ifndef _AMD
|
||||
long long n1;
|
||||
QueryPerformanceCounter((LARGE_INTEGER*)&(n1));
|
||||
n = (double)n1;
|
||||
#else
|
||||
// AMD Window Timing
|
||||
#endif
|
||||
|
||||
#else
|
||||
// General Linux timing method
|
||||
#ifndef _AMD
|
||||
struct timeval s;
|
||||
gettimeofday(&s, 0);
|
||||
n = s.tv_sec * 1.0E3 + (double)(s.tv_usec / 1.0E3);
|
||||
#else
|
||||
// AMD Linux timing
|
||||
unsigned int unused;
|
||||
n = __rdtscp(&unused);
|
||||
#endif
|
||||
#endif
|
||||
|
||||
n -= timers_[index]->start;
|
||||
timers_[index]->start = 0;
|
||||
|
||||
#ifndef _AMD
|
||||
timers_[index]->clocks += n;
|
||||
#else
|
||||
// timers_[index]->clocks += 10 * n / freq_in_100mhz_; // unit is ns
|
||||
timers_[index]->clocks += 1.0E-6 * 10 * n / freq_in_100mhz_; // convert to ms
|
||||
#endif
|
||||
|
||||
return SUCCESS;
|
||||
}
|
||||
|
||||
void PerfTimer::Error(std::string str) { std::cout << str << std::endl; }
|
||||
|
||||
|
||||
double PerfTimer::ReadTimer(int index) {
|
||||
if (index >= (int)timers_.size()) {
|
||||
Error("Cannot read timer. Invalid handle.");
|
||||
return FAILURE;
|
||||
}
|
||||
|
||||
double reading = double(timers_[index]->clocks);
|
||||
|
||||
reading = double(reading / timers_[index]->freq);
|
||||
|
||||
return reading;
|
||||
}
|
||||
|
||||
|
||||
uint64_t PerfTimer::CoarseTimestampUs() {
|
||||
#ifdef _WIN32
|
||||
uint64_t freqHz, ticks;
|
||||
QueryPerformanceFrequency((LARGE_INTEGER*)&freqHz);
|
||||
QueryPerformanceCounter((LARGE_INTEGER*)&ticks);
|
||||
|
||||
// Scale numerator and divisor until (ticks * 1000000) fits in uint64_t.
|
||||
while (ticks > (1ULL << 44)) {
|
||||
ticks /= 16;
|
||||
freqHz /= 16;
|
||||
}
|
||||
|
||||
return (ticks * 1000000) / freqHz;
|
||||
#else
|
||||
struct timespec ts;
|
||||
clock_gettime(CLOCK_MONOTONIC_RAW, &ts);
|
||||
return uint64_t(ts.tv_sec) * 1000000 + ts.tv_nsec / 1000;
|
||||
#endif
|
||||
}
|
||||
|
||||
uint64_t PerfTimer::MeasureTSCFreqHz() {
|
||||
// Make a coarse interval measurement of TSC ticks for 1 gigacycles.
|
||||
unsigned int unused;
|
||||
uint64_t tscTicksEnd;
|
||||
|
||||
uint64_t coarseBeginUs = CoarseTimestampUs();
|
||||
uint64_t tscTicksBegin = __rdtscp(&unused);
|
||||
do {
|
||||
tscTicksEnd = __rdtscp(&unused);
|
||||
} while (tscTicksEnd - tscTicksBegin < 1000000000);
|
||||
|
||||
uint64_t coarseEndUs = CoarseTimestampUs();
|
||||
|
||||
// Compute the TSC frequency and round to nearest 100MHz.
|
||||
uint64_t coarseIntervalNs = (coarseEndUs - coarseBeginUs) * 1000;
|
||||
uint64_t tscIntervalTicks = tscTicksEnd - tscTicksBegin;
|
||||
return (tscIntervalTicks * 10 + (coarseIntervalNs / 2)) / coarseIntervalNs;
|
||||
}
|
||||
@@ -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:
|
||||
|
||||
<95> Redistributions of source code must retain the above copyright notice, this list of
|
||||
conditions and the following disclaimer.
|
||||
<95> 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_UTIL_PERF_TIMER_H_
|
||||
#define TEST_UTIL_PERF_TIMER_H_
|
||||
|
||||
// Will use AMD timer or general Linux timer based on compilation flag
|
||||
// Need to consider platform is Windows or Linux
|
||||
|
||||
#include <stdint.h>
|
||||
#include <stdio.h>
|
||||
#include <stdlib.h>
|
||||
|
||||
#if defined(_MSC_VER)
|
||||
#include <intrin.h>
|
||||
#include <time.h>
|
||||
#include <windows.h>
|
||||
#else
|
||||
#if defined(__GNUC__)
|
||||
#include <sys/time.h>
|
||||
#include <x86intrin.h>
|
||||
#endif // __GNUC__
|
||||
#endif // _MSC_VER
|
||||
|
||||
#include <iostream>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
class PerfTimer {
|
||||
public:
|
||||
enum { SUCCESS = 0, FAILURE = 1 };
|
||||
|
||||
PerfTimer();
|
||||
~PerfTimer();
|
||||
|
||||
// General Linux timing method
|
||||
int CreateTimer();
|
||||
int StartTimer(int index);
|
||||
int StopTimer(int index);
|
||||
|
||||
// retrieve time
|
||||
double ReadTimer(int index);
|
||||
// write into a file
|
||||
double WriteTimer(int index);
|
||||
|
||||
private:
|
||||
struct Timer {
|
||||
std::string name; /* name of time object */
|
||||
long long freq; /* frequency */
|
||||
double clocks; /* number of ticks at end */
|
||||
double start; /* start point ticks */
|
||||
};
|
||||
|
||||
std::vector<Timer*> timers_; /* vector to Timer objects */
|
||||
double freq_in_100mhz_;
|
||||
|
||||
// AMD timing method
|
||||
uint64_t CoarseTimestampUs();
|
||||
uint64_t MeasureTSCFreqHz();
|
||||
|
||||
void Error(std::string str);
|
||||
};
|
||||
|
||||
#endif // TEST_UTIL_PERF_TIMER_H_
|
||||
@@ -0,0 +1,52 @@
|
||||
/******************************************************************************
|
||||
|
||||
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_ASSERT_H_
|
||||
#define TEST_CTRL_TEST_ASSERT_H_
|
||||
|
||||
#define TEST_ASSERT(cond) \
|
||||
{ \
|
||||
if (!(cond)) { \
|
||||
std::cerr << "Assert failed(" << #cond << ") at " << __FILE__ << ", line " << __LINE__ \
|
||||
<< std::endl; \
|
||||
exit(-1); \
|
||||
} \
|
||||
}
|
||||
|
||||
#define TEST_STATUS(cond) \
|
||||
{ \
|
||||
if (!(cond)) { \
|
||||
std::cerr << "Test error at " << __FILE__ << ", line " << __LINE__ \
|
||||
<< std::endl; \
|
||||
const char* message; \
|
||||
rocprofiler_error_string(&message); \
|
||||
std::cerr << "ERROR: " << message << std::endl; \
|
||||
exit(-1); \
|
||||
} \
|
||||
}
|
||||
|
||||
#endif // TEST_CTRL_TEST_ASSERT_H_
|
||||
@@ -0,0 +1,221 @@
|
||||
#ifndef TEST_UTIL_XML_H_
|
||||
#define TEST_UTIL_XML_H_
|
||||
|
||||
#include <fcntl.h>
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
#include <sys/stat.h>
|
||||
#include <sys/types.h>
|
||||
#include <unistd.h>
|
||||
|
||||
#include <fstream>
|
||||
#include <iostream>
|
||||
#include <map>
|
||||
#include <vector>
|
||||
|
||||
namespace xml {
|
||||
|
||||
class Xml {
|
||||
public:
|
||||
typedef std::vector<char> token_t;
|
||||
struct level_t {
|
||||
std::string tag;
|
||||
std::vector<level_t*> nodes;
|
||||
std::map<std::string, std::string> opts;
|
||||
};
|
||||
typedef std::vector<level_t*> nodes_vec_t;
|
||||
|
||||
enum {
|
||||
DECL_STATE,
|
||||
BODY_STATE
|
||||
};
|
||||
|
||||
Xml(const char* file_name) :
|
||||
file_name_(file_name),
|
||||
file_line_(0),
|
||||
data_size_(0),
|
||||
index_(0),
|
||||
state_(BODY_STATE),
|
||||
level_(NULL),
|
||||
comment_(false)
|
||||
{
|
||||
AddLevel("top");
|
||||
|
||||
fd_ = open(file_name, O_RDONLY);
|
||||
if (fd_ == -1) {
|
||||
std::cout << "XML file not found: " << file_name << std::endl;
|
||||
return;
|
||||
}
|
||||
|
||||
token_t remainder;
|
||||
while (1) {
|
||||
token_t token = (remainder.size()) ? remainder : NextToken();
|
||||
remainder.clear();
|
||||
// End of file
|
||||
if (token.size() == 0) break;
|
||||
|
||||
// token_t token1 = token;
|
||||
// token1.push_back('\0');
|
||||
// std::cout << "> " << &token1[0] << std::endl;
|
||||
|
||||
switch(state_) {
|
||||
case BODY_STATE:
|
||||
if (token[0] == '<') {
|
||||
bool node_begin = true;
|
||||
unsigned ind = 1;
|
||||
if (token[1] == '/') {
|
||||
node_begin = false;
|
||||
++ind;
|
||||
}
|
||||
|
||||
unsigned i = ind;
|
||||
while (i < token.size()) { if (token[i] == '>') break; ++i; }
|
||||
for (unsigned j = i + 1; j < token.size(); ++j) remainder.push_back(token[j]);
|
||||
|
||||
if (i == token.size()) {
|
||||
if (node_begin) state_ = DECL_STATE;
|
||||
else BadFormat(token);
|
||||
token.push_back('\0');
|
||||
} else token[i] = '\0';
|
||||
|
||||
const char* tag = strdup(&token[ind]);
|
||||
if (node_begin) {
|
||||
AddLevel(tag);
|
||||
} else {
|
||||
if (strncmp(CurrentLevel().c_str(), tag, strlen(tag))) {
|
||||
token.back() = '>';
|
||||
BadFormat(token);
|
||||
}
|
||||
UpLevel();
|
||||
}
|
||||
} else BadFormat(token);
|
||||
break;
|
||||
case DECL_STATE:
|
||||
if (token[0] == '>') {
|
||||
state_ = BODY_STATE;
|
||||
for (unsigned j = 1; j < token.size(); ++j) remainder.push_back(token[j]);
|
||||
continue;
|
||||
} else {
|
||||
token.push_back('\0');
|
||||
unsigned j = 0;
|
||||
for (j = 0; j < token.size(); ++j) if (token[j] == '=') break;
|
||||
if (j == token.size()) BadFormat(token);
|
||||
token[j] = '\0';
|
||||
const char* key = &token[0];
|
||||
const char* value = &token[j + 1];
|
||||
AddOption(key, value);
|
||||
}
|
||||
break;
|
||||
default:
|
||||
std::cout << "Wrong state: " << state_ << std::endl;
|
||||
exit(1);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
std::vector<level_t*> GetNodes(std::string global_tag) {
|
||||
return map_[global_tag];
|
||||
}
|
||||
|
||||
void Print() const {
|
||||
for(auto& elem : map_) {
|
||||
for (auto node : elem.second) {
|
||||
if (node->opts.size()) {
|
||||
std::cout << elem.first << ":" << std::endl;
|
||||
for (auto& opt : node->opts) {
|
||||
std::cout << " " << opt.first << " = " << opt.second << std::endl;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
bool LineEndCheck() {
|
||||
bool found = false;
|
||||
if (buffer_[index_] == '\n') {
|
||||
buffer_[index_] = ' ';
|
||||
++file_line_;
|
||||
found = true;
|
||||
comment_ = false;
|
||||
} else if (comment_ || (buffer_[index_] == '#')) {
|
||||
found = true;
|
||||
comment_ = true;
|
||||
}
|
||||
return found;
|
||||
}
|
||||
|
||||
token_t NextToken() {
|
||||
token_t token;
|
||||
|
||||
while (1) {
|
||||
if (data_size_ == 0) {
|
||||
data_size_ = read(fd_, buffer_, buf_size_);
|
||||
if (data_size_ <= 0) break;
|
||||
}
|
||||
if (token.empty()) while ((index_ < data_size_) && ((buffer_[index_] == ' ') || LineEndCheck())) {
|
||||
++index_;
|
||||
}
|
||||
while ((index_ < data_size_) && (buffer_[index_] != ' ') && !LineEndCheck()) {
|
||||
token.push_back(buffer_[index_++]);
|
||||
}
|
||||
if (index_ == data_size_) {
|
||||
index_ = 0;
|
||||
data_size_ = 0;
|
||||
} else break;
|
||||
}
|
||||
|
||||
return token;
|
||||
}
|
||||
|
||||
void BadFormat(token_t token) {
|
||||
token.push_back('\0');
|
||||
std::cout << "Error: " << file_name_ << ", line " << file_line_ << ", bad XML token '" << &token[0] << "'" << std::endl;
|
||||
exit(1);
|
||||
}
|
||||
|
||||
void AddLevel(const std::string& tag) {
|
||||
level_t* level = new level_t;
|
||||
level->tag = tag;
|
||||
if (level_) {
|
||||
level_->nodes.push_back(level);
|
||||
stack_.push_back(level_);
|
||||
}
|
||||
level_ = level;
|
||||
|
||||
std::string global_tag;
|
||||
for (level_t* level : stack_) { global_tag += level->tag + "."; }
|
||||
global_tag += tag;
|
||||
map_[global_tag].push_back(level_);
|
||||
}
|
||||
|
||||
void UpLevel() {
|
||||
level_ = stack_.back();
|
||||
stack_.pop_back();
|
||||
}
|
||||
|
||||
std::string CurrentLevel() const {
|
||||
return level_->tag;
|
||||
}
|
||||
|
||||
void AddOption(const std::string& key, const std::string& value) {
|
||||
level_->opts[key] = value;
|
||||
}
|
||||
|
||||
const char* file_name_;
|
||||
unsigned file_line_;
|
||||
int fd_;
|
||||
static const unsigned buf_size_ = 256;
|
||||
char buffer_[buf_size_];
|
||||
unsigned data_size_;
|
||||
unsigned index_;
|
||||
unsigned state_;
|
||||
level_t* level_;
|
||||
std::vector<level_t*> stack_;
|
||||
std::map<std::string, nodes_vec_t> map_;
|
||||
bool comment_;
|
||||
};
|
||||
|
||||
} // namespace xml
|
||||
|
||||
#endif // TEST_UTIL_XML_H_
|
||||
Reference in New Issue
Block a user