Pr3 shared memory fwt rework - #1396
Conversation
This commit removes the naive dense matrix fallback from the C++ kernels to exclusively rely on standard SIMT implicit FWT for all sizes. It also includes comprehensive correctness tests comparing against PyTorch eager baselines.
|
This pull request has been automatically marked as stale because it has not had any activity for 30 days. It will be closed in another 7 days if no further activity occurs. Thank you for your contribution. If you'd like to keep this open, leave any comment and the stale label will be removed. You can always ask for help on the Mahout dev mailing list or in GitHub Discussions. |
|
This pull request has been automatically marked as stale because it has not had any activity for 30 days. It will be closed in another 7 days if no further activity occurs. Thank you for your contribution. If you'd like to keep this open, leave any comment and the stale label will be removed. You can always ask for help on the Mahout dev mailing list or in GitHub Discussions. |
2 participants
Related Issues
related #1385
Changes
Why
For smaller qubit counts ($N \le 12$ ), calculating the Fast Walsh-Hadamard Transform (FWT) using multiple global memory kernel launches becomes heavily DRAM bandwidth-bound (launch overhead and global memory roundtrips dominate the execution time). By keeping the entire state vector within the GPU's Shared Memory (which is much faster and has lower latency than DRAM), we can fuse the operations into a single kernel launch.
How
iqp_phase_fwt_normalize_tc_kernel): Created a new fused kernel that handles three steps entirely within Shared Memory:extern __shared__ cuDoubleComplex shared_state[].launch_iqp_encode_tcto dynamically allocate Shared Memory and dispatch to this fused kernel whennum_qubits <= FWT_SHARED_MEM_THRESHOLD.Checklist