Add HIPRTC, glorious ersatz for NVRTC (#1097)

* Add ersatz for NVRTC.

* Fix extraneous paren and use correct namespace.

* Use lowerCamelCase (yuck, yuck) consistently.

* Link against FS when building hiprtc lib.

* Correctly mark Manipulators. Fix dual compile.

* Add unit tests. Extend HIT to accept linker options.

* Make sure the HIPRTC library is installed.

* Better logging. Try to auto-detect the target.

* Stop specifying the target explicitly.

* Add missing flavour of `hipModuleLaunchKernel`.

* Program was already destroyed.

* Don't use `--genco`. Fix mangled name trimming.

* Fix HIPRTC breakage due to upstream noise.

* [dtests] Replace RUN -> TEST in hiprtc tests

Change-Id: Ie499e92dfe4e5c94634b1c2b76cf52d241bcfea3

* [hit] Set HIP_PATH to HIP_ROOT_DIR for all tests

Change-Id: Ib0ad1f99bc71c03e363e055dd508a7a4a210680a


[ROCm/clr commit: a538eb705a]

projects/clr/hipamd/src/hip_hcc_internal.h

@@ -371,6 +371,8 @@ struct ihipModule_t {
     hsa_executable_t executable = {};
     hsa_code_object_reader_t coReader = {};
     std::string hash;
+    std::unordered_map<
+        std::string, std::vector<std::pair<std::size_t, std::size_t>>> kernargs;
 
     ~ihipModule_t() {
         if (executable.handle) hsa_executable_destroy(executable);

projects/clr/hipamd/src/hip_module.cpp

@@ -91,6 +91,7 @@ struct ihipModuleSymbol_t {
     uint64_t _object{};  // The kernel object.
     amd_kernel_code_t const* _header{};
     string _name;  // TODO - review for performance cost.  Name is just used for debug.
+    vector<pair<size_t, size_t>> _kernarg_layout{};
 };
 
 template <>
@@ -132,6 +133,8 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
                                   uint32_t localWorkSizeZ, size_t sharedMemBytes,
                                   hipStream_t hStream, void** kernelParams, void** extra,
                                   hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags) {
+    using namespace hip_impl;
+
     auto ctx = ihipGetTlsDefaultCtx();
     hipError_t ret = hipSuccess;
 
@@ -146,19 +149,26 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
         void* config[5] = {0};
         size_t kernArgSize;
 
-        if (kernelParams != NULL) {
-            std::string name = f->_name;
-            struct ihipKernArgInfo pl = kernelArguments[name];
-            char* argBuf = (char*)malloc(pl.totalSize);
-            memset(argBuf, 0, pl.totalSize);
-            int index = 0;
-            for (int i = 0; i < pl.Size.size(); i++) {
-                memcpy(argBuf + index, kernelParams[i], pl.Size[i]);
-                index += pl.Align[i];
+        std::vector<char> tmp{};
+        if (kernelParams) {
+            if (extra) return hipErrorInvalidValue;
+
+            for (auto&& x : f->_kernarg_layout) {
+                const auto p{static_cast<const char*>(*kernelParams)};
+
+                tmp.insert(
+                    tmp.cend(),
+                    round_up_to_next_multiple_nonnegative(
+                        tmp.size(), x.second) - tmp.size(),
+                    '\0');
+                tmp.insert(tmp.cend(), p, p + x.first);
+
+                ++kernelParams;
             }
-            config[1] = (void*)argBuf;
-            kernArgSize = pl.totalSize;
-        } else if (extra != NULL) {
+            config[1] = static_cast<void*>(tmp.data());
+
+            kernArgSize = tmp.size();
+        } else if (extra) {
             memcpy(config, extra, sizeof(size_t) * 5);
             if (config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER &&
                 config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END) {
@@ -236,10 +246,6 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX,
             stopEvent->attachToCompletionFuture(&cf, hStream, hipEventTypeStopCommand);
         }
 
-
-        if (kernelParams != NULL) {
-            free(config[1]);
-        }
         ihipPostLaunchKernel(f->_name.c_str(), hStream, lp);
     }
 
@@ -461,7 +467,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const ch
     // TODO: refactor the whole ihipThisThat, which is a mess and yields the
     //       below, due to hipFunction_t being a pointer to ihipModuleSymbol_t.
     func[0][0] = *static_cast<hipFunction_t>(
-        Kernel_descriptor{kernel_object(kernel), name});
+        Kernel_descriptor{kernel_object(kernel), name, hmod->kernargs[name]});
 
     return hipSuccess;
 }
@@ -549,6 +555,11 @@ hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) {
     (*module)->executable = get_program_state().load_executable(
                                             content.data(), content.size(), (*module)->executable,
                                             this_agent());
+    istringstream elf{content};
+    ELFIO::elfio reader;
+    if (reader.load(elf)) {
+        program_state_impl::read_kernarg_metadata(reader, (*module)->kernargs);
+    }
 
     // compute the hash of the code object
     (*module)->hash = checksum(content.length(), content.data());

projects/clr/hipamd/src/hiprtc.cpp

@@ -0,0 +1,601 @@
+/*
+Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "../include/hip/hiprtc.h"
+#include "../include/hip/hcc_detail/code_object_bundle.hpp"
+#include "../include/hip/hcc_detail/elfio/elfio.hpp"
+#include "../include/hip/hcc_detail/program_state.hpp"
+
+#include "../lpl_ca/pstreams/pstream.h"
+
+#include <hsa/hsa.h>
+
+#include <cxxabi.h>
+
+#include <algorithm>
+#include <cassert>
+#include <cstdio>
+#include <cstdlib>
+#include <experimental/filesystem>
+#include <fstream>
+#include <future>
+#include <iterator>
+#include <memory>
+#include <mutex>
+#include <stdexcept>
+#include <string>
+#include <unordered_map>
+#include <utility>
+#include <vector>
+
+#include <iostream>
+
+const char* hiprtcGetErrorString(hiprtcResult x)
+{
+    switch (x) {
+    case HIPRTC_SUCCESS:
+        return "HIPRTC_SUCCESS";
+    case HIPRTC_ERROR_OUT_OF_MEMORY:
+        return "HIPRTC_ERROR_OUT_OF_MEMORY";
+    case HIPRTC_ERROR_PROGRAM_CREATION_FAILURE:
+        return "HIPRTC_ERROR_PROGRAM_CREATION_FAILURE";
+    case HIPRTC_ERROR_INVALID_INPUT:
+        return "HIPRTC_ERROR_INVALID_INPUT";
+    case HIPRTC_ERROR_INVALID_PROGRAM:
+        return "HIPRTC_ERROR_INVALID_PROGRAM";
+    case HIPRTC_ERROR_INVALID_OPTION:
+        return "HIPRTC_ERROR_INVALID_OPTION";
+    case HIPRTC_ERROR_COMPILATION:
+        return "HIPRTC_ERROR_COMPILATION";
+    case HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE:
+        return "HIPRTC_ERROR_BUILTIN_OPERATION_FAILURE";
+    case HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION:
+        return "HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION";
+    case HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION:
+        return "HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION";
+    case HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID:
+        return "HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID";
+    case HIPRTC_ERROR_INTERNAL_ERROR:
+        return "HIPRTC_ERROR_INTERNAL_ERROR";
+    default: throw std::logic_error{"Invalid HIPRTC result."};
+    };
+}
+
+namespace
+{
+    struct Symbol {
+        std::string name;
+        ELFIO::Elf64_Addr value = 0;
+        ELFIO::Elf_Xword size = 0;
+        ELFIO::Elf_Half sect_idx = 0;
+        std::uint8_t bind = 0;
+        std::uint8_t type = 0;
+        std::uint8_t other = 0;
+    };
+
+    inline
+    Symbol read_symbol(const ELFIO::symbol_section_accessor& section,
+                    unsigned int idx) {
+        assert(idx < section.get_symbols_num());
+
+        Symbol r;
+        section.get_symbol(
+            idx, r.name, r.value, r.size, r.bind, r.type, r.sect_idx, r.other);
+
+        return r;
+    }
+} // Unnamed namespace.
+
+struct _hiprtcProgram {
+    // DATA - STATICS
+    static std::vector<std::unique_ptr<_hiprtcProgram>> programs;
+    static std::mutex mtx;
+
+    // DATA
+    std::vector<std::pair<std::string, std::string>> headers;
+    std::vector<std::pair<std::string, std::string>> names;
+    std::vector<std::string> loweredNames;
+    std::vector<char> elf;
+    std::string source;
+    std::string name;
+    std::string log;
+    bool compiled;
+
+    // STATICS
+    static
+    hiprtcResult destroy(_hiprtcProgram* p)
+    {
+        using namespace std;
+
+        lock_guard<mutex> lck{mtx};
+
+        const auto it{find_if(programs.cbegin(), programs.cend(),
+                              [=](const unique_ptr<_hiprtcProgram>& x) {
+            return x.get() == p;
+        })};
+
+        if (it == programs.cend()) return HIPRTC_ERROR_INVALID_PROGRAM;
+
+        return HIPRTC_SUCCESS;
+    }
+
+    static
+    std::string handleMangledName(std::string name)
+    {
+        using namespace std;
+
+        name = hip_impl::demangle(name.c_str());
+
+        if (name.empty()) return name;
+
+        if (name.find("void ") == 0) name.erase(0, strlen("void "));
+
+        auto dx{name.find_first_of("(<")};
+
+        if (dx == string::npos) return name;
+
+        if (name[dx] == '<') {
+            auto cnt{1u};
+            do {
+                ++dx;
+                cnt += (name[dx] == '<') ? 1 : ((name[dx] == '>') ? -1 : 0);
+            } while (cnt);
+
+            name.erase(++dx);
+        }
+        else name.erase(dx);
+
+        return name;
+    }
+
+    static
+    _hiprtcProgram* make(std::string s, std::string n,
+                         std::vector<std::pair<std::string, std::string>> h)
+    {
+        using namespace std;
+
+        unique_ptr<_hiprtcProgram> tmp{new _hiprtcProgram{move(h), {}, {}, {},
+                                                          move(s), move(n), {},
+                                                          false}};
+
+        lock_guard<mutex> lck{mtx};
+
+        programs.push_back(move(tmp));
+
+        return programs.back().get();
+    }
+
+    static
+    bool isValid(_hiprtcProgram* p) noexcept
+    {
+        return std::find_if(programs.cbegin(), programs.cend(),
+                            [=](const std::unique_ptr<_hiprtcProgram>& x) {
+            return x.get() == p;
+        }) != programs.cend();
+    }
+
+    // MANIPULATORS
+    bool compile(const std::vector<std::string>& args,
+                 const std::experimental::filesystem::path& program_folder)
+    {
+        using namespace ELFIO;
+        using namespace redi;
+        using namespace std;
+
+        ipstream compile{args.front(), args, pstreambuf::pstderr};
+
+        constexpr const auto tmp_size{1024u};
+        char tmp[tmp_size]{};
+        while (!compile.eof()) {
+            log.append(tmp, tmp + compile.readsome(tmp, tmp_size));
+        }
+
+        compile.close();
+
+        if (compile.rdbuf()->exited() &&
+            compile.rdbuf()->status() != EXIT_SUCCESS) return false;
+
+        elfio reader;
+        if (!reader.load(args.back())) return false;
+
+        const auto it{find_if(reader.sections.begin(), reader.sections.end(),
+                              [](const section* x) {
+            return x->get_name() == ".kernel";
+        })};
+
+        if (it == reader.sections.end()) return false;
+
+        hip_impl::Bundled_code_header h{(*it)->get_data()};
+
+        if (bundles(h).empty()) return false;
+
+        elf.assign(bundles(h).back().blob.cbegin(),
+                   bundles(h).back().blob.cend());
+
+        return true;
+    }
+
+    bool readLoweredNames()
+    {
+        using namespace ELFIO;
+        using namespace hip_impl;
+        using namespace std;
+
+        if (names.empty()) return true;
+
+        istringstream blob{string{elf.cbegin(), elf.cend()}};
+
+        elfio reader;
+
+        if (!reader.load(blob)) return false;
+
+        const auto it{find_if(reader.sections.begin(), reader.sections.end(),
+                              [](const section* x) {
+            return x->get_type() == SHT_SYMTAB;
+        })};
+
+        ELFIO::symbol_section_accessor symbols{reader, *it};
+
+        auto n{symbols.get_symbols_num()};
+
+        if (n < loweredNames.size()) return false;
+
+        while (n--) {
+            const auto tmp{read_symbol(symbols, n)};
+
+            auto it{find_if(names.cbegin(), names.cend(),
+                            [&](const pair<string, string>& x) {
+                return x.second == tmp.name;
+            })};
+
+            if (it == names.cend()) {
+                const auto name{handleMangledName(tmp.name)};
+
+                if (name.empty()) continue;
+
+                it = find_if(names.cbegin(), names.cend(),
+                             [&](const pair<string, string>& x) {
+                    return x.second == name;
+                });
+
+                if (it == names.cend()) continue;
+            }
+
+            loweredNames[distance(names.cbegin(), it)] = tmp.name;
+        }
+
+        return true;
+    }
+
+    // ACCESSORS
+    std::experimental::filesystem::path writeTemporaryFiles(
+        const std::experimental::filesystem::path& programFolder) const
+    {
+        using namespace std;
+
+        vector<future<void>> fut{headers.size()};
+        transform(headers.cbegin(), headers.cend(), begin(fut),
+                  [&](const pair<string, string>& x) {
+            return async([&]() {
+                ofstream h{programFolder / x.first};
+                h.write(x.second.data(), x.second.size());
+            });
+        });
+
+        auto tmp{(programFolder / name).replace_extension(".cpp")};
+        ofstream{tmp}.write(source.data(), source.size());
+
+        return tmp;
+    }
+
+
+};
+std::vector<std::unique_ptr<_hiprtcProgram>> _hiprtcProgram::programs{};
+std::mutex _hiprtcProgram::mtx{};
+
+namespace
+{
+    inline
+    bool isValidProgram(const hiprtcProgram p)
+    {
+        if (!p) return false;
+
+        std::lock_guard<std::mutex> lck{_hiprtcProgram::mtx};
+
+        return _hiprtcProgram::isValid(p);
+    }
+} // Unnamed namespace.
+
+hiprtcResult hiprtcAddNameExpression(hiprtcProgram p, const char* n)
+{
+    if (!n) return HIPRTC_ERROR_INVALID_INPUT;
+    if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
+    if (p->compiled) return HIPRTC_ERROR_NO_NAME_EXPRESSIONS_AFTER_COMPILATION;
+
+    const auto id{p->names.size()};
+
+    p->names.emplace_back(n, n);
+    p->loweredNames.emplace_back();
+
+    if (p->names.back().second.back() == ')') {
+        p->names.back().second.pop_back();
+        p->names.back().second.erase(0, p->names.back().second.find('('));
+    }
+    if (p->names.back().second.front() == '&') {
+        p->names.back().second.erase(0, 1);
+    }
+
+    const auto var{"__hiprtc_" + std::to_string(id)};
+    p->source.append("\nextern \"C\" constexpr auto " + var + " = " + n + ';');
+
+    return HIPRTC_SUCCESS;
+}
+
+namespace
+{
+    class Unique_temporary_path {
+        // DATA
+        std::experimental::filesystem::path path_{};
+    public:
+        // CREATORS
+        Unique_temporary_path() : path_{std::tmpnam(nullptr)}
+        {
+            while (std::experimental::filesystem::exists(path_)) {
+                path_ = std::tmpnam(nullptr);
+            }
+        }
+        Unique_temporary_path(const std::string& extension)
+            : Unique_temporary_path{}
+        {
+            path_.replace_extension(extension);
+        }
+
+        Unique_temporary_path(const Unique_temporary_path&) = default;
+        Unique_temporary_path(Unique_temporary_path&&) = default;
+
+        ~Unique_temporary_path() noexcept
+        {
+            std::experimental::filesystem::remove_all(path_);
+        }
+
+        // MANIPULATORS
+        Unique_temporary_path& operator=(
+            const Unique_temporary_path&) = default;
+        Unique_temporary_path& operator=(Unique_temporary_path&&) = default;
+
+        // ACCESSORS
+        const std::experimental::filesystem::path& path() const noexcept
+        {
+            return path_;
+        }
+    };
+} // Unnamed namespace.
+
+namespace hip_impl
+{
+    inline
+    std::string demangle(const char* x)
+    {
+        if (!x) return {};
+
+        int s{};
+        std::unique_ptr<char, decltype(std::free)*> tmp{
+            abi::__cxa_demangle(x, nullptr, nullptr, &s), std::free};
+
+        if (s != 0) return {};
+
+        return tmp.get();
+    }
+} // Namespace hip_impl.
+
+namespace
+{
+    const std::string& defaultTarget()
+    {
+        using namespace std;
+
+        static string r{"gfx900"};
+        static once_flag f{};
+
+        call_once(f, []() {
+            static hsa_agent_t a{};
+            hsa_iterate_agents([](hsa_agent_t x, void*) {
+                hsa_device_type_t t{};
+                hsa_agent_get_info(x, HSA_AGENT_INFO_DEVICE, &t);
+
+                if (t != HSA_DEVICE_TYPE_GPU) return HSA_STATUS_SUCCESS;
+
+                a = x;
+
+                return HSA_STATUS_INFO_BREAK;
+            }, nullptr);
+
+            if (!a.handle) return;
+
+            hsa_agent_iterate_isas(a, [](hsa_isa_t x, void*){
+                uint32_t n{};
+                hsa_isa_get_info_alt(x, HSA_ISA_INFO_NAME_LENGTH, &n);
+
+                if (n == 0) return HSA_STATUS_SUCCESS;
+
+                r.resize(n);
+                hsa_isa_get_info_alt(x, HSA_ISA_INFO_NAME, &r[0]);
+
+                r.erase(0, r.find("gfx"));
+
+                return HSA_STATUS_INFO_BREAK;
+            }, nullptr);
+        });
+
+        return r;
+    }
+
+    inline
+    void handleTarget(std::vector<std::string>& args)
+    {
+        using namespace std;
+
+        bool hasTarget{false};
+        for (auto&& x : args) {
+            const auto dx{x.find("--gpu-architecture")};
+            const auto dy{(dx == string::npos) ? x.find("-arch")
+                                               : string::npos};
+
+            if (dx == dy) continue;
+
+            x.replace(0, x.find('=', min(dx, dy)), "--amdgpu-target");
+            hasTarget = true;
+
+            break;
+        }
+        if (!hasTarget) args.push_back("--amdgpu-target=" + defaultTarget());
+    }
+} // Unnamed namespace.
+
+hiprtcResult hiprtcCompileProgram(hiprtcProgram p, int n, const char** o)
+{
+    using namespace std;
+
+    if (n && !o) return HIPRTC_ERROR_INVALID_INPUT;
+    if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
+    if (p->compiled) return HIPRTC_ERROR_COMPILATION;
+
+    static const string hipcc{
+        getenv("HIP_PATH") ? (getenv("HIP_PATH") + string{"/bin/hipcc"})
+                           : "/opt/rocm/bin/hipcc"};
+
+    if (!experimental::filesystem::exists(hipcc)) {
+        return HIPRTC_ERROR_INTERNAL_ERROR;
+    }
+
+    Unique_temporary_path tmp{};
+    experimental::filesystem::create_directory(tmp.path());
+
+    const auto src{p->writeTemporaryFiles(tmp.path())};
+
+    vector<string> args{hipcc, "-shared"};
+    if (n) args.insert(args.cend(), o, o + n);
+
+    handleTarget(args);
+
+    args.emplace_back(src);
+    args.emplace_back("-o");
+    args.emplace_back(tmp.path() / "hiprtc.out");
+
+    if (!p->compile(args, tmp.path())) return HIPRTC_ERROR_INTERNAL_ERROR;
+    if (!p->readLoweredNames()) return HIPRTC_ERROR_INTERNAL_ERROR;
+
+    p->compiled = true;
+
+    return HIPRTC_SUCCESS;
+}
+
+hiprtcResult hiprtcCreateProgram(hiprtcProgram* p, const char* src,
+                                 const char* name, int n, const char** hdrs,
+                                 const char** incs)
+{
+    using namespace std;
+
+    if (!p) return HIPRTC_ERROR_INVALID_PROGRAM;
+    if (n < 0) return HIPRTC_ERROR_INVALID_INPUT;
+    if (n && (!hdrs || !incs)) return HIPRTC_ERROR_INVALID_INPUT;
+
+    vector<pair<string, string>> h;
+    for (auto i = 0; i != n; ++i) h.emplace_back(incs[i], hdrs[i]);
+
+    *p = _hiprtcProgram::make(src, name ? name : "default_name", move(h));
+
+    return HIPRTC_SUCCESS;
+}
+
+hiprtcResult hiprtcDestroyProgram(hiprtcProgram* p)
+{
+    if (!p) return HIPRTC_SUCCESS;
+
+    return _hiprtcProgram::destroy(*p);
+}
+
+hiprtcResult hiprtcGetLoweredName(hiprtcProgram p, const char* n,
+                                  const char** ln)
+{
+    using namespace std;
+
+    if (!n || !ln) return HIPRTC_ERROR_INVALID_INPUT;
+    if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
+    if (!p->compiled) return HIPRTC_ERROR_NO_LOWERED_NAMES_BEFORE_COMPILATION;
+
+    const auto it{find_if(p->names.cbegin(), p->names.cend(),
+                          [=](const pair<string, string>& x) {
+        return x.first == n;
+    })};
+
+    if (it == p->names.cend()) return HIPRTC_ERROR_NAME_EXPRESSION_NOT_VALID;
+
+    *ln = p->loweredNames[distance(p->names.cbegin(), it)].c_str();
+
+    return HIPRTC_SUCCESS;
+}
+
+hiprtcResult hiprtcGetProgramLog(hiprtcProgram p, char* l)
+{
+    if (!l) return HIPRTC_ERROR_INVALID_INPUT;
+    if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
+    if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM;
+
+    l = std::copy_n(p->log.data(), p->log.size(), l);
+    *l = '\0';
+
+    return HIPRTC_SUCCESS;
+}
+
+hiprtcResult hiprtcGetProgramLogSize(hiprtcProgram p, std::size_t* sz)
+{
+    if (!sz) return HIPRTC_ERROR_INVALID_INPUT;
+    if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
+    if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM;
+
+    *sz = p->log.empty() ? 0 : p->log.size() + 1;
+
+    return HIPRTC_SUCCESS;
+}
+
+hiprtcResult hiprtcGetCode(hiprtcProgram p, char* c)
+{
+    if (!c) return HIPRTC_ERROR_INVALID_INPUT;
+    if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
+    if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM;
+
+    std::copy_n(p->elf.data(), p->elf.size(), c);
+
+    return HIPRTC_SUCCESS;
+}
+
+hiprtcResult hiprtcGetCodeSize(hiprtcProgram p, std::size_t* sz)
+{
+    if (!sz) return HIPRTC_ERROR_INVALID_INPUT;
+    if (!isValidProgram(p)) return HIPRTC_ERROR_INVALID_PROGRAM;
+    if (!p->compiled) return HIPRTC_ERROR_INVALID_PROGRAM;
+
+    *sz = p->elf.size();
+
+    return HIPRTC_SUCCESS;
+}

projects/clr/hipamd/src/program_state.inl

@@ -66,10 +66,17 @@ class Kernel_descriptor {
     std::uint64_t kernel_object_{};
     amd_kernel_code_t const* kernel_header_{nullptr};
     std::string name_{};
+    std::vector<std::pair<std::size_t, std::size_t>> kernarg_layout_{};
 public:
     Kernel_descriptor() = default;
-    Kernel_descriptor(std::uint64_t kernel_object, const std::string& name)
-        : kernel_object_{kernel_object}, name_{name}
+    Kernel_descriptor(
+        std::uint64_t kernel_object,
+        const std::string& name,
+        std::vector<std::pair<std::size_t, std::size_t>> kernarg_layout = {})
+        :
+        kernel_object_{kernel_object},
+        name_{name},
+        kernarg_layout_{std::move(kernarg_layout)}
     {
         bool supported{false};
         std::uint16_t min_v{UINT16_MAX};
@@ -548,6 +555,7 @@ public:
         return functions[agent].second;
     }
 
+    static
     std::size_t parse_args(
             const std::string& metadata,
             std::size_t f,
@@ -576,6 +584,7 @@ public:
         } while (true);
     }
 
+    static
     void read_kernarg_metadata(
             ELFIO::elfio& reader,
             std::unordered_map<