Skip to content

fix(tensilelite): align topology filtering with selection #11946

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

Closed
davidd-amd wants to merge 3 commits into from
Closed

fix(tensilelite): align topology filtering with selection #11946

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
e71edb7
fix(tensilelite): align topology filtering with selection
davidd-amd Sep 10, 2026
5602a27
fix(tensilelite): honour the resolved SK5 mode on unknown hardware
davidd-amd Sep 11, 2026
cf40154
fix(tensilelite): link tensilelite-tests in the default shared build
davidd-amd Sep 11, 2026
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
23 changes: 9 additions & 14 deletions projects/hipblaslt/tensilelite/client/src/SolutionIterator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -104,20 +104,6 @@ namespace TensileLite
return false;
}

// The all-solutions benchmark path selects kernels by index instead
// of going through a SolutionLibrary search. Apply the dynamic
// StreamK topology predicate explicitly so this path has the same
// support boundary as normal library selection. In particular, a
// gfx942 MI300A reports six XCDs while these kernels bake eight
// per-XCD queues; launching one is invalid, but static StreamK and
// non-StreamK solutions in the same config remain runnable.
if(!solution.streamKDynamicQueueSupported(problem, *m_hardware))
{
if(isReportValid)
m_reporter->report(ResultKey::Validation, "UNSUPPORTED_XCD_TOPOLOGY");
return false;
}

// Test if the persistent kernel is eligible for the current hw and solution
problem.checkPersistentKernelEligibility(solution, *m_hardware);
Task task(*m_hardware, problem, solution);
Expand All @@ -139,6 +125,15 @@ namespace TensileLite
return false;
}

// Explicit index selection bypasses SolutionLibrary; apply its
// StreamK topology predicate after the same problem/task checks.
if(!solution.streamKDynamicQueueSupported(problem, *m_hardware))
{
if(isReportValid)
m_reporter->report(ResultKey::Validation, "UNSUPPORTED_XCD_TOPOLOGY");
return false;
}

if(solution.requiredHostWorkspaceSizePerProblem == static_cast<size_t>(-1))
{
solution.requiredHostWorkspaceSizePerProblem
Expand Down
14 changes: 10 additions & 4 deletions projects/hipblaslt/tensilelite/include/Tensile/Debug.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,6 +47,13 @@ namespace TensileLite
class TENSILELITEHOST_EXPORT Debug : public LazySingleton<Debug>
{
public:
// Defined out of line rather than inherited from LazySingleton. That base
// holds the instance in a function-local static of an inline template
// member, which VISIBILITY_INLINES_HIDDEN gives hidden visibility, so a
// shared tensilelite-host and its consumer each get their own Debug and
// one side's reloadDebugBitsForTest() cannot be seen by the other.
static Debug& Instance();

bool printPropertyEvaluation() const;
bool printPredicateEvaluation() const;
bool printPredicateEvaluationVerbose() const;
Expand Down Expand Up @@ -185,11 +192,10 @@ namespace TensileLite
bool m_printMarker = false;
bool m_disableStaggerU = false;
// -1 = unset (use API attribute); 0 = force static SK3; 1 = force dynamic SK4
int m_streamK5ForceMode = -1;
bool m_usePreciseSMTarget = false;
StringSet m_excludedFromGetAll;
int m_streamK5ForceMode = -1;
bool m_usePreciseSMTarget = false;
StringSet m_excludedFromGetAll;

Debug();
};
} // namespace TensileLite

11 changes: 5 additions & 6 deletions projects/hipblaslt/tensilelite/include/Tensile/KernelLanguageTypes.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,8 +59,8 @@ namespace TensileLite
Count
};

TENSILELITEHOST_EXPORT std::string ToString(KernelLanguage d);
TENSILELITEHOST_EXPORT std::string TypeAbbrev(KernelLanguage d);
TENSILELITEHOST_EXPORT std::string ToString(KernelLanguage d);
TENSILELITEHOST_EXPORT std::string TypeAbbrev(KernelLanguage d);
TENSILELITEHOST_EXPORT std::ostream& operator<<(std::ostream& stream, KernelLanguage const& t);
TENSILELITEHOST_EXPORT std::istream& operator>>(std::istream& stream, KernelLanguage& t);

Expand All @@ -70,9 +70,9 @@ namespace TensileLite
*/
struct KernelLanguageTypeInfo
{
static KernelLanguageTypeInfo const& Get(int index);
static KernelLanguageTypeInfo const& Get(KernelLanguage t);
static KernelLanguageTypeInfo const& Get(std::string const& str);
TENSILELITEHOST_EXPORT static KernelLanguageTypeInfo const& Get(int index);
TENSILELITEHOST_EXPORT static KernelLanguageTypeInfo const& Get(KernelLanguage t);
TENSILELITEHOST_EXPORT static KernelLanguageTypeInfo const& Get(std::string const& str);

KernelLanguage m_kernelLanguage;
std::string name;
Expand Down Expand Up @@ -148,4 +148,3 @@ namespace std
}
};
} // namespace std

11 changes: 6 additions & 5 deletions projects/hipblaslt/tensilelite/include/Tensile/MLFeatures.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -59,11 +59,13 @@ namespace TensileLite
float waveScale;
};

float tilesPerCU(ContractionProblemGemm const& problem,
CUGranularityScaleFactors const& cuFactors);
TENSILELITEHOST_EXPORT float tilesPerCU(ContractionProblemGemm const& problem,
CUGranularityScaleFactors const& cuFactors);

