Reverse-Engineering hipBLASLt TensileLite Kernels: From Solution Name to a Tuning Config

27 min read. | 6721 total words.

In a previous blog, Customizing Kernels with hipBLASLt TensileLite GEMM Tuning, we showed how TensileLite Tuning generates brand-new GEMM kernels by searching a parameter space and selecting the fastest valid candidate for a given problem size. The search space is defined entirely by the tuning configuration YAML you provide. Its ForkParameters lists are expanded into a Cartesian product, and the tuning run can only ever return a kernel that is expressible as one combination within that product.

That property cuts both ways. TensileLite can generate genuinely new kernels, but it is only as good as the search space you hand it. If the parameter combinations in your configuration cannot reproduce the best kernel you already have, that kernel is not even a candidate, and the run can return something slower than where you started. The strongest kernel you currently hold is usually the one Offline Tuning selected from the existing solution pool for that shape, so the real question is sharper than “can tuning find a fast kernel?” It is “can tuning guarantee a kernel at least as fast as the pool’s current best?”

Reverse-engineer the current best kernel into a pinned configuration. Decode the pool-best kernel’s solution name back into its exact parameters and write them as single-value ForkParameters. The resulting search space is exactly one point: that kernel. This is the anchor the run can never fall below.

Reverse-engineer the current best kernel into a pinned configuration. Decode the pool-best kernel’s solution name back into its exact parameters and write them as single-value ForkParameters. The resulting search space is exactly one point: that kernel. This is the anchor the run can never fall below.

Expand the search space outward from that anchor. Once the pinned configuration is in hand, you can widen the parameters to look for something faster, while keeping the anchor kernel inside the search space so the floor never drops.

Expand the search space outward from that anchor. Once the pinned configuration is in hand, you can widen the parameters to look for something faster, while keeping the anchor kernel inside the search space so the floor never drops.

Together they guarantee the result never regresses: the search space always contains the pool’s best kernel, so the kernel TensileLite generates is at least as fast as the pool’s current best, and usually faster. This blog focuses on the first idea, the reverse-engineering that establishes the anchor; the expansion strategy is a topic for a follow-up.

Treat the tuning configuration as the definition of a set: its ForkParameters, with the problem type, enumerate a finite set of candidate kernels, and a run keeps the fastest valid member. One hard constraint governs everything:

A TensileLite run can only ever return a kernel that is a member of the search space its configuration defines.

A TensileLite run can only ever return a kernel that is a member of the search space its configuration defines.

Offline Tuning, by contrast, generates nothing; it selects the best existing kernel from the pool for your shape, which is by definition your current floor. Combine the two and the requirement is unavoidable:

If your current best kernel is not a member of the configuration’s search space, the TensileLite run cannot produce it, and there is no guarantee the winner will be at least as fast. The outcome could easily be a regression.

If your current best kernel is not a member of the configuration’s search space, the TensileLite run cannot produce it, and there is no guarantee the winner will be at least as fast. The outcome could easily be a regression.

If your current best kernel is a member of the search space, then in the worst case the run reproduces it and ties, and in the better case it finds a neighbor that is faster. Either way you cannot lose.

If your current best kernel is a member of the search space, then in the worst case the run reproduces it and ties, and in the better case it finds a neighbor that is faster. Either way you cannot lose.

Figure 1 makes this contrast concrete and previews the contribution of the method. A generic tuning configuration is built from a fixed template, blind to your golden kernel, so the golden can fall outside its search space (left); the run then returns the fastest kernel it happens to contain, which may be slower than the golden. Reverse-engineering instead builds the search space around the golden and keeps it inside (right), so the run can only match or beat it.

Figure 1. A TensileLite run can only return a kernel inside the search space its configuration defines. Left: a generic configuration is built from a template without reference to the golden, so the pool’s best kernel can fall outside the search space and be unreachable, and the best kernel found may be slower than the golden. Right: reverse-engineering builds the search space around the golden and keeps it inside, so the tuned result can only match or beat it.#

Figure 1. A TensileLite run can only return a kernel inside the search space its configuration defines. Left: a generic configuration is built from a template without reference to the golden, so the pool’s best kernel can fall outside the search space and be unreachable, and the best kernel found may be slower than the golden. Right: reverse-engineering builds the search space around the golden and keeps it inside, so the tuned result can only match or beat it.#

