Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP into feature_get_alignment_and_size_from_metadata

2018-11-01 23:17:27 +00:00
commit f28cde3f98
@@ -207,9 +207,6 @@ if(HIP_PLATFORM STREQUAL "hcc")
        src/env.cpp
        src/program_state.cpp)

-    set(SOURCE_FILES_DEVICE
-        src/device_util.cpp)
-
    execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS)
    set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic")
    set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} --amdgpu-target=gfx701 --amdgpu-target=gfx803 --amdgpu-target=gfx900 --amdgpu-target=gfx906")
@@ -222,19 +219,18 @@ if(HIP_PLATFORM STREQUAL "hcc")
      target_link_libraries(hip_hcc PRIVATE hc_am)
      target_link_libraries(hip_hcc_static PRIVATE hc_am)
    endif()
-    add_library(hip_device STATIC ${SOURCE_FILES_DEVICE})

    string(REPLACE " " ";" HCC_CXX_FLAGS_LIST ${HCC_CXX_FLAGS})
-    foreach(TARGET hip_hcc hip_hcc_static hip_device)
+    foreach(TARGET hip_hcc hip_hcc_static)
        target_include_directories(${TARGET} SYSTEM INTERFACE $<INSTALL_INTERFACE:$<INSTALL_PREFIX>/include>;${HSA_PATH}/include)
    endforeach()
    add_library(host INTERFACE)
    target_link_libraries(host INTERFACE hip_hcc)
    add_library(device INTERFACE)
    if(HIP_COMPILER STREQUAL "hcc")
-        target_link_libraries(device INTERFACE host hip_device hcc::hccrt hcc::hc_am)
-    elseif(HIP_COMPILER STREQUAL "clang")
-        target_link_libraries(device INTERFACE host hip_device)
+        target_link_libraries(device INTERFACE host hcc::hccrt hcc::hc_am)
+    else()
+        target_link_libraries(device INTERFACE host)
    endif()

    # Generate .hipInfo
@@ -256,7 +252,7 @@ endif()
 #############################
 # Install hip_hcc if platform is hcc
 if(HIP_PLATFORM STREQUAL "hcc")
-    install(TARGETS hip_hcc_static hip_hcc hip_device DESTINATION lib)
+    install(TARGETS hip_hcc_static hip_hcc DESTINATION lib)

    # Install .hipInfo
    install(FILES ${PROJECT_BINARY_DIR}/.hipInfo DESTINATION lib)
@@ -284,7 +280,7 @@ set(BIN_INSTALL_DIR ${CMAKE_INSTALL_PREFIX}/bin)
 set(CONFIG_PACKAGE_INSTALL_DIR ${LIB_INSTALL_DIR}/cmake/hip)

 if(HIP_PLATFORM STREQUAL "hcc")
-    install(TARGETS hip_hcc_static hip_hcc hip_device host device EXPORT hip-targets DESTINATION ${LIB_INSTALL_DIR})
+    install(TARGETS hip_hcc_static hip_hcc host device EXPORT hip-targets DESTINATION ${LIB_INSTALL_DIR})
    install(EXPORT hip-targets DESTINATION ${CONFIG_PACKAGE_INSTALL_DIR} NAMESPACE hip::)
    include(CMakePackageConfigHelpers)

@@ -353,7 +349,7 @@ add_custom_target(pkg_hip_hcc  COMMAND ${CMAKE_COMMAND} .
    COMMAND cp *.rpm ${PROJECT_BINARY_DIR}
    COMMAND cp *.tar.gz ${PROJECT_BINARY_DIR}
    WORKING_DIRECTORY ${BUILD_DIR}
-    DEPENDS hip_hcc hip_device hip_hcc_static)
+    DEPENDS hip_hcc hip_hcc_static)

 # Package: hip_nvcc
 set(BUILD_DIR ${CMAKE_CURRENT_BINARY_DIR}/packages/hip_nvcc)
