Skip to content

feat(tensilelite): add gfx950 mixed MX custom kernels #12965

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
jaopaulolc merged 3 commits into from
Oct 17, 2026
Merged

feat(tensilelite): add gfx950 mixed MX custom kernels #12965

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
2bdf0dd
feat(tensilelite): add gfx950 mixed MX rocRoller custom kernels
jaopaulolc Oct 2, 2026
9b18e9a
style(tensilelite): use the short SPDX header on new rocRoller files
jaopaulolc Oct 13, 2026
6818222
fix(tensilelite): make mixed MX custom kernels run in the client
jaopaulolc Oct 15, 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
22 changes: 21 additions & 1 deletion projects/hipblaslt/tensilelite/client/include/TypedId.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -330,7 +330,7 @@ namespace TensileLite
using TypedGemm_B6F8_S_S = TypedGemm<BFloat6x32, Float8, float, float, float, float, BFloat6x32, Float8>;
using TypedGemm_B8B6_S_S = TypedGemm<BFloat8, BFloat6x32, float, float, float, float, BFloat8, BFloat6x32>;
using TypedGemm_B6B8_S_S = TypedGemm<BFloat6x32, BFloat8, float, float, float, float, BFloat6x32, BFloat8>;
// Half/BFloat16 dest for pure 6-bit and F6/B6 mixed gemms (RR custom kernels)
// Half/BFloat16 dest for 6-bit and 6-bit/other mixed gemms (RR custom kernels)
using TypedGemm_F6_H_S = TypedGemm<Float6x32, Float6x32, Half, Half, float, float, Float6x32, Float6x32>;
using TypedGemm_F6_B_S = TypedGemm<Float6x32, Float6x32, BFloat16, BFloat16, float, float, Float6x32, Float6x32>;
using TypedGemm_BF6_H_S = TypedGemm<BFloat6x32, BFloat6x32, Half, Half, float, float, BFloat6x32, BFloat6x32>;
Expand All @@ -339,5 +339,25 @@ namespace TensileLite
using TypedGemm_F6B6_B_S = TypedGemm<Float6x32, BFloat6x32, BFloat16, BFloat16, float, float, Float6x32, BFloat6x32>;
using TypedGemm_B6F6_H_S = TypedGemm<BFloat6x32, Float6x32, Half, Half, float, float, BFloat6x32, Float6x32>;
using TypedGemm_B6F6_B_S = TypedGemm<BFloat6x32, Float6x32, BFloat16, BFloat16, float, float, BFloat6x32, Float6x32>;
using TypedGemm_F6F4_H_S = TypedGemm<Float6x32, Float4x2, Half, Half, float, float, Float6x32, Float4x2>;
using TypedGemm_F4F6_H_S = TypedGemm<Float4x2, Float6x32, Half, Half, float, float, Float4x2, Float6x32>;
using TypedGemm_F6F8_H_S = TypedGemm<Float6x32, Float8, Half, Half, float, float, Float6x32, Float8>;
using TypedGemm_F6F8_B_S = TypedGemm<Float6x32, Float8, BFloat16, BFloat16, float, float, Float6x32, Float8>;
using TypedGemm_F8F6_H_S = TypedGemm<Float8, Float6x32, Half, Half, float, float, Float8, Float6x32>;
using TypedGemm_F8F6_B_S = TypedGemm<Float8, Float6x32, BFloat16, BFloat16, float, float, Float8, Float6x32>;
using TypedGemm_F6B8_H_S = TypedGemm<Float6x32, BFloat8, Half, Half, float, float, Float6x32, BFloat8>;
using TypedGemm_F6B8_B_S = TypedGemm<Float6x32, BFloat8, BFloat16, BFloat16, float, float, Float6x32, BFloat8>;
using TypedGemm_B8F6_H_S = TypedGemm<BFloat8, Float6x32, Half, Half, float, float, BFloat8, Float6x32>;
using TypedGemm_B8F6_B_S = TypedGemm<BFloat8, Float6x32, BFloat16, BFloat16, float, float, BFloat8, Float6x32>;
using TypedGemm_B6F4_H_S = TypedGemm<BFloat6x32, Float4x2, Half, Half, float, float, BFloat6x32, Float4x2>;
using TypedGemm_F4B6_H_S = TypedGemm<Float4x2, BFloat6x32, Half, Half, float, float, Float4x2, BFloat6x32>;
using TypedGemm_B6F8_H_S = TypedGemm<BFloat6x32, Float8, Half, Half, float, float, BFloat6x32, Float8>;
using TypedGemm_B6F8_B_S = TypedGemm<BFloat6x32, Float8, BFloat16, BFloat16, float, float, BFloat6x32, Float8>;
using TypedGemm_F8B6_H_S = TypedGemm<Float8, BFloat6x32, Half, Half, float, float, Float8, BFloat6x32>;
using TypedGemm_F8B6_B_S = TypedGemm<Float8, BFloat6x32, BFloat16, BFloat16, float, float, Float8, BFloat6x32>;
using TypedGemm_B6B8_H_S = TypedGemm<BFloat6x32, BFloat8, Half, Half, float, float, BFloat6x32, BFloat8>;
using TypedGemm_B6B8_B_S = TypedGemm<BFloat6x32, BFloat8, BFloat16, BFloat16, float, float, BFloat6x32, BFloat8>;
using TypedGemm_B8B6_H_S = TypedGemm<BFloat8, BFloat6x32, Half, Half, float, float, BFloat8, BFloat6x32>;
using TypedGemm_B8B6_B_S = TypedGemm<BFloat8, BFloat6x32, BFloat16, BFloat16, float, float, BFloat8, BFloat6x32>;
#endif // !_WIN32
} // namespace TensileLite
102 changes: 101 additions & 1 deletion projects/hipblaslt/tensilelite/client/src/Reference.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3085,7 +3085,7 @@ namespace TensileLite
problem, inputs, elementsToValidate);
}
#ifndef _WIN32
// Half/BFloat16 dest for pure 6-bit and F6/B6 mixed gemms
// Half/BFloat16 dest for 6-bit and 6-bit/other mixed gemms
case TypedGemm_F6_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_F6_H_S, float>::SolveCPU(
Expand Down Expand Up @@ -3126,6 +3126,106 @@ namespace TensileLite
return ReferenceSolution<TypedGemm_B6F6_B_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_F6F4_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_F6F4_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_F4F6_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_F4F6_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_F6F8_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_F6F8_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_F6F8_B_S::TypeId():
{
return ReferenceSolution<TypedGemm_F6F8_B_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_F8F6_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_F8F6_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_F8F6_B_S::TypeId():
{
return ReferenceSolution<TypedGemm_F8F6_B_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_F6B8_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_F6B8_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_F6B8_B_S::TypeId():
{
return ReferenceSolution<TypedGemm_F6B8_B_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_B8F6_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_B8F6_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_B8F6_B_S::TypeId():
{
return ReferenceSolution<TypedGemm_B8F6_B_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_B6F4_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_B6F4_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_F4B6_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_F4B6_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_B6F8_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_B6F8_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_B6F8_B_S::TypeId():
{
return ReferenceSolution<TypedGemm_B6F8_B_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_F8B6_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_F8B6_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_F8B6_B_S::TypeId():
{
return ReferenceSolution<TypedGemm_F8B6_B_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_B6B8_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_B6B8_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_B6B8_B_S::TypeId():
{
return ReferenceSolution<TypedGemm_B6B8_B_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_B8B6_H_S::TypeId():
{
return ReferenceSolution<TypedGemm_B8B6_H_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
case TypedGemm_B8B6_B_S::TypeId():
{
return ReferenceSolution<TypedGemm_B8B6_B_S, float>::SolveCPU(
problem, inputs, elementsToValidate);
}
#endif // !_WIN32
default:;
}
Expand Down
Loading
Loading