So the first goal is precise: make the configuration’s search space a superset that includes your current best kernel. The catch is that you usually only have that kernel’s solution name from the hipBLASLt logs, not the configuration that produced it, and user-tuned kernels are frequently absent from the committed library logic. The only general way to recover its parameters is to decode the name itself, and because a TensileLite solution name is a deterministic encoding of the kernel’s solution parameters rather than an opaque hash, the decode is exact. That decode is the subject of the rest of this blog.

Figure 2 summarizes the two-stage idea: reverse-engineering pins the pool’s best kernel inside the search space, and any later expansion keeps it there, so the result can only match or beat it.

Figure 2. The tuning configuration defines a search space (the outer region): the set of all kernels expressible as a combination of its ForkParameters. A TensileLite run can only return a kernel inside this region. Stage 1 places the pool’s best kernel inside the region by reverse-engineering it; Stage 2 grows the region outward while keeping that kernel inside, so the run can only match or beat it.#

Figure 2. The tuning configuration defines a search space (the outer region): the set of all kernels expressible as a combination of its ForkParameters. A TensileLite run can only return a kernel inside this region. Stage 1 places the pool’s best kernel inside the region by reverse-engineering it; Stage 2 grows the region outward while keeping that kernel inside, so the run can only match or beat it.#

The goal is to turn the pool’s best kernel, which you only know as a solution name, into a pinned tuning configuration whose search space is exactly that one kernel. This is the heart of the method and the part you must get exactly right, because pinning the current best inside the search space is precisely what makes the guarantee hold. The following sections work through the decode end to end: the structure of a solution name, the rules that map each token back to a parameter, the reconstruction of the matrix instruction, and the recovery of the problem type. The output is a configuration you can hand to TensileLite as-is, or use as the starting point for the expansion step, which is a topic for a follow-up.

A TensileLite solution name (the SolutionNameMin string) is produced by the naming logic in tensilelite/Tensile/SolutionStructs/Naming.py. Consider this example:

It has three logical regions.

1. The Problem-Type Prefix#

The leading tokens describe the operation and its features, not the tunable kernel knobs. They come from ProblemType.__str__ in tensilelite/Tensile/SolutionStructs/Problem.py:

Cijk_Alik_Bljk encodes the index assignments (the transpose layout).

Cijk_Alik_Bljk encodes the index assignments (the transpose layout).

BBS is the datatype signature. Each letter is a datatype character: B for bfloat16, H for float16, S for float32, and so on. The trailing letters capture destination and compute types.

BBS is the datatype signature. Each letter is a datatype character: B for bfloat16, H for float16, S for float32, and so on. The trailing letters capture destination and compute types.

These prefix tokens map back to ProblemType fields, which belong in the BenchmarkProblems problem type, not in ForkParameters.

2. The Tile and Matrix-Instruction Anchors#

MI{M}x{N}x{B}, for example MI16x16x1.

MI{M}x{N}x{B}, for example MI16x16x1.

Note that the MI token encodes only M, N, and B. It does not carry the matrix-instruction K dimension. We reconstruct K from the datatype and the MI dimensions, as described below.

3. The Solution-Parameter Tokens#

Everything after the anchors is a sequence of underscore-separated tokens, one per non-default solution parameter. Each token is a parameter-name abbreviation immediately followed by a value abbreviation, so decoding means reversing two rules in turn.

PrefetchGlobalRead keeps P, G, R, giving PGR.

PrefetchGlobalRead keeps P, G, R, giving PGR.

ScheduleIterAlg keeps S, I, A, giving SIA.

ScheduleIterAlg keeps S, I, A, giving SIA.

The rule also explains cryptic tokens. UseSgprForGRO keeps the uppercase letters of every word, including the F of For, giving USFGRO. Use64bShadowLimit gives USL (digits and the lowercase b drop). A leading digit counts as non-uppercase, so 1LDSBuffer reduces to LDSB.

Because the rule is lossy, different names can collapse to the same prefix; see disambiguation below.

The suffix of a token encodes the value. The encoding depends on the value’s type, summarized in Table 1:

Table 1. Value-abbreviation encodings by value type.

