standalone/sampling mode

Change-Id: I24c358bfebacb5678fb3d31f2f94ae867a1c9ac9

test/CMakeLists.txt

@@ -35,19 +35,35 @@ endif ()
 ## Util sources
 file( GLOB UTIL_SRC "${TEST_DIR}/util/*.cpp" )
 
+## Standalone test sources
+set ( STEXE_NAME "standalone_test" )
+set ( STST_SRC
+  ${TEST_DIR}/app/standalone_test.cpp
+  ${TEST_DIR}/ctrl/test_hsa.cpp
+)
+
 ## Test control sources
 set ( CTRL_SRC
   ${TEST_DIR}/app/test.cpp
   ${TEST_DIR}/ctrl/test_hsa.cpp
 )
 
-## Test kernel sources
+## Dummy kernel
+set ( DUMMY_NAME dummy_kernel )
+execute_process ( COMMAND sh -xc "cp ${TEST_DIR}/${DUMMY_NAME}/*.hsaco ${PROJECT_BINARY_DIR}" )
+
+## Test kernel
 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} )
+## Building standalone test executable
+add_executable ( ${STEXE_NAME} ${STST_SRC} ${UTIL_SRC} ${KERN_SRC} )
+target_include_directories ( ${STEXE_NAME} PRIVATE ${TEST_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH} )
+target_link_libraries( ${STEXE_NAME} ${ROCPROFILER_TARGET} ${HSA_RUNTIME_LIB} ${HSA_KMT_LIB} c stdc++ dl pthread rt )
+
+## Building ctrl test executable
+add_executable ( ${EXE_NAME} ${CTRL_SRC} ${UTIL_SRC} ${KERN_SRC} )
 target_include_directories ( ${EXE_NAME} PRIVATE ${TEST_DIR} ${ROOT_DIR} ${HSA_RUNTIME_INC_PATH} )
 target_link_libraries( ${EXE_NAME} ${HSA_RUNTIME_LIB} ${HSA_KMT_LIB} c stdc++ dl pthread rt )
 execute_process ( COMMAND sh -xc "cp ${TEST_DIR}/run.sh ${PROJECT_BINARY_DIR}" )

test/app/standalone_test.cpp

@@ -22,97 +22,25 @@ THE SOFTWARE.
 
 #include <hsa.h>
 #include <string.h>
+#include <unistd.h>
 #include <iostream>
 
 #include "ctrl/run_kernel.h"
 #include "ctrl/test_aql.h"
 #include "ctrl/test_hsa.h"
 #include "inc/rocprofiler.h"
