diff --git a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp index 884bc2490ad6..75ed2c4b39a7 100644 --- a/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp +++ b/projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp @@ -2064,7 +2064,13 @@ namespace { using MSL = TensileLite::MasterSolutionLibrary; - m_library = std::dynamic_pointer_cast(lib); + m_library = std::dynamic_pointer_cast(lib); + if(!m_library->initLibraryMapping(tensileLibPath)) + { + std::cerr << "\nrocblaslt error: Could not initialize Tensile library " + "mapping" + << std::endl; + } m_tensileLibPath = tensileLibPath.string(); } return 0; @@ -2078,19 +2084,6 @@ namespace // rocblaslt_abort(); } } - -#if ROCBLASLT_TENSILE_LAZY_LOAD - // A workaround for getSolutionsFromIndex and isSolutionSupported with lazy_lib_load. - // preload() shouldn't be called more than once. - void preload() - { - auto lib = TensileLite::LoadLibraryFilePreload( - m_tensileLibPath, - std::vector{m_deviceSet.begin(), m_deviceSet.end()}); - using MSL = TensileLite::MasterSolutionLibrary; - m_library = std::dynamic_pointer_cast(lib); - } -#endif }; // Return the library and adapter for the current HIP device @@ -2099,12 +2092,7 @@ namespace library = nullptr, std::shared_ptr* deviceProp = nullptr, - int device = -1 -#if ROCBLASLT_TENSILE_LAZY_LOAD - , - bool isPreload = false -#endif - ) + int device = -1) try { // TensileHost is initialized on the first call @@ -2137,16 +2125,6 @@ namespace } } -#if ROCBLASLT_TENSILE_LAZY_LOAD - // A workaround for getSolutionsFromIndex and isSolutionSupported when lazy_lib_load is on. - // preload() shouldn't be called more than once. - if(isPreload) - static int once = [&] { - host.preload(); - *library = host.get_library(); - return 0; - }(); -#endif // If an adapter is found, it is assumed that the library is initialized if(library) *library = host.get_library(); @@ -3494,12 +3472,7 @@ rocblaslt_status std::shared_ptr deviceProp; std::shared_ptr hardware; -#if ROCBLASLT_TENSILE_LAZY_LOAD - // isPreload = true is to load placeholder libraries except code objects - auto adapter = get_library_and_adapter(&library, &deviceProp, handle->device, true); -#else auto adapter = get_library_and_adapter(&library, &deviceProp, handle->device); -#endif if(!library) { @@ -3508,9 +3481,8 @@ rocblaslt_status hardware = TensileLite::hip::GetDevice(*deviceProp); - int lastSolutionIndex = library->solutions.rbegin()->first; - bool isOutOfBound = true; - int i = 0; + bool isOutOfBound = false; + int i = 0; for(auto index : solutionIndex) { #ifdef HIPBLASLT_USE_ROCROLLER @@ -3522,10 +3494,12 @@ rocblaslt_status } #endif - isOutOfBound = isOutOfBound && (index > lastSolutionIndex); auto solution = library->getSolutionByIndex(*hardware, index); if(!solution) + { + isOutOfBound = true; continue; + } rocblaslt_matmul_heuristic_result result; memset(&result, 0, sizeof(rocblaslt_matmul_heuristic_result)); memset(result.algo.data, 0, sizeof(result.algo.data)); @@ -3556,12 +3530,7 @@ rocblaslt_status isSolutionSupported(rocblaslt_handle handle, std::shared_ptr deviceProp; std::shared_ptr hardware; -#if ROCBLASLT_TENSILE_LAZY_LOAD - // isPreload = true is a workaround for lazy_lib_load - auto adapter = get_library_and_adapter(&library, &deviceProp, handle->device, true); -#else auto adapter = get_library_and_adapter(&library, &deviceProp, handle->device); -#endif if(!library) { diff --git a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/MasterSolutionLibrary.hpp b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/MasterSolutionLibrary.hpp index 4a1356d23c9e..6191a2ff7705 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/MasterSolutionLibrary.hpp +++ b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/MasterSolutionLibrary.hpp @@ -51,6 +51,9 @@ namespace TensileLite // If lazy loading is used, this may be updated in const functions SolutionMap* solutions; std::mutex* solutionsGuard; + std::set* loadedFiles; + + void* indexLoadedLibraries; }; /** @@ -84,13 +87,95 @@ namespace TensileLite ")"); } + std::string libraryDirectory; + std::string suffix; + std::map libraryMapping; + mutable std::set loadedFiles; + + mutable std::map>> + indexLoadedLibraries; + std::shared_ptr> library; - SolutionMap solutions; + mutable SolutionMap solutions; std::string version; mutable std::mutex solutionsGuard; MasterSolutionLibrary() = default; + bool initLibraryMapping(const std::string& tensileLibPath) + { + libraryDirectory = tensileLibPath; + size_t directoryPos = tensileLibPath.rfind('/'); + if(directoryPos != std::string::npos) + libraryDirectory.resize(directoryPos + 1); + else + libraryDirectory = '.'; + + //Extract file extension + size_t periodPos = tensileLibPath.rfind('.'); + suffix = tensileLibPath.substr(periodPos); + + // libraryMapping + libraryMapping + = LoadLibraryMapping(libraryDirectory + "/TensileLiteLibrary_lazy_Mapping.dat"); + if(libraryMapping.empty()) + { + std::cout << "No library mapping found in " << libraryDirectory << std::endl; + return false; + } + return true; + } + + void loadLibrary(const int index) const + { + auto it = libraryMapping.upper_bound(index); + if(it != libraryMapping.begin()) + { + --it; + std::string filePrefix = it->second; + // load the file here directly and push the library for later use. + { + std::lock_guard lock(solutionsGuard); + if(loadedFiles.find(filePrefix) != loadedFiles.end()) + return; + } + if(Debug::Instance().printDataInit()) + std::cout << "Loading library for index " << index + << " from file: " << filePrefix << std::endl; + + std::string path = (libraryDirectory + "/" + filePrefix + suffix).c_str(); + auto newLibrary = LoadLibraryFile(path); + auto mLibrary + = static_cast*>(newLibrary.get()); + + using std::begin; + using std::end; + + std::lock_guard lock(solutionsGuard); + if(loadedFiles.find(filePrefix) != loadedFiles.end()) + return; + // Push to cache + indexLoadedLibraries[filePrefix] = mLibrary->library; + + std::transform(begin(mLibrary->solutions), + end(mLibrary->solutions), + std::inserter(solutions, end(solutions)), + [this, filePrefix](auto& i) { + i.second->codeObjectFilename = filePrefix + ".co"; + return i; + }); + loadedFiles.insert(filePrefix); + + if(Debug::Instance().printCodeObjectInfo()) + std::cout << "load placeholder library " << path << std::endl + << mLibrary->solutions.size() << " solutions loaded" << std::endl; + } + else + { + std::cerr << "No library found for index " << index << std::endl; + } + } + virtual std::shared_ptr getSolutionByIndex(MyProblem const& problem, Hardware const& hardware, const int index) const override @@ -111,6 +196,7 @@ namespace TensileLite virtual std::shared_ptr getSolutionByIndex(Hardware const& hardware, const int index) const override { + loadLibrary(index); if(solutions.find(index) == solutions.end()) { return std::shared_ptr(); diff --git a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/PlaceholderLibrary.hpp b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/PlaceholderLibrary.hpp index c40f3bd65aa1..2ab89b2912fa 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/PlaceholderLibrary.hpp +++ b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/PlaceholderLibrary.hpp @@ -2,7 +2,7 @@ * * MIT License * - * Copyright (C) 2022-2024 Advanced Micro Devices, Inc. + * Copyright (C) 2022-2025 Advanced Micro Devices, Inc. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -32,6 +32,7 @@ #include #include +#include namespace TensileLite { @@ -120,7 +121,7 @@ namespace TensileLite case LazyLoadingInit::gfx1150: return "TensileLibrary_*_gfx1150"; case LazyLoadingInit::gfx1151: - return "TensileLibrary_*_gfx1151"; + return "TensileLibrary_*_gfx1151"; case LazyLoadingInit::gfx1200: return "TensileLibrary_*_gfx1200"; case LazyLoadingInit::gfx1201: @@ -135,6 +136,7 @@ namespace TensileLite struct PlaceholderLibrary : public SolutionLibrary { mutable std::shared_ptr> library; + mutable std::set* loadedFiles; mutable SolutionMap* masterSolutions; mutable std::mutex* solutionsGuard; mutable std::mutex lazyLoadingGuard; @@ -142,6 +144,9 @@ namespace TensileLite std::string suffix; std::string libraryDirectory; + mutable std::map>>* + indexLoadedLibraries; + PlaceholderLibrary() = default; bool loadPlaceholderLibrary() const @@ -157,6 +162,20 @@ namespace TensileLite library = mLibrary->library; std::lock_guard lock(*solutionsGuard); + if(loadedFiles->find(filePrefix) != loadedFiles->end()) + { + if(indexLoadedLibraries->find(filePrefix) == indexLoadedLibraries->end()) + return true; + library = (*indexLoadedLibraries)[filePrefix]; + indexLoadedLibraries->erase(filePrefix); + + // If the library is already loaded, just return it + if(Debug::Instance().printCodeObjectInfo()) + std::cout << "load placeholder library " << filePrefix + << " from MasterLibrary cache" << std::endl; + return true; + } + using std::begin; using std::end; @@ -167,6 +186,7 @@ namespace TensileLite i.second->codeObjectFilename = getCodeObjectFileName(); return i; }); + loadedFiles->insert(filePrefix); if(Debug::Instance().printCodeObjectInfo()) std::cout << "load placeholder library " << path << std::endl diff --git a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/PlaceholderLibrary.hpp b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/PlaceholderLibrary.hpp index 194925b99fce..26328290677f 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/PlaceholderLibrary.hpp +++ b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/PlaceholderLibrary.hpp @@ -2,7 +2,7 @@ * * MIT License * - * Copyright (C) 2022-2024 Advanced Micro Devices, Inc. + * Copyright (C) 2022-2025 Advanced Micro Devices, Inc. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -56,8 +56,13 @@ namespace TensileLite if(!iot::outputting(io)) { auto ctx = static_cast*>(iot::getContext(io)); - lib.masterSolutions = ctx->solutions; - lib.solutionsGuard = ctx->solutionsGuard; + lib.masterSolutions = ctx->solutions; + lib.solutionsGuard = ctx->solutionsGuard; + lib.loadedFiles = ctx->loadedFiles; + lib.indexLoadedLibraries = static_cast< + std::map>>*>( + ctx->indexLoadedLibraries); //Extract directory where TensileLibrary.dat/yaml file is located //TODO: Consider refactoring this to use filesystem::path handling instead diff --git a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/SolutionLibrary.hpp b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/SolutionLibrary.hpp index c170062b38d1..4107e12b5e34 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/SolutionLibrary.hpp +++ b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/SolutionLibrary.hpp @@ -41,11 +41,10 @@ #include #include +#include #include #include #include -#include - namespace TensileLite { @@ -163,8 +162,10 @@ namespace TensileLite lib.solutions[s->index] = s; auto ctx = static_cast*>(iot::getContext(io)); - ctx->solutions = &lib.solutions; - ctx->solutionsGuard = &lib.solutionsGuard; + ctx->solutions = &lib.solutions; + ctx->solutionsGuard = &lib.solutionsGuard; + ctx->loadedFiles = &lib.loadedFiles; + ctx->indexLoadedLibraries = (void*)&lib.indexLoadedLibraries; } std::shared_ptr> innerLibrary; diff --git a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Tensile.hpp b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Tensile.hpp index 0affde5bc022..9e8023b03928 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Tensile.hpp +++ b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Tensile.hpp @@ -187,6 +187,8 @@ namespace TensileLite */ enum class LazyLoadingInit; + TENSILE_API std::map LoadLibraryMapping(std::string const& filename); + template TENSILE_API std::shared_ptr> LoadLibraryFile(std::string const& filename); diff --git a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/msgpack/Loading.hpp b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/msgpack/Loading.hpp index cbed8d6bb028..3536a594260e 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/msgpack/Loading.hpp +++ b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/msgpack/Loading.hpp @@ -2,7 +2,7 @@ * * MIT License * - * Copyright (C) 2022 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (C) 2022-2025 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -34,6 +34,8 @@ namespace TensileLite { + std::map MessagePackLoadLibraryMapping(std::string const& filename); + template std::shared_ptr> MessagePackLoadLibraryFile(std::string const& filename, diff --git a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/msgpack/MessagePack.hpp b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/msgpack/MessagePack.hpp index 8d2e855b1a16..81687deaa70f 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/msgpack/MessagePack.hpp +++ b/projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/msgpack/MessagePack.hpp @@ -2,7 +2,7 @@ * * MIT License * - * Copyright (C) 2022-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (C) 2022-2025 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -90,7 +90,7 @@ namespace TensileLite static const bool value = true; }; - void objectToMap(msgpack::object& object, + void objectToMap(const msgpack::object& object, std::unordered_map& map); struct MessagePackInput diff --git a/projects/hipblaslt/tensilelite/Tensile/Source/lib/source/Tensile.cpp b/projects/hipblaslt/tensilelite/Tensile/Source/lib/source/Tensile.cpp index 7b110ff43314..753d9d62dc4c 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Source/lib/source/Tensile.cpp +++ b/projects/hipblaslt/tensilelite/Tensile/Source/lib/source/Tensile.cpp @@ -2,7 +2,7 @@ * * MIT License * - * Copyright (C) 2022-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (C) 2022-2025 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -49,6 +49,11 @@ namespace TensileLite TENSILE_API SolutionAdapter::~SolutionAdapter() = default; #ifdef TENSILE_DEFAULT_SERIALIZATION + std::map LoadLibraryMapping(std::string const& filename) + { + return MessagePackLoadLibraryMapping(filename); + } + template std::shared_ptr> LoadLibraryFile(std::string const& filename) diff --git a/projects/hipblaslt/tensilelite/Tensile/Source/lib/source/msgpack/MessagePack.cpp b/projects/hipblaslt/tensilelite/Tensile/Source/lib/source/msgpack/MessagePack.cpp index 4a094df4847b..2246cdc936e7 100644 --- a/projects/hipblaslt/tensilelite/Tensile/Source/lib/source/msgpack/MessagePack.cpp +++ b/projects/hipblaslt/tensilelite/Tensile/Source/lib/source/msgpack/MessagePack.cpp @@ -2,7 +2,7 @@ * * MIT License * - * Copyright (C) 2022-2023 Advanced Micro Devices, Inc. All rights reserved. + * Copyright (C) 2022-2025 Advanced Micro Devices, Inc. All rights reserved. * * Permission is hereby granted, free of charge, to any person obtaining a copy * of this software and associated documentation files (the "Software"), to deal @@ -34,7 +34,7 @@ namespace TensileLite { namespace Serialization { - void objectToMap(msgpack::object& object, + void objectToMap(const msgpack::object& object, std::unordered_map& result) { if(object.type != msgpack::type::object_type::MAP) @@ -67,13 +67,10 @@ namespace TensileLite } } - template - std::shared_ptr> - MessagePackLoadLibraryFile(std::string const& filename, - const std::vector& preloaded) + inline bool fileToMsgObject(std::string const& filename, msgpack::object_handle& result) { // parse file into a msgpack::object_handle - msgpack::object_handle result; + try { std::ifstream in(filename, std::ios::in | std::ios::binary); @@ -83,7 +80,7 @@ namespace TensileLite std::cout << "Error loading " << filename << " (msgpack):\nFailed to open file" << std::endl; - return nullptr; + return false; } msgpack::unpacker unp; @@ -107,7 +104,7 @@ namespace TensileLite << error_str << std::endl; } - return nullptr; + return false; } } catch(std::runtime_error const& exc) @@ -115,8 +112,52 @@ namespace TensileLite if(Debug::Instance().printDataInit()) std::cout << "Error loading msgpack data:\n" << exc.what() << std::endl; - return nullptr; + return false; } + return true; + } + + std::map MessagePackLoadLibraryMapping(std::string const& filename) + { + if(Debug::Instance().printDataInit()) + std::cout << "Loading library mapping from file: " << filename << std::endl; + msgpack::object_handle result; + if(!fileToMsgObject(filename, result)) + return {}; + + std::map libraryMapping; + try + { + std::unordered_map objectMap; + Serialization::objectToMap(result.get(), objectMap); + + for(auto const& pair : objectMap) + { + int key = std::stoi(pair.first); + std::string value; + pair.second.convert(value); + libraryMapping[key] = value; + } + } + catch(std::runtime_error const& exc) + { + if(Debug::Instance().printDataInit()) + std::cout << "Error loading library mapping: " << exc.what() << std::endl; + + return {}; + } + + return libraryMapping; + } + + template + std::shared_ptr> + MessagePackLoadLibraryFile(std::string const& filename, + const std::vector& preloaded) + { + msgpack::object_handle result; + if(!fileToMsgObject(filename, result)) + return nullptr; // copy data from msgpack::object_handle into MasterSolutionLibrary try diff --git a/projects/hipblaslt/tensilelite/Tensile/TensileCreateLibrary/Run.py b/projects/hipblaslt/tensilelite/Tensile/TensileCreateLibrary/Run.py index 16c06bf8640f..4764331ce4b6 100644 --- a/projects/hipblaslt/tensilelite/Tensile/TensileCreateLibrary/Run.py +++ b/projects/hipblaslt/tensilelite/Tensile/TensileCreateLibrary/Run.py @@ -572,13 +572,30 @@ def libraryIter(lib: MasterSolutionLibrary): if key != "fallback": value.merge(masterLibraries["fallback"]) masterLibraries.pop("fallback") + solIndex = [] for _, masterLibrary in masterLibraries.items(): for _, sol in masterLibrary.solutions.items(): solutions.append(sol.originalSolution) + solIndex.append(sol.index) for name, lib in masterLibrary.lazyLibraries.items(): for _, sol in lib.solutions.items(): sol.originalSolution._state["codeObjectFile"] = name solutions.append(sol.originalSolution) + solIndex.append(sol.index) + + # Get the solution index and it's codeObjectFile name + codeObjectFilesIndex = {} + for solution, index in zip(solutions, solIndex): + if "codeObjectFile" in solution._state and solution._state["codeObjectFile"] is not None: + if solution._state["codeObjectFile"] in codeObjectFilesIndex: + codeObjectFilesIndex[solution._state["codeObjectFile"]] = min(index, codeObjectFilesIndex[solution._state["codeObjectFile"]]) + else: + codeObjectFilesIndex[solution._state["codeObjectFile"]] = index + + # Reorder to int: name format + codeObjectFilesIndex = {v: k for k, v in codeObjectFilesIndex.items()} + # Reorder to maintain ascending order by index + codeObjectFilesIndex = dict(sorted(codeObjectFilesIndex.items())) # remove duplicates while preserving order numSoln = len(solutions) @@ -587,7 +604,7 @@ def libraryIter(lib: MasterSolutionLibrary): print1(f"Number of solutions parsed: {numSoln}") print1(f"Number of unique solutions: {len(solutions)}") - return solutions, masterLibraries + return solutions, masterLibraries, codeObjectFilesIndex ################################################################################ @@ -682,7 +699,7 @@ def validLogicFile(p: Path): print2("# %s" % logicFile) start_glds = timer() - solutions, masterLibraries = generateLogicDataAndSolutions( + solutions, masterLibraries, libraryMapping = generateLogicDataAndSolutions( logicFiles, arguments, asmToolchain.assembler, isaInfoMap ) stop_glds = timer() @@ -731,6 +748,9 @@ def writeMsl(name, lib): lib.applyNaming(splitGSU) LibraryIO.write(filename, state(lib), arguments["LibraryFormat"]) + filename = os.path.join(newLibraryDir, "TensileLiteLibrary_lazy_Mapping") + LibraryIO.write(filename, libraryMapping, "msgpack") + start_msl = timer() for archName, newMasterLibrary in masterLibraries.items(): if archName in archs: