rpl_run w/o input file; queue create callback; test for n gpus and n threads
Change-Id: I37157c49cf6454de591cae97b5cc43287ea95956
Этот коммит содержится в:
Evgeny
+15
-10
@@ -276,16 +276,17 @@ if [ "$ARG_CK" = "-" ] ; then
|
||||
fi
|
||||
|
||||
if [ -z "$INPUT_FILE" ] ; then
|
||||
fatal "Need input file"
|
||||
input_base="results"
|
||||
input_type="none"
|
||||
else
|
||||
input_base=`echo "$INPUT_FILE" | sed "s/^\(.*\)\.\([^\.]*\)$/\1/"`
|
||||
input_type=`echo "$INPUT_FILE" | sed "s/^\(.*\)\.\([^\.]*\)$/\2/"`
|
||||
if [ -z "${input_base}" -o -z "${input_type}" ] ; then
|
||||
fatal "Bad input file '$INPUT_FILE'"
|
||||
fi
|
||||
input_base=`basename $input_base`
|
||||
fi
|
||||
|
||||
input_base=`echo "$INPUT_FILE" | sed "s/^\(.*\)\.\([^\.]*\)$/\1/"`
|
||||
input_type=`echo "$INPUT_FILE" | sed "s/^\(.*\)\.\([^\.]*\)$/\2/"`
|
||||
if [ -z "${input_base}" -o -z "${input_type}" ] ; then
|
||||
fatal "Bad input file '$INPUT_FILE'"
|
||||
fi
|
||||
input_base=`basename $input_base`
|
||||
|
||||
if [ "$OUTPUT_DIR" = "--" ] ; then
|
||||
fatal "Bad output dir '$OUTPUT_DIR'"
|
||||
fi
|
||||
@@ -309,7 +310,7 @@ input_list=""
|
||||
RES_DIR=""
|
||||
if [ "$input_type" = "xml" ] ; then
|
||||
input_list=$INPUT_FILE
|
||||
elif [ "$input_type" = "txt" ] ; then
|
||||
elif [ "$input_type" = "txt" -o "$input_type" = "none" ] ; then
|
||||
OUTPUT_DIR="-"
|
||||
RES_DIR=$DATA_PATH/$DATA_DIR
|
||||
if [ -e $RES_DIR ] ; then
|
||||
@@ -317,7 +318,11 @@ elif [ "$input_type" = "txt" ] ; then
|
||||
fi
|
||||
mkdir -p $RES_DIR
|
||||
echo "RPL: output dir '$RES_DIR'"
|
||||
$BIN_DIR/txt2xml.sh $INPUT_FILE $RES_DIR
|
||||
if [ "$input_type" = "txt" ] ; then
|
||||
$BIN_DIR/txt2xml.sh $INPUT_FILE $RES_DIR
|
||||
else
|
||||
echo "<metric></metric>" > $RES_DIR/input.xml
|
||||
fi
|
||||
input_list=`/bin/ls $RES_DIR/input*.xml`
|
||||
export ROCPROFILER_SESS=$RES_DIR
|
||||
else
|
||||
|
||||
@@ -46,7 +46,7 @@ THE SOFTWARE.
|
||||
#include <hsa_ven_amd_aqlprofile.h>
|
||||
#include <stdint.h>
|
||||
|
||||
#define ROCPROFILER_VERSION_MAJOR 4
|
||||
#define ROCPROFILER_VERSION_MAJOR 5
|
||||
#define ROCPROFILER_VERSION_MINOR 0
|
||||
|
||||
#ifdef __cplusplus
|
||||
@@ -234,6 +234,7 @@ typedef hsa_status_t (*rocprofiler_callback_t)(
|
||||
// Queue callbacks
|
||||
typedef struct {
|
||||
rocprofiler_callback_t dispatch; // dispatch callback
|
||||
hsa_status_t (*create)(hsa_queue_t* queue, void* data); // create callback
|
||||
hsa_status_t (*destroy)(hsa_queue_t* queue, void* data); // destroy callback
|
||||
} rocprofiler_queue_callbacks_t;
|
||||
|
||||
|
||||
@@ -30,13 +30,14 @@ void InterceptQueue::HsaIntercept(HsaApiTable* table) {
|
||||
|
||||
InterceptQueue::mutex_t InterceptQueue::mutex_;
|
||||
rocprofiler_callback_t InterceptQueue::dispatch_callback_ = NULL;
|
||||
InterceptQueue::queue_callback_t InterceptQueue::create_callback_ = NULL;
|
||||
InterceptQueue::queue_callback_t InterceptQueue::destroy_callback_ = NULL;
|
||||
void* InterceptQueue::callback_data_ = NULL;
|
||||
InterceptQueue::obj_map_t* InterceptQueue::obj_map_ = NULL;
|
||||
const char* InterceptQueue::kernel_none_ = "";
|
||||
Tracker* InterceptQueue::tracker_ = NULL;
|
||||
bool InterceptQueue::tracker_on_ = false;
|
||||
bool InterceptQueue::in_constr_call_ = false;
|
||||
bool InterceptQueue::in_create_call_ = false;
|
||||
InterceptQueue::queue_id_t InterceptQueue::current_queue_id = 0;
|
||||
|
||||
} // namespace rocprofiler
|
||||
|
||||
@@ -63,8 +63,8 @@ class InterceptQueue {
|
||||
std::lock_guard<mutex_t> lck(mutex_);
|
||||
hsa_status_t status = HSA_STATUS_ERROR;
|
||||
|
||||
if (in_constr_call_) EXC_ABORT(status, "recursive InterceptQueueCreate()");
|
||||
in_constr_call_ = true;
|
||||
if (in_create_call_) EXC_ABORT(status, "recursive InterceptQueueCreate()");
|
||||
in_create_call_ = true;
|
||||
|
||||
ProxyQueue* proxy = ProxyQueue::Create(agent, size, type, queue_event_callback, data, private_segment_size,
|
||||
group_segment_size, queue, &status);
|
||||
@@ -84,7 +84,11 @@ class InterceptQueue {
|
||||
obj->queue_id = current_queue_id;
|
||||
++current_queue_id;
|
||||
|
||||
in_constr_call_ = false;
|
||||
if (create_callback_ != NULL) {
|
||||
status = create_callback_(*queue, callback_data_);
|
||||
}
|
||||
|
||||
in_create_call_ = false;
|
||||
return status;
|
||||
}
|
||||
|
||||
@@ -208,10 +212,15 @@ class InterceptQueue {
|
||||
}
|
||||
}
|
||||
|
||||
static void SetCallbacks(rocprofiler_callback_t dispatch_callback, queue_callback_t destroy_callback, void* data) {
|
||||
static void SetCallbacks(rocprofiler_callback_t dispatch_callback,
|
||||
queue_callback_t create_callback,
|
||||
queue_callback_t destroy_callback,
|
||||
void* data)
|
||||
{
|
||||
std::lock_guard<mutex_t> lck(mutex_);
|
||||
callback_data_ = data;
|
||||
dispatch_callback_ = dispatch_callback;
|
||||
create_callback_ = create_callback;
|
||||
destroy_callback_ = destroy_callback;
|
||||
}
|
||||
|
||||
@@ -303,13 +312,14 @@ class InterceptQueue {
|
||||
static mutex_t mutex_;
|
||||
static const packet_word_t header_type_mask = (1ul << HSA_PACKET_HEADER_WIDTH_TYPE) - 1;
|
||||
static rocprofiler_callback_t dispatch_callback_;
|
||||
static queue_callback_t create_callback_;
|
||||
static queue_callback_t destroy_callback_;
|
||||
static void* callback_data_;
|
||||
static obj_map_t* obj_map_;
|
||||
static const char* kernel_none_;
|
||||
static Tracker* tracker_;
|
||||
static bool tracker_on_;
|
||||
static bool in_constr_call_;
|
||||
static bool in_create_call_;
|
||||
static queue_id_t current_queue_id;
|
||||
|
||||
hsa_queue_t* const queue_;
|
||||
|
||||
@@ -523,14 +523,14 @@ PUBLIC_API hsa_status_t rocprofiler_get_metrics(const rocprofiler_t* handle) {
|
||||
// Set/remove queue callbacks
|
||||
PUBLIC_API hsa_status_t rocprofiler_set_queue_callbacks(rocprofiler_queue_callbacks_t callbacks, void* data) {
|
||||
API_METHOD_PREFIX
|
||||
rocprofiler::InterceptQueue::SetCallbacks(callbacks.dispatch, callbacks.destroy, data);
|
||||
rocprofiler::InterceptQueue::SetCallbacks(callbacks.dispatch, callbacks.create, callbacks.destroy, data);
|
||||
API_METHOD_SUFFIX
|
||||
}
|
||||
|
||||
// Remove queue callbacks
|
||||
PUBLIC_API hsa_status_t rocprofiler_remove_queue_callbacks() {
|
||||
API_METHOD_PREFIX
|
||||
rocprofiler::InterceptQueue::SetCallbacks(NULL, NULL, NULL);
|
||||
rocprofiler::InterceptQueue::SetCallbacks(NULL, NULL, NULL, NULL);
|
||||
API_METHOD_SUFFIX
|
||||
}
|
||||
|
||||
|
||||
@@ -145,8 +145,7 @@ int main() {
|
||||
TEST_STATUS(status == HSA_STATUS_SUCCESS);
|
||||
|
||||
// Test initialization
|
||||
TestHsa::SetQueue(prof_queue);
|
||||
TestHsa::HsaInstantiate(0);
|
||||
TestHsa::HsaInstantiate();
|
||||
|
||||
// Dispatching profiled kernel n-times to collect all counter groups data
|
||||
const unsigned group_n = 0;
|
||||
@@ -157,9 +156,9 @@ int main() {
|
||||
for (unsigned ind = 0; ind < 3; ++ind) {
|
||||
#if 1
|
||||
const unsigned queue_ind = ind % queue_count;
|
||||
TestHsa::SetQueue(queue[queue_ind]);
|
||||
// ret_val = RunKernel<DummyKernel, TestAql>();
|
||||
ret_val = RunKernel<SimpleConvolution, TestAql>();
|
||||
hsa_queue_t* prof_queue = queue[queue_ind];
|
||||
//ret_val = RunKernel<DummyKernel, TestAql>(0, NULL, NULL, prof_queue);
|
||||
ret_val = RunKernel<SimpleConvolution, TestAql>(0, NULL, NULL, prof_queue);
|
||||
std::cout << "run kernel, queue " << queue_ind << std::endl;
|
||||
#else
|
||||
sleep(3);
|
||||
|
||||
@@ -21,20 +21,66 @@ THE SOFTWARE.
|
||||
*******************************************************************************/
|
||||
|
||||
#include <hsa.h>
|
||||
#include <stdio.h>
|
||||
#include <string.h>
|
||||
#include <iostream>
|
||||
#include <thread>
|
||||
|
||||
#include "ctrl/run_kernel.h"
|
||||
#include "ctrl/test_aql.h"
|
||||
#include "dummy_kernel/dummy_kernel.h"
|
||||
#include "simple_convolution/simple_convolution.h"
|
||||
|
||||
void thread_fun(const int kiter, const int diter, const uint32_t agents_number) {
|
||||
const AgentInfo* agent_info[agents_number];
|
||||
hsa_queue_t* queue[agents_number];
|
||||
HsaRsrcFactory* rsrc = &HsaRsrcFactory::Instance();
|
||||
|
||||
for (uint32_t n = 0; n < agents_number; ++n) {
|
||||
uint32_t agent_id = n % rsrc->GetCountOfGpuAgents();
|
||||
if (rsrc->GetGpuAgentInfo(agent_id, &agent_info[n]) == false) {
|
||||
fprintf(stderr, "AgentInfo failed\n");
|
||||
abort();
|
||||
}
|
||||
if (rsrc->CreateQueue(agent_info[n], 128, &queue[n]) == false) {
|
||||
fprintf(stderr, "CreateQueue failed\n");
|
||||
abort();
|
||||
}
|
||||
}
|
||||
|
||||
for (int i = 0; i < kiter; ++i) {
|
||||
for (uint32_t n = 0; n < agents_number; ++n) {
|
||||
// RunKernel<DummyKernel, TestAql>(0, NULL, agent_info[n], queue[n], diter);
|
||||
RunKernel<SimpleConvolution, TestAql>(0, NULL, agent_info[n], queue[n], diter);
|
||||
}
|
||||
}
|
||||
|
||||
for (uint32_t n = 0; n < agents_number; ++n) {
|
||||
hsa_queue_destroy(queue[n]);
|
||||
}
|
||||
}
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
const char* kiter_s = getenv("ROCP_KITER");
|
||||
const char* diter_s = getenv("ROCP_DITER");
|
||||
const char* agents_s = getenv("ROCP_AGENTS");
|
||||
const char* thrs_s = getenv("ROCP_THRS");
|
||||
|
||||
const int kiter = (kiter_s != NULL) ? atol(kiter_s) : 1;
|
||||
const int diter = (diter_s != NULL) ? atol(diter_s) : 1;
|
||||
const uint32_t agents_number = (agents_s != NULL) ? (uint32_t)atol(agents_s) : 1;
|
||||
const int thrs = (thrs_s != NULL) ? atol(thrs_s) : 1;
|
||||
|
||||
TestHsa::HsaInstantiate();
|
||||
for (int i = 0; i < kiter; ++i) RunKernel<SimpleConvolution, TestAql>(argc, argv, diter);
|
||||
|
||||
std::thread t[thrs];
|
||||
for (int n = 0; n < thrs; ++n) {
|
||||
t[n] = std::thread(thread_fun, kiter, diter, agents_number);
|
||||
}
|
||||
for (int n = 0; n < thrs; ++n) {
|
||||
t[n].join();
|
||||
}
|
||||
|
||||
TestHsa::HsaShutdown();
|
||||
return 0;
|
||||
}
|
||||
|
||||
@@ -26,7 +26,7 @@ THE SOFTWARE.
|
||||
#include "ctrl/test_hsa.h"
|
||||
#include "util/test_assert.h"
|
||||
|
||||
template <class Kernel, class Test> bool RunKernel(int argc = 0, char* argv[] = NULL, int count = 1) {
|
||||
template <class Kernel, class Test> bool RunKernel(int argc = 0, char* argv[] = NULL, const AgentInfo* agent_info = NULL, hsa_queue_t* queue = NULL, int count = 1) {
|
||||
bool ret_val = false;
|
||||
|
||||
if (getenv("ROC_TEST_TRACE") == NULL) std::clog.rdbuf(NULL);
|
||||
@@ -34,8 +34,12 @@ template <class Kernel, class Test> bool RunKernel(int argc = 0, char* argv[] =
|
||||
|
||||
// Create test kernel object
|
||||
Kernel test_kernel;
|
||||
TestAql* test_aql = new TestHsa(&test_kernel);
|
||||
test_aql = new Test(test_aql);
|
||||
|
||||
TestHsa* test_hsa = new TestHsa(&test_kernel);
|
||||
test_hsa->SetAgentInfo(agent_info);
|
||||
test_hsa->SetQueue(queue);
|
||||
|
||||
TestAql* test_aql = new Test(test_hsa);
|
||||
TEST_ASSERT(test_aql != NULL);
|
||||
if (test_aql == NULL) return 1;
|
||||
|
||||
|
||||
@@ -29,60 +29,54 @@ THE SOFTWARE.
|
||||
#include "util/hsa_rsrc_factory.h"
|
||||
|
||||
HsaRsrcFactory* TestHsa::hsa_rsrc_ = NULL;
|
||||
const 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) {
|
||||
HsaRsrcFactory* TestHsa::HsaInstantiate() {
|
||||
// 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;
|
||||
}
|
||||
return hsa_rsrc_;
|
||||
}
|
||||
|
||||
void TestHsa::HsaShutdown() {
|
||||
if (hsa_queue_ != NULL) {
|
||||
hsa_queue_destroy(hsa_queue_);
|
||||
hsa_queue_ = NULL;
|
||||
}
|
||||
if (hsa_rsrc_) hsa_rsrc_->Destroy();
|
||||
}
|
||||
|
||||
bool TestHsa::Initialize(int /*arg_cnt*/, char** /*arg_list*/) {
|
||||
std::clog << "TestHsa::Initialize :" << std::endl;
|
||||
|
||||
// Create an instance of Aql Queue
|
||||
if (hsa_queue_ == NULL) {
|
||||
uint32_t num_pkts = 128;
|
||||
if (hsa_rsrc_->CreateQueue(agent_info_, num_pkts, &hsa_queue_) == false) {
|
||||
hsa_queue_ = NULL;
|
||||
TEST_ASSERT(false);
|
||||
}
|
||||
}
|
||||
|
||||
// Instantiate a Timer object
|
||||
setup_timer_idx_ = hsa_timer_.CreateTimer();
|
||||
dispatch_timer_idx_ = hsa_timer_.CreateTimer();
|
||||
|
||||
if (HsaInstantiate(agent_id_) == NULL) {
|
||||
if (hsa_rsrc_ == NULL) {
|
||||
TEST_ASSERT(false);
|
||||
return false;
|
||||
}
|
||||
|
||||
// Create an instance of Gpu agent
|
||||
if (agent_info_ == NULL) {
|
||||
const uint32_t agent_id = 0;
|
||||
if (!hsa_rsrc_->GetGpuAgentInfo(agent_id, &agent_info_)) {
|
||||
agent_info_ = NULL;
|
||||
std::cerr << "> error: agent[" << agent_id << "] is not found" << std::endl;
|
||||
return NULL;
|
||||
}
|
||||
}
|
||||
std::clog << "> Using agent[" << agent_info_->dev_index << "] : " << agent_info_->name << std::endl;
|
||||
|
||||
// Create an instance of Aql Queue
|
||||
if (hsa_queue_ == NULL) {
|
||||
const uint32_t num_pkts = 128;
|
||||
if (hsa_rsrc_->CreateQueue(agent_info_, num_pkts, &hsa_queue_) == false) {
|
||||
hsa_queue_ = NULL;
|
||||
TEST_ASSERT(false);
|
||||
}
|
||||
my_queue_ = true;
|
||||
}
|
||||
|
||||
// Obtain handle of signal
|
||||
hsa_rsrc_->CreateSignal(1, &hsa_signal_);
|
||||
|
||||
@@ -283,7 +277,8 @@ void TestHsa::PrintTime() {
|
||||
bool TestHsa::Cleanup() {
|
||||
hsa_executable_destroy(hsa_exec_);
|
||||
hsa_signal_destroy(hsa_signal_);
|
||||
hsa_queue_destroy(hsa_queue_);
|
||||
if (my_queue_) hsa_queue_destroy(hsa_queue_);
|
||||
hsa_queue_ = NULL;
|
||||
agent_info_ = NULL;
|
||||
return true;
|
||||
}
|
||||
|
||||
+12
-11
@@ -32,24 +32,27 @@ THE SOFTWARE.
|
||||
class TestHsa : public TestAql {
|
||||
public:
|
||||
// Instantiate HSA resources
|
||||
static HsaRsrcFactory* HsaInstantiate(const uint32_t agent_ind = agent_id_);
|
||||
static HsaRsrcFactory* HsaInstantiate();
|
||||
static void HsaShutdown();
|
||||
static void SetQueue(hsa_queue_t* queue) { hsa_queue_ = queue; }
|
||||
static hsa_agent_t HsaAgent() { return agent_info_->dev_id; }
|
||||
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;
|
||||
agent_info_ = NULL;
|
||||
hsa_queue_ = NULL;
|
||||
my_queue_ = false;
|
||||
hsa_exec_ = {};
|
||||
}
|
||||
|
||||
// Get methods for Agent Info, HAS queue, HSA Resourcse Manager
|
||||
const AgentInfo* GetAgentInfo() { return agent_info_; }
|
||||
hsa_queue_t* GetQueue() { return hsa_queue_; }
|
||||
HsaRsrcFactory* GetRsrcFactory() { return hsa_rsrc_; }
|
||||
hsa_agent_t HsaAgent() { return agent_info_->dev_id; }
|
||||
const AgentInfo* GetAgentInfo() { return agent_info_; }
|
||||
void SetAgentInfo(const AgentInfo* agent_info) { agent_info_ = agent_info; }
|
||||
hsa_queue_t* GetQueue() { return hsa_queue_; }
|
||||
void SetQueue(hsa_queue_t* queue) { hsa_queue_ = queue; }
|
||||
|
||||
// Initialize application environment including setting
|
||||
// up of various configuration parameters based on
|
||||
@@ -106,14 +109,12 @@ class TestHsa : public TestAql {
|
||||
// Instance of Hsa Resources Factory
|
||||
static HsaRsrcFactory* hsa_rsrc_;
|
||||
|
||||
// GPU id
|
||||
static uint32_t agent_id_;
|
||||
|
||||
// Handle to an Hsa Gpu Agent
|
||||
static const AgentInfo* agent_info_;
|
||||
const AgentInfo* agent_info_;
|
||||
|
||||
// Handle to an Hsa Queue
|
||||
static hsa_queue_t* hsa_queue_;
|
||||
hsa_queue_t* hsa_queue_;
|
||||
bool my_queue_;
|
||||
|
||||
// Test kernel name
|
||||
std::string name_;
|
||||
|
||||
+4
-2
@@ -60,8 +60,10 @@ export ROCP_DITER=4
|
||||
export ROCP_INPUT=input1.xml
|
||||
eval ./test/ctrl
|
||||
|
||||
export ROCP_KITER=100
|
||||
export ROCP_DITER=100
|
||||
export ROCP_KITER=50
|
||||
export ROCP_DITER=50
|
||||
export ROCP_AGENTS=1
|
||||
export ROCP_THRS=1
|
||||
export ROCP_INPUT=input.xml
|
||||
eval ./test/ctrl
|
||||
|
||||
|
||||
@@ -1020,10 +1020,10 @@ extern "C" PUBLIC_API void OnUnloadTool() {
|
||||
// Dump stored profiling output data
|
||||
fflush(stdout);
|
||||
if (result_file_opened) {
|
||||
printf("\nROCPRofiler: %u contexts collected", context_collected); fflush(stdout);
|
||||
printf("\nROCPRofiler:"); fflush(stdout);
|
||||
dump_context_array(NULL);
|
||||
fclose(result_file_handle);
|
||||
printf(", output directory %s\n", result_prefix);
|
||||
printf(" %u contexts collected, output directory %s\n", context_collected, result_prefix);
|
||||
} else {
|
||||
if (context_collected != context_count) {
|
||||
results_output_break();
|
||||
|
||||
Ссылка в новой задаче
Block a user