Skip to content

v1.8.1 — CUTLASS v4.6.2 Submodule Sync & Package Version 2.9.2.post2

Latest
Latest

Choose a tag to compare

View all tags
@ussoewwin ussoewwin released this 05 Sep 04:38
· 1 commit to main since this release
v1.8.1
6f4ac8d

Overview

Fork release v1.8.1 synchronizes the bundled C++ CUTLASS submodule (csrc/cutlass) from v4.3.4 to the official tag v4.6.2 (commit 6c65a17), bringing the C++ template headers into strict alignment with FlashAttention-4's nvidia-cutlass-dsl>=4.6.2 requirement while ensuring 100% backward compatibility with FlashAttention-2 CUDA/C++ kernels. Package version is incremented to 2.9.2.post2 across flash_attn/__init__.py and setup2.py. Furthermore, legacy compilation artifacts (build/) have been purged to guarantee a clean build baseline under Windows MSVC.

1. CUTLASS Submodule Synchronization to v4.6.2

Background & Motivation

In recent upstream commits (specifically commit ce088ab, Dao-AILab#2798), the Python DSL dependency floor was raised to nvidia-cutlass-dsl>=4.6.2 to support QuACK packed subtraction compatibility and upcoming CUTLASS 4.7 capabilities. However, the Git submodule in csrc/cutlass had remained at v4.3.4 (commit 7127592). This release updates the submodule to tag v4.6.2 (6c65a175668952f09bcbf66cb97a8de1b734b4a0).

Technical Compatibility with FlashAttention-2

FlashAttention-2 (csrc/flash_attn/) does not instantiate CUTLASS 3.x/4.x Collective Mainloops or CuTe C++ GEMM pipelines; its usage of CUTLASS is strictly scoped to fundamental data types and utility headers:

  • <cutlass/cutlass.h>
  • <cutlass/array.h>
  • <cutlass/numeric_types.h>
  • <cutlass/numeric_conversion.h>

The numeric primitives (cutlass::half_t, cutlass::bfloat16_t, cutlass::Array<T, N>) and conversion routines retain identical semantics and binary layouts in CUTLASS 4.6.2. Under MSVC with /std:c++20, the FA2 compilation units build cleanly without symbol collisions or deprecation regressions.

2. Package Version Bump (2.9.2.post2)

Package metadata and runtime indicators have been updated from 2.9.2.post1 to 2.9.2.post2:

  • flash_attn/__init__.py: Default fallback version updated to "2.9.2.post2".
  • setup2.py: public_version = "2.9.2.post2" in get_package_version().

This ensures that binary wheels produced from this revision are unambiguously distinguishable from earlier post1 builds.

3. Build Tree Hygiene & Clean Build Assurance

  • Build Directory Purge: The legacy build/ tree (temp.win-amd64-*, lib.win-amd64-*) was completely wiped. When building with Ninja/MSVC, cached .obj files can bypass recompilation if source file timestamps appear unchanged, risking subtle ABI or header inconsistencies. Removing build/ guarantees that every object file is freshly compiled against CUTLASS 4.6.2 headers.
  • Git Ignore Hygiene: Added explicit rules in .gitignore for build logs (*.log), temporary directories (temp/, tmp/), and backup files (*.bak*, *.tmp), eliminating unintended staging of build side-effects.

4. Preserved Fork Architecture & Windows Optimizations

All fork-specific Windows enhancements remain active and verified:

  • 24 split_align Forward Kernels: Independent translation units (flash_fwd_split_align_*.cu) for num_splits == 1 forward pass prevent ptxas compiler timeouts.
  • MSVC 2GB COFF Linker Limit Bypass (3:3 SM Split):
    • WindowsWhlBuilder_cuda_2.bat: Blackwell targets (100;120;121).
    • WindowsWhlBuilder_cuda_3.bat: Legacy targets (80;89;90), with MAX_JOBS=5.
  • Windows MSVC Compatibility Overrides:
    • Standard flags: /std:c++20 enabled for modern C++ initializers.
    • C10 ArrayRef operator overrides in namespace c10 (resolves MSVC C2666 ambiguity).
    • Clean linking against PyTorch 2.14.0's 6-argument c10::cuda::c10_cuda_check_implementation.
  • PEP 427 Compliant Wheel Naming: Employs .blackwell and .legacy dot-suffixes.
  • CI Workflow Suppression: .github and tools/ci remain permanently suppressed.

5. Pre-built Windows Wheels

Pre-compiled wheels will be published on Hugging Face:
👉 ussoewwin/Flash-Attention-2_for_Windows

Quick Install Command

pip install https://huggingface.co/ussoewwin/Flash-Attention-2_for_Windows/resolve/main/<matching_wheel_file>.whl

6. Verification and Validation Guide

Detailed accuracy matrix, functional sweep (7 API paths), backward pass gradients, and latency benchmarks on NVIDIA GeForce RTX 5060 Ti (sm_120) are fully documented in:
👉 md/2.9.2.post2_COMPLETE_TEST_AND_VALIDATION_GUIDE.md

Assets 2
Loading