This switches HIP from its currently convoluted macro + pfe based dispatch mechanism to a more natural one partially based on the existing module API. The basic idea is that HCC will always correctly emit __global__ functions: as empty-bodied stubs, on host, and as kernels, on device. It then becomes trivial to obtain the mangled name on host, at dispatch, from the function's address, and then to use the mangled name to retrieve the kernel. This should address all problems stemming from serialisation, dubious mismatches due to the manufactured functor, macro-isms et al. It also immediately enables support for generalised globals as a consequence of that being available in the module API. Finally, it will make debug much easier, since the actual names of the __global__ functions will automatically be used in traces etc. One detail is that due to how dispatch works now (hipLaunchKernel and hipLaunchKernelGGL are themselves variadic function templates which deduce the function type of the callee), in certain cases it may be necesssary to insert explicit casts to ensure that the variadic argument list selects a viable overload - this can be observed in some unit tests. Eventually we may be able to remove this limitation, but for now it does not appear terribly onerous. The code is not extremely HIPpie, nor is it fully optimised, but rather is intended as a starting point for the HIP team to make its own.

hipamd/src/code_object_bundle.cpp

@@ -0,0 +1,39 @@
+#include "../include/hip/hcc_detail/code_object_bundle.hpp"
+
+#include <hsa/hsa.h>
+
+#include <cstdint>
+#include <string>
+#include <vector>
+
+hsa_isa_t hip_impl::triple_to_hsa_isa(const std::string& triple)
+{
+    static constexpr const char prefix[] = "hcc-amdgcn--amdhsa-gfx";
+    static constexpr std::size_t prefix_sz = sizeof(prefix) - 1;
+
+    hsa_isa_t r = {};
+
+    auto idx = triple.find(prefix);
+
+    if (idx != std::string::npos) {
+        idx += prefix_sz;
+        std::string tmp = "AMD:AMDGPU";
+        while (idx != triple.size()) {
+            tmp.push_back(':');
+            tmp.push_back(triple[idx++]);
+        }
+
+        hsa_isa_from_name(tmp.c_str(), &r);
+    }
+
+    return r;
+}
+
+// DATA - STATICS
+constexpr const char hip_impl::Bundled_code_header::magic_string_[];
+
+// CREATORS
+hip_impl::Bundled_code_header::Bundled_code_header(
+    const std::vector<std::uint8_t>& x)
+    : Bundled_code_header{x.cbegin(), x.cend()}
+{}

hipamd/src/grid_launch.cpp

@@ -21,76 +21,118 @@ THE SOFTWARE.
 */
 
 #include "hip/hcc_detail/grid_launch_GGL.hpp"
+#include "hip/hcc_detail/program_state.hpp"
+
+#include "hip/hip_runtime_api.h"
 
 // Internal header, do not percolate upwards.
 #include "hip_hcc_internal.h"
 #include "hc.hpp"
 #include "trace_helper.h"
 
+#include <cassert>
+#include <cstddef>
+#include <cstdint>
+#include <stdexcept>
+
 #include <iostream>
