diff --git a/projects/rocr-runtime/rocrtst/common/base_rocr_utils.cc b/projects/rocr-runtime/rocrtst/common/base_rocr_utils.cc index bba6391419..2630a24bb6 100755 --- a/projects/rocr-runtime/rocrtst/common/base_rocr_utils.cc +++ b/projects/rocr-runtime/rocrtst/common/base_rocr_utils.cc @@ -57,7 +57,6 @@ #include "common/helper_funcs.h" #include "common/os.h" #include "hsa/hsa.h" -#include "hsa/hsa_ext_profiler.h" namespace rocrtst { @@ -331,7 +330,7 @@ hsa_status_t LoadKernelFromObjFile(BaseRocR* test) { } hsa_status_t CreateQueue(hsa_agent_t device, hsa_queue_t** queue, - uint32_t num_pkts, bool do_profile) { + uint32_t num_pkts) { hsa_status_t err; if (num_pkts == 0) { @@ -340,16 +339,9 @@ hsa_status_t CreateQueue(hsa_agent_t device, hsa_queue_t** queue, RET_IF_HSA_UTILS_ERR(err); } - if (do_profile) { - err = hsa_ext_tools_queue_create_profiled(device, - num_pkts, HSA_QUEUE_TYPE_SINGLE, NULL, - NULL, UINT32_MAX, UINT32_MAX, queue); - RET_IF_HSA_UTILS_ERR(err); - } else { - err = hsa_queue_create(device, num_pkts, HSA_QUEUE_TYPE_MULTI, NULL, - NULL, UINT32_MAX, UINT32_MAX, queue); - RET_IF_HSA_UTILS_ERR(err); - } + err = hsa_queue_create(device, num_pkts, HSA_QUEUE_TYPE_MULTI, NULL, + NULL, UINT32_MAX, UINT32_MAX, queue); + RET_IF_HSA_UTILS_ERR(err); return HSA_STATUS_SUCCESS; } diff --git a/projects/rocr-runtime/rocrtst/common/base_rocr_utils.h b/projects/rocr-runtime/rocrtst/common/base_rocr_utils.h index d083608314..40c0b50026 100755 --- a/projects/rocr-runtime/rocrtst/common/base_rocr_utils.h +++ b/projects/rocr-runtime/rocrtst/common/base_rocr_utils.h @@ -78,7 +78,7 @@ hsa_status_t SetDefaultAgents(BaseRocR* test); /// be created /// \returns HSA_STATUS_SUCCESS if no errors encountered hsa_status_t CreateQueue(hsa_agent_t device, hsa_queue_t** queue, - uint32_t num_pkts = 0, bool do_profile = false); + uint32_t num_pkts = 0); /// This function sets some reasonable default values for an AQL packet. /// Override any field as necessary after calling this function. diff --git a/projects/rocr-runtime/rocrtst/common/hsa_perf_cntrs.cc b/projects/rocr-runtime/rocrtst/common/hsa_perf_cntrs.cc deleted file mode 100755 index fc2a79c66b..0000000000 --- a/projects/rocr-runtime/rocrtst/common/hsa_perf_cntrs.cc +++ /dev/null @@ -1,229 +0,0 @@ -/* - * ============================================================================= - * ROC Runtime Conformance Release License - * ============================================================================= - * The University of Illinois/NCSA - * Open Source License (NCSA) - * - * Copyright (c) 2017, Advanced Micro Devices, Inc. - * All rights reserved. - * - * Developed by: - * - * AMD Research and AMD ROC Software Development - * - * Advanced Micro Devices, Inc. - * - * www.amd.com - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal with the Software without restriction, including without limitation - * the rights to use, copy, modify, merge, publish, distribute, sublicense, - * and/or sell copies of the Software, and to permit persons to whom the - * Software is furnished to do so, subject to the following conditions: - * - * - Redistributions of source code must retain the above copyright notice, - * this list of conditions and the following disclaimers. - * - Redistributions in binary form must reproduce the above copyright - * notice, this list of conditions and the following disclaimers in - * the documentation and/or other materials provided with the distribution. - * - Neither the names of , - * nor the names of its contributors may be used to endorse or promote - * products derived from this Software without specific prior written - * permission. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL - * THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR - * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, - * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER - * DEALINGS WITH THE SOFTWARE. - * - */ - -#include "common/hsa_perf_cntrs.h" -#include -#include -#include -#include -#include -#include -#include -#include -#include "hsa/hsa.h" -#include "hsa/hsa_ext_profiler.h" -#include "hsa/amd_hsa_tools_interfaces.h" - -namespace rocrtst { - - -static void -PreDispatchCallback(const hsa_dispatch_callback_t* dispParam, void* usrArg) { - assert((dispParam->pre_dispatch) && - "Pre Dispatch Callback Param is Malformed"); - - hsa_ext_tools_pmu_t* perfMgr = - reinterpret_cast(usrArg); - hsa_status_t status = hsa_ext_tools_pmu_begin(*perfMgr, dispParam->queue, - dispParam->aql_translation_handle, true); - assert((status == HSA_STATUS_SUCCESS) && - "Error in beginning Perf Cntr Session"); - (void)status; // Avoid warning -} - -static void -PostDispatchCallback(const hsa_dispatch_callback_t* dispParam, void* usrArg) { - assert((!dispParam->pre_dispatch) && - "Post Dispatch Callback Param is Malformed"); - - hsa_ext_tools_pmu_t* perfMgr = reinterpret_cast(usrArg); - hsa_status_t status = hsa_ext_tools_pmu_end(*perfMgr, dispParam->queue, - dispParam->aql_translation_handle); - assert((status == HSA_STATUS_SUCCESS) && - "Error in endning Perf Cntr Session"); - (void)status; // Avoid warning -} - -/// Constructor of the class -RocrPerfCntrApp::RocrPerfCntrApp() : perfMgr_(NULL) { -} - -/// Destructor of the class. Ideally it should delete the -/// PMU and its counters -RocrPerfCntrApp::~RocrPerfCntrApp() { -} - -/// Return the number of perf counters -uint32_t RocrPerfCntrApp::GetNumPerfCntrs() { - return uint32_t(cntrList_.size()); -} - -/// Return the handle of perf counter at specified index -CntrInfo* RocrPerfCntrApp::GetPerfCntr(uint32_t idx) { - return cntrList_[idx]; -} - -/// Print the various fields of Perf Cntrs being programmed. -bool RocrPerfCntrApp::PrintCntrs() { - CntrInfo* info; - int size = uint32_t(cntrList_.size()); - - for (int idx = 0; idx < size; idx++) { - info = cntrList_[idx]; - std::cout << std::endl; - std::cout << "Rocr Perf Cntr Id: " << info->cntrId << std::endl; - std::cout << "Rocr Perf Cntr Name: " << info->cntrName << std::endl; - std::cout << "Rocr Perf Cntr Blk Id: " << info->blkId << std::endl; - std::cout << "Rocr Perf Cntr Value: " << info->cntrResult << std::endl; - std::cout << "Rocr Perf Cntr Validation: " << info->cnfType << std::endl; - std::cout << std::endl; - } - - return true; -} - -// Initialize the list of perf counters -// block id of kHsaAiCounterBlockSQ = 14 == 0x0E -hsa_status_t RocrPerfCntrApp::Init(hsa_agent_t agent) { - // Initialize the list of Perf Cntrs - // Add SQ counter for number of waves - CntrInfo* info = NULL; - cntrList_.reserve(23); - - // Event for number of Waves - info = new CntrInfo(0x4, "SQ_SQ_PERF_SEL_WAVES", NULL, - 0x0E, NULL, 0x00, 0xFFFFFFFF, CntrValCnf_Exact); - cntrList_.push_back(info); - - // Event for number of Threads - info = new CntrInfo(0xE, "SQ_SQ_PERF_SEL_ITEMS", NULL, - 0x0E, NULL, 0x00, 0xFFFFFFFF, CntrValCnf_Exact); - cntrList_.push_back(info); - - - // Create an instance of Perf Mgr - hsa_status_t status; - status = hsa_ext_tools_create_pmu(agent, &perfMgr_); - assert((status == HSA_STATUS_SUCCESS) && "Error in creating Perf Cntr Mgr"); - - // Process each counter from the list as necessary - // each counter descriptor with its perf block handle - // and create an instance of counter in that block - uint32_t size = GetNumPerfCntrs(); - - for (uint32_t idx = 0; idx < size; idx++) { - info = GetPerfCntr(idx); - - // Obtain the handle of perf block - if (info->blkHndl == NULL) { - status = hsa_ext_tools_get_counter_block_by_id(perfMgr_, - info->blkId, &info->blkHndl); - assert((status == HSA_STATUS_SUCCESS) && - "Error in getting Perf Cntr Blk Hndl"); - } - - // Create an instance of counter in the perf block - status = hsa_ext_tools_create_counter(info->blkHndl, &info->cntrHndl); - assert((status == HSA_STATUS_SUCCESS) && - "Error in creating Perf Cntr in Perf Blk"); - - // Update the Event Index property of counter - uint32_t cntrProp = HSA_EXT_TOOLS_COUNTER_PARAMETER_EVENT_INDEX; - status = hsa_ext_tools_set_counter_parameter(info->cntrHndl, cntrProp, - sizeof(uint32_t), static_cast(&info->cntrId)); - assert((status == HSA_STATUS_SUCCESS) && - "Error in updating Perf Cntr Property Event Index"); - - // Enable the updated perf counter - status = hsa_ext_tools_set_counter_enabled(info->cntrHndl, true); - assert((status == HSA_STATUS_SUCCESS) && "Error in enabing Perf Cntr"); - } - - return status; -} - -// Register Pre and Post dispatch callbacks -void RocrPerfCntrApp::RegisterCallbacks(hsa_queue_t* queue) { - hsa_status_t status; - status = hsa_ext_tools_set_callback_functions(queue, PreDispatchCallback, - PostDispatchCallback); - assert((status == HSA_STATUS_SUCCESS) && - "Error in registering Pre & Post Dispatch Callbacks"); - status = hsa_ext_tools_set_callback_arguments(queue, &perfMgr_, &perfMgr_); - assert((status == HSA_STATUS_SUCCESS) && - "Error in registering Pre & Post Dispatch Callback Params"); - - (void)status; // Avoid warning - return; -} - -// Wait for perf counter collection to complete -hsa_status_t RocrPerfCntrApp::Wait() { - hsa_status_t status; - status = hsa_ext_tools_pmu_wait_for_completion(perfMgr_, 5000); - assert((status == HSA_STATUS_SUCCESS) && - "Error in Waiting for Perf Cntr Completion"); - return status; -} - -// Validate perf counter values -hsa_status_t RocrPerfCntrApp::Validate() { - // Retrieve the results of the different Perf Cntrs - // and validate them as configured - CntrInfo* info = NULL; - hsa_status_t status = HSA_STATUS_SUCCESS; - uint32_t size = GetNumPerfCntrs(); - - for (uint32_t idx = 0; idx < size; idx++) { - info = GetPerfCntr(idx); - status = hsa_ext_tools_get_counter_result(info->cntrHndl, - &info->cntrResult); - std::cout << "Value of Perf Cntr is: " << info->cntrResult << std::endl; - } - - return status; -} - -} // namespace rocrtst diff --git a/projects/rocr-runtime/rocrtst/common/hsa_perf_cntrs.h b/projects/rocr-runtime/rocrtst/common/hsa_perf_cntrs.h deleted file mode 100755 index f1fee9e35e..0000000000 --- a/projects/rocr-runtime/rocrtst/common/hsa_perf_cntrs.h +++ /dev/null @@ -1,159 +0,0 @@ -/* - * ============================================================================= - * ROC Runtime Conformance Release License - * ============================================================================= - * The University of Illinois/NCSA - * Open Source License (NCSA) - * - * Copyright (c) 2017, Advanced Micro Devices, Inc. - * All rights reserved. - * - * Developed by: - * - * AMD Research and AMD ROC Software Development - * - * Advanced Micro Devices, Inc. - * - * www.amd.com - * - * Permission is hereby granted, free of charge, to any person obtaining a copy - * of this software and associated documentation files (the "Software"), to - * deal with the Software without restriction, including without limitation - * the rights to use, copy, modify, merge, publish, distribute, sublicense, - * and/or sell copies of the Software, and to permit persons to whom the - * Software is furnished to do so, subject to the following conditions: - * - * - Redistributions of source code must retain the above copyright notice, - * this list of conditions and the following disclaimers. - * - Redistributions in binary form must reproduce the above copyright - * notice, this list of conditions and the following disclaimers in - * the documentation and/or other materials provided with the distribution. - * - Neither the names of , - * nor the names of its contributors may be used to endorse or promote - * products derived from this Software without specific prior written - * permission. - * - * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR - * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, - * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL - * THE CONTRIBUTORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR - * OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, - * ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER - * DEALINGS WITH THE SOFTWARE. - * - */ - -/// \file -/// Contains counter related functionality that can be used by samples and -/// tests. -#ifndef ROCRTST_COMMON_HSA_PERF_CNTRS_H_ -#define ROCRTST_COMMON_HSA_PERF_CNTRS_H_ - -#include "hsa/hsa.h" -#include "hsa/hsa_ext_profiler.h" -#include -#include -#include -#include -#include -#include -#include - -namespace rocrtst { - - -typedef enum CntrValCnfType { - ///< no counter value validation should be performed - CntrValCnf_None, - - ///< counter value should be an exact match to expectedResult - CntrValCnf_Exact, - - ///< counter value should be greater than expectedResult - CntrValCnf_GreaterThan, - - ///< counter value should be less than expectedResult - CntrValCnf_LessThan -} CntrValCnfType; - -/// Struct used to encapsulate Counter Info -typedef struct CntrInfo { - ///< Id of counter in hardware block - uint32_t cntrId; - - ///< Name of counter - char cntrName[72]; - - ///< Handle of perf counter - hsa_ext_tools_counter_t cntrHndl; - - ///< Id of hardware block containing the counter - uint32_t blkId; - - ///< Handle of counter block - hsa_ext_tools_counter_block_t blkHndl; - - ///< Expected value of perf counte - uint64_t expectedResult; - - ///< Value of perf counter expected - uint64_t cntrResult; - - ///< Type of validation upon completion of dispatch - CntrValCnfType cnfType; - - CntrInfo(uint32_t cntrId, const char* cntrName, void* cntrHndl, - uint32_t blkId, void* blkHndl, - uint64_t expResult, uint64_t result, CntrValCnfType cnfType) { - this->cntrId = cntrId; - this->cntrHndl = cntrHndl; - this->blkId = blkId; - this->blkHndl = blkHndl; - this->expectedResult = expResult; - this->cntrResult = result; - this->cnfType = cnfType; - memcpy(this->cntrName, cntrName, strlen(cntrName)); - } -} CntrInfo; - -class RocrPerfCntrApp { - public: - // Constructor of the class. Will initialize the list of perf counters - // that will be used to program the device - RocrPerfCntrApp(); - - // Destructor of the class - ~RocrPerfCntrApp(); - - // Return the number of perf counters - uint32_t GetNumPerfCntrs(); - - // Return the handle of perf counter at specified index - CntrInfo* GetPerfCntr(uint32_t idx); - - // Print the list of perf counters - bool PrintCntrs(); - - // Initialize the list of perf counters - hsa_status_t Init(hsa_agent_t agent); - - // Register Pre and Post dispatch callbacks - void RegisterCallbacks(hsa_queue_t* queue); - - // Wait for perf counter collection to complete - hsa_status_t Wait(); - - // Validate perf counter values - hsa_status_t Validate(); - - private: - // Number of queues to create - std::vector cntrList_; - - // Handle of Perf Cntr Manager - hsa_ext_tools_pmu_t perfMgr_; -}; - -} // namespace rocrtst - -#endif // ROCRTST_COMMON_HSA_PERF_CNTRS_H_ diff --git a/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt b/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt index 6618fade3f..1fc58413d6 100755 --- a/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt +++ b/projects/rocr-runtime/rocrtst/suites/test_common/CMakeLists.txt @@ -203,8 +203,8 @@ set(ROCRTST_ROOT ${CMAKE_CURRENT_SOURCE_DIR}/../..) # set(GOOGLE_TEST_FRWK_NAME "google-test-frwk${ONLY64STR}") add_subdirectory(${ROCRTST_ROOT}/gtest "${PROJECT_BINARY_DIR}/gtest") -set (ROCRTST_LIBS ${ROCRTST_LIBS} ${GOOGLE_TEST_FRWK_NAME} - hsa-runtime-tools${ONLY64STR}) +set (ROCRTST_LIBS ${ROCRTST_LIBS} ${GOOGLE_TEST_FRWK_NAME}) + MESSAGE("ROCRTST_LIBS SET STEP:Done") # #