@@ -723,9 +723,9 @@ if ($setStdLib eq 0 and $HIP_PLATFORM eq 'hcc')

 if ($needHipHcc) {
    if ($linkType eq 0) {
-        substr($HIPLDFLAGS,0,0) = "  $HIP_PATH/lib/libhip_hcc_static.a $HIP_PATH/lib/libhip_device.a " ;
+        substr($HIPLDFLAGS,0,0) = "  $HIP_PATH/lib/libhip_hcc_static.a " ;
    } else {
-        substr($HIPLDFLAGS,0,0) = "  -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so $HIP_PATH/lib/libhip_device.a ";
+        substr($HIPLDFLAGS,0,0) = "  -Wl,--rpath=$HIP_PATH/lib $HIP_PATH/lib/libhip_hcc.so ";
    }
 }

@@ -11,6 +11,7 @@
 - [Build and install](#build-and-install)
     * [Building](#building)
     * [Testing](#testing)
+     * [Linux](#linux)
     * [Windows](#windows)
 - [Running and using hipify-clang](#running-and-using-hipify-clang)
 - [Disclaimer](#disclaimer)
@@ -80,7 +81,7 @@ Debug build type `-DCMAKE_BUILD_TYPE=Debug` is also supported and tested; `LLVM+

 The binary can then be found at `./dist/bin/hipify-clang`.

-### <a name="testing"></a> Test
+### <a name="testing"></a> Testing

 `hipify-clang` has unit tests using LLVM [`lit`](https://llvm.org/docs/CommandGuide/lit.html)/[`FileCheck`](https://llvm.org/docs/CommandGuide/FileCheck.html).

@@ -88,61 +89,183 @@ The binary can then be found at `./dist/bin/hipify-clang`.

 To run it:
 1. Download [`LLVM`](http://releases.llvm.org/6.0.1/llvm-6.0.1.src.tar.xz)+[`CLANG`](http://releases.llvm.org/6.0.1/cfe-6.0.1.src.tar.xz) sources.
-2. Build [`LLVM+CLANG`](http://llvm.org/docs/CMake.html).
-    For instance:
-    ```shell
+2. Build [`LLVM+CLANG`](http://llvm.org/docs/CMake.html):
+   ```shell
   cd llvm
   mkdir build dist
   cd build
+   ```
+     - **Linux**:

-   cmake \
-     -DCMAKE_INSTALL_PREFIX=../dist \
-     -DLLVM_SOURCE_DIR=../llvm \
-     -DCMAKE_BUILD_TYPE=Release \
-     -Thost=x64 \
-     ../llvm
+   ```shell
+        cmake \
+         -DCMAKE_INSTALL_PREFIX=../dist \
+         -DLLVM_SOURCE_DIR=../llvm \
+         -DCMAKE_BUILD_TYPE=Release \
+         ../llvm
+        make -j install
+   ```
+     - **Windows**:
+
+```shell
+        cmake \
+         -G "Visual Studio 15 2017 Win64" \
+         -DCMAKE_INSTALL_PREFIX=../dist \
+         -DLLVM_SOURCE_DIR=../llvm \
+         -DCMAKE_BUILD_TYPE=Release \
+         -Thost=x64 \
+         ../llvm
+```
+
+&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp;&nbsp; Run `Visual Studio 15 2017`, open the generated `LLVM.sln`, build all, build project `INSTALL`.

-    make -j install
-    ```
-    On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"`; the generated `LLVM.sln` should be built by `Visual Studio 15 2017` instead of `make`.

 3. Ensure [`CUDA`](https://developer.nvidia.com/cuda-toolkit-archive) of minimum version 7.5 is installed.

-    * Having multiple CUDA installations, in order to choose a particular version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified:
+    * Having multiple CUDA installations to choose a particular version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified:

-      `-DCUDA_TOOLKIT_ROOT_DIR="C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"`
+        - Linux: `-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-8.0`

-      * On Windows `CUDA_SDK_ROOT_DIR` option should be specified as well:
+        - Windows: `-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"`

-      `-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0"`
+          `-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v9.0"`

 4. Ensure [`cuDNN`](https://developer.nvidia.com/rdp/cudnn-archive) of version corresponding to CUDA's version is installed.

    * Path to cuDNN should be specified by the `CUDA_DNN_ROOT_DIR` option:

-      `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1`
+        - Linux: `-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-8.0-v7.1`
+
+        - Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1`

 5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed.

 6. Ensure `lit` and `FileCheck` are installed - these are distributed with LLVM.

-    * installing `lit` into `python` might be required:
+    * Install `lit` into `python`:

-      `python f:/LLVM/6.0.1/llvm/utils/lit/setup.py install`,
+        - Linux: `python /srv/git/LLVM/6.0.1/llvm/utils/lit/setup.py install`

-      where `f:/LLVM/6.0.1/llvm` is LLVM sources root directory.
+        - Windows: `python f:/LLVM/6.0.1/llvm/utils/lit/setup.py install`

-    * Starting with LLVM 6.0.1 path to llvm-lit.py script should be specified by the `LLVM_EXTERNAL_LIT` option:
+    * Starting with LLVM 6.0.1 path to `llvm-lit` python script should be specified by the `LLVM_EXTERNAL_LIT` option:

-      `-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py`,
+        - Linux: `-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/6.0.1/build/bin/llvm-lit`

-      where `f:/LLVM/6.0.1/build/Release` is LLVM build directory.
+        - Windows: `-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py`

-7. Build with the `HIPIFY_CLANG_TESTS` option turned on: -DHIPIFY_CLANG_TESTS=1.
+7. Set `HIPIFY_CLANG_TESTS` option turned on: `-DHIPIFY_CLANG_TESTS=1`.

-8. `make test-hipify`
+8. Run `cmake`:
+     * [Linux](#linux)
+     * [Windows](#windows)

-    On Windows after `cmake` the project `test-hipify` in the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make test-hipify`.
+9. Run tests:
+
+     - Linux: `make test-hipify`.
+
+     - Windows: run `Visual Studio 15 2017`, open the generated `hipify-clang.sln`, build project `test-hipify`.
+
+### <a name="linux"></a >Linux
+
+On Linux (Ubuntu 14-18) the following configurations are tested:
+
+LLVM 5.0.0 - 6.0.1, CUDA 8.0, cudnn-8.0
+
+Build system for the above configurations:
+
+Python 2.7 (min), cmake 3.5.2 (min), GNU C/C++ 5.4.0 (min).
+
+Here is an example of building `hipify-clang` with testing support on `Ubuntu 16.04`:
+
+```shell
+cmake
+ -DHIPIFY_CLANG_TESTS=1 \
+ -DCMAKE_BUILD_TYPE=Release \
+ -DCMAKE_INSTALL_PREFIX=../dist \
+ -DCMAKE_PREFIX_PATH=/srv/git/LLVM/6.0.1/dist \
+ -DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-8.0 \
+ -DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-8.0-v7.1 \
+ -DLLVM_EXTERNAL_LIT=/srv/git/LLVM/6.0.1/build/bin/llvm-lit \
+ ..
+```
+*A corresponding successful output:*
+```shell
+-- The C compiler identification is GNU 5.4.0
+-- The CXX compiler identification is GNU 5.4.0
+-- Check for working C compiler: /usr/bin/cc
+-- Check for working C compiler: /usr/bin/cc -- works
+-- Detecting C compiler ABI info
+-- Detecting C compiler ABI info - done
+-- Detecting C compile features
+-- Detecting C compile features - done
+-- Check for working CXX compiler: /usr/bin/c++
+-- Check for working CXX compiler: /usr/bin/c++ -- works
+-- Detecting CXX compiler ABI info
+-- Detecting CXX compiler ABI info - done
+-- Detecting CXX compile features
+-- Detecting CXX compile features - done
+-- Found LLVM 6.0.1:
+--    - CMake module path: /srv/git/LLVM/6.0.1/dist/lib/cmake/llvm
+--    - Include path     : /srv/git/LLVM/6.0.1/dist/include
+--    - Binary path      : /srv/git/LLVM/6.0.1/dist/bin
+-- Linker detection: GNU ld
+-- Found PythonInterp: /usr/bin/python2.7 (found suitable version "2.7.12", minimum required is "2.7")
+-- Found lit: /usr/local/bin/lit
+-- Found FileCheck: /srv/git/LLVM/6.0.1/dist/bin/FileCheck
+-- Looking for pthread.h
+-- Looking for pthread.h - found
+-- Looking for pthread_create
+-- Looking for pthread_create - not found
+-- Looking for pthread_create in pthreads
+-- Looking for pthread_create in pthreads - not found
+-- Looking for pthread_create in pthread
+-- Looking for pthread_create in pthread - found
+-- Found Threads: TRUE
+-- Found CUDA: /usr/local/cuda-8.0 (found version "8.0")
+-- Configuring done
+-- Generating done
+-- Build files have been written to: /srv/git/HIP/hipify-clang/build
+```
+```shell
+make test-hipify
+```
+*A corresponding successful output:*
+```shell
+[100%] Running HIPify regression tests
+-- Testing: 28 tests, 12 threads --
+PASS: hipify :: allocators.cu (1 of 28)
+PASS: hipify :: coalescing.cu (2 of 28)
+PASS: hipify :: cuDNN/cudnn_softmax.cu (3 of 28)
+PASS: hipify :: cuFFT/simple_cufft.cu (4 of 28)
+PASS: hipify :: cuComplex/cuComplex_Julia.cu (5 of 28)
+PASS: hipify :: cuBLAS/cublas_sgemm_matrix_multiplication.cu (6 of 28)
+PASS: hipify :: cuBLAS/cublas_1_based_indexing.cu (7 of 28)
+PASS: hipify :: cuBLAS/cublas_0_based_indexing.cu (8 of 28)
+PASS: hipify :: axpy.cu (9 of 28)
+PASS: hipify :: dynamic_shared_memory.cu (10 of 28)
+PASS: hipify :: headers_test_01.cu (11 of 28)
+PASS: hipify :: headers_test_02.cu (12 of 28)
+PASS: hipify :: headers_test_03.cu (13 of 28)
+PASS: hipify :: headers_test_05.cu (14 of 28)
+PASS: hipify :: cuDNN/cudnn_convolution_forward.cu (15 of 28)
+PASS: hipify :: cuRAND/poisson_api_example.cu (16 of 28)
+PASS: hipify :: cudaRegister.cu (17 of 28)
+PASS: hipify :: headers_test_06.cu (18 of 28)
+PASS: hipify :: headers_test_04.cu (19 of 28)
+PASS: hipify :: intro.cu (20 of 28)
+PASS: hipify :: headers_test_07.cu (21 of 28)
+PASS: hipify :: square.cu (22 of 28)
+PASS: hipify :: static_shared_memory.cu (23 of 28)
+PASS: hipify :: vec_add.cu (24 of 28)
+PASS: hipify :: headers_test_08.cu (25 of 28)
+PASS: hipify :: cuRAND/benchmark_curand_generate.cpp (26 of 28)
+PASS: hipify :: cuRAND/benchmark_curand_kernel.cpp (27 of 28)
+PASS: hipify :: headers_test_09.cu (28 of 28)
+Testing Time: 1.71s
+  Expected Passes    : 28
+[100%] Built target test-hipify
+```

 ### <a name="windows"></a >Windows

@@ -172,7 +295,7 @@ cmake
 -Thost=x64
 ..
 ```
-A corresponding successful output:
+*A corresponding successful output:*
 ```shell
 -- Found LLVM 6.0.1:
 --    - CMake module path: F:/LLVM/6.0.1/dist/lib/cmake/llvm
@@ -194,12 +317,13 @@ To process a file, `hipify-clang` needs access to the same headers that would be
 For example:

 ```shell
-hipify-clang square.cu -- \
+./hipify-clang \
+  square.cu \
+  -- \
  -x cuda \
-  --cuda-path=/opt/cuda \
-  --cuda-gpu-arch=sm_30 \
-  -isystem /opt/cuda/samples/common/inc
-  -I /opt/cuda/cuDNN
+  --cuda-path=/usr/local/cuda-8.0 \
+  --cuda-gpu-arch=sm_50 \
+  -isystem /usr/local/cuda-8.0/samples/common/inc
 ```

 `hipify-clang` arguments are given first, followed by a separator, and then the arguments you'd pass to `clang` if you
@@ -14,7 +14,7 @@ const char *counterNames[NUM_CONV_TYPES] = {
  "addressing", // CONV_ADDRESSING
  "stream", // CONV_STREAM
  "event", // CONV_EVENT
-  "external_resource_interop" // CONV_EXT_RES
+  "external_resource_interop", // CONV_EXT_RES
  "stream_memory", // CONV_STREAM_MEMORY
  "execution", // CONV_EXECUTION
  "graph", // CONV_GRAPH
@@ -74,7 +74,7 @@ void printStat(std::ostream *csv, llvm::raw_ostream* printOut, const std::string

 } // Anonymous namespace

-void StatCounter::incrementCounter(const hipCounter& counter, std::string name) {
+void StatCounter::incrementCounter(const hipCounter& counter, const std::string& name) {
  counters[name]++;
  apiCounters[(int) counter.apiType]++;
  convTypeCounters[(int) counter.type]++;
@@ -100,7 +100,7 @@ int StatCounter::getConvSum() {
  return acc;
 }

-void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix) {
+void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& prefix) {
  conditionalPrint(csv, printOut, "\nCUDA ref type;Count\n", "[HIPIFY] info: " + prefix + " refs by type:\n");
  for (int i = 0; i < NUM_CONV_TYPES; i++) {
    if (convTypeCounters[i] > 0) {
@@ -117,7 +117,7 @@ void StatCounter::print(std::ostream* csv, llvm::raw_ostream* printOut, std::str
  }
 }

-Statistics::Statistics(std::string name): fileName(name) {
+Statistics::Statistics(const std::string& name): fileName(name) {
  // Compute the total bytes/lines in the input file.
  std::ifstream src_file(name, std::ios::binary | std::ios::ate);
  src_file.clear();
@@ -129,7 +129,7 @@ Statistics::Statistics(std::string name): fileName(name) {

 ///////// Counter update routines //////////

-void Statistics::incrementCounter(const hipCounter &counter, std::string name) {
+void Statistics::incrementCounter(const hipCounter &counter, const std::string& name) {
  if (counter.unsupported) {
    unsupported.incrementCounter(counter, name);
  } else {
@@ -218,7 +218,7 @@ Statistics& Statistics::current() {
  return *Statistics::currentStatistics;
 }

-void Statistics::setActive(std::string name) {
+void Statistics::setActive(const std::string& name) {
  stats.emplace(std::make_pair(name, Statistics{name}));
  Statistics::currentStatistics = &stats.at(name);
 }
@@ -22,7 +22,7 @@ enum ConvTypes {
  // Driver API : 5.5. Device Management, 5.6. Device Management [DEPRECATED]
  // Runtime API: 5.1. Device Management
  CONV_DEVICE,
-  // Driver API : 5.8.Context Management, 5.9. Context Management [DEPRECATED]
+  // Driver API : 5.7. Primary Context Management, 5.8.Context Management, 5.9. Context Management [DEPRECATED]
  CONV_CONTEXT,
  // Driver API : 5.10. Module Management
  CONV_MODULE,
@@ -134,11 +134,11 @@ private:
  int convTypeCounters[NUM_CONV_TYPES] = {};

 public:
-  void incrementCounter(const hipCounter& counter, std::string name);
+  void incrementCounter(const hipCounter& counter, const std::string& name);
  // Add the counters from `other` onto the counters of this object.
  void add(const StatCounter& other);
  int getConvSum();
-  void print(std::ostream* csv, llvm::raw_ostream* printOut, std::string prefix);
+  void print(std::ostream* csv, llvm::raw_ostream* printOut, const std::string& prefix);
 };

 /**
@@ -156,8 +156,8 @@ class Statistics {
  chr::steady_clock::time_point completionTime;

 public:
-  Statistics(std::string name);
-  void incrementCounter(const hipCounter &counter, std::string name);
+  Statistics(const std::string& name);
+  void incrementCounter(const hipCounter &counter, const std::string& name);
  // Add the counters from `other` onto the counters of this object.
  void add(const Statistics &other);
  void lineTouched(int lineNumber);
@@ -192,5 +192,5 @@ public:
    * Set the active Statistics object to the named one, creating it if necessary, and write the completion
    * timestamp into the currently active one.
    */
-  static void setActive(std::string name);
+  static void setActive(const std::string& name);
 };
@@ -7,7 +7,7 @@ llvm::StringRef unquoteStr(llvm::StringRef s) {
  return s;
 }

-void removePrefixIfPresent(std::string &s, std::string prefix) {
+void removePrefixIfPresent(std::string &s, const std::string& prefix) {
  if (s.find(prefix) != 0) {
    return;
  }
@@ -11,4 +11,4 @@ llvm::StringRef unquoteStr(llvm::StringRef s);
 /**
  * If `s` starts with `prefix`, remove it. Otherwise, does nothing.
  */
-void removePrefixIfPresent(std::string &s, std::string prefix);
+void removePrefixIfPresent(std::string &s, const std::string& prefix);
@@ -84,6 +84,9 @@ class Bundled_code_header {
                std::copy_n(f + y.header.offset, y.header.bundle_sz, std::back_inserter(y.blob));

                it += y.header.triple_sz;
+
+                x.bundled_code_size = std::max(x.bundled_code_size, 
+                                               y.header.offset + y.header.bundle_sz);
            }

            return true;
@@ -123,6 +126,8 @@ class Bundled_code_header {
    // MANIPULATORS
    Bundled_code_header& operator=(const Bundled_code_header&) = default;
    Bundled_code_header& operator=(Bundled_code_header&&) = default;
+
+    size_t bundled_code_size = 0;
 };

 // CREATORS
@@ -876,6 +876,15 @@ static void __threadfence_system()
  __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_all_svm_devices);
 }

+// abort
+__device__
+inline
+__attribute__((weak))
+void abort() {
+  return __builtin_trap();
+}
+
+
 #endif // __HCC_OR_HIP_CLANG__

 #ifdef __HCC__
@@ -181,8 +181,7 @@ extern int HIP_TRACE_API;
 #define __HCC_C__
 #endif

-// abort
-__device__ void abort();
+__host__ inline void* __get_dynamicgroupbaseptr() { return nullptr; }

 #if __HIP_ARCH_GFX701__ == 0

@@ -94,7 +94,7 @@ typedef struct ihipModule_t* hipModule_t;

 typedef struct ihipModuleSymbol_t* hipFunction_t;

-struct hipFuncAttributes {
+typedef struct hipFuncAttributes {
    int binaryVersion;
    int cacheModeCA;
    size_t constSizeBytes;
@@ -105,7 +105,7 @@ struct hipFuncAttributes {
    int preferredShmemCarveout;
    int ptxVersion;
    size_t sharedSizeBytes;
-};
+} hipFuncAttributes;

 typedef struct ihipEvent_t* hipEvent_t;

@@ -41,15 +41,5 @@ endif()
 list(APPEND _IMPORT_CHECK_TARGETS hip::hip_hcc )
 list(APPEND _IMPORT_CHECK_FILES_FOR_hip::hip_hcc "/opt/rocm/hip/lib/libhip_hcc.so" )

-# Import target "hip::hip_device" for configuration "Release"
-set_property(TARGET hip::hip_device APPEND PROPERTY IMPORTED_CONFIGURATIONS RELEASE)
-set_target_properties(hip::hip_device PROPERTIES
-  IMPORTED_LINK_INTERFACE_LANGUAGES_RELEASE "CXX"
-  IMPORTED_LOCATION_RELEASE "/opt/rocm/hip/lib/libhip_device.a"
-  )
-
-list(APPEND _IMPORT_CHECK_TARGETS hip::hip_device )
-list(APPEND _IMPORT_CHECK_FILES_FOR_hip::hip_device "/opt/rocm/hip/lib/libhip_device.a" )
-
 # Commands beyond this point should not need to know the version.
 set(CMAKE_IMPORT_FILE_VERSION)
@@ -16,7 +16,7 @@ set(CMAKE_IMPORT_FILE_VERSION 1)
 set(_targetsDefined)
 set(_targetsNotDefined)
 set(_expectedTargets)
-foreach(_expectedTarget hip::hip_hcc_static hip::hip_hcc hip::hip_device hip::host hip::device)
+foreach(_expectedTarget hip::hip_hcc_static hip::hip_hcc hip::host hip::device)
  list(APPEND _expectedTargets ${_expectedTarget})
  if(NOT TARGET ${_expectedTarget})
    list(APPEND _targetsNotDefined ${_expectedTarget})
@@ -57,14 +57,6 @@ set_target_properties(hip::hip_hcc PROPERTIES
  INTERFACE_SYSTEM_INCLUDE_DIRECTORIES "${_IMPORT_PREFIX}/include;/opt/rocm/hsa/include"
 )

-# Create imported target hip::hip_device
-add_library(hip::hip_device STATIC IMPORTED)
-
-set_target_properties(hip::hip_device PROPERTIES
-  INTERFACE_INCLUDE_DIRECTORIES "${_IMPORT_PREFIX}/include;/opt/rocm/hsa/include"
-  INTERFACE_SYSTEM_INCLUDE_DIRECTORIES "${_IMPORT_PREFIX}/include;/opt/rocm/hsa/include"
-)
-
 # Create imported target hip::host
 add_library(hip::host INTERFACE IMPORTED)

@@ -75,13 +67,13 @@ set_target_properties(hip::host PROPERTIES
 # Create imported target hip::device
 add_library(hip::device INTERFACE IMPORTED)

-if(HIP_COMPILER STREQUAL "clang")
+if(HIP_COMPILER STREQUAL "hcc")
 set_target_properties(hip::device PROPERTIES
-  INTERFACE_LINK_LIBRARIES "hip::host;hip::hip_device"
+  INTERFACE_LINK_LIBRARIES "hip::host;hcc::hccrt;hcc::hc_am"
 )
 else()
 set_target_properties(hip::device PROPERTIES
-  INTERFACE_LINK_LIBRARIES "hip::host;hip::hip_device;hcc::hccrt;hcc::hc_am"
+  INTERFACE_LINK_LIBRARIES "hip::host"
 )
 endif()

@@ -3,7 +3,6 @@ project(hip_hcc)

 install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib)
 install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib)
-install(FILES @PROJECT_BINARY_DIR@/libhip_device.a DESTINATION lib)
 install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib)
 install(FILES @PROJECT_BINARY_DIR@/hip-config.cmake @PROJECT_BINARY_DIR@/hip-config-version.cmake DESTINATION lib/cmake/hip)
 install(FILES @hip_SOURCE_DIR@/packaging/hip-targets.cmake @hip_SOURCE_DIR@/packaging/hip-targets-release.cmake DESTINATION lib/cmake/hip)
@@ -1,33 +0,0 @@
-/*
-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.
-*/
-
-#include <hc.hpp>
-#include "device_util.h"
-#include "hip/hcc_detail/device_functions.h"
-#include "hip/hcc_detail/grid_launch.h"
-#include "hip/hip_runtime.h"
-#include <atomic>
-
-// abort
-__device__ void abort() { return hc::abort(); }
-
-__host__ void* __get_dynamicgroupbaseptr() { return nullptr; }
@@ -574,24 +574,8 @@ hipError_t hipMallocArray(hipArray** array, const hipChannelFormatDesc* desc, si
                size = size * height;
            }

-            size_t allocSize = 0;
-            switch (desc->f) {
-                case hipChannelFormatKindSigned:
-                    allocSize = size * sizeof(int);
-                    break;
-                case hipChannelFormatKindUnsigned:
-                    allocSize = size * sizeof(unsigned int);
-                    break;
-                case hipChannelFormatKindFloat:
-                    allocSize = size * sizeof(float);
-                    break;
-                case hipChannelFormatKindNone:
-                    allocSize = size * sizeof(size_t);
-                    break;
-                default:
-                    hip_status = hipErrorUnknown;
-                    break;
-            }
+            const size_t allocSize = size * ((desc->x + desc->y + desc->z + desc->w) / 8);
+
            hc::accelerator acc = ctx->getDevice()->_acc;
            hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());

@@ -800,24 +784,7 @@ hipError_t hipMalloc3DArray(hipArray** array, const struct hipChannelFormatDesc*
        const unsigned am_flags = 0;
        const size_t size = extent.width * extent.height * extent.depth;

-        size_t allocSize = 0;
-        switch (desc->f) {
-            case hipChannelFormatKindSigned:
-                allocSize = size * sizeof(int);
-                break;
-            case hipChannelFormatKindUnsigned:
-                allocSize = size * sizeof(unsigned int);
-                break;
-            case hipChannelFormatKindFloat:
-                allocSize = size * sizeof(float);
-                break;
-            case hipChannelFormatKindNone:
-                allocSize = size * sizeof(size_t);
-                break;
-            default:
-                hip_status = hipErrorUnknown;
-                break;
-        }
+        const size_t allocSize = size * ((desc->x + desc->y + desc->z + desc->w) / 8);

        hc::accelerator acc = ctx->getDevice()->_acc;
        hsa_agent_t* agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
@@ -1688,7 +1655,7 @@ hipError_t hipMemcpy2DAsync(void* dst, size_t dpitch, const void* src, size_t sp
    } else {
        try {
            if(!isLocked){
-                for (int i = 0; i < height; ++i) 
+                for (int i = 0; i < height; ++i)
                    e = hip_internal::memcpyAsync((unsigned char*)dst + i * dpitch,
                                          (unsigned char*)src + i * spitch, width, kind, stream);
            } else{
@@ -1738,7 +1705,7 @@ hipError_t hipMemset(void* dst, int value, size_t sizeBytes) {
        stream->locked_wait();
    } else {
        e = hipErrorInvalidValue;
-    } 
+    }
    return ihipLogStatus(e);
 }

@@ -209,10 +209,16 @@ const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs(bool reb
            nullptr);

        for (auto&& blob : blobs) {
-            Bundled_code_header tmp{blob};
-            if (valid(tmp)) {
-                for (auto&& bundle : bundles(tmp)) {
-                    r[triple_to_hsa_isa(bundle.triple)].push_back(bundle.blob);
+            for (auto sub_blob = blob.begin(); sub_blob != blob.end(); ) {
+                Bundled_code_header tmp(sub_blob, blob.end());
+                if (valid(tmp)) {
+                    for (auto&& bundle : bundles(tmp)) {
+                        r[triple_to_hsa_isa(bundle.triple)].push_back(bundle.blob);
+                    }
+                    sub_blob+=tmp.bundled_code_size;
+                }
+                else {
+                    break;
                }
            }
        }