From 0cae3e06c1f2b1ce1a3508066daed1f8773fe872 Mon Sep 17 00:00:00 2001
From: Siu Chi Chan <siuchi.chan@amd.com>
Date: Sat, 4 May 2019 12:57:36 -0400
Subject: [PATCH 1/6] moving agent_globals_impl into hip_module

[ROCm/clr commit: 6852be819f95e50961da29019c2eaa902b1f4a79]
---
 .../include/hip/hcc_detail/hip_runtime_api.h  | 147 +++-------------
 projects/clr/hipamd/src/hip_module.cpp        | 165 ++++++++++++++++--
 2 files changed, 177 insertions(+), 135 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
index d870963101..17a42e2a16 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
@@ -2590,6 +2590,7 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con
 
 hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func);
 
+#if 0
 struct Agent_global {
 
     Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {}
@@ -2629,6 +2630,7 @@ struct Agent_global {
     hipDeviceptr_t address;
     uint32_t byte_cnt;
 };
+#endif
 
 #if !__HIP_VDI__
 #if defined(__cplusplus)
@@ -2636,134 +2638,35 @@ struct Agent_global {
 #endif
 
 namespace hip_impl {
-hsa_executable_t executable_for(hipModule_t);
-const char* hash_for(hipModule_t);
+    class agent_globals_impl;
+    class agent_globals {
+        public:
+            agent_globals();
+            ~agent_globals();
+            agent_globals(const agent_globals&) = delete;
 
-template<typename ForwardIterator>
-std::pair<hipDeviceptr_t, std::size_t> read_global_description(
-    ForwardIterator f, ForwardIterator l, const char* name) {
-    const auto it = std::find_if(f, l, [=](const Agent_global& x) {
-        return strcmp(x.name, name) == 0;
-    });
+            hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
+                    hipModule_t hmod, const char* name);
+            hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
+                    const char* name);
+        private:
+            agent_globals_impl* impl;
+    };
 
-    return it == l ?
-        std::make_pair(nullptr, 0u) : std::make_pair(it->address, it->byte_cnt);
-}
-
-std::vector<Agent_global> read_agent_globals(hsa_agent_t agent,
-                                             hsa_executable_t executable);
-hsa_agent_t this_agent();
-
-
-class agent_globals_impl {
-private:
-    std::pair<
-        std::mutex,
-        std::unordered_map<
-            std::string, std::vector<Agent_global>>> globals_from_module;
-
-    std::unordered_map<
-        hsa_agent_t,
-        std::pair<
-            std::once_flag,
-            std::vector<Agent_global>>> globals_from_process;
-
-public:
-
-    hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
-            hipModule_t hmod, const char* name) {
-        // the key of the map would the hash of code object associated with the
-        // hipModule_t instance
-        std::string key(hash_for(hmod));
-
-        if (globals_from_module.second.count(key) == 0) {
-            std::lock_guard<std::mutex> lck{globals_from_module.first};
-
-            if (globals_from_module.second.count(key) == 0) {
-                globals_from_module.second.emplace(
-                        key, read_agent_globals(this_agent(), executable_for(hmod)));
-            }
-        }
-
-        const auto it0 = globals_from_module.second.find(key);
-        if (it0 == globals_from_module.second.cend()) {
-            hip_throw(
-                    std::runtime_error{"agent_globals data structure corrupted."});
-        }
-
-        std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(),
-                it0->second.cend(), name);
-
-        return *dptr ? hipSuccess : hipErrorNotFound;
+    inline
+    __attribute__((visibility("hidden")))
+    agent_globals& get_agent_globals() {
+        static agent_globals ag;
+        return ag;
     }
 
+    extern "C"
+    inline
+    __attribute__((visibility("hidden")))
     hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
-            const char* name) {
-
-        auto agent = this_agent();
-
-        std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) {
-            std::vector<Agent_global> tmp0;
-            for (auto&& executable : hip_impl::get_program_state().executables(aa)) {
-                auto tmp1 = read_agent_globals(aa, executable);
-                tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()),
-                            make_move_iterator(tmp1.end()));
-            }
-            globals_from_process[aa].second = move(move(tmp0));
-        }, agent);
-
-        const auto it = globals_from_process.find(agent);
-
-        if (it == globals_from_process.cend()) return hipErrorNotInitialized;
-
-        std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(),
-                it->second.second.cend(), name);
-
-        return *dptr ? hipSuccess : hipErrorNotFound;
+        const char* name) {
+        return get_agent_globals().read_agent_global_from_process(dptr, bytes, name);
     }
