[zjin-lcf]

Summary

• implement gerqf and getri with native rocSOLVER entry points for buffer and USM APIs
• implement strided and grouped batched geqrf, getrf, getrs, getri, potrf, potrs, and orgqr/ungqr routines with scratchpad queries
• preserve SYCL dependencies, convert oneMath 64-bit pivots safely, report per-matrix batch errors, and clean up temporary USM allocations on failures
• leave ormrq/unmrq, trtrs, and hetrf unimplemented because rocSOLVER has no native equivalents

Test plan

• Build the rocSOLVER LAPACK backend with DPC++ and ROCm 7.2.4
• Run compile-time dispatch LAPACK tests on AMD GPU
• Run runtime dispatch LAPACK tests on AMD GPU
• Verify grouped potrs_batch accuracy and dependency tests with multiple right-hand sides
• Run clang-format 19.1.0 and git diff --check

Portability

The implementation uses architecture-neutral rocSOLVER APIs and basic SYCL kernels.

[melonakos]

I can't review this as a single PR, and I don't think anyone else can either. That's the main message. But there's real quality in here and a couple of specific findings, so let me be concrete about both.

This is at least twelve routines in one PR

Counting the unimplemented throws you remove, this implements:

getrf_batch, getrs_batch, getri_batch, geqrf_batch, potrf_batch, potrs_batch, orgqr_batch, ungqr_batch, getri, gerqf — plus roughly ten corresponding *_scratchpad_size queries.

At +1631/−933 with rocsolver_batch.cpp alone at +1383/−847, that's twelve independent numerical-library features sharing nothing but a filename. Every one of them is a place where a transposed argument, a wrong leading dimension, a mis-mapped stride or a misread info convention hides — and none of it can be checked by CI, because there's no AMD runner.

I spot-checked for the systemic problems I'd found in your other PRs rather than attempting twelve routine-by-routine reviews, because a review that skims twelve features is worse than no review: it produces an approval that means nothing. This PR needs two approvals, and asking two reviewers each to spend a day on it is why it will sit here indefinitely.

Split it — ideally one PR per routine family, e.g. the LU batch group (getrf/getrs/getri), the QR batch group (geqrf/orgqr/ungqr), the Cholesky batch group (potrf/potrs), and the two non-batched routines. Four PRs of ~400 lines each are individually reviewable in an hour, and you'd start getting them merged instead of holding all twelve hostage to the hardest one.

What I checked, and what's good

Your exception handling here is right, and I want to highlight it because it contrasts with the older PRs:

try {
    lapack_info_check_batch(queue, devInfo, __func__, func_name, batch_size);
}
catch (...) {
    sycl::free(ipiv32, queue);
    sycl::free(devInfo, queue);
    throw;
}
sycl::free(ipiv32, queue);
sycl::free(devInfo, queue);

Both allocations released on both paths. That's exactly the pattern #748, #759 and #762 are missing — they call lapack_info_check and then free, so they leak whenever a matrix is singular. Your newer PRs (#766, #767, #768) all get this right and the older ones don't, which suggests you worked it out along the way. Worth going back and applying this to the three older ones; better still, a small RAII holder so it's not written out by hand at every site.

lapack_info_check_batch is a genuinely necessary addition. The existing get_rocsolver_devinfo only ever returns dev_info_[0], so before this there was no way to check per-matrix info at all on the rocSOLVER side — a batched call would have reported success while matrices 1..N−1 failed silently. Adding a batch-aware check rather than reusing the single-value one is the right call.

Finding: the frees depend on an unrelated global wait

auto done_casting = copy_ipiv_to_64(queue, ipiv32, ipiv, ipiv_len, stride_ipiv, batch_size, { done });

try {
    lapack_info_check_batch(queue, devInfo, __func__, func_name, batch_size);
}
...
sycl::free(ipiv32, queue);

done_casting reads ipiv32, and you free ipiv32 without waiting on it. This is safe today only because lapack_info_check_batch performs a queue.wait() internally — a global sync that happens to drain done_casting too.

Two problems with relying on that:

1. It's a fragile coupling. Nothing at this call site says the free is ordered behind the copy, and the guarantee lives inside a function whose job is error checking, not synchronisation. If anyone later narrows that queue.wait() to a targeted wait — which would be a reasonable optimisation — this silently becomes a use-after-free on the pivot array.
2. It makes the USM entry point synchronous. You return done_casting, but by then it has already completed, so callers get an event that can't overlap with anything. Same async-contract issue I raised on #748.

Cleanest fix is to stop relying on the side effect: either done_casting.wait() explicitly before freeing, which at least makes the dependency visible, or attach the frees to a host_task that depends on done_casting so the memory is released when the work is actually done and the function stays asynchronous.

Interaction with #762 and #759 worth deciding

This PR implements the batched rocSOLVER routines using a 32-bit pivot temporary plus a copy_ipiv_to_64 widening pass. Meanwhile #762 migrates the non-batched rocSOLVER getrf/getrs to native 64-bit pivots, and #759 migrates the cuSOLVER batched routines to native 64-bit pivots.

Land all three as they stand and you get: cuSOLVER batched on native 64-bit pivots, rocSOLVER batched on 32-bit round trips, rocSOLVER non-batched on native 64-bit pivots. Three different strategies across two backends for the same operation.

Given rocSOLVER's *_64 strided-batched entry points exist from 3.26 alongside the non-batched ones you're already gating on in #762, it's worth checking whether the batched routines here can use them directly and skip the pivot round trip entirely. That might make this PR meaningfully smaller as well as more consistent — which is a good reason to settle it before splitting rather than after.

I'm happy to review the split-out pieces promptly, and I'd suggest starting with whichever group you most want merged.

[zjin-lcf]

Split into focused PRs for independent review:

• [lapack][rocSOLVER] Report info for every matrix in a batch #778 — shared rocSOLVER version detection and batched info reporting
• [lapack][rocSOLVER] Implement Cholesky batch routines #779 — Cholesky batch routines (potrf_batch, potrs_batch)
• [lapack][rocSOLVER] Implement LU batch routines #780 — LU batch routines (getrf_batch, getrs_batch, getri_batch)
• [lapack][rocSOLVER] Implement QR batch routines #781 — QR batch routines (geqrf_batch, orgqr_batch, ungqr_batch)
• [lapack][rocSOLVER] Implement gerqf and getri #782 — non-batched gerqf and getri

Closing this combined PR in favor of those feature-focused changes.