2026-08-19 - Portable Claude session sync (pull/push, box-agnostic project-dir naming, gz store, prune, size guard) + separate logs branch setup - CP / .claude/scripts/{lib-sessions,pull,push,logs-pull,logs-push}.sh, sessions.conf, .gitignore, README.md
2026-08-19 - Session sync bootstrapped on H200: repo reset 845MB->48MB .git, logs branch created, 4 bugs fixed (missing .gitattributes aborted squash, --init inherited main's tree as 1.9GB clone, squash unrepeatable when stranded on scratch branch, prune orphaned subagent sidecar dirs) - CP / .claude/scripts/.sh, .gitattributes, sessions.conf
2026-07-18 - Hybrid Ring x Ulysses mem-scaling sweep: new mem_scaling_sweep_hybrid.sh (CP_ATTN_MODE=hybrid, TOPOLOGIES matrix of ring x ulysses factorizations, T-doubling-to-OOM, per-topology label so table keeps rows distinct, fused-tile T guard, RING_OUTER for NUMA); wired CP_FFN (8/3d_model default) + CP_WEIGHT_TYING (were ignored, froze ffn/tying) so scaling configs are faithful; verified snapshots + mem_scaling_table.py parse - CP / Tests/bluscriptcp/mem_scaling_sweep_hybrid.sh, Scripts/Blutrain/bluscriptCP.cpp
2026-07-15 - Hybrid CP server validation + placement: U=2 hybrid forward BIT-EXACT vs fused-kernel ground truth (sdpa_fused_forward_rope anchor); found degree-4 pure-ring diverges (cos 0.91, separate ring issue); added CP_RING_OUTER mesh-order toggle ({ring,dp,ulysses}) so Ring crosses NUMA/nodes while Ulysses+DDP stay intra-domain; fused-tile CP_T guard - CP / Scripts/Blutrain/bluscriptCP.cpp, Tests/cp_hybrid_parity.cpp, README.md
2026-07-15 - 3D hybrid CP {DDP,CP-Ring,CP-Ulysses}: Ulysses-inner all-to-all nested in Ring-outer via enable_hybrid/forward_hybrid + composed HeadTail-over-ring/contiguous-over-ulysses sharder + combine/partition autograd nodes; bluscriptCP CP_ATTN_MODE=hybrid 3D mesh wiring; cp-hybrid parity test (MHA bit-exact vs pure-ring) + hybrid_cp_smoke.sh; README - CP / context_parallel/ContextParallel.h, HybridUlyssesBackward.h, Scripts/Blutrain/bluscriptCP.cpp, Makefile, Tests/cp_hybrid_parity.cpp, Tests/hybrid_cp_smoke.sh, README.md
2026-04-20 - CP-aware TF32 WMMA attention fwd/bwd: AttentionForward/Backward CP copies modified with T_q/T_k/q_offset/k_offset, exp11 bwd dQ/dK/dV atomicAdd for ring accumulation, namespace OwnTensor::cp - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/AttentionForward.cu, AttentionBackward.cu
2026-04-18 - Throughput fixes: remove cudaDeviceSynchronize (Fix1), zeros->empty bwd (Fix2), P2PRingRotator grouped sendrecv (Fix4) - TensorParallelismBeta / FusedSDPAOp.h, RingRotator.h
2026-04-16 - SDPA LB fwd+bwd implemented but grad explosion unresolved; LB disabled for causal, TC backward re-enabled, backward deadlock fix (step_skipped), AlltoAllRingRotator rewritten to sparse alltoallv (PyTorch parity), alltoallv_async added to ProcessGroupNCCL - TensorParallelismBeta / RingRotator.h, ContextParallel.h, ContextParallelBackward.h, FusedSDPABackwardKernel.cu, ProcessGroupNCCL.h, processGroupNCCL.cpp
2026-04-15 - Grad explosion fix: TC backward if (false) guard re-applied in launch_flash_attn_bwd_f32 - TensorParallelismBeta / FusedSDPABackwardKernel.cu
2026-04-10 - CP throughput gap debug (55.5k vs 64k PyTorch) - TC kernel T_q/T_k/K_head_stride, LEFT_HALF/BOT_HALF TC dispatch, clone overhead identified - TensorParallelismBeta / AttentionKernels.cu, AttentionKernels.h, FusedSDPAOp.h, ContextParallel.h, ContextParallelBackward.h
2026-03-20 - Context Parallelism implementation (RingRotator, SDPA, Merger, Forward/Backward) - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/RingRotator.h, SDPAOp.h, SDPAMerger.h, ContextParallel.h, ContextParallelBackward.h
2026-03-20 - Context Parallel test file and Makefile target - TensorParallelismBeta / DTensor/gpt2_cp_test/gpt2_cp_test.cpp, DTensor/Makefile
2026-03-21 - CP backward rotator fix, load_balance causal fix, gpt2_cp_test.cpp full rewrite with data loader/timing/logging/token-gen - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/ContextParallelBackward.h, ContextParallel.h, gpt2_cp_test.cpp
2026-03-21 - gpt2_cp_test.cpp weight tying (NoGradGuard+transpose), global_batch/grad_accum/max_steps/warmup_steps from gpt2_tp_test pattern, grad_scale backward - TensorParallelismBeta / DTensor/gpt2_cp_test/gpt2_cp_test.cpp
2026-03-21 - OOM fix: flatten_concat returns [1,N] not [N]; added .flatten() before narrow() in ContextParallel ring loop to prevent 1.15TB allocation - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/ContextParallel.h
2026-03-28 - Debug: wrong gradients in CP - root cause: saved_k_chunks[i] aliases recv_buffer_ (P2P/AlltoAll rotators), fix: clone before saving in ContextParallel.h:216-217 - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/ContextParallel.h, RingRotator.h, ContextParallelBackward.h, SDPAOp.h
2026-03-28 - Debug: CP vs normal wpe gradient mismatch (10-17x smaller for positions 0..511) - root cause: Phase 3 sendrecv in ContextParallelBackward not delivering dK/dV from rank 1 back to rank 0 - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/ContextParallelBackward.h, gpt2_cp_test.cpp
2026-03-28 - Fix: Phase 3 backward now dispatches on rotator_type_: AlltoAll uses alltoall_async (build send_buf per dest slot, sum recv slots), P2P keeps sendrecv - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/ContextParallelBackward.h
2026-03-28 - New test: attn_block_grad_test.cpp - compares CP vs standard full attention block gradients (ln, c_attn, c_proj, x.grad) with exact same weights/config - TensorParallelismBeta / DTensor/gpt2_cp_test/attn_block_grad_test.cpp, DTensor/Makefile
2026-03-21 - CP correctness fixes (3 bugs: non-causal attention, backward grad shape, backward merger rescaling) + optimizations (GPU HeadTail, pre-alloc buffers, ring loop cleanup) - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/ContextParallel.h, ContextParallelBackward.h, SDPAOp.h, RingRotator.h, gpt2_cp_test.cpp, tensor/dtensor.cpp
2026-03-23 - HeadTail permutation rewrite: replaced wrong stride permutation with correct head-tail CUDA kernel [0,T-1,1,T-2,...], backward clone fix, k/v edge fix - TensorParallelismBeta / DTensor/tensor/headtail_kernel.cu, headtail_kernel.cuh, dtensor.cpp, ContextParallel.h, ContextParallelBackward.h
2026-03-24 - CP backward Di correction debug & revert, added test_sdpa_backward.cpp - TensorParallelismBeta / SDPAOp.h, ContextParallelBackward.h, ContextParallel.h, tests/test_sdpa_backward.cpp
2026-03-26 - CP vs Normal loss divergence root cause analysis (rcp.approx.f32 P_local/LSE inconsistency) - TensorParallelismBeta / FusedKernels.cu, SDPAOp.h, SDPAMerger.h
2026-03-27 - Context Parallel contiguous() stride bug fix - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/ContextParallel.h
2026-03-30 16:19 - megatronCP.py rewrite using Megatron TEDotProductAttention + TE install - TensorParallelismBeta / megatronCP.py
2026-03-31 - Fused SDPA (FlashAttention tiled kernel) impl + test - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/FusedSDPAKernel.cu, FusedSDPAKernel.h, FusedSDPAOp.h, gpt2_cp_test/test_fused_sdpa.cpp, DTensor/Makefile
2026-03-31 - Fused SDPA (FlashAttention tiled kernel) + CP fair timing test vs MegatronLM - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/FusedSDPAKernel.cu, FusedSDPAKernel.h, FusedSDPAOp.h, gpt2_cp_test/test_fused_sdpa.cpp, gpt2_cp_test/cp_sdpa_compare_test.cpp, DTensor/Makefile
2026-03-31 - Wire fused SDPA into ContextParallel::forward_cp replacing unfused sdpa_forward; CP time 5.547ms->2.659ms (2.1x), gap vs Megatron 5.4x->2.6x - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/ContextParallel.h, FusedSDPAOp.h, DTensor/Makefile
2026-04-01 - Component optimizations (FusedAdamW 7.86x, VectorizedLayerNorm 1.89x fwd/1.75x bwd, FusedLMHead TF32 1.62x fwd) + test files - TensorParallelismBeta / dnn/FusedAdamWKernel.cu, dnn/VectorizedLayerNormKernel.cu, dnn/FusedLMHeadKernel.cu, gpt2_cp_test/test_fused_adamw.cpp, test_vectorized_layernorm.cpp, test_lm_head_opt.cpp
2026-04-02 - Wire FusedLayerNorm into gpt2_cp_test (3 sites); revert FusedAdamW wiring (543MB regression); FusedLayerNorm confirmed no memory regression - TensorParallelismBeta / DTensor/gpt2_cp_test/gpt2_cp_test.cpp, dnn/FusedLayerNormOp.h, dnn/FusedLayerNormOp.cpp
2026-04-04 - CP training norm explosion fix (TC backward kernel bug) - TensorParallelismBeta / FusedSDPABackwardKernel.cu, FusedSDPAKernel.cu
2026-04-04 - PyTorch SDPA backward timing regression debug & fix - TensorParallelismBeta / pytorch_attn_bench.py
2026-04-04 - PyTorch SDPA backward timing regression debug & fix (2-step ring simulation) - TensorParallelismBeta / pytorch_attn_bench.py
2026-04-07 - CP backward graph fix (autograd::mul + autograd::contiguous) + generation shape fix (generation_mode_ in CPAttention) - TensorParallelismBeta / gpt2_cp_test.cpp, ContextParallel.h
2026-04-10 - TDD verification tests for CP Round-Robin + Tiling - TensorParallelismBeta / DTensor/test_roundrobin_verify.py
2026-04-10 - Debug: CP 56k vs PyTorch 64k gap root cause = scalar FP32 kernel (no TC), WMMA TC path disabled in FusedSDPAOp.h; cudaStreamSync fix was no-op, reverted - TensorParallelismBeta / FusedSDPAOp.h, FusedSDPAKernel.cu, ContextParallel.h, AttentionKernels.cu
2026-04-10 - Wire TC forward kernel for CAUSAL+NONE ring steps (self+future), scalar fallback for LEFT_HALF/BOT_HALF; namespace fix OwnTensor::cuda:: - TensorParallelismBeta / FusedSDPAOp.h
2026-04-14 - BOT_HALF TC forward/backward mismatch root cause analysis & fix - TensorParallelismBeta / FusedSDPAOp.h, FusedSDPABackwardKernel.cu
2026-04-27 - MultiTensorKernels scale/grad_norm rewritten to chunked+float4 (sm_86) - TensorParallelismBeta / MultiTensorKernels.cu
2026-04-27 - 124M scaling diagnosis: dual optim.step + cudaSync + clip kernels - TensorParallelismBeta / gpt2_cp_test.cpp, MultiTensorKernels.cu
2026-04-27 - 124M bottleneck rediagnosis: optim 0.06% of step, real cost is attn_cp+bwd compute - TensorParallelismBeta / gpt2_cp_test.cpp
2026-04-28 - CP attention forward+backward kernels updated to support strided Q/K/V/dO/O/LSE/dQ/dK/dV; dropped autograd::contiguous(q) in ContextParallel.h - TensorParallelismBeta / AttentionForward.cu, AttentionBackward.cu, AttentionForward.h, AttentionBackward.h, FusedSDPAOp.h, ContextParallel.h
2026-04-28 - Added strided KV pack kernel; CP forward ring rotator now sends from non-contig K/V; dropped autograd::contiguous on K and V in pre_sharded path - TensorParallelismBeta / KVPackKernel.cu, KVPackKernel.h, ContextParallel.h
2026-04-29 - CP causal LB enable: drop !is_causal_ guard on lb_active + gate inner loadbalance on !pre_sharded (Zigzag round-robin sub-chunking now active for causal training) - TensorParallelismBeta / ContextParallel.h
2026-04-29 - CP LB sub-chunking: drop .contiguous() on strided make_shards halves; rely on strided kernel dispatch (fixes silent data corruption on transposed Q with stride[1]<stride[2]) - TensorParallelismBeta / ContextParallel.h, ContextParallelBackward.h
2026-04-30 - Reverted LB-causal enable: lb_active back to load_balance_ && !is_causal_; LB sub-chunking diverges/explodes norm at 44M correctness check. Skip-future path is the stable baseline. - TensorParallelismBeta / ContextParallel.h
2026-04-30 - LB-causal correctness harness scaffold (TDD RED): cp_lb_causal_correctness_test.cpp + force_lb_active knob in CP ctor + Makefile target - TensorParallelismBeta / cp_lb_causal_correctness_test.cpp, ContextParallel.h, Makefile
2026-05-07 - Step 2 forward activation parity debug - TensorParallelismBeta / gpt2_cp_headtail_fp32.py, gpt2_cp_test.cpp, diff_act_checksums.py - Fixed token dump placement (C++ was pre-sharding, PT post-sharding), found embedding divergence starts with sharded tokens
2026-05-07 - Embedding divergence was missing LOAD_INIT_WEIGHTS env var; embedding parity confirmed bit-exact, real divergence at block_0_out (attention) - TensorParallelismBeta / gpt2_cp_test.cpp, gpt2_cp_headtail_fp32.py
2026-05-07 - CP backward dK/dV ring accumulation fix (memset wiping rotating accumulator) - TensorParallelismBeta / AttentionBackward.cu
2026-05-08 - CP attn kernel HD=64 + T_q!=T_k divergence diagnosed (not yet fixed) - TensorParallelismBeta / AttentionForward.cu
2026-05-08 - HD=64 TC kernel bypass + repro mismatch fix - TensorParallelismBeta / AttentionForward.cu, ContextParallel.h
2026-05-08 - Gradient checksum logging added to both PT and C++ (DUMP_GRADS=1) - TensorParallelismBeta / gpt2_cp_test.cpp, gpt2_cp_headtail_fp32.py
2026-05-09 - CatBackward autograd node + autograd::cat fix restoring wte/wpe gradient chain in shard_sequence - TensorParallelismBeta / CatBackward.h, CatBackward.cpp, ContextParallel.h
2026-05-13 - CP utility-based feature comparison table (Ours vs Megatron vs PyTorch) - TensorParallelismBeta / misc/comparison_table_prompt.md
2026-05-15 - Debug printing restored for all CP ranks at every step (all parameter/gradient values) - TensorParallelismBeta / gpt2_cp_headtail_fp32.py
2026-05-15 - Pre-embedding sequence sharding (contiguous + HeadTail) for CP - TensorParallelismBeta/DTensor/gpt2_cp_test/context_parallel/ContextParallel.h, gpt2_cp_test.cpp
2026-05-16 - LB+causal gate untangled (external_balanced + sub_chunk_active) - TensorParallelismBeta/DTensor/gpt2_cp_test/context_parallel/ContextParallel.h, ContextParallelBackward.h
2026-05-16 - HeadTail kernel rewritten to chunk-level (PyTorch parity, Option 2) - TensorParallelismBeta/DTensor/tensor/headtail_kernel.cu, headtail_kernel.cuh, dtensor.cpp, gpt2_cp_test/context_parallel/ContextParallel.h
2026-05-18 - Re-introduced LOAD_INIT_WEIGHTS pipeline (PT writer + C++ reader) for true bit-parity init between PyTorch and C++ CP training - TensorParallelismBeta/DTensor / Pytorch/gpt2_cp_headtail_fp32.py, gpt2_cp_test/gpt2_cp_test.cpp
2026-05-19 - CP backward switched from merged_out/merged_lse to per-step saved_out/saved_lse to fix LB partial-step gradient corruption - TensorParallelismBeta/DTensor / gpt2_cp_test/context_parallel/ContextParallel.h, ContextParallelBackward.h
2026-05-20 - CP-bwd isolated test built (PT+C++); localized LB-true training bug to exp11 SDPA backward kernel producing ~1.35x over-amplified dQ at head index >= 1 in chunk_3 tail region - TensorParallelismBeta/DTensor / gpt2_cp_test/cp_bwd_isolated_test.cpp, Pytorch/cp_bwd_isolated_test.py, tensor_lib_partial_update_test.cpp, AttentionBackward.cu (suspect)
2026-05-20 - Isolated CP-bwd test LB-OFF run confirms kernel is innocent: all dQ/dK/dV match PT within ~1e-5 (TF32 noise) under non-LB; LB-ON has 100x errors localized to LB-specific accumulation paths (partial dQ on rank 0, head-half pad-then-add on rank 1). Tensor library ops verified correct in isolation. Bug is in LIVE invocation of CPB.h accumulation paths under LB - TensorParallelismBeta/DTensor
2026-05-20 - SDPAMerger partial-merge narrow_view stride bug fix - TensorParallelismBeta / SDPAMerger.h, ContextParallel.h, cp_bwd_isolated_test.{py,cpp}, cp_bwd_diff_deep.py
2026-05-20 - EOD writeup on loss deviation / step-33 norm explosion debug (root cause: SDPAMerger partial-branch narrow_view fed to flat-fast-path elementwise op) - TensorParallelismBeta / SDPAMerger.h, ContextParallelBackward.h, cp_bwd_isolated_test
2026-05-23 - Link DTensor Makefile with BluBlas via WITH_BLUBLAS flag - TensorParallelismBeta / DTensor/Makefile
2026-06-05 - CP step-0 grad parity: root-caused orthogonal grads to LOAD_INIT_WEIGHTS scrambling weights (C++ vs PT param-order mismatch); added name+transpose-aware LOAD_INIT_NAMED loader; forward parity confirmed (emb cos 1.0), residual TF32/GELU divergence left for next session - TensorParallelismBeta / DTensor/gpt2_cp_test.cpp, Pytorch/gpt2_cp_attnstyle_fp32.py, Pytorch/compare_step0_grads.py, Pytorch/compare_fwd_parity.py
2026-06-05 - Backported TF32 round-to-nearest (round_tf32_wmma_frag, +0x1000u half-ULP) into CP attention kernels to remove biased truncation suspected of driving C++ grad-norm drift (->2 vs PT 0.8); 2 fwd + 7 bwd WMMA sites - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/AttentionForward.cu, AttentionBackward.cu
2026-06-05 - Verified TF32 round-to-nearest backport: NO effect (blk0 cos 0.9475 unchanged, grad cos unchanged) -> attention WMMA NOT the divergence source; suspect Linear-GEMM matmul precision / LayerNorm / structural pos-0 anomaly; next = intra-block bisect - TensorParallelismBeta / misc/compareoutputenwattnkernel.md, DTensor/gpt2_cp_test/context_parallel/Attention{Forward,Backward}.cu
2026-06-05 - Compared latest-of-latest TI attn kernels vs TxT/2: GPU is RTX 3060 (sm86) so sm89/Ada + 3xTF32 precision paths cannot run; on sm86 generic path TxT/2 is numerically equivalent to latest -> no port needed, attention RULED OUT as divergence source; suspect Linear-GEMM/LayerNorm/pos-0 - TensorParallelismBeta / misc/compareoutputenwattnkernel.md
2026-06-05 - Reran ws=1 parity with new latest-of-latest TI lib: ZERO change (blk0 cos 0.9475 identical); added PT_TF32 toggle, PT-in-TF32 also ZERO change -> precision RULED OUT, divergence is STRUCTURAL (prime suspect: causal position-0 edge, cos 0.17); grad magnitudes still match (1.0026) - TensorParallelismBeta / Pytorch/gpt2_cp_attnstyle_fp32.py, misc/compareoutputenwattnkernel.md
2026-06-05 - LOCALIZED the C++<->PT divergence to the ATTENTION sub-block via intra-block-0 bisect: emb cos 1.0 -> post-attn 0.941 (gap appears here) -> post-MLP 0.948 (MLP adds nothing); prime cause = C++ TF32-WMMA SDPA vs PT fp32 SDPA (PT_TF32 didnt fix bc allow_tf32 != SDPA backend); added ATTN_FP32=1 toggle to force scalar fp32 attn kernel to confirm - TensorParallelismBeta / gpt2_cp_test/context_parallel/FusedSDPAOp.h, AttentionForward.h, Pytorch/gpt2_cp_attnstyle_fp32.py, compare_fwd_parity.py
2026-06-05 - ATTN_FP32 test: C++ scalar-fp32 vs TF32 attn give DIFFERENT bytes (md5 e7fe29 vs 2d3312) but BOTH cos 0.9408 vs PT -> attention gap is STRUCTURAL not precision (precision fully ruled out). Clue: per-token pos-0 cos 0.18 (attends only to self) -> causal-mask/first-row or CP-ring-vs-PT-SDPA algorithmic diff. NOTE grad mags match + training works (may be benign). Next: dump raw SDPA output pre-c_proj - TensorParallelismBeta / gpt2_cp_test/context_parallel/FusedSDPAOp.h
2026-06-05 - ROOT CAUSE of entire CP grad-orthogonality saga: raw SDPA output matches PT exactly (cos 1.0 incl token0), divergence is 100% attn.c_proj. LOAD_INIT_NAMED loader transposed Linear weights by SHAPE-detection, which silently fails for the only SQUARE Linear (attn.c_proj 768x768) -> loaded W instead of W^T (proved: cos(c_proj_cpp, y@W)=1.0 vs y@W^T=0.003). Fixed loader to transpose by ROLE (explicit needs_transpose per param). Only attn.c_proj behavior changes. Rebuild+rerun pending - TensorParallelismBeta / gpt2_cp_test/gpt2_cp_test.cpp
2026-06-05 - CONFIRMED FIX: after attn.c_proj transpose correction, full ws=1 parity achieved - all forward checkpoints (emb/sdpa/blk0attn/blk0/lnf) cosine 1.000000, step-0 grads cosine 1.000000 (ln_f & c_attn, was 0.42/0.033), TOTAL ratio 0.9999. Entire CP grad-orthogonality investigation CLOSED: root cause was parity-harness square-Linear transpose bug, not model/CP/kernel - TensorParallelismBeta / gpt2_cp_test/gpt2_cp_test.cpp
2026-06-05 - FULL backward parity verified: extended step-0 grad dump to all 149 params per-element cosine (C++ raw blob, Linear weights transposed to PT layout). Result: cosine min=mean=1.000000, #(<0.999)=0, TOTAL ratio 0.9999. C++ CP forward AND backward now bit-faithful to PyTorch at matched init - TensorParallelismBeta / gpt2_cp_test/gpt2_cp_test.cpp, Pytorch/compare_step0_grads.py
2026-06-05 - ws=2 ring parity: C++ ws=2 grad == PT ws=1 (true single-pass full-batch grad) at cosine min 0.9998/mean 0.99998 -> C++ CP RING IS CORRECT. PT attnstyle ws=2 DEVIATES from its own ws=1 (wpe cos 0.80, early-MLP lower) -> PyTorch's load-balanced ring changes the gradient (position-routing fingerprint on wpe). So C++ ws=2 'mismatch' vs PT ws=2 is PT deviating, not C++. Likely root of original Q1 (attnstyle behaves differently) - TensorParallelismBeta / gpt2_cp_test + Pytorch attnstyle
2026-06-05 - DUG OUT Q1 ROOT CAUSE: PT attnstyle CP position-embedding bug. gpt2_cp_attnstyle_fp32.py:423 sets pos=arange(0,T_local) -> at ws=2 BOTH ranks use pos 0..511, so wpe[512:1024] gets ZERO grad (measured: same-pos cos=0.000 for p>=512, 0.10 for 256:512, 0.886 for 0:256) and rank1 tokens get wrong positional encoding. Should be GLOBAL positions per rank (LB-reorder-aware). C++ does this correctly (C++ ws=2 == PT ws=1 single-pass at cos 0.9998). This is why attnstyle loss differed/worsened. Fix = compute global pos per rank - TensorParallelismBeta / Pytorch/gpt2_cp_attnstyle_fp32.py:423
2026-06-05 - Q1 FIX CONFIRMED: sharded global pos through CP LB machinery (gpt2_cp_attnstyle_fp32.py:413, pos_full added to _context_parallel_buffers). After fix: PT ws=2 vs PT ws=1 cosine min=mean=1.000000 (#<0.999=0), wpe[512:] now nonzero; PT ws=2(fixed) vs C++ ws=2 cosine 0.9998. The attnstyle CP pos-embedding bug (pos=arange(0,T_local)) was THE root of Q1. All now in full parity: ws=1 & ws=2, PT & C++, forward & backward - TensorParallelismBeta / Pytorch/gpt2_cp_attnstyle_fp32.py
2026-06-05 - Added PT MODEL_44M env flag mirroring ALL C++ fourtyfour changes (n_embd384/n_layer3/n_head6, global_batch65536, max_steps6768); SAVE_INIT_NAMED now writes config-tagged files. Generated init_weights_named_161M.bin (149 params, 652MB) and init_weights_named_44M.bin (41 params, 177MB), seed 1234, loadable via LOAD_INIT_NAMED with matching config - TensorParallelismBeta / Pytorch/gpt2_cp_attnstyle_fp32.py
2026-06-08 - ROOT CAUSE of C++ CP grad-norm explosion / loss saturation (161M, ~step 800+): param-gradient all-reduce across CP ranks was COMMENTED OUT in training loop -> divergent replicas (each rank steps on its own sequence-shard grad). Ring only sums attn dQ/dK/dV, not Linear/LN/embed weight grads. Fixed: enabled AVG all-reduce (not stale SUM; loss is NOT 1/ws-scaled, AVG verified == single-GPU grad cos 1.0 at step0) after micro-accum, before clip; CP_NO_GRAD_ALLREDUCE=1 to A/B. Independent of overlap work. - TensorParallelismBeta / gpt2_cp_test/gpt2_cp_test.cpp:1278
2026-06-08 - REAL root cause of no compute/comm overlap: DeviceMesh creates the NCCL comm stream with cudaStreamCreate (DEFAULT/blocking flags) at tensor/device_mesh.cpp:190 and passes it to init_process_group -> else-branch uses it as-is. A default-flagged stream implicitly serializes with the legacy NULL stream (stream 0) where ALL compute runs (kernels launch with no stream arg; setCurrentStream never called) -> zero overlap. The cudaStreamNonBlocking edit in processGroupNCCL.cpp was DEAD CODE (stream==0 branch never taken from mesh, so its print never fired). FIX: cudaStreamCreateWithFlags(&comm_stream, cudaStreamNonBlocking) at device_mesh.cpp:190. Algorithm already matches PyTorch (post-exchange-before-compute, copy-to-send, streamWait consume); PyTorch overlaps same ncclDevKernel_SendRecv because its streams are non-blocking. P2P disabled on 2x3060 (canAccessPeer=0) but asyncEngineCount=2. - TensorParallelismBeta / tensor/device_mesh.cpp
2026-06-10 - Overlap run unstable AND CP_NO_OVERLAP=1 gave SAME unstable curve. Cause: comm stream was made non-blocking GLOBALLY+UNCONDITIONALLY in device_mesh.cpp -> removed the implicit stream0<->comm ordering the whole codebase relied on (e.g. K/V pack on stream0 then NCCL send on comm with NO explicit event in non-overlap path) -> both paths race -> identical instability. CP_NO_OVERLAP only gated ring scheduling, not the stream flag. FIX: gate comm-stream flag on CP_NO_OVERLAP in device_mesh.cpp (blocking when off=stable original; non-blocking when on). Overlap-ON path still has unmasked races to fix (validate via ws=2 step0 parity). Cleaner long-term: dedicated CP-ring stream, not global comm-stream flip. - TensorParallelismBeta / tensor/device_mesh.cpp
2026-06-10 - VERIFIED via nsys (ws=2, this 2x3060 box): dedicated-CP-ring-stream overlap WORKS and toggles correctly. CP_NO_OVERLAP=0 -> ring on dedicated non-blocking stream(18), 2324 NCCL-vs-compute(stream7) overlap pairs (1966ms incl bwd-SDPA overlapping SendRecv). CP_NO_OVERLAP=1 -> dedicated stream never created, ring on shared blocking stream(14), 0 overlap pairs. Shared comm stream stays blocking (DP/TP safe). Step-0 grad parity ovl-vs-noovl cos 0.99998. Next: full multi-step =0 run to confirm training stability (global-flip destabilizer removed) + measure throughput. - TensorParallelismBeta / ProcessGroupNCCL, device_mesh.cpp, RingRotator.h2026-06-10 - CP ring-attention overlap race audit: CUDA stream ordering bugs in forward send_buf lifetime, bwd dkv_send lifetime, shared pack_event destroy-while-pending, and cross-step recv slot aliasing - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/RingRotator.h, ContextParallel.h, ContextParallelBackward.h
2026-06-10 - CP ring-attention overlap race audit: CUDA stream ordering bugs in forward send_buf lifetime, bwd dkv_send lifetime, shared pack_event destroy-while-pending, and cross-step recv slot aliasing - TensorParallelismBeta / DTensor/gpt2_cp_test/context_parallel/RingRotator.h, ContextParallel.h, ContextParallelBackward.h
2026-06-10 - Fixed the step-145 overlap explosion: two races found by cpp-review. (1) pack_event was created+destroyed PER CALL and shared across ring iters -> at full pipeline depth cudaEventDestroy fired while the ring stream's pending cudaStreamWaitEvent still referenced it (UB) -> send stopped waiting for pack. (2) stack-local send_buf/dkv_send freed at function return while the FINAL slot's send still in flight (reuse guard never covers last slot). FIX: rotator now OWNS persistent per-slot pack events (pack_ev_[2], never per-call destroy); exchange_buffers takes pack_stream (compute stream) not an event; added post-loop streamWait drain of all pending sends before staging buffers free (fwd + bwd). Step-0 A/B overlap-vs-nooverlap now cos 1.000000. NEEDS multi-step (>145) run to confirm explosion gone. Added CP_NO_OVERLAP_FWD/BWD bisection gates. - TensorParallelismBeta / RingRotator.h, ContextParallel.h, ContextParallelBackward.h
2026-06-10 - Addressed all 4 overlap_fix.md review points: (1) comments recast event-destroy race as unconditional UB not 'deep pipeline'; (2) pack_ev_[2] invariant documented + rotator SELF-GUARDS (streamWait prior slot send before re-record); (3) audited drain coverage - in-loop guard + both-slot drain = complete; (4) drain tail-serialization ceiling documented w/ PyTorch upgrade path (recordStream/TensorShelf, verified from shipped ProcessGroupNCCL.hpp). FIXED sentinel bug: pack stream==NULL stream(0) made overlap silently degrade to blocking; exchange_buffers now takes explicit bool overlap. Verified: toggle flips fwd+bwd correctly, step-0 A/B cos 0.99998 - TensorParallelismBeta / RingRotator.h, ContextParallel.h, ContextParallelBackward.h, misc/overlap_fix.md
2026-06-11 - PyTorch-parity fix: gave the CP ring its OWN NCCL communicator (cp_comm_) in ProcessGroupNCCL, separate from comm_ (loss/param all-reduces). Constructor inits cp_comm_ (2nd ncclGetUniqueId on rank0 + MPI_Bcast + ncclCommInitRank); sendrecv_async_stream/alltoallv_async_stream (the overlap-only ring methods) now use cp_comm_; launch_work_collectives takes work_comm; destructor frees cp_comm_+cp_ring_stream_. Zero plumbing through model/mesh (ring already used only the stream methods). Matches PyTorch (separate comm per parallelism dim) -> ring + all-reduces never share a comm. Also reconciled overlapfixplanfix.md review: its criticals (per-slot send-work streamWait guard exch_work/dkv_work; events absorbed into rotator as pack_ev; .data()+k_numel) were ALREADY in code. Verified: toggle ok, no deadlock, step-0 A/B cos 0.99999. PENDING: full >150-step run to confirm slow divergence fixed. - TensorParallelismBeta / process_group/ProcessGroupNCCL.{h,cpp}
2026-06-11 - Separate cp_comm_ did NOT fix the overlap explosion (still diverges ~step140) -> NOT the shared comm. Found the REAL bug: dkv_rotater is a LOCAL unique_ptr in ContextParallelBackward (line 142), destroyed at function return; grad_key/grad_value were returned to autograd as VIEWS into its recv buffer (final_flat.narrow().reshape()) -> dangling/use-after-free. No-overlap hides it (CPU-blocking + caching allocator delays reuse); overlap's async CPU-run-ahead lets the freed buffer be reused before downstream c_attn backward consumes grad_key -> rare corruption -> grad spikes at depth. Explains overlap-only + depth-dependent + spiky + survived-separate-comm. FIX: .clone() grad_key/grad_value at the LB final collect (ContextParallelBackward.h:543). Step-0 A/B cos 0.999999 (no regression). Running full 160-step confirmation. - TensorParallelismBeta / gpt2_cp_test/context_parallel/ContextParallelBackward.h
2026-06-11 - CP overlap divergence root-caused & fixed (cross-stream allocator-reuse race at ring tail; GPU-side drain -> CPU wait) - gpt2_cp_test / ContextParallel.h, ContextParallelBackward.h
2026-06-15 - CP overlap race-fix docs review (both false-positive); fixed misleading drain comment & AllGather sync - OldPush / ContextParallel.h, RingRotator.h
2026-06-16 - PyTorch 2.9.0+cu128 TF32-vs-FP32 precision research (Linear/SDPA/ring-attn/LayerNorm/GELU on RTX 3060 sm86); confirmed allow_tf32 default off + SDPA independent of allow_tf32 vs v2.9.0 source - TensorParallelismBeta/DTensor (research only, no files modified)
2026-06-16 - Audited C++ vs PyTorch CP FP32/TF32 placement parity: MISMATCH at Linear GEMMs (C++ TF32 hardcoded vs PT FP32 default) and SDPA attention (C++ TF32-WMMA vs PT FP32 mem-efficient; PT_TF32 doesn't affect SDPA, ATTN_FP32 doesn't affect C++ Linear); MATCH on LayerNorm/GELU/softmax/CE/embeddings; full FP32 parity needs PT-default+ATTN_FP32=1+C++ GEMM code change, full TF32 impossible (PT SDPA has no TF32 path) - OldPush / gpt2_cp_test.cpp, gpt2_cp_attnstyle_fp32.py, GenMatmul.cu, AttentionForward.cu, AttentionBackward.cu, FusedSDPAOp.h
2026-06-16 - PyTorch 2.9.0 mem-efficient SDPA fp32-on-sm86 MMA-op source audit: confirmed DefaultGemmType<ArchTag,float,kMinComputeCapability>=80> uses OpClassTensorOp + cutlass::arch::OpMultiplyAddFastF32 (3xTF32 tensor cores), compile-time fixed, NOT gated by allow_tf32; both Q.K^T (MM0) and P.V (MM1) GEMMs use it - TensorParallelismBeta/DTensor (research only, no files modified)
2026-06-17 - CP TxT/2 kernel question: confirmed PyTorch 2.9 needs NO specialized TxT/2 kernel - zigzag LB produces 3 shapes (2cx2c causal diag, 2cxc + cx2c dense rectangular) all dispatched to the SAME SDPA op via tensor slicing in templated_ring_attention; rectangular cases are dense (is_causal=False, no MFU waste); only requirement = one kernel handling Tq!=Tk - TensorParallelismBeta / torch attention.py:442-466, AttentionForward.cu
2026-06-18 - sm89 drop-in FusedSDPAOp variant: calls Ada PTX-MMA fwd/bwd kernels directly (no arch gate), throws on CP offsets/T_q!=T_k - TensorParallelismBeta / gpt2_cp_test/context_parallel/FusedSDPAOp_sm89.h
2026-06-18 - CP-aware sm89 attention kernels: ported arch/AttentionForward_sm89 & AttentionBackward_sm89 (exp12) to OwnTensor::cp with T_q/T_k+q_offset/k_offset global-causal masking; wired into FusedSDPAOp_sm89.h - TensorParallelismBeta / gpt2_cp_test/context_parallel/AttentionForward_sm89.cu, AttentionBackward_sm89.cu, AttentionForward_sm89.h, AttentionBackward_sm89.h, FusedSDPAOp_sm89.h
2026-06-18 - Fixed sm89 link error (undefined mem_efficient_attn_forward_strided): removed ATTN_FP32 scalar-fp32 fallback from FusedSDPAOp_sm89.h for pure-sm89 build; updated makefile_in_ada.md (arch sm_89, sm89 CP kernel objs+rules+linking) - TensorParallelismBeta / FusedSDPAOp_sm89.h, makefile_in_ada.md
2026-06-18 - Ported the CP forward cp.async race fix into sm89 kernel: wait_group<1> on final KV iter (no next-K group) waits for nothing -> P@V reads unarrived V under co-exec; now wait_group<0> when !has_next - TensorParallelismBeta / gpt2_cp_test/context_parallel/AttentionForward_sm89.cu
2026-06-18 - Ada grad-norm spike debug: exp12 abandoned (catastrophic), exp11 leaves a transient spike; gated CP forward sm89-arch dispatch behind CP_FWD_SM89 env (default OFF -> race-fixed WMMA fwd fused_attn_forward_kernel_tc used everywhere) to A/B test arch sm89 forward (unfixed wait_group race) as spike cause - TensorParallelismBeta / gpt2_cp_test/context_parallel/AttentionForward.cu2026-06-20 - Replaced hand-rolled DataLoaderLite with native torch.utils.data DataLoader (byte-identical batches, per-shard cycling, drop_last) in attnstyle CP parity script; audited remaining non-native impls (CPAttention CP bracketing, get_lr, CudaTimer, shard_for_cp) - OldPush / Pytorch/gpt2_cp_attnstyle_fp32_dl.py
2026-06-22 - Root-caused AllGather-only CP validation-loss degradation: backward dK/dV travelling accumulator used AllGatherRingRotator, whose exchange_buffers is a no-op after the first one-shot gather, so per-step cross-rank dK/dV grads were dropped (forward/step-0 val correct, gap grew monotonically over training: +0.16 by step 6767 vs P2P/AlltoAll). Fix: force dkv_rotater to AlltoAll when rotator_type==AllGather, mirroring PyTorch's forced ALL_TO_ALL backward dkv_rotater - OldPush / gpt2_cp_test/context_parallel/ContextParallelBackward.h
2026-06-22 - Found 2nd AllGather-only bug (forward): persistent kv_rotator reused across steps, but AllGatherRingRotator gathers once and never re-gathers -> serves step-0 stale KV with drifting (rank-idx) shard index from step 2 on (P2P/AlltoAll re-communicate so immune; step-0 matches; grows/noisier over training; survives the backward dkv fix). Fix: recreate AllGather rotator each forward in ContextParallel.h - OldPush / gpt2_cp_test/context_parallel/ContextParallel.h
2026-06-22 - Made AllGather CP rotator non-blocking/overlapped like PyTorch: added ProcessGroupNCCL::all_gather_async_stream (async ncclAllGather on cp_comm/cpRingStream, returns Work), rewrote AllGatherRingRotator to issue async gather on ring stream w/ pack-event ordering + recordStream + GPU streamWait (overlaps step-0 local SDPA), added RingRotatorBase::reset() and switched forward to reset() the persistent rotator each call (re-gathers fresh, no alloc churn / no event-destroy UB) instead of recreating it - OldPush / process_group/ProcessGroupNCCL.h+.cpp, gpt2_cp_test/context_parallel/RingRotator.h, ContextParallel.h
2026-06-22 - Added one-time "[CP ring AllGather] OVERLAP/NO-OVERLAP" stderr print in AllGatherRingRotator::exchange_buffers to confirm the active forward rotator path (mirrors the existing [CP ring A2A] backward print) - OldPush / gpt2_cp_test/context_parallel/RingRotator.h
2026-06-22 - AllGather rotator micro-opt: index_chunk now returns a zero-copy narrow_view into the gathered buffer instead of allocating+cudaMemcpy a fresh chunk per ring step (trims per-step overhead/jitter; safe via existing wait/streamWait + persistent buffer + pack-event ordering) - OldPush / gpt2_cp_test/context_parallel/RingRotator.h
2026-06-23 - 6000 Ada server build port: LDFLAGS/MPI path -> /usr/lib/x86_64-linux-gnu/openmpi/lib, arch sm_86->sm_89, drop -lnvToolsExt (NVTX3 header-only), nvlink missing gpt2_cp_test.o root-caused to -j/-k OOM on -O3 giant TU; serial build fixed - TensorParallelismBeta / DTensor/Makefile, gpt2_cp_test/gpt2_cp_test.cpp
2026-06-23 - Generalized sequence length T in attnstyle CP script: derive block_size=max(1024,T) (was hardcoded 1024) so wpe table fits any T; added BLOCK_SIZE env override + fast asserts for T%cp_world_size and BT|total_batch_size; noted parity break vs init_weights/C++ ref - OldPush / Pytorch/gpt2_cp_attnstyle_fp32.py
2026-06-24 - Memory-scaling sweep harness: env-parametrized T/config/weight-tying/rotator + 2-step nvidia-smi probe mode in both CP scripts; portable batch sweep + table aggregator (collapsed peak_mb column) - OldPush / gpt2_cp_test/gpt2_cp_test.cpp, Pytorch/gpt2_cp_attnstyle_fp32.py, mem_scaling_sweep.sh, mem_scaling_table.py
2026-06-24 - Mem-scaling: added allocator-breakdown instrumentation to C++ probe (reserved/active/requested/frag - found frag only 6.4%, footprint is peak working set held reserved); fine-grained binary-search T-cutoff mode in sweep + limits summary (max_T_ok/first_T_oom) in table gen - OldPush / gpt2_cp_test/gpt2_cp_test.cpp, mem_scaling_sweep.sh, mem_scaling_table.py
2026-06-25 - CP GPT2 memory: found nsys_report=1 forced probes to 1 step (step-0, no optimizer state) invalidating sweep; fixed precedence; gated per-step out/lse save (dead weight in default backward) -> -925MB reserved, +11% max-T, zero throughput cost, exact parity - OldPush / gpt2_cp_test.cpp, context_parallel/ContextParallel.h, mem_scaling_sweep.sh
2026-06-27 - CP GPT2 shared forward rotator (CP_SHARE_FWD_ROTATOR, env-gated): one rotator+staging across all layers instead of per-layer -> -1100MB reserved at 124M T2048 (9393->8293), throughput unchanged (9600 tok/s), step0 loss byte-identical all rotators, 160-step grad-norm stable - OldPush / gpt2_cp_test/context_parallel/ContextParallel.h, gpt2_cp_test/gpt2_cp_test.cpp2026-06-27 - Wrote CP repo Makefile in Tensor-Implementations template style (auto-detect SM/CUDA, source discovery, run/run-snippet/run-folder/print_sm/help/clean/rebuild, libtensor delegate) - CP / Makefile
2026-06-27 - CP include-path fix (gpt2_cp_test/context_parallel -> context_parallel, 8 files); confirmed tensor/dtensor.h, dnn/DistributedNN.h, dnn/FusedLayerNormOp.h still required by DeviceMesh/DTensor/fused_layer_norm usage - CP / context_parallel/, Scripts/BluTrain/gpt2_cp_test.cpp
2026-06-27 - CP self-contained: repointed DeviceMesh include tensor/dtensor.h -> process_group/device_mesh.h (script + ContextParallel.h + ContextParallelBackward.h), dropped unused dnn/DistributedNN.h, swapped dnn::fused_layer_norm -> OwnTensor autograd::layer_norm (3 sites, signature-identical), fixed dnn/AttentionKernels.h -> ops/helpers/AttentionKernels.h and device_mesh.cpp self-include - CP / Scripts/BluTrain/gpt2_cp_test.cpp, context_parallel/, process_group/device_mesh.cpp
2026-06-27 - Wrote CP README.md (build via make libtensor + make, run via make run NP=, env-var config table CP_*/LOAD_INIT_*, hardcoded data_root pre-run edit caveat, make targets, PyTorch parity ref) - CP / README.md
2026-06-27 - Refactor: moved CP load-balancing out of core dtensor into gpt2_cp_test/context_parallel/ (new LoadBalancer.h = LoadBalancer/HeadTail/SDPA_Merger decls; new ContextParallel.cpp = DTensor::context_parallel_shard/unshard + HeadTail defs). dtensor.h keeps LoadBalancer fwd decl; CP users include LoadBalancer.h; Makefile GPT2_CP_TEST_SRCS += ContextParallel.cpp. Builds clean - OldPush / tensor/dtensor.{h,cpp}, gpt2_cp_test/context_parallel/{LoadBalancer.h,ContextParallel.cpp}, Makefile2026-06-27 - CP build wiring: Makefile .cpp discovery now spans context_parallel/+process_group/ (picks up ContextParallel.cpp, headtail_kernel.cu auto-discovered); re-fixed include paths (gpt2_cp_test/->context_parallel/, tensor/dtensor.h->process_group/device_mesh.h, tensor/headtail_kernel.cuh->context_parallel/, dropped unused dnn/DistributedNN.h + using dnn). Remaining blocker: device/AllocationTracker.h / OwnTensor::AllocationTracker absent from submodule - CP / Makefile, context_parallel/, Scripts/BluTrain/gpt2_cp_test.cpp
2026-06-27 - CP Makefile: adapted to BluTrain monorepo layout (Tensor-Implementations -> BluTrain/Tensor-Implementations); added libprofiler build (BluTrain/Profiler/src/.cpp -> libprofiler.so, pure host C++), -IBluTrain/Profiler/include, -lprofiler + rpath, libtensor now depends on profiler; run targets LD_LIBRARY_PATH include Profiler/lib; fixed script include device/AllocationTracker.h -> profiler/AllocationTracker.h. libprofiler.so builds clean - CP / Makefile, Scripts/BluTrain/gpt2_cp_test.cpp, .gitignore
2026-06-29 - CP made independent of dnn layer: re-applied dnn::fused_layer_norm -> OwnTensor::autograd::layer_norm (3 sites, signature-identical) and removed dangling 'using namespace OwnTensor::dnn' in gpt2_cp_test.cpp + ContextParallel.h (dnn includes already commented). No DistributedNN/dnn dependency remains - CP / Scripts/BluTrain/gpt2_cp_test.cpp, context_parallel/ContextParallel.h
2026-06-29 - CP dnn-independence completed: ContextParallel base class DModule->dropped (no base-class API used; forward_cp called directly), replaced commented tensor/dtensor.h with process_group/device_mesh.h in ContextParallel.h + script (fixes DeviceMesh undeclared), re-swapped dnn::fused_layer_norm->autograd::layer_norm x3 + removed using dnn, AllocationTracker include -> profiler/. Main TU now -fsyntax-only clean - CP / context_parallel/ContextParallel.h, Scripts/BluTrain/gpt2_cp_test.cpp
2026-06-29 - Adapted Tests/mem_scaling_sweep.sh to CP layout: REPO_ROOT=script/.. (Tests/), PT_SCRIPT->Scripts/Pytorch/, CPP_EXEC->build/gpt2_cp_test_exec, build 'make gpt2_cp_test' in DTENSOR_DIR -> 'make' in REPO_ROOT, OUT_DIR/table msg under Tests/. mem_scaling_table.py unchanged (OUT_DIR arg, layout-agnostic). bash -n clean - CP / Tests/mem_scaling_sweep.sh
2026-06-29 - CP attention-kernel compatibility analysis: traced production fused GQA(QK-norm+RoPE+SDPA sm103) vs CP ring needs; wrote requirements doc (dual q_offset/k_offset global-causal mask, separate T_q/T_k, LSE-as-output, pure-fn backward, de-fuse RoPE/QK-norm for HeadTail non-contig positions, keep GQA-grouped KV in ring) - CP / .mdfiles/cp_attention_kernel_requirements.md
2026-06-29 - CP fused-RoPE load-balance research: verified PyTorch torchtitan + Megatron-LM (get_pos_emb_on_this_cp_rank, index_select [r,2cp-1-r]) + flash-attn -> NO framework fuses RoPE into the attn kernel for CP; they shard the rotary cos/sin cache with the same zigzag perm as tokens. To keep RoPE fused, kernel must take per-token q_pos_ids/k_pos_ids (k travels ring) not scalar pos_offset; added section 4.2 - CP / .mdfiles/cp_attention_kernel_requirements.md
2026-06-29 - CP kernel-requirements doc revised after adversarial review: corrected fabricated "CP already de-fuses RoPE" claim (CP layer is GPT-2/learned-wpe + de-fused sm89 SDPA, NO RoPE anywhere); split backward contract into target-sm89 vs fused-production (internal RoPE/QK-norm/D recompute, dq/dk_gamma must accumulate across ring steps); fixed merged-LSE-backward claim (corrupts LB partial steps, USE_PER_STEP_LSE); distinguished production bf16 vs CP sm89 TF32 + per-block-O rounding; flipped recommendation to fused 4-delta per no-de-fuse constraint - CP / .mdfiles/cp_attention_kernel_requirements.md
2026-06-29 - Wrote offset-based RoPE fused-attn CP implementation plan (4-delta interface, contract-first parallel strategy, de-fused stand-in oracle, 5 phases incl gamma-grad accumulation + per-step LSE) - CP / .mdfiles/cp_rope_offset_implementation_plan.md
2026-06-29 - CP fused-RoPE Phase 1a (TDD GREEN): added pure-CPU RopeDeltas.h (compute_deltas 4-delta offset helper: non-CP/contiguous/HeadTail-LB + sub-chunk collapse) + standalone cp_rope_deltas_test.cpp (19/19 checks pass, incl N=2 {0,1,6,7} global-pos reconstruction). Additive only, no existing files touched. Next: Phase 0 interface freeze before editing ContextParallel.h/FusedSDPAOp.h - CP / context_parallel/RopeDeltas.h, Tests/cp_rope_deltas_test.cpp
2026-06-29 - CP fused-RoPE Verification 1b (GREEN): added pure-CPU RopeGammaAccum.h (gamma_accumulate_step = LOCAL add no rotation; gamma_reduce_reference = AVG) + cp_gamma_accum_test.cpp (12/12, incl negative SUM-discriminator proving the test catches the AVG-vs-SUM factor-N bug). Encodes replicated-param gamma-grad spec (local sum over ring steps -> existing op_t::avg all-reduce). Additive only. Next gate: Phase 0 interface freeze before existing-file edits - CP / context_parallel/RopeGammaAccum.h, Tests/cp_gamma_accum_test.cpp
2026-06-29 - CP fused-RoPE Phase 1b/1c (forward plumbing, additive): new FusedRoPESDPA.h (frozen Phase-0 contract extern decls gqa_fused_rope_cp_forward/backward + sdpa_fused_forward_rope/backward_rope wrappers, body gated by -DCP_FUSED_ROPE so default build needs no team symbol); ContextParallel.h gained use_rope members + enable_rope() setter + gated RoPE branch at fwd call site (else = exact existing call). Verified: FusedRoPESDPA.h + ContextParallel.h syntax-clean (mpic++) in BOTH default and CP_FUSED_ROPE modes; default GPT-2/wpe path byte-identical (use_rope_=false); both CPU unit tests still pass. Backward (1d) not yet wired - CP / context_parallel/FusedRoPESDPA.h, ContextParallel.h
2026-06-29 - CP fused-RoPE Phase 1d (backward gamma wiring, additive/gated Option A): ContextParallelBackward node Node(3)->Node(has_gamma?5:3); added use_rope_/q_gamma_/k_gamma_/cos_sin_cache_/rope_eps_/has_gamma_ members + ctor params; apply() ring-loop gated branch -> sdpa_fused_backward_rope with matching 4-delta derivation, dq/dk_gamma accumulated LOCALLY (no rotation); both return sites append gamma grads when has_gamma_; release clears gamma. ContextParallel.h node ctor passes RoPE args + wires gamma edges 3,4 (gated). Verified syntax-clean (mpic++) default + CP_FUSED_ROPE; GPT-2/wpe path byte-identical (Node(3), 3-grad, edges 0-2); both CPU tests pass - CP / context_parallel/ContextParallelBackward.h, ContextParallel.h
2026-06-29 - CP fused-RoPE Phase 2 stand-in: wrote Tests/cp_rope_standin_parity.cpp (de-fused rms_norm+rope+pure-SDPA CP oracle; ws=2 forward + dQ/dK/dV/dgamma parity vs single-GPU ref, contiguous + HeadTail) + additive Makefile target cp-rope-standin/run-cp-rope-standin. Harness COMPILES clean (mpic++). RUNTIME parity BLOCKED: libtensor.a unbuildable in this env - framework NvmlMonitor.cpp uses CUDA-12 nvmlDeviceGetComputeRunningProcesses_v2 (CUDA-13 renamed to v3); lib/objects holds only host .o (GPU .cu objects e.g. ReductionImplGPU.o absent), so manual archive link-fails on dispatch_reduction_gpu. CP RoPE code complete+static-verified; runtime needs working framework build. Removed my incomplete host-only libtensor.a - CP / Tests/cp_rope_standin_parity.cpp, Makefile
2026-06-29 - CP fused-RoPE Phase 2 RUNTIME PASS (2x3060, ws=2): NVML v3 fix built libtensor cleanly; cp_rope_standin_parity_exec ran -> de-fused RoPE CP ring matches single-GPU ref on ALL gates (forward + dQ/dK/dV/dq_gamma/dk_gamma), BOTH contiguous and HeadTail LB, cos=1.0000000 maxdiff 1e-4..1e-6 (TF32 noise). Validates CP orchestration under RoPE + HeadTail + gamma grad flow; confirms additive ContextParallel/Backward edits do NOT break existing pure-SDPA CP path at runtime. Oracle established for Phase 3 - CP / Tests/cp_rope_standin_parity.cpp, BluTrain/Tensor-Implementations/src/device/NvmlMonitor.cpp
2026-06-30 - CP Ulysses (DeepSpeed) seq-parallel attention, additive opt-in alongside ring: new UlyssesAttention.h (combine/partition all-to-all layout swaps [B,H,Tl,D]<->[B,Hl,T,D] + plain rank-order gather, raw ncclAllToAll on cp_comm, NO unloadbalance de-zigzag) + UlyssesAttentionBackward.h (Node(3); reverse all-to-alls + ONE sdpa_fused_backward, reuses existing CP kernel, branches on PERSISTED unshard flag); ContextParallel gained use_ulysses_/enable_ulysses()/forward_ulysses gate at top of forward_cp (ring path byte-identical when off; pre_sharded rejected, MHA-only, contiguous shard). Reuses sdpa_fused_forward/backward (no new kernel; sq causal T_q==T_k q_off=k_off=0). RUNTIME PASS 2x3060 ws=2: cp_ulysses_parity forward+dQ/dK/dV cos=1.0000000 (unshard true+false); ring rope-standin regression still ALL PASS rebuilt against edited header. ws=4 axis-coverage pending 4-GPU box (test device-mapping oversubscription-safe) - CP / context_parallel/UlyssesAttention.h, UlyssesAttentionBackward.h, ContextParallel.h, Tests/cp_ulysses_parity.cpp, Makefile
2026-06-30 - CP Ulysses wired into gpt2_cp_test full model (CP_ATTN_MODE=ulysses): main() forces load_balancing=false (contiguous shard for both pre-embed shard_sequence_pre_embed AND CP layer); CPAttention ctor calls cp_->enable_ulysses(). Relaxed forward_ulysses to accept pre_sharded=true when !load_balance_ (training feeds local contiguous [B,H,Tl,D]; zig-zag still rejected via checkable load_balance_ guard); UlyssesAttentionBackward gained pre_sharded_ flag -> returns LOCAL grads (no gather) when input was pre-sharded, full gather otherwise. SMOKE PASS 2x3060 (44M T=256 ws=2): trains end-to-end, val 10.90, loss 10.907->10.896 decreasing, finite norms, generation OK, no NaN/throw. ws=2 parity still ALL PASS - CP / Scripts/BluTrain/gpt2_cp_test.cpp, context_parallel/ContextParallel.h, UlyssesAttentionBackward.h
2026-07-01 - CP Ulysses GQA/MQA (DeepSpeed Strategy B, strictly additive): new UlyssesGQAAttentionBackward.h node + forward_ulysses_gqa method (reached via 1-line guard in forward_ulysses when nkv<nq; MHA body byte-unchanged); KV all-to-all carries only nkv heads (comm x1), partial-replicate by rep=P/nkv only when nkv<P, grouped broadcast done LOCALLY (kv_local->nq_local) since CP SDPA is MHA-only, backward = two group-reduces (adjoint of local broadcast g_local, then of replication rep). Added head_repeat_interleave (narrow+cat) + head_group_reduce (narrow + raw Tensor operator+) to UlyssesAttention.h. Bug found+fixed during bringup: OwnTensor::reduce_sum over a middle axis of a 5D tensor gave wrong dK/dV (dQ passed since no reduce) -> reimplemented group-reduce as explicit narrow+add. RUNTIME PASS 2x3060 ws=2: cp_ulysses_parity forward+dQ/dK/dV cos=1.0000000 for GQA nkv=4,2 + MQA nkv=1 + MHA(unshard t/f); ring rope-standin regression still ALL PASS; gpt2_cp_test TU parses clean. ws=4 GQA configs (nq=16 nkv=4/1) pending 4-GPU box. GQA is CP-layer capability (parity test); gpt2 model stays MHA - CP / context_parallel/UlyssesGQAAttentionBackward.h, UlyssesAttention.h, ContextParallel.h, Tests/cp_ulysses_parity.cpp
2026-07-02 - CP Ulysses FUSED RoPE+QKnorm+GQA path (Llama-style, additive): new enable_ulysses_fused()+forward_ulysses_fused()+UlyssesFusedGQAAttentionBackward.h calling the team's bf16 gqa_fused_flash_attn_forward/backward (kernel is GQA-native so NO local head-broadcast; only nkv<P replication remains; pos_offset=0 since Ulysses gathers full contiguous seq). Mirrors production sdpa_gqa_fused pack (as_type bf16 + cat) / narrow_view split; gamma grads SUM-all-reduced over CP group (op_t::sum). Reused v2 head_repeat_interleave/head_group_reduce. Parity run_fused_mode (rms_norm+rope+GQA ref, bf16 tol). FOUND KERNEL BUG: fused GQA backward* returns wrong dQ+dq_gamma for true GQA 1<Nkv<Nq (nkv=2 g=4: dQ cos0.12, dq_gamma cos0.27) while fwd/dK/dV/dk_gamma correct; reproduced at ws=1 (kernel called directly, no CP) -> it's GQA_fused_bwd_sm103.cu dQ path, NOT CP. MHA(nkv=8,G=1) + MQA(nkv=1,G=8) fused PASS all grads (cos>0.9999) ws=1&2; marked intermediate-GQA dQ/dq_gamma XFAIL. All other gates (v1 MHA, v2 plain-GQA/MQA incl nkv=2, fused MHA/MQA) + ring rope-standin regression PASS. gpt2 TU clean - CP / context_parallel/UlyssesFusedGQAAttentionBackward.h, ContextParallel.h, Tests/cp_ulysses_parity.cpp
2026-07-02 - Wrote EOD report (EOD.md first-person narrative style) covering the GQA-based Ulysses attention implementation - CP / .mdfiles/EOD_2026-07-02_gqa_ulysses.md
2026-07-03 - Ada CP grad explosion localized to FORWARD compute/comm overlap race (CP_NO_OVERLAP_FWD=1 clean to 111 steps; full CP_NO_OVERLAP clean). Next probe CP_FWD_SYNC_RECV=1 to split recv-consume event ordering (next_buffer_streamordered work->streamWait) vs pack/allocator path. Interim stable config: CP_NO_OVERLAP_FWD=1 - TensorParallelismBeta / gpt2_cp_test/context_parallel/ContextParallel.h, RingRotator.h
2026-07-03 - CP fwd-overlap race fix: CP_DRAIN_AT=1 pinned the backward-save clone (saved_k_chunks[i]=curr_k.clone()) as the victim racing the concurrent ring exchange; hoisted Step3 skip/causal + the conditional clone ABOVE exchange_buffers (so pack_ev orders the clone before ring send/recv), moved skip-continue below exchange, preserved backward is_valid() skip detection - TensorParallelismBeta / gpt2_cp_test/context_parallel/ContextParallel.h2026-07-06 - CP Ulysses fused-GQA bug: sharpened root cause via new run_fused_vs_seq_mode() isolator in cp_ulysses_parity.cpp - compares FUSED kernel directly vs the team's OWN plain (non-fused) GQA kernel (gqa_flash_attn_forward/backward), fed identical RoPE+QKnorm'd Q/K (no hand-rolled matmul/softmax ref). Non-causal (plain kernel has no causal support): dQ/dK/dV/dq_gamma/dk_gamma ALL match fused perfectly (cos>0.9999) for nkv=2, nkv=4(nq=12), and nkv=8(MHA) - so GQA group-indexing (hkv=hq/G) is NOT the bug. Combined with prior causal finding (dQ/dq_gamma wrong only for 1<nkv<nq when causal), isolates the defect to the fused backward's CAUSAL dQ path interacting with G>1 grouping specifically (dK/dV use a separate KV-tile-centric path, unaffected). Verified at ws=1 and ws=2, deterministic - CP / Tests/cp_ulysses_parity.cpp
2026-07-07 - CP Ulysses GQA bug MAJOR reframe: added 3rd leg to run_fused_vs_seq_mode() comparing the REAL plain (non-fused) GQA kernel (gqa_flash_attn_forward/backward, non-causal) directly against ref_sdpa - it does NOT match for true GQA (1<nkv<nq): dQ cos=0.09 (nkv=2) and cos=-0.07 (nkv=4, sign-flipped!), dq_gamma also fails; MHA (nkv=8) passes cos=0.99. Same failure signature as the fused kernel. This proves the bug is NOT fused-specific and NOT causal-specific - it's shared by BOTH plain and fused GQA backward dQ kernels for true GQA, non-causally too. Ruled out ref_sdpa/expansion methodology as the culprit since the identical construction matches a 3rd unrelated kernel (our own CP sm89 path via run_gqa_mode) at cos=1.0 for the same G values. dQ is per-query-head independent (no cross-head sum needed, unlike dK/dV which DO match) so there's no legitimate algorithmic reason for it to diverge - strong signature of a real, shared dQ-kernel defect (GQA_fused_bwd_sm103.cu documents itself as a "sibling" of GQA_bwd_sm103.cu, likely inherited bug) rather than a causal-masking or fusion-specific issue - CP / Tests/cp_ulysses_parity.cpp
2026-07-08 - MAX_LR sweep script (MultiStep sched) + MAX_LR env override - BluTrain / gpt2_fmha_ddp_sched.cpp, run_maxlr_sweep.sh
2026-07-08 - Fixed MAX_LR sweep env: rebuilt libdist (GLIBCXX mismatch), fixed MPI 4.1.2/5.0.9 mismatch in Makefile, local data_root + INIT_FROM_BIN; smoke-verified all 9 LRs - BluTrain / gpt2_fmha_ddp_sched.cpp, Makefile, run_maxlr_sweep.sh
2026-07-09 - MAX_LR sweep complete (9/9 runs, 6767 steps each); compared via compare_sched_fmha.py, best val loss at 1.5e-3 - BluTrain / compare_maxlr_*.png, maxlr_sweep_logs/
2026-07-09 - Ported disk-state checkpointing (CheckpointManager save/resume) into CP training script, run-numbered filenames tied to log index, preserved across resume; env CP_CHECKPOINT/CP_CKPT_FREQ/CP_CKPT_NEW_RUN/CP_CKPT_RESUME; auto-resume latest incomplete run, new run on complete, targeted resume; rank-0-only save (shard_dir=false shared path); Makefile host std -> c++2a for Checkpointing.h C++20 API. 2x3060 verified all paths - CP / Scripts/BluTrain/gpt2_cp_test.cpp, Makefile
2026-07-09 - Implemented canonical (HF/Qwen3) YaRN in build_rope_cache: NTK-by-parts freq blend via find_correction_dim dim-index ramp + attention temperature m=0.1ln(s)+1 baked into cos/sin (m^2 logit temp, no kernel change); env-overridable YARN_SCALE/ORIG_MAXPOS/BETA_FAST/BETA_SLOW (scale=1 => bit-identical std RoPE); build-on-CPU-then-.to(device) bypasses uniform-freq CUDA kernel; robustness guards (R2 scale/seqlen warn, R3 init log, R5 bad-range throw). YARNOps.o symbol wins link over libtensor RopeOps. 2x3060 verified: standin parity 100% PASS cos=1.0 both no-env and YARN_SCALE=8, m=1.2079 matches, hd=128 low=11/high=36; ulysses fails only in pre-existing GQA-bwd dQ/dq_gamma bug (identical at scale=1) - CP / context_parallel/YARNOps.cpp2026-07-09 - CORRECTION to 2026-07-09 YaRN entry ("pre-existing GQA-bwd dQ/dq_gamma bug"): this was WRONG, not a kernel bug. Root-caused via precision-injection test: took the G=1 (MHA, no GQA at all) case, bf16-round-tripped K on the ref_sdpa side only (no kernel change), and dQ cos collapsed to 0.108 - identical magnitude to the "GQA bug" signature seen for 1<Nkv<Nq. Proves dQ is simply sensitive to bf16-vs-fp32 K precision mismatch between the fused/normal GQA kernels (bf16-internal) and the fp32 ref_sdpa reference; nothing GQA-specific. Kernel is correct. Retracting all prior "fused/normal GQA backward dQ/dq_gamma bug" claims (2026-07-02, 2026-07-07 entries) - CP / Tests/cp_ulysses_parity.cpp
2026-07-16 - bluscriptCP re-warmup anchor fix + long-context PPL eval: (1) LR re-warmup was anchored at start_step (current resume step) so every resume of a context-extension run restarted the warmup ramp; added CP_REWARMUP_ANCHOR= to pin the ramp to a fixed absolute step (default -1 = old start_step behavior), so resuming mid/post-ramp continues the same schedule; log now prints anchor+resume+phase (mid-ramp/past-ramp). (2) Added CP_EVAL_PPL forward-only eval mode: NP=1/CP_SIZE=1 (T_local==T, no cross-rank gather), loads ckpt via CP_CKPT_RESUME, teacher-forces CP_EVAL_WINDOWS windows on held-out val split, per-position NLL=logsumexp-logit[tgt] -> pos,count,mean_nll,ppl CSV (CP_EVAL_OUT); skips training + log/config writes (/dev/null). Base-ckpt+YARN_SCALE=1 (naive extrapolation) vs extended-ckpt+YARN_SCALE=4 curves show blow-up-vs-flat past orig ctx. (3) Notebook cell in graph_bluscript_cp.ipynb plots the CSVs + before/after-ctx PPL ratio verdict. Compile-verified (mpic++ -c, exit 0); server build/run pending - CP / Scripts/Blutrain/bluscriptCP.cpp, graph_bluscript_cp.ipynb
2026-07-10 - Cosine MAX_LR sweep complete (9/9) + compare_sched_fmha.py; best val 1.5e-3, cosine 0.03 worse than multistep in basin - BluTrain / compare_maxlr_cosine_*.png, maxlr_cosine_sweep_logs/
2026-07-10 - CP fused-RoPE Phase 3: implemented 4-delta offset-parameterized fused GQA kernels (fwd+bwd) in GQA_fused_{fwd,bwd}sm103_cp.cu -> cp::cuda::gqa_fused_rope_cp{forward,backward}. Scalar pos_offset replaced by per-side head/tail deltas (pos=local+(local>=len/2?d1:d0), matches RopeDeltas.h) for HeadTail-LB non-contiguous [head|tail] shards; separate T_q/T_k + Q/K/V ptrs; bwd recomputes rstd->smem (drops saved q/k_rstd); causal UNCHANGED (local-monotonic sub-chunks; deltas index RoPE only). Whole TU gated #if CP_FUSED_ROPE (default=empty TU, no libtensor collision); FusedRoPESDPA.h decls+wrappers reconciled to bf16 (as_type casts); Makefile CP_FUSED_ROPE/CP_ROPE_DEBUG knobs; monotonicity-contract comment at ContextParallel.h call site. 2x3060 sm86 VERIFIED: default build unchanged (standin 12/12 PASS); CP_FUSED_ROPE=1 new Tests/cp_rope_fused_parity - fwd+dQ/dK/dV BIT-IDENTICAL to production kernel at deltas=0 (cos=1.0 maxdiff=0), dq/dk_gamma cos=1.0 maxdiff3e-7, HeadTail seam case (deltas 0,32) cos=0.99999 + seam rows exact, OOR counter zero. Makefile hardened so NO make clean is ever needed: (1) flag-stamp $(OBJDIR)/.cp_rope_flags (rewritten only on change, prereq of all obj rules) makes CP_FUSED_ROPE/CP_ROPE_DEBUG a real build dep; (2) added -MMD -MP + -include .d for automatic header-dependency tracking (was absent). Verified: same-flags=0 rebuilds, flag-flip auto-rebuilds 13 TUs, touching FusedRoPESDPA.h selectively rebuilds only its 1 includer. GOTCHA: pass the flags to the RUN target too (make CP_FUSED_ROPE=1 run-cp-rope-fused), else run-only re-evaluates flag-off and rebuilds to the empty default -> throws - CP / context_parallel/GQA_fused_fwd_sm103_cp.cu, GQA_fused_bwd_sm103_cp.cu, FusedRoPESDPA.h, ContextParallel.h, Makefile, Tests/cp_rope_fused_parity.cpp
2026-07-10 - Fixed AdamW::load_state segfault on resume-from-step-0 checkpoint (m_/v_ empty vectors + initialized_=true left step() indexing OOB); count==0 now stays uninitialized so lazy-init runs normally. Same latent bug also present in gpt2_fmha_ddp.cpp (shared library) - CP / BluTrain/Tensor-Implementations/src/nn/optimizer/Optim.cpp
2026-07-10 - Process-group UNIFICATION (prereq for DataParallel+ContextParallel coexisting in bluscriptCP): made BluTrain dist/communication ProcessGroupNCCL the single canonical PG for CP+DP. (A) Extended BluTrain PG (additive) with the CP ring-overlap surface ported to run on its own comm_: cpRingStream() + cp_ring_stream_, sendrecv_async_stream, all_gather_async_stream, alltoallv_async, alltoallv_async_stream; + comm-adopting ctor (world_size,rank,ncclComm_t,...) + get_comm(). (B) Rewrote CP process_group/device_mesh.cpp to build sub-groups via ncclCommSplit(world_comm,color=other-coords,key=axis-coord) instead of MPI_Comm_split+WORLD id-bcast -> NCCL-native sub-groups, fixes the 2D-mesh wrong-id bug (world PG still 1x MPI_Bcast); device_mesh.h now includes canonical PG + holds world_pg_. (C) Repointed 8 CP files' #include process_group/ProcessGroupNCCL.h -> "ProcessGroupNCCL.h" (canonical). Collision: device_mesh.h's quote-include resolved CP's same-dir header first -> RENAMED (not deleted, per user) process_group/ProcessGroupNCCL.h -> ProcessGroupNCCL_cplegacy.h (dormant). (D) Makefile: +INCLUDES -IBluTrain/dist/communication/include -IBluTrain/dist/Data-Parallel/include; +BLUTRAIN_PG_SRC (processGroupNccl.cpp + Data-Parallel/src/Error_logs.cpp); CPP_EXCLUDE += process_group/processGroupNCCL.cpp. 2x3060 sm86 VERIFIED: builds clean (one ProcessGroupNCCL/Work/NCCLCHECK, no dup symbols); cp-rope-standin 12/12 PASS cos=1.0000000 (ring fwd+bwd+dgamma + ported overlap methods + ncclCommSplit 1D world identical to old CP PG); cp-ulysses 72 PASS, only the pre-existing bf16-ref-precision dQ/dq_gamma FAILs (unchanged by migration). Follow-ups: 2-D ncclCommSplit per-axis all-reduce test (needs 4 ranks), then bluscriptCP resume (wire DataParallel.cpp) - CP / BluTrain/dist/communication/{include/ProcessGroupNCCL.h,src/processGroupNccl.cpp}, process_group/device_mesh.{h,cpp}, context_parallel/.h, Makefile2026-07-10 - bluscriptCP RESUME (CP Llama training on the unified PG): DataParallel + ContextParallel now coexist. Included DataParallel.hpp DIRECTLY (not the dist/distributed.h umbrella, which pulls the TP stack incl. a different class DeviceMesh -> collision); init_process_group comes from the canonical PG. Fixed double-init crash ("process group instance already created"): the DeviceMesh builds the world PG once (ncclCommSplit parent) and bluscriptCP now reuses it via new DeviceMesh::world_pg() for DataParallel instead of a 2nd init_process_group. Makefile: bluscript-cp/run-bluscript-cp targets; DataParallel.cpp + Data-Parallel/src/profiler.cpp scoped to this target only. 2x3060 sm86 VERIFIED (CP_FUSED_ROPE=1, tiny cfg L2/d256/hd64/GQA kv2/T256, edufineweb data in /Downloads via CP_DATA_ROOT): builds clean; world=1 trains loss 10.89->10.23 decreasing, finite grad-norm, fused-RoPE CP ring ran (first real training integration of the new kernels); ws=2 CP_SIZE=2 ring trains AND loss-parity vs world=1 (step0 10.884 vs 10.890, step4 10.2269 vs 10.2280) -> CP sequence-sharding + RoPE positions + ring correct end-to-end, and overlap-ON matching the serial reference proves the ring overlap is race-free (plan concern d). FOLLOW-UP: CP_NO_OVERLAP=1 diagnostic path crashes ("convert_type_cuda: invalid resource handle" — stream-handle issue under the unified PG; default overlap path unaffected + verified). Also pending: 2-D DPxCP run (dp>1), Ulysses-mode bluscriptCP (CP_ATTN_MODE=ulysses), full-scale config - CP / Scripts/BluTrain/bluscriptCP.cpp, process_group/device_mesh.{h,cpp}, Makefile
2026-07-10 - bluscriptCP follow-ups #1/#3/#4 + ROOT-CAUSED a real sub-group-PG device bug. #1 CP_NO_OVERLAP=1 crash ("convert_type/NCCL enqueue: invalid resource handle") AND #3 Ulysses crash ("Sync failed line 484 / Cuda 400 invalid resource handle") had the SAME root cause introduced by the ncclCommSplit migration: BluTrain PG derived its CUDA device from rank_%gpus_per_node, and DeviceMesh passed the SUB-GROUP rank (my_group_rank); a size-1 DP axis has my_group_rank=0 -> its adopt-ctor set device 0, so the NEXT axis (CP) created its comm_stream while device=0 while its ncclCommSplit comm_ lived on device 1 -> stream/comm device mismatch (blocking collectives on communication_stream_ failed; overlap ring escaped it because cpRingStream is created lazily at ring-time on the correct device). FIX: adopt-ctor now cudaGetDevice(&local_rank_) (adopt caller's global physical device, don't derive from sub-group rank); DeviceMesh re-asserts cudaSetDevice(global_device) at the start of each axis iteration before ncclCommSplit/stream create. 2x3060 VERIFIED: cp-rope-standin 12/12 still PASS (1D regression); ws=2 ring no-overlap now trains + A/B matches overlap (step0 identical 10.884422, step4 diff 2e-5 => overlap race-free, plan concern d); #3 Ulysses ws=2 trains loss 10.890->10.228 with step0 parity vs world=1 (10.889622 vs 10.889627); #4 full-scale d768/L12/q12-kv4/T4096 114M params, CP_SIZE=2 (T_local=2048/rank), trains loss 11.02->10.21 no OOM/NaN on 2x3060. (#2 2D DPxCP deferred: needs 4 ranks.) - CP / BluTrain/dist/communication/src/processGroupNccl.cpp, process_group/device_mesh.cpp
2026-07-10 - Documented bluscriptCP in README.md: new "Context-parallel Llama training (bluscriptCP)" section (Ring vs Ulysses build+run commands, CP_FUSED_ROPE=1-to-both-build-and-run + checkpoint-resume gotchas, full env-var reference incl. topology CP_SIZE/CP_ATTN_MODE/CP_ROTATOR, model/data CP_N_EMBD/LAYER/HEAD/KVHEAD/T/B/GLOBAL_BATCH/MAX_STEPS/WARMUP/DATA_ROOT, YaRN YARN_*, diagnostics CP_NO_OVERLAP*/RING_BLOCKING/SYNC_RING/NO_GPUS_PER_NODE, build flags CP_FUSED_ROPE/CP_ROPE_DEBUG, full-scale example), + bluscript-cp/run-bluscript-cp make targets, intro + repo-layout mentions - CP / README.md
2026-07-10 - Aligned bluscriptCP logging/terminal prints with gpt2_cp_test.cpp: added banner, ULYSSES attn-mode note, gpt2-style params block (Parameters/per-GPU/max_steps/warmup), per-step line (step|loss|lr|norm|dt|tok/sec|mem|Time Left) + [TIMING] breakdown via ported CudaTimer (event-based data/fwd/loss/bwd/clip/optim), validation-loss print, and the RUN-NUMBERED logging system (CP_Training_logs/CP_Training_log{N}.csv + _config.txt with GPU mem, run_number tied to ckpt prefix blumodelcp_run{N}, MPI_Bcast, CP_CKPT/CP_CKPT_FREQ/CP_CKPT_NEW_RUN/CP_CKPT_RESUME env, auto-resume-latest/complete->new-run scan of checkpoints_bluscriptcp). Kept the boxed Configuration block as-is per user. 2x3060 sm86 VERIFIED (world=1 tiny cfg): builds clean, all 4 log paths correct (fresh run1, complete->new run2, resume run2 from step3, CP_CKPT_NEW_RUN->run3), CSV/config.txt written, checkpoints run-numbered - CP / Scripts/BluTrain/bluscriptCP.cpp
2026-07-11 - bluscriptCP: added 2-stage re-warmup for YaRN context-extension resume + checkpoint-visibility prints. get_lr gained a warmup_start anchor (ramp [warmup_start, warmup_start+warmup) then cosine-decay; warmup_start=0 reproduces old schedule exactly) since the base schedule's warmup is absolute-from-0 and cannot re-ramp on resume. New env: CP_REWARMUP= (0=off, only when start_step>0) re-anchors a fresh warmup at the resume step, CP_REWARMUP_PEAK=<frac of max_lr, default 1.0> for a reduced long-context peak; prints "[re-warmup] anchored at step S ...". Also added checkpointing/ckpt_freq/dir to the boxed Configuration block and a "Checkpointing: ON ->