From 43bb490947bbae7a10de30598c1a699315d620fa Mon Sep 17 00:00:00 2001
From: Yaxun Sam Liu <yaxun.liu@amd.com>
Date: Tue, 31 Jul 2018 18:20:13 -0400
Subject: [PATCH 01/16] Add HCC compatibility mode

---
 bin/hipcc                            |  5 +++++
 include/hip/hcc_detail/hip_runtime.h | 20 ++++++++++++++++++++
 2 files changed, 25 insertions(+)

diff --git a/bin/hipcc b/bin/hipcc
index 9df3459fd1..1eebb571d5 100755
--- a/bin/hipcc
+++ b/bin/hipcc
@@ -57,6 +57,7 @@ $HIP_PATH=$ENV{'HIP_PATH'} // dirname (dirname $0); # use parent directory of hi
 $HIP_VDI_HOME=$ENV{'HIP_VDI_HOME'};
 $HIP_CLANG_PATH=$ENV{'HIP_CLANG_PATH'};
 $DEVICE_LIB_PATH=$ENV{'DEVICE_LIB_PATH'};
+$HIP_CLANG_HCC_COMPAT_MODE=$ENV{'HIP_CLANG_HCC_COMPAT_MODE'}; # HCC compatibility mode
 
 #---
 # Read .hipInfo
@@ -158,6 +159,10 @@ if ($HIP_PLATFORM eq "clang") {
 
     $HIPCXXFLAGS .= " -std=c++11 -isystem $HIP_CLANG_INCLUDE_PATH";
     $HIPLDFLAGS .= " --hip-device-lib-path=$DEVICE_LIB_PATH -L$HIP_LIB_PATH -Wl,--rpath=$HIP_LIB_PATH -lhip_hcc";
+    if ($HIP_CLANG_HCC_COMPAT_MODE) {
+        ## Allow __fp16 as function parameter and return type.
+        $HIPCXXFLAGS .= " -Xclang -fallow-half-arguments-and-returns -D__HIP_HCC_COMPAT_MODE__=1";
+    }
 } elsif ($HIP_PLATFORM eq "hcc") {
     $HIP_INCLUDE_PATH = "$HIP_PATH/include";
     $HSA_PATH=$ENV{'HSA_PATH'} // "/opt/rocm/hsa";
diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index c2ae6e8e4f..61a614ee15 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -438,6 +438,26 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim;
 #define hipGridDim_y gridDim.y
 #define hipGridDim_z gridDim.z
 
+#if __HIP_HCC_COMPAT_MODE__
+// Define HCC work item functions in terms of HIP builtin variables.
+#pragma push_macro("__DEFINE_HCC_FUNC")
+#define __DEFINE_HCC_FUNC(hc_fun,hip_var) \
+inline __device__ __attribute__((always_inline)) uint hc_get_##hc_fun(uint i) { \
+  if (i==0) \
+    return hip_var.x; \
+  else if(i==1) \
+    return hip_var.y; \
+  else \
+    return hip_var.z; \
+}
+
+__DEFINE_HCC_FUNC(workitem_id, threadIdx)
+__DEFINE_HCC_FUNC(group_id, blockIdx)
+__DEFINE_HCC_FUNC(group_size, blockDim)
+__DEFINE_HCC_FUNC(num_groups, gridDim)
+#pragma pop_macro("__DEFINE_HCC_FUNC")
+#endif
+
 // Support std::complex.
 #pragma push_macro("__CUDA__")
 #define __CUDA__

From ef1d1d57f7ac0a935a0f1b22d37d049b3cac61f3 Mon Sep 17 00:00:00 2001
From: Yaxun Sam Liu <yaxun.liu@amd.com>
Date: Wed, 1 Aug 2018 15:13:07 -0400
Subject: [PATCH 02/16] Add hc_get_workitem_absolute_id for hip-clang HCC
 compatibility mode

---
 include/hip/hcc_detail/hip_runtime.h | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h
index 61a614ee15..0b44c5f9c0 100644
--- a/include/hip/hcc_detail/hip_runtime.h
+++ b/include/hip/hcc_detail/hip_runtime.h
@@ -456,6 +456,14 @@ __DEFINE_HCC_FUNC(group_id, blockIdx)
 __DEFINE_HCC_FUNC(group_size, blockDim)
 __DEFINE_HCC_FUNC(num_groups, gridDim)
 #pragma pop_macro("__DEFINE_HCC_FUNC")
+
+extern "C" __device__ __attribute__((const)) size_t __ockl_get_global_id(uint);
+inline __device__ __attribute__((always_inline)) uint
+hc_get_workitem_absolute_id(int dim)
+{
+  return (uint)__ockl_get_global_id(dim);
+}
+
 #endif
 
 // Support std::complex.

From 704f236c9a59658c185c707f259404f458ec9458 Mon Sep 17 00:00:00 2001
From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Wed, 1 Aug 2018 20:37:46 +0000
Subject: [PATCH 03/16] Enable HIP_COMPILER to choose HIP-Clang lib links

---
 CMakeLists.txt | 21 ++++++++++++++++++++-
 1 file changed, 20 insertions(+), 1 deletion(-)

diff --git a/CMakeLists.txt b/CMakeLists.txt
index a4da3b1920..818c75783e 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -58,6 +58,21 @@ if(NOT DEFINED HIP_PLATFORM)
 endif()
 message(STATUS "HIP Platform: " ${HIP_PLATFORM})
 
+# Determine HIP_COMPILER
+# Either hcc or clang; default is hcc
+if(NOT DEFINED HIP_COMPILER)
+    if(NOT DEFINED ENV{HIP_COMPILER})
+        set(HIP_COMPILER "hcc" CACHE STRING "HIP Compiler")
+    else()
+        set(HIP_COMPILER $ENV{HIP_COMPILER} CACHE STRING "HIP Compiler")
+    endif()
+endif()
+if(NOT (HIP_COMPILER STREQUAL "hcc" OR HIP_COMPILER STREQUAL "clang"))
+    message(FATAL_ERROR "Must use HIP_COMPILER as hcc or clang")
+endif()
+message(STATUS "HIP Compiler: " ${HIP_COMPILER})
+
+
 # If HIP_PLATFORM is hcc, we need HCC_HOME and HSA_PATH to be defined
 if(HIP_PLATFORM STREQUAL "hcc")
     # Determine HCC_HOME
@@ -214,7 +229,11 @@ if(HIP_PLATFORM STREQUAL "hcc")
     add_library(host INTERFACE)
     target_link_libraries(host INTERFACE hip_hcc)
     add_library(device INTERFACE)
-    target_link_libraries(device INTERFACE host hip_device hcc::hccrt hcc::hc_am)
+    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)
+    endif()
 
     # Generate .hipInfo
     file(WRITE "${PROJECT_BINARY_DIR}/.hipInfo" ${_buildInfo})

From f6316a95b88355770ab99f0528c0dcefc6fefed4 Mon Sep 17 00:00:00 2001
From: Yaxun Sam Liu <yaxun.liu@amd.com>
Date: Wed, 1 Aug 2018 17:01:39 -0400
Subject: [PATCH 04/16] Let hipcc handle static library for hip-clang

---
 bin/hipcc | 47 +++++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 47 insertions(+)

diff --git a/bin/hipcc b/bin/hipcc
index 9df3459fd1..c277fc42aa 100755
--- a/bin/hipcc
+++ b/bin/hipcc
@@ -455,6 +455,53 @@ foreach $arg (@ARGV)
         close $in;
         close $out;
         $arg = "$new_arg -Wl,\@$new_file";
