diff --git a/hipamd/src/hip_code_object.cpp b/hipamd/src/hip_code_object.cpp index bbc8fd0017..92130aa077 100644 --- a/hipamd/src/hip_code_object.cpp +++ b/hipamd/src/hip_code_object.cpp @@ -35,18 +35,14 @@ 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 { -size_t constexpr strLiteralLength(char const* str) { - return *str ? 1 + strLiteralLength(str + 1) : 0; -} -constexpr char const* CLANG_OFFLOAD_BUNDLER_MAGIC_STR = "__CLANG_OFFLOAD_BUNDLE__"; -constexpr char const* OFFLOAD_KIND_HIP = "hip"; -constexpr char const* OFFLOAD_KIND_HIPV4 = "hipv4"; -constexpr char const* OFFLOAD_KIND_HCC = "hcc"; -constexpr char const* AMDGCN_TARGET_TRIPLE = "amdgcn-amd-amdhsa-"; +constexpr char kOffloadBundleMagicStr[] = "__CLANG_OFFLOAD_BUNDLE__"; +constexpr char kOffloadKindHip[] = "hip"; +constexpr char kOffloadKindHipv4[] = "hipv4"; +constexpr char kOffloadKindHcc[] = "hcc"; +constexpr char kAmdgcnTargetTriple[] = "amdgcn-amd-amdhsa-"; // ClangOFFLOADBundle info. -static constexpr size_t bundle_magic_string_size = - strLiteralLength(CLANG_OFFLOAD_BUNDLER_MAGIC_STR); +static constexpr size_t kOffloadBundleMagicStrSize = sizeof(kOffloadBundleMagicStr); // Clang Offload bundler description & Header. struct __ClangOffloadBundleInfo { @@ -57,7 +53,7 @@ struct __ClangOffloadBundleInfo { }; struct __ClangOffloadBundleHeader { - const char magic[bundle_magic_string_size - 1]; + const char magic[kOffloadBundleMagicStrSize - 1]; uint64_t numOfCodeObjects; __ClangOffloadBundleInfo desc[1]; }; @@ -66,8 +62,8 @@ struct __ClangOffloadBundleHeader { namespace hip { bool CodeObject::IsClangOffloadMagicBundle(const void* data) { - std::string magic(reinterpret_cast(data), bundle_magic_string_size); - return magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) ? false : true; + 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); } @@ -276,7 +272,7 @@ static bool getTripleTargetIDFromCodeObject(const void* code_object, std::string std::string proc_name; if (!getProcName(ehdr->e_flags, proc_name, isXnackSupported, isSramEccSupported)) return false; - target_id = std::string(AMDGCN_TARGET_TRIPLE) + '-' + proc_name; + target_id = std::string(kAmdgcnTargetTriple) + '-' + proc_name; switch (ehdr->e_ident[EI_ABIVERSION]) { case ELFABIVERSION_AMDGPU_HSA_V2: { @@ -377,11 +373,11 @@ static bool getTargetIDValue(std::string& input, std::string& processor, char& s static bool getTripleTargetID(std::string bundled_co_entry_id, const void* code_object, std::string& co_triple_target_id) { std::string offload_kind = trimName(bundled_co_entry_id, '-'); - if (offload_kind != OFFLOAD_KIND_HIPV4 && offload_kind != OFFLOAD_KIND_HIP && - offload_kind != OFFLOAD_KIND_HCC) + if (offload_kind != kOffloadKindHipv4 && offload_kind != kOffloadKindHip && + offload_kind != kOffloadKindHcc) return false; - if (offload_kind != OFFLOAD_KIND_HIPV4) + if (offload_kind != kOffloadKindHipv4) return getTripleTargetIDFromCodeObject(code_object, co_triple_target_id); // For code object V4 onwards the bundled code object entry ID correctly @@ -396,7 +392,7 @@ static bool isCodeObjectCompatibleWithDevice(std::string co_triple_target_id, if (co_triple_target_id == agent_triple_target_id) return true; // Parse code object triple target id - if (!consume(co_triple_target_id, std::string(AMDGCN_TARGET_TRIPLE) + '-')) { + if (!consume(co_triple_target_id, std::string(kAmdgcnTargetTriple) + '-')) { return false; } @@ -409,7 +405,7 @@ static bool isCodeObjectCompatibleWithDevice(std::string co_triple_target_id, if (!co_triple_target_id.empty()) return false; // Parse agent isa triple target id - if (!consume(agent_triple_target_id, std::string(AMDGCN_TARGET_TRIPLE) + '-')) { + if (!consume(agent_triple_target_id, std::string(kAmdgcnTargetTriple) + '-')) { return false; } @@ -469,8 +465,8 @@ hipError_t CodeObject::ExtractCodeObjectFromMemory( hipError_t CodeObject::extractCodeObjectFromFatBinary( const void* data, const std::vector& agent_triple_target_ids, std::vector>& code_objs) { - std::string magic((const char*)data, bundle_magic_string_size); - if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR)) { + std::string magic((const char*)data, kOffloadBundleMagicStrSize); + if (magic.compare(kOffloadBundleMagicStr)) { return hipErrorInvalidKernelFile; } diff --git a/hipamd/src/hiprtc/hiprtcComgrHelper.cpp b/hipamd/src/hiprtc/hiprtcComgrHelper.cpp index c0b709796e..485c44516a 100644 --- a/hipamd/src/hiprtc/hiprtcComgrHelper.cpp +++ b/hipamd/src/hiprtc/hiprtcComgrHelper.cpp @@ -535,8 +535,9 @@ bool extractByteCodeBinary(const amd_comgr_data_set_t inDataSet, amd::Comgr::release_data(binaryData); - bin.reserve(binarySize); - bin.assign(binary, binary + binarySize); + std::vector temp_bin; + temp_bin.assign(binary, binary + binarySize); + bin = temp_bin; delete[] binary; return true; diff --git a/hipamd/src/hiprtc/hiprtcInternal.cpp b/hipamd/src/hiprtc/hiprtcInternal.cpp index 446379184f..f7de2315b2 100644 --- a/hipamd/src/hiprtc/hiprtcInternal.cpp +++ b/hipamd/src/hiprtc/hiprtcInternal.cpp @@ -706,9 +706,7 @@ bool RTCLinkProgram::LinkComplete(void** bin_out, size_t* size_out) { } *size_out = executable_.size(); - char* bin_out_c = new char[*size_out]; - std::copy(executable_.begin(), executable_.end(), bin_out_c); - *bin_out = reinterpret_cast(bin_out_c); + *bin_out = executable_.data(); return true; } diff --git a/rocclr/device/rocm/rocdevice.cpp b/rocclr/device/rocm/rocdevice.cpp index 7f18e220bd..86a102d441 100644 --- a/rocclr/device/rocm/rocdevice.cpp +++ b/rocclr/device/rocm/rocdevice.cpp @@ -88,6 +88,10 @@ inline bool getIsaMeta(std::string isaName, amd_comgr_metadata_node_t& isaMeta) return (status == AMD_COMGR_STATUS_SUCCESS) ? true : false; } +inline bool releaseIsaMeta(amd_comgr_metadata_node_t& isaMeta) { + return AMD_COMGR_STATUS_SUCCESS == amd::Comgr::destroy_metadata(isaMeta); +} + bool getValueFromIsaMeta(amd_comgr_metadata_node_t& isaMeta, const char* key, std::string& retValue) { amd_comgr_status_t status; @@ -102,6 +106,9 @@ bool getValueFromIsaMeta(amd_comgr_metadata_node_t& isaMeta, const char* key, retValue.resize(size - 1); status = amd::Comgr::get_metadata_string(valMeta, &size, &(retValue[0])); } + if (status == AMD_COMGR_STATUS_SUCCESS) { + status = amd::Comgr::destroy_metadata(valMeta); + } return (status == AMD_COMGR_STATUS_SUCCESS) ? true : false; } @@ -1726,6 +1733,9 @@ bool Device::populateOCLDeviceConstants() { info_.availableSGPRs_ = (getValueFromIsaMeta(isaMeta, "AddressableNumSGPRs", sgprValue)) ? (atoi(sgprValue.c_str())) : 0; + if (!releaseIsaMeta(isaMeta)) { + LogInfo("Can not release the isa meta node"); + } } // Generic support for HMM interfaces