Repository navigation
Conversation
Implements accurate FP64 GEMM emulation by decomposing double-precision
inputs into up to 14 INT8 moduli and accumulating partial products via
the Chinese Remainder Theorem (Ozaki Scheme II).
New files:
fp64_emulation.cpp / fp64_emulation.hpp - core emulation logic
Integration:
CMakeLists.txt updated to build the new source
rocblaslt_mat.cpp gates each FP64 GEMM through fp64EmulationIsEnabled()
and an arithmetic-intensity check before dispatching to fp64EmulatedGemm().
Tuning via environment variables:
HIPBLASLT_EMULATE_DOUBLE_PRECISION=1 - enable emulation
HIPBLASLT_FIXEDPOINT_EMULATION_MANTISSA_BIT_COUNT - requested precision in
mantissa bits, which determines the minimum number of INT8 moduli needed
HIPBLASLT_EMULATION_SPECIAL_VALUES_SUPPORT_MASK - Inf/NaN detection mask
HIPBLASLT_EMULATION_STRATEGY=eager - bypass AI threshold check
Callers can supply a pre-allocated workspace via the handle-level API
added to hipblaslt.h to avoid per-call device allocations.
- Extend OZ2_S_MAX from 14 to 20, adding all GEMMul8 tables for s=15..20 (moduli 197..173, ~155 bits CRT capacity) - Fix oz2_accum_kernel: reduce C32i mod m_t before multiplying by qPi so the product is exact in double (matches GEMMul8's conv_32i_2_8i step) - Add third float correction pass (ITER=3) in oz2_scaleA/B_kernel for s≥19, matching GEMMul8's threshold
Allow sweep to skip error check. Disable inf/nan check in sweep
…ces, and one for the GEMM and acuumulation.
…input. This should improve performance of the scale and prelim extract kernels. Originally, the int8 GEMM was always in TN since other transpose type were not supported by cuBLAS.
…he FP64 input." This reverts commit 118dbd1.
This was referenced Oct 14, 2026
geomin12
added a commit
that referenced
this pull request
Oct 14, 2026
The bot cleared baseline-run-id in #13623 alongside the TheRock ref bump to 2a55cec, and this branch inherited the empty value when develop was merged. An empty baseline makes therock-multi-arch-ci.yml pass prebuilt_stages="" and stage_reuse_mode=off, so every stage rebuilds and the run looks like changed_projects failed to narrow anything. It did not: setup resolved all 34 rocm-libraries projects and put rccl/rocshmem/core-hiptests in reusable_artifacts, exactly as #8514 expects. Restore 38220501023, the baseline develop carried before the bump and the one the stacked rocPRIM PR still uses, so the two test PRs differ only in their diffs. Co-Authored-By: Claude Opus 4.6 <noreply@anthropic.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.
Aligns hipBLASLt FP64 emulation env/API behavior with the cuBLAS-style control surface while keeping runtime ADP deferred.
Adds strict parsing for HIPBLASLT_ FP64 emulation env vars, including explicit
EMULATE_DOUBLE_PRECISION=0,EMULATION_STRATEGY=performant|eager, special-values mask validation, and fixed mantissa bit-count validation. Replaces bool-only emulation gating with a status-returning decision path so invalid env values return invalid_value instead of silently falling back to native FP64. Updates precedence rules: env overrides API for strategy and special-values; enable remains API-over-env; mantissa bit-count env forces FIXED only from the default handle state. Keeps DYNAMIC on the existing 16-moduli non-ADP path with a one-time warning, and defers ADP/mantissa offset work. Hardens API validation, updates docs, and adds parser/decision/API validation tests.