feat: add build graph schemas and CI diagnostic tools with updated FA… #53
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
| name: build-wheels | |
| # Build all custom CUDA wheels (TE 2.16.0.dev0, flash-attn, mamba_ssm, | |
| # causal_conv1d, fast_hadamard, tilelang, qoptim_cuda) on the repository's | |
| # persistent self-hosted Linux x86_64 runner. | |
| # | |
| # Each matrix entry caches its own ~/.ccache between runs, keyed on source | |
| # git commit. Wheels are uploaded as a single GH Release `wheels-${{ inputs.tag }}`. | |
| # | |
| # Trigger manually after bumping STACK.lock; or on push to STACK.lock / | |
| # upstream_prs/*.patch. | |
| on: | |
| workflow_dispatch: | |
| inputs: | |
| tag: | |
| description: "Release tag suffix (default: short SHA of triggering commit)" | |
| required: false | |
| type: string | |
| push: | |
| branches: [main] | |
| paths: | |
| - STACK.lock | |
| - upstream_prs/*.patch | |
| - .github/workflows/build-wheels.yml | |
| permissions: | |
| contents: write # needed to create the Release that holds the wheel assets | |
| concurrency: | |
| # Cancel any in-flight build-wheels run when a new one is queued — two | |
| # parallel matrix runs (7 jobs each) compete for the runner pool and | |
| # GitHub starts SIGTERM-ing long-running jobs (flash_attn) at exit 143. | |
| group: build-wheels-${{ github.ref }} | |
| cancel-in-progress: true | |
| env: | |
| CUDA_VERSION: "13.2.1" | |
| PYTHON_VERSION: "3.13" | |
| # Keep exact: C++ extension wheels, especially TransformerEngine, are ABI | |
| # tied to the precise stable CUDA torch wheel they were compiled against. | |
| TORCH_VERSION: "2.13.0+cu132" | |
| TORCH_INDEX: "https://download.pytorch.org/whl/cu132" | |
| TORCH_CUDA_ARCH_LIST: "9.0;10.0;12.1" | |
| jobs: | |
| resolve: | |
| if: github.repository == 'DatasunriseOU/cppmega' | |
| runs-on: [self-hosted, Linux, X64, cppmega] | |
| outputs: | |
| tag: ${{ steps.t.outputs.tag }} | |
| steps: | |
| - id: t | |
| run: | | |
| if [[ -n "${{ inputs.tag }}" ]]; then | |
| echo "tag=${{ inputs.tag }}" >> "$GITHUB_OUTPUT" | |
| else | |
| echo "tag=$(echo ${{ github.sha }} | cut -c1-7)" >> "$GITHUB_OUTPUT" | |
| fi | |
| build: | |
| if: github.repository == 'DatasunriseOU/cppmega' | |
| needs: [resolve] | |
| runs-on: [self-hosted, Linux, X64, cppmega] | |
| timeout-minutes: 350 # 6h hard limit; flash_attn at MAX_JOBS=1 ~= 4h | |
| strategy: | |
| fail-fast: false | |
| matrix: | |
| include: | |
| - name: transformer_engine | |
| repo: NVIDIA/TransformerEngine | |
| ref: v2.16 | |
| submodules: "true" | |
| # MAX_JOBS=1 NVCC_THREADS=2 — without these, ninja saturates the | |
| # 4-core / 7 GB runner and the job gets cancelled at ~1h30m with | |
| # nothing useful in the logs. TE has 100s of .cu files * 3 archs. | |
| extra_env: "NVTE_FRAMEWORK=pytorch MAX_JOBS=1 NVCC_THREADS=2" | |
| - name: flash_attn | |
| # Public upstream builds the FA2 cuda kernels + FA4 cute python | |
| # module in one wheel. cppmega-specific build trimming stays below | |
| # in this workflow so the build recipe is owned by DatasunriseOU/cppmega. | |
| # | |
| # MAX_JOBS=1 NVCC_THREADS=1 is the only RAM-stable combo on the | |
| # 7-GB runner. NVCC_THREADS=2 OOM-killed at 58 min (single nvcc | |
| # peaked >6 GB). MAX_JOBS=2 NVCC_THREADS=1 also OOM-killed at | |
| # 13 min (two parallel nvccs went over the cap). Wall time at | |
| # MAX_JOBS=1/THREADS=1 is the constraint; we keep it under | |
| # 350 min by trimming the kernel matrix to 8 files (drop splitkv, | |
| # patch dispatcher in trim step to avoid undefined references). | |
| # | |
| # arch_override is sm_120 only. cppmega architecture priority: | |
| # - H100/H200: covered by flash_attn_3 (FA3 hopper) | |
| # - B200: primary path is FA4 cute python (same wheel, | |
| # JIT-compiled at runtime — no arch list needed) | |
| # FA2 is therefore only the legacy fallback; sm_120 is enough for | |
| # this wheel build. GB10-specific sm_120f/121 fixes must live as | |
| # local patches in this repo if we re-enable that path. | |
| repo: Dao-AILab/flash-attention | |
| ref: main | |
| submodules: "false" | |
| patch: upstream_prs/flash_attn_setup_sm120f.patch | |
| arch_override: "12.1" | |
| extra_env: "MAX_JOBS=1 NVCC_THREADS=1" | |
| - name: flash_attn_3 | |
| # Hopper FA3 path from upstream flash-attention (subdir hopper/). Trimmed to | |
| # training essentials only: bf16+fp8, hdim 64/128, varlen+backward. | |
| # Drop fp16, hdim 96/192/256, paged/append KV (inference), local, | |
| # softcap, packgqa, cluster, hdimdiff64/192. Even at sm_90 + MAX_JOBS=1 | |
| # the full kernel matrix runs ~6h+ on free runners — cancelled at | |
| # 5h51m on previous runs. The trimmed set fits in <1h. | |
| # | |
| # DISABLE_SM80=TRUE is critical: even though we set arch_override=9.0, | |
| # leaving SM8x enabled pulls sm_80 .cu sources AND causes setup.py to | |
| # instantiate \`run_mha_fwd_constexpr<Arch=80, ..., PackGQA=true, ...>\`. | |
| # That template's e4m3 dispatch branch hardcodes \`<90, e4m3, ...>\` in | |
| # the run_mha_fwd_<> call — so the compiler emits an undefined | |
| # reference to \`run_mha_fwd_<90, e4m3, 64, 64, ..., PackGQA=true>\` | |
| # even though we built with DISABLE_PACKGQA. Result: \`_C.abi3.so\` | |
| # ImportError on \`undefined symbol\` at smoke test time. Locking | |
| # ARCH_SWITCH to 90 (via DISABLE_SM8x → defined when SM80=TRUE) | |
| # prevents the sm80 instantiation entirely, and PackGQA collapses to | |
| # false everywhere for arch=90. | |
| repo: Dao-AILab/flash-attention | |
| ref: main | |
| submodules: "false" | |
| subdir: hopper | |
| arch_override: "9.0" | |
| extra_env: >- | |
| MAX_JOBS=1 NVCC_THREADS=2 FLASH_ATTENTION_FORCE_BUILD=TRUE | |
| FLASH_ATTENTION_DISABLE_SM80=TRUE | |
| FLASH_ATTENTION_DISABLE_FP16=TRUE | |
| FLASH_ATTENTION_DISABLE_HDIM96=TRUE | |
| FLASH_ATTENTION_DISABLE_HDIM192=TRUE | |
| FLASH_ATTENTION_DISABLE_HDIM256=TRUE | |
| FLASH_ATTENTION_DISABLE_PAGEDKV=TRUE | |
| FLASH_ATTENTION_DISABLE_APPENDKV=TRUE | |
| FLASH_ATTENTION_DISABLE_LOCAL=TRUE | |
| FLASH_ATTENTION_DISABLE_SOFTCAP=TRUE | |
| FLASH_ATTENTION_DISABLE_PACKGQA=TRUE | |
| FLASH_ATTENTION_DISABLE_CLUSTER=TRUE | |
| FLASH_ATTENTION_DISABLE_HDIMDIFF64=TRUE | |
| FLASH_ATTENTION_DISABLE_HDIMDIFF192=TRUE | |
| FLASH_ATTENTION_DISABLE_SPLIT=TRUE | |
| - name: mamba_ssm | |
| repo: state-spaces/mamba | |
| ref: 31f3d7b | |
| submodules: "false" | |
| extra_env: "MAX_JOBS=1 NVCC_THREADS=2" | |
| - name: causal_conv1d | |
| repo: Dao-AILab/causal-conv1d | |
| ref: v1.6.1 | |
| submodules: "false" | |
| extra_env: "MAX_JOBS=1 NVCC_THREADS=2" | |
| - name: fast_hadamard_transform | |
| repo: Dao-AILab/fast-hadamard-transform | |
| ref: master | |
| submodules: "false" | |
| extra_env: "MAX_JOBS=1 NVCC_THREADS=2" | |
| - name: tilelang | |
| # DatasunriseOU/tilelang fork (not upstream tile-ai/tilelang) because: | |
| # - 5952468a carries upstream PR #2071 (removes apache-tvm-ffi<0.1.10 cap) | |
| # - vendored TVM submodule (DatasunriseOU/tvm@9b0a1667) includes | |
| # apache/tvm#18938 (TVMDerivedObject.__slots__ fix) → no | |
| # '_NestedLoopCheckVisitor' import crash under tvm-ffi >=0.1.12 | |
| # - upstream tile-ai/tilelang HEAD still caps apache-tvm-ffi<0.1.12, | |
| # incompatible with FA4 beta23's >=0.1.12 floor. | |
| # Must match STACK.lock wheels.tilelang.ref. | |
| repo: DatasunriseOU/tilelang | |
| ref: 5952468a | |
| submodules: "true" | |
| extra_env: "" | |
| - name: qoptim_cuda | |
| repo: NVlabs/COAT | |
| ref: main | |
| submodules: "false" | |
| subdir: coat/optimizer/kernels | |
| extra_env: "MAX_JOBS=1 NVCC_THREADS=2" | |
| steps: | |
| - name: Checkout cppmega (for patches) | |
| uses: actions/checkout@v4 | |
| - name: Verify persistent runner capacity | |
| shell: bash | |
| run: | | |
| set -euo pipefail | |
| df -h "$GITHUB_WORKSPACE" | |
| available_kib=$(df -Pk "$GITHUB_WORKSPACE" | awk 'NR == 2 {print $4}') | |
| minimum_kib=$((120 * 1024 * 1024)) | |
| if (( available_kib < minimum_kib )); then | |
| echo "::error::CUDA wheel build requires at least 120 GiB free on the persistent runner" | |
| exit 1 | |
| fi | |
| - name: Resolve source SHA | |
| id: src | |
| run: | | |
| set -euo pipefail | |
| REPO_URL="https://github.com/${{ matrix.repo }}.git" | |
| SHA=$(git ls-remote "$REPO_URL" ${{ matrix.ref }} | awk '{print $1}' | head -1) | |
| if [ -z "$SHA" ]; then SHA="${{ matrix.ref }}"; fi | |
| echo "sha=$SHA" >> "$GITHUB_OUTPUT" | |
| echo "Resolved ${{ matrix.repo }}@${{ matrix.ref }} → $SHA" | |
| - name: Set up Python | |
| uses: actions/setup-python@v5 | |
| with: | |
| python-version: ${{ env.PYTHON_VERSION }} | |
| - name: Install gcc-15 + ccache (vendored debs) | |
| # We pulled gcc-15 .debs from Ubuntu questing (25.10) main archive | |
| # once and uploaded them as a release asset on this repo. This avoids | |
| # the chronically flaky Launchpad PPA (api.launchpad.net + ppa.launchpad | |
| # content.net both regularly time out / return 503 from CI runners, | |
| # blocking every wheel build for hours). | |
| # Source: archive.ubuntu.com/ubuntu/pool/main/g/gcc-15/ (15.2.0-4ubuntu4) | |
| # Release: github.com/${{ github.repository }}/releases/tag/${TOOLCHAIN_TAG} | |
| env: | |
| TOOLCHAIN_TAG: toolchain-gcc-15-15.2.0-4ubuntu4 | |
| GH_TOKEN: ${{ secrets.GITHUB_TOKEN }} | |
| run: | | |
| set -euo pipefail | |
| mkdir -p /tmp/gcc15 | |
| cd /tmp/gcc15 | |
| # Download from our own release (GitHub's CDN, very stable). One blob, | |
| # no PPA, no add-apt-repository, no Launchpad anywhere in the path. | |
| if ! gh release download "$TOOLCHAIN_TAG" \ | |
| --repo "${{ github.repository }}" \ | |
| --pattern 'gcc-15-noble-amd64.tar.gz'; then | |
| echo "::warning::vendored gcc-15 release missing in ${{ github.repository }}; falling back to ubuntu-toolchain PPA" | |
| sudo apt-get -o Acquire::Retries=5 update | |
| sudo apt-get install -y software-properties-common | |
| sudo add-apt-repository -y ppa:ubuntu-toolchain-r/test | |
| sudo apt-get -o Acquire::Retries=5 update | |
| sudo apt-get install -y gcc-15 g++-15 ccache ninja-build | |
| sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-15 90 | |
| sudo update-alternatives --install /usr/bin/g++ g++ /usr/bin/g++-15 90 | |
| gcc --version | |
| g++ --version | |
| exit 0 | |
| fi | |
| tar xzf gcc-15-noble-amd64.tar.gz | |
| sha256sum -c SHA256SUMS | |
| # Install ccache + ninja from noble main (no PPA needed for these). | |
| sudo apt-get -o Acquire::Retries=5 update | |
| sudo apt-get install -y ccache ninja-build | |
| # dpkg -i everything in one shot (dpkg topo-sorts internally). The | |
| # tarball includes both the compiler binaries and all gcc-15 runtime | |
| # libs (libgcc-s1, libstdc++6, libquadmath0, libgomp1, libasan8, | |
| # liblsan0, libtsan2, libubsan1, libitm1, libatomic1, libhwasan0, | |
| # libcc1-0) — these MUST be upgraded together because libgcc-15-dev | |
| # has internal version-locked deps on every one of them. Other deps | |
| # (libgmp10, libisl23, libmpc3, libmpfr6, libzstd1, zlib1g, binutils) | |
| # are already in noble main and noble's versions are new enough. | |
| sudo dpkg -i *.deb | |
| sudo update-alternatives --install /usr/bin/gcc gcc /usr/bin/gcc-15 90 | |
| sudo update-alternatives --install /usr/bin/g++ g++ /usr/bin/g++-15 90 | |
| gcc --version | |
| g++ --version | |
| - name: Install CUDA ${{ env.CUDA_VERSION }} (apt from NVIDIA repo) | |
| run: | | |
| set -euo pipefail | |
| wget -q https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2404/x86_64/cuda-keyring_1.1-1_all.deb | |
| sudo dpkg -i cuda-keyring_1.1-1_all.deb | |
| sudo apt-get update | |
| # cuda-toolkit-13-2 pulls latest 13.2.x (13.2.1 by 2026-04). cudnn9-cuda-13 = cuDNN 9. | |
| # libnccl-dev/libnccl2 needed by TE (nccl.h #include in common/util/logging.h). | |
| sudo apt-get install -y --no-install-recommends \ | |
| cuda-toolkit-13-2 cudnn9-cuda-13 libnccl-dev libnccl2 | |
| echo "CUDA_HOME=/usr/local/cuda-13.2" >> "$GITHUB_ENV" | |
| echo "/usr/local/cuda-13.2/bin" >> "$GITHUB_PATH" | |
| /usr/local/cuda-13.2/bin/nvcc --version | |
| # Keep real driver libraries intact on persistent GPU runners. Add | |
| # toolkit stubs only when the host has no driver library at all. | |
| if ! ldconfig -p | grep -q 'libcuda\.so\.1'; then | |
| sudo ln -sf /usr/local/cuda-13.2/lib64/stubs/libcuda.so /usr/local/cuda-13.2/lib64/libcuda.so | |
| sudo ln -sf /usr/local/cuda-13.2/lib64/stubs/libcuda.so /usr/local/cuda-13.2/lib64/libcuda.so.1 | |
| fi | |
| if ! ldconfig -p | grep -q 'libnvidia-ml\.so\.1'; then | |
| sudo ln -sf /usr/local/cuda-13.2/lib64/stubs/libnvidia-ml.so /usr/local/cuda-13.2/lib64/libnvidia-ml.so | |
| sudo ln -sf /usr/local/cuda-13.2/lib64/stubs/libnvidia-ml.so /usr/local/cuda-13.2/lib64/libnvidia-ml.so.1 | |
| fi | |
| - name: Wire ccache as gcc/g++/nvcc via PATH | |
| # CC="ccache gcc-15" breaks setuptools/ninja (it splits the value and | |
| # treats "gcc-15" as a positional arg). Use PATH-based wrappers instead: | |
| # /usr/lib/ccache/<tool> -> ccache; ccache then exec's the real binary. | |
| run: | | |
| sudo mkdir -p /usr/lib/ccache | |
| for tool in cc gcc g++ c++ gcc-15 g++-15 gcc-14 g++-14 nvcc; do | |
| sudo ln -sf /usr/bin/ccache /usr/lib/ccache/$tool | |
| done | |
| echo "/usr/lib/ccache" >> "$GITHUB_PATH" | |
| - name: Configure nvcc → gcc-15 | |
| run: | | |
| echo "NVCC_PREPEND_FLAGS=--compiler-bindir=/usr/bin/g++-15 --allow-unsupported-compiler" >> "$GITHUB_ENV" | |
| # Per-matrix arch_override wins over the global TORCH_CUDA_ARCH_LIST | |
| # (e.g. flash_attn_3 hopper-only sm_90, no 10.0/12.1). | |
| ARCH="${{ matrix.arch_override || env.TORCH_CUDA_ARCH_LIST }}" | |
| echo "TORCH_CUDA_ARCH_LIST=${ARCH}" >> "$GITHUB_ENV" | |
| echo "LD_LIBRARY_PATH=/usr/local/cuda-13.2/lib64:${LD_LIBRARY_PATH:-}" >> "$GITHUB_ENV" | |
| echo "Using TORCH_CUDA_ARCH_LIST=${ARCH}" | |
| - name: Restore ccache | |
| uses: actions/cache@v4 | |
| with: | |
| path: ~/.ccache | |
| key: ccache-${{ matrix.name }}-${{ steps.src.outputs.sha }}-cuda${{ env.CUDA_VERSION }}-torch${{ env.TORCH_VERSION }} | |
| restore-keys: | | |
| ccache-${{ matrix.name }}- | |
| - name: Configure ccache | |
| run: | | |
| ccache -M 5G | |
| ccache -o compiler_check=content | |
| ccache -z | |
| - name: Install torch ${{ env.TORCH_VERSION }} | |
| # download.pytorch.org sometimes drops the connection mid-stream on | |
| # the big CUDA wheels (IncompleteRead on nvidia_cublas-*.whl, ~400 MB). | |
| # --retries 10 + --timeout 60 makes pip retry the whole package on | |
| # ProtocolError. Wrap in a shell loop too in case pip exits before retrying. | |
| run: | | |
| set -euo pipefail | |
| pip install --upgrade pip setuptools wheel | |
| n=0; until pip install --retries 10 --timeout 60 \ | |
| --extra-index-url ${{ env.TORCH_INDEX }} \ | |
| "torch==${{ env.TORCH_VERSION }}"; do | |
| n=$((n+1)) | |
| [ "$n" -ge 5 ] && { echo "::error::torch install failed after 5 attempts"; exit 1; } | |
| echo "::warning::torch install attempt $n failed; retrying in $((30*n))s" | |
| sleep $((30 * n)) | |
| done | |
| python -c "import torch; print(torch.__version__, torch.version.cuda)" | |
| - name: Install build deps | |
| # scikit-build-core: PEP 517 backend used by tilelang/CMake-based wheels. | |
| # setuptools-scm: many wheels resolve their version from git tags. | |
| # numpy must be present for torch's tensor_numpy.cpp at build time. | |
| # z3-solver: tilelang's CMake calls 'import z3' during configure (FindZ3.cmake). | |
| # cython: tilelang ships a .pyx module (jit/adapter/cython/cython_wrapper.pyx). | |
| run: pip install --retries 10 --timeout 60 ninja packaging pybind11 cmake numpy einops scikit-build-core setuptools-scm "z3-solver>=4.13.0" cython | |
| - name: Clone source (any ref — branch / tag / SHA) | |
| run: | | |
| set -euo pipefail | |
| cd "$GITHUB_WORKSPACE" | |
| rm -rf src | |
| REPO_URL="https://github.com/${{ matrix.repo }}.git" | |
| git clone "$REPO_URL" src | |
| cd src | |
| git checkout ${{ matrix.ref }} | |
| if [[ "${{ matrix.submodules }}" == "true" ]]; then | |
| git submodule update --init --recursive --depth 1 | |
| fi | |
| git log -1 --format='Source HEAD: %h %s' | |
| - name: Apply patch | |
| if: matrix.patch | |
| run: | | |
| set -euo pipefail | |
| cd src | |
| PATCH="$GITHUB_WORKSPACE/${{ matrix.patch }}" | |
| if git apply --check "$PATCH"; then | |
| git apply --verbose "$PATCH" | |
| elif [[ "${{ matrix.name }}" == "flash_attn" ]] && grep -q 'arch=compute_120f,code=sm_120' setup.py; then | |
| echo "${{ matrix.patch }} is already present in upstream ${{ matrix.repo }}@${{ matrix.ref }}" | |
| else | |
| echo "::error::${{ matrix.patch }} does not apply and no known upstream marker was found." | |
| git apply --check "$PATCH" || true | |
| exit 1 | |
| fi | |
| - name: Trim flash_attn FA2 kernel matrix | |
| # FA2 setup.py has no DISABLE_* knobs — kernel list is hardcoded as | |
| # 6 hdims × 2 dtypes (fp16/bf16) × 2 causal × 2 (fwd/bwd) = 48 .cu | |
| # per arch. Even at sm_120 only / MAX_JOBS=1 / NVCC_THREADS=1 this | |
| # blew the 350-min job timeout three runs in a row. | |
| # cppmega's training path uses bf16 + hdim 64 (Mamba3) and hdim 128 | |
| # (TE attention). Drop fp16 entirely + drop hdims 32/96/192/256 → | |
| # 8 .cu files instead of 48, ~6x faster. | |
| # | |
| # We must ALSO trim HEADDIM_SWITCH in static_switch.h: the macro | |
| # expands all 6 hdim if-branches, each referencing run_mha_fwd_<T, | |
| # hdim, ...>. With kernels missing, those become undefined refs in | |
| # the .so → ImportError at runtime. Same gotcha as FA3. | |
| if: matrix.name == 'flash_attn' | |
| run: | | |
| set -euo pipefail | |
| cd src | |
| # 1) Drop unwanted kernel sources from setup.py's hardcoded list. | |
| # FA2 has fwd, bwd, AND fwd_split families. The fwd_split is the | |
| # splitkv path (dispatched at runtime when num_splits > 1, used by | |
| # mha_fwd_kvcache). cppmega's training only needs num_splits==1 | |
| # fwd+bwd; keeping fwd_split makes the build overflow the 350-min | |
| # timeout. Drop it AND patch the call site (step 3) so the | |
| # dispatcher stops referencing run_mha_fwd_splitkv_dispatch<>. | |
| # Survivors: flash_{fwd,bwd}_hdim{64,128}_bf16{,_causal}_sm80.cu = 8. | |
| before=$(grep -c '_sm80\.cu' setup.py || true) | |
| sed -i -E '/csrc\/flash_attn\/src\/flash_(fwd|bwd|fwd_split)_hdim[0-9]+_fp16(_causal)?_sm80\.cu/d' setup.py | |
| sed -i -E '/csrc\/flash_attn\/src\/flash_(fwd|bwd|fwd_split)_hdim(32|96|192|256)_bf16(_causal)?_sm80\.cu/d' setup.py | |
| sed -i -E '/csrc\/flash_attn\/src\/flash_fwd_split_hdim(64|128)_bf16(_causal)?_sm80\.cu/d' setup.py | |
| after=$(grep -c '_sm80\.cu' setup.py || true) | |
| echo "trimmed $((before - after)) kernel sources ($before → $after)" | |
| # 2) Restrict HEADDIM_SWITCH to {64,128} and FP16_SWITCH to bf16 | |
| # only. Append #undef + new definitions at end of static_switch.h | |
| # — any subsequent #include picks up the override. Without these, | |
| # the original macros expand all 6 hdim and both fp16/bf16 | |
| # branches, generating run_mha_fwd_<T, hdim, ...> references with | |
| # no matching template instantiation → ImportError on undefined | |
| # symbol. | |
| cat >> csrc/flash_attn/src/static_switch.h <<'EOF' | |
| // cppmega CI trim: kernel sources only include bf16 + hdim 64/128. | |
| #undef HEADDIM_SWITCH | |
| #define HEADDIM_SWITCH(HEADDIM, ...) \ | |
| [&] { \ | |
| if (HEADDIM <= 64) { \ | |
| constexpr static int kHeadDim = 64; \ | |
| return __VA_ARGS__(); \ | |
| } else if (HEADDIM <= 128) { \ | |
| constexpr static int kHeadDim = 128; \ | |
| return __VA_ARGS__(); \ | |
| } \ | |
| }() | |
| #undef FP16_SWITCH | |
| #define FP16_SWITCH(COND, ...) \ | |
| [&] { \ | |
| using elem_type = cutlass::bfloat16_t; \ | |
| return __VA_ARGS__(); \ | |
| }() | |
| EOF | |
| # 3) Disable the splitkv dispatch path in flash_api.cpp. Without | |
| # this, run_mha_fwd_splitkv_dispatch<T, kHeadDim, IsCausal> is | |
| # referenced from the dispatcher → undefined symbol in .so since | |
| # we dropped the splitkv .cu sources. cppmega doesn't use the | |
| # kvcache path (inference-only), so a runtime check is fine. | |
| sed -i 's|run_mha_fwd_splitkv_dispatch<elem_type, kHeadDim, Is_causal>(params, stream);|TORCH_CHECK(false, "flash_attn splitkv path disabled in cppmega CI build");|' \ | |
| csrc/flash_attn/flash_api.cpp | |
| # 4) Keep FA4 out of this wheel. Upstream excludes flash_attn.cute | |
| # because FA4 ships as the separate flash-attn-4 package with its | |
| # own nvidia-cutlass-dsl/quack dependency contract. Forcing cute into | |
| # flash-attn produced an import surface that looked present but | |
| # failed at runtime against modern CUTLASS Python. | |
| echo "HEADDIM_SWITCH restricted to {64,128}; splitkv stubbed; survivors:" | |
| grep '_sm80\.cu' setup.py | sed 's/^/ /' | |
| - name: Build wheel | |
| run: | | |
| set -euo pipefail | |
| BUILD_DIR="src${{ matrix.subdir && format('/{0}', matrix.subdir) || '' }}" | |
| mkdir -p wheels logs | |
| LOG="$GITHUB_WORKSPACE/logs/build-${{ matrix.name }}.log" | |
| cd "$BUILD_DIR" | |
| # Tee build output to a file. Runner's per-step blob expires fast | |
| # (Azure cleanup hits within ~30 min for some jobs); the artifact | |
| # below keeps it for 30 days regardless. | |
| env ${{ matrix.extra_env }} \ | |
| pip wheel . --no-build-isolation --no-deps -w "$GITHUB_WORKSPACE/wheels" -v 2>&1 | tee "$LOG" | |
| - name: ccache stats | |
| if: always() | |
| run: ccache -s | |
| - name: List wheels | |
| if: success() | |
| run: ls -lh wheels/ | |
| - name: Smoke TransformerEngine wheel ABI | |
| if: success() && matrix.name == 'transformer_engine' | |
| run: | | |
| set -euo pipefail | |
| python -m pip install --no-deps wheels/transformer_engine-*.whl | |
| python - <<'PY' | |
| import torch | |
| import transformer_engine | |
| import transformer_engine.pytorch | |
| print("torch", torch.__version__, torch.version.cuda) | |
| print("transformer_engine", transformer_engine.__version__) | |
| PY | |
| - name: Upload build log (always, even on failure/cancel) | |
| if: always() | |
| uses: actions/upload-artifact@v4 | |
| with: | |
| name: log-${{ matrix.name }} | |
| path: logs/ | |
| retention-days: 30 | |
| if-no-files-found: warn | |
| - name: Upload wheels artifact | |
| if: success() | |
| uses: actions/upload-artifact@v4 | |
| with: | |
| name: wheel-${{ matrix.name }} | |
| path: wheels/*.whl | |
| retention-days: 7 | |
| release: | |
| if: github.repository == 'DatasunriseOU/cppmega' | |
| needs: [resolve, build] | |
| runs-on: [self-hosted, Linux, X64, cppmega] | |
| steps: | |
| - uses: actions/checkout@v4 | |
| - name: Download all wheel artifacts | |
| uses: actions/download-artifact@v4 | |
| with: | |
| pattern: wheel-* | |
| merge-multiple: true | |
| path: wheels | |
| - name: List collected wheels | |
| run: ls -lh wheels/ | |
| - name: Verify required wheels | |
| run: | | |
| set -euo pipefail | |
| required=( | |
| 'transformer_engine-*.whl' | |
| 'flash_attn-*.whl' | |
| 'flash_attn_3-*.whl' | |
| 'mamba_ssm-*.whl' | |
| 'causal_conv1d-*.whl' | |
| 'fast_hadamard_transform-*.whl' | |
| 'tilelang-*.whl' | |
| 'qoptim_cuda-*.whl' | |
| ) | |
| for pattern in "${required[@]}"; do | |
| matches=(wheels/$pattern) | |
| if [[ ! -e "${matches[0]}" ]]; then | |
| echo "::error::missing required wheel asset: $pattern" | |
| ls -lh wheels/ || true | |
| exit 1 | |
| fi | |
| done | |
| - name: Create / update Release | |
| env: | |
| GH_TOKEN: ${{ secrets.GITHUB_TOKEN }} | |
| run: | | |
| TAG="wheels-${{ needs.resolve.outputs.tag }}" | |
| gh release view "$TAG" >/dev/null 2>&1 || \ | |
| gh release create "$TAG" \ | |
| --title "cppmega wheels ${{ needs.resolve.outputs.tag }}" \ | |
| --notes "Auto-built by build-wheels.yml from commit ${{ github.sha }}." \ | |
| --prerelease | |
| gh release upload "$TAG" wheels/*.whl --clobber |