[blas][rocblas] Support int8 inputs with float output in gemm_batch - #763
Open
zjin-lcf wants to merge 3 commits into
Open
[blas][rocblas] Support int8 inputs with float output in gemm_batch#763zjin-lcf wants to merge 3 commits into
zjin-lcf wants to merge 3 commits into
Conversation
cuBLAS reaches this combination through cublasGemmStridedBatchedEx and cublasGemmBatchedEx, which already accept the datatypes the existing launchers forward, so the column-major buffer, USM strided and USM group entry points only needed routing to the implementation instead of throwing unimplemented. The int32 output combination stays unimplemented because cuBLAS produces it only under CUBLAS_COMPUTE_32I, which takes int32 alpha and beta, whereas oneMath specifies float scalars. A float output accumulated from int8 inputs is rounded at the magnitude of the terms summed rather than at the magnitude of the output entry, so an entry whose sum cancels cannot meet any relative bound. The shared checker takes an optional absolute tolerance, defaulted to zero so that existing callers are unaffected, and the int8-to-float gemm_batch tests pass eps times k * 128 * 128, an upper bound on the accumulated magnitude sum|a*b|. Int8Int8SinglePrecisionErrorModel covers that path with fixed data whose leading rows and columns cancel exactly. Co-authored-by: Cursor <cursoragent@cursor.com>
…del bound An entry whose terms and stored C value are all zero gives a zero model bound, which the reported usage ratio would divide by. Co-authored-by: Cursor <cursoragent@cursor.com>
rocBLAS reaches int8 inputs only with an int32 output and compute type, which takes int32 alpha and beta, whereas oneMath specifies a float output and float scalars for this combination. Accumulate the products exactly in an int32 workspace and apply the float scalars in a scaling kernel afterwards, for the buffer strided, USM strided and USM grouped entry points in both layouts. Bound k so the int32 accumulator cannot wrap, reject sizes whose workspace or kernel range would overflow, and reject an ldc below the row count, which the scaling kernel would otherwise fold onto the next column. Grouped scaling stays a single kernel by locating the group that owns each entry with a binary search over per-group metadata. The scaling kernels launch over a flat range, since HIP cannot map every large multi-dimensional range onto its grid, and the added regression holds a large prime in n to cover that. 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.
Dependency
This is stacked on #761 and contains its commits. Only the last commit,
[blas][rocblas] Support int8 inputs with float output in gemm_batch, is new here; thetest-side tolerance work and the cancellation regression belong to that PR. Please review
#761 first, after which I will rebase this onto
developso the diff reduces to the rocBLASbackend and one added regression.
Summary
rocBLAS reaches int8 inputs only with an int32 output and compute type, which takes int32
alpha and beta, whereas oneMath specifies a float output and float scalars for this
combination. This accumulates the products exactly in an int32 workspace and applies the
float scalars in a scaling kernel afterwards, covering the buffer strided, USM strided and
USM grouped entry points in both layouts.
kis bounded so the int32 accumulator cannot wrap. Every int8 magnitude is at most 128,so
128 * 128bounds a product andkaboveINT32_MAX / (128 * 128)reportsunimplemented.allocated, using checked multiplication and addition rather than a check on the product.
ldcbelow the row count is rejected. rocBLAS rejects such a call itself, but onlyfrom its host task, which does not hold back the separately submitted scaling kernel, and
the kernel would fold one column of C onto the next.
locating the group that owns each entry with a binary search over per-group metadata.
multi-dimensional range onto its grid. The added
Int8Int8SinglePrecisionLargePrimeRangeregression holds a large prime in
nto cover that.Verification
Built with a DPC++ build against ROCm 7.2.4 and run on an AMD Instinct MI300A (gfx942).
The GEMM batch suite passes in both dispatch modes and both layouts:
GemmBatchUsmTests.Complex*is excluded because it is flaky on this machine independently ofthis change: it hangs under compile-time dispatch and dies with
hipErrorIllegalAddressunderruntime dispatch on an unmodified
developcheckout as well.The int8 cases are excluded because they cannot execute in this environment. Any SYCL kernel
launched from a backend shared library trips an assertion in the DPC++ build I have:
This is not specific to this change. The untouched
Gercrow-major test fails identically,since
rocblas_level2.cppalso launches aparallel_forfrom the backend library. I wouldappreciate a run of
*Int8*in the batch suite on a machine with a released toolchain.To cover the paths the unit tests cannot reach here, I exercised the backend directly from
standalone programs linked against
libonemath_blas_rocblas.so, which do run on this device:These check the strided and grouped fallback against an exact integer reference for every
combination of zero and nonzero alpha and beta in both layouts, that grouped scaling submits
one kernel, and that the rejections above report
invalid_argument.Made with Cursor