The string rule stacks two abbreviations: GlobalSplitUAlgorithm gives GSUA and its value MultipleBuffer gives MB, so the token is GSUAMB (likewise GSUAMBSK for MultipleBufferSingleKernel and GSUASB for SingleBuffer). Read the value by matching the trailing letters against that parameter’s known values, not a generic dictionary.

Because prefix and value are concatenated with no separator, split a token by longest-prefix matching against known parameter abbreviations. In GSUAMBSK, a naive match stops at GSU (GlobalSplitU) and leaves the meaningless AMBSK; the correct split takes the longer GSUA (GlobalSplitUAlgorithm) and reads MBSK as MultipleBufferSingleKernel. Always prefer the longest valid abbreviation, then read the remainder as the value.

Some parameters share a prefix because the rule is lossy. AssertFree0ElementMultiple and AssertFree1ElementMultiple both reduce to AFEM, so a name carries two, for example AFEM1_AFEM1. They are emitted in sorted key order, so the first is AssertFree0ElementMultiple and the second AssertFree1ElementMultiple; preserve that order when decoding.

The name lists only parameters whose value differs from the default in the reduced (“min”) set used for naming; everything absent takes its default. That is exactly the shape a tuning configuration wants, so a faithful decode lists only the non-default knobs and does not pad with defaults, which would add noise without changing the kernel.

Table 2 groups the abbreviations you will most often encounter on gfx942 GEMM kernels. It is not exhaustive, but combined with the name-abbreviation rule it lets you decode essentially any token by inspection. Each row lists the token prefix and the parameter it denotes.

LdsPadA / ...B / ...Metadata

LdsPadA / ...B / ...Metadata

VectorWidthA / ...B

VectorWidthA / ...B

NonTemporalA / ...B / ...C / ...D / ...Metadata

NonTemporalA / ...B / ...C / ...D / ...Metadata

Table 2. Common token prefixes on gfx942 GEMM kernels and the parameters they denote.

When you meet a token that is not in this table, apply the name-abbreviation rule in reverse: list the candidate parameters whose uppercase letters match the prefix, pick the longest valid match, and read the remainder as the value with the type-based rules above.

The single most important field to recover correctly is MatrixInstruction, because it pins the tile shape and the wave tiling. TensileLite accepts the canonical nine-element form:

M, N, and B come from the MI token.

M, N, and B come from the MI token.

MIBlockM defaults to 1.

MIBlockM defaults to 1.

K is the one value not in the name; it is fixed by the hardware MFMA instruction and must be looked up, as the next subsection covers.

K is the one value not in the name; it is fixed by the hardware MFMA instruction and must be looked up, as the next subsection covers.

One detail is easy to miss: the nine-element form does not carry LocalSplitU. TensileLite reads it from the third element of WorkGroup (see matrixInstructionToMIParameters in tensilelite/Tensile/SolutionStructs/Validators/MatrixInstruction.py), so the recovered configuration must keep WorkGroup; dropping it silently loses LocalSplitU for any kernel that splits the summation locally.

K is fixed by the MFMA (Matrix Fused Multiply-Add) instruction the kernel issues, which is fully determined by the architecture, the tile dimensions M and N, and the input datatype. The authoritative source is the AMD CDNA ISA documentation: each MFMA is named V_MFMA_<accumulator>_<M>X<N>X<K>_<input-type>, spelling out M, N, K, and the datatype. A lookup keyed by (architecture, M, N, datatype), built from that instruction set, yields an exact K for every supported shape.

The critical subtlety is that K depends on M and N, not the datatype alone: the 16x16 and 32x32 instructions of one datatype use different K. Table 3 lists the dense MFMA K values for gfx942 (CDNA3):

Table 3. Matrix-instruction K by instruction shape and datatype on gfx942.

Reusing the 16x16 row for 32x32 is an easy mistake; the values really do differ. Newer architectures make this harder: on gfx950 (CDNA4), double-rate variants let one (M, N, datatype) map to two valid K values (a 16x16 bfloat16 instruction can use K=16 or K=32), and microscaling formats add more shapes, so there you need an extra signal such as DepthU (a multiple of K) to disambiguate. For unique-K architectures like gfx942, the lookup is unambiguous.

The hipblaslt-bench command provides the datatypes, transpose flags, and problem size. For example, --a_type bf16_r maps to B, --transA T --transB N maps to a TN layout, and -m 96 -n 12 -k 5120 --batch_count 1 becomes the exact size [96, 12, 1, 5120].

