|
1 | 1 | # DTWC++ Development TODO |
2 | 2 |
|
3 | | -**Last Updated:** 2026-04-02 (GPU optimization session) |
| 3 | +**Last Updated:** 2026-04-02 |
4 | 4 |
|
5 | | -## Completed (This Session) |
| 5 | +## Completed |
6 | 6 |
|
| 7 | +### Waves 1A/1B/2A/2B (2026-04-02) |
| 8 | +- [x] Wave 1A: Metrics + Missing Data (missing_utils, MissingStrategy, DTW-AROW, 5 scoring metrics) |
| 9 | +- [x] Wave 1B: Multivariate Foundation (ndim, MVL1/MVSquaredL2, MV DTW, MV DDTW) |
| 10 | +- [x] Wave 2A: Clustering Algorithms (deferred alloc, medoid utils, hierarchical, CLARANS, FastCLARA fixes) |
| 11 | +- [x] Wave 2B: MV Variants + Lower Bounds (MV WDTW/ADTW/DDTW, per-channel LB_Keogh, MV missing DTW) |
| 12 | + |
| 13 | +### GPU Optimization (2026-04-02) |
7 | 14 | - [x] CUDA anti-diagonal wavefront kernel (85x kernel speedup) |
8 | 15 | - [x] FP32/FP64 auto-precision by GPU architecture |
9 | | -- [x] Register-tiled kernel (L<=256, cuDTW++-inspired) |
10 | | -- [x] Warp-level kernel (L<=32, 8 pairs/block) |
| 16 | +- [x] Register-tiled kernel (L<=256), warp-level kernel (L<=32) |
11 | 17 | - [x] Persistent kernel with atomic work queue |
12 | | -- [x] Double-buffer for long series (L>1024) |
13 | | -- [x] Series preloading into shared memory (L<=256) |
14 | | -- [x] On-device pair index computation (eliminate pair arrays) |
15 | | -- [x] GPU-side NxN matrix write (eliminate host fill loop) |
16 | | -- [x] Integer band boundary precomputation |
17 | | -- [x] CUDA streams + pinned memory + event-based timing |
18 | | -- [x] GPU LB_Keogh lower bound kernels |
19 | | -- [x] GPU pruned distance matrix with skip_threshold |
20 | | -- [x] 1-vs-N and K-vs-N GPU kernels for clustering |
| 18 | +- [x] GPU LB_Keogh, 1-vs-N and K-vs-N kernels |
21 | 19 | - [x] Parallel CPU pruned distance matrix (OpenMP + atomic NN) |
22 | 20 | - [x] DistanceMatrixStrategy enum (Auto/BruteForce/Pruned/GPU) |
23 | | -- [x] Python `device='cpu'|'cuda'` API |
24 | | -- [x] CLI `--device cuda --precision auto` |
25 | | -- [x] Cross-platform CUDA/MPI/OpenMP detection |
26 | | -- [x] CI workflow for CUDA/MPI smoke testing |
27 | | -- [x] 3 adversarial reviews (GPU kernel, Python API, build system) |
28 | | -- [x] 41/41 tests, 312/312 MPI tests |
| 21 | +- [x] Python `device='cpu'|'cuda'` API, CLI `--device cuda` |
| 22 | +- [x] Cross-platform CUDA/MPI/OpenMP detection + CI workflow |
29 | 23 |
|
30 | 24 | ## Remaining Work |
31 | 25 |
|
32 | | -### High Priority |
33 | | -- [ ] Wire K-vs-N kernel into fast_pam clustering loop (C++ side) |
34 | | -- [ ] Wire GPU LB_Keogh into clustering (prune before DTW in iterations) |
35 | | -- [ ] Fix register-tiled kernel for banded DTW edge cases (some precision issues at boundaries) |
36 | | -- [ ] CUDA streams: multi-stream pipelining for very large N (currently single stream) |
37 | | -- [ ] Profile register pressure with `--ptxas-options=-v`, add `__launch_bounds__` |
| 26 | +### MIP Solver Improvements (from UNIMODULAR.md analysis) |
| 27 | +- [ ] MIP start from PAM (warm start for both HiGHS and Gurobi, ~10 lines each) |
| 28 | +- [ ] Gurobi: reduce NumericFocus 3->1, add MIPFocus=2, branching priority on A[i,i] diagonals |
| 29 | +- [ ] Benders decomposition for N > 200 (master: N binary vars, subproblem: O(Nk) assignment) |
| 30 | +- [ ] Odd-cycle cutting planes ({0,1/2}-CG cuts) as lazy constraints |
| 31 | + |
| 32 | +### CUDA Next Phase (see .claude/superpowers/plans/2026-04-02-cuda-next-phase.md) |
| 33 | +- [ ] Device-side pruning: stop launching DTW for LB-pruned pairs |
| 34 | +- [ ] Architecture-aware dispatch (DispatchProfile by compute capability) |
| 35 | +- [ ] Wire K-vs-N kernel into fast_pam clustering loop |
| 36 | +- [ ] Wire GPU LB_Keogh into clustering iterations |
| 37 | +- [ ] Benchmark expansion: standalone LB, pruned matrix, 1-vs-N, K-vs-N |
| 38 | + |
| 39 | +### CUDA Medium Priority |
| 40 | +- [ ] Fix register-tiled kernel for banded DTW edge cases |
| 41 | +- [ ] Multi-stream pipelining for very large N |
| 42 | +- [ ] Profile register pressure, add `__launch_bounds__` |
| 43 | +- [ ] Hilbert-curve pair ordering for L2 cache locality |
| 44 | +- [ ] GPU early-abandon within DTW kernels |
| 45 | +- [ ] Template kernels on `use_squared_l2` |
| 46 | + |
| 47 | +### Algorithms & Scale |
| 48 | +- [ ] Condensed distance matrix (half memory for symmetric storage) |
| 49 | +- [ ] Two-phase clustering for pre-categorized data (within-group + cross-group) |
| 50 | +- [ ] Lazy loading (FileBackedDataSource, CachedDataSource) |
| 51 | +- [ ] Binary distance matrix storage (HDF5 + mmap'd flat binary) |
| 52 | +- [ ] Algorithm auto-selection based on cost = N^2 * min(L, band) * ndim |
38 | 53 |
|
39 | | -### Medium Priority |
40 | | -- [ ] Hilbert-curve pair ordering for L2 cache locality (helps when N*L > L2) |
41 | | -- [ ] GPU early-abandon within DTW kernels (periodic threshold check) |
42 | | -- [ ] Template kernels on `use_squared_l2` (compile-time metric dispatch) |
43 | | -- [ ] HIPify for AMD GPU support (~1-2 days for 300 LOC) |
44 | | -- [ ] MATLAB MEX bindings with GPU support |
| 54 | +### Bindings |
| 55 | +- [ ] MATLAB MEX bindings (skill drafted in .claude/skills/matlab-wrapper-skill.md) |
| 56 | +- [ ] Python binding updates (skill drafted in .claude/skills/python-wrapper-skill.md) |
| 57 | +- [ ] HIPify for AMD GPU support |
| 58 | + |
| 59 | +### Technical Debt |
| 60 | +- [ ] Clean up wavefront kernel dead code (unreachable preload branch for L<=256) |
| 61 | +- [ ] Unify kernel dispatch logic |
| 62 | +- [ ] Add `DeviceContext` abstraction (device_id, stream, workspace pool) |
45 | 63 |
|
46 | 64 | ### Low Priority / Research |
47 | 65 | - [ ] Multi-GPU support (data sharding across devices) |
48 | | -- [ ] TMA (Tensor Memory Accelerator) for Hopper GPUs |
49 | | -- [ ] `cp.async` for compute-transfer overlap on Ampere+ |
| 66 | +- [ ] TMA for Hopper GPUs |
50 | 67 | - [ ] GPU DTW variants (DDTW, WDTW, ADTW on GPU) |
51 | 68 | - [ ] Soft-DTW GPU kernel |
52 | | - |
53 | | -### Technical Debt |
54 | | -- [ ] Clean up wavefront kernel dead code (preload branch for L<=256 is now unreachable) |
55 | | -- [ ] Unify kernel dispatch logic (currently scattered across launch_dtw_kernel) |
56 | | -- [ ] Add `DeviceContext` abstraction (carries device_id, stream, workspace pool) |
| 69 | +- [ ] SIMD via Google Highway (LB_Keogh, z_normalize, multi-pair DTW) |
0 commit comments