+    } elsif (($arg =~ m/\.a$/ || $arg =~ m/\.lo$/) &&
+             $HIP_PLATFORM eq 'clang') {
+        ## process static library for hip-clang
+        ## extract object files from static library and pass them directly to
+        ## hip-clang.
+        ## ToDo: Remove this after hip-clang switch to lto and lld is able to
+        ## handle clang-offload-bundler bundles.
+        
+        my $new_arg = "";
+        my $tmpdir = get_temp_dir ();
+        my $libFile = $arg;
+        my $path = abs_path($arg);
+        my @objs = split ('\n', `cd $tmpdir; ar xv $path`);
+        ## Check if all files in .a are object files.
+        my $allIsObj = 1;
+        my $realObjs = "";
+        foreach my $obj (@objs) {
+            chomp $obj;
+            $obj =~ s/^x - //;
+            $obj = "$tmpdir/$obj";
+            my $fileType = `file $obj`;
+            my $isObj = ($fileType =~ m/ELF/ or $fileType =~ m/COFF/);
+            $allIsObj = ($allIsObj and $isObj);
+            if ($isObj) {
+                $realObjs = ($realObjs . " " . $obj);
+            } else {
+                push (@inputs, $obj);
+                if ($new_arg ne "") {
+                    $new_arg .= " ";
+                }
+                $new_arg .= "$obj";
+            }
+        }
+        chomp $realObjs;
+        if ($allIsObj) {
+            $new_arg = $arg;
+        } elsif ($realObjs) {
+            my($libBaseName, $libDir, $libExt) = fileparse($libFile);
+            $libBaseName = mktemp($libBaseName . "XXXX") . $libExt;
+            system("cd $tmpdir; ar c $libBaseName $realObjs");
+            $new_arg .= " $tmpdir/$libBaseName";
+        }
+        $arg = "$new_arg";
+        if ($toolArgs =~ m/-Xlinker$/) {
+            $toolArgs = substr $toolArgs, 0, -8;
+            chomp $toolArgs;
+        }
     } elsif ($arg =~ m/^-/) {
         # options start with -
 

From 07f8f09aff211b2e67422ef70043813895359e1d Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Thu, 2 Aug 2018 15:16:36 +0530
Subject: [PATCH 05/16] Remove adipose extension from genco output

---
 lpl_ca/lpl.hpp | 6 +++---
 1 file changed, 3 insertions(+), 3 deletions(-)

diff --git a/lpl_ca/lpl.hpp b/lpl_ca/lpl.hpp
index cbd7fe8386..c3992e43c0 100644
--- a/lpl_ca/lpl.hpp
+++ b/lpl_ca/lpl.hpp
@@ -80,7 +80,7 @@ inline void copy_kernel_section_to_fat_binary(const std::string& tmp, const std:
         std::find_if(reader.sections.begin(), reader.sections.end(),
                      [](const ELFIO::section* x) { return x->get_name() == kernel_section(); });
 
-    std::ofstream out{output + fat_binary_extension()};
+    std::ofstream out{output};
 
     if (it == reader.sections.end()) {
         std::cerr << "Warning: no kernels were generated; fat binary shall "
@@ -95,8 +95,8 @@ inline void generate_fat_binary(const std::vector<std::string>& sources,
                                 const std::vector<std::string>& targets, const std::string& flags,
                                 const std::string& output) {
     static const auto d = [](const std::string* f) { remove(f->c_str()); };
-
-    std::unique_ptr<const std::string, decltype(d)> tmp{&output, d};
+    std::string temp_str = output + ".tmp";
+    std::unique_ptr<const std::string, decltype(d)> tmp{&temp_str, d};
 
     redi::ipstream hipcc{make_hipcc_call(sources, targets, flags, *tmp), redi::pstream::pstderr};
 

From 136bcc298100a698a2e0e451842ebc2895398ee8 Mon Sep 17 00:00:00 2001
From: "Wen-Heng (Jack) Chung" <whchung@gmail.com>
Date: Fri, 3 Aug 2018 17:02:50 -0500
Subject: [PATCH 06/16] Revert "Keep the map which tracks GPU kernel symbols to
 grow monotonically"

This reverts commit 32789a8b7dee21e8e4f320e8d673591b01113ced.
---
 src/program_state.cpp | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/src/program_state.cpp b/src/program_state.cpp
index fab71b1def..3eba84f12e 100644
--- a/src/program_state.cpp
+++ b/src/program_state.cpp
@@ -444,9 +444,7 @@ const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& fu
 
     auto cons = [rebuild]() {
         if (rebuild) {
-            // do NOT clear r so we reuse instances of pair<hsa_agent_t, Kernel_descriptor>
-            // created previously
-
+            r.clear();
             function_names(rebuild);
             kernels(rebuild);
             globals(rebuild);

From 3426f151718142f76c81dfbc4084f48701d88314 Mon Sep 17 00:00:00 2001
From: "Wen-Heng (Jack) Chung" <whchung@gmail.com>
Date: Fri, 3 Aug 2018 17:02:58 -0500
Subject: [PATCH 07/16] Revert "Improve performance of re-initialization logic"

This reverts commit ece4539c1da091c749b24a73aaba9dc40ec99567.
---
 src/program_state.cpp | 17 ++++-------------
 1 file changed, 4 insertions(+), 13 deletions(-)

diff --git a/src/program_state.cpp b/src/program_state.cpp
index 3eba84f12e..3ad9285b8c 100644
--- a/src/program_state.cpp
+++ b/src/program_state.cpp
@@ -180,29 +180,21 @@ const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs(bool reb
     static once_flag f;
 
     auto cons = [rebuild]() {
-        // names of shared libraries who .kernel sections already loaded
-        static unordered_set<string> lib_names;
         static vector<vector<char>> blobs{code_object_blob_for_process()};
 
         if (rebuild) {
-            r.clear();
             blobs.clear();
+            blobs.push_back(code_object_blob_for_process());
         }
 
         dl_iterate_phdr(
             [](dl_phdr_info* info, std::size_t, void*) {
                 elfio tmp;
-                if ((lib_names.find(info->dlpi_name) == lib_names.end()) &&
-                    (tmp.load(info->dlpi_name))) {
+                if (tmp.load(info->dlpi_name)) {
                     const auto it = find_section_if(
                         tmp, [](const section* x) { return x->get_name() == ".kernel"; });
 
-                    if (it) {
-                        blobs.emplace_back(
-                            it->get_data(), it->get_data() + it->get_size());
-                        // register the shared library as already loaded
-                        lib_names.emplace(info->dlpi_name);
-                    }
+                    if (it) blobs.emplace_back(it->get_data(), it->get_data() + it->get_size());
                 }
                 return 0;
             },
@@ -346,8 +338,7 @@ executables(bool rebuild) {  // TODO: This leaks the hsa_executable_ts, it shoul
         static const auto accelerators = hc::accelerator::get_all();
 
         if (rebuild) {
-            // do NOT clear r so we reuse instances of hsa_executable_t
-            // created previously
+            r.clear();
             code_object_blobs(rebuild);
         }
 

From 2604f3393003281b0fee6731d32f50ddeb814326 Mon Sep 17 00:00:00 2001
From: "Wen-Heng (Jack) Chung" <whchung@gmail.com>
Date: Fri, 3 Aug 2018 17:03:04 -0500
Subject: [PATCH 08/16] Revert "HIP program state re-initialization logic"

This reverts commit 379b7a224138acd54c76dac8fe23f40c5d1b53cc.
---
 include/hip/hcc_detail/program_state.hpp |   9 +-
 src/functional_grid_launch.inl           |  16 +--
 src/program_state.cpp                    | 131 ++++-------------------
 3 files changed, 32 insertions(+), 124 deletions(-)

diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp
index bdb87b3509..ac689fdb89 100644
--- a/include/hip/hcc_detail/program_state.hpp
+++ b/include/hip/hcc_detail/program_state.hpp
@@ -93,12 +93,11 @@ public:
     }
 };
 
-const std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>& executables(
-    bool rebuild = false);
+const std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>& executables();
 const std::unordered_map<std::uintptr_t, std::vector<std::pair<hsa_agent_t, Kernel_descriptor>>>&
-functions(bool rebuild = false);
-const std::unordered_map<std::uintptr_t, std::string>& function_names(bool rebuild = false);
-std::unordered_map<std::string, void*>& globals(bool rebuild = false);
+functions();
+const std::unordered_map<std::uintptr_t, std::string>& function_names();
+std::unordered_map<std::string, void*>& globals();
 
 hsa_executable_t load_executable(const std::string& file, hsa_executable_t executable,
                                  hsa_agent_t agent);
diff --git a/src/functional_grid_launch.inl b/src/functional_grid_launch.inl
index 6283d1aaba..704cf9ffeb 100644
--- a/src/functional_grid_launch.inl
+++ b/src/functional_grid_launch.inl
@@ -92,19 +92,13 @@ namespace hip_impl
         hipStream_t stream,
         void** kernarg)
     {
-        auto it0 = functions().find(function_address);
+        const auto it0 = functions().find(function_address);
 
         if (it0 == functions().cend()) {
-            // Re-init device code maps once again to help locate kernels
-            // loaded after HIP runtime initialization via means such as
-            // dlopen().
-            it0 = functions(true).find(function_address);
-            if (it0 == functions().cend()) {
-                throw runtime_error{
-                    "No device code available for function: " +
-                    name(function_address)
-                };
-            }
+            throw runtime_error{
+                "No device code available for function: " +
+                name(function_address)
+            };
         }
 
         auto agent = target_agent(stream);
diff --git a/src/program_state.cpp b/src/program_state.cpp
index 3ad9285b8c..c4478bec2f 100644
--- a/src/program_state.cpp
+++ b/src/program_state.cpp
@@ -74,15 +74,11 @@ vector<string> copy_names_of_undefined_symbols(const symbol_section_accessor& se
 }
 
 const std::unordered_map<std::string, std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword>>&
-symbol_addresses(bool rebuild = false) {
+symbol_addresses() {
     static unordered_map<string, pair<Elf64_Addr, Elf_Xword>> r;
     static once_flag f;
 
-    auto cons = [rebuild]() {
-        if (rebuild) {
-            r.clear();
-        }
-
+    call_once(f, []() {
         dl_iterate_phdr(
             [](dl_phdr_info* info, size_t, void*) {
                 static constexpr const char self[] = "/proc/self/exe";
@@ -112,12 +108,7 @@ symbol_addresses(bool rebuild = false) {
                 return 0;
             },
             nullptr);
-    };
-
-    call_once(f, cons);
-    if (rebuild) {
-        cons();
-    }
+    });
 
     return r;
 }
@@ -175,18 +166,13 @@ vector<char> code_object_blob_for_process() {
     return r;
 }
 
-const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs(bool rebuild = false) {
+const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs() {
     static unordered_map<hsa_isa_t, vector<vector<char>>> r;
     static once_flag f;
 
-    auto cons = [rebuild]() {
+    call_once(f, []() {
         static vector<vector<char>> blobs{code_object_blob_for_process()};
 
-        if (rebuild) {
-            blobs.clear();
-            blobs.push_back(code_object_blob_for_process());
-        }
-
         dl_iterate_phdr(
             [](dl_phdr_info* info, std::size_t, void*) {
                 elfio tmp;
@@ -208,13 +194,7 @@ const unordered_map<hsa_isa_t, vector<vector<char>>>& code_object_blobs(bool reb
                 }
             }
         }
-    };
-
-
-    call_once(f, cons);
-    if (rebuild) {
-        cons();
-    }
+    });
 
     return r;
 }
@@ -236,13 +216,13 @@ vector<pair<uintptr_t, string>> function_names_for(const elfio& reader, section*
     return r;
 }
 
-const vector<pair<uintptr_t, string>>& function_names_for_process(bool rebuild = false) {
+const vector<pair<uintptr_t, string>>& function_names_for_process() {
     static constexpr const char self[] = "/proc/self/exe";
 
     static vector<pair<uintptr_t, string>> r;
     static once_flag f;
 
-    auto cons = [rebuild]() {
+    call_once(f, []() {
         elfio reader;
 
         if (!reader.load(self)) {
@@ -253,26 +233,16 @@ const vector<pair<uintptr_t, string>>& function_names_for_process(bool rebuild =
             find_section_if(reader, [](const section* x) { return x->get_type() == SHT_SYMTAB; });
 
         if (symtab) r = function_names_for(reader, symtab);
-    };
-
-    call_once(f, cons);
-    if (rebuild) {
-        cons();
-    }
+    });
 
     return r;
 }
 
-const unordered_map<string, vector<hsa_executable_symbol_t>>& kernels(bool rebuild = false) {
+const unordered_map<string, vector<hsa_executable_symbol_t>>& kernels() {
     static unordered_map<string, vector<hsa_executable_symbol_t>> r;
     static once_flag f;
 
-    auto cons = [rebuild]() {
-        if (rebuild) {
-            r.clear();
-            executables(rebuild);
-        }
-
+    call_once(f, []() {
         static const auto copy_kernels = [](hsa_executable_t, hsa_agent_t,
                                             hsa_executable_symbol_t s, void*) {
             if (type(s) == HSA_SYMBOL_KIND_KERNEL) r[name(s)].push_back(s);
@@ -286,12 +256,7 @@ const unordered_map<string, vector<hsa_executable_symbol_t>>& kernels(bool rebui
                                                      copy_kernels, nullptr);
             }
         }
-    };
-
-    call_once(f, cons);
-    if (rebuild) {
-        cons();
-    }
+    });
 
     return r;
 }
@@ -330,18 +295,13 @@ void load_code_object_and_freeze_executable(
 
 namespace hip_impl {
 const unordered_map<hsa_agent_t, vector<hsa_executable_t>>&
-executables(bool rebuild) {  // TODO: This leaks the hsa_executable_ts, it should use RAII.
+executables() {  // TODO: This leaks the hsa_executable_ts, it should use RAII.
     static unordered_map<hsa_agent_t, vector<hsa_executable_t>> r;
     static once_flag f;
 
-    auto cons = [rebuild]() {
+    call_once(f, []() {
         static const auto accelerators = hc::accelerator::get_all();
 
-        if (rebuild) {
-            r.clear();
-            code_object_blobs(rebuild);
-        }
-
         for (auto&& acc : accelerators) {
             auto agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
 
@@ -375,29 +335,17 @@ executables(bool rebuild) {  // TODO: This leaks the hsa_executable_ts, it shoul
                                    },
                                    agent);
         }
-    };
-
-    call_once(f, cons);
-    if (rebuild) {
-        cons();
-    }
+    });
 
     return r;
 }
 
-const unordered_map<uintptr_t, string>& function_names(bool rebuild) {
+const unordered_map<uintptr_t, string>& function_names() {
     static unordered_map<uintptr_t, string> r{function_names_for_process().cbegin(),
                                               function_names_for_process().cend()};
     static once_flag f;
 
-    auto cons = [rebuild]() {
-        if (rebuild) {
-            r.clear();
-            function_names_for_process(rebuild);
-            r.insert(function_names_for_process().cbegin(),
-                     function_names_for_process().cend());
-        }
-
+    call_once(f, []() {
         dl_iterate_phdr(
             [](dl_phdr_info* info, size_t, void*) {
                 elfio tmp;
@@ -417,30 +365,16 @@ const unordered_map<uintptr_t, string>& function_names(bool rebuild) {
                 return 0;
             },
             nullptr);
-    };
-
-    call_once(f, cons);
-    if (rebuild) {
-        static mutex mtx;
-        lock_guard<mutex> lck{mtx};
-        cons();
-    }
+    });
 
     return r;
 }
 
-const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& functions(bool rebuild) {
+const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& functions() {
     static unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>> r;
     static once_flag f;
 
-    auto cons = [rebuild]() {
-        if (rebuild) {
-            r.clear();
-            function_names(rebuild);
-            kernels(rebuild);
-            globals(rebuild);
-        }
-
+    call_once(f, []() {
         for (auto&& function : function_names()) {
             const auto it = kernels().find(function.second);
 
@@ -452,34 +386,15 @@ const unordered_map<uintptr_t, vector<pair<hsa_agent_t, Kernel_descriptor>>>& fu
                 }
             }
         }
-    };
-
-    call_once(f, cons);
-    if (rebuild) {
-        static mutex mtx;
-        lock_guard<mutex> lck{mtx};
-        cons();
-    }
+    });
 
     return r;
 }
 
-unordered_map<string, void*>& globals(bool rebuild) {
+unordered_map<string, void*>& globals() {
     static unordered_map<string, void*> r;
     static once_flag f;
-    auto cons =[rebuild]() {
-        if (rebuild) {
-            r.clear();
-            symbol_addresses(rebuild);
-        }
-
-        r.reserve(symbol_addresses().size());
-    };
-
-    call_once(f, cons);
-    if (rebuild) {
-        cons();
-    }
+    call_once(f, []() { r.reserve(symbol_addresses().size()); });
 
     return r;
 }

From 3a68ab49193389773c873843e0120d4f00ca7d50 Mon Sep 17 00:00:00 2001
From: sunway513 <peng.sun@amd.com>
Date: Sat, 4 Aug 2018 01:48:06 +0000
Subject: [PATCH 09/16] Add startup loader under HIP_STARTUP_LOADER env var,
 disable by default

---
 src/program_state.cpp | 13 ++++++++++---
 1 file changed, 10 insertions(+), 3 deletions(-)

diff --git a/src/program_state.cpp b/src/program_state.cpp
index fab71b1def..17f127fe87 100644
--- a/src/program_state.cpp
+++ b/src/program_state.cpp
@@ -513,8 +513,9 @@ hsa_executable_t load_executable(const string& file, hsa_executable_t executable
     return executable;
 }
 
-// To force HIP to load the kernels and to setup the function
-// symbol map on program startup
+// HIP startup kernel loader logic
+// When enabled HIP_STARTUP_LOADER, HIP will load the kernels and setup
+// the function symbol map on program startup
 class startup_kernel_loader {
    private:
     startup_kernel_loader() { functions(); }
@@ -522,6 +523,12 @@ class startup_kernel_loader {
     startup_kernel_loader& operator=(const startup_kernel_loader&) = delete;
     static startup_kernel_loader skl;
 };
-startup_kernel_loader startup_kernel_loader::skl;
+
+extern "C" void __attribute__((constructor)) __startup_kernel_loader_init() {
+    if (atoi(std::getenv("HIP_STARTUP_LOADER"))) functions();
+}
+
+extern "C" void __attribute__((destructor)) __startup_kernel_loader_fini() {
+}
 
 }  // Namespace hip_impl.

From 30dfa6f129baa91576ef86d1125d49111097d86f Mon Sep 17 00:00:00 2001
From: sunway513 <peng.sun@amd.com>
Date: Sat, 4 Aug 2018 01:52:27 +0000
Subject: [PATCH 10/16] Add more check to ensure the startup loader only be
 enabled with the env var set to 1

---
 src/program_state.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/program_state.cpp b/src/program_state.cpp
index 17f127fe87..4c14eb2953 100644
--- a/src/program_state.cpp
+++ b/src/program_state.cpp
@@ -525,7 +525,7 @@ class startup_kernel_loader {
 };
 
 extern "C" void __attribute__((constructor)) __startup_kernel_loader_init() {
-    if (atoi(std::getenv("HIP_STARTUP_LOADER"))) functions();
+    if (atoi(std::getenv("HIP_STARTUP_LOADER")) == 1) functions();
 }
 
 extern "C" void __attribute__((destructor)) __startup_kernel_loader_fini() {

From dacb18414ee7f4b1e4c877e8a18b69795a3b11fc Mon Sep 17 00:00:00 2001
From: Yaxun Sam Liu <yaxun.liu@amd.com>
Date: Tue, 7 Aug 2018 15:44:59 -0400
Subject: [PATCH 11/16] Fix declaration conflict when hip/math_functions.h is
 included first

This fixes build failure in TensorFlow 1.8 for HCC
---
 include/hip/hcc_detail/math_functions.h    |   6 +
 tests/src/deviceLib/hipTestIncludeMath.cpp | 136 +++++++++++++++++++++
 2 files changed, 142 insertions(+)
 create mode 100644 tests/src/deviceLib/hipTestIncludeMath.cpp

diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h
index 629ecf7c50..d9125e1654 100644
--- a/include/hip/hcc_detail/math_functions.h
+++ b/include/hip/hcc_detail/math_functions.h
@@ -31,6 +31,12 @@ THE SOFTWARE.
 #include <limits>
 #include <stdint.h>
 
+// HCC's own math functions should be included first, otherwise there will
+// be conflicts when hip/math_functions.h is included before hip/hip_runtime.h.
+#ifdef __HCC__
+#include "kalmar_math.h"
+#endif
+
 #pragma push_macro("__DEVICE__")
 #pragma push_macro("__RETURN_TYPE")
 
diff --git a/tests/src/deviceLib/hipTestIncludeMath.cpp b/tests/src/deviceLib/hipTestIncludeMath.cpp
new file mode 100644
index 0000000000..1097c23166
--- /dev/null
+++ b/tests/src/deviceLib/hipTestIncludeMath.cpp
@@ -0,0 +1,136 @@
+/*
+Copyright (c) 2015-2016 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: %t %s ../test_common.cpp
+ * RUN: %t
+ * HIT_END
+ */
+
+// Test include math_functions.h then hip_runtime.h.
+// Incorrect implementation causes compilation failure due to conflict
+// declartions.
+
+#include <hip/math_functions.h>
+#include <hip/hip_runtime.h>
+#include "test_common.h"
+
+__global__ void FloatMathPrecise() {
+    int iX;
+    float fX, fY;
+
+    acosf(1.0f);
+    acoshf(1.0f);
+    asinf(0.0f);
+    asinhf(0.0f);
+    atan2f(0.0f, 1.0f);
+    atanf(0.0f);
+    atanhf(0.0f);
+    cbrtf(0.0f);
+    fX = ceilf(0.0f);
+    fX = copysignf(1.0f, -2.0f);
+    cosf(0.0f);
+    coshf(0.0f);
+    cospif(0.0f);
+    cyl_bessel_i0f(0.0f);
+    cyl_bessel_i1f(0.0f);
+    erfcf(0.0f);
+    erfcinvf(2.0f);
+    erfcxf(0.0f);
+    erff(0.0f);
+    erfinvf(1.0f);
+    exp10f(0.0f);
+    exp2f(0.0f);
+    expf(0.0f);
+    expm1f(0.0f);
+    fX = fabsf(1.0f);
+    fdimf(1.0f, 0.0f);
+    fdividef(0.0f, 1.0f);
+    fX = floorf(0.0f);
+    fmaf(1.0f, 2.0f, 3.0f);
+    fX = fmaxf(0.0f, 0.0f);
+    fX = fminf(0.0f, 0.0f);
+    fmodf(0.0f, 1.0f);
+    frexpf(0.0f, &iX);
+    hypotf(1.0f, 0.0f);
+    ilogbf(1.0f);
+    isfinite(0.0f);
+    fX = isinf(0.0f);
+    fX = isnan(0.0f);
+    j0f(0.0f);
+    j1f(0.0f);
+    jnf(-1.0f, 1.0f);
+    ldexpf(0.0f, 0);
+    lgammaf(1.0f);
+    llrintf(0.0f);
+    llroundf(0.0f);
+    log10f(1.0f);
+    log1pf(-1.0f);
+    log2f(1.0f);
+    logbf(1.0f);
+    logf(1.0f);
+    lrintf(0.0f);
+    lroundf(0.0f);
+    modff(0.0f, &fX);
+    fX = nanf("1");
+    fX = nearbyintf(0.0f);
+    nextafterf(0.0f, 0.0f);
+    norm3df(1.0f, 0.0f, 0.0f);
+    norm4df(1.0f, 0.0f, 0.0f, 0.0f);
+    normcdff(0.0f);
+    normcdfinvf(1.0f);
+    fX = 1.0f;
+    normf(1, &fX);
+    powf(1.0f, 0.0f);
+    rcbrtf(1.0f);
+    remainderf(2.0f, 1.0f);
+    remquof(1.0f, 2.0f, &iX);
+    rhypotf(0.0f, 1.0f);
+    fY = rintf(1.0f);
+    rnorm3df(0.0f, 0.0f, 1.0f);
+    rnorm4df(0.0f, 0.0f, 0.0f, 1.0f);
+    fX = 1.0f;
+    rnormf(1, &fX);
+    fY = roundf(0.0f);
+    rsqrtf(1.0f);
+    scalblnf(0.0f, 1);
+    scalbnf(0.0f, 1);
+    signbit(1.0f);
+    sincosf(0.0f, &fX, &fY);
+    sincospif(0.0f, &fX, &fY);
+    sinf(0.0f);
+    sinhf(0.0f);
+    sinpif(0.0f);
+    sqrtf(0.0f);
+    tanf(0.0f);
+    tanhf(0.0f);
+    tgammaf(2.0f);
+    fY = truncf(0.0f);
+    y0f(1.0f);
+    y1f(1.0f);
+    ynf(1, 1.0f);
+}
+
+int main() {
+    hipLaunchKernelGGL(FloatMathPrecise, dim3(1, 1, 1), dim3(1, 1, 1), 0, 0);
+    passed();
+}

From 450ea49df0bc01140d3ed1ae6f2bf2018bbbe87c Mon Sep 17 00:00:00 2001
From: Yaxun Sam Liu <yaxun.liu@amd.com>
Date: Tue, 7 Aug 2018 17:14:39 -0400
Subject: [PATCH 12/16] Fix __HIP_DEVICE_COMPILE__ not defined when
 hip/math_functions.h is included

This fixes build failure in TensorFlow 1.8 for HCC
---
 include/hip/math_functions.h               |  1 +
 tests/src/deviceLib/hipTestIncludeMath.cpp | 14 ++++++++++++++
 2 files changed, 15 insertions(+)

diff --git a/include/hip/math_functions.h b/include/hip/math_functions.h
index 23dddf9f9d..2dfec4551b 100644
--- a/include/hip/math_functions.h
+++ b/include/hip/math_functions.h
@@ -27,6 +27,7 @@ THE SOFTWARE.
 // paths to provide a consistent include env and avoid "missing symbol" errors that only appears
 // on NVCC path:
 
+#include <hip/hip_common.h>
 
 #if defined(__HIP_PLATFORM_HCC__) && !defined(__HIP_PLATFORM_NVCC__)
 #include <hip/hcc_detail/math_functions.h>
diff --git a/tests/src/deviceLib/hipTestIncludeMath.cpp b/tests/src/deviceLib/hipTestIncludeMath.cpp
index 1097c23166..fd9bff1ad9 100644
--- a/tests/src/deviceLib/hipTestIncludeMath.cpp
+++ b/tests/src/deviceLib/hipTestIncludeMath.cpp
@@ -31,6 +31,20 @@ THE SOFTWARE.
 // declartions.
 
 #include <hip/math_functions.h>
+
+// Test __HIP_DEVICE_COMPILE__ is defined after math_functions.h
+// is included.
+//
+__device__ __host__ inline void throw_std_bad_alloc()
+{
+  #ifndef __HIP_DEVICE_COMPILE__
+    throw std::bad_alloc();
+  #else
+    std::size_t huge = static_cast<std::size_t>(-1);
+    new int[huge];
+  #endif
+}
+
 #include <hip/hip_runtime.h>
 #include "test_common.h"
 

From b8a0c14171a1fd6744ea4d82e065c62524237fea Mon Sep 17 00:00:00 2001
From: Yaxun Sam Liu <yaxun.liu@amd.com>
Date: Wed, 8 Aug 2018 08:55:28 -0400
Subject: [PATCH 13/16] Fix __HIP_ARCH_* not defined after including
 math_functions.h

hcc_detail/math_functions.h used to include hcc_detail/hip_runtime.h.

Removing it has caused regression in TensorFlow 1.8.

Put it back for backward compatibiliity.
---
 include/hip/hcc_detail/math_functions.h    | 5 +++++
 tests/src/deviceLib/hipTestIncludeMath.cpp | 8 ++++++++
 2 files changed, 13 insertions(+)

diff --git a/include/hip/hcc_detail/math_functions.h b/include/hip/hcc_detail/math_functions.h
index d9125e1654..10967605d4 100644
--- a/include/hip/hcc_detail/math_functions.h
+++ b/include/hip/hcc_detail/math_functions.h
@@ -1338,3 +1338,8 @@ __HIP_OVERLOAD2(double, min)
 #pragma pop_macro("__HIP_OVERLOAD2")
 #pragma pop_macro("__DEVICE__")
 #pragma pop_macro("__RETURN_TYPE")
+
+// For backward compatibility.
+// There are HIP applications e.g. TensorFlow, expecting __HIP_ARCH_* macros
+// defined after including math_functions.h.
+#include <hip/hcc_detail/hip_runtime.h>
diff --git a/tests/src/deviceLib/hipTestIncludeMath.cpp b/tests/src/deviceLib/hipTestIncludeMath.cpp
index fd9bff1ad9..6063eee76c 100644
--- a/tests/src/deviceLib/hipTestIncludeMath.cpp
+++ b/tests/src/deviceLib/hipTestIncludeMath.cpp
@@ -45,6 +45,14 @@ __device__ __host__ inline void throw_std_bad_alloc()
   #endif
 }
 
+// Test __HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ and __HIP_ARCH_HAS_DYNAMIC_PARALLEL__
+// is defined. Eigen HIP/hcc/Half.h __ldg depends on this.
+#if !defined(__HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__) ||                            \
+    !defined(__HIP_ARCH_HAS_DYNAMIC_PARALLEL__)
+#error                                                                         \
+    "__HIP_ARCH_HAS_WARP_FUNNEL_SHIFT__ or __HIP_ARCH_HAS_DYNAMIC_PARALLEL__ not defined"
+#endif
+
 #include <hip/hip_runtime.h>
 #include "test_common.h"
 

From 368977f75bfdfb478dab169a635148e0fbc372b9 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <Evgeniy.Mankov@amd.com>
Date: Wed, 8 Aug 2018 18:34:57 +0300
Subject: [PATCH 14/16] [HIPIFY][Complex] Add cuComplex support

+ Add API_COMPLEX support (data types and functions)
+ Add cuComplex_API_supported_by_HIP.md
+ Add cuComplex_Julia.cu test
+ Update README.md
---
 .../cuComplex_API_supported_by_HIP.md         | 37 ++++++++++++
 hipify-clang/README.md                        |  1 +
 hipify-clang/src/CUDA2HipMap.cpp              | 38 ++++++++++--
 hipify-clang/src/HipifyAction.cpp             |  4 ++
 hipify-clang/src/HipifyAction.h               |  1 +
 hipify-clang/src/Statistics.cpp               |  4 +-
 hipify-clang/src/Statistics.h                 |  2 +
 .../hipify-clang/cuComplex/cuComplex_Julia.cu | 58 +++++++++++++++++++
 8 files changed, 138 insertions(+), 7 deletions(-)
 create mode 100644 docs/markdown/cuComplex_API_supported_by_HIP.md
 create mode 100644 tests/hipify-clang/cuComplex/cuComplex_Julia.cu

diff --git a/docs/markdown/cuComplex_API_supported_by_HIP.md b/docs/markdown/cuComplex_API_supported_by_HIP.md
new file mode 100644
index 0000000000..621973ba61
--- /dev/null
+++ b/docs/markdown/cuComplex_API_supported_by_HIP.md
@@ -0,0 +1,37 @@
+# cuComplex API supported by HIP
+
+## **1. cuComplex Data types**
+
+| **type**     |   **CUDA**                                                    |   **HIP**                                                  |**HIP value** (if differs) |
+|-------------:|---------------------------------------------------------------|------------------------------------------------------------|---------------------------|
+| float2       |***`cuFloatComplex`***                                         |***`hipFloatComplex`***                                     | struct                    |
+| double2      |***`cuDoubleComplex`***                                        |***`hipDoubleComplex`***                                    | struct                    |
+| float2       |***`cuComplex`***                                              |***`hipComplex`***                                          | struct                    |
+
+## **2. cuComplex API functions**
+
+|   **CUDA**                                                |   **HIP**                                       |
+|-----------------------------------------------------------|-------------------------------------------------|
+|`cuCrealf`                                                 |`hipCrealf`                                      |
+|`cuCimagf`                                                 |`hipCimagf`                                      |
+|`make_cuFloatComplex`                                      |`make_hipFloatComplex`                           |
+|`cuConjf`                                                  |`hipConjf`                                       |
+|`cuCaddf`                                                  |`hipCaddf`                                       |
+|`cuCsubf`                                                  |`hipCsubf`                                       |
+|`cuCmulf`                                                  |`hipCmulf`                                       |
+|`cuCdivf`                                                  |`hipCdivf`                                       |
+|`cuCabsf`                                                  |`hipCabsf`                                       |
+|`cuCreal`                                                  |`hipCreal`                                       |
+|`cuCimag`                                                  |`hipCimag`                                       |
+|`make_cuDoubleComplex`                                     |`make_hipDoubleComplex`                          |
+|`cuConj`                                                   |`hipConj`                                        |
+|`cuCadd`                                                   |`hipCadd`                                        |
+|`cuCsub`                                                   |`hipCsub`                                        |
+|`cuCmul`                                                   |`hipCmul`                                        |
+|`cuCdiv`                                                   |`hipCdiv`                                        |
+|`cuCabs`                                                   |`hipCabs`                                        |
+|`make_cuComplex`                                           |`make_hipComplex`                                |
+|`cuComplexFloatToDouble`                                   |`hipComplexFloatToDouble`                        |
+|`cuComplexDoubleToFloat`                                   |`hipComplexDoubleToFloat`                        |
+|`cuCfmaf`                                                  |`hipCfmaf`                                       |
+|`cuCfma`                                                   |`hipCfma`                                        |
diff --git a/hipify-clang/README.md b/hipify-clang/README.md
index d5a825033d..7f55dea800 100644
--- a/hipify-clang/README.md
+++ b/hipify-clang/README.md
@@ -21,6 +21,7 @@
 
 - [Runtime API](../docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md)
 - [Driver API](../docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md)
+- [cuComplex API](../docs/markdown/cuComplex_API_supported_by_HIP.md)
 - [cuBLAS](../docs/markdown/CUBLAS_API_supported_by_HIP.md)
 - [cuRAND](../docs/markdown/CURAND_API_supported_by_HIP.md)
 - [cuDNN](../docs/markdown/CUDNN_API_supported_by_HIP.md)
diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp
index ca53738abb..4bd7263bad 100644
--- a/hipify-clang/src/CUDA2HipMap.cpp
+++ b/hipify-clang/src/CUDA2HipMap.cpp
@@ -366,7 +366,6 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
     {"channel_descriptor.h",      {"hip/channel_descriptor.h", CONV_INCLUDE,             API_RUNTIME}},
     {"device_functions.h",        {"hip/device_functions.h",   CONV_INCLUDE,             API_RUNTIME}},
     {"driver_types.h",            {"hip/driver_types.h",       CONV_INCLUDE,             API_RUNTIME}},
-    {"cuComplex.h",               {"hip/hip_complex.h",        CONV_INCLUDE,             API_RUNTIME}},
     {"cuda_fp16.h",               {"hip/hip_fp16.h",           CONV_INCLUDE,             API_RUNTIME}},
     {"cuda_texture_types.h",      {"hip/hip_texture_types.h",  CONV_INCLUDE,             API_RUNTIME}},
     {"vector_types.h",            {"hip/hip_vector_types.h",   CONV_INCLUDE,             API_RUNTIME}},
@@ -397,12 +396,11 @@ const std::map <llvm::StringRef, hipCounter> CUDA_INCLUDE_MAP{
     // CUDNN includes
     {"cudnn.h",                   {"hipDNN.h",                 CONV_INCLUDE_CUDA_MAIN_H, API_DNN}},
 
-    // CUDNN includes
+    // CUFFT includes
     {"cufft.h",                   {"hipfft.h",                 CONV_INCLUDE_CUDA_MAIN_H, API_FFT}},
 
-    // HIP includes
-    // TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h
-    //    {"cudacommon.h", {"hipcommon.h", CONV_INCLUDE, API_RUNTIME}},
+    // cuComplex includes
+    {"cuComplex.h",               {"hip/hip_complex.h",        CONV_INCLUDE_CUDA_MAIN_H, API_COMPLEX}},
 };
 
 /// All other identifiers: function and macro names.
@@ -1387,6 +1385,36 @@ const std::map<llvm::StringRef, hipCounter> CUDA_IDENTIFIER_MAP{
     {"cuGraphicsEGLRegisterImage",          {"hipGraphicsEGLRegisterImage",          CONV_EGL, API_DRIVER, HIP_UNSUPPORTED}},    // API_Runtime ANALOGUE (cudaGraphicsEGLRegisterImage)
     {"cuGraphicsResourceGetMappedEglFrame", {"hipGraphicsResourceGetMappedEglFrame", CONV_EGL, API_DRIVER, HIP_UNSUPPORTED}},    // API_Runtime ANALOGUE (cudaGraphicsResourceGetMappedEglFrame)
 
+
+////////////////////////////// cuComplex API //////////////////////////////
+    {"cuFloatComplex",                               {"hipFloatComplex",                                     CONV_TYPE, API_COMPLEX}},
+    {"cuDoubleComplex",                              {"hipDoubleComplex",                                    CONV_TYPE, API_COMPLEX}},
+    {"cuComplex",                                    {"hipComplex",                                          CONV_TYPE, API_COMPLEX}},
+
+    {"cuCrealf",                                     {"hipCrealf",                                           CONV_COMPLEX, API_COMPLEX}},
+    {"cuCimagf",                                     {"hipCimagf",                                           CONV_COMPLEX, API_COMPLEX}},
+    {"make_cuFloatComplex",                          {"make_hipFloatComplex",                                CONV_COMPLEX, API_COMPLEX}},
+    {"cuConjf",                                      {"hipConjf",                                            CONV_COMPLEX, API_COMPLEX}},
+    {"cuCaddf",                                      {"hipCaddf",                                            CONV_COMPLEX, API_COMPLEX}},
+    {"cuCsubf",                                      {"hipCsubf",                                            CONV_COMPLEX, API_COMPLEX}},
+    {"cuCmulf",                                      {"hipCmulf",                                            CONV_COMPLEX, API_COMPLEX}},
+    {"cuCdivf",                                      {"hipCdivf",                                            CONV_COMPLEX, API_COMPLEX}},
+    {"cuCabsf",                                      {"hipCabsf",                                            CONV_COMPLEX, API_COMPLEX}},
+    {"cuCreal",                                      {"hipCreal",                                            CONV_COMPLEX, API_COMPLEX}},
+    {"cuCimag",                                      {"hipCimag",                                            CONV_COMPLEX, API_COMPLEX}},
+    {"make_cuDoubleComplex",                         {"make_hipDoubleComplex",                               CONV_COMPLEX, API_COMPLEX}},
+    {"cuConj",                                       {"hipConj",                                             CONV_COMPLEX, API_COMPLEX}},
+    {"cuCadd",                                       {"hipCadd",                                             CONV_COMPLEX, API_COMPLEX}},
+    {"cuCsub",                                       {"hipCsub",                                             CONV_COMPLEX, API_COMPLEX}},
+    {"cuCmul",                                       {"hipCmul",                                             CONV_COMPLEX, API_COMPLEX}},
+    {"cuCdiv",                                       {"hipCdiv",                                             CONV_COMPLEX, API_COMPLEX}},
+    {"cuCabs",                                       {"hipCabs",                                             CONV_COMPLEX, API_COMPLEX}},
+    {"make_cuComplex",                               {"make_hipComplex",                                     CONV_COMPLEX, API_COMPLEX}},
+    {"cuComplexFloatToDouble",                       {"hipComplexFloatToDouble",                             CONV_COMPLEX, API_COMPLEX}},
+    {"cuComplexDoubleToFloat",                       {"hipComplexDoubleToFloat",                             CONV_COMPLEX, API_COMPLEX}},
+    {"cuCfmaf",                                      {"hipCfmaf",                                            CONV_COMPLEX, API_COMPLEX}},
+    {"cuCfma",                                       {"hipCfma",                                             CONV_COMPLEX, API_COMPLEX}},
+
 /////////////////////////////// CUDA RT API ///////////////////////////////
     // Data types
     {"cudaDataType_t", {"hipDataType_t", CONV_TYPE, API_RUNTIME, HIP_UNSUPPORTED}},
diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp
index 20bc9dcc2c..a9bd1aa085 100644
--- a/hipify-clang/src/HipifyAction.cpp
+++ b/hipify-clang/src/HipifyAction.cpp
@@ -168,6 +168,10 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) {
                   if (insertedFFTHeader) { return true; }
                   insertedFFTHeader = true;
                   return false;
+                case API_COMPLEX:
+                  if (insertedComplexHeader) { return true; }
+                  insertedComplexHeader = true;
+                  return false;
                 default:
                     return false;
             }
diff --git a/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h
index ad987c921e..1262142cfc 100644
--- a/hipify-clang/src/HipifyAction.h
+++ b/hipify-clang/src/HipifyAction.h
@@ -29,6 +29,7 @@ private:
     bool insertedRAND_kernelHeader = false;
     bool insertedDNNHeader = false;
     bool insertedFFTHeader = false;
+    bool insertedComplexHeader = false;
     bool firstHeader = false;
     bool pragmaOnce = false;
     clang::SourceLocation firstHeaderLoc;
diff --git a/hipify-clang/src/Statistics.cpp b/hipify-clang/src/Statistics.cpp
index 595f243857..2115c567a5 100644
--- a/hipify-clang/src/Statistics.cpp
+++ b/hipify-clang/src/Statistics.cpp
@@ -8,13 +8,13 @@ const char *counterNames[NUM_CONV_TYPES] = {
     "version", "init", "device", "mem", "kern", "coord_func", "math_func", "device_func",
     "special_func", "stream", "event", "occupancy", "ctx", "peer", "module",
     "cache", "exec", "err", "def", "tex", "gl", "graphics",
-    "surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl",
+    "surface", "jit", "d3d9", "d3d10", "d3d11", "vdpau", "egl", "complex",
     "thread", "other", "include", "include_cuda_main_header", "type", "literal",
     "numeric_literal"
 };
 
 const char *apiNames[NUM_API_TYPES] = {
-    "CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API", "CUDNN API", "CUFFT API"
+    "CUDA Driver API", "CUDA RT API", "CUBLAS API", "CURAND API", "CUDNN API", "CUFFT API", "cuComplex API"
 };
 
 namespace {
diff --git a/hipify-clang/src/Statistics.h b/hipify-clang/src/Statistics.h
index 00ecce3eda..1ded45f0e4 100644
--- a/hipify-clang/src/Statistics.h
+++ b/hipify-clang/src/Statistics.h
@@ -40,6 +40,7 @@ enum ConvTypes {
     CONV_D3D11,
     CONV_VDPAU,
     CONV_EGL,
+    CONV_COMPLEX,
     CONV_THREAD,
     CONV_OTHER,
     CONV_INCLUDE,
@@ -58,6 +59,7 @@ enum ApiTypes {
     API_RAND,
     API_DNN,
     API_FFT,
+    API_COMPLEX,
     API_LAST
 };
 constexpr int NUM_API_TYPES = (int) ApiTypes::API_LAST;
diff --git a/tests/hipify-clang/cuComplex/cuComplex_Julia.cu b/tests/hipify-clang/cuComplex/cuComplex_Julia.cu
new file mode 100644
index 0000000000..8bf9587b94
--- /dev/null
+++ b/tests/hipify-clang/cuComplex/cuComplex_Julia.cu
@@ -0,0 +1,58 @@
+// RUN: %run_test hipify "%s" "%t" %cuda_args
+
+// CHECK: #include <hip/hip_runtime.h>
+// CHECK: #include "hip/hip_complex.h"
+#include "cuComplex.h"
+
+#define TYPEFLOAT
+#define DIMX 100
+#define DIMY 40
+#define moveX 2
+#define moveY 1
+
+#define MAXITERATIONS 10
+
+#ifdef TYPEFLOAT
+#define TYPE float
+// CHECK: #define cTYPE hipFloatComplex
+#define cTYPE cuFloatComplex
+// CHECK: #define cMakecuComplex(re,i) make_hipFloatComplex(re,i)
+#define cMakecuComplex(re,i) make_cuFloatComplex(re,i)
+#endif
+#ifdef TYPEDOUBLE
+// CHECK: #define TYPE hipDoubleComplex
+#define TYPE cuDoubleComplex
+// CHECK: #define cMakecuComplex(re,i) make_hipDoubleComplex(re,i)
+#define cMakecuComplex(re,i) make_cuDoubleComplex(re,i)
+#endif
+
+__device__ cTYPE juliaFunctor(cTYPE p, cTYPE c) {
+  // CHECK: return hipCaddf(hipCmulf(p, p), c);
+  return cuCaddf(cuCmulf(p, p), c);
+}
+
+__device__ cTYPE convertToComplex(int x, int y, float zoom) {
+  TYPE jx = 1.5 * (x - DIMX / 2) / (0.5 * zoom * DIMX) + moveX;
+  TYPE jy = (y - DIMY / 2) / (0.5 * zoom * DIMY) + moveY;
+  return cMakecuComplex(jx, jy);
+}
+
+__device__ int evolveComplexPoint(cTYPE p, cTYPE c) {
+  int it = 1;
+  // CHECK: while (it <= MAXITERATIONS && hipCabsf(p) <= 4) {
+  while (it <= MAXITERATIONS && cuCabsf(p) <= 4) {
+    p = juliaFunctor(p, c);
+    it++;
+  }
+  return it;
+}
+
+__global__ void computeJulia(int* data, cTYPE c, float zoom) {
+  int i = blockIdx.x * blockDim.x + threadIdx.x;
+  int j = blockIdx.y * blockDim.y + threadIdx.y;
+
+  if (i<DIMX && j<DIMY) {
+    cTYPE p = convertToComplex(i, j, zoom);
+    data[i*DIMY + j] = evolveComplexPoint(p, c);
+  }
+}

From 8f0d6849884411f53581e7082257331f42426433 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Wed, 8 Aug 2018 22:28:13 +0530
Subject: [PATCH 15/16] Clean up module api samples

---
 samples/0_Intro/module_api/runKernel.cpp      | 25 +++--------------
 .../0_Intro/module_api_global/runKernel.cpp   | 25 +++--------------
 .../11_texture_driver/texture2dDrv.cpp        | 27 +++----------------
 3 files changed, 10 insertions(+), 67 deletions(-)

diff --git a/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp
index 129c458dd0..ed74a7e4eb 100644
--- a/samples/0_Intro/module_api/runKernel.cpp
+++ b/samples/0_Intro/module_api/runKernel.cpp
@@ -30,7 +30,7 @@ THE SOFTWARE.
 #define LEN 64
 #define SIZE LEN << 2
 
-#define fileName "vcpy_kernel.code.adipose"
+#define fileName "vcpy_kernel.code"
 #define kernel_name "hello_world"
 
 #define HIP_CHECK(status)                                                                          \
@@ -66,32 +66,13 @@ int main() {
     HIP_CHECK(hipModuleLoad(&Module, fileName));
     HIP_CHECK(hipModuleGetFunction(&Function, Module, kernel_name));
 
-#ifdef __HIP_PLATFORM_HCC__
-    uint32_t len = LEN;
-    uint32_t one = 1;
-
     struct {
         void* _Ad;
         void* _Bd;
     } args;
 
-    args._Ad = Ad;
-    args._Bd = Bd;
-
-#endif
-
-#ifdef __HIP_PLATFORM_NVCC__
-    struct {
-        uint32_t _hidden[1];
-        void* _Ad;
-        void* _Bd;
-    } args;
-
-    args._hidden[0] = 0;
-    args._Ad = Ad;
-    args._Bd = Bd;
-#endif
-
+    args._Ad = (void*) Ad;
+    args._Bd = (void*) Bd;
 
     size_t size = sizeof(args);
 
diff --git a/samples/0_Intro/module_api_global/runKernel.cpp b/samples/0_Intro/module_api_global/runKernel.cpp
index ebae029a16..0047f79b7f 100644
--- a/samples/0_Intro/module_api_global/runKernel.cpp
+++ b/samples/0_Intro/module_api_global/runKernel.cpp
@@ -30,7 +30,7 @@ THE SOFTWARE.
 #define LEN 64
 #define SIZE LEN * sizeof(float)
 
-#define fileName "vcpy_kernel.code.adipose"
+#define fileName "vcpy_kernel.code"
 float myDeviceGlobal;
 float myDeviceGlobalArray[16];
 #define HIP_CHECK(cmd)                                                                             \
@@ -79,32 +79,13 @@ int main() {
         myDeviceGlobalArray[i] = i * 1000.0f;
     }
 
-#ifdef __HIP_PLATFORM_HCC__
-    uint32_t len = LEN;
-    uint32_t one = 1;
-
     struct {
         void* _Ad;
         void* _Bd;
     } args;
 
-    args._Ad = Ad;
-    args._Bd = Bd;
-
-#endif
-
-#ifdef __HIP_PLATFORM_NVCC__
-    struct {
-        uint32_t _hidden[1];
-        void* _Ad;
-        void* _Bd;
-    } args;
-
-    args._hidden[0] = 0;
-    args._Ad = Ad;
-    args._Bd = Bd;
-#endif
-
+    args._Ad = (void*) Ad;
+    args._Bd = (void*) Bd;
 
     size_t size = sizeof(args);
 
diff --git a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
index 2485d455ef..385ed4775a 100644
--- a/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
+++ b/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp
@@ -21,13 +21,13 @@ THE SOFTWARE.
 */
 
 #include "hip/hip_runtime.h"
-#include "hip/hip_runtime_api.h"
+//#include "hip/hip_runtime_api.h"
 #include <iostream>
 #include <fstream>
 #include <vector>
-#include <hip/hip_hcc.h>
+//#include <hip/hip_hcc.h>
 
-#define fileName "tex2dKernel.code.adipose"
+#define fileName "tex2dKernel.code"
 
 texture<float, 2, hipReadModeElementType> tex;
 bool testResult = false;
@@ -87,34 +87,15 @@ bool runTest(int argc, char** argv) {
     float* dData = NULL;
     hipMalloc((void**)&dData, size);
 
-#ifdef __HIP_PLATFORM_HCC__
-
     struct {
         void* _Ad;
         unsigned int _Bd;
         unsigned int _Cd;
     } args;
-    args._Ad = dData;
+    args._Ad = (void*) dData;
     args._Bd = width;
     args._Cd = height;
 
-#endif
-
-#ifdef __HIP_PLATFORM_NVCC__
-    struct {
-        uint32_t _hidden[1];
-        void* _Ad;
-        unsigned int _Bd;
-        unsigned int _Cd;
-    } args;
-
-    args._hidden[0] = 0;
-    args._Ad = dData;
-    args._Bd = width;
-    args._Cd = height;
-#endif
-
-
     size_t sizeTemp = sizeof(args);
 
     void* config[] = {HIP_LAUNCH_PARAM_BUFFER_POINTER, &args, HIP_LAUNCH_PARAM_BUFFER_SIZE,

From 17f38937e07f3839368735453de28edf4514cfe0 Mon Sep 17 00:00:00 2001
From: sunway513 <peng.sun@amd.com>
Date: Thu, 9 Aug 2018 16:40:26 +0000
Subject: [PATCH 16/16] resolve a segfault bug when env var not set; remove
 startup_kernel_loader class

---
 src/program_state.cpp | 13 ++++---------
 1 file changed, 4 insertions(+), 9 deletions(-)

diff --git a/src/program_state.cpp b/src/program_state.cpp
index 4c14eb2953..8766134582 100644
--- a/src/program_state.cpp
+++ b/src/program_state.cpp
@@ -516,16 +516,11 @@ hsa_executable_t load_executable(const string& file, hsa_executable_t executable
 // HIP startup kernel loader logic
 // When enabled HIP_STARTUP_LOADER, HIP will load the kernels and setup
 // the function symbol map on program startup
-class startup_kernel_loader {
-   private:
-    startup_kernel_loader() { functions(); }
-    startup_kernel_loader(const startup_kernel_loader&) = delete;
-    startup_kernel_loader& operator=(const startup_kernel_loader&) = delete;
-    static startup_kernel_loader skl;
-};
-
 extern "C" void __attribute__((constructor)) __startup_kernel_loader_init() {
-    if (atoi(std::getenv("HIP_STARTUP_LOADER")) == 1) functions();
+    int hip_startup_loader=0;
+    if (std::getenv("HIP_STARTUP_LOADER"))
+        hip_startup_loader = atoi(std::getenv("HIP_STARTUP_LOADER"));
+    if (hip_startup_loader) functions();
 }
 
 extern "C" void __attribute__((destructor)) __startup_kernel_loader_fini() {