[lapack][rocsolver] Use the 64-bit pivot API for getrf and getrs - #762
Open
zjin-lcf wants to merge 1 commit into
Open
[lapack][rocsolver] Use the 64-bit pivot API for getrf and getrs#762zjin-lcf wants to merge 1 commit into
zjin-lcf wants to merge 1 commit into
Conversation
oneMath types ipiv as int64_t while the legacy rocSOLVER entry points take rocblas_int, so getrf and getrs each allocated a temporary 32-bit array and ran a cast kernel around the call. In the USM getrs that temporary also forced a queue.wait() before it could be released. rocSOLVER has accepted int64_t pivots since 3.26 (ROCm 6.2), so both routines now pass the user's ipiv straight through, which removes the four temporary allocations, their cast kernels and the blocking wait. The 64-bit getrf reports info as int64_t rather than int, so the devinfo readback is templated on the info type. Earlier rocSOLVER versions keep the conversion path, leaving the minimum supported ROCm unchanged. Because the 64-bit entry points take the dimensions as int64_t, getrf and getrs no longer need the overflow_check that rejected sizes above the 32-bit limit. Note that actually factorizing a matrix with more than 2^28 rows additionally requires the rocBLAS fix from ROCm/rocm-libraries#10831: its 64-bit scal launcher advances the alpha scalar pointer once per n-chunk, so getrf_64 either faults or returns a silently wrong factorization for longer columns. With that fix applied, an m = 2^31 + 1024 factorization returns the correct pivot index 2147484161, which is the case the legacy 32-bit array could not represent. Co-authored-by: Cursor <cursoragent@cursor.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
The rocSOLVER counterpart of #759, which does the same for cuSOLVER (#230).
oneMath types
ipivasint64_t, while the legacy rocSOLVER entry points takerocblas_int.getrfandgetrstherefore each allocated a temporary 32-bit array and ran a cast kernel around the call, and in the USMgetrsthat temporary also forced aqueue.wait()before it could be released.rocSOLVER has accepted
int64_tpivots since 3.26 (ROCm 6.2), so both routines now pass the user'sipivstraight through. That removes four temporary allocations, their cast kernels and the blocking wait. The 64-bitgetrfreportsinfoasint64_trather thanint, soget_rocsolver_devinfoandlapack_info_checkare templated on the info type.ONEMATH_ROCSOLVER_HAS_64BIT_PIVOTSkeeps the previous conversion path for rocSOLVER older than 3.26, so the minimum supported ROCm is unchanged.Because the 64-bit entry points take the dimensions as
int64_t,getrfandgetrsno longer need theoverflow_checkthat rejected sizes above the 32-bit limit.Scope, and how this differs from #759
rocSOLVER has no
sytrf_64, andgetriand thegetrf_batch,getrs_batchandgetri_batchvariants areunimplementedin this backend, sogetrfandgetrsare the only routines affected. There is nothing here matching #759's scratchpad-tail change, andsytrfkeeps its existing conversion path.As in #759, this reduces synchronization without making the USM paths fully asynchronous:
getrfstill blocks inlapack_info_check, which is pre-existing behaviour shared by the rest of the backend and left for separate work.Upstream rocBLAS dependency above 2^28 rows
The API now accepts dimensions past 2^31, but actually factorizing a matrix with more than 2^28 rows also needs the rocBLAS fix from ROCm/rocm-libraries#10831, which I filed while testing this.
rocblas_internal_scal_launcher_64advances the alpha scalar pointer once per n-chunk, sogetrf_64either faults withhipErrorIllegalAddressor returns a silently wrong factorization for any column longer than 2^28.This is not a regression introduced here: it affects every rocSOLVER release that has the 64-bit API (verified on 6.2.4, 6.3.4, 6.4.4, 7.0.1, 7.1.1, 7.2.0 and 7.2.4), the fix is one line in rocBLAS, and nothing at or below 2^28 rows is affected. Reviewers should just know that the "dimensions beyond 2^31" part of this change is not usable until that fix lands. I am happy to gate the 64-bit path on a fixed rocBLAS version, or to cap rows at 2^28 with the legacy path above it, if you would prefer this not to advertise a size range that upstream cannot yet deliver.
Test plan
Built against ROCm 7.2.4 with DPC++ and tested on an AMD Instinct MI300A (gfx942).
developon the same machine gives 180 passed, 0 failed, 364 skipped. The skips are the batch variants this backend does not implement.GetrfandGetrsaccuracy and dependency suites for all four types in buffer and USM form. To be precise about why they were not passing before: this DPC++ build trips a pre-existing assertion (It != m_DeviceKernelInfoMap.end()) on kernels compiled into a backend shared library, which the legacy path's cast kernels hit and the 64-bit path avoids by launching no kernels at all. On a toolchain without that problem both paths pass, so this is not a correctness change — the routines are simply now exercised locally.Sytrfsuites still abort from that same pre-existing assertion, unchanged from baseline;sytrfis untouched by this PR.int64_tpivots verified end to end atm = 2^31 + 1024(8 GiB,n = 1, largest entry at rowm - 512):ipiv[0] = 2147484161, which is 514 pastINT32_MAXand would have truncated to-2147483135in the old 32-bit array.a[0]and the scaled column are correct. This run needs [rocBLAS] scal_64: alpha pointer advanced by the n-chunk offset corrupts n > 2^28 with stride_alpha != 0 (breaks rocsolver_?getrf_64) ROCm/rocm-libraries#10831; I validated it with that one-line fix applied at runtime through anLD_PRELOADshim, cross-checked by confirming the shim also makes a standalone rocBLAS reproducer of the bug pass.clang-formatclean against the repository's_clang-format.