+#include "dummy_kernel/dummy_kernel.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;
+void print_features(rocprofiler_feature_t* feature, uint32_t feature_count) {
+    for (rocprofiler_feature_t* p = feature; p < feature + feature_count; ++p) {
+      std::cout << (p - feature) << ": " << p->name;
       switch (p->data.kind) {
         case ROCPROFILER_DATA_KIND_INT64:
-          std::cout << std::dec << " result64 (" << p->data.result64 << ")" << std::endl;
+          std::cout << std::dec << " result64 (" << p->data.result_int64 << ")" << std::endl;
           break;
-        case ROCPROFILER_BYTES: {
+        case ROCPROFILER_DATA_KIND_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) {
@@ -130,15 +58,119 @@ int main(int argc, char** argv) {
           TEST_ASSERT(false);
       }
     }
+}
 
-    ++n;
-    if (n < invocation) {
-      status = rocprofiler_next(context);
-      TEST_STATUS(status == HSA_STATUS_SUCCESS);
-      continue;
-    }
-    break;
+void read_features(uint32_t n, rocprofiler_t* context, rocprofiler_feature_t* feature, const unsigned feature_count) {
+    std::cout << "read features" << std::endl;
+    hsa_status_t status = rocprofiler_read(context, n);
+    TEST_STATUS(status == HSA_STATUS_SUCCESS);
+    std::cout << "read issue" << std::endl;
+    status = rocprofiler_get_data(context, n);
+    TEST_STATUS(status == HSA_STATUS_SUCCESS);
+    status = rocprofiler_get_metrics(context);
+    TEST_STATUS(status == HSA_STATUS_SUCCESS);
+    print_features(feature, feature_count);
+}
+
+int main() {
+  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;
+
+  // Profiling feature objects
+  const unsigned feature_count = 9;
+  rocprofiler_feature_t feature[feature_count];
+  // PMC events
+  memset(feature, 0, sizeof(feature));
+  feature[0].kind = ROCPROFILER_FEATURE_KIND_METRIC;
+  feature[0].name = "GRBM_COUNT";
+  feature[1].kind = ROCPROFILER_FEATURE_KIND_METRIC;
+  feature[1].name = "GRBM_GUI_ACTIVE";
+  feature[2].kind = ROCPROFILER_FEATURE_KIND_METRIC;
+  feature[2].name = "GPUBusy";
+  feature[3].kind = ROCPROFILER_FEATURE_KIND_METRIC;
+  feature[3].name = "SQ_WAVES";
+  feature[4].kind = ROCPROFILER_FEATURE_KIND_METRIC;
+  feature[4].name = "SQ_INSTS_VALU";
+  feature[5].kind = ROCPROFILER_FEATURE_KIND_METRIC;
+  feature[5].name = "VALUInsts";
+  feature[6].kind = ROCPROFILER_FEATURE_KIND_METRIC;
+  feature[6].name = "TCC_HIT_sum";
+  feature[7].kind = ROCPROFILER_FEATURE_KIND_METRIC;
+  feature[7].name = "TCC_MISS_sum";
+  feature[8].kind = ROCPROFILER_FEATURE_KIND_METRIC;
+  feature[8].name = "WRITE_SIZE";
+//  feature[8].kind = ROCPROFILER_FEATURE_KIND_METRIC;
+//  feature[8].name = "TCC_EA_WRREQ_sum";
+//  feature[9].kind = ROCPROFILER_FEATURE_KIND_METRIC;
+//  feature[9].name = "TCC_EA_WRREQ_64B_sum";
+#if 0
+  // Tracing parameters
+  const unsigned parameter_count = 2;
+  rocprofiler_parameter_t parameters[parameter_count];
+  feature[2].name = "THREAD_TRACE";
+  feature[2].kind = ROCPROFILER_FEATURE_KIND_TRACE;
+  feature[2].parameters = parameters;
+  feature[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
+
+  // Instantiate HSA resources
+  HsaRsrcFactory::Create();
+
+  // Getting GPU device info
+  const AgentInfo* agent_info = NULL;
+  if (HsaRsrcFactory::Instance().GetGpuAgentInfo(0, &agent_info) == false) abort();
+
+  // Creating the queues pool
+  const unsigned queue_count = 16;
+  hsa_queue_t* queue[queue_count];
+  for (unsigned queue_ind = 0; queue_ind < queue_count; ++queue_ind) {
+    if (HsaRsrcFactory::Instance().CreateQueue(agent_info, 128, &queue[queue_ind]) == false) abort();
   }
+  hsa_queue_t* prof_queue = queue[0];
+
+  // Creating profiling context
+  properties = {};
+  properties.queue = prof_queue;
+  status = rocprofiler_open(agent_info->dev_id, feature, feature_count, &context,
+                            ROCPROFILER_MODE_STANDALONE, &properties);
+  TEST_STATUS(status == HSA_STATUS_SUCCESS);
+
+  // Test initialization
+  TestHsa::SetQueue(prof_queue);
+  TestHsa::HsaInstantiate(0);
+
+  // Dispatching profiled kernel n-times to collect all counter groups data
+  const unsigned group_n = 0;
+  status = rocprofiler_start(context, group_n);
+  TEST_STATUS(status == HSA_STATUS_SUCCESS);
+  std::cout << "start" << std::endl;
+
+  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>();
+    std::cout << "run kernel, queue " << queue_ind << std::endl;
+#else
+    sleep(3);
+#endif
+    read_features(group_n, context, feature, feature_count);
+  }
+
+  // Stop counters
+  status = rocprofiler_stop(context, group_n);
+  TEST_STATUS(status == HSA_STATUS_SUCCESS);
+  std::cout << "stop" << std::endl;
 
   // Finishing cleanup
   // Deleting profiling context will delete all allocated resources

test/ctrl/run_kernel.h

@@ -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, char* argv[], int count = 1) {
+template <class Kernel, class Test> bool RunKernel(int argc = 0, char* argv[] = NULL, int count = 1) {
   bool ret_val = false;
 
   // Create test kernel object

test/ctrl/test_hsa.cpp

@@ -71,7 +71,7 @@ void TestHsa::HsaShutdown() {
   if (hsa_rsrc_) hsa_rsrc_->Destroy();
 }
 
-bool TestHsa::Initialize(int arg_cnt, char** arg_list) {
+bool TestHsa::Initialize(int /*arg_cnt*/, char** /*arg_list*/) {
   std::clog << "TestHsa::Initialize :" << std::endl;
 
   // Instantiate a Timer object
@@ -119,6 +119,8 @@ bool TestHsa::Setup() {
   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;
+    if (des.size == 0) continue;
+
     switch (des.id) {
       case TestKernel::LOCAL_DES_ID:
         des.ptr = hsa_rsrc_->AllocateLocalMemory(agent_info_, des.size);
@@ -245,6 +247,8 @@ bool TestHsa::VerifyResults() {
   const uint32_t size = test_->GetOutputSize();
   bool suc = false;
 
+  if (size == 0) return true;
+
   // Copy local kernel output buffers from local memory into host memory
   if (test_->IsOutputLocal()) {
     output = hsa_rsrc_->AllocateSysMemory(agent_info_, size);

test/ctrl/test_hsa.h

@@ -35,6 +35,7 @@ class TestHsa : public TestAql {
   static HsaRsrcFactory* HsaInstantiate(const uint32_t agent_ind = agent_id_);
   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

test/dummy_kernel/dummy_kernel.cl

@@ -0,0 +1,28 @@
+/******************************************************************************
+Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in 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:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+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
+AUTHORS 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 IN
+THE SOFTWARE.
+*******************************************************************************/
+
+/**
+ dummy kernel
+ */
+__kernel void DummyKernel() {
+  uint tid   = get_global_id(0);
+}

test/dummy_kernel/dummy_kernel.h

@@ -0,0 +1,71 @@
+/******************************************************************************
+Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in 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:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+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
+AUTHORS 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 IN
+THE SOFTWARE.
+*******************************************************************************/
+
+#ifndef TEST_DUMMY_KERNEL_DUMMY_KERNEL_H_
+#define TEST_DUMMY_KERNEL_DUMMY_KERNEL_H_
+
+#include <map>
+#include <vector>
+
+#include "ctrl/test_kernel.h"
+
+// Class implements DummyKernel kernel parameters
+class DummyKernel : public TestKernel {
+ public:
+  // Kernel buffers IDs
+  enum { KERNARG_BUF_ID, LOCAL_BUF_ID };
+
+  // Constructor
+  DummyKernel() :
+    width_(64),
+    height_(64)
+  {
+    SetInDescr(KERNARG_BUF_ID, KERNARG_DES_ID, 0);
+    SetOutDescr(LOCAL_BUF_ID, LOCAL_DES_ID, 0);
+  }
+
+  // Initialize method
+  void Init() {}
+
+  // Return compute grid size
+  uint32_t GetGridSize() const { return width_ * height_; }
+
+  // Print output
+  void PrintOutput(const void* ptr) const {}
+
+  // Return name
+  std::string Name() const { return std::string("DummyKernel"); }
+
+ private:
+  // Reference CPU implementation
+  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) { return true; }
+
+  // Width of the Input array
+  const uint32_t width_;
+
+  // Height of the Input array
+  const uint32_t height_;
+};
+
+#endif  // TEST_DUMMY_KERNEL_DUMMY_KERNEL_H_

test/run.sh

@@ -22,21 +22,21 @@
 # THE SOFTWARE.
 ################################################################################
 
-test_bin_dflt=./test/ctrl
-
 # paths to ROC profiler and oher libraries
 export LD_LIBRARY_PATH=$PWD
-# enable error messages logging to '/tmp/rocprofiler_log.txt'
-export ROCPROFILER_LOG=1
-
 # ROC profiler library loaded by HSA runtime
 export HSA_TOOLS_LIB=librocprofiler64.so
-# tool library loaded by ROC profiler
-export ROCP_TOOL_LIB=libtool.so
+# enable error messages logging to '/tmp/rocprofiler_log.txt'
+export ROCPROFILER_LOG=1
 # ROC profiler metrics config file
 unset ROCP_PROXY_QUEUE
 # ROC profiler metrics config file
 export ROCP_METRICS=metrics.xml
+
+eval ./test/standalone_test
+
+# tool library loaded by ROC profiler
+export ROCP_TOOL_LIB=libtool.so
 # ROC profiler kernels timing
 export ROCP_TIMESTAMP_ON=1
 # output directory for the tool library, for metrics results file 'results.txt'
@@ -47,21 +47,15 @@ if [ ! -e $ROCP_TOOL_LIB ] ; then
   export ROCP_TOOL_LIB=test/libtool.so
 fi
 
-if [ -n "$1" ] ; then
-  tbin="$*"
-else
-  tbin=$test_bin_dflt
-fi
-
 export ROCP_KITER=100
 export ROCP_DITER=100
 export ROCP_INPUT=input.xml
-eval $tbin
+eval ./test/ctrl
 
-#export ROCP_KITER=1
-#export ROCP_DITER=4
-#export ROCP_INPUT=input1.xml
-#eval $tbin
+export ROCP_KITER=1
+export ROCP_DITER=4
+export ROCP_INPUT=input1.xml
+eval ./test/ctrl
 
 #valgrind --leak-check=full $tbin
 #valgrind --tool=massif $tbin