[icavan]

Summary

• remove the unused raw-gate SM100 intra entry points and keep the production FP32-gk exact path only
• reduce the exact intra kernel from 31 to 15 warps and remove the obsolete K1 shared-memory/barrier protocol
• initialize manual persistent-scheduler mbarriers from exactly one elected CTA thread
• replace Lightning persistent scheduling across divergent warp-role PCs with a 256-thread completion mbarrier
• add AST regression checks for single-threaded mbarrier initialization, the exact-only intra API, and the Lightning scheduler barrier rule

The removed raw-gate APIs were experimental internal entry points and were not used by the public dispatch.

Validation

• GB200 GPU 1: exact intra/full-port tests: 12 passed
• GB200 GPU 1: delta-h varlen plus Lightning persistent varlen tests: 5 passed
• compute-sanitizer racecheck: exact intra, delta-h persistent scheduler, and Lightning persistent scheduler: 0 hazards
• compute-sanitizer synccheck: exact intra, delta-h persistent scheduler, and Lightning persistent scheduler: 0 errors
• static synchronization/API regression tests: 3 passed
• pre-commit run --all-files: passed

Synchronization audit

• exact intra double buffers use matching TMA, input-ready, MMA-done, stage-reuse, and store-done arrival counts and 0/1 phases
• delta-h publish/consume barriers retain 32 producer and 224 consumer arrivals with double-buffered phase tracking
• Lightning work-unit completion receives one arrival from every CTA thread before the load warp publishes the next work index

Summary by CodeRabbit

• Updates
  • The SM100 KDA intra-chunk path now uses precomputed gates and supports fixed-length and packed variable-length inputs.
• Bug Fixes
  • Improved synchronization in persistent variable-length scheduling for KDA and Lightning prefill.
• Compatibility
  • The experimental equal-length and variable-length KDA intra entry points are no longer available; use the precomputed-gate path instead.

[coderabbitai]

No actionable comments were generated in the recent review.

Recent review info

Run configuration

• Source mode: Auto
• Repository: inclusionAI/cuLA
• Branch: icavan/fix-cutedsl-sync-races
• File: .coderabbit.yaml
• Commit: 043ebc0
• Layers: Repository YAML
• Review profile: CHILL
• Plan: Advanced
• Run ID: cfd2fb10-e5c1-4503-8c07-b134be60ce8e

Commits

Reviewing files that changed from the base of the PR and between 5a9caf3 and 043ebc0.

Files selected for processing: 9

• benchmarks/bench_kda_sm100_intra_fused.py
• benchmarks/stress_kda_sm100_varlen_determinism.py
• cula/ops/kda/sm100/delta_h.py
• cula/ops/kda/sm100/full_port.py
• cula/ops/kda/sm100/intra_fused.py
• cula/ops/lightning/prefill_sm100.py
• docs/kda_sm100_cutedsl_fwd_status.md
• tests/test_cutedsl_mbarrier_init.py
• tests/test_kda_sm100_intra_fused_cutedsl.py

Files with no reviewable changes: 3

• benchmarks/bench_kda_sm100_intra_fused.py
• benchmarks/stress_kda_sm100_varlen_determinism.py
• cula/ops/kda/sm100/full_port.py

Included review availability: This review used your included allowance. Your plan provides up to 4 included reviews per hour; 3 remain after this review.

Walkthrough

The SM100 KDA intra path now consumes precomputed FP32 cumulative gates and uses a revised kernel and host interface. Persistent varlen barrier handling changes in two kernels. Related benchmark scripts and legacy entry points are removed, and tests and documentation are updated.

Changes: SM100 KDA intra path

Layer / File(s)	Summary
Intra kernel and synchronization cula/ops/kda/sm100/intra_fused.py, cula/ops/kda/sm100/delta_h.py	The intra kernel now uses 10 MMA warps, a dedicated TMA producer, and four store/inverse warps. It loads FP32 cumulative gates and beta, applies varlen bounds, and initializes persistent scheduling barriers with one elected lane.
Exact-gate host and port wiring cula/ops/kda/sm100/intra_fused.py, cula/ops/kda/sm100/full_port.py	The exact-gate path passes beta and Akk outputs directly and removes dummy K1 and gate-fusion arguments. The kernel and wrapper cache is renamed _INTRA_CACHE.
Legacy path removal and validation benchmarks/bench_kda_sm100_intra_fused.py, benchmarks/stress_kda_sm100_varlen_determinism.py, docs/kda_sm100_cutedsl_fwd_status.md, tests/test_kda_sm100_intra_fused_cutedsl.py, tests/test_cutedsl_mbarrier_init.py	The raw-gate wrappers and benchmark scripts are removed. Documentation and tests no longer cover the removed entry points or packed-varlen comparison.