std::ostream& operator<<(std::ostream& stream, CUGranularityScaleFactors const& cugsf);
std::ostream& operator<<(std::ostream& stream, WaveGranularityScaleFactors const& wgsf);
TENSILELITEHOST_EXPORT std::ostream& operator<<(std::ostream& stream,
CUGranularityScaleFactors const& cugsf);
TENSILELITEHOST_EXPORT std::ostream& operator<<(std::ostream& stream,
WaveGranularityScaleFactors const& wgsf);

/**
* @brief A Property whose value is of type `float`.
Expand Down Expand Up @@ -257,4 +259,3 @@ namespace TensileLite
*/
} // namespace MLFeatures
} // namespace TensileLite

9 changes: 5 additions & 4 deletions projects/hipblaslt/tensilelite/include/Tensile/MLPClassification.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,7 +73,8 @@ namespace TensileLite
struct DenseLayer
{
DenseLayer() = default;
DenseLayer(const std::vector<float>& weights, const std::vector<float>& bias);
TENSILELITEHOST_EXPORT DenseLayer(const std::vector<float>& weights,
const std::vector<float>& bias);

std::vector<dtype> operator()(const std::vector<dtype>& F) const;

Expand Down Expand Up @@ -101,9 +102,10 @@ namespace TensileLite
{
MLPNet() = default;

std::vector<dtype> predict(std::vector<float> const& probkey) const;
TENSILELITEHOST_EXPORT std::vector<dtype>
predict(std::vector<float> const& probkey) const;

bool valid(bool verbose = false) const;
TENSILELITEHOST_EXPORT bool valid(bool verbose = false) const;

std::string description() const
{
Expand All @@ -117,4 +119,3 @@ namespace TensileLite

} // namespace MLPClassification
} // namespace TensileLite

11 changes: 5 additions & 6 deletions projects/hipblaslt/tensilelite/include/Tensile/ScalarValueTypes.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -61,8 +61,8 @@ namespace TensileLite
Count
};

TENSILELITEHOST_EXPORT std::string ToString(ScalarValue d);
TENSILELITEHOST_EXPORT std::string TypeAbbrev(ScalarValue d);
TENSILELITEHOST_EXPORT std::string ToString(ScalarValue d);
TENSILELITEHOST_EXPORT std::string TypeAbbrev(ScalarValue d);
TENSILELITEHOST_EXPORT std::ostream& operator<<(std::ostream& stream, ScalarValue const& t);
TENSILELITEHOST_EXPORT std::istream& operator>>(std::istream& stream, ScalarValue& t);

Expand All @@ -72,9 +72,9 @@ namespace TensileLite
*/
struct ScalarValueTypeInfo
{
static ScalarValueTypeInfo const& Get(int index);
static ScalarValueTypeInfo const& Get(ScalarValue t);
static ScalarValueTypeInfo const& Get(std::string const& str);
TENSILELITEHOST_EXPORT static ScalarValueTypeInfo const& Get(int index);
TENSILELITEHOST_EXPORT static ScalarValueTypeInfo const& Get(ScalarValue t);
TENSILELITEHOST_EXPORT static ScalarValueTypeInfo const& Get(std::string const& str);

ScalarValue m_value;
std::string name;
Expand Down Expand Up @@ -189,4 +189,3 @@ namespace std
}
};
} // namespace std

13 changes: 10 additions & 3 deletions projects/hipblaslt/tensilelite/src/Debug.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -205,6 +205,13 @@ namespace TensileLite
m_excludedFromGetAll = excludedSet;
}

Debug& Debug::Instance()
{
static Debug instance;

return instance;
}

Debug::Debug()
: m_value(DEBUG_SM)
, m_value2(DEBUG_SM2)
Expand All @@ -230,7 +237,7 @@ namespace TensileLite
const char* sk5Force = std::getenv("TENSILE_STREAMK5_FORCE_MODE");
if(sk5Force)
{
char* end = nullptr;
char* end = nullptr;
const long val = strtol(sk5Force, &end, 0);
if(end != sk5Force && *end == '\0' && val >= -1 && val <= 1)
m_streamK5ForceMode = static_cast<int>(val);
Expand Down Expand Up @@ -296,8 +303,8 @@ namespace TensileLite
const char* db2 = std::getenv("TENSILE_DB2");
m_value2 = db2 ? static_cast<int>(strtol(db2, nullptr, 0)) : DEBUG_SM2;

const char* sk5Force = std::getenv("TENSILE_STREAMK5_FORCE_MODE");
m_streamK5ForceMode = -1;
const char* sk5Force = std::getenv("TENSILE_STREAMK5_FORCE_MODE");
m_streamK5ForceMode = -1;
if(sk5Force)
{
char* end = nullptr;
Expand Down
9 changes: 9 additions & 0 deletions projects/hipblaslt/tensilelite/tests/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -82,6 +82,15 @@ endif()

target_link_libraries(tensilelite-tests PUBLIC GTest::gtest)

# In YAML builds the serialization templates are instantiated in these test
# translation units, not only inside tensilelite-host. tensilelite-host links
# LLVMObjectYAML PRIVATE and, when shared, is linked with --exclude-libs,ALL,
# so it does not re-export those symbols. Link them here as well.
if(HIPBLASLT_ENABLE_YAML)
target_link_libraries(tensilelite-tests PRIVATE LLVMObjectYAML)
target_include_directories(tensilelite-tests PRIVATE ${LLVM_INCLUDE_DIRS})
endif()

if(HIPBLASLT_ENABLE_OPENMP)
target_link_libraries(tensilelite-tests PRIVATE OpenMP::OpenMP_CXX)
endif()
Expand Down
Loading
Loading