From 0c8ce3dd18a308477f8960d3fd877c64cd4caff5 Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Wed, 15 Nov 2017 18:24:51 -0500
Subject: [PATCH] P4 to Git Change 1482981 by skeely@skeely_HSA_linux on
 2017/11/15 18:17:38

	SWDEV-136633 - Remove correlation handle.

	Not used by debugger or profiler, dead code.

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#47 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.hpp#15 edit


[ROCm/clr commit: 88783fe59b41ea80719e8923e874d3daf5ef377d]
---
 .../rocclr/runtime/device/rocm/rocvirtual.cpp | 40 -------------------
 .../rocclr/runtime/device/rocm/rocvirtual.hpp |  1 -
 2 files changed, 41 deletions(-)

diff --git a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp
index 5b902c9742..377c58a0da 100644
--- a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp
+++ b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.cpp
@@ -80,23 +80,6 @@ static const hsa_barrier_and_packet_t kBarrierReleasePacket = {
 
 double Timestamp::ticksToTime_ = 0;
 
-/**
-* Set the ocl correlation handle (essentially the cl_event handle)
-* to correlate the cl kernel launch and HSA kernel dispatch
-*/
-typedef hsa_status_t (*hsa_ext_tools_set_correlation_handle)(const hsa_agent_t agent,
-                                                             void* correlation_handle);
-static void SetOclCorrelationHandle(void* tools_lib, const hsa_agent_t agent, void* handle) {
-  hsa_ext_tools_set_correlation_handle func =
-      (hsa_ext_tools_set_correlation_handle)amd::Os::getSymbol(
-          tools_lib, "hsa_ext_tools_set_correlation_handler");
-  if (func) {
-    func(agent, handle);
-  }
-
-  return;
-}
-
 bool VirtualGPU::MemoryDependency::create(size_t numMemObj) {
   if (numMemObj > 0) {
     // Allocate the array of memory objects for dependency tracking
@@ -467,7 +450,6 @@ VirtualGPU::VirtualGPU(Device& device)
   // Initialize the last signal and dispatch flags
   timestamp_ = nullptr;
   hasPendingDispatch_ = false;
-  tools_lib_ = nullptr;
 
   kernarg_pool_base_ = nullptr;
   kernarg_pool_size_ = 0;
@@ -489,23 +471,10 @@ VirtualGPU::~VirtualGPU() {
     printfdbg_ = nullptr;
   }
 
-  tools_lib_ = nullptr;
   --roc_device_.numOfVgpus_;  // Virtual gpu unique index decrementing
 }
 
 bool VirtualGPU::create(bool profilingEna) {
-  // Set the event handle to the tools lib if the env var
-  // Load the library using its advertised "soname"
-  std::string lib_name = amd::Os::getEnvironment("HSA_TOOLS_LIB");
-  if (lib_name != "") {
-#if defined(_WIN32) || defined(__CYGWIN__)
-    const char* tools_lib_name = "hsa-runtime-tools" LP64_SWITCH("", "64") ".dll";
-#else
-    const char* tools_lib_name = "libhsa-runtime-tools" LP64_SWITCH("", "64") ".so.1";
-#endif
-    tools_lib_ = amd::Os::loadLibrary(tools_lib_name);
-  }
-
   // Checking Virtual gpu unique index for ROCm backend
   if (index() > device().settings().commandQueues_) {
     return false;
@@ -589,11 +558,6 @@ bool VirtualGPU::terminate() {
     hsa_signal_destroy(barrier_signal_);
   }
 
-  if (tools_lib_) {
-    amd::Os::unloadLibrary(tools_lib_);
-    tools_lib_ = nullptr;
-  }
-
   destroyPool();
 
   return true;
@@ -1492,10 +1456,6 @@ static void fillSampleDescriptor(hsa_ext_sampler_descriptor_t& samplerDescriptor
 
 bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const amd::Kernel& kernel,
                                       const_address parameters, void* eventHandle) {
-  if (tools_lib_) {
-    SetOclCorrelationHandle(tools_lib_, this->gpu_device_, eventHandle);
-  }
-
   device::Kernel* devKernel = const_cast<device::Kernel*>(kernel.getDeviceKernel(dev()));
   Kernel& gpuKernel = static_cast<Kernel&>(*devKernel);
 
diff --git a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp
index d8c64d3f17..148699c172 100644
--- a/projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp
+++ b/projects/clr/rocclr/runtime/device/rocm/rocvirtual.hpp
@@ -277,7 +277,6 @@ class VirtualGPU : public device::VirtualDevice {
   hsa_signal_t barrier_signal_;
   uint32_t dispatch_id_;  //!< This variable must be updated atomically.
   Device& roc_device_;    //!< roc device object
-  void* tools_lib_;
   PrintfDbg* printfdbg_;
   MemoryDependency memoryDependency_;  //!< Memory dependency class
   uint16_t aqlHeader_;                 //!< AQL header for dispatch