diff --git a/projects/hip/src/code_object_bundle.cpp b/projects/hip/src/code_object_bundle.cpp index e2d47e05e2..ede7090a52 100644 --- a/projects/hip/src/code_object_bundle.cpp +++ b/projects/hip/src/code_object_bundle.cpp @@ -10,23 +10,55 @@ using namespace std; -hsa_isa_t hip_impl::triple_to_hsa_isa(const std::string& triple) { - static constexpr const char prefix[] = "hcc-amdgcn--amdhsa-gfx"; - static constexpr size_t prefix_sz = sizeof(prefix) - 1; +inline +std::string transmogrify_triple(const std::string& triple) +{ + static constexpr const char old_prefix[]{"hcc-amdgcn--amdhsa-gfx"}; + static constexpr const char new_prefix[]{"hcc-amdgcn-amd-amdhsa--gfx"}; - hsa_isa_t r = {}; + if (triple.find(old_prefix) == 0) { + return new_prefix + triple.substr(sizeof(old_prefix) - 1); + } - auto idx = triple.find(prefix); + return (triple.find(new_prefix) == 0) ? triple : ""; +} - if (idx != string::npos) { - idx += prefix_sz; - string tmp = "AMD:AMDGPU"; - while (idx != triple.size()) { - tmp.push_back(':'); - tmp.push_back(triple[idx++]); +inline +std::string isa_name(std::string triple) +{ + static constexpr const char offload_prefix[]{"hcc-"}; + + triple = transmogrify_triple(triple); + if (triple.empty()) return {}; + + triple.erase(0, sizeof(offload_prefix) - 1); + + static hsa_isa_t tmp{}; + static const bool is_old_rocr{ + hsa_isa_from_name(triple.c_str(), &tmp) != HSA_STATUS_SUCCESS}; + + if (is_old_rocr) { + auto tmp{triple.substr(triple.rfind('x') + 1)}; + triple.replace(0, std::string::npos, "AMD:AMDGPU"); + + for (auto&& x : tmp) { + triple.push_back(':'); + triple.push_back(x); } + } - hsa_isa_from_name(tmp.c_str(), &r); + return triple; +} + +hsa_isa_t hip_impl::triple_to_hsa_isa(const std::string& triple) { + const auto isa{isa_name(std::move(triple))}; + + if (isa.empty()) return hsa_isa_t({}); + + hsa_isa_t r{}; + + if(HSA_STATUS_SUCCESS != hsa_isa_from_name(isa.c_str(), &r)) { + r.handle = 0; } return r; @@ -59,4 +91,4 @@ hip_impl::Bundled_code_header::Bundled_code_header( } read(static_cast(p), static_cast(p) + sz, *this); -} \ No newline at end of file +} diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp index 34e164be8d..84451a17a0 100644 --- a/projects/hip/src/hip_module.cpp +++ b/projects/hip/src/hip_module.cpp @@ -429,30 +429,27 @@ string code_object_blob_for_agent(const void* maybe_bundled_code, hsa_agent_t ag } // namespace hipError_t ihipModuleGetFunction(hipFunction_t* func, hipModule_t hmod, const char* name) { - HIP_INIT_API(func, hmod, name); - if (!func || !name) return ihipLogStatus(hipErrorInvalidValue); + if (!func || !name) return hipErrorInvalidValue; auto ctx = ihipGetTlsDefaultCtx(); - if (!ctx) return ihipLogStatus(hipErrorInvalidContext); - - hipError_t ret = hipSuccess; + if (!ctx) return hipErrorInvalidContext; *func = new ihipModuleSymbol_t; - if (!*func) return ihipLogStatus(hipErrorInvalidValue); + if (!*func) return hipErrorInvalidValue; auto kernel = find_kernel_by_name(hmod->executable, name); - if (kernel.handle == 0u) return ihipLogStatus(hipErrorNotFound); + if (kernel.handle == 0u) return hipErrorNotFound; (*func)->_object = kernel_object(kernel); (*func)->_groupSegmentSize = group_size(kernel); (*func)->_privateSegmentSize = private_size(kernel); (*func)->_name = name; - return ihipLogStatus(hipSuccess); + return hipSuccess; } hipError_t hipModuleGetFunction(hipFunction_t* hfunc, hipModule_t hmod, const char* name) { @@ -474,6 +471,31 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t* dptr, size_t* bytes, hipModule_t h return ihipLogStatus(r); } +hipError_t ihipModuleLoadData(hipModule_t* module, const void* image) { + + if (!module) return hipErrorInvalidValue; + + *module = new ihipModule_t; + + auto ctx = ihipGetTlsDefaultCtx(); + if (!ctx) return hipErrorInvalidContext; + + hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr, + &(*module)->executable); + + auto tmp = code_object_blob_for_agent(image, this_agent()); + + (*module)->executable = hip_impl::load_executable( + tmp.empty() ? read_elf_file_as_string(image) : tmp, (*module)->executable, this_agent()); + + return (*module)->executable.handle ? hipSuccess : hipErrorUnknown; +} + +hipError_t hipModuleLoadData(hipModule_t* module, const void* image) { + HIP_INIT_API(module, image); + return ihipLogStatus(ihipModuleLoadData(module,image)); +} + hipError_t hipModuleLoad(hipModule_t* module, const char* fname) { HIP_INIT_API(module, fname); @@ -485,33 +507,13 @@ hipError_t hipModuleLoad(hipModule_t* module, const char* fname) { vector tmp{istreambuf_iterator{file}, istreambuf_iterator{}}; - return hipModuleLoadData(module, tmp.data()); -} - -hipError_t hipModuleLoadData(hipModule_t* module, const void* image) { - HIP_INIT_API(module, image); - - if (!module) return ihipLogStatus(hipErrorInvalidValue); - - *module = new ihipModule_t; - - auto ctx = ihipGetTlsDefaultCtx(); - if (!ctx) return ihipLogStatus(hipErrorInvalidContext); - - hsa_executable_create_alt(HSA_PROFILE_FULL, HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT, nullptr, - &(*module)->executable); - - auto tmp = code_object_blob_for_agent(image, this_agent()); - - (*module)->executable = hip_impl::load_executable( - tmp.empty() ? read_elf_file_as_string(image) : tmp, (*module)->executable, this_agent()); - - return ihipLogStatus((*module)->executable.handle ? hipSuccess : hipErrorUnknown); + return ihipLogStatus(ihipModuleLoadData(module, tmp.data())); } hipError_t hipModuleLoadDataEx(hipModule_t* module, const void* image, unsigned int numOptions, hipJitOption* options, void** optionValues) { - return hipModuleLoadData(module, image); + HIP_INIT_API(module, image, numOptions, options, optionValues); + return ihipLogStatus(ihipModuleLoadData(module, image)); } hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const char* name) {