Merge 'amd-master-next' into 'amd-npi-next'

Change-Id: I88aec6ceb6b6e49caf48a0287ee2379b8ed9b296


[ROCm/hip commit: 85e4233a23]

projects/hip/include/hip/hcc_detail/hip_runtime.h

@@ -478,24 +478,36 @@ struct __hip_builtin_threadIdx_t {
   __HIP_DEVICE_BUILTIN(x,__hip_get_thread_idx_x());
   __HIP_DEVICE_BUILTIN(y,__hip_get_thread_idx_y());
   __HIP_DEVICE_BUILTIN(z,__hip_get_thread_idx_z());
+#ifdef __cplusplus
+  __device__ operator dim3() const { return dim3(x, y, z); }
+#endif
 };
 
 struct __hip_builtin_blockIdx_t {
   __HIP_DEVICE_BUILTIN(x,__hip_get_block_idx_x());
   __HIP_DEVICE_BUILTIN(y,__hip_get_block_idx_y());
   __HIP_DEVICE_BUILTIN(z,__hip_get_block_idx_z());
+#ifdef __cplusplus
+  __device__ operator dim3() const { return dim3(x, y, z); }
+#endif
 };
 
 struct __hip_builtin_blockDim_t {
   __HIP_DEVICE_BUILTIN(x,__hip_get_block_dim_x());
   __HIP_DEVICE_BUILTIN(y,__hip_get_block_dim_y());
   __HIP_DEVICE_BUILTIN(z,__hip_get_block_dim_z());
+#ifdef __cplusplus
+  __device__ operator dim3() const { return dim3(x, y, z); }
+#endif
 };
 
 struct __hip_builtin_gridDim_t {
   __HIP_DEVICE_BUILTIN(x,__hip_get_grid_dim_x());
   __HIP_DEVICE_BUILTIN(y,__hip_get_grid_dim_y());
   __HIP_DEVICE_BUILTIN(z,__hip_get_grid_dim_z());
+#ifdef __cplusplus
+  __device__ operator dim3() const { return dim3(x, y, z); }
+#endif
 };
 
 #undef __HIP_DEVICE_BUILTIN

projects/hip/lpl_ca/CMakeLists.txt

@@ -1,4 +1,10 @@
 #-------------------------------------LPL--------------------------------------#
+# Look for ROCclr which contains elfio
+find_package(ROCclr REQUIRED CONFIG
+  PATHS
+    /opt/rocm
+    /opt/rocm/rocclr)
+
 add_executable(lpl lpl.cpp)
 set_target_properties(
     lpl PROPERTIES
@@ -6,7 +12,12 @@ set_target_properties(
         CXX_STANDARD_REQUIRED ON
         CXX_EXTENSIONS OFF
         RUNTIME_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR})
-target_include_directories(lpl PUBLIC ${PROJECT_SOURCE_DIR}/src)
+target_include_directories(lpl
+  PUBLIC
+    ${PROJECT_SOURCE_DIR}/src
+  PRIVATE
+    $<TARGET_PROPERTY:amdrocclr_static,INTERFACE_INCLUDE_DIRECTORIES>)
+
 target_compile_options(lpl PUBLIC -Wall)
 target_link_libraries(lpl PUBLIC pthread)
 

projects/hip/lpl_ca/lpl.hpp

@@ -4,7 +4,7 @@
 
 #include "clara/clara.hpp"
 #include "pstreams/pstream.h"
-#include "hip/hcc_detail/elfio/elfio.hpp"
+#include <elfio/elfio.hpp>
 
 #include <unistd.h>
 
@@ -71,14 +71,14 @@ inline std::string make_hipcc_call(const std::vector<std::string>& sources,
 }
 
 inline void copy_kernel_section_to_fat_binary(const std::string& tmp, const std::string& output) {
-    ELFIO::elfio reader;
+    amd::ELFIO::elfio reader;
     if (!reader.load(tmp)) {
         throw std::runtime_error{"The result of the compilation is inaccessible."};
     }
 
     const auto it =
         std::find_if(reader.sections.begin(), reader.sections.end(),
-                     [](const ELFIO::section* x) { return x->get_name() == kernel_section(); });
+                     [](const amd::ELFIO::section* x) { return x->get_name() == kernel_section(); });
 
     std::ofstream out{output};
 

projects/hip/rocclr/CMakeLists.txt