The hipblaslt-bench command provides the datatypes, transpose flags, and problem size. For example, --a_type bf16_r maps to B, --transA T --transB N maps to a TN layout, and -m 96 -n 12 -k 5120 --batch_count 1 becomes the exact size [96, 12, 1, 5120].

The solution-name prefix provides the features that the bench command alone cannot always convey, such as SwizzleTensorA, UseScaleAB, UseScaleAlphaVec, bias support, and the activation set.

The solution-name prefix provides the features that the bench command alone cannot always convey, such as SwizzleTensorA, UseScaleAB, UseScaleAlphaVec, bias support, and the activation set.

On gfx942, the 8-bit float types are the FNUZ variants, so f8_r maps to the F8N character. When the input datatype differs from the compute datatype, HighPrecisionAccumulate is enabled.

One nuance is worth calling out, because it changes which kernel you target. The epilogue features encoded in the name, such as the activation set or the scale modes, often describe a general-purpose UserArgs kernel that can switch those features off at runtime. In many solution pools there is no narrower kernel for a specific epilogue. For instance, a bias kernel may always be emitted with the activation and scale-alpha-vector features attached, so a pure bias-only name simply does not exist. This gives you two legitimate targets:

Reproduce the exact golden kernel. Enrich the problem type with every epilogue feature encoded in the name. The configuration then pins the precise kernel you observed, including features the benchmark may not actually exercise.

Reproduce the exact golden kernel. Enrich the problem type with every epilogue feature encoded in the name. The configuration then pins the precise kernel you observed, including features the benchmark may not actually exercise.

Target the leanest kernel for the real workload. Build the problem type from the benchmark command alone and ignore the name’s epilogue tokens. Tuning then generates the kernel whose prefix matches the benchmark’s true needs, which can be leaner than the fat UserArgs kernel.

Target the leanest kernel for the real workload. Build the problem type from the benchmark command alone and ignore the name’s epilogue tokens. Tuning then generates the kernel whose prefix matches the benchmark’s true needs, which can be leaner than the fat UserArgs kernel.

Either choice keeps the core solution parameters (tile, matrix instruction, wave tile, and the rest) from the name; they differ only in which epilogue features the problem type declares.

The rules above are enough to decode a real kernel end to end with nothing but the name and the benchmark command. Consider this benchmark command and the solution name it selected on gfx942:

Step 1, read the prefix into a problem type. Cijk_Alik_Bljk is the TN layout (--transA T --transB N); BBS is bfloat16 input with bfloat16 output and float32 compute (--a_type bf16_r ... --compute_type f32_r); STA sets SwizzleTensorA (--swizzleA); BH sets UseBeta and HighPrecisionAccumulate. With no Bias, SAV, or SABV token, no bias or scaling is enabled. The size comes from the command as Exact: [96, 12, 1, 5120] in [M, N, batch, K] order.

Step 2, read the anchors. MT32x16x256 gives MacroTile0 = 32, MacroTile1 = 16, and DepthU = 256. MI16x16x1 gives matrix-instruction M = 16, N = 16, B = 1.

Step 3, rebuild the matrix instruction. The wave tile token is absent, so MIWaveTile is [1, 1] and MIBlockM is 1. K is not in the name; a 16x16 bfloat16 instruction on gfx942 gives K = 16 (a 32x32 one would use K = 8, which is why the lookup keys on M and N, not the datatype alone). From the macro tile, WaveM = 32 / (16 * 1) = 2 and WaveN = 16 / (16 * 1) = 1, giving MatrixInstruction = [16, 16, 16, 1, 1, 1, 1, 2, 1]; the checks 16 * 1 * 2 = 32 and 16 * 1 * 1 = 16 match the macro tile.

Step 4, decode the remaining tokens with the rules. Walking left to right and applying longest-prefix matching plus the value rules:

PGR2 is PrefetchGlobalRead = 2.

PGR2 is PrefetchGlobalRead = 2.

PLR1 is PrefetchLocalRead = 1.

PLR1 is PrefetchLocalRead = 1.

SIA3 is ScheduleIterAlg = 3.

SIA3 is ScheduleIterAlg = 3.

SS1 is SourceSwap = true.

SS1 is SourceSwap = true.

