diff --git a/hipamd/src/hip_code_object.cpp b/hipamd/src/hip_code_object.cpp index 477e9a8113..1a3963af84 100644 --- a/hipamd/src/hip_code_object.cpp +++ b/hipamd/src/hip_code_object.cpp @@ -30,34 +30,21 @@ 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 { -// 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 kOffloadBundleMagicStr[] = "__CLANG_OFFLOAD_BUNDLE__"; 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. -// Clang Offload bundler description & Header in uncompressed mode. +// ClangOFFLOADBundle info. +static constexpr size_t kOffloadBundleMagicStrSize = sizeof(kOffloadBundleMagicStr); + +// Clang Offload bundler description & Header. struct __ClangOffloadBundleInfo { uint64_t offset; uint64_t size; @@ -65,37 +52,16 @@ struct __ClangOffloadBundleInfo { const char bundleEntryId[1]; }; -struct __ClangOffloadBundleUncompressedHeader { - const char magic[kOffloadBundleUncompressedMagicStrSize - 1]; +struct __ClangOffloadBundleHeader { + const char magic[kOffloadBundleMagicStrSize - 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, 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; +bool CodeObject::IsClangOffloadMagicBundle(const void* data) { + std::string magic(reinterpret_cast(data), kOffloadBundleMagicStrSize - 1); + return magic.compare(kOffloadBundleMagicStr) ? false : true; } uint64_t CodeObject::ElfSize(const void* emi) { return amd::Elf::getElfSize(emi); } @@ -390,7 +356,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; @@ -403,18 +369,6 @@ 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))) { @@ -493,353 +447,111 @@ static bool isCodeObjectCompatibleWithDevice(std::string co_triple_target_id, return true; } -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; +// 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; } + + // 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); } -/** - * @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, size_t size, const std::vector& agent_triple_target_ids, - std::vector>& code_objs) { - hipError_t hipStatus = hipSuccess; - amd_comgr_status_t comgrStatus = AMD_COMGR_STATUS_SUCCESS; +// 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; + } - 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 extractCodeObjectFromFatBinary(data, device_names, code_objs); +} + +// This will be moved to COMGR eventually +hipError_t CodeObject::extractCodeObjectFromFatBinary( + const void* data, const std::vector& agent_triple_target_ids, + std::vector>& code_objs) { + std::string magic((const char*)data, kOffloadBundleMagicStrSize); + if (magic.compare(kOffloadBundleMagicStr)) { return hipErrorInvalidKernelFile; } - 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]); + // 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)); } - for (auto& device : devicesSet) { - bundleEntryIDs.push_back(device.c_str()); - } + 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; - 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; - } + if (num_code_objs == 0) break; + std::string bundleEntryId{desc->bundleEntryId, desc->bundleEntryIdSize}; - // 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; - } + std::string co_triple_target_id; + if (!getTripleTargetID(bundleEntryId, image, co_triple_target_id)) continue; - 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()); - } + 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; } } } - - // 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; + 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")); } - } - 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; - } - } + 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); - 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; - } - } + std::string co_triple_target_id; + bool valid_co = getTripleTargetID(bundleEntryId, image, co_triple_target_id); - 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 (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()); + } } + return hipErrorNoBinaryForGpu; } - - 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 d5b6df0a8c..c83ec37b24 100644 --- a/hipamd/src/hip_code_object.hpp +++ b/hipamd/src/hip_code_object.hpp @@ -48,22 +48,31 @@ 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, 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); + static bool IsClangOffloadMagicBundle(const void* data); protected: - CodeObject() {} + //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 a66b1106dc..562f193861 100644 --- a/hipamd/src/hip_fatbin.cpp +++ b/hipamd/src/hip_fatbin.cpp @@ -50,46 +50,57 @@ 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; } } - 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); + 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); + } } - if (!PlatformState::instance().CloseUniqueFileHandle(ufd_)) { - LogPrintfError("Cannot close file for fdesc: %d", ufd_->fdesc_); - 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"); } - } - fname_ = std::string(); - fdesc_ = amd::Os::FDescInit(); - fsize_ = 0; - image_ = nullptr; - uri_ = std::string(); + } 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); + } + } - if (0 == PlatformState::instance().UfdMapSize()) { - LogError("All Unique FDs are closed"); + fname_ = std::string(); + fdesc_ = amd::Os::FDescInit(); + fsize_ = 0; + image_ = nullptr; + uri_ = std::string(); } } @@ -103,8 +114,11 @@ void ListAllDeviceWithNoCOFromBundle(const std::unordered_map& devices) { +hipError_t FatBinaryInfo::ExtractFatBinaryUsingCOMGR(const std::vector& devices) { + amd_comgr_data_t data_object {0}; + amd_comgr_status_t comgr_status = AMD_COMGR_STATUS_SUCCESS; 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) { @@ -149,52 +163,107 @@ hipError_t FatBinaryInfo::ExtractFatBinary(const std::vector& devi fname_.c_str()); do { - 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); - 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 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()); if (fatbin_dev_info_[devices[dev_idx]->deviceId()]->program_ == nullptr) { hip_status = hipErrorOutOfMemory; break; } } + break; } - else { - LogPrintfError( - "CodeObject::extractCodeObjectFromFatBinary failed with status %d\n", - hip_status); + + // 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())); } } while(0); @@ -217,9 +286,115 @@ hipError_t FatBinaryInfo::ExtractFatBinary(const std::vector& devi fsize_ = 0; } } + + if (data_object.handle) { + if ((comgr_status = amd_comgr_release_data(data_object)) != AMD_COMGR_STATUS_SUCCESS) { + LogPrintfError("Releasing COMGR data failed with status %d ", comgr_status); + return hipErrorInvalidValue; + } + } + return hip_status; } +hipError_t FatBinaryInfo::ExtractFatBinary(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 ac6d64e19e..f9057a4b5a 100644 --- a/hipamd/src/hip_fatbin.hpp +++ b/hipamd/src/hip_fatbin.hpp @@ -64,6 +64,7 @@ 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 5d458e3f45..8707486424 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -66,6 +66,10 @@ 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 9079d84024..9f902c071a 100644 --- a/hipamd/src/hiprtc/hiprtcComgrHelper.cpp +++ b/hipamd/src/hiprtc/hiprtcComgrHelper.cpp @@ -441,6 +441,39 @@ 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 aaa428ac70..05e1c013d0 100644 --- a/hipamd/src/hiprtc/hiprtcComgrHelper.hpp +++ b/hipamd/src/hiprtc/hiprtcComgrHelper.hpp @@ -31,6 +31,8 @@ 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 9b88aa0347..de597272aa 100644 --- a/hipamd/src/hiprtc/hiprtcInternal.cpp +++ b/hipamd/src/hiprtc/hiprtcInternal.cpp @@ -545,7 +545,8 @@ 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 = AMD_COMGR_DATA_KIND_BC_BUNDLE; + data_kind = + HIPRTC_USE_RUNTIME_UNBUNDLER ? AMD_COMGR_DATA_KIND_BC : AMD_COMGR_DATA_KIND_BC_BUNDLE; break; case HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE: data_kind = AMD_COMGR_DATA_KIND_AR_BUNDLE; @@ -560,13 +561,32 @@ 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_, link_data, link_file_name, data_kind)) { + if (!addCodeObjData(link_input_, llvm_bitcode, 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 5b825d0db0..1a8afe60a9 100644 --- a/rocclr/device/comgrctx.cpp +++ b/rocclr/device/comgrctx.cpp @@ -119,7 +119,6 @@ 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 03ea2eeee3..d8dd5e17f5 100644 --- a/rocclr/device/comgrctx.hpp +++ b/rocclr/device/comgrctx.hpp @@ -76,7 +76,6 @@ 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; @@ -130,7 +129,6 @@ 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 @@ -312,11 +310,7 @@ 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 5d894e4983..6eee4e8e48 100644 --- a/rocclr/utils/flags.hpp +++ b/rocclr/utils/flags.hpp @@ -229,6 +229,10 @@ 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, \