opt(gemm): add hipBLASLt algorithm cache and thread-local workspace

[jasainio]

Cache hipblasLtMatmulAlgo_t across repeat GEMM calls to skip expensive heuristic lookups on recurring shapes. Pre-allocate workspace via a thread-local tensor to avoid per-call CUDA malloc overhead.

Changes:

• Add hipblaslt_algo_cache.h with FNV-1a hashed lookup by GEMM config
• Integrate cache into hipblaslt_gemm_impl (store on miss, reuse on hit)
• Replace per-call at::empty workspace with thread-local pre-allocation
• Add unit tests for bit-exactness, multi-shape, and layout correctness
• Add benchmark script for cold-cache vs warm-cache latency comparison

Description

Cache hipblasLtMatmulAlgo_t results across repeated GEMM calls so that expensive hipBLASLt heuristic lookups are only performed once per unique shape/config. On subsequent calls with the same GEMM configuration, the cached algorithm is reused directly. Additionally, the per-call workspace allocation (at::empty) is replaced with a thread-local pre-allocated tensor, eliminating repeated CUDA malloc/free overhead on every GEMM invocation.

Type of change

• Documentation change (change only to the documentation, either a fix or a new content)
• Bug fix (non-breaking change which fixes an issue)
• New feature (non-breaking change which adds functionality)
• Breaking change (fix or feature that would cause existing functionality to not work as expected)
• Infra/Build change
• Code refactoring

Changes

Please list the changes introduced in this PR:

• Add csrc/kernels/gemm/hipblaslt_algo_cache.h implementing a thread-safe algorithm cache with FNV-1a hashing keyed on GEMM config (M, N, K, dtypes, layouts, compute type)
• Integrate the cache into csrc/kernels/gemm/hipblaslt_gemm.cu and csrc/pytorch/gemm/hipblaslt_gemm.cpp -- store algorithm on cache miss, reuse on hit
• Replace per-call at::empty workspace allocation with a thread-local pre-allocated tensor that grows as needed
• Add tests/pytorch/ops/test_algo_cache.py with tests for bit-exact repeatability, multi-shape correctness (no key collisions), BF16 support, and NT/NN layout handling
• Add benchmark/ops/bench_algo_cache.py for measuring cold-cache vs warm-cache GEMM latency across model configurations

Checklist:

• The functionality is complete
• I have commented my code, particularly in hard-to-understand areas
• I have made corresponding changes to the documentation
• My changes generate no new warnings
• I have added tests that prove my fix is effective or that my feature works
• New and existing unit tests pass locally with my changes

[jasainio]

Need to check if this feat is required once #277 is merged

[RuibinCheung]

Generally LGTM. Could you share the performance improvement of hipblaslt gemm after applying this PR ?

[RuibinCheung]

Could you create a new directory named hipblaslt in kernels/gemm and put this header file into this directory?

[jasainio]

Verified end-to-end as part of the integration branch dev/mxfp4-aiter-preshuffle-min @ 026e136 (Flux 12B MXFP4 1000-step exit 0, FP8 delayed 1000-step exit 0). Tip @ 9737552.

[jasainio]

Thanks @RuibinCheung. Addressed both comments:

1. Directory restructure: Done. Moved hipblaslt_algo_cache.h into csrc/kernels/gemm/hipblaslt/.

2. Performance numbers (MI355X, FP8 GEMM, hipBLASLt backend):

The cache eliminates hipblasLtMatmulAlgoGetHeuristic on recurring shapes. The benchmark measures the first call per unique shape (cold = heuristic runs) vs. median of subsequent calls (warm = heuristic skipped):

Model	Shape (M x N x K)	Cold (ms)	Warm (ms)	Saved (us)
Llama-2-7B	4096 x 12288 x 4096	0.70	0.23	469
Llama-3.1-8B	8192 x 6144 x 4096	0.33	0.22	115
Llama-2-70B	4096 x 8192 x 8192	0.40	0.28	121
Llama-2-70B	8192 x 8192 x 8192	0.49	0.48	12

The benefit is most visible on shapes where the heuristic lookup is a meaningful fraction of total dispatch time (small/medium GEMMs: 100-470 us saved per call). For large compute-bound GEMMs, the kernel runtime dominates and warm/cold are nearly identical.

Note: The very first hipBLASLt call for a new dimension-class incurs a one-shot ~1.5s penalty (internal solution library load from disk). This happens once per process regardless of caching. The cache eliminates the heuristic overhead on all subsequent calls for recurring shapes.

The thread-local workspace eliminates a 64-128 MiB at::empty per GEMM call, removing allocator overhead on every call regardless of shape. BF16 GEMMs benefit from both optimizations equally (test suite covers FP8 + BF16 paths).

Additional improvements in this push:

• Added kMaxEntries = 256 cap to prevent unbounded growth under dynamic-shape workloads (e.g., MoE without capacity factor).
• Added TestAlgoCacheOverflow test verifying correctness across cache clear events (300 unique shapes).
• Added hipblaslt_algo_cache_clear() binding for test isolation.
• Added single-stream-per-thread safety comment on the thread-local workspace.

Reproduce:

python benchmark/ops/bench_algo_cache.py --warmup 5 --repeat 50