-  
-};
-
-class agent_globals {
-public:
-    agent_globals() : impl(new agent_globals_impl()) { 
-        if (!impl) 
-            hip_throw(
-                std::runtime_error{"Error when constructing agent global data structures."});
-    }
-    ~agent_globals() { delete impl; }
-
-    hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
-                                             hipModule_t hmod, const char* name) {
-        return impl->read_agent_global_from_module(dptr, bytes, hmod, name);
-    }
-
-    hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
-                                              const char* name) {
-        return impl->read_agent_global_from_process(dptr, bytes, name);
-    }
-
-private:
-    agent_globals_impl* impl;
-};
-
-inline
-__attribute__((visibility("hidden")))
-agent_globals& get_agent_globals() {
-    static agent_globals ag;
-    return ag;
-}
-
-
-extern "C"
-inline
-__attribute__((visibility("hidden")))
-hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
-                                          const char* name) {
-    return get_agent_globals().read_agent_global_from_process(dptr, bytes, name);
-}
-
-
 } // Namespace hip_impl.
 
 #if defined(__cplusplus)
diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp
index e8a8801e98..2b41988ac5 100644
--- a/projects/clr/hipamd/src/hip_module.cpp
+++ b/projects/clr/hipamd/src/hip_module.cpp
@@ -289,16 +289,6 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
         localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0));
 }
 
-hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
-                              hipModule_t hmod, const char* name) {
-    HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name);
-    if (!dptr || !bytes || !hmod) return hipErrorInvalidValue;
-
-    if (!name) return hipErrorNotInitialized;
-
-    return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name);
-}
-
 namespace hip_impl {
     hsa_executable_t executable_for(hipModule_t hmod) {
         return hmod->executable;
@@ -323,10 +313,159 @@ namespace hip_impl {
 
         return currentDevice->_hsaAgent;
     }
+
+    struct Agent_global {
+        Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {}
+        Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt) 
+            : name(nullptr), address(address), byte_cnt(byte_cnt) {
+                if (name)
+                    this->name = strdup(name);
+        }
+
+        Agent_global& operator=(Agent_global&& t) {
+            if (this == &t) return *this;
+
+            if (name) free(name);
+            name = t.name;
+            address = t.address;
+            byte_cnt = t.byte_cnt;
+
+            t.name = nullptr;
+            t.address = nullptr;
+            t.byte_cnt = 0;
+
+            return *this;
+        }
+
+        Agent_global(Agent_global&& t) 
+            : name(nullptr), address(nullptr), byte_cnt(0) {
+                *this = std::move(t);
+        }
+
+        // not needed, delete them to prevent bugs
+        Agent_global(const Agent_global&) = delete;
+        Agent_global& operator=(Agent_global& t) = delete;
+
+        ~Agent_global() { if (name) free(name); }
+
+        char* name;
+        hipDeviceptr_t address;
+        uint32_t byte_cnt;
+    };
+
+    template<typename ForwardIterator>
+    std::pair<hipDeviceptr_t, std::size_t> read_global_description(
+        ForwardIterator f, ForwardIterator l, const char* name) {
+        const auto it = std::find_if(f, l, [=](const Agent_global& x) {
+            return strcmp(x.name, name) == 0;
+        });
+
+        return it == l ?
+            std::make_pair(nullptr, 0u) : std::make_pair(it->address, it->byte_cnt);
+    }
+
+    std::vector<Agent_global> read_agent_globals(hsa_agent_t agent,
+                                            hsa_executable_t executable);
+    class agent_globals_impl {
+        private:
+            std::pair<
+                std::mutex,
+                std::unordered_map<
+                    std::string, std::vector<Agent_global>>> globals_from_module;
+
+            std::unordered_map<
+                hsa_agent_t,
+                std::pair<
+                    std::once_flag,
+                    std::vector<Agent_global>>> globals_from_process;
+
+        public:
+
+            hipError_t read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
+                    hipModule_t hmod, const char* name) {
+                // the key of the map would the hash of code object associated with the
+                // hipModule_t instance
+                std::string key(hash_for(hmod));
+
+                if (globals_from_module.second.count(key) == 0) {
+                    std::lock_guard<std::mutex> lck{globals_from_module.first};
+
+                    if (globals_from_module.second.count(key) == 0) {
+                        globals_from_module.second.emplace(
+                                key, read_agent_globals(this_agent(), executable_for(hmod)));
+                    }
+                }
+
+                const auto it0 = globals_from_module.second.find(key);
+                if (it0 == globals_from_module.second.cend()) {
+                    hip_throw(
+                            std::runtime_error{"agent_globals data structure corrupted."});
+                }
+
+                std::tie(*dptr, *bytes) = read_global_description(it0->second.cbegin(),
+                        it0->second.cend(), name);
+                
+                return *dptr ? hipSuccess : hipErrorNotFound;
+            }
+
+            hipError_t read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
+                    const char* name) {
+
+                auto agent = this_agent();
+
+                std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) {
+                    std::vector<Agent_global> tmp0;
+                    for (auto&& executable : hip_impl::get_program_state().executables(aa)) {
+                        auto tmp1 = read_agent_globals(aa, executable);
+                        tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()),
+                            make_move_iterator(tmp1.end()));
+                    }
+                    globals_from_process[aa].second = move(move(tmp0));
+                }, agent);
+
+                const auto it = globals_from_process.find(agent);
+
+                if (it == globals_from_process.cend()) return hipErrorNotInitialized;
+
+                std::tie(*dptr, *bytes) = read_global_description(it->second.second.cbegin(),
+                        it->second.second.cend(), name);
+
+                return *dptr ? hipSuccess : hipErrorNotFound;
+            }
+
+    };
+
+    agent_globals::agent_globals() : impl(new agent_globals_impl()) { 
+        if (!impl) 
+            hip_throw(
+                std::runtime_error{"Error when constructing agent global data structures."});
+    }
+    agent_globals::~agent_globals() { delete impl; }
+
+    hipError_t agent_globals::read_agent_global_from_module(hipDeviceptr_t* dptr, size_t* bytes,
+                                             hipModule_t hmod, const char* name) {
+        return impl->read_agent_global_from_module(dptr, bytes, hmod, name);
+    }
+
+    hipError_t agent_globals::read_agent_global_from_process(hipDeviceptr_t* dptr, size_t* bytes,
+                                              const char* name) {
+        return impl->read_agent_global_from_process(dptr, bytes, name);
+    }
+
 } // Namespace hip_impl.
 
+hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes,
+                              hipModule_t hmod, const char* name) {
+    HIP_INIT_API(hipModuleGetGlobal, dptr, bytes, hmod, name);
+    if (!dptr || !bytes || !hmod) return hipErrorInvalidValue;
+
+    if (!name) return hipErrorNotInitialized;
+
+    return hip_impl::get_agent_globals().read_agent_global_from_module(dptr, bytes, hmod, name);
+}
+
 namespace {
-inline void track(const Agent_global& x, hsa_agent_t agent) {
+inline void track(const hip_impl::Agent_global& x, hsa_agent_t agent) {
     tprintf(DB_MEM, "  add variable '%s' with ptr=%p size=%u to tracker\n", x.name,
             x.address, x.byte_cnt);
 
@@ -347,7 +486,7 @@ inline void track(const Agent_global& x, hsa_agent_t agent) {
 
 }
 
-template <typename Container = vector<Agent_global>>
+template <typename Container = vector<hip_impl::Agent_global>>
 inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t agent,
                                                 hsa_executable_symbol_t x, void* out) {
     using namespace hip_impl;
@@ -358,7 +497,7 @@ inline hsa_status_t copy_agent_global_variables(hsa_executable_t, hsa_agent_t ag
     hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &t);
 
     if (t == HSA_SYMBOL_KIND_VARIABLE) {
-        Agent_global tmp(name(x).c_str(), address(x), size(x));
+        hip_impl::Agent_global tmp(name(x).c_str(), address(x), size(x));
         static_cast<Container*>(out)->push_back(std::move(tmp));
 
         track(static_cast<Container*>(out)->back(),agent);

From 305eb4239ec18e29714ccd4efc46b824e31142fd Mon Sep 17 00:00:00 2001
From: Siu Chi Chan <siuchi.chan@amd.com>
Date: Sat, 4 May 2019 14:35:27 -0400
Subject: [PATCH 2/6] remove executables() from program_state

[ROCm/clr commit: e2c01228927531008158387c98ec79b08fd948d1]
---
 .../hipamd/include/hip/hcc_detail/program_state.hpp    | 10 ++--------
 projects/clr/hipamd/src/hip_module.cpp                 |  2 +-
 projects/clr/hipamd/src/program_state.cpp              |  4 ----
 3 files changed, 3 insertions(+), 13 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp
index 19db8f9e0d..5741c22b4b 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp
+++ b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp
@@ -31,9 +31,6 @@ THE SOFTWARE.
 #include <cstdint>
 #include <cstdlib>
 #include <stdexcept>
-#include <string>
-#include <unordered_map>
-#include <vector>
 
 struct ihipModuleSymbol_t;
 using hipFunction_t = ihipModuleSymbol_t*;
@@ -72,6 +69,7 @@ class program_state {
 public:
     program_state();
     ~program_state();
+    program_state(const program_state&) = delete;
 
     hipFunction_t kernel_descriptor(std::uintptr_t,
                                     hsa_agent_t);
@@ -83,12 +81,8 @@ public:
 
     void* global_addr_by_name(const char* name);
 
-    // to fix later
-    const std::vector<hsa_executable_t>& executables(hsa_agent_t agent);
-
-    program_state(const program_state&) = delete;
-
 private:
+    friend class agent_globals_impl;
     program_state_impl* impl;
 };
 
diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp
index 2b41988ac5..4ff6a987fd 100644
--- a/projects/clr/hipamd/src/hip_module.cpp
+++ b/projects/clr/hipamd/src/hip_module.cpp
@@ -415,7 +415,7 @@ namespace hip_impl {
 
                 std::call_once(globals_from_process[agent].first, [this](hsa_agent_t aa) {
                     std::vector<Agent_global> tmp0;
-                    for (auto&& executable : hip_impl::get_program_state().executables(aa)) {
+                    for (auto&& executable : hip_impl::get_program_state().impl->get_executables(aa)) {
                         auto tmp1 = read_agent_globals(aa, executable);
                         tmp0.insert(tmp0.end(), make_move_iterator(tmp1.begin()),
                             make_move_iterator(tmp1.end()));
diff --git a/projects/clr/hipamd/src/program_state.cpp b/projects/clr/hipamd/src/program_state.cpp
index 6783f85b9b..e226e1af66 100644
--- a/projects/clr/hipamd/src/program_state.cpp
+++ b/projects/clr/hipamd/src/program_state.cpp
@@ -45,10 +45,6 @@ hsa_executable_t program_state::load_executable(const char* data,
     return impl->load_executable(data, data_size, executable, agent);
 }
 
-const std::vector<hsa_executable_t>& program_state::executables(hsa_agent_t agent) {
-    return impl->get_executables(agent);
-}
-
 hipFunction_t program_state::kernel_descriptor(std::uintptr_t function_address,
                                                           hsa_agent_t agent) {
     auto& kd = impl->kernel_descriptor(function_address, agent);

From 304a1e2dbefadea5adcdc8fd2b9fddb9e187e8fb Mon Sep 17 00:00:00 2001
From: Siu Chi Chan <siuchi.chan@amd.com>
Date: Sat, 4 May 2019 15:08:28 -0400
Subject: [PATCH 3/6] move executable_cache into program_state.cpp

[ROCm/clr commit: 1a2d332e761d2c4f647421b2f2585446a2f53287]
---
 .../include/hip/hcc_detail/program_state.hpp  | 27 ------
 projects/clr/hipamd/src/hip_hcc.cpp           |  9 --
 projects/clr/hipamd/src/program_state.cpp     | 87 ++++++++++---------
 projects/clr/hipamd/src/program_state.inl     | 24 +++++
 4 files changed, 71 insertions(+), 76 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp
index 5741c22b4b..b3656e1e48 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp
+++ b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp
@@ -30,39 +30,12 @@ THE SOFTWARE.
 #include <cstddef>
 #include <cstdint>
 #include <cstdlib>
-#include <stdexcept>
 
 struct ihipModuleSymbol_t;
 using hipFunction_t = ihipModuleSymbol_t*;
 
-namespace std {
-template<>
-struct hash<hsa_agent_t> {
-    size_t operator()(hsa_agent_t x) const {
-        return hash<decltype(x.handle)>{}(x.handle);
-    }
-};
-
-template<>
-struct hash<hsa_isa_t> {
-    size_t operator()(hsa_isa_t x) const {
-        return hash<decltype(x.handle)>{}(x.handle);
-    }
-};
-}  // namespace std
-
-inline constexpr bool operator==(hsa_agent_t x, hsa_agent_t y) {
-    return x.handle == y.handle;
-}
-inline constexpr bool operator==(hsa_isa_t x, hsa_isa_t y) {
-    return x.handle == y.handle;
-}
-
 namespace hip_impl {
 
-[[noreturn]]
-void hip_throw(const std::exception&);
-
 class kernargs_size_align;
 class program_state_impl;
 class program_state {
diff --git a/projects/clr/hipamd/src/hip_hcc.cpp b/projects/clr/hipamd/src/hip_hcc.cpp
index 648b53dab7..583017a30a 100644
--- a/projects/clr/hipamd/src/hip_hcc.cpp
+++ b/projects/clr/hipamd/src/hip_hcc.cpp
@@ -2492,13 +2492,4 @@ namespace hip_impl {
         #endif
     }
 
-    std::mutex executables_cache_mutex;
-
-    std::vector<hsa_executable_t>& executables_cache(
-            std::string elf, hsa_isa_t isa, hsa_agent_t agent) {
-        static std::unordered_map<std::string,
-            std::unordered_map<hsa_isa_t,
-                std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>>> cache;
-        return cache[elf][isa][agent];
-    }
 } // Namespace hip_impl.
diff --git a/projects/clr/hipamd/src/program_state.cpp b/projects/clr/hipamd/src/program_state.cpp
index e226e1af66..906584bc38 100644
--- a/projects/clr/hipamd/src/program_state.cpp
+++ b/projects/clr/hipamd/src/program_state.cpp
@@ -1,4 +1,6 @@
 #include "../include/hip/hcc_detail/program_state.hpp"
+// contains implementation of program_state_impl
+#include "program_state.inl"
 
 #include <hsa/hsa.h>
 
@@ -7,53 +9,58 @@
 #include <unordered_map>
 #include <vector>
 
-// contains implementation of program_state_impl
-#include "program_state.inl"
-
 namespace hip_impl {
 
-std::size_t kernargs_size_align::kernargs_size_align::size(std::size_t n) const{
-    return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].first;
-}
+    std::size_t kernargs_size_align::kernargs_size_align::size(std::size_t n) const{
+        return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].first;
+    }
 
-std::size_t kernargs_size_align::alignment(std::size_t n) const{
-    return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].second;
-}
+    std::size_t kernargs_size_align::alignment(std::size_t n) const{
+        return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].second;
+    }
 
-program_state::program_state() :
-    impl(new program_state_impl) {
-    if (!impl) hip_throw(std::runtime_error {
-                         "Unknown error when constructing program state."});
-}
+    program_state::program_state() : impl(new program_state_impl) {
+        if (!impl) hip_throw(std::runtime_error {
+            "Unknown error when constructing program state."});
+    }
 
-program_state::~program_state() {
-    delete(impl);
-}
+    program_state::~program_state() {
+        delete(impl);
+    }
 
-void* program_state::global_addr_by_name(const char* name) {
-    const auto it = impl->get_globals().find(name);
-    if (it == impl->get_globals().end())
-      return nullptr;
-    else
-      return it->second;
-}
+    void* program_state::global_addr_by_name(const char* name) {
+        const auto it = impl->get_globals().find(name);
+        if (it == impl->get_globals().end())
+            return nullptr;
+        else
+            return it->second;
+    }
 
-hsa_executable_t program_state::load_executable(const char* data,
-                                                const size_t data_size,
-                                                hsa_executable_t executable,
-                                                hsa_agent_t agent) {
-    return impl->load_executable(data, data_size, executable, agent);
-}
+    hsa_executable_t program_state::load_executable(const char* data,
+        const size_t data_size,
+        hsa_executable_t executable,
+        hsa_agent_t agent) {
+        return impl->load_executable(data, data_size, executable, agent);
+    }
 
-hipFunction_t program_state::kernel_descriptor(std::uintptr_t function_address,
-                                                          hsa_agent_t agent) {
-    auto& kd = impl->kernel_descriptor(function_address, agent);
-    return kd;
-}
+    hipFunction_t program_state::kernel_descriptor(std::uintptr_t function_address,
+        hsa_agent_t agent) {
+        auto& kd = impl->kernel_descriptor(function_address, agent);
+        return kd;
+    }
 
-kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t kernel) {
-  kernargs_size_align t;
-  t.handle = reinterpret_cast<const void*>(&impl->kernargs_size_align(kernel));
-  return t;
-}
+    kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t kernel) {
+        kernargs_size_align t;
+        t.handle = reinterpret_cast<const void*>(&impl->kernargs_size_align(kernel));
+        return t;
+    }
+
+    std::mutex executables_cache_mutex;
+    std::vector<hsa_executable_t>& executables_cache(
+        std::string elf, hsa_isa_t isa, hsa_agent_t agent) {
+        static std::unordered_map<std::string,
+            std::unordered_map<hsa_isa_t,
+                std::unordered_map<hsa_agent_t, std::vector<hsa_executable_t>>>> cache;
+        return cache[elf][isa][agent];
+    }
 };
diff --git a/projects/clr/hipamd/src/program_state.inl b/projects/clr/hipamd/src/program_state.inl
index f1397b3fe9..92b2c74597 100644
--- a/projects/clr/hipamd/src/program_state.inl
+++ b/projects/clr/hipamd/src/program_state.inl
@@ -26,12 +26,36 @@
 #include <cstdio>
 #include <memory>
 #include <mutex>
+#include <stdexcept>
 #include <string>
 #include <sstream>
 #include <unordered_map>
 #include <utility>
 #include <vector>
 
+namespace std {
+template<>
+struct hash<hsa_agent_t> {
+    size_t operator()(hsa_agent_t x) const {
+        return hash<decltype(x.handle)>{}(x.handle);
+    }
+};
+
+template<>
+struct hash<hsa_isa_t> {
+    size_t operator()(hsa_isa_t x) const {
+        return hash<decltype(x.handle)>{}(x.handle);
+    }
+};
+}  // namespace std
+
+inline constexpr bool operator==(hsa_agent_t x, hsa_agent_t y) {
+    return x.handle == y.handle;
+}
+inline constexpr bool operator==(hsa_isa_t x, hsa_isa_t y) {
+    return x.handle == y.handle;
+}
+
 namespace hip_impl {
 
 [[noreturn]]

From bafd29662c5c1a88556cf044605dab4c9e4e2e04 Mon Sep 17 00:00:00 2001
From: Siu Chi Chan <siuchi.chan@amd.com>
Date: Sat, 4 May 2019 17:51:13 -0400
Subject: [PATCH 4/6] replace std::vector for kernarg

[ROCm/clr commit: 1fb9ab2d44c5c3d44c73d411a073f8687b31f48b]
---
 .../hip/hcc_detail/functional_grid_launch.hpp | 12 ++++----
 .../include/hip/hcc_detail/program_state.hpp  | 14 ++++++++++
 projects/clr/hipamd/src/program_state.cpp     | 28 +++++++++++++++++++
 projects/clr/hipamd/src/program_state.inl     |  5 ++++
 4 files changed, 53 insertions(+), 6 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp
index 8082188489..3e464a3326 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp
+++ b/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp
@@ -57,10 +57,10 @@ template <
     std::size_t n,
     typename... Ts,
     typename std::enable_if<n == sizeof...(Ts)>::type* = nullptr>
-inline std::vector<std::uint8_t> make_kernarg(
+inline hip_impl::kernarg make_kernarg(
     const std::tuple<Ts...>&,
     const kernargs_size_align&,
-    std::vector<std::uint8_t> kernarg) {
+    hip_impl::kernarg kernarg) {
     return kernarg;
 }
 
@@ -68,10 +68,10 @@ template <
     std::size_t n,
     typename... Ts,
     typename std::enable_if<n != sizeof...(Ts)>::type* = nullptr>
-inline std::vector<std::uint8_t> make_kernarg(
+inline hip_impl::kernarg make_kernarg(
     const std::tuple<Ts...>& formals,
     const kernargs_size_align& size_align,
-    std::vector<std::uint8_t> kernarg) {
+    hip_impl::kernarg kernarg) {
     using T = typename std::tuple_element<n, std::tuple<Ts...>>::type;
 
     static_assert(
@@ -96,7 +96,7 @@ inline std::vector<std::uint8_t> make_kernarg(
 }
 
 template <typename... Formals, typename... Actuals>
-inline std::vector<std::uint8_t> make_kernarg(
+inline hip_impl::kernarg make_kernarg(
     void (*kernel)(Formals...), std::tuple<Actuals...> actuals) {
     static_assert(sizeof...(Formals) == sizeof...(Actuals),
         "The count of formal arguments must match the count of actuals.");
@@ -104,7 +104,7 @@ inline std::vector<std::uint8_t> make_kernarg(
     if (sizeof...(Formals) == 0) return {};
 
     std::tuple<Formals...> to_formals{std::move(actuals)};
-    std::vector<std::uint8_t> kernarg;
+    hip_impl::kernarg kernarg;
     kernarg.reserve(sizeof(to_formals));
 
     auto& ps = hip_impl::get_program_state();
diff --git a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp
index b3656e1e48..c64b64fde8 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp
+++ b/projects/clr/hipamd/include/hip/hcc_detail/program_state.hpp
@@ -36,6 +36,20 @@ using hipFunction_t = ihipModuleSymbol_t*;
 
 namespace hip_impl {
 
+struct kernarg_impl;
+class kernarg {
+public:
+    kernarg();
+    kernarg(kernarg&&);
+    ~kernarg();
+    std::uint8_t* data();
+    std::size_t   size();
+    void reserve(std::size_t);
+    void resize(std::size_t);
+private:
+    kernarg_impl* impl;
+};
+
 class kernargs_size_align;
 class program_state_impl;
 class program_state {
diff --git a/projects/clr/hipamd/src/program_state.cpp b/projects/clr/hipamd/src/program_state.cpp
index 906584bc38..dbd7d3ebc4 100644
--- a/projects/clr/hipamd/src/program_state.cpp
+++ b/projects/clr/hipamd/src/program_state.cpp
@@ -10,6 +10,34 @@
 #include <vector>
 
 namespace hip_impl {
+    
+    kernarg::kernarg() : impl(new kernarg_impl) { 
+    }
+
+    kernarg::kernarg(kernarg&& k) : impl(k.impl) {
+        k.impl = nullptr;
+    }
+
+    kernarg::~kernarg() {
+        if (impl)
+            delete(impl);
+    }
+
+    std::uint8_t* kernarg::data() {
+        return impl->v.data();
+    }
+
+    std::size_t kernarg::size() {
+        return impl->v.size();
+    }
+
+    void kernarg::reserve(std::size_t c) {
+        impl->v.reserve(c);
+    }
+
+    void kernarg::resize(std::size_t c) {
+        impl->v.resize(c);
+    }
 
     std::size_t kernargs_size_align::kernargs_size_align::size(std::size_t n) const{
         return (*reinterpret_cast<const std::vector<std::pair<std::size_t, std::size_t>>*>(handle))[n].first;
diff --git a/projects/clr/hipamd/src/program_state.inl b/projects/clr/hipamd/src/program_state.inl
index 92b2c74597..9fc590ef75 100644
--- a/projects/clr/hipamd/src/program_state.inl
+++ b/projects/clr/hipamd/src/program_state.inl
@@ -743,4 +743,9 @@ public:
     }
 };  // class program_state_impl
 
+struct kernarg_impl {
+    std::vector<std::uint8_t> v;
+};
+
+
 };

From f2cdf536f616045d49a45089300af9f42efac788 Mon Sep 17 00:00:00 2001
From: Siu Chi Chan <siuchi.chan@amd.com>
Date: Mon, 6 May 2019 12:03:09 -0400
Subject: [PATCH 5/6] fix breakage due to compiling in C++17

[ROCm/clr commit: 2bc9455928352dce6471ee4aad99e9d569cd4fcb]
---
 .../clr/hipamd/include/hip/hcc_detail/helpers.hpp     | 11 ++++++++++-
 1 file changed, 10 insertions(+), 1 deletion(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp b/projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp
index 1916945c1a..9366885f8b 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp
+++ b/projects/clr/hipamd/include/hip/hcc_detail/helpers.hpp
@@ -84,12 +84,21 @@ struct is_callable_impl<F(Ts...), 4u, void_t_<decltype(std::declval<F>()(std::de
 // Not callable.
 template <FunctionalProcedure F>
 struct is_callable_impl<F, 5u> : std::false_type {};
-#else
+#elif  (__cplusplus < 201703L)
 template <typename, typename = void>
 struct is_callable_impl : std::false_type {};
 
 template <FunctionalProcedure F, typename... Ts>
 struct is_callable_impl<F(Ts...), void_t_<std::result_of_t<F(Ts...)> > > : std::true_type {};
+#else
+
+// C++17
+
+template <typename, typename = void>
+struct is_callable_impl : std::false_type {};
+
+template <FunctionalProcedure F, typename... Ts>
+struct is_callable_impl<F(Ts...), void_t_<std::invoke_result<F(Ts...)> > > : std::true_type {};
 #endif
 template <typename Call>
 struct is_callable : is_callable_impl<Call> {};

From cf2e2c507793dcec4f4e259abf698d0723fbdefe Mon Sep 17 00:00:00 2001
From: Siu Chi Chan <siuchi.chan@amd.com>
Date: Fri, 10 May 2019 17:23:23 -0400
Subject: [PATCH 6/6] remove code_object_bundle.hpp, clean up the old
 Agent_global

[ROCm/clr commit: d40a4627c3989c47a82354710147ba2be15d0f2f]
---
 .../hip/hcc_detail/functional_grid_launch.hpp |  1 -
 .../include/hip/hcc_detail/hip_runtime_api.h  | 42 -------------------
 2 files changed, 43 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp b/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp
index 3e464a3326..b09d0d2c6c 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp
+++ b/projects/clr/hipamd/include/hip/hcc_detail/functional_grid_launch.hpp
@@ -22,7 +22,6 @@ THE SOFTWARE.
 
 #pragma once
 
-#include "code_object_bundle.hpp"
 #include "concepts.hpp"
 #include "helpers.hpp"
 #include "program_state.hpp"
diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
index 17a42e2a16..5dd7227f61 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
+++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h
@@ -2590,48 +2590,6 @@ hipError_t hipModuleGetFunction(hipFunction_t* function, hipModule_t module, con
 
 hipError_t hipFuncGetAttributes(struct hipFuncAttributes* attr, const void* func);
 
-#if 0
-struct Agent_global {
-
-    Agent_global() : name(nullptr), address(nullptr), byte_cnt(0) {}
-    Agent_global(const char* name, hipDeviceptr_t address, uint32_t byte_cnt) 
-      : name(nullptr), address(address), byte_cnt(byte_cnt) {
-      if (name)
-        this->name = strdup(name);
-    }
-
-    Agent_global& operator=(Agent_global&& t) {
-      if (this == &t) return *this;
-
-      if (name) free(name);
-      name = t.name;
-      address = t.address;
-      byte_cnt = t.byte_cnt;
-
-      t.name = nullptr;
-      t.address = nullptr;
-      t.byte_cnt = 0;
-
-      return *this;
-    }
-
-    Agent_global(Agent_global&& t) 
-      : name(nullptr), address(nullptr), byte_cnt(0) {
-      *this = std::move(t);
-    }
-
-    // not needed, delete them to prevent bugs
-    Agent_global(const Agent_global&) = delete;
-    Agent_global& operator=(Agent_global& t) = delete;
-
-    ~Agent_global() { if (name) free(name); }
-    
-    char* name;
-    hipDeviceptr_t address;
-    uint32_t byte_cnt;
-};
-#endif
-
 #if !__HIP_VDI__
 #if defined(__cplusplus)
 } // extern "C"