[QDP] Complete Tensor Core IQP path remediation

[aloha1357]

Tensor Core IQP path remediation

This PR updates the Tensor Core IQP path from the earlier scaffold into a merge-ready implementation addressing the review feedback on #1387.

The main goals are:

• remove unused experimental AdaptiveOzakiEngine code
• fix Tensor Core launcher error propagation
• validate encode_batch_tc directly with deterministic correctness tests
• keep Tensor Core behavior explicit opt-in and separate from the default IQP path

What changed

Dead code removal

The previous branch carried an unused AdaptiveOzakiEngine implementation and an unused launch_adaptive_ozaki_gemm FFI surface. That code was not part of the IQP Tensor Core path and made the PR broader than necessary.

This PR removes that path entirely:

• deleted qdp/qdp-kernels/src/AdaptiveOzaki.cu
• deleted qdp/qdp-kernels/src/AdaptiveOzaki.h
• removed AdaptiveOzaki.cu from build.rs
• removed launch_adaptive_ozaki_gemm FFI declaration and non-CUDA stub from lib.rs
• extracted the minimal shared Ozaki config into ozaki_config.h

CUDA error handling

The N <= 12 Tensor Core path now checks CUDA failures instead of silently returning success:

• captures and checks cudaFuncSetAttribute
• launches the shared-memory TC kernel
• calls cudaStreamSynchronize(stream)
• returns any synchronization error
• otherwise returns the final cudaGetLastError()

Test coverage

The test suite now exercises encode_batch_tc directly instead of only the default encode path.

Test area	Coverage
Deterministic inputs	seeds 42 and 137
Independent oracle	torch_iqp_encode_ref CPU formula reference
Small-N correctness	N = 6, 8, 10, 12, max_err < 1e-9
Large-N correctness	N = 14, 16, max_err < 1e-5
Existing TC tests	preserved and moved to deterministic inputs
Binding smoke tests	unchanged behavior verified

Reviewer feedback addressed

Review point	Resolution
TC API returns invalid IQP state	Direct TC path tests validate against an independent CPU formula oracle
Full-ZZ TC batches read wrong sample data	TC path accounts for the full IQP parameter count
Test did not exercise TC path	Tests now call encode_batch_tc directly
TC launcher hid CUDA failures	cudaFuncSetAttribute is checked; stream is synchronized before final error check
Random-only tests / tolerance unclear	Deterministic seeds and documented 1e-9 / 1e-5 tolerances
Scaffolding should not merge to main	Branch is now a complete TC-path remediation, not scaffold-only

Path separation

This PR does not introduce the later Native/Fused FP32 Hadamard work. That will be proposed separately as a new explicit opt-in path.

Path	Status in this PR
Default IQP encode path	preserved
Tensor Core encode_batch_tc path	remediated and directly tested
Native/Fused FP32-FP64 path	separate future PR

Validation

Validated locally in the repository WSL environment:

Command	Result
cargo build -p qdp-kernels	passed
cargo test -p qdp-kernels -q	passed
cargo clippy -p qdp-kernels -- -D warnings	passed
pytest testing/qdp/test_iqp_tc_path.py -v	25 passed
pytest testing/qdp/test_iqp_tc_path.py -q	passed
pytest testing/qdp/test_bindings.py -q	80 passed, 2 skipped
pytest testing/qdp/ -q	165 passed, 2 skipped
ruff check qdp/	passed
ty check qdp/	passed

Notes on large N

The implementation is intended to support larger Tensor Core IQP workloads, but the automated CPU-oracle correctness tests intentionally stop at N = 16.

For N = 27, a full state has 2^27 = 134,217,728 amplitudes, so a full CPU-oracle max-error test is not appropriate for this PR gate. Larger-N coverage should use smoke/invariant tests and benchmark-specific validation rather than materializing a full CPU reference in CI.

[ryankert01]

There's several issue found:

1. TC API returns an invalid IQP state
2. Full-ZZ TC batches read the wrong sample data
3. The new test does not exercise the TC path
4. TC launcher hides CUDA allocation and launch failures
5. and some test enhancements

