SWDEV-236178 - Remove use of old routines functions()/modules().
Change-Id: I3a1ed967227c91b0d8cdf39e1360ade685e0bc73
[ROCm/clr commit: fd783c1e88]
Tá an tiomantas seo le fáil i:
kjayapra-amd
tiomanta ag
Karthik Jayaprakash
Karthik Jayaprakash
tuismitheoir
7aa68bc3cb
tiomantas
db1fbf23d2
@@ -561,202 +561,6 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(int* numBlocks,
|
||||
|
||||
namespace hip_impl {
|
||||
|
||||
struct dl_phdr_info {
|
||||
ELFIO::Elf64_Addr dlpi_addr;
|
||||
const char *dlpi_name;
|
||||
const ELFIO::Elf64_Phdr *dlpi_phdr;
|
||||
ELFIO::Elf64_Half dlpi_phnum;
|
||||
};
|
||||
|
||||
extern "C" int dl_iterate_phdr(
|
||||
int (*callback) (struct dl_phdr_info *info, size_t size, void *data), void *data
|
||||
);
|
||||
|
||||
struct Symbol {
|
||||
std::string name;
|
||||
ELFIO::Elf64_Addr value = 0;
|
||||
ELFIO::Elf_Xword size = 0;
|
||||
ELFIO::Elf_Half sect_idx = 0;
|
||||
uint8_t bind = 0;
|
||||
uint8_t type = 0;
|
||||
uint8_t other = 0;
|
||||
};
|
||||
|
||||
inline Symbol read_symbol(const ELFIO::symbol_section_accessor& section, unsigned int idx) {
|
||||
assert(idx < section.get_symbols_num());
|
||||
|
||||
Symbol r;
|
||||
section.get_symbol(idx, r.name, r.value, r.size, r.bind, r.type, r.sect_idx, r.other);
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
template <typename P>
|
||||
inline ELFIO::section* find_section_if(ELFIO::elfio& reader, P p) {
|
||||
const auto it = find_if(reader.sections.begin(), reader.sections.end(), std::move(p));
|
||||
|
||||
return it != reader.sections.end() ? *it : nullptr;
|
||||
}
|
||||
|
||||
std::vector<std::pair<uintptr_t, std::string>> function_names_for(const ELFIO::elfio& reader,
|
||||
ELFIO::section* symtab) {
|
||||
std::vector<std::pair<uintptr_t, std::string>> r;
|
||||
ELFIO::symbol_section_accessor symbols{reader, symtab};
|
||||
|
||||
for (auto i = 0u; i != symbols.get_symbols_num(); ++i) {
|
||||
auto tmp = read_symbol(symbols, i);
|
||||
|
||||
if (tmp.type == STT_FUNC && tmp.sect_idx != SHN_UNDEF && !tmp.name.empty()) {
|
||||
r.emplace_back(tmp.value, tmp.name);
|
||||
}
|
||||
}
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
const std::vector<std::pair<uintptr_t, std::string>>& function_names_for_process() {
|
||||
static constexpr const char self[] = "/proc/self/exe";
|
||||
|
||||
static std::vector<std::pair<uintptr_t, std::string>> r;
|
||||
static std::once_flag f;
|
||||
|
||||
std::call_once(f, []() {
|
||||
ELFIO::elfio reader;
|
||||
|
||||
if (reader.load(self)) {
|
||||
const auto it = find_section_if(
|
||||
reader, [](const ELFIO::section* x) { return x->get_type() == SHT_SYMTAB; });
|
||||
|
||||
if (it) r = function_names_for(reader, it);
|
||||
}
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
|
||||
const std::unordered_map<uintptr_t, std::string>& function_names()
|
||||
{
|
||||
static std::unordered_map<uintptr_t, std::string> r{
|
||||
function_names_for_process().cbegin(),
|
||||
function_names_for_process().cend()};
|
||||
static std::once_flag f;
|
||||
|
||||
std::call_once(f, []() {
|
||||
dl_iterate_phdr([](dl_phdr_info* info, size_t, void*) {
|
||||
ELFIO::elfio reader;
|
||||
|
||||
if (reader.load(info->dlpi_name)) {
|
||||
const auto it = find_section_if(
|
||||
reader, [](const ELFIO::section* x) { return x->get_type() == SHT_SYMTAB; });
|
||||
|
||||
if (it) {
|
||||
auto n = function_names_for(reader, it);
|
||||
|
||||
for (auto&& f : n) f.first += info->dlpi_addr;
|
||||
|
||||
r.insert(make_move_iterator(n.begin()), make_move_iterator(n.end()));
|
||||
}
|
||||
}
|
||||
return 0;
|
||||
},
|
||||
nullptr);
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
std::vector<char> bundles_for_process() {
|
||||
static constexpr const char self[] = "/proc/self/exe";
|
||||
static constexpr const char kernel_section[] = ".kernel";
|
||||
std::vector<char> r;
|
||||
|
||||
ELFIO::elfio reader;
|
||||
|
||||
if (reader.load(self)) {
|
||||
auto it = find_section_if(
|
||||
reader, [](const ELFIO::section* x) { return x->get_name() == kernel_section; });
|
||||
|
||||
if (it) r.insert(r.end(), it->get_data(), it->get_data() + it->get_size());
|
||||
}
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
const std::vector<hipModule_t>& modules() {
|
||||
static std::vector<hipModule_t> r;
|
||||
static std::once_flag f;
|
||||
|
||||
std::call_once(f, []() {
|
||||
static std::vector<std::vector<char>> bundles{bundles_for_process()};
|
||||
|
||||
dl_iterate_phdr(
|
||||
[](dl_phdr_info* info, std::size_t, void*) {
|
||||
ELFIO::elfio tmp;
|
||||
if (tmp.load(info->dlpi_name)) {
|
||||
const auto it = find_section_if(
|
||||
tmp, [](const ELFIO::section* x) { return x->get_name() == ".kernel"; });
|
||||
|
||||
if (it) bundles.emplace_back(it->get_data(), it->get_data() + it->get_size());
|
||||
}
|
||||
return 0;
|
||||
},
|
||||
nullptr);
|
||||
|
||||
for (auto&& bundle : bundles) {
|
||||
if (bundle.empty()) {
|
||||
continue;
|
||||
}
|
||||
std::string magic(&bundle[0], sizeof(CLANG_OFFLOAD_BUNDLER_MAGIC_STR) - 1);
|
||||
if (magic.compare(CLANG_OFFLOAD_BUNDLER_MAGIC_STR))
|
||||
continue;
|
||||
|
||||
const auto obheader = reinterpret_cast<const hip::CodeObject::__ClangOffloadBundleHeader*>(&bundle[0]);
|
||||
const auto* desc = &obheader->desc[0];
|
||||
for (uint64_t i = 0; i < obheader->numBundles; ++i,
|
||||
desc = reinterpret_cast<const hip::CodeObject::__ClangOffloadBundleDesc*>(
|
||||
reinterpret_cast<uintptr_t>(&desc->triple[0]) + desc->tripleSize)) {
|
||||
|
||||
std::string triple(desc->triple, sizeof(HCC_AMDGCN_AMDHSA_TRIPLE) - 1);
|
||||
if (triple.compare(HCC_AMDGCN_AMDHSA_TRIPLE))
|
||||
continue;
|
||||
|
||||
std::string target(desc->triple + sizeof(HCC_AMDGCN_AMDHSA_TRIPLE),
|
||||
desc->tripleSize - sizeof(HCC_AMDGCN_AMDHSA_TRIPLE));
|
||||
|
||||
if (isCompatibleCodeObject(target, hip::getCurrentDevice()->devices()[0]->info().name_)) {
|
||||
hipModule_t module;
|
||||
if (hipSuccess == hipModuleLoadData(&module, reinterpret_cast<const void*>(
|
||||
reinterpret_cast<uintptr_t>(obheader) + desc->offset)))
|
||||
r.push_back(module);
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
const std::unordered_map<uintptr_t, hipFunction_t>& functions()
|
||||
{
|
||||
static std::unordered_map<uintptr_t, hipFunction_t> r;
|
||||
static std::once_flag f;
|
||||
|
||||
std::call_once(f, []() {
|
||||
for (auto&& function : function_names()) {
|
||||
for (auto&& module : modules()) {
|
||||
hipFunction_t f;
|
||||
if (hipSuccess == hipModuleGetFunction(&f, module, function.second.c_str())) {
|
||||
r[function.first] = f;
|
||||
}
|
||||
}
|
||||
}
|
||||
});
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
void hipLaunchKernelGGLImpl(
|
||||
uintptr_t function_address,
|
||||
const dim3& numBlocks,
|
||||
@@ -767,11 +571,19 @@ void hipLaunchKernelGGLImpl(
|
||||
{
|
||||
HIP_INIT();
|
||||
|
||||
const auto it = functions().find(function_address);
|
||||
if (it == functions().cend())
|
||||
assert(0);
|
||||
hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
|
||||
int deviceId = (s != nullptr)? s->DeviceId() : ihipGetDevice();
|
||||
if (deviceId == -1) {
|
||||
DevLogPrintfError("Wrong Device Id: %d \n", deviceId);
|
||||
}
|
||||
|
||||
hipModuleLaunchKernel(it->second,
|
||||
hipFunction_t func = nullptr;
|
||||
hipError_t hip_error = PlatformState::instance().getStatFunc(&func, reinterpret_cast<void*>(function_address), deviceId);
|
||||
if ((hip_error != hipSuccess) || (func == nullptr)) {
|
||||
DevLogPrintfError("Cannot find the static function: 0x%x", function_address);
|
||||
}
|
||||
|
||||
hipModuleLaunchKernel(func,
|
||||
numBlocks.x, numBlocks.y, numBlocks.z,
|
||||
dimBlocks.x, dimBlocks.y, dimBlocks.z,
|
||||
sharedMemBytes, stream, nullptr, kernarg);
|
||||
@@ -815,16 +627,7 @@ hipError_t ihipLaunchKernel(const void* hostFunction,
|
||||
hipFunction_t func = nullptr;
|
||||
hipError_t hip_error = PlatformState::instance().getStatFunc(&func, hostFunction, deviceId);
|
||||
if ((hip_error != hipSuccess) || (func == nullptr)) {
|
||||
#ifdef ATI_OS_LINUX
|
||||
const auto it = hip_impl::functions().find(reinterpret_cast<uintptr_t>(hostFunction));
|
||||
if (it == hip_impl::functions().cend()) {
|
||||
DevLogPrintfError("Cannot find function: 0x%x \n", hostFunction);
|
||||
HIP_RETURN(hipErrorInvalidDeviceFunction);
|
||||
}
|
||||
func = it->second;
|
||||
#else
|
||||
HIP_RETURN(hipErrorInvalidDeviceFunction);
|
||||
#endif
|
||||
}
|
||||
HIP_RETURN(ihipModuleLaunchKernel(func, (gridDim.x * blockDim.x), (gridDim.y * blockDim.y),
|
||||
(gridDim.z * blockDim.z), blockDim.x, blockDim.y, blockDim.z,
|
||||
|
||||
Tagairt in Eagrán Nua
Cuir bac ar úsáideoir