WS64 is WavefrontSize = 64.

WS64 is WavefrontSize = 64.

WGM1 is WorkGroupMapping = 1.

WGM1 is WorkGroupMapping = 1.

Step 5, assemble the configuration. Place the problem-type fields under the BenchmarkProblems problem type, the exact size under BenchmarkFinalParameters, and every decoded solution parameter as a single-value list under ForkParameters. Crucially, keep WorkGroup alongside the nine-element MatrixInstruction, because its third element carries LocalSplitU. The result is a configuration whose entire search space is exactly one kernel: the one you started from.

Because the ForkParameters are single-valued, feeding this configuration into the TensileLite workflow reproduces the exact kernel. To explore its neighborhood while preserving the guarantee, widen a few axes into multi-value lists, for example DepthU: [128, 256] or GlobalSplitU: [1, 2, 4], and leave the rest fixed. The original kernel remains a member of the enlarged search space, so the run can only match or beat it.

This technique is deterministic, but a few practical limits are worth stating clearly.

K is recovered from the ISA, not stored in the name. The matrix-instruction K dimension is fixed by the hardware MFMA instruction, so it is recovered by looking up (architecture, M, N, datatype) against the architecture’s matrix-instruction set. On unique-K architectures such as gfx942 this is exact. On gfx950, double-rate and microscaling instructions make K ambiguous from the name alone, so an extra signal such as DepthU is required to disambiguate. When the kernel exists in the library logic, the logic file also gives the authoritative value.

K is recovered from the ISA, not stored in the name. The matrix-instruction K dimension is fixed by the hardware MFMA instruction, so it is recovered by looking up (architecture, M, N, datatype) against the architecture’s matrix-instruction set. On unique-K architectures such as gfx942 this is exact. On gfx950, double-rate and microscaling instructions make K ambiguous from the name alone, so an extra signal such as DepthU is required to disambiguate. When the kernel exists in the library logic, the logic file also gives the authoritative value.

Version drift in tokens. Names produced by a different TensileLite version can contain tokens that the current parameter set does not recognize, for example a stray marker that no longer maps to any parameter. Treat unrecognized tokens as a signal to check the version, skip the ones with no mapping, and fold derived load tokens back into their source parameter where possible.

Version drift in tokens. Names produced by a different TensileLite version can contain tokens that the current parameter set does not recognize, for example a stray marker that no longer maps to any parameter. Treat unrecognized tokens as a signal to check the version, skip the ones with no mapping, and fold derived load tokens back into their source parameter where possible.

Defaults are intentional. A decoded configuration lists only the non-default knobs. Do not pad it with default values. That subset is the faithful tuning input.

Defaults are intentional. A decoded configuration lists only the non-default knobs. Do not pad it with default values. That subset is the faithful tuning input.

Architecture-specific datatype mapping. The 8-bit float mapping is architecture dependent. On gfx942 the FNUZ variants apply; other architectures use different encodings.

Architecture-specific datatype mapping. The 8-bit float mapping is architecture dependent. On gfx942 the FNUZ variants apply; other architectures use different encodings.

gfx942 is not a single device. The gfx942 ISA string is shared by several SKUs (for example MI300A, MI300X, MI308X, and MI325X) and partition modes, and Tensile dispatches by CU count and DeviceNames, not by the ISA string. So the pinned configuration must set its LibraryLogic CUCount and DeviceNames to the SKU and partition you actually deploy on, and you should re-check WorkGroupMappingXCC (whose valid values are powers of two that divide the CU count, which changes with the partition). Otherwise you recover a valid but poorly tuned kernel.

gfx942 is not a single device. The gfx942 ISA string is shared by several SKUs (for example MI300A, MI300X, MI308X, and MI325X) and partition modes, and Tensile dispatches by CU count and DeviceNames, not by the ISA string. So the pinned configuration must set its LibraryLogic CUCount and DeviceNames to the SKU and partition you actually deploy on, and you should re-check WorkGroupMappingXCC (whose valid values are powers of two that divide the CU count, which changes with the partition). Otherwise you recover a valid but poorly tuned kernel.

To check the guarantee in practice, we applied the method to a matrix of 72 GEMMs (eight shapes times three layouts times three precisions) on an AMD Instinct MI325X (gfx942, 304 CU, SPX partition). Every case was measured in the same tuning client on the same device, median of three runs. The sweep is now complete, so all 72 cases are reported.