@@ -96,8 +96,6 @@ find_package(amd_comgr REQUIRED CONFIG
 
 message(STATUS "Code Object Manager found at ${amd_comgr_DIR}.")
 
-add_definitions(-DBSD_LIBELF)
-
 add_library(hip64 OBJECT
  hip_context.cpp
  hip_code_object.cpp
@@ -129,10 +127,8 @@ target_include_directories(hip64
     ${PROJECT_SOURCE_DIR}/include
     ${PROJECT_BINARY_DIR}/include
   PRIVATE
-    ${CMAKE_SOURCE_DIR}/elfio
     ${PROJECT_SOURCE_DIR}
     ${PROJECT_SOURCE_DIR}/amdocl
-    ${PROJECT_SOURCE_DIR}/include/hip/hcc_detail/elfio
     ${ROCR_INCLUDES}
     $<TARGET_PROPERTY:amdrocclr_static,INTERFACE_INCLUDE_DIRECTORIES>
     $<TARGET_PROPERTY:amd_comgr,INTERFACE_INCLUDE_DIRECTORIES>)

projects/hip/rocclr/hip_code_object.cpp

@@ -6,27 +6,12 @@
 #include "hip/hip_runtime.h"
 #include "hip_internal.hpp"
 #include "platform/program.hpp"
+#include <elf/elf.hpp>
 
 namespace hip {
 
 uint64_t CodeObject::ElfSize(const void *emi) {
-  const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi;
-  const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
-
-  uint64_t max_offset = ehdr->e_shoff;
-  uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum;
-
-  for (uint16_t i=0; i < ehdr->e_shnum; ++i){
-    uint64_t cur_offset = static_cast<uint64_t>(shdr[i].sh_offset);
-    if (max_offset < cur_offset) {
-      max_offset = cur_offset;
-      total_size = max_offset;
-      if(SHT_NOBITS != shdr[i].sh_type) {
-        total_size += static_cast<uint64_t>(shdr[i].sh_size);
-      }
-    }
-  }
-  return total_size;
+  return amd::Elf::getElfSize(emi);
 }
 
 bool CodeObject::isCompatibleCodeObject(const std::string& codeobj_target_id,

projects/hip/rocclr/hip_event.cpp

@@ -90,12 +90,10 @@ hipError_t Event::elapsedTime(Event& eStop, float& ms) {
     return hipErrorNotReady;
   }
 
-  // For certain HIP API's that take start and stop event
-  // and no hipEventRecord needs to be called
-  if (event_ == eStop.event_ && !recorded_  && !eStop.recorded_) {
+  if (event_ != eStop.event_ && recorded_ && eStop.recorded_) {
     ms = static_cast<float>(static_cast<int64_t>(eStop.event_->profilingInfo().end_ -
-                          event_->profilingInfo().start_))/1000000.f;
-  } else if (event_ == eStop.event_) {
+                          event_->profilingInfo().end_))/1000000.f;
+  } else if (event_ == eStop.event_ && (recorded_ || eStop.recorded_)) {
     // Events are the same, which indicates the stream is empty and likely
     // eventRecord is called on another stream. For such cases insert and measure a
     // marker.
@@ -107,8 +105,10 @@ hipError_t Event::elapsedTime(Event& eStop, float& ms) {
     command->release();
 
   } else {
+  // For certain HIP API's that take both start and stop event
+  // or scenarios where HIP API takes one of the events and the other event is recorded with hipEventRecord
     ms = static_cast<float>(static_cast<int64_t>(eStop.event_->profilingInfo().end_ -
-                          event_->profilingInfo().end_))/1000000.f;
+                          event_->profilingInfo().start_))/1000000.f;
   }
   return hipSuccess;
 }

projects/hip/rocclr/hip_memory.cpp

@@ -150,12 +150,14 @@ hipError_t ihipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKin
   size_t dOffset = 0;
   amd::Memory *dstMemory = getMemoryObject(dst, dOffset);
   amd::Device* queueDevice = &queue.device();
-
-  if (((srcMemory == nullptr) && (dstMemory == nullptr)) ||
-      (kind == hipMemcpyHostToHost)) {
-    queue.finish();
-    memcpy(dst, src, sizeBytes);
-    return hipSuccess;
+  if ((srcMemory == nullptr) && (dstMemory == nullptr)) {
+    if ((kind == hipMemcpyHostToHost) || (kind == hipMemcpyDefault)) {
+      queue.finish();
+      memcpy(dst, src, sizeBytes);
+      return hipSuccess;
+    } else {
+      return hipErrorInvalidValue;
+    }
   } else if ((srcMemory == nullptr) && (dstMemory != nullptr)) {
     amd::HostQueue* pQueue = &queue;
     if (queueDevice != dstMemory->getContext().devices()[0]) {

projects/hip/rocclr/hip_module.cpp

@@ -19,7 +19,7 @@
  THE SOFTWARE. */
 
 #include <hip/hip_runtime.h>
-#include <libelf.h>
+#include <elf/elf.hpp>
 #include <fstream>
 
 #include "hip_internal.hpp"
@@ -45,23 +45,7 @@ const std::string& FunctionName(const hipFunction_t f) {
 
 static uint64_t ElfSize(const void *emi)
 {
-  const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi;
-  const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
-
-  uint64_t max_offset = ehdr->e_shoff;
-  uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum;
-
-  for (uint16_t i=0; i < ehdr->e_shnum; ++i){
-    uint64_t cur_offset = static_cast<uint64_t>(shdr[i].sh_offset);
-    if (max_offset < cur_offset) {
-      max_offset = cur_offset;
-      total_size = max_offset;
-      if(SHT_NOBITS != shdr[i].sh_type) {
-        total_size += static_cast<uint64_t>(shdr[i].sh_size);
-      }
-    }
-  }
-  return total_size;
+  return amd::Elf::getElfSize(emi);
 }
 
 hipError_t hipModuleUnload(hipModule_t hmod) {
@@ -303,7 +287,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
     }
   }
 
-  profileNDRange = (startEvent != nullptr && stopEvent != nullptr);
+  profileNDRange = (startEvent != nullptr || stopEvent != nullptr);
 
   // Flag set to 1 signifies that kernel can be launched in anyorder
   if (flags & hipExtAnyOrderLaunch) {

projects/hip/rocclr/hip_platform.cpp

@@ -26,7 +26,6 @@
 #include "platform/runtime.hpp"
 
 #include <unordered_map>
-#include "elfio.hpp"
 
 constexpr unsigned __hipFatMAGIC2 = 0x48495046; // "HIPF"
 

projects/hip/rocclr/hip_stream.cpp

@@ -159,10 +159,14 @@ void iHipWaitActiveStreams(amd::HostQueue* blocking_queue, bool wait_null_stream
           (stream->Null() == wait_null_stream)) {
         // Get the last valid command
         amd::Command* command = active_queue->getLastQueuedCommand(true);
-        if ((command != nullptr) &&
-            // Check the current active status
-            (command->status() != CL_COMPLETE)) {
-          eventWaitList.push_back(command);
+        if (command != nullptr) {
+          // Check the current active status
+          if (command->status() != CL_COMPLETE) {
+            command->notifyCmdQueue();
+            eventWaitList.push_back(command);
+          } else {
+            command->release();
+          }
         }
       }
     }

projects/hip/tests/performance/module/hipPerfModuleLoad.cpp

@@ -0,0 +1,244 @@
+/*
+Copyright (c) 2015-2017 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.
+*/
+
+/* HIT_START
+ * BUILD_CMD: hipPerfModuleLoad %hc -I%S/../../src %S/%s %S/../../src/test_common.cpp -o %T/%t EXCLUDE_HIP_PLATFORM nvcc
+ * TEST: %t
+ * HIT_END
+ */
+
+#include "test_common.h"
+
+#include <vector>
+#include <unordered_map>
+#include <fstream>
+#include <unistd.h>
+#include <chrono>
+
+#ifdef __unix__
+
+#include <dirent.h>
+
+//List of Download files
+std::unordered_map<std::string, bool> TL_contents {
+  {"Kernels.so", true},
+  {"TensileLibrary.yaml", true},
+  {"TensileLibrary_gfx803.co", true},
+  {"TensileLibrary_gfx900.co", true},
+  {"TensileLibrary_gfx906.co", true},
+  {"TensileLibrary_gfx908.co", true},
+  {"kernel_names.txt", true}
+};
+
+bool GetDirectoryContents(std::unordered_map<std::string, bool>& dir_contents) {
+  DIR* dir = nullptr;
+  struct dirent* ent = nullptr;
+
+  //Open Current Directory
+  if ((dir = opendir(".")) == 0) {
+    std::cout<<"Failed to open current working directory, check permissions"<<std::endl;
+    return false;
+  }
+
+  // Read the contents of the directory
+  while ((ent = readdir(dir)) != nullptr) {
+    dir_contents.insert({ent->d_name, true});
+  }
+
+  closedir(dir);
+  return true;
+}
+
+bool ContentsAvailable() {
+  std::unordered_map<std::string, bool> dir_contents;
+
+  //Get recent directory contents
+  if(!GetDirectoryContents(dir_contents)) {
+    std::cout<<"Failed to get directory Contents"<<std::endl;
+    return false;
+  }
+
+  //If the Tensile Library content is not present, then fail
+  for (auto& TL_elem : TL_contents) {
+    if (dir_contents.end() == dir_contents.find(TL_elem.first)) {
+      std::cout<<"Failed to find the Tensile Library file: "<<TL_elem.first<<std::endl;
+      return false;
+    }
+  }
+  return true;
+}
+
+bool DownloadContents() {
+  //Download the contents from TC repo
+  std::cout<<"Downloading conents .... "<<std::endl;
+  std::string wget_str = "wget -nH -q -N -r -np -R \" index.html* \" --cut-dirs=3 ";
+  wget_str += "http://ocltc-backup.amd.com/hiptest/TensileLibrary/";
+  system(wget_str.c_str());
+
+  return true;
+}
+
+bool PreProcessContents() {
+  // If Contents already available no other action needed
+  if (ContentsAvailable()) {
+    std::cout<<"Contents already available"<<std::endl;
+    return true;
+  }
+
+  // Download the TL(Tensile Library) contents from TC
+  if (!DownloadContents()) {
+    std::cout<<"Failed to download contents"<<std::endl;
+    return false;
+  }
+
+  //Check if downloaded contents are available
+  if (!ContentsAvailable()) {
+    std::cout<<"Failed to find TL contents even after download in CWD"<<std::endl;
+    return false;
+  }
+
+  return true;
+}
+
+//Get Tensile Library File name, changes wrt target
+bool getTLFileName(int device_id, std::string& tlf_name) {
+  hipDeviceProp_t props;
+  HIPCHECK(hipGetDeviceProperties(&props, device_id));
+  std::string archName = props.gcnArchName;
+  if (archName.size() <= 3) {
+    std::cout<<"ArchName too small, Exiting"<<std::endl;
+    HIPASSERT(false);
+  }
+  archName = archName.substr(3, (archName.size()-1));
+
+  tlf_name = "TensileLibrary_gfx" + archName;
+  tlf_name += ".co";
+  return true;
+}
+
+bool RunTest(int device_id) {
+
+  std::cout<<"For Device: "<<device_id<<std::endl;
+
+  //Get Tensile Library File name, changes wrt target
+  std::string tlf_name;
+  if (!getTLFileName(device_id, tlf_name)) {
+    return false;
+  }
+
+  //Measure Time taken for hipModuleLoad
+  hipModule_t Module;
+  auto mload_clock_start = std::chrono::steady_clock::now();
+  HIPCHECK(hipModuleLoad(&Module, tlf_name.c_str()));
+  auto mload_clock_stop = std::chrono::steady_clock::now();
+  std::chrono::duration<double, std::nano> mload_duration = (mload_clock_stop - mload_clock_start);
+  std::cout<<"Time taken for hipModuleLoad : " <<std::chrono::duration_cast<std::chrono::nanoseconds>
+    (mload_duration).count()<<" nanoseconds "<<std::endl;
+
+  //Read kernels from a pre-populated text file
+  std::string kernel_file_name = "kernel_names.txt";
+  std::ifstream kernel_file(kernel_file_name);
+  if (!kernel_file.is_open()) {
+     std::cout<<"Failed to open Kernel File: "<<kernel_file_name<<std::endl;
+     return false;
+  }
+
+  std::string kernel_line;
+  std::vector<std::string> kernel_vec;
+  while (std::getline(kernel_file, kernel_line)) {
+    kernel_line.erase(std::remove(kernel_line.begin(), kernel_line.end(), '\r'),
+                      kernel_line.end());
+    kernel_vec.push_back(kernel_line);
+  }
+
+  //Measure the first hipModuleGetFunction
+  hipFunction_t hfunc = nullptr;
+  auto mgetf_clock_start = std::chrono::steady_clock::now();
+  HIPCHECK(hipModuleGetFunction(&hfunc, Module, kernel_vec[0].c_str()));
+  auto mgetf_clock_stop = std::chrono::steady_clock::now();
+  std::chrono::duration<double, std::nano> mgetf_duration = (mgetf_clock_stop - mgetf_clock_start);
+  std::cout<<"Time taken to fetch a function via hipModuleGetFunction : "
+    <<std::chrono::duration_cast<std::chrono::nanoseconds>
+    (mgetf_duration).count()<<" nanoseconds "<<std::endl;
+
+  //Measure the second hipModuleGetFunction
+  hfunc = nullptr;
+  mgetf_clock_start = std::chrono::steady_clock::now();
+  HIPCHECK(hipModuleGetFunction(&hfunc, Module, kernel_vec[0].c_str()));
+  mgetf_clock_stop = std::chrono::steady_clock::now();
+  mgetf_duration = (mgetf_clock_stop - mgetf_clock_start);
+  std::cout<<"Time taken fetch the same function via hipModuleGetFunction : "
+    <<std::chrono::duration_cast<std::chrono::nanoseconds>
+    (mgetf_duration).count()<<" nanoseconds "<<std::endl;
+
+  double all_duration = 0;
+  for (auto& kernel : kernel_vec) {
+    hfunc = nullptr;
+    mgetf_clock_start = std::chrono::steady_clock::now();
+    HIPCHECK(hipModuleGetFunction(&hfunc, Module, kernel.c_str()));
+    mgetf_clock_stop = std::chrono::steady_clock::now();
+    mgetf_duration = (mgetf_clock_stop - mgetf_clock_start);
+    all_duration += mgetf_duration.count();
+  }
+
+  if (kernel_vec.size() > 0) {
+    std::cout << "Time taken for Average hipModuleGetFunction : "
+      << (static_cast<double>(all_duration) / static_cast<double>(kernel_vec.size()))<<" nanoseconds "<<std::endl;
+  }
+
+  std::cout<<std::endl<<std::endl;
+
+  HIPCHECK(hipModuleUnload(Module));
+  return true;
+}
+#endif //__unix__
+
+int main() {
+  bool test_passed = true;
+
+  do {
+#ifdef __unix__
+    //Preprocess contents for the test
+    if (!PreProcessContents()) {
+      std::cout<<"Failed in PreProcessContents step"<<std::endl;
+      test_passed = false;
+      break;
+    }
+
+    //Run the test for all devices
+    int num_devices = 0;
+    HIPCHECK(hipGetDeviceCount(&num_devices));
+    for (int dev_idx = 0; dev_idx < num_devices; ++dev_idx) {
+      if (!RunTest(dev_idx)) {
+        test_passed = false;
+        break;
+      }
+    }
+#else
+    std::cout<<"Detected non-linux Os. Skipping the test"<<std::endl;
+#endif // __unix__
+  } while(0);
+
+  if (test_passed) {
+    passed();
+  }
+
+  return 0;
+}
+

projects/hip/tests/src/printf/printf_common.h

@@ -21,20 +21,17 @@ struct CaptureStream {
     orig_fd = fileno(original);
     saved_fd = dup(orig_fd);
 
-    temp_fd = mkstemp(tempname);
-    if (errno) {
+    if ((temp_fd = mkstemp(tempname)) == -1) {
       error(0, errno, "Error");
       assert(false);
     }
 
     fflush(nullptr);
-    dup2(temp_fd, orig_fd);
-    if (errno) {
+    if (dup2(temp_fd, orig_fd) == -1) {
       error(0, errno, "Error");
       assert(false);
     }
-    close(temp_fd);
-    if (errno) {
+    if (close(temp_fd) != 0) {
       error(0, errno, "Error");
       assert(false);
     }
@@ -44,13 +41,11 @@ struct CaptureStream {
     if (saved_fd == -1)
       return;
     fflush(nullptr);
-    dup2(saved_fd, orig_fd);
-    if (errno) {
+    if (dup2(saved_fd, orig_fd) == -1) {
       error(0, errno, "Error");
       assert(false);
     }
-    close(saved_fd);
-    if (errno) {
+    if (close(saved_fd) != 0) {
       error(0, errno, "Error");
       assert(false);
     }
@@ -65,8 +60,7 @@ struct CaptureStream {
 
   ~CaptureStream() {
     restoreStream();
-    remove(tempname);
-    if (errno) {
+    if (remove(tempname) != 0) {
       error(0, errno, "Error");
       assert(false);
     }