Lightning persistent varlen scheduler

Layer / File(s)	Summary
Completion barrier and checks cula/ops/lightning/prefill_sm100.py, tests/test_cutedsl_mbarrier_init.py	Persistent varlen work-unit boundaries use sched_done_mbar arrivals and a phase-based wait instead of the CTA-wide named barrier. Tests check single-threaded barrier initialization and the absence of the named barrier call.

Priority: Low

Estimated code review effort: 4 (Complex) | ~45 minutes

Change: Bug fix

Sequence Diagram(s)

sequenceDiagram
  participant TMALoadWarp
  participant InputReadyBarrier
  participant MMAWarps
  participant StoreWarps
  TMALoadWarp->>InputReadyBarrier: Signal staged beta and input readiness
  InputReadyBarrier->>MMAWarps: Release after TMA data and beta are ready
  MMAWarps->>StoreWarps: Provide intra accumulator results
  StoreWarps->>StoreWarps: Apply masks and write output rows

Loading

sequenceDiagram
  participant LoadWarp
  participant SchedDoneMbar
  participant MMAWarp
  participant CUDACoreWarps
  participant EpilogueWarp
  LoadWarp->>SchedDoneMbar: Wait for work-unit completion
  MMAWarp->>SchedDoneMbar: Arrive at work-unit boundary
  CUDACoreWarps->>SchedDoneMbar: Arrive at work-unit boundary
  EpilogueWarp->>SchedDoneMbar: Arrive at work-unit boundary
  SchedDoneMbar->>LoadWarp: Release Load warp to publish next work index

Loading

Possibly related PRs

• inclusionAI/cuLA#131: Introduced the SM100 CuTeDSL intra path and raw-gate entry points that this change revises.

Suggested reviewers: cherhh, kevinzeng08

Merge Risk: Minimal · up to 043eb

This change removes unused raw-gate intra entry points and fixes synchronization races in the SM100 kernels. The review found no concrete defect, and the author reports clean sanitizer and test runs, so it looks ready to merge.

Pre-merge checks: 4

Check name	Status	Explanation
Title check		The title "fix(cutedsl): close synchronization races" is a short, specific conventional-commit style title. The PR does center on synchronization fixes: single-threaded mbarrier initialization, a Ligh…
Description check		The description follows the template's main sections: Summary, Validation, and a synchronization audit. It gives concrete validation results: GPU test counts, compute-sanitizer racecheck and synccheck…
Linked Issues check		Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check		Check skipped because no linked issues were found for this pull request.

Full details: Title check

Explanation

The title "fix(cutedsl): close synchronization races" is a short, specific conventional-commit style title. The PR does center on synchronization fixes: single-threaded mbarrier initialization, a Lightning completion mbarrier replacing a CTA-wide barrier, and matching arrival counts. It does not mention the main removal of the raw-gate SM100 intra entry points, which is the largest change by line count. The title is therefore partially related to the changeset: it names a real aspect but not the largest part.

Full details: Description check

Explanation

The description follows the template's main sections: Summary, Validation, and a synchronization audit. It gives concrete validation results: GPU test counts, compute-sanitizer racecheck and synccheck results, and pre-commit status. It does not fill the template's 'Related Issues' section, and the pre-commit and test checklist boxes are not ticked. The 'Performance' section is optional and is omitted, which is acceptable. The description is mostly complete. The missing related-issue link is a non-critical gap.

Finishing Touches

📝 Generate docstrings

• Commit to this branch
• Create a new PR

🧪 Generate unit tests (beta)

• Commit to this branch
• Create a new PR

• Autofix · Keep fixing CodeRabbit findings and required CI, and resolving merge conflicts

Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out.

❤️ Share

• X
• Mastodon
• Reddit
• LinkedIn

_{Comment @coderabbitai help to get the list of available commands.}

[cherhh]

LGTM

[KevinZeng08]

LGTM