Skip to content

Don't load all dat for index related API #579

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
KKyang merged 1 commit into from
Jul 11, 2025
Merged

Don't load all dat for index related API #579

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
57 changes: 13 additions & 44 deletions projects/hipblaslt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2064,7 +2064,13 @@ namespace
{
using MSL
= TensileLite::MasterSolutionLibrary<TensileLite::ContractionProblemGemm>;
m_library = std::dynamic_pointer_cast<MSL>(lib);
m_library = std::dynamic_pointer_cast<MSL>(lib);
if(!m_library->initLibraryMapping(tensileLibPath))

@ScottTodd ScottTodd Jul 17, 2025

Copy link
Copy Markdown
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Detected here: ROCm/TheRock#1064 (comment)

This line does not compile on MSVC:

[hipBLASLt] FAILED: CMakeFiles/hipblaslt.dir/C_/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp.obj 
[hipBLASLt] ccache B:\build\core\clr\dist\lib\llvm\bin\clang++.exe -DMSGPACK_DEFAULT_API_VERSION=3 -DMSGPACK_NO_BOOST -DROCBLASLT_TENSILE_LAZY_LOAD -DROCM_USE_FLOAT16 -DTENSILE_DEFAULT_SERIALIZATION -DTENSILE_USE_HIP -DUSE_PROF_API=1 -D__HIP_PLATFORM_AMD__ -D__HIP_PLATFORM_AMD__=1 -Dhipblaslt_EXPORTS -IC:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/next-cmake/../library/src/include -IC:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/next-cmake/../library/include -IB:/build/math-libs/BLAS/hipBLASLt/build/host-library/include -IC:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/next-cmake/../library/src/amd_detail/rocblaslt/include -IC:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/next-cmake/../library/src/amd_detail/rocblaslt/src/include -IB:/build/math-libs/BLAS/hipBLASLt/build/include -IC:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/next-cmake/../tensilelite/Tensile/Source/lib/include -IC:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/next-cmake/rocisa/rocisa-cpp/../../../tensilelite/rocisa -isystem B:/build/math-libs/BLAS/hipBLAS-common/dist/include -isystem B:/build/core/clr/dist/include -isystem B:/build/third-party/msgpack-cxx/dist/include -DWIN32 -DWIN32_LEAN_AND_MEAN -D_CRT_SECURE_NO_WARNINGS -DNOMINMAX -fms-extensions -fms-compatibility -D_ENABLE_EXTENDED_ALIGNED_STORAGE  -Wno-documentation-unknown-command -Wno-documentation-pedantic -Wno-unused-command-line-argument -Wno-explicit-specialization-storage-class -Wno-ignored-attributes -Wno-unknown-attributes -Wno-duplicate-decl-specifier --hip-path=B:/build/core/clr/dist --hip-device-lib-path=B:/build/core/clr/dist/lib/llvm/amdgcn/bitcode -O3 -DNDEBUG -std=gnu++17 -D_DLL -D_MT -Xclang --dependent-lib=msvcrt -x hip --offload-arch=gfx1100 --offload-arch=gfx1101 -MD -MT CMakeFiles/hipblaslt.dir/C_/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp.obj -MF CMakeFiles\hipblaslt.dir\C_\home\runner\_work\TheRock\TheRock\math-libs\BLAS\hipBLASLt\library\src\amd_detail\rocblaslt\src\tensile_host.cpp.obj.d -o CMakeFiles/hipblaslt.dir/C_/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp.obj -c C:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp
[hipBLASLt] In file included from C:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp:38:
[hipBLASLt] In file included from C:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/next-cmake/../library/src/amd_detail/rocblaslt/src/include\rocblaslt_mat_utils.hpp:32:
[hipBLASLt] In file included from C:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/next-cmake/../library/src/amd_detail/rocblaslt/src/include\utility.hpp:33:
[hipBLASLt] C:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/next-cmake/../library/src/amd_detail/rocblaslt/src/include\logging.h:190:65: warning: 'getpid' is deprecated: The POSIX name for this item is deprecated. Instead, use the ISO C and C++ conformant name: _getpid. See online help for details. [-Wdeprecated-declarations]
[hipBLASLt]   190 |                 logfile_pathname.replace(pos, 2, std::to_string(getpid()));
[hipBLASLt]       |                                                                 ^
[hipBLASLt] C:\Program Files (x86)\Windows Kits\10\include\10.0.26100.0\ucrt\process.h:363:9: note: 'getpid' has been explicitly marked deprecated here
[hipBLASLt]   363 |         _CRT_NONSTDC_DEPRECATE(_getpid)
[hipBLASLt]       |         ^
[hipBLASLt] C:\Program Files (x86)\Windows Kits\10\include\10.0.26100.0\ucrt\corecrt.h:428:50: note: expanded from macro '_CRT_NONSTDC_DEPRECATE'
[hipBLASLt]   428 |         #define _CRT_NONSTDC_DEPRECATE(_NewName) _CRT_DEPRECATE_TEXT(             \
[hipBLASLt]       |                                                  ^
[hipBLASLt] C:\Program Files\Microsoft Visual Studio\2022\Community\VC\Tools\MSVC\14.44.35207\include\vcruntime.h:358:47: note: expanded from macro '_CRT_DEPRECATE_TEXT'
[hipBLASLt]   358 | #define _CRT_DEPRECATE_TEXT(_Text) __declspec(deprecated(_Text))
[hipBLASLt]       |                                               ^
[hipBLASLt] C:/home/runner/_work/TheRock/TheRock/math-libs/BLAS/hipBLASLt/library/src/amd_detail/rocblaslt/src/tensile_host.cpp:2078:55: error: no viable conversion from 'std::filesystem::path' to 'const std::string' (aka 'const basic_string<char, char_traits<char>, allocator<char>>')
[hipBLASLt]  2078 |                     if(!m_library->initLibraryMapping(tensileLibPath))
[hipBLASLt]       |                                                       ^~~~~~~~~~~~~~

Missing a conversion like so?

Suggested change
if(!m_library->initLibraryMapping(tensileLibPath))
if(!m_library->initLibraryMapping(tensileLibPath.string()))

{
std::cerr << "\nrocblaslt error: Could not initialize Tensile library "
"mapping"
<< std::endl;
}
m_tensileLibPath = tensileLibPath.string();
}
return 0;
Expand All @@ -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<TensileLite::ContractionProblemGemm>(
m_tensileLibPath,
std::vector<TensileLite::LazyLoadingInit>{m_deviceSet.begin(), m_deviceSet.end()});
using MSL = TensileLite::MasterSolutionLibrary<TensileLite::ContractionProblemGemm>;
m_library = std::dynamic_pointer_cast<MSL>(lib);
}
#endif
};