[ryankert01]

with different precision can have different tolerances(f32,f64)

[ryankert01]

Let's not use truely random numbers because it introduces flackiness.

[ryankert01]

Add weird n_qubits and batch_size for robustness.

[ryankert01]

A few comments

[ryankert01]

This TC launcher currently writes the pre-FWT phase vector rather than an IQP-encoded state. iqp_phase_split_kernel computes exp(i theta(x)), and this recombine writes it directly to the output; the FWT + normalization is still marked as future work. For all-zero params, this returns all ones instead of |0...0>. Can we either keep this as private scaffolding for now, or complete the transform before exposing it through encode_batch_tc?

[ryankert01]

For full ZZ batches, this hardcodes data_len = num_qubits, but Rust validates sample_size as num_qubits + num_qubits * (num_qubits - 1) / 2. That means sample 1 starts inside sample 0's ZZ parameters when enable_zz=true. Can we pass the actual sample_size/data length into launch_iqp_encode_tc, matching the existing batch launcher?

[ryankert01]

This test currently exercises engine.encode(...), which routes through the existing standard IQP batch path, not the new encode_batch_tc path. So it can pass even if the TC scaffold is incorrect. Can we add a direct CUDA/Rust test for IqpEncoder::encode_batch_tc, or expose an explicit opt-in Python API and test that path?

[ryankert01]

This launcher ignores both cudaMalloc return values and then returns cudaSuccess unconditionally. On OOM or launch failure, Rust will report success while the output may be invalid; if the second allocation fails, the first allocation can also leak. Can we return CUDA errors consistently with the existing IQP launchers and free any partial allocations on failure?

[ryankert01]

Btw, @aloha1357 Thanks for your patch!

[aloha1357]

@ryankert01 Thanks for the detailed review — these are exactly the right things to flag on this PR.

Context: intentional scaffolding

#1387 (PR2) is scoped as TC-path scaffolding only: encode_batch_tc API wiring,
phase split (real/imag), and bank-conflict-free batch transpose. The full Hadamard /
Kronecker / Tensor Core pipeline is delivered incrementally in the stacked PRs below,
so each PR stays reviewable on its own.

Where your points are addressed
Reviewer issues vs. PR stack

#	Reviewer Issue	PR2 #1387	PR4 #1391	PR5 #1390	PR6 #1389 / PR7
Q1	TC API returns invalid IQP state (no FWT / not a proper quantum state)	Scaffolding only — phase split → recombine; no Hadamard	Partially fixed — shared-mem fused FWT (N≤12) + Kronecker naive-GEMM pipeline (N>12); still broken for full-ZZ until PR5	Fixed — full pipeline; Ozaki implicit Hadamard replaces naive GEMM	Fixed — same as PR5; WMMA / TC kernels hardened
Q2	Full-ZZ TC batches read wrong sample data (data_len = num_qubits only)	Not fixed — hardcoded data_len = num_qubits	Not fixed — same bug in launch_iqp_encode_tc	Fixed — data_len accounts for ZZ params (commit 806a419)	Fixed — inherited from PR5
Q3	Tests do not exercise the TC path (engine.encode() instead of encode_batch_tc)	Not fixed — test_batch_throughput.py calls engine.encode() only	Not fixed — no TC-specific test; Rust encode_batch_tc exists but Python API not exposed yet	Fixed — test_iqp_tc_path.py calls encode_batch_tc()	Fixed — extended (e.g. N=14 FWT vs TC smoke / loose agreement)
Q4	TC launcher hides CUDA alloc/launch failures (always return cudaSuccess)	Not fixed — hardcoded cudaSuccess; cudaMalloc unchecked	Not fixed — same	Partially fixed — cudaFuncSetAttribute error check + cudaGetLastError() at exit; cudaMalloc return values still unchecked	Partially fixed — same as PR5; cudaMalloc checks + buffer pool planned for PR6/PR7
Q5	Test quality (random data; tolerance should follow precision)	Not addressed — torch.randn + assert_close on wrong API (encode)	Not addressed — no TC tests	Partially addressed — encode_batch_tc normalization smoke tests; still uses torch.randn; no torch_ref baseline for TC correctness	Partially addressed — adds FWT-vs-TC comparison (loose tolerance); deterministic data + precision-aware torch_ref baseline still planned for PR6/PR7

