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: 5d790f0cc0]

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);

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_;

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)
    *