// Return the library and adapter for the current HIP device
Expand All @@ -2099,12 +2092,7 @@ namespace
library
= nullptr,
std::shared_ptr<hipDeviceProp_t>* 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
Expand Down Expand Up @@ -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();
Expand Down Expand Up @@ -3494,12 +3472,7 @@ rocblaslt_status
std::shared_ptr<hipDeviceProp_t> deviceProp;
std::shared_ptr<TensileLite::Hardware> 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)
{
Expand All @@ -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
Expand All @@ -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));
Expand Down Expand Up @@ -3556,12 +3530,7 @@ rocblaslt_status isSolutionSupported(rocblaslt_handle handle,
std::shared_ptr<hipDeviceProp_t> deviceProp;
std::shared_ptr<TensileLite::Hardware> 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)
{
Expand Down
88 changes: 87 additions & 1 deletion projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/MasterSolutionLibrary.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -51,6 +51,9 @@ namespace TensileLite
// If lazy loading is used, this may be updated in const functions
SolutionMap<MySolution>* solutions;
std::mutex* solutionsGuard;
std::set<std::string>* loadedFiles;

void* indexLoadedLibraries;
};

/**
Expand Down Expand Up @@ -84,13 +87,95 @@ namespace TensileLite
")");
}

std::string libraryDirectory;
std::string suffix;
std::map<int, std::string> libraryMapping;
mutable std::set<std::string> loadedFiles;

mutable std::map<std::string, std::shared_ptr<SolutionLibrary<MyProblem, MySolution>>>
indexLoadedLibraries;

std::shared_ptr<SolutionLibrary<MyProblem, MySolution>> library;
SolutionMap<MySolution> solutions;
mutable SolutionMap<MySolution> 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");
Comment thread
KKyang marked this conversation as resolved.
Resolved
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<std::mutex> 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<MyProblem, MySolution>(path);
auto mLibrary
= static_cast<MasterSolutionLibrary<MyProblem, MySolution>*>(newLibrary.get());

using std::begin;
using std::end;

std::lock_guard<std::mutex> 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<MySolution> getSolutionByIndex(MyProblem const& problem,
Hardware const& hardware,
const int index) const override
Expand All @@ -111,6 +196,7 @@ namespace TensileLite
virtual std::shared_ptr<MySolution> getSolutionByIndex(Hardware const& hardware,
const int index) const override
{
loadLibrary(index);
if(solutions.find(index) == solutions.end())
{
return std::shared_ptr<MySolution>();
Expand Down
24 changes: 22 additions & 2 deletions projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/PlaceholderLibrary.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -32,6 +32,7 @@
#include <Tensile/Tensile.hpp>

#include <algorithm>
#include <map>

namespace TensileLite
{
Expand Down Expand Up @@ -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:
Expand All @@ -135,13 +136,17 @@ namespace TensileLite
struct PlaceholderLibrary : public SolutionLibrary<MyProblem, MySolution>
{
mutable std::shared_ptr<SolutionLibrary<MyProblem, MySolution>> library;
mutable std::set<std::string>* loadedFiles;
mutable SolutionMap<MySolution>* masterSolutions;
mutable std::mutex* solutionsGuard;
mutable std::mutex lazyLoadingGuard;
std::string filePrefix;
std::string suffix;
std::string libraryDirectory;

mutable std::map<std::string, std::shared_ptr<SolutionLibrary<MyProblem, MySolution>>>*
indexLoadedLibraries;

PlaceholderLibrary() = default;

bool loadPlaceholderLibrary() const
Expand All @@ -157,6 +162,20 @@ namespace TensileLite
library = mLibrary->library;

std::lock_guard<std::mutex> 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;

Expand All @@ -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
Expand Down
11 changes: 8 additions & 3 deletions ...laslt/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/PlaceholderLibrary.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -56,8 +56,13 @@ namespace TensileLite
if(!iot::outputting(io))
{
auto ctx = static_cast<LibraryIOContext<MySolution>*>(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<std::string,
std::shared_ptr<SolutionLibrary<MyProblem, MySolution>>>*>(
ctx->indexLoadedLibraries);

//Extract directory where TensileLibrary.dat/yaml file is located
//TODO: Consider refactoring this to use filesystem::path handling instead
Expand Down
9 changes: 5 additions & 4 deletions ...ipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Serialization/SolutionLibrary.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,11 +41,10 @@

#include <Tensile/Serialization/ExactLogicLibrary.hpp>
#include <Tensile/Serialization/GranularitySelectionLibrary.hpp>
#include <Tensile/Serialization/MLPClassificationLibrary.hpp>
#include <Tensile/Serialization/MapLibrary.hpp>
#include <Tensile/Serialization/MatchingLibrary.hpp>
#include <Tensile/Serialization/PlaceholderLibrary.hpp>
#include <Tensile/Serialization/MLPClassificationLibrary.hpp>


namespace TensileLite
{
Expand Down Expand Up @@ -163,8 +162,10 @@ namespace TensileLite
lib.solutions[s->index] = s;

auto ctx = static_cast<LibraryIOContext<MySolution>*>(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<SolutionLibrary<MyProblem, MySolution>> innerLibrary;
Expand Down
2 changes: 2 additions & 0 deletions projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/Tensile.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -187,6 +187,8 @@ namespace TensileLite
*/
enum class LazyLoadingInit;

TENSILE_API std::map<int, std::string> LoadLibraryMapping(std::string const& filename);

template <typename MyProblem, typename MySolution = typename MyProblem::Solution>
TENSILE_API std::shared_ptr<SolutionLibrary<MyProblem, MySolution>>
LoadLibraryFile(std::string const& filename);
Expand Down
4 changes: 3 additions & 1 deletion projects/hipblaslt/tensilelite/Tensile/Source/lib/include/Tensile/msgpack/Loading.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -34,6 +34,8 @@

namespace TensileLite
{
std::map<int, std::string> MessagePackLoadLibraryMapping(std::string const& filename);

template <typename MyProblem, typename MySolution>
std::shared_ptr<SolutionLibrary<MyProblem, MySolution>>
MessagePackLoadLibraryFile(std::string const& filename,
Expand Down
Loading