rich7420 commented on PR #1389:
URL: https://github.com/apache/mahout/pull/1389#issuecomment-5031186281
Reviewed and tested on real Turing (sm_75, RTX 2080 Ti) hardware. One
blocking regression, a few smaller issues.
**🔴 Blocking: `qdp/qdp-kernels/build.rs:29-31` drops sm_75, breaking ALL
existing GPU kernels (not just the new TC path) on Turing GPUs**
`DEFAULT_CUBIN_ARCHES`/`DEFAULT_PTX_CANDIDATES`/`LEGACY_FALLBACK_ARCHES` no
longer include `"75"`. Since all kernel files share one `cc::Build`, this
silently drops sm_75 SASS/PTX for unrelated pre-existing kernels too (`iqp.cu`,
`amplitude.cu`, `phase.cu`, ...), reverting Turing support that was explicitly
added in #731.
Verified on this branch vs `main` on an actual sm_75 card:
- `main`: `cargo test -p qdp-core --test gpu_iqp_encoding` → 22/22 pass
- this PR: same test → 12/22 fail with `CUDA error code: 500`
(`cudaErrorInvalidDeviceFunction`)
- Even the PR's own claimed verification command, `cargo test --workspace
--exclude qdp-python --lib`, fails on this hardware
(`test_stream_encode_end_to_end`), contradicting the "0 failures" claim in the
description.
Root cause: `ImplicitHadamardOzaki.cu` uses inline PTX
`mma.sync.aligned.m16n8k32...s8` which requires sm_80+ (confirmed via `nvcc
-arch=sm_75`: `Feature '.m16n8k32' requires .target sm_80 or higher`). That's a
real constraint for the *new* Ozaki files, but the fix should scope sm_80+ to
those files specifically, not drop sm_75 for the whole crate.
**🟡 No Rust-level test coverage for
`encode_batch_tc`/`launch_iqp_encode_tc`**
Only `testing/qdp/test_iqp_tc_path.py` exists, which `pytest.skip`s whenever
`encode_batch_tc` isn't present — silently masks build/arch failures like the
one above if CI doesn't build with CUDA.
**🟡 `qdp/qdp-kernels/src/AdaptiveOzaki.cu` / `ImplicitHadamardOzaki.cu`:
unchecked `cudaMalloc` return values**
First place in this crate with raw `cudaMalloc` calls and no error checking
— allocation failure would silently proceed with a null/garbage pointer instead
of a clean error.
**🟡 `launch_adaptive_ozaki_gemm` (`qdp/qdp-kernels/src/lib.rs:428`) ignores
the `stream` param it accepts**
`AdaptiveOzakiEngine::execute()` uses its own internal streams, so the
caller-provided stream's ordering guarantee is broken. Low practical impact
today since nothing calls this from Rust yet, but the FFI contract is
misleading as-is.
**🟡 N>12 Kronecker TC path has known accuracy issues but no guard/warning on
the public API**
`testing/qdp/test_iqp_tc_path.py` tolerates up to 0.1 abs error for N=14
with a comment admitting the scaffold "may diverge," but `encode_batch_tc`
accepts any N with no warning to callers.
Nit: PR title has a double space and inconsistent casing ("Pr6") — `[QDP]
[feature] Pr6 tensor core acceleration`.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
