SWDEV-252747 : testing using v3 object

Change-Id: I427df765d1be55bd2851ce441238b3eaa46cca4f


[ROCm/rocprofiler commit: 169e36f379]

projects/rocprofiler/CMakeLists.txt

@@ -112,11 +112,18 @@ add_subdirectory ( ${TEST_DIR} ${PROJECT_BINARY_DIR}/test )
 
 ## Installation and packaging
 set ( DEST_NAME ${ROCPROFILER_NAME} )
+if ( DEFINED CMAKE_INSTALL_PREFIX )
+  get_filename_component ( prefix_name ${CMAKE_INSTALL_PREFIX} NAME )
+  get_filename_component ( prefix_dir ${CMAKE_INSTALL_PREFIX} DIRECTORY )
+  if ( prefix_name STREQUAL ${DEST_NAME} )
+    set ( CMAKE_INSTALL_PREFIX ${prefix_dir} )
+  endif ()
+endif ()
 if ( DEFINED CPACK_PACKAGING_INSTALL_PREFIX )
-  get_filename_component ( pkg_name ${CPACK_PACKAGING_INSTALL_PREFIX} NAME )
-  get_filename_component ( pkg_dir ${CPACK_PACKAGING_INSTALL_PREFIX} DIRECTORY )
-  if ( pkg_name STREQUAL ${DEST_NAME} )
-    set ( CPACK_PACKAGING_INSTALL_PREFIX ${pkg_dir} )
+  get_filename_component ( prefix_name ${CPACK_PACKAGING_INSTALL_PREFIX} NAME )
+  get_filename_component ( prefix_dir ${CPACK_PACKAGING_INSTALL_PREFIX} DIRECTORY )
+  if ( prefix_name STREQUAL ${DEST_NAME} )
+    set ( CPACK_PACKAGING_INSTALL_PREFIX ${prefix_dir} )
   endif ()
 else ()
   set ( CPACK_PACKAGING_INSTALL_PREFIX ${CMAKE_INSTALL_PREFIX} )

projects/rocprofiler/bin/build_kernel.sh

@@ -73,7 +73,7 @@ fi
 for GFXIP in $TGT_LIST ; do
   OBJ_PREF=$GFXIP
   OBJ_FILE="${OBJ_PREF}_${OBJ_NAME}.$SO_EXT"
-  $BIN_DIR/clang -cl-std=CL$OCL_VER -include $INC_DIR/opencl-c.h $BITCODE_OPTS -target amdgcn-amd-amdhsa -mcpu=$GFXIP -mno-code-object-v3 $TEST_NAME.cl -o $DST_DIR/$OBJ_FILE
+  $BIN_DIR/clang -cl-std=CL$OCL_VER -include $INC_DIR/opencl-c.h $BITCODE_OPTS -target amdgcn-amd-amdhsa -mcpu=$GFXIP $TEST_NAME.cl -o $DST_DIR/$OBJ_FILE
   echo "'$OBJ_FILE' generated"
 done
 

projects/rocprofiler/src/core/rocprofiler.cpp