-#include <sstream>
+
+using namespace hc;
+using namespace std;
 
 namespace hip_impl
 {
-    hc::accelerator_view lock_stream_hip_(
-        hipStream_t& stream, void*& locked_stream)
-    {   // This allocated but does not take ownership of locked_stream. If it is
-        // not deleted elsewhere it will leak.
-        using L = decltype(stream->lockopen_preKernelCommand());
-
-        HIP_INIT();
-
-        stream = ihipSyncAndResolveStream(stream);
-        locked_stream = new L{stream->lockopen_preKernelCommand()};
-        return (*static_cast<L*>(locked_stream))->_av;
-    }
-
-    void print_prelaunch_trace_(
-        const char* kernel_name,
-        dim3 num_blocks,
-        dim3 dim_blocks,
-        int group_mem_bytes,
-        hipStream_t stream)
+    namespace
     {
-        if ((HIP_TRACE_API & (1 << TRACE_KCMD)) ||
-            HIP_PROFILE_API ||
-            (COMPILE_HIP_DB && (HIP_TRACE_API & (1<<TRACE_ALL)))) {
-            std::stringstream os;
-            os  << tls_tidInfo.tid() << "." << tls_tidInfo.apiSeqNum()
-                << " hipLaunchKernel '" << kernel_name << "'"
-                << " gridDim:"  << num_blocks
-                << " groupDim:" << dim_blocks
-                << " sharedMem:+" << group_mem_bytes
-                << " " << *stream;
+        inline
+        string name(uintptr_t function_address)
+        {
+            const auto it = function_names().find(function_address);
 
-            if (HIP_PROFILE_API == 0x1) {
-                std::string shortAtpString("hipLaunchKernel:");
-                shortAtpString += kernel_name;
-                MARKER_BEGIN(shortAtpString.c_str(), "HIP");
-            } else if (HIP_PROFILE_API == 0x2) {
-                MARKER_BEGIN(os.str().c_str(), "HIP");
+            if (it == function_names().cend())  {
+                throw runtime_error{
+                    "Invalid function passed to hipLaunchKernelGGL."};
             }
 
-            if (COMPILE_HIP_DB && HIP_TRACE_API) {
-                std::string fullStr;
-                recordApiTrace(&fullStr, os.str());
+            return it->second;
+        }
+
+        inline
+        string name(hsa_agent_t agent)
+        {
+            char n[64] = {};
+            hsa_agent_get_info(agent, HSA_AGENT_INFO_NAME, n);
+
+            return string{n};
+        }
+
+        inline
+        hsa_agent_t target_agent(hipStream_t stream)
+        {
+            if (stream) {
+                return *static_cast<hsa_agent_t*>(
+                    stream->locked_getAv()->get_hsa_agent());
+            }
+            else if (
+                ihipGetTlsDefaultCtx() && ihipGetTlsDefaultCtx()->getDevice()) {
+                return ihipGetDevice(
+                    ihipGetTlsDefaultCtx()->getDevice()->_deviceId)->_hsaAgent;
+            }
+            else {
+                return *static_cast<hsa_agent_t*>(
+                    accelerator{}.get_default_view().get_hsa_agent());
             }
         }
     }
 
-    void unlock_stream_hip_(
+    void hipLaunchKernelGGLImpl(
+        uintptr_t function_address,
+        const dim3& numBlocks,
+        const dim3& dimBlocks,
+        uint32_t sharedMemBytes,
         hipStream_t stream,
-        void* locked_stream,
-        const char* kernel_name,
-        hc::accelerator_view* acc_v)
-    {   // Precondition: acc_v is the accelerator_view associated with stream
-        //               which is guarded by locked_stream;
-        //               locked_stream is deletable.
-        using L = decltype(stream->lockopen_preKernelCommand());
+        void** kernarg)
+    {
+        const auto it0 = functions().find(function_address);
 
-        stream->lockclose_postKernelCommand(kernel_name, acc_v);
+        if (it0 == functions().cend()) {
+            throw runtime_error{
+                "No device code available for function: " +
+                name(function_address)
+            };
+        }
 
-        delete static_cast<L*>(locked_stream);
-        locked_stream = nullptr;
+        auto agent = target_agent(stream);
+
+        const auto it1 = find_if(
+            it0->second.cbegin(),
+            it0->second.cend(),
+            [=](const pair<hsa_agent_t, Kernel_descriptor>& x) {
+                return x.first.handle == agent.handle;
+            });
+
+        if (it1 == it0->second.cend()) {
+            throw runtime_error{
+                "No code available for function: " + name(function_address) +
+                ", for agent: " + name(agent)
+            };
+        }
+
+        for (auto&& agent_kernel : it0->second) {
+            if (agent.handle == agent_kernel.first.handle) {
+                hipModuleLaunchKernel(
+                    agent_kernel.second,
+                    numBlocks.x,
+                    numBlocks.y,
+                    numBlocks.z,
+                    dimBlocks.x,
+                    dimBlocks.y,
+                    dimBlocks.z,
+                    sharedMemBytes,
+                    stream,
+                    nullptr,
+                    kernarg);
+            }
+        }
     }
 }

hipamd/src/hip_hcc_internal.h

@@ -233,7 +233,7 @@ static const DbName dbName [] =
 #if COMPILE_HIP_DB
 #define tprintf(trace_level, ...) {\
     if (HIP_DB & (1<<(trace_level))) {\
-        char msgStr[1000];\
+        char msgStr[2000];\
         snprintf(msgStr, 2000, __VA_ARGS__);\
         fprintf (stderr, "  %ship-%s tid:%d:%s%s", dbName[trace_level]._color, dbName[trace_level]._shortName, tls_tidInfo.tid(), msgStr, KNRM); \
     }\

hipamd/src/hip_memory.cpp

@@ -65,7 +65,7 @@ int sharePtr(void *ptr, ihipCtx_t *ctx, bool shareWithAll, unsigned hipFlags)
 
     if (shareWithAll) {
         hsa_status_t s = hsa_amd_agents_allow_access(g_deviceCnt+1, g_allAgents, NULL, ptr);
-        tprintf (DB_MEM, "    allow access to CPU + all %d GPUs (shareWithAll)\n", g_deviceCnt); 
+        tprintf (DB_MEM, "    allow access to CPU + all %d GPUs (shareWithAll)\n", g_deviceCnt);
         if (s != HSA_STATUS_SUCCESS) {
             ret = -1;
         }
@@ -122,7 +122,7 @@ void * allocAndSharePtr(const char *msg, size_t sizeBytes, ihipCtx_t *ctx, bool
     if (HIP_INIT_ALLOC != -1) {
         // TODO , dont' call HIP API directly here:
         hipMemset(ptr, HIP_INIT_ALLOC, sizeBytes);
-    }	
+    }
 
     if (ptr != nullptr) {
         int r = sharePtr(ptr, ctx, shareWithAll, hipFlags);
@@ -251,7 +251,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes)
             hip_status = hipErrorMemoryAllocation;
         }
 
-    } 
+    }
 
 
     return ihipLogStatus(hip_status);
@@ -284,10 +284,10 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
         }
 
 
-        const unsigned supportedFlags = hipHostMallocPortable 
-                                      | hipHostMallocMapped 
-                                      | hipHostMallocWriteCombined 
-                                      | hipHostMallocCoherent 
+        const unsigned supportedFlags = hipHostMallocPortable
+                                      | hipHostMallocMapped
+                                      | hipHostMallocWriteCombined
+                                      | hipHostMallocCoherent
                                       | hipHostMallocNonCoherent;
 
 
@@ -300,7 +300,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
             hip_status = hipErrorInvalidValue;
         } else {
             auto device = ctx->getWriteableDevice();
-            
+
             unsigned amFlags = 0;
             if (flags & hipHostMallocCoherent) {
                 amFlags = amHostCoherent;
@@ -581,7 +581,7 @@ hipError_t hipMalloc3DArray(hipArray_t *array,
         hsa_ext_image_data_info_t imageInfo;
         hsa_status_t status = hsa_ext_image_data_get_info(*agent, &imageDescriptor, permission, &imageInfo);
         size_t alignment = imageInfo.alignment <= allocGranularity ? 0 : imageInfo.alignment;
- 
+
         *ptr = hip_internal::allocAndSharePtr("device_array", allocSize, ctx, false, am_flags, 0, alignment);
 
         if (size && (*ptr == NULL)) {
@@ -1585,7 +1585,7 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr){
     HIP_INIT_API ( handle, devPtr);
     hipError_t hipStatus = hipSuccess;
     // Get the size of allocated pointer
-    size_t psize;
+    size_t psize = 0u;
     hc::accelerator acc;
     if((handle == NULL) || (devPtr == NULL)) {
         hipStatus = hipErrorInvalidResourceHandle;

hipamd/src/hip_module.cpp

@@ -119,15 +119,18 @@ namespace hipdrv {
 uint64_t PrintSymbolSizes(const void *emi, const char *name){
     using namespace ELFIO;
 
-    const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi;
+    const ELFIO::Elf64_Ehdr *ehdr = (const ELFIO::Elf64_Ehdr*)emi;
     if(NULL == ehdr || EV_CURRENT != ehdr->e_version){}
-    const Elf64_Shdr * shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
+    const ELFIO::Elf64_Shdr * shdr =
+        (const ELFIO::Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
     for(uint16_t i=0;i<ehdr->e_shnum;++i){
         if(shdr[i].sh_type == SHT_SYMTAB){
-            const Elf64_Sym *syms = (const Elf64_Sym*)((char*)emi + shdr[i].sh_offset);
+            const ELFIO::Elf64_Sym *syms =
+                (const ELFIO::Elf64_Sym*)((char*)emi + shdr[i].sh_offset);
             assert(syms);
             uint64_t numSyms = shdr[i].sh_size/shdr[i].sh_entsize;
-            const char* strtab = (const char*)((char*)emi + shdr[shdr[i].sh_link].sh_offset);
+            const char* strtab =
+                (const char*)((char*)emi + shdr[shdr[i].sh_link].sh_offset);
             assert(strtab);
             for(uint64_t i=0;i<numSyms;++i){
                 const char *symname = strtab + syms[i].st_name;
@@ -145,8 +148,8 @@ uint64_t PrintSymbolSizes(const void *emi, const char *name){
 uint64_t ElfSize(const void *emi){
     using namespace ELFIO;
 
-    const Elf64_Ehdr *ehdr = (const Elf64_Ehdr*)emi;
-    const Elf64_Shdr *shdr = (const Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
+    const ELFIO::Elf64_Ehdr *ehdr = (const ELFIO::Elf64_Ehdr*)emi;
+    const ELFIO::Elf64_Shdr *shdr = (const ELFIO::Elf64_Shdr*)((char*)emi + ehdr->e_shoff);
 
     uint64_t max_offset = ehdr->e_shoff;
     uint64_t total_size = max_offset + ehdr->e_shentsize * ehdr->e_shnum;
@@ -164,156 +167,8 @@ uint64_t ElfSize(const void *emi){
     return total_size;
 }
 
-namespace
-{
-    template<typename P>
-    inline
-    ELFIO::section* find_section_if(ELFIO::elfio& reader, P p)
-    {
-        using namespace std;
-
-        const auto it = find_if(
-            reader.sections.begin(), reader.sections.end(), move(p));
-
-        return it != reader.sections.end() ? *it : nullptr;
-    }
-
-    inline
-    std::vector<std::string> copy_names_of_undefined_symbols(
-        const ELFIO::symbol_section_accessor& section)
-    {
-        using namespace ELFIO;
-        using namespace std;
-
-        vector<string> r;
-
-        for (auto i = 0u; i != section.get_symbols_num(); ++i) {
-            // TODO: this is boyscout code, caching the temporaries
-            //       may be of worth.
-            string name;
-            Elf64_Addr value = 0;
-            Elf_Xword size = 0;
-            Elf_Half sect_idx = 0;
-            uint8_t bind = 0;
-            uint8_t type = 0;
-            uint8_t other = 0;
-
-            section.get_symbol(
-                i, name, value, size, bind, type, sect_idx, other);
-
-            if (sect_idx == SHN_UNDEF && !name.empty()) {
-                r.push_back(std::move(name));
-            }
-        }
-
-        return r;
-    }
-
-    inline
-    std::pair<ELFIO::Elf64_Addr, ELFIO::Elf_Xword> find_symbol_address(
-        const ELFIO::symbol_section_accessor& section,
-        const std::string& symbol_name)
-    {
-        using namespace ELFIO;
-        using namespace std;
-
-        static constexpr pair<Elf64_Addr, Elf_Xword> r{0, 0};
-
-        for (auto i = 0u; i != section.get_symbols_num(); ++i) {
-            // TODO: this is boyscout code, caching the temporaries
-            //       may be of worth.
-            string name;
-            Elf64_Addr value = 0;
-            Elf_Xword size = 0;
-            Elf_Half sect_idx = 0;
-            uint8_t bind = 0;
-            uint8_t type = 0;
-            uint8_t other = 0;
-
-            section.get_symbol(
-                i, name, value, size, bind, type, sect_idx, other);
-
-            if (name == symbol_name) return make_pair(value, size);
-        }
-
-        return r;
-    }
-
-    inline
-    void associate_code_object_symbols_with_host_allocation(
-        const ELFIO::elfio& reader,
-        const ELFIO::elfio& self_reader,
-        ELFIO::section* code_object_dynsym,
-        ELFIO::section* process_symtab,
-        hsa_agent_t agent,
-        hsa_executable_t executable)
-    {
-        using namespace ELFIO;
-        using namespace std;
-
-        if (!code_object_dynsym || !process_symtab) return;
-
-        const auto undefined_symbols = copy_names_of_undefined_symbols(
-            symbol_section_accessor{reader, code_object_dynsym});
-
-        for (auto&& x : undefined_symbols) {
-            const auto tmp = find_symbol_address(
-                symbol_section_accessor{self_reader, process_symtab}, x);
-
-            assert(tmp.first);
-
-            void* p = nullptr;
-            hsa_amd_memory_lock(
-                reinterpret_cast<void*>(tmp.first), tmp.second, &agent, 1, &p);
-
-            hsa_executable_agent_global_variable_define(
-                executable, agent, x.c_str(), p);
-
-            static vector<
-                unique_ptr<void, decltype(hsa_amd_memory_unlock)*>> globals;
-            static mutex mtx;
-
-            lock_guard<std::mutex> lck{mtx};
-            globals.emplace_back(p, hsa_amd_memory_unlock);
-        }
-    }
-
-    inline
-    void load_code_object_and_freeze_executable(
-        const char* file, hsa_agent_t agent, hsa_executable_t executable)
-    {   // TODO: the following sequence is inefficient, should be refactored
-        //       into a single load of the file and subsequent ELFIO
-        //       processing.
-        using namespace std;
-
-        static const auto cor_deleter = [](hsa_code_object_reader_t* p) {
-            hsa_code_object_reader_destroy(*p);
-        };
-
-        using RAII_code_reader = unique_ptr<
-            hsa_code_object_reader_t, decltype(cor_deleter)>;
-
-        unique_ptr<FILE, decltype(fclose)*> cobj{fopen(file, "r"), fclose};
-        RAII_code_reader tmp{new hsa_code_object_reader_t, cor_deleter};
-        hsa_code_object_reader_create_from_file(fileno(cobj.get()), tmp.get());
-
-        hsa_executable_load_agent_code_object(
-            executable, agent, *tmp, nullptr, nullptr);
-
-        hsa_executable_freeze(executable, nullptr);
-
-        static vector<RAII_code_reader> code_readers;
-        static mutex mtx;
-
-        lock_guard<mutex> lck{mtx};
-        code_readers.push_back(move(tmp));
-    }
-}
-
 hipError_t hipModuleLoad(hipModule_t *module, const char *fname)
 {
-    using namespace ELFIO;
-
     HIP_INIT_API(module, fname);
     hipError_t ret = hipSuccess;
     *module = new ihipModule_t;
@@ -336,36 +191,14 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname)
             nullptr,
             &(*module)->executable);
 
-        elfio reader;
-        if (!reader.load(fname)) {
+        std::ifstream file{fname};
+
+        if (!file.is_open()) {
             return ihipLogStatus(hipErrorFileNotFound);
         }
-        else {
-            // TODO: this may benefit from caching as well.
-            elfio self_reader;
-            self_reader.load("/proc/self/exe");
-
-            const auto symtab =
-                find_section_if(self_reader, [](const ELFIO::section* x) {
-                    return x->get_type() == SHT_SYMTAB;
-                });
-
-            const auto code_object_dynsym =
-                find_section_if(reader, [](const ELFIO::section* x) {
-                    return x->get_type() == SHT_DYNSYM;
-                });
-
-            associate_code_object_symbols_with_host_allocation(
-                reader,
-                self_reader,
-                code_object_dynsym,
-                symtab,
-                currentDevice->_hsaAgent,
-                (*module)->executable);
-
-            load_code_object_and_freeze_executable(
-                fname, currentDevice->_hsaAgent, (*module)->executable);
-        }
+        (*module)->executable = hip_impl::load_executable(
+            (*module)->executable, currentDevice->_hsaAgent, file);
+        ret = (*module)->executable.handle ? hipSuccess : hipErrorUnknown;
     }
 
     return ihipLogStatus(ret);

hipamd/src/program_state.cpp

@@ -0,0 +1,498 @@
+#include "../include/hip/hcc_detail/program_state.hpp"
+
+#include "../include/hip/hcc_detail/code_object_bundle.hpp"
+
+#include "hip_hcc_internal.h"
+#include "trace_helper.h"
+
+#include "elfio/elfio.hpp"
+
+#include <link.h>
+
+#include <hsa/hsa.h>
+#include <hsa/hsa_ext_amd.h>
+
+#include <cassert>
+#include <cstddef>
+#include <cstdint>
+#include <memory>
+#include <mutex>
+#include <stdexcept>
+#include <string>
+#include <unordered_map>
+#include <utility>
+#include <vector>
+
+using namespace ELFIO;
+using namespace hip_impl;
+using namespace std;
+
+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);
+        }
+    };
+}
+
+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
+{
+    vector<string> copy_names_of_undefined_symbols(
+        const symbol_section_accessor& section)
+    {
+        vector<string> r;
+
+        for (auto i = 0u; i != section.get_symbols_num(); ++i) {
+            // TODO: this is boyscout code, caching the temporaries
+            //       may be of worth.
+            string name;
+            Elf64_Addr value = 0;
+            Elf_Xword size = 0;
+            Elf_Half sect_idx = 0;
+            uint8_t bind = 0;
+            uint8_t type = 0;
+            uint8_t other = 0;
+
+            section.get_symbol(
+                i, name, value, size, bind, type, sect_idx, other);
+
+            if (sect_idx == SHN_UNDEF && !name.empty()) {
+                r.push_back(std::move(name));
+            }
+        }
+
+        return r;
+    }
+
+    pair<Elf64_Addr, Elf_Xword> find_symbol_address(
+        const symbol_section_accessor& section,
+        const string& symbol_name)
+    {
+        static constexpr pair<Elf64_Addr, Elf_Xword> r{0, 0};
+
+        for (auto i = 0u; i != section.get_symbols_num(); ++i) {
+            // TODO: this is boyscout code, caching the temporaries
+            //       may be of worth.
+            string name;
+            Elf64_Addr value = 0;
+            Elf_Xword size = 0;
+            Elf_Half sect_idx = 0;
+            uint8_t bind = 0;
+            uint8_t type = 0;
+            uint8_t other = 0;
+
+            section.get_symbol(
+                i, name, value, size, bind, type, sect_idx, other);
+
+            if (name == symbol_name) return make_pair(value, size);
+        }
+
+        return r;
+    }
+
+    void associate_code_object_symbols_with_host_allocation(
+        const elfio& reader,
+        const elfio& self_reader,
+        section* code_object_dynsym,
+        section* process_symtab,
+        hsa_agent_t agent,
+        hsa_executable_t executable)
+    {
+        if (!code_object_dynsym || !process_symtab) return;
+
+        const auto undefined_symbols = copy_names_of_undefined_symbols(
+            symbol_section_accessor{reader, code_object_dynsym});
+
+        for (auto&& x : undefined_symbols) {
+            const auto tmp = find_symbol_address(
+                symbol_section_accessor{self_reader, process_symtab}, x);
+
+            assert(tmp.first);
+
+            void* p = nullptr;
+            hsa_amd_memory_lock(
+                reinterpret_cast<void*>(tmp.first), tmp.second, &agent, 1, &p);
+
+            hsa_executable_agent_global_variable_define(
+                executable, agent, x.c_str(), p);
+
+            static vector<
+                unique_ptr<void, decltype(hsa_amd_memory_unlock)*>> globals;
+            static mutex mtx;
+
+            lock_guard<std::mutex> lck{mtx};
+            globals.emplace_back(p, hsa_amd_memory_unlock);
+        }
+    }
+
+    template<typename P>
+    inline
+    section* find_section_if(elfio& reader, P p)
+    {
+        const auto it = find_if(
+            reader.sections.begin(), reader.sections.end(), std::move(p));
+
+        return it != reader.sections.end() ? *it : nullptr;
+    }
+
+    vector<uint8_t> code_object_blob_for_process()
+    {
+        static constexpr const char self[] = "/proc/self/exe";
+        static constexpr const char kernel_section[] = ".kernel";
+
+        elfio reader;
+
+        if (!reader.load(self)) {
+            throw runtime_error{"Failed to load ELF file for current process."};
+        }
+
+        auto kernels = find_section_if(reader, [](const section* x) {
+            return x->get_name() == kernel_section;
+        });
+
+        vector<uint8_t> r;
+        if (kernels) {
+            r.insert(
+                r.end(),
+                kernels->get_data(),
+                kernels->get_data() + kernels->get_size());
+        }
+
+        return r;
+    }
+
+    const unordered_map<hsa_isa_t, vector<vector<uint8_t>>>& code_object_blobs()
+    {
+        static unordered_map<hsa_isa_t, vector<vector<uint8_t>>> r;
+        static once_flag f;
+
+        call_once(f, []() {
+            static vector<vector<uint8_t>> blobs{
+                code_object_blob_for_process()};
+
+            dl_iterate_phdr([](dl_phdr_info* i, std::size_t, void*) {
+                elfio tmp;
+                if (tmp.load(i->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());
+                }
+                return 0;
+            }, 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);
+                    }
+                }
+            }
+        });
+
+        return r;
+    }
+
+    const unordered_map<hsa_agent_t, vector<hsa_executable_t>>& executables()
+    {
+        static unordered_map<hsa_agent_t, vector<hsa_executable_t>> r;
+        static once_flag f;
+
+        call_once(f, []() {
+            static const auto accelerators = hc::accelerator::get_all();
+
+            for (auto&& acc : accelerators) {
+                auto agent = static_cast<hsa_agent_t*>(acc.get_hsa_agent());
+
+                if (!agent) continue;
+
+                hsa_agent_iterate_isas(*agent, [](hsa_isa_t x, void* pa) {
+                    const auto it = code_object_blobs().find(x);
+
+                    if (it != code_object_blobs().cend()) {
+                        hsa_agent_t a = *static_cast<hsa_agent_t*>(pa);
+
+                        for (auto&& blob : it->second) {
+                            hsa_executable_t tmp = {};
+
+                            hsa_executable_create_alt(
+                                HSA_PROFILE_FULL,
+                                HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
+                                nullptr,
+                                &tmp);
+
+                            // TODO: this is massively inefficient and only
+                            //       meant for illustration.
+                            string blob_to_str{blob.cbegin(), blob.cend()};
+                            stringstream istr{blob_to_str};
+                            tmp = load_executable(tmp, a, istr);
+
+                            if (tmp.handle) r[a].push_back(tmp);
+                        }
+                    }
+
+                    return HSA_STATUS_SUCCESS;
+                }, agent);
+            }
+        });
+
+        cout << r.size() << endl;
+        return r;
+    }
+
+    inline
+    hsa_agent_t agent(hsa_executable_symbol_t x)
+    {
+        hsa_agent_t r = {};
+        hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_AGENT, &r);
+
+        return r;
+    }
+
+    inline
+    uint32_t group_size(hsa_executable_symbol_t x)
+    {
+        uint32_t r = 0u;
+        hsa_executable_symbol_get_info(
+            x, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &r);
+
+        return r;
+    }
+
+    inline
+    uint64_t kernel_object(hsa_executable_symbol_t x)
+    {
+        uint64_t r = 0u;
+        hsa_executable_symbol_get_info(
+            x, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &r);
+
+        return r;
+    }
+
+    inline
+    string name(hsa_executable_symbol_t x)
+    {
+        uint32_t sz = 0u;
+        hsa_executable_symbol_get_info(
+            x, HSA_EXECUTABLE_SYMBOL_INFO_NAME_LENGTH, &sz);
+
+        string r(sz, '\0');
+        hsa_executable_symbol_get_info(
+            x, HSA_EXECUTABLE_SYMBOL_INFO_NAME, &r.front());
+
+        return r;
+    }
+
+    inline
+    uint32_t private_size(hsa_executable_symbol_t x)
+    {
+        uint32_t r = 0u;
+        hsa_executable_symbol_get_info(
+            x, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &r);
+
+        return r;
+    }
+
+    inline
+    hsa_symbol_kind_t type(hsa_executable_symbol_t x)
+    {
+        hsa_symbol_kind_t r = {};
+        hsa_executable_symbol_get_info(x, HSA_EXECUTABLE_SYMBOL_INFO_TYPE, &r);
+
+        return r;
+    }
+
+    const unordered_map<string, vector<hsa_executable_symbol_t>>& kernels()
+    {
+        static unordered_map<string, vector<hsa_executable_symbol_t>> r;
+        static once_flag f;
+
+        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);
+
+                return HSA_STATUS_SUCCESS;
+            };
+
+            for (auto&& agent_executables : executables()) {
+                for (auto&& executable : agent_executables.second) {
+                    hsa_executable_iterate_agent_symbols(
+                        executable,
+                        agent_executables.first,
+                        copy_kernels,
+                        nullptr);
+                }
+            }
+        });
+
+        return r;
+    }
+
+    void load_code_object_and_freeze_executable(
+        istream& file, hsa_agent_t agent, hsa_executable_t executable)
+    {   // TODO: the following sequence is inefficient, should be refactored
+        //       into a single load of the file and subsequent ELFIO
+        //       processing.
+        static const auto cor_deleter = [](hsa_code_object_reader_t* p) {
+            hsa_code_object_reader_destroy(*p);
+        };
+
+        using RAII_code_reader = unique_ptr<
+            hsa_code_object_reader_t, decltype(cor_deleter)>;
+
+        file.seekg(0);
+
+        vector<uint8_t> blob{
+            istreambuf_iterator<char>{file}, istreambuf_iterator<char>{}};
+        RAII_code_reader tmp{new hsa_code_object_reader_t, cor_deleter};
+        hsa_code_object_reader_create_from_memory(
+            blob.data(), blob.size(), tmp.get());
+
+        hsa_executable_load_agent_code_object(
+            executable, agent, *tmp, nullptr, nullptr);
+
+        hsa_executable_freeze(executable, nullptr);
+
+        static vector<RAII_code_reader> code_readers;
+        static mutex mtx;
+
+        lock_guard<mutex> lck{mtx};
+        code_readers.push_back(move(tmp));
+    }
+}
+
+namespace hip_impl
+{
+    const unordered_map<uintptr_t, string>& function_names()
+    {
+        static constexpr const char self[] = "/proc/self/exe";
+
+        static unordered_map<uintptr_t, string> r;
+        static once_flag f;
+
+        call_once(f, []() {
+            elfio reader;
+
+            if (!reader.load(self)) {
+                throw runtime_error{
+                    "Failed to load the ELF file for the current process."};
+            }
+
+            auto symtab = find_section_if(reader, [](const section* x) {
+                return x->get_type() == SHT_SYMTAB;
+            });
+
+            symbol_section_accessor symbols{reader, symtab};
+
+            for (auto i = 0u; i != symbols.get_symbols_num(); ++i) {
+                // TODO: this is boyscout code, caching the temporaries
+                //       may be of worth.
+                string name;
+                Elf64_Addr value = 0;
+                Elf_Xword size = 0;
+                Elf_Half sect_idx = 0;
+                uint8_t bind = 0;
+                uint8_t type = 0;
+                uint8_t other = 0;
+
+                symbols.get_symbol(
+                    i, name, value, size, bind, type, sect_idx, other);
+
+                if (type == STT_FUNC && sect_idx != SHN_UNDEF && !name.empty()) {
+                    r.emplace(value, name);
+                }
+            }
+        });
+
+        return r;
+    }
+
+    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;
+
+        call_once(f, []() {
+            for (auto&& function : function_names()) {
+                const auto it = kernels().find(function.second);
+
+                if (it != kernels().cend()) {
+                    for (auto&& kernel_symbol : it->second) {
+                        r[function.first].emplace_back(
+                            agent(kernel_symbol),
+                            Kernel_descriptor{
+                                kernel_object(kernel_symbol),
+                                group_size(kernel_symbol),
+                                private_size(kernel_symbol),
+                                it->first});
+                    }
+                }
+            }
+        });
+
+        return r;
+    }
+
+    hsa_executable_t load_executable(
+        hsa_executable_t executable, hsa_agent_t agent, istream& file)
+    {
+        elfio reader;
+        if (!reader.load(file)) {
+            return hsa_executable_t{};
+        }
+        else {
+            // TODO: this may benefit from caching as well.
+            elfio self_reader;
+            self_reader.load("/proc/self/exe");
+
+            const auto symtab =
+                find_section_if(self_reader, [](const ELFIO::section* x) {
+                    return x->get_type() == SHT_SYMTAB;
+                });
+
+            const auto code_object_dynsym =
+                find_section_if(reader, [](const ELFIO::section* x) {
+                    return x->get_type() == SHT_DYNSYM;
+                });
+
+            associate_code_object_symbols_with_host_allocation(
+                reader, self_reader, code_object_dynsym, symtab, agent, executable);
+
+            load_code_object_and_freeze_executable(file, agent, executable);
+
+            return executable;
+        }
+    }
+} // Namespace hip_impl.