[FEATURE] Reduce compute bubbles in backward skip path

[Copilot]

This PR implements a comprehensive set of optimizations to reduce compute bubbles in the backward kernel's skip branch when handling fully masked blocks. The optimizations target the key inefficiencies identified in sparse attention scenarios where many tiles are inactive.

Problem

The original backward kernel skip branch showed substantial compute bubbles even when successfully skipping mathematically null GEMMs. Key inefficiencies included:

• Unnecessary global loads (K/V, dO) issued before mask activity decisions
• Unconditional barrier synchronization even when no work follows
• Idle periods where SMs cannot schedule useful instructions
• Resource over-reservation limiting occupancy
• Coarse-grained skipping (only whole tiles, not sub-tile optimization)

Solution

Implemented a 4-phase optimization strategy:

Phase 1: Early Mask Prefetch

Moved mask loading and activity checking before heavy K/V/dO async loads, enabling skip decisions before expensive memory operations complete.

// Before: Wait for all loads, then check mask
cute::cp_async_wait<0>();
__syncthreads();
// Copy mask and check activity...

// After: Check mask early, skip loads if inactive
Tensor tSrMask = make_tensor<Element>(shape(acc_s));
cute::copy(smem_tiled_copy_PdS, tSsMask, tSrMask_copy_view);
bool any_active = FLASH_NAMESPACE::check_mask_activity_early(tSrMask);

Phase 2: Conditional Synchronization

Bypass unnecessary __syncthreads() barriers for fully masked tiles when safe for pipeline correctness.

Phase 3: Next-Tile Look-Ahead

Added infrastructure for prefetching subsequent mask/bias tiles during skip cycles to hide future operation latency.

Phase 4: Adaptive Density Mode

Dynamically disable skip optimization when active tile density exceeds 85%, preventing overhead in high-density scenarios.

// Track density and adapt
float current_density = float(active_tiles) / float(total_tiles);
use_skip_optimization = (current_density <= DENSITY_THRESHOLD);

Performance Impact

Expected speedups based on sparsity level:

• High sparsity (≥70%): 2-4x speedup for block-sparse patterns
• Medium sparsity (40-70%): 1.2-2x speedup for mixed patterns
• Low sparsity (<15%): No overhead due to adaptive fallback

Key Features

• Automatic activation: No API changes required - optimizations apply transparently
• Numerical equivalence: All optimizations preserve mathematical correctness
• Architecture compatibility: Works with existing SM 8.0+ requirements
• Comprehensive testing: Full test suite and performance benchmarks included

Testing

Added comprehensive validation:

• Unit tests for mask activity logic and adaptive thresholds
• Integration tests for various sparse patterns (block, random, structured)
• Performance benchmarks demonstrating expected characteristics
• Numerical equivalence verification

The optimizations maintain backward compatibility while delivering significant performance improvements for sparse attention workloads common in long-context models, mixture-of-experts architectures, and structured document attention.

Fixes #132.

---

✨ Let Copilot coding agent set things up for you — coding agent works faster and does higher quality work when set up for your repo.