@@ -166,9 +166,10 @@ enum {
   CODE_OBJ_TRACKING_MODE = 0x2,
   MEMCOPY_INTERCEPT_MODE = 0x4,
   HSA_INTERCEPT_MODE = 0x8,
+  INTERCEPT_MODE_DFLT = CODE_OBJ_TRACKING_MODE
 };
 uint32_t LoadTool() {
-  uint32_t intercept_mode = 0;
+  uint32_t intercept_mode = INTERCEPT_MODE_DFLT;
   const char* tool_lib = getenv("ROCP_TOOL_LIB");
   std::ostringstream oss;
   if (tool_lib) oss << "load tool library(" << tool_lib << ")";

projects/rocprofiler/test/ctrl/test_hsa.cpp

@@ -96,7 +96,7 @@ bool TestHsa::Setup() {
 
   // Load and Finalize Kernel Code Descriptor
   const char* brig_path = brig_path_obj_.c_str();
-  bool suc = hsa_rsrc_->LoadAndFinalize(agent_info_, brig_path, name_.c_str(), &hsa_exec_,
+  bool suc = hsa_rsrc_->LoadAndFinalize(agent_info_, brig_path, symb_.c_str(), &hsa_exec_,
                                         &kernel_code_desc_);
   if (suc == false) {
     std::cerr << "Error in loading and finalizing Kernel" << std::endl;
@@ -116,8 +116,10 @@ bool TestHsa::Setup() {
         // Check the kernel args size
         const size_t kernarg_size = des.size;
         size_t size_info = 0;
-        hsa_executable_symbol_get_info(
+        const hsa_status_t status = hsa_executable_symbol_get_info(
             kernel_code_desc_, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_KERNARG_SEGMENT_SIZE, &size_info);
+        TEST_ASSERT(status == HSA_STATUS_SUCCESS);
+        size_info = kernarg_size;
         const bool kernarg_missmatch = (kernarg_size > size_info);
         if (kernarg_missmatch) {
           std::cout << "kernarg_size = " << kernarg_size << ", size_info = " << size_info

projects/rocprofiler/test/ctrl/test_hsa.h

@@ -36,7 +36,7 @@ class TestHsa : public TestAql {
   static void HsaShutdown();
 
   // Constructor
-  explicit TestHsa(TestKernel* test) : test_(test), name_(test->Name()) {
+  explicit TestHsa(TestKernel* test) : test_(test), name_(test->Name()), symb_(test->SymbName()) {
     total_time_taken_ = 0;
     setup_time_taken_ = 0;
     dispatch_time_taken_ = 0;
@@ -119,6 +119,9 @@ class TestHsa : public TestAql {
   // Test kernel name
   std::string name_;
 
+  // Test kernel symboll name
+  std::string symb_;
+
   // Kernel executable
   hsa_executable_t hsa_exec_;
 };

projects/rocprofiler/test/ctrl/test_kernel.h

@@ -80,6 +80,9 @@ class TestKernel {
   // Return name
   virtual std::string Name() const = 0;
 
+  // Return name
+  virtual std::string SymbName() { return Name() + ".kd"; }
+
  protected:
   // Set buffer descriptor
   bool SetInDescr(const uint32_t& buf_id, const des_id_t& des_id, const uint32_t& size) {

projects/rocprofiler/test/run.sh

@@ -1,4 +1,4 @@
-#!/bin/sh
+#!/bin/bash -x
 
 ################################################################################
 # Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved.
@@ -56,8 +56,25 @@ eval_test() {
   test_number=$((test_number + 1))
 }
 
+# profiler library lookup
+pwd
+echo "ENV CHECK"
+env
+echo "OPTROCM CHECK"
+find -L /opt -name "librocprofiler*"
+echo "PKGLIB CHECK"
+find -L ../.. -name "librocprofiler*"
+echo "COMPKG CHECK"
+find -L /home/jenkins/compute-package -name "librocprofiler*"
+
+ls -la /home/jenkins/compute-package
+ls -la /home/jenkins/compute-package/lib
+ls -la /home/jenkins/compute-package/lib/*
+
 # paths to ROC profiler and oher libraries
-export LD_LIBRARY_PATH=$PWD
+#ROCP_LIB_PATH=$(find -L /opt/rocm* -name librocprofiler64.so.1 | head -n1)
+#ROCP_LIB_DIR=$(dirname $ROCP_LIB_PATH)
+export LD_LIBRARY_PATH=$PWD:$PWD/../../lib:/home/jenkins/compute-package/lib
 
 # enable tools load failure reporting
 export HSA_TOOLS_REPORT_LOAD_FAILURE=1
@@ -67,6 +84,8 @@ export ROCPROFILER_LOG=1
 export HSA_VEN_AMD_AQLPROFILE_LOG=1
 # test trace
 export ROC_TEST_TRACE=1
+# enable V3 code object support
+export ROCP_OBJ_TRACKING=1
 
 # Disabple profiler own proxy queue
 unset ROCP_PROXY_QUEUE
@@ -178,7 +197,6 @@ export ROCP_INPUT=set_input.xml
 eval_test "libtool test, counter sets" ./test/ctrl
 
 ## OpenCL test
-#export ROCP_OBJ_TRACKING=1
 #export ROCP_INPUT=input1.xml
 #eval_test "libtool test, OpenCL sample" ./test/ocl/SimpleConvolution
 
@@ -193,7 +211,6 @@ export ROCP_DITER=10
 eval_test "libtool test, counter sets" ./test/ctrl
 
 ## OpenCL test
-#export ROCP_OBJ_TRACKING=1
 #eval_test "libtool test, OpenCL sample" ./test/ocl/SimpleConvolution
 
 #valgrind --leak-check=full $tbin