From 29aabcf956a49b1af3cd65b613dc38e3cfd00b0c Mon Sep 17 00:00:00 2001
From: foreman <dl.swbuild@amd.com>
Date: Mon, 18 Mar 2019 18:44:55 -0400
Subject: [PATCH] P4 to Git Change 1757948 by kjayapra@1_HIPWS_SL_IPC on
 2019/03/18 18:29:24

	       SWDEV-144570 - Implementation of
	               hipMemcpyToSymbol, hipMemcpyFromSymbol,
	               hipMemcpyToSymbolAsync, hipMemcpyFromSymbolAsync,
	               hipGetSymbolAddress, hipModuleGetGlobal

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.def.in#12 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#13 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#23 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#45 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_module.cpp#21 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#22 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/devprogram.hpp#20 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.cpp#101 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocprogram.hpp#45 edit


[ROCm/clr commit: 5d790f0cc025c28fb33b575638b1a22cc9875c6d]
---
 .../clr/rocclr/runtime/device/devprogram.hpp  |  6 +-
 .../rocclr/runtime/device/rocm/rocprogram.cpp | 60 +++++++++++++++++++
 .../rocclr/runtime/device/rocm/rocprogram.hpp |  2 +
 3 files changed, 67 insertions(+), 1 deletion(-)

diff --git a/projects/clr/rocclr/runtime/device/devprogram.hpp b/projects/clr/rocclr/runtime/device/devprogram.hpp
index 7730c92774..92fdd7d239 100644
--- a/projects/clr/rocclr/runtime/device/devprogram.hpp
+++ b/projects/clr/rocclr/runtime/device/devprogram.hpp
@@ -106,7 +106,7 @@ class Program : public amd::HeapObject {
   aclBinary* binaryElf_;            //!< Binary for the new compiler library
 
   std::string lastBuildOptionsArg_;
-  std::string buildLog_;            //!< build log.
+  mutable std::string buildLog_;    //!< build log.
   cl_int buildStatus_;              //!< build status.
   cl_int buildError_;               //!< build error
 
@@ -224,6 +224,10 @@ class Program : public amd::HeapObject {
   //! Check if SRAM ECC is enable
   const bool sramEccEnable() const { return (sramEccEnabled_ == 1); }
 
+  virtual bool findGlobalSymbols(void** dptr, size_t* bytes, const char* globalName) const {
+    ShouldNotReachHere();
+  }
+
  protected:
   //! pre-compile setup
   bool initBuild(amd::option::Options* options);
diff --git a/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp b/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp
index e5ec876a17..350182c056 100644
--- a/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp
+++ b/projects/clr/rocclr/runtime/device/rocm/rocprogram.cpp
@@ -113,6 +113,66 @@ bool Program::initClBinary(char* binaryIn, size_t size) {
   return clBinary()->setBinary(bin, sz, (decryptedBin != nullptr));
 }
 
+bool Program::findGlobalSymbols(void** device_pptr, size_t* bytes, const char* global_name) const {
+  hsa_status_t status = HSA_STATUS_SUCCESS;
+  hsa_agent_t hsa_device;
+  hsa_symbol_kind_t sym_type;
+  hsa_executable_symbol_t global_symbol;
+
+  hsa_device= dev().getBackendDevice();
+
+  /* Find HSA Symbol by name */
+  status = hsa_executable_get_symbol_by_name(hsaExecutable_, global_name, &hsa_device,
+                                             &global_symbol);
+  if (status != HSA_STATUS_SUCCESS) {
+    buildLog_ += "Error: Failed to find the Symbol by Name: ";
+    buildLog_ += hsa_strerror(status);
+    buildLog_ += "\n";
+    return false;
+  }
+
+  /* Find HSA Symbol Type */
+  status = hsa_executable_symbol_get_info(global_symbol, HSA_EXECUTABLE_SYMBOL_INFO_TYPE,
+                                          &sym_type);
+  if (status != HSA_STATUS_SUCCESS) {
+    buildLog_ += "Error: Failed to find the Symbol Type : ";
+    buildLog_ += hsa_strerror(status);
+    buildLog_ += "\n";
+    return false;
+  }
+
+  /* Make sure symbol type is VARIABLE */
+  if (sym_type != HSA_SYMBOL_KIND_VARIABLE) {
+    buildLog_ += "Error: Symbol is not of type VARIABLE : ";
+    buildLog_ += hsa_strerror(status);
+    buildLog_ += "\n";
+    return false;
+  }
+
+  /* Retrieve the size of the variable */
+  status = hsa_executable_symbol_get_info(global_symbol, HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_SIZE,
+                                          bytes);
+
+  if (status != HSA_STATUS_SUCCESS) {
+    buildLog_ += "Error: Failed to retrieve the Symbol Size : ";
+    buildLog_ += hsa_strerror(status);
+    buildLog_ += "\n";
+    return false;
+  }
+
+  /* Find HSA Symbol Address */
+  status = hsa_executable_symbol_get_info(global_symbol,
+                                          HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, device_pptr);
+  if (status != HSA_STATUS_SUCCESS) {
+    buildLog_ += "Error: Failed to find the Symbol Address : ";
+    buildLog_ += hsa_strerror(status);
+    buildLog_ += "\n";
+    return false;
+  }
+
+  return true;
+}
+
 HSAILProgram::HSAILProgram(roc::NullDevice& device) : roc::Program(device) {
   xnackEnabled_ = dev().settings().enableXNACK_;
   sramEccEnabled_ = dev().info().sramEccEnabled_;
diff --git a/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp b/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp
index 98e86f222a..60ba2a7f38 100644
--- a/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp
+++ b/projects/clr/rocclr/runtime/device/rocm/rocprogram.hpp
@@ -39,6 +39,8 @@ class Program : public device::Program {
 
   hsa_executable_t hsaExecutable() const { return hsaExecutable_; }
 
+  virtual bool findGlobalSymbols(void** dptr, size_t* bytes, const char* globalName) const;
+
  protected:
   /*! \brief Compiles LLVM binary to HSAIL code (compiler backend: link+opt+codegen)
    *