diff --git a/bin/rpl_run.sh b/bin/rpl_run.sh index 123fcaf411..7cf183b57d 100755 --- a/bin/rpl_run.sh +++ b/bin/rpl_run.sh @@ -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 "" > $RES_DIR/input.xml + fi input_list=`/bin/ls $RES_DIR/input*.xml` export ROCPROFILER_SESS=$RES_DIR else diff --git a/inc/rocprofiler.h b/inc/rocprofiler.h index cc21b76f0c..4a25919e6a 100644 --- a/inc/rocprofiler.h +++ b/inc/rocprofiler.h @@ -46,7 +46,7 @@ THE SOFTWARE. #include #include -#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; diff --git a/src/core/intercept_queue.cpp b/src/core/intercept_queue.cpp index 2b901767dc..91028f739d 100644 --- a/src/core/intercept_queue.cpp +++ b/src/core/intercept_queue.cpp @@ -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 diff --git a/src/core/intercept_queue.h b/src/core/intercept_queue.h index 0420ccb81f..1f31b0d9c9 100644 --- a/src/core/intercept_queue.h +++ b/src/core/intercept_queue.h @@ -63,8 +63,8 @@ class InterceptQueue { std::lock_guard 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 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_; diff --git a/src/core/rocprofiler.cpp b/src/core/rocprofiler.cpp index 4b7fc9a452..6042e59ef7 100644 --- a/src/core/rocprofiler.cpp +++ b/src/core/rocprofiler.cpp @@ -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 } diff --git a/test/app/standalone_test.cpp b/test/app/standalone_test.cpp index 28a34ea0d0..b173c4d38c 100644 --- a/test/app/standalone_test.cpp +++ b/test/app/standalone_test.cpp @@ -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(); - ret_val = RunKernel(); + hsa_queue_t* prof_queue = queue[queue_ind]; + //ret_val = RunKernel(0, NULL, NULL, prof_queue); + ret_val = RunKernel(0, NULL, NULL, prof_queue); std::cout << "run kernel, queue " << queue_ind << std::endl; #else sleep(3); diff --git a/test/app/test.cpp b/test/app/test.cpp index 9e6948339e..796ba1eb1d 100644 --- a/test/app/test.cpp +++ b/test/app/test.cpp @@ -21,20 +21,66 @@ THE SOFTWARE. *******************************************************************************/ #include +#include #include #include +#include #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(0, NULL, agent_info[n], queue[n], diter); + RunKernel(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(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; } diff --git a/test/ctrl/run_kernel.h b/test/ctrl/run_kernel.h index 23d7ea0385..846e0b6825 100644 --- a/test/ctrl/run_kernel.h +++ b/test/ctrl/run_kernel.h @@ -26,7 +26,7 @@ THE SOFTWARE. #include "ctrl/test_hsa.h" #include "util/test_assert.h" -template bool RunKernel(int argc = 0, char* argv[] = NULL, int count = 1) { +template 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 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; diff --git a/test/ctrl/test_hsa.cpp b/test/ctrl/test_hsa.cpp index e321b72cac..d006d19c4e 100644 --- a/test/ctrl/test_hsa.cpp +++ b/test/ctrl/test_hsa.cpp @@ -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; } diff --git a/test/ctrl/test_hsa.h b/test/ctrl/test_hsa.h index a7d5398808..b5df8b697e 100644 --- a/test/ctrl/test_hsa.h +++ b/test/ctrl/test_hsa.h @@ -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_; diff --git a/test/run.sh b/test/run.sh index abb1ba0725..528d5f5ab9 100755 --- a/test/run.sh +++ b/test/run.sh @@ -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 diff --git a/test/tool/tool.cpp b/test/tool/tool.cpp index 2317595f5d..ac9b7a28cd 100644 --- a/test/tool/tool.cpp +++ b/test/tool/tool.cpp @@ -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();