Library default: the kernel the library heuristic picks out of the box.

Library default: the kernel the library heuristic picks out of the box.

Anchor (pool best): the fastest existing kernel for that shape, recovered by reverse-engineering and pinned as described above. Because the library default is always one of the candidates, the anchor is by construction no slower than the default.

Anchor (pool best): the fastest existing kernel for that shape, recovered by reverse-engineering and pinned as described above. Because the library default is always one of the candidates, the anchor is by construction no slower than the default.

Tuned (final): the kernel produced by expanding the search space outward from the anchor.

Tuned (final): the kernel produced by expanding the search space outward from the anchor.

The central prediction of Figure 2 is that latency is monotone: tuned is at least as fast as the anchor, which is at least as fast as the default. This held in all 72 cases, with zero regressions.

Table 4 aggregates the tuned-versus-default latency reduction by precision.

Table 4. Latency reduction of the tuned kernel versus the library default, by precision.

Table 5 breaks the same reductions down by problem shape.

Table 5. Latency reduction of the tuned kernel versus the library default, by problem shape.

Across all 72 cases, the tuned kernel is 18.0% faster than the library default at the median (20.2% faster on average and up to 66.2% faster). Reverse-engineering the anchor already recovers a mean 10.6% over the default, and expanding from that anchor adds a further mean 10.6% on top, improving on the anchor in 70 of the 72 cases while never falling below it. The largest gains appear for float16 and for transpose-N GEMMs with A swizzling; the biggest single case, a small square float16 GEMM with swizzling, drops from 12.20 us to 4.12 us (66.2%).

The complete matrix follows in Table 6, and it doubles as the data for the next section. Latencies are microseconds. Reverse vs default is the tuned kernel’s reduction over the library default (higher is better); Generic vs reverse is how much slower the generic one-shot config, analyzed in the next section, is than the reverse-anchored kernel (positive means slower). The five tiny sqS generic winners were re-measured median-of-five.

Table 6. Full 72-GEMM results on MI325X. Default, golden (reverse-engineered anchor), reverse-anchored tuned, and generic one-shot winner, with the reverse win and the generic gap in the last two columns.

The guarantee above hinges on one thing: the pool’s best kernel being inside the search space. To show what happens without it, we compared reverse-anchored tuning against a generic one-shot config on the same 72-GEMM matrix on MI325X. The generic config is built straight from the official hipBLASLt template: one config per precision and layout, covering all eight shapes with a single shared grid (25 representative power-of-two MacroTiles times DepthU, GlobalSplitU, StaggerU, and WorkGroupMapping), about 10,000 candidate solutions per shape, which is roughly ten times the budget of the reverse-anchored run. The point is that this grid is generic: it is not guaranteed to contain each shape’s golden kernel. (Winner latencies on the tiny 512-cubed shape were re-measured median-of-five to remove a min-over-many-runs bias.)

As Figure 1 anticipated, the search space is a hard ceiling: a generic config that does not contain the golden cannot return it, no matter how large its budget. The results are decisive:

Even with about ten times the budget, the generic config is slower than reverse-anchored tuning in 70 of 72 cases (median +35%, mean +37%, up to +96%). The single exception is a 0.9% noise-level tie.

Even with about ten times the budget, the generic config is slower than reverse-anchored tuning in 70 of 72 cases (median +35%, mean +37%, up to +96%). The single exception is a 0.9% noise-level tie.

It is slower than the golden kernel it was supposed to beat in 61 of 72 cases. This is direct evidence that a generic search space often does not contain the golden.

It is slower than the golden kernel it was supposed to beat in 61 of 72 cases. This is direct evidence that a generic search space often does not contain the golden.

It is not even monotone against the library default: faster on 22 cases but slower on 49 (up to +87%). A generic config provides no floor at all.

It is not even monotone against the library default: faster on 22 cases but slower on 49 (up to +87%). A generic config provides no floor at all.

The gap is largest for swizzled transpose-N GEMMs, the most tile-shape-sensitive case.

The gap is largest for swizzled transpose-N GEMMs, the most tile-shape-sensitive case.

Table 7 counts how often the generic config is slower than each reference point.

Table 7. Number of cases (of 72) where the generic config is slower than each reference point.

