[sstamenk]

Linked issue or supporting maintainer

Supported by @jeffdaily

BC-breaking?

No

Summary (human written only)

The issue

TestCuda::test_graph_capture_tunableop_cublas_workspace fails consistently on gfx103X (failed every retry, FAILED CONSISTENTLY). It raises a hipBLASLt runtime error on a plain matrix multiply.

Context: ROCm/TheRock#8782

Agent details (if applicable)

Drafted with an AI assistant (Claude); reviewed by the author.

Problem. With TunableOp enabled, the first matmul raises on GPUs that hipBLASLt ships no kernels for (e.g. gfx103X):

RuntimeError: hipblaslt error: HIPBLAS_STATUS_INVALID_VALUE when calling `hipblaslt_ext::getAllAlgos(...)`

TunableOp lists hipBLASLt candidates with getAllAlgos and throws on any error, before the rocBLAS and default candidates are registered. The non-tunable path already avoids hipBLASLt on these GPUs via Context::blasPreferredBackend(). On gfx1036, 29 of the 30 non-skipped test_linalg -k tunable tests fail this way, as does test_graph_capture_tunableop_cublas_workspace in CI.

Fix:

• INVALID_VALUE from getAllAlgos means no hipBLASLt candidates; other statuses still raise. In hipBLASLt this status only comes from a missing kernel library; an uncovered type returns success with zero algorithms. The device is remembered so it isn't queried again.
• getAllAlgos uses the pooled handle HipblasltGemmOp::Call already uses, instead of a private handle that leaked when the check threw.
• A handle can't be created mid-capture, so during stream capture TunableOp only asks the pool for an existing one: the thread's own handle for that stream, or a spare that CUDAGraph::capture_begin stocked in advance. If neither exists, it offers no hipBLASLt candidates and the GEMM is tuned with rocBLAS. For this, DeviceThreadHandlePool::reserve gets a create flag (default true), exposed as getCurrentCUDABlasLtHandleIfAvailable(). Outside capture nothing changes.
• A warning is raised only if PYTORCH_TUNABLEOP_HIPBLASLT_ENABLED=1 was set explicitly, as in blasPreferredBackend().

Gating on getHipblasltSupportedArchs() was the alternative, both for the fix and for the capture check. It was rejected because that list doesn't match what hipBLASLt ships, and it would drop hipBLASLt tuning on a supported GPU whenever an unsupported one is visible, rather than only when the unsupported GPU is the first to use a GEMM type.

Testing (ROCm 10.1; gfx1036, gfx1100 and gfx1201 in one system):

	gfx1036	gfx1100	gfx1201
test_linalg -k tunable, before	29 failed	29 passed	34 passed
test_linalg -k tunable, after	29 passed	29 passed	34 passed
test_graph_capture_tunableop_cublas_workspace, after	passed	passed	skipped (hipBLASLt wins)

test_call_count_tunableop fails on all three, before and after. Supported GPUs match an unpatched nightly, including the fp8 scaled-GEMM TunableOp tests on gfx1201.

The failure reproduces on any GPU by pointing hipBLASLt at an empty kernel directory, which could serve as a CI regression test:

HIPBLASLT_TENSILE_LIBPATH=/tmp/empty PYTORCH_TUNABLEOP_ENABLED=1 \
  python -c "import torch; a = torch.randn(512, 512, device='cuda'); a @ a"

🤖 Generated with Claude Code

Pull Request resolved: pytorch#200850
Approved by: https://github.com/jeffdaily

Co-authored-by: Claude Opus 5.5 noreply@anthropic.com
(cherry picked from commit be1b713)