One-Line Summary per PR

PR	Q1–Q5 Coverage
PR2	Scaffolding; all five issues are visible — expected at this stage
PR4	Q1 partially fixed; Q2–Q5 still open
PR5	Q1–Q3 fixed; Q4–Q5 partial
PR6/PR7	Q1–Q3 done; Q4–Q5 to be completed here

---

Note: PR3 (#1388) is a separate line (shared-mem FWT in iqp.cu only; no iqp_tc.cu). It does not address Q1–Q5 for the TC path discussed in #1387.
Stack for review

• PR4: [QDP] [feature] Pr4 kronecker fwt #1391
• PR5: [QDP] [feature] Pr5 implicit hadamard engine #1390
• PR6: [QDP] [feature] Pr6 tensor core acceleration #1389

Question for you

If the PR2 scaffolding looks acceptable in principle, would you prefer to:

1. Approve PR2 as scaffolding and continue reviewing PR4 → PR5 → PR6 in order, or
2. Hold final feedback until PR6/PR7, where we will land the remaining Q4/Q5 fixes
   (cudaMalloc error propagation, deterministic/precision-aware tests)?

Option 2 avoids repeatedly growing the same PR and keeps the review surface stable.
Happy to adjust based on what works best for your review workflow.

[ryankert01]

Im object to merge any scaffolding pr into main. We can do a dev branch dedicate for this or a big pr.

[aloha1357]

Thanks, that makes sense. I agree that merging a
scaffolding-only PR into main is not the right direction.

I reworked this branch following the “big PR” approach you
suggested, so this is no longer just TC scaffolding. The
branch now contains the complete remediation for the TC path
review issues.

What changed:

• Removed the unused AdaptiveOzakiEngine path entirely:

  • deleted AdaptiveOzaki.cu
  • deleted AdaptiveOzaki.h
  • removed AdaptiveOzaki.cu from build.rs
  • removed the unused launch_adaptive_ozaki_gemm FFI
    declaration and non-CUDA stub

• Extracted only the shared Ozaki config types needed by
  ImplicitHadamardOzakiEngine into ozaki_config.h.

• Fixed the TC launcher error handling:

  • cudaFuncSetAttribute return value is now checked
  • the N <= 12 path now calls
    cudaStreamSynchronize(stream) before the final
    cudaGetLastError()
  • the larger path continues to use checked CUDA calls

• Fixed the Full-ZZ sample data issue:

  • TC data length now accounts for ZZ parameters instead of
    assuming data_len = num_qubits

• Added direct TC-path tests:

  • tests now call encode_batch_tc
  • deterministic seeds: 42 and 137
  • added independent CPU formula oracle
    torch_iqp_encode_ref
  • small-N checks: N = 6, 8, 10, 12, max_err < 1e-9
  • large-N checks: N = 14, 16, max_err < 1e-5
  • existing TC smoke/agreement tests are preserved and now
    use deterministic inputs

  Validation run locally in the repo WSL environment:

cargo build -p qdp-kernels
cargo test -p qdp-kernels -q
cargo clippy -p qdp-kernels -- -D warnings
pytest testing/qdp/test_iqp_tc_path.py -v
pytest testing/qdp/test_iqp_tc_path.py -q
pytest testing/qdp/test_bindings.py -q
pytest testing/qdp/ -q
ruff check qdp/
ty check qdp/

Results:

pytest testing/qdp/test_iqp_tc_path.py -v: 25 passed
pytest testing/qdp/test_bindings.py -q: 80 passed, 2 skipped
pytest testing/qdp/ -q: 165 passed, 2 skipped
ruff check qdp/: passed
ty check qdp/: passed
cargo build/test/clippy: passed

So the branch is now intended to be reviewed as the complete
TC-path PR, not as scaffolding.