From e53df57ffe7bbee5267aeffd49a5e2a401b97753 Mon Sep 17 00:00:00 2001 From: taosang2 Date: Thu, 11 Apr 2024 08:57:47 -0400 Subject: [PATCH] SWDEV-433371 - use comgr to unbundle code objects 1.Make runtime use comgr to unbundle code objects 2.Support compressed/uncompressed modes 3.Remove HIP_USE_RUNTIME_UNBUNDLER and HIPRTC_USE_RUNTIME_UNBUNDLER to simplify logics 4.Add comgr wrapper for amd_comgr_action_info_set_bundle_entry_ids() Change-Id: Ic41b1ad1b64cca1e31986437983a5146d52a7329 --- hipamd/src/hip_code_object.cpp | 480 +++++++++++++++++++----- hipamd/src/hip_code_object.hpp | 31 +- hipamd/src/hip_fatbin.cpp | 313 ++++----------- hipamd/src/hip_fatbin.hpp | 1 - hipamd/src/hip_module.cpp | 4 - hipamd/src/hiprtc/hiprtcComgrHelper.cpp | 33 -- hipamd/src/hiprtc/hiprtcComgrHelper.hpp | 2 - hipamd/src/hiprtc/hiprtcInternal.cpp | 24 +- rocclr/device/comgrctx.cpp | 1 + rocclr/device/comgrctx.hpp | 8 +- rocclr/utils/flags.hpp | 4 - 11 files changed, 474 insertions(+), 427 deletions(-) diff --git a/hipamd/src/hip_code_object.cpp b/hipamd/src/hip_code_object.cpp index 6b3a6d08f7..2b834f3cc2 100644 --- a/hipamd/src/hip_code_object.cpp +++ b/hipamd/src/hip_code_object.cpp @@ -30,21 +30,34 @@ THE SOFTWARE. #include "hip_internal.hpp" #include "platform/program.hpp" #include +#include "comgrctx.hpp" + namespace hip { hipError_t ihipFree(void* ptr); // forward declaration of methods required for managed variables hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align = 0); namespace { -constexpr char kOffloadBundleMagicStr[] = "__CLANG_OFFLOAD_BUNDLE__"; +// In uncompressed mode +constexpr char kOffloadBundleUncompressedMagicStr[] = "__CLANG_OFFLOAD_BUNDLE__"; +static constexpr size_t kOffloadBundleUncompressedMagicStrSize = + sizeof(kOffloadBundleUncompressedMagicStr); + +//In compressed mode +constexpr char kOffloadBundleCompressedMagicStr[] = "CCOB"; +static constexpr size_t kOffloadBundleCompressedMagicStrSize = + sizeof(kOffloadBundleCompressedMagicStr); + constexpr char kOffloadKindHip[] = "hip"; constexpr char kOffloadKindHipv4[] = "hipv4"; constexpr char kOffloadKindHcc[] = "hcc"; constexpr char kAmdgcnTargetTriple[] = "amdgcn-amd-amdhsa-"; - +constexpr char kHipFatBinName[] = "hipfatbin"; +constexpr char kHipFatBinName_[] = "hipfatbin-"; +constexpr char kOffloadKindHipv4_[] = "hipv4-"; // bundled code objects need the prefix +constexpr char kOffloadHipV4FatBinName_[] = "hipfatbin-hipv4-"; // ClangOFFLOADBundle info. -static constexpr size_t kOffloadBundleMagicStrSize = sizeof(kOffloadBundleMagicStr); -// Clang Offload bundler description & Header. +// Clang Offload bundler description & Header in uncompressed mode. struct __ClangOffloadBundleInfo { uint64_t offset; uint64_t size; @@ -52,16 +65,37 @@ struct __ClangOffloadBundleInfo { const char bundleEntryId[1]; }; -struct __ClangOffloadBundleHeader { - const char magic[kOffloadBundleMagicStrSize - 1]; +struct __ClangOffloadBundleUncompressedHeader { + const char magic[kOffloadBundleUncompressedMagicStrSize - 1]; uint64_t numOfCodeObjects; __ClangOffloadBundleInfo desc[1]; }; + +struct __ClangOffloadBundleCompressedHeader { + const char magic[kOffloadBundleCompressedMagicStrSize - 1]; + uint16_t versionNumber; + uint16_t compressionMethod; + uint32_t totalSize; + uint32_t uncompressedBinarySize; + uint64_t Hash; + const char compressedBinarydesc[1]; +}; } // namespace -bool CodeObject::IsClangOffloadMagicBundle(const void* data) { - std::string magic(reinterpret_cast(data), kOffloadBundleMagicStrSize - 1); - return magic.compare(kOffloadBundleMagicStr) ? false : true; +bool CodeObject::IsClangOffloadMagicBundle(const void* data, bool &isCompressed) { + std::string magic(reinterpret_cast(data), + kOffloadBundleUncompressedMagicStrSize - 1); + if (!magic.compare(kOffloadBundleUncompressedMagicStr)) { + isCompressed = false; + return true; + } + std::string magic1(reinterpret_cast(data), + kOffloadBundleCompressedMagicStrSize - 1); + if (!magic1.compare(kOffloadBundleCompressedMagicStr)) { + isCompressed = true; + return true; + } + return false; } uint64_t CodeObject::ElfSize(const void* emi) { return amd::Elf::getElfSize(emi); } @@ -356,7 +390,7 @@ static bool consume(std::string& input, std::string consume_) { // Trim String till character, will be used to get gpuname // example: input is gfx908:sram-ecc+ and trim char is : -// input will become sram-ecc+. +// input will become :sram-ecc+ static std::string trimName(std::string& input, char trim) { auto pos_ = input.find(trim); auto res = input; @@ -369,6 +403,18 @@ static std::string trimName(std::string& input, char trim) { return res; } +// Trim String till character, will be used to get bundle entry ID. +// example: input is amdgcn-amd-amdhsa--gfx1035.bc and trim char is . +// input will become amdgcn-amd-amdhsa--gfx1035 +static bool trimNameTail(std::string& input, char trim) { + auto pos_ = input.rfind(trim); + if (pos_ == std::string::npos) { + return false; + } + input = input.substr(0, pos_); + return true; +} + static char getFeatureValue(std::string& input, std::string feature) { char res = ' '; if (consume(input, std::move(feature))) { @@ -447,111 +493,353 @@ static bool isCodeObjectCompatibleWithDevice(std::string co_triple_target_id, return true; } -// This will be moved to COMGR eventually -hipError_t CodeObject::ExtractCodeObjectFromFile( - amd::Os::FileDesc fdesc, size_t fsize, const void** image, - const std::vector& device_names, - std::vector>& code_objs) { - if (!amd::Os::isValidFileDesc(fdesc)) { - return hipErrorFileNotFound; +size_t CodeObject::getFatbinSize(const void* data, const bool isCompressed) { + if (isCompressed) { + const auto obheader = reinterpret_cast(data); + return obheader->totalSize; + } else { + const auto obheader = reinterpret_cast(data); + const __ClangOffloadBundleInfo* desc = &obheader->desc[0]; + uint64_t i = 0; + while (++i < obheader->numOfCodeObjects) { + desc = reinterpret_cast( + reinterpret_cast(&desc->bundleEntryId[0]) + desc->bundleEntryIdSize); + } + return desc->offset + desc->size; } - - // Map the file to memory, with offset 0. - // file will be unmapped in ModuleUnload - // const void* image = nullptr; - if (!amd::Os::MemoryMapFileDesc(fdesc, fsize, 0, image)) { - return hipErrorInvalidValue; - } - - // retrieve code_objs{binary_image, binary_size} for devices - return extractCodeObjectFromFatBinary(*image, device_names, code_objs); } -// This will be moved to COMGR eventually -hipError_t CodeObject::ExtractCodeObjectFromMemory( - const void* data, const std::vector& device_names, - std::vector>& code_objs, std::string& uri) { - // Get the URI from memory - if (!amd::Os::GetURIFromMemory(data, 0, uri)) { - return hipErrorInvalidValue; - } - - return extractCodeObjectFromFatBinary(data, device_names, code_objs); -} - -// This will be moved to COMGR eventually +/** + * @brief Extract code object from fatbin using comgr + * + * @param[in] data the bundle data(fatbin or loaded module data) + * @param[in] size the size of the bundle data + * @param[in] agent_triple_target_ids isa names of concerned devices + * @param[out] code_objs the buffer address and size pairs of extracted code objects of + * concerned devices + * Returned error code + * + * @return #hipSuccess, #hipErrorInvalidKernelFile, #hipErrorInvalidValue, #hipErrorNoBinaryForGpu + * + * @see FatBinaryInfo::ExtractFatBinaryUsingCOMGR + */ hipError_t CodeObject::extractCodeObjectFromFatBinary( - const void* data, const std::vector& agent_triple_target_ids, + const void* data, size_t size, const std::vector& agent_triple_target_ids, std::vector>& code_objs) { - std::string magic((const char*)data, kOffloadBundleMagicStrSize); - if (magic.compare(kOffloadBundleMagicStr)) { + hipError_t hipStatus = hipSuccess; + amd_comgr_status_t comgrStatus = AMD_COMGR_STATUS_SUCCESS; + + const size_t num_devices = agent_triple_target_ids.size(); + size_t num_code_objs = num_devices; + bool isCompressed = false; + if (!IsClangOffloadMagicBundle(data, isCompressed)) { + LogPrintfInfo("IsClangOffloadMagicBundle(%p) return false", data); + // hipModuleLoadData() will possibly call here return hipErrorInvalidKernelFile; } - // Initialize Code objects - code_objs.reserve(agent_triple_target_ids.size()); - for (size_t i = 0; i < agent_triple_target_ids.size(); i++) { - code_objs.push_back(std::make_pair(nullptr, 0)); + if (size == 0) size = getFatbinSize(data, isCompressed); + + amd_comgr_data_t dataCodeObj{0}; + amd_comgr_data_set_t dataSetBundled{0}; + amd_comgr_data_set_t dataSetUnbundled{0}; + amd_comgr_action_info_t actionInfoUnbundle{0}; + amd_comgr_data_t item{0}; + + + std::set devicesSet{}; // To make sure device is unique + std::vector bundleEntryIDs{}; + static const std::string hipv4 = kOffloadKindHipv4_; // bundled code objects need the prefix + for (size_t i = 0; i < num_devices; i++) { + devicesSet.insert(hipv4 + agent_triple_target_ids[i]); } - const auto obheader = reinterpret_cast(data); - const auto* desc = &obheader->desc[0]; - size_t num_code_objs = code_objs.size(); - for (uint64_t i = 0; i < obheader->numOfCodeObjects; ++i, - desc = reinterpret_cast( - reinterpret_cast(&desc->bundleEntryId[0]) + - desc->bundleEntryIdSize)) { - const void* image = - reinterpret_cast(reinterpret_cast(obheader) + desc->offset); - const size_t image_size = desc->size; + for (auto& device : devicesSet) { + bundleEntryIDs.push_back(device.c_str()); + } - if (num_code_objs == 0) break; - std::string bundleEntryId{desc->bundleEntryId, desc->bundleEntryIdSize}; + do { + // Create Bundled dataset + comgrStatus = amd::Comgr::create_data_set(&dataSetBundled); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::create_data_set() failed with status 0x%xh", comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } - std::string co_triple_target_id; - if (!getTripleTargetID(bundleEntryId, image, co_triple_target_id)) continue; + // CodeObject + comgrStatus = amd::Comgr::create_data(AMD_COMGR_DATA_KIND_OBJ_BUNDLE, &dataCodeObj); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError( + "amd::Comgr::create_data(AMD_COMGR_DATA_KIND_OBJ_BUNDLE) failed with status 0x%xh", + comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } - for (size_t dev = 0; dev < agent_triple_target_ids.size(); ++dev) { - if (code_objs[dev].first) continue; - if (isCodeObjectCompatibleWithDevice(co_triple_target_id, agent_triple_target_ids[dev])) { - code_objs[dev] = std::make_pair(image, image_size); - --num_code_objs; + comgrStatus = amd::Comgr::set_data(dataCodeObj, size, static_cast(data)); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::set_data(size=%zu, data=%p) failed with status 0x%xh", size, data, + comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + + comgrStatus = amd::Comgr::set_data_name(dataCodeObj, kHipFatBinName); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError( + "amd::Comgr::set_data_name("") failed with status 0x%xh", comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + comgrStatus = amd::Comgr::data_set_add(dataSetBundled, dataCodeObj); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::data_set_add() failed with status 0x%xh", comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + // Set up ActionInfo + comgrStatus = amd::Comgr::create_action_info(&actionInfoUnbundle); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::create_action_info() failed with status 0x%xh", comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + + comgrStatus = amd::Comgr::action_info_set_language(actionInfoUnbundle, AMD_COMGR_LANGUAGE_HIP); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::action_info_set_language(HIP) failed with status 0x%xh", + comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + + comgrStatus = amd::Comgr::action_info_set_bundle_entry_ids( + actionInfoUnbundle, bundleEntryIDs.data(), bundleEntryIDs.size()); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::action_info_set_bundle_entry_ids(%p, %zu) failed with status 0x%xh", + bundleEntryIDs.data(), bundleEntryIDs.size(), comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + + // Unbundle + comgrStatus = amd::Comgr::create_data_set(&dataSetUnbundled); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::create_data_set(&dataSetUnbundled) failed with status 0x%xh", + comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + comgrStatus = amd::Comgr::do_action(AMD_COMGR_ACTION_UNBUNDLE, actionInfoUnbundle, dataSetBundled, + dataSetUnbundled); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::do_action(AMD_COMGR_ACTION_UNBUNDLE) failed with status 0x%xh", + comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + + // Check CodeObject count + size_t count = 0; + comgrStatus = + amd::Comgr::action_data_count(dataSetUnbundled, AMD_COMGR_DATA_KIND_EXECUTABLE, &count); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::action_data_count() failed with status 0x%xh", comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + + // Initialize Code objects + code_objs.reserve(num_code_objs); + for (size_t i = 0; i < num_code_objs; i++) { + code_objs.push_back(std::make_pair(nullptr, 0)); + } + + for (size_t i = 0; i < count; i++) { + if (num_code_objs == 0) break; + + size_t itemSize = 0; + comgrStatus = amd::Comgr::action_data_get_data(dataSetUnbundled, + AMD_COMGR_DATA_KIND_EXECUTABLE, i, &item); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::action_data_get_data(%zu/%zu) failed with 0x%xh", i, count, + comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + + comgrStatus = amd::Comgr::get_data_name(item, &itemSize, nullptr); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::get_data_name(%zu/%zu) failed with 0x%xh", i, count, + comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + std::string bundleEntryId(itemSize, 0); + comgrStatus = amd::Comgr::get_data_name(item, &itemSize, bundleEntryId.data()); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::get_data_name(%zu/%zu, %d) failed with 0x%xh", i, count, + itemSize, comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + // Remove bundleEntryId_ + if (!consume(bundleEntryId, kOffloadHipV4FatBinName_)) { + // This is behavour in comgr unbundling which is subject to change. + // So just give info. + LogPrintfInfo("bundleEntryId=%s isn't prefixed with %s", bundleEntryId.c_str(), + kOffloadHipV4FatBinName_); + } + trimNameTail(bundleEntryId, '.'); // Remove .fileExtention + + char* itemData = nullptr; + for (size_t dev = 0; dev < num_devices; ++dev) { + if (code_objs[dev].first) continue; + //LogPrintfError("agent_triple_target_ids[%zu]=%s, bundleEntryId=%s", dev, + // agent_triple_target_ids[dev].c_str(), bundleEntryId.c_str()); + + if (bundleEntryId == agent_triple_target_ids[dev]) { + if (itemData == nullptr) { + itemSize = 0; + comgrStatus = amd::Comgr::get_data(item, &itemSize, nullptr); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::get_data(%zu/%zu) failed with 0x%xh", i, count, + comgrStatus); + hipStatus = hipErrorInvalidValue; + break; + } + + if (itemSize == 0) { + // If there isn't a code object for this device, + // amd::Comgr::do_action(AMD_COMGR_ACTION_UNBUNDLE) still returns item with + // valid name but no data. We need continue searching for other devices + LogPrintfInfo( + "amd::Comgr::get_data() return 0 size for agent_triple_target_ids[%zu]=%s", + dev, agent_triple_target_ids[dev].c_str()); + continue; + } + + // itemData should be deleted in fatbin's destructor + itemData = new char[itemSize]; + if (itemData == nullptr) { + LogError("no enough memory"); + hipStatus = hipErrorOutOfMemory; + break; + } + comgrStatus = amd::Comgr::get_data(item, &itemSize, itemData); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::get_data(%zu/%zu, %d) failed with 0x%xh", i, count, + itemSize, comgrStatus); + hipStatus = hipErrorInvalidValue; + delete []itemData; + itemData = nullptr; + break; + } + + } + code_objs[dev] = std::make_pair(reinterpret_cast(itemData), itemSize); + --num_code_objs; + LogPrintfInfo( + "Found agent_triple_target_ids[%zu]=%s: item: Data=%p(%s), " + "Size=%zu, num_code_objs=%zu", + dev, agent_triple_target_ids[dev].c_str(), itemData, + isCompressed ? "compressed" : "uncompressed", itemSize, num_code_objs); + } + } + + comgrStatus = amd::Comgr::release_data(item); + item.handle = 0; + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::release_data(item) failed with status 0x%xh", comgrStatus); + hipStatus = hipErrorInvalidValue; + } + if (hipStatus != hipSuccess) break; + } + } while(0); + + if (hipStatus == hipSuccess && num_code_objs != 0) { + hipStatus = hipErrorNoBinaryForGpu; + + // Leave it for debug purpose in uncompressed mode. + if (!isCompressed) { + LogPrintfError("%s", + "hipErrorNoBinaryForGpu: Unable to find code object for all current devices!"); + LogPrintfError("%s", " Devices:"); + for (size_t i = 0; i < agent_triple_target_ids.size(); i++) { + LogPrintfError(" %s - [%s]", agent_triple_target_ids[i].c_str(), + ((code_objs[i].first) ? "Found" : "Not Found")); + } + const auto obheader = reinterpret_cast(data); + const auto* desc = &obheader->desc[0]; + LogPrintfError("%s", " Bundled Code Objects:"); + for (uint64_t i = 0; i < obheader->numOfCodeObjects; ++i, + desc = reinterpret_cast( + reinterpret_cast(&desc->bundleEntryId[0]) + + desc->bundleEntryIdSize)) { + std::string bundleEntryId{desc->bundleEntryId, desc->bundleEntryIdSize}; + const void* image = + reinterpret_cast(reinterpret_cast(obheader) + desc->offset); + + std::string co_triple_target_id; + bool valid_co = getTripleTargetID(bundleEntryId, image, co_triple_target_id); + + if (valid_co) { + LogPrintfError(" %s - [Code object targetID is %s]", bundleEntryId.c_str(), + co_triple_target_id.c_str()); + } else { + LogPrintfError(" %s - [Unsupported]", bundleEntryId.c_str()); + } } } } - if (num_code_objs == 0) { - return hipSuccess; - } else { - LogPrintfError("%s", - "hipErrorNoBinaryForGpu: Unable to find code object for all current devices!"); - LogPrintfError("%s", " Devices:"); - for (size_t i = 0; i < agent_triple_target_ids.size(); i++) { - LogPrintfError(" %s - [%s]", agent_triple_target_ids[i].c_str(), - ((code_objs[i].first) ? "Found" : "Not Found")); - } - const auto obheader = reinterpret_cast(data); - const auto* desc = &obheader->desc[0]; - LogPrintfError("%s", " Bundled Code Objects:"); - for (uint64_t i = 0; i < obheader->numOfCodeObjects; ++i, - desc = reinterpret_cast( - reinterpret_cast(&desc->bundleEntryId[0]) + - desc->bundleEntryIdSize)) { - std::string bundleEntryId{desc->bundleEntryId, desc->bundleEntryIdSize}; - const void* image = - reinterpret_cast(reinterpret_cast(obheader) + desc->offset); - std::string co_triple_target_id; - bool valid_co = getTripleTargetID(bundleEntryId, image, co_triple_target_id); - - if (valid_co) { - LogPrintfError(" %s - [Code object targetID is %s]", bundleEntryId.c_str(), - co_triple_target_id.c_str()); - } else { - LogPrintfError(" %s - [Unsupported]", bundleEntryId.c_str()); - } + // Cleanup + if (actionInfoUnbundle.handle) { + comgrStatus = amd::Comgr::destroy_action_info(actionInfoUnbundle); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::destroy_action_info(actionInfoUnbundle) failed with status 0x%xh", + comgrStatus); + hipStatus = hipErrorInvalidValue; } - return hipErrorNoBinaryForGpu; } + if (dataSetBundled.handle) { + comgrStatus = amd::Comgr::destroy_data_set(dataSetBundled); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::destroy_data_set(dataSetBundled) failed with status 0x%xh", + comgrStatus); + hipStatus = hipErrorInvalidValue; + } + } + + if (dataSetUnbundled.handle) { + comgrStatus = amd::Comgr::destroy_data_set(dataSetUnbundled); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::destroy_data_set(dataSetUnbundled) failed with status 0x%xh", + comgrStatus); + hipStatus = hipErrorInvalidValue; + } + } + + if (dataCodeObj.handle) { + comgrStatus = amd::Comgr::release_data(dataCodeObj); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::release_data(dataCodeObj) failed with status 0x%xh", + comgrStatus); + hipStatus = hipErrorInvalidValue; + } + } + + if (item.handle) { + comgrStatus = amd::Comgr::release_data(item); + if (comgrStatus != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("amd::Comgr::release_data(item) failed with status 0x%xh", + comgrStatus); + hipStatus = hipErrorInvalidValue; + } + } + + return hipStatus; } hipError_t DynCO::loadCodeObject(const char* fname, const void* image) { diff --git a/hipamd/src/hip_code_object.hpp b/hipamd/src/hip_code_object.hpp index c83ec37b24..d5b6df0a8c 100644 --- a/hipamd/src/hip_code_object.hpp +++ b/hipamd/src/hip_code_object.hpp @@ -48,31 +48,22 @@ class CodeObject { size_t binary_size); static hipError_t build_module(hipModule_t hmod, const std::vector& devices); - // Given an file desc and file size, extracts to code object for corresponding devices, - // return code_objs{binary_ptr, binary_size}, which could be used to determine foffset - static hipError_t ExtractCodeObjectFromFile(amd::Os::FileDesc fdesc, size_t fsize, - const void ** image, const std::vector& device_names, - std::vector>& code_objs); - - // Given an ptr to memory, extracts to code object for corresponding devices, - // returns code_objs{binary_ptr, binary_size} and uniform resource indicator - static hipError_t ExtractCodeObjectFromMemory(const void* data, - const std::vector& device_names, - std::vector>& code_objs, - std::string& uri); - static uint64_t ElfSize(const void* emi); - static bool IsClangOffloadMagicBundle(const void* data); + static bool IsClangOffloadMagicBundle(const void* data, bool& isCompressed); + + // Given an ptr to image or file, extracts to code object + // for corresponding devices + static hipError_t extractCodeObjectFromFatBinary( + const void*, size_t, const std::vector&, + std::vector>&); + + // Return size of fat bin + static size_t getFatbinSize(const void* data, const bool isCompressed = false); protected: - //Given an ptr to image or file, extracts to code object - //for corresponding devices - static hipError_t extractCodeObjectFromFatBinary(const void*, - const std::vector&, - std::vector>&); - CodeObject() {} + private: friend const std::vector& modules(); }; diff --git a/hipamd/src/hip_fatbin.cpp b/hipamd/src/hip_fatbin.cpp index 562f193861..a66b1106dc 100644 --- a/hipamd/src/hip_fatbin.cpp +++ b/hipamd/src/hip_fatbin.cpp @@ -50,57 +50,46 @@ FatBinaryInfo::FatBinaryInfo(const char* fname, const void* image) : fdesc_(amd: } FatBinaryInfo::~FatBinaryInfo() { - + // Different devices in the same model have the same binary_image_ + std::set toDelete; // Release per device fat bin info. for (auto* fbd: fatbin_dev_info_) { if (fbd != nullptr) { + if (fbd->binary_image_ && fbd->binary_offset_ == 0 && fbd->binary_image_ != image_) { + // binary_image_ was allocated in CodeObject::extractCodeObjectFromFatBinary + toDelete.insert(fbd->binary_image_); + } delete fbd; } } - if (!HIP_USE_RUNTIME_UNBUNDLER) { - // Using COMGR Unbundler - if (ufd_ && amd::Os::isValidFileDesc(ufd_->fdesc_)) { - // Check for ufd_ != nullptr, since sometimes, we never create unique_file_desc. - if (ufd_->fsize_ && image_mapped_ - && !amd::Os::MemoryUnmapFile(image_, ufd_->fsize_)) { - LogPrintfError("Cannot unmap file for fdesc: %d fsize: %d", ufd_->fdesc_, ufd_->fsize_); - assert(false); - } - if (!PlatformState::instance().CloseUniqueFileHandle(ufd_)) { - LogPrintfError("Cannot close file for fdesc: %d", ufd_->fdesc_); - assert(false); - } + for (auto itemData : toDelete) { + LogPrintfInfo("~FatBinaryInfo(%p) will delete binary_image_ %p", this, itemData); + delete[] reinterpret_cast(itemData); + } + + // Using COMGR Unbundler + if (ufd_ && amd::Os::isValidFileDesc(ufd_->fdesc_)) { + // Check for ufd_ != nullptr, since sometimes, we never create unique_file_desc. + if (ufd_->fsize_ && image_mapped_ + && !amd::Os::MemoryUnmapFile(image_, ufd_->fsize_)) { + LogPrintfError("Cannot unmap file for fdesc: %d fsize: %d", ufd_->fdesc_, ufd_->fsize_); + assert(false); } - - fname_ = std::string(); - fdesc_ = amd::Os::FDescInit(); - fsize_ = 0; - image_ = nullptr; - uri_ = std::string(); - - if (0 == PlatformState::instance().UfdMapSize()) { - LogError("All Unique FDs are closed"); + if (!PlatformState::instance().CloseUniqueFileHandle(ufd_)) { + LogPrintfError("Cannot close file for fdesc: %d", ufd_->fdesc_); + assert(false); } + } - } else { - // Using Runtime Unbundler - if (amd::Os::isValidFileDesc(fdesc_)) { - if (fsize_ && !amd::Os::MemoryUnmapFile(image_, fsize_)) { - LogPrintfError("Cannot unmap file for fdesc: %d fsize: %d", fdesc_, fsize_); - assert(false); - } - if (!amd::Os::CloseFileHandle(fdesc_)) { - LogPrintfError("Cannot close file for fdesc: %d", fdesc_); - assert(false); - } - } + fname_ = std::string(); + fdesc_ = amd::Os::FDescInit(); + fsize_ = 0; + image_ = nullptr; + uri_ = std::string(); - fname_ = std::string(); - fdesc_ = amd::Os::FDescInit(); - fsize_ = 0; - image_ = nullptr; - uri_ = std::string(); + if (0 == PlatformState::instance().UfdMapSize()) { + LogError("All Unique FDs are closed"); } } @@ -114,11 +103,8 @@ void ListAllDeviceWithNoCOFromBundle(const std::unordered_map& devices) { - amd_comgr_data_t data_object {0}; - amd_comgr_status_t comgr_status = AMD_COMGR_STATUS_SUCCESS; +hipError_t FatBinaryInfo::ExtractFatBinary(const std::vector& devices) { hipError_t hip_status = hipSuccess; - // If image was passed as a pointer to our hipMod* api, we can try to extract the file name // if it was mapped by the app. Otherwise use the COMGR data API. if (fname_.size() == 0) { @@ -163,107 +149,52 @@ hipError_t FatBinaryInfo::ExtractFatBinaryUsingCOMGR(const std::vector> code_objs; + // Copy device names + std::vector device_names; + device_names.reserve(devices.size()); + for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) { + device_names.push_back(devices[dev_idx]->devices()[0]->isa().isaName()); + } + hip_status = CodeObject::extractCodeObjectFromFatBinary( + image_, 0, device_names, code_objs); + if (hip_status == hipErrorNoBinaryForGpu || hip_status == hipSuccess) { + for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) { + if (code_objs[dev_idx].first) { + fatbin_dev_info_[devices[dev_idx]->deviceId()] + = new FatBinaryDeviceInfo(code_objs[dev_idx].first, code_objs[dev_idx].second, 0); - // If the image ptr is not clang offload bundle then just directly point the image. - if (!CodeObject::IsClangOffloadMagicBundle(image_)) { - for (size_t dev_idx=0; dev_idx < devices.size(); ++dev_idx) { - fatbin_dev_info_[devices[dev_idx]->deviceId()] - = new FatBinaryDeviceInfo(image_, CodeObject::ElfSize(image_), 0); - fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_ - = new amd::Program(*devices[dev_idx]->asContext()); + fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_ + = new amd::Program(*devices[dev_idx]->asContext()); + if (fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_ == NULL) { + break; + } + } + else { + // This is the case of hipErrorNoBinaryForGpu which will finally fail app + LogPrintfError("Cannot find CO in the bundle %s for ISA: %s", fname_.c_str(), + device_names[dev_idx].c_str()); + } + } + } + else if (hip_status == hipErrorInvalidKernelFile) { + hip_status = hipSuccess; + // If the image ptr is not clang offload bundle then just directly point the image. + for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) { + fatbin_dev_info_[devices[dev_idx]->deviceId()] = + new FatBinaryDeviceInfo(image_, CodeObject::ElfSize(image_), 0); + fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_ = + new amd::Program(*devices[dev_idx]->asContext()); if (fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_ == nullptr) { hip_status = hipErrorOutOfMemory; break; } } - break; } - - // Create a data object, if it fails return error - if ((comgr_status = amd_comgr_create_data(AMD_COMGR_DATA_KIND_FATBIN, &data_object)) - != AMD_COMGR_STATUS_SUCCESS) { - LogPrintfError("Creating data object failed with status %d ", comgr_status); - hip_status = hipErrorInvalidValue; - break; - } - -#if !defined(_WIN32) - // Using the file descriptor and file size, map the data object. - if (amd::Os::isValidFileDesc(fdesc_)) { - guarantee(fsize_ > 0, "Cannot have a file size of 0, fdesc: %d fname: %s", - fdesc_, fname_.c_str()); - if ((comgr_status = amd_comgr_set_data_from_file_slice(data_object, fdesc_, foffset_, - fsize_)) != AMD_COMGR_STATUS_SUCCESS) { - LogPrintfError("Setting data from file slice failed with status %d ", comgr_status); - hip_status = hipErrorInvalidValue; - break; - } - } else -#endif - if (image_ != nullptr) { - // Using the image ptr, map the data object. - if ((comgr_status = amd_comgr_set_data(data_object, 4096, - reinterpret_cast(image_))) != AMD_COMGR_STATUS_SUCCESS) { - LogPrintfError("Setting data from file slice failed with status %d ", comgr_status); - hip_status = hipErrorInvalidValue; - break; - } - } else { - guarantee(false, "Cannot have both fname_ and image_ as nullptr"); - } - - // Find the unique number of ISAs needed for this COMGR query. - std::unordered_map> unique_isa_names; - for (auto device : devices) { - std::string device_name = device->devices()[0]->isa().isaName(); - unique_isa_names.insert({device_name, std::make_pair(0,0)}); - } - - // Create a query list using COMGR info for unique ISAs. - std::vector query_list_array; - query_list_array.reserve(unique_isa_names.size()); - for (const auto &isa_name : unique_isa_names) { - auto &item = query_list_array.emplace_back(); - item.isa = isa_name.first.c_str(); - item.size = 0; - item.offset = 0; - } - - // Look up the code object info passing the query list. - if ((comgr_status = amd_comgr_lookup_code_object(data_object, query_list_array.data(), - unique_isa_names.size())) != AMD_COMGR_STATUS_SUCCESS) { - LogPrintfError("Setting data from file slice failed with status %d ", comgr_status); - hip_status = hipErrorInvalidValue; - break; - } - - for (const auto &item : query_list_array) { - auto unique_it = unique_isa_names.find(item.isa); - guarantee(unique_isa_names.cend() != unique_it, "Cannot find unique isa "); - unique_it->second = std::pair - (static_cast(item.size), - static_cast(item.offset)); - } - - for (auto device : devices) { - std::string device_name = device->devices()[0]->isa().isaName(); - auto dev_it = unique_isa_names.find(device_name); - // If the size is 0, then COMGR API could not find the CO for this GPU device/ISA - if (dev_it->second.first == 0) { - LogPrintfError("Cannot find CO in the bundle %s for ISA: %s", - fname_.c_str(), device_name.c_str()); - hip_status = hipErrorNoBinaryForGpu; - ListAllDeviceWithNoCOFromBundle(unique_isa_names); - break; - } - guarantee(unique_isa_names.cend() != dev_it, - "Cannot find the device name in the unique device name"); - fatbin_dev_info_[device->deviceId()] - = new FatBinaryDeviceInfo(reinterpret_cast
(const_cast(image_)) - + dev_it->second.second, dev_it->second.first, - dev_it->second.second); - fatbin_dev_info_[device->deviceId()]->program_ - = new amd::Program(*(device->asContext())); + else { + LogPrintfError( + "CodeObject::extractCodeObjectFromFatBinary failed with status %d\n", + hip_status); } } while(0); @@ -286,115 +217,9 @@ hipError_t FatBinaryInfo::ExtractFatBinaryUsingCOMGR(const std::vector& devices) { - if (!HIP_USE_RUNTIME_UNBUNDLER) { - return ExtractFatBinaryUsingCOMGR(devices); - } - - hipError_t hip_error = hipSuccess; - std::vector> code_objs; - - // Copy device names for Extract Code object File - std::vector device_names; - device_names.reserve(devices.size()); - for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) { - device_names.push_back(devices[dev_idx]->devices()[0]->isa().isaName()); - } - - // We are given file name, get the file desc and file size - if (fname_.size() > 0) { - // Get File Handle & size of the file. - if (!amd::Os::GetFileHandle(fname_.c_str(), &fdesc_, &fsize_)) { - return hipErrorFileNotFound; - } - if (fsize_ == 0) { - return hipErrorInvalidImage; - } - - // Extract the code object from file - hip_error = CodeObject::ExtractCodeObjectFromFile(fdesc_, fsize_, &image_, - device_names, code_objs); - - } else if (image_ != nullptr) { - // We are directly given image pointer directly, try to extract file desc & file Size - hip_error = CodeObject::ExtractCodeObjectFromMemory(image_, - device_names, code_objs, uri_); - } else { - return hipErrorInvalidValue; - } - - if (hip_error == hipErrorNoBinaryForGpu) { - if (fname_.size() > 0) { - LogPrintfError("hipErrorNoBinaryForGpu: Couldn't find binary for file: %s", fname_.c_str()); - } else { - LogPrintfError("hipErrorNoBinaryForGpu: Couldn't find binary for ptr: 0x%x", image_); - } - - // For the condition: unable to find code object for all devices, - // still extract available images to those devices owning them. - // This helps users to work with ROCm if there is any supported - // GFX on system. - for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) { - if (code_objs[dev_idx].first) { - // Calculate the offset wrt binary_image and the original image - size_t offset_l - = (reinterpret_cast
(const_cast(code_objs[dev_idx].first)) - - reinterpret_cast
(const_cast(image_))); - - fatbin_dev_info_[devices[dev_idx]->deviceId()] - = new FatBinaryDeviceInfo(code_objs[dev_idx].first, code_objs[dev_idx].second, offset_l); - - fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_ - = new amd::Program(*devices[dev_idx]->asContext()); - if (fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_ == NULL) { - break; - } - } - } - - return hip_error; - } - - if (hip_error == hipErrorInvalidKernelFile) { - for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) { - // the image type is no CLANG_OFFLOAD_BUNDLER, image for current device directly passed - fatbin_dev_info_[devices[dev_idx]->deviceId()] - = new FatBinaryDeviceInfo(image_, CodeObject::ElfSize(image_), 0); - } - } else if(hip_error == hipSuccess) { - for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) { - // Calculate the offset wrt binary_image and the original image - size_t offset_l - = (reinterpret_cast
(const_cast(code_objs[dev_idx].first)) - - reinterpret_cast
(const_cast(image_))); - - fatbin_dev_info_[devices[dev_idx]->deviceId()] - = new FatBinaryDeviceInfo(code_objs[dev_idx].first, code_objs[dev_idx].second, offset_l); - } - } - - for (size_t dev_idx = 0; dev_idx < devices.size(); ++dev_idx) { - fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_ - = new amd::Program(*devices[dev_idx]->asContext()); - if (fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_ == NULL) { - return hipErrorOutOfMemory; - } - } - - return hipSuccess; -} - hipError_t FatBinaryInfo::AddDevProgram(const int device_id) { // Device Id bounds Check DeviceIdCheck(device_id); diff --git a/hipamd/src/hip_fatbin.hpp b/hipamd/src/hip_fatbin.hpp index f9057a4b5a..ac6d64e19e 100644 --- a/hipamd/src/hip_fatbin.hpp +++ b/hipamd/src/hip_fatbin.hpp @@ -64,7 +64,6 @@ public: ~FatBinaryInfo(); // Loads Fat binary from file or image, unbundles COs for devices. - hipError_t ExtractFatBinaryUsingCOMGR(const std::vector& devices); hipError_t ExtractFatBinary(const std::vector& devices); hipError_t AddDevProgram(const int device_id); hipError_t BuildProgram(const int device_id); diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 8707486424..5d458e3f45 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -66,10 +66,6 @@ hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned HIP_RETURN(PlatformState::instance().loadModule(module, 0, image)); } -extern hipError_t __hipExtractCodeObjectFromFatBinary( - const void* data, const std::vector& devices, - std::vector>& code_objs); - hipError_t hipModuleGetFunction(hipFunction_t* hfunc, hipModule_t hmod, const char* name) { HIP_INIT_API(hipModuleGetFunction, hfunc, hmod, name); diff --git a/hipamd/src/hiprtc/hiprtcComgrHelper.cpp b/hipamd/src/hiprtc/hiprtcComgrHelper.cpp index 9f902c071a..9079d84024 100644 --- a/hipamd/src/hiprtc/hiprtcComgrHelper.cpp +++ b/hipamd/src/hiprtc/hiprtcComgrHelper.cpp @@ -441,39 +441,6 @@ bool isCodeObjectCompatibleWithDevice(std::string co_triple_target_id, return true; } -bool UnbundleBitCode(const std::vector& bundled_llvm_bitcode, const std::string& isa, - size_t& co_offset, size_t& co_size) { - std::string magic(bundled_llvm_bitcode.begin(), - bundled_llvm_bitcode.begin() + bundle_magic_string_size); - if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) { - // Handle case where the whole file is unbundled - return true; - } - - std::string bundled_llvm_bitcode_s(bundled_llvm_bitcode.begin(), - bundled_llvm_bitcode.begin() + bundled_llvm_bitcode.size()); - const void* data = reinterpret_cast(bundled_llvm_bitcode_s.c_str()); - const auto obheader = reinterpret_cast(data); - const auto* desc = &obheader->desc[0]; - for (uint64_t idx = 0; idx < obheader->numOfCodeObjects; ++idx, - desc = reinterpret_cast( - reinterpret_cast(&desc->bundleEntryId[0]) + - desc->bundleEntryIdSize)) { - const void* image = - reinterpret_cast(reinterpret_cast(obheader) + desc->offset); - const size_t image_size = desc->size; - std::string bundleEntryId{desc->bundleEntryId, desc->bundleEntryIdSize}; - - // Check if the device id and code object id are compatible - if (isCodeObjectCompatibleWithDevice(bundleEntryId, isa)) { - co_offset = (reinterpret_cast(image) - reinterpret_cast(data)); - co_size = image_size; - break; - } - } - return true; -} - bool addCodeObjData(amd_comgr_data_set_t& input, const std::vector& source, const std::string& name, const amd_comgr_data_kind_t type) { amd_comgr_data_t data; diff --git a/hipamd/src/hiprtc/hiprtcComgrHelper.hpp b/hipamd/src/hiprtc/hiprtcComgrHelper.hpp index 05e1c013d0..aaa428ac70 100644 --- a/hipamd/src/hiprtc/hiprtcComgrHelper.hpp +++ b/hipamd/src/hiprtc/hiprtcComgrHelper.hpp @@ -31,8 +31,6 @@ THE SOFTWARE. namespace hiprtc { namespace helpers { -bool UnbundleBitCode(const std::vector& bundled_bit_code, const std::string& isa, - size_t& co_offset, size_t& co_size); bool addCodeObjData(amd_comgr_data_set_t& input, const std::vector& source, const std::string& name, const amd_comgr_data_kind_t type); bool extractBuildLog(amd_comgr_data_set_t dataSet, std::string& buildLog); diff --git a/hipamd/src/hiprtc/hiprtcInternal.cpp b/hipamd/src/hiprtc/hiprtcInternal.cpp index de597272aa..9b88aa0347 100644 --- a/hipamd/src/hiprtc/hiprtcInternal.cpp +++ b/hipamd/src/hiprtc/hiprtcInternal.cpp @@ -545,8 +545,7 @@ amd_comgr_data_kind_t RTCLinkProgram::GetCOMGRDataKind(hiprtcJITInputType input_ data_kind = AMD_COMGR_DATA_KIND_BC; break; case HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE: - data_kind = - HIPRTC_USE_RUNTIME_UNBUNDLER ? AMD_COMGR_DATA_KIND_BC : AMD_COMGR_DATA_KIND_BC_BUNDLE; + data_kind = AMD_COMGR_DATA_KIND_BC_BUNDLE; break; case HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE: data_kind = AMD_COMGR_DATA_KIND_AR_BUNDLE; @@ -561,32 +560,13 @@ amd_comgr_data_kind_t RTCLinkProgram::GetCOMGRDataKind(hiprtcJITInputType input_ bool RTCLinkProgram::AddLinkerDataImpl(std::vector& link_data, hiprtcJITInputType input_type, std::string& link_file_name) { - std::vector llvm_bitcode; - // If this is bundled bitcode then unbundle this. - if (HIPRTC_USE_RUNTIME_UNBUNDLER && input_type == HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE) { - if (!findIsa()) { - return false; - } - - size_t co_offset = 0; - size_t co_size = 0; - if (!UnbundleBitCode(link_data, isa_, co_offset, co_size)) { - LogError("Error in hiprtc: unable to unbundle the llvm bitcode"); - return false; - } - - llvm_bitcode.assign(link_data.begin() + co_offset, link_data.begin() + co_offset + co_size); - } else { - llvm_bitcode.assign(link_data.begin(), link_data.end()); - } - amd_comgr_data_kind_t data_kind; if ((data_kind = GetCOMGRDataKind(input_type)) == AMD_COMGR_DATA_KIND_UNDEF) { LogError("Cannot find the correct COMGR data kind"); return false; } - if (!addCodeObjData(link_input_, llvm_bitcode, link_file_name, data_kind)) { + if (!addCodeObjData(link_input_, link_data, link_file_name, data_kind)) { LogError("Error in hiprtc: unable to add linked code object"); return false; } diff --git a/rocclr/device/comgrctx.cpp b/rocclr/device/comgrctx.cpp index 1a8afe60a9..5b825d0db0 100644 --- a/rocclr/device/comgrctx.cpp +++ b/rocclr/device/comgrctx.cpp @@ -119,6 +119,7 @@ bool Comgr::LoadLib(bool is_versioned) { GET_COMGR_SYMBOL(amd_comgr_get_mangled_name) GET_COMGR_SYMBOL(amd_comgr_populate_name_expression_map) GET_COMGR_SYMBOL(amd_comgr_map_name_expression_to_symbol_name) + GET_COMGR_SYMBOL(amd_comgr_action_info_set_bundle_entry_ids) is_ready_ = true; return true; } diff --git a/rocclr/device/comgrctx.hpp b/rocclr/device/comgrctx.hpp index d8dd5e17f5..03ea2eeee3 100644 --- a/rocclr/device/comgrctx.hpp +++ b/rocclr/device/comgrctx.hpp @@ -76,6 +76,7 @@ typedef amd_comgr_status_t (*t_amd_comgr_populate_mangled_names)(amd_comgr_data_ typedef amd_comgr_status_t (*t_amd_comgr_get_mangled_name)(amd_comgr_data_t data, size_t index, size_t *size, char *mangled_name); typedef amd_comgr_status_t (*t_amd_comgr_populate_name_expression_map)(amd_comgr_data_t data, size_t *count); typedef amd_comgr_status_t (*t_amd_comgr_map_name_expression_to_symbol_name)(amd_comgr_data_t data, size_t *size, char *name_expression, char* symbol_name); +typedef amd_comgr_status_t (*t_amd_comgr_action_info_set_bundle_entry_ids)(amd_comgr_action_info_t action_info, const char* bundle_entry_ids[], size_t count); struct ComgrEntryPoints { void* handle; @@ -129,6 +130,7 @@ struct ComgrEntryPoints { t_amd_comgr_get_mangled_name amd_comgr_get_mangled_name; t_amd_comgr_populate_name_expression_map amd_comgr_populate_name_expression_map; t_amd_comgr_map_name_expression_to_symbol_name amd_comgr_map_name_expression_to_symbol_name; + t_amd_comgr_action_info_set_bundle_entry_ids amd_comgr_action_info_set_bundle_entry_ids; }; #ifdef COMGR_DYN_DLL @@ -310,7 +312,11 @@ public: static amd_comgr_status_t map_name_expression_to_symbol_name(amd_comgr_data_t data, size_t *size, char *name_expression, char* symbol_name) { return COMGR_DYN(amd_comgr_map_name_expression_to_symbol_name)(data, size, name_expression, symbol_name); } - + static amd_comgr_status_t action_info_set_bundle_entry_ids(amd_comgr_action_info_t action_info, + const char* bundle_entry_ids[], size_t count) { + return COMGR_DYN(amd_comgr_action_info_set_bundle_entry_ids)(action_info, bundle_entry_ids, + count); + } private: static ComgrEntryPoints cep_; diff --git a/rocclr/utils/flags.hpp b/rocclr/utils/flags.hpp index 6eee4e8e48..5d894e4983 100644 --- a/rocclr/utils/flags.hpp +++ b/rocclr/utils/flags.hpp @@ -229,10 +229,6 @@ release(bool, ROC_SKIP_KERNEL_ARG_COPY, false, \ "If true, then runtime can skip kernel arg copy") \ release(bool, GPU_STREAMOPS_CP_WAIT, false, \ "Force the stream wait memory operation to wait on CP.") \ -release(bool, HIP_USE_RUNTIME_UNBUNDLER, false, \ - "Force this to use Runtime code object unbundler.") \ -release(bool, HIPRTC_USE_RUNTIME_UNBUNDLER, false, \ - "Set this to true to force runtime unbundler in hiprtc.") \ release(size_t, HIP_INITIAL_DM_SIZE, 8 * Mi, \ "Set initial heap size for device malloc.") \ release(bool, HIP_FORCE_DEV_KERNARG, true, \