[QDP] [feature] Pr6 tensor core acceleration#1389
Open
aloha1357 wants to merge 13 commits into
Open
Conversation
aloha1357
requested review from
400Ping,
guan404ming and
ryankert01
as code owners
…tests - Expose encode_batch_tc through Rust core, PyO3, and Python backend - Fix IQP TC kernel: correct batch stride for ZZ params and raise FWT_SHARED_MEM_THRESHOLD to 12 for fused shared-memory path at N<=12 - Align ImplicitHadamardOzaki.cu with PR6 ldmatrix/alignment fixes - Add benchmark_pr5.py with --path fwt|tc|both (GPU-vs-GPU, no PyTorch) - Add test_iqp_tc_path.py smoke and normalization tests
…C GEMM on non-Hadamard logic (PR6). Tests: wsl cargo test passed (0 failures). PR6 comments added.
…r fusion and Kronecker decomposition This commit introduces the foundational architecture for Tensor Core FWT, including: - Fused shared-memory kernel for N <= 12 (1.68x speedup). - Matrix-Free Kronecker decomposition for N > 12 using Adaptive Ozaki GEMM. - Fixed critical ldmatrix alignment bugs in CUDA kernels. - Exposed encode_batch_tc to Python API. - Comprehensive benchmark script and performance report.
aloha1357
force-pushed
the
pr6-tensor-core-acceleration
branch
from
8aca5e6 to
543586a
Compare
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
1 participant
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.
Related Issues
related #1385
Changes
Why
PR1–PR5 delivered the specialized
ImplicitHadamardOzakiEngine(matrix-free, +/-1 perfect quantization, Kronecker-blocked FWT for IQP). However, the full AdaptiveGEMM research engine (AdaptiveOzakiEngine) providing mixed-precision graded-ring (Ozaki + CRT over 7 primes, hybrid FP64/INT8 TC, Phase26 persistent kernels, general A @ B for arbitrary matrices) was only present in the final research snapshot and standalone pybind (adaptive_gemm_py).To finalize the pipeline, we must hook the general engine for "non-Hadamard logic" — i.e., any case where the second operand is not the special structured Hadamard matrix. This enables future general Tensor Core accelerated linear algebra inside QDP (beyond pure IQP FWT) while reusing the same Ozaki INT8 TC machinery.
How
AdaptiveOzaki.cu(full hybrid/persistent/general GEMM implementation) frompr-final-versioninto the clean PR chain.qdp/qdp-kernels/build.rsto compileAdaptiveOzaki.cualongside the Implicit path.launch_adaptive_ozaki_gemmC FFI entry point (wrapsAdaptiveOzakiEngine::executewith default Phase26Hybrid config) + matching declaration and no-cuda stub inlib.rs.// PR6:inline English comments on all changed sites.wsl -e bash -ic 'export PATH=/usr/local/cuda/bin:$PATH && cd .../qdp && cargo test --workspace --exclude qdp-python --lib'passes with 0 failures (builds the new CUDA symbols successfully).The new public kernel symbol
launch_adaptive_ozaki_gemmcan now be called from Rust (qdp-core) or exposed upward, providing the general non-Hadamard TC path that complementslaunch_iqp_encode_tc(Hadamard-specialized).Benchmark Results
Environment: Dev Machine (NVIDIA GeForce RTX 4060 Laptop GPU)
Configuration: Batch size 1024, 50 iterations, GPU-vs-GPU only (no PyTorch reference)
Script:
qdp/qdp-python/benchmark/benchmark_pr6.pyMeasured: 2026-06-10
Branch:
pr6-tensor-core-acceleration(8aca5e69b)Encoding path options
Micro-benchmark (
benchmark_pr6.py): same--pathchoices as PR5 (fwt|tc|both).End-to-end benchmark (
benchmark_e2e.py):--encoding-methodiqp,iqp-z, …--encode-pathfwt,tc,bothencode_batch_tcinside the disk→GPU pipelineE2E intentionally omits PyTorch from the hot path when comparing Mahout encode paths — only Mahout frameworks are timed for encoding throughput.
FWT vs Tensor Core path (after PR6)
Checklist