Table 8 breaks the generic-versus-reverse gap down by layout.

Table 8. Mean gap of the generic config versus reverse-anchored tuning, by layout.

The per-case numbers are already in the full results table above: its Generic column against the Default, Golden, and Reverse columns. Every positive Generic vs reverse entry is a case where the generic config, despite far more search budget, loses to the reverse-anchored kernel.

Reverse-anchored tuning avoids all of this by construction. Because the golden kernel is pinned into the search space by reverse-engineering, the tuned result satisfies tuned <= golden <= default in every case: strict, monotone, and never worse than where you started. The generic config, lacking that anchor, routinely performs worse than the golden and even worse than the library default. More search budget does not fix this; only putting the golden inside the search space does.

In this blog, you learned why a TensileLite run can only return a kernel its configuration can express, which makes the search space the most important design choice in tuning. If your current best kernel, typically the one Offline Tuning selected from the pool, is not inside that space, a re-tune has no guaranteed floor and may regress.

You then learned a deterministic recipe that removes this risk: decode a kernel’s solution name back into its exact parameters, rebuild the nine-element MatrixInstruction while keeping WorkGroup so that LocalSplitU survives, recover the ProblemType from the name prefix and the hipblaslt-bench command, and pin the result as single-value ForkParameters. Pinning the pool’s best kernel as an anchor and expanding outward without ever dropping it guarantees a tuned kernel at least as fast as the pool’s best, and usually faster. On the 72-GEMM matrix on MI325X, this held with zero regressions and a median latency reduction of 18.0% over the library default.

The decode rests on three ideas worth remembering. First, the name’s prefix maps back to ProblemType fields, while the remaining tokens map back to solution parameters. Second, MatrixInstruction must be rebuilt as the nine-element form, with WorkGroup retained so that LocalSplitU survives. Third, because user-tuned kernels are often absent from the committed library logic, decoding the name directly is the most general recovery path.

This blog focused on establishing the anchor. In a follow-up, we will cover the second half of the method: how to expand the search space outward from the anchor efficiently, deciding which parameters to widen and by how much to reach faster kernels without ever losing the guarantee. Follow the AMD ROCm Blogs for that next installment and for more hipBLASLt and TensileLite tuning deep dives from our team.

We would like to express our thanks to our colleagues Brian Chang, Eveline Chen, Bobo Fang, Bill Ku, Kaiping Lu, and Menghsuan Yang for their insightful feedback and technical assistance.

The information presented in this document is for informational purposes only and may contain technical inaccuracies, omissions, and typographical errors. The information contained herein is subject to change and may be rendered inaccurate for many reasons, including but not limited to product and roadmap changes, component and motherboard version changes, new model and/or product releases, product differences between differing manufacturers, software changes, BIOS flashes, firmware upgrades, or the like. Any computer system has risks of security vulnerabilities that cannot be completely prevented or mitigated. AMD assumes no obligation to update or otherwise correct or revise this information. However, AMD reserves the right to revise this information and to make changes from time to time to the content hereof without obligation of AMD to notify any person of such revisions or changes. THIS INFORMATION IS PROVIDED ‘AS IS.” AMD MAKES NO REPRESENTATIONS OR WARRANTIES WITH RESPECT TO THE CONTENTS HEREOF AND ASSUMES NO RESPONSIBILITY FOR ANY INACCURACIES, ERRORS, OR OMISSIONS THAT MAY APPEAR IN THIS INFORMATION. AMD SPECIFICALLY DISCLAIMS ANY IMPLIED WARRANTIES OF NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR ANY PARTICULAR PURPOSE. IN NO EVENT WILL AMD BE LIABLE TO ANY PERSON FOR ANY RELIANCE, DIRECT, INDIRECT, SPECIAL, OR OTHER CONSEQUENTIAL DAMAGES ARISING FROM THE USE OF ANY INFORMATION CONTAINED HEREIN, EVEN IF AMD IS EXPRESSLY ADVISED OF THE POSSIBILITY OF SUCH DAMAGES. AMD, the AMD Arrow logo, and combinations thereof are trademarks of Advanced Micro Devices, Inc. Other product names used in this publication are for identification purposes only and may be trademarks of their respective companies. © 2026 Advanced Micro Devices, Inc. All rights reserved