Skip to content

feat(fast): add an fp32 TileLang CPU lowering for the kernels - #896

Merged
NandhaKishorM merged 2 commits into
NandhaKishorM:mainfrom
cklxx:feat/tilelang-cpu
Oct 3, 2026
Merged

NandhaKishorM merged 2 commits into
NandhaKishorM:mainfrom
cklxx:feat/tilelang-cpu

Conversation

@cklxx

@cklxx cklxx commented Oct 3, 2026

Copy link
Copy Markdown
Contributor

laya/tl_kernels.py has always been GPU-only: the kernels use fragments, fragment reductions and bf16, and none of that lowers for TileLang's CPU ("c") target. docs/fast-backends.md recorded that as prose. This PR makes the CPU target work for all five kernels and pins, with tests, what still does not work.

What changed

  • The CPU specialization replaces fragment and shared allocations with local buffers, uses TileLang's scalar path where vectorized loading is unavailable, and promotes 16-bit activations to fp32 for the CPU forward. test_fast_cpu.py covers each kernel against an fp32 torch reference.
  • The additive, keyword-only cpu= selector on the kernel factories is the only API surface added. The GPU dtype default is unchanged.

What is still not supported

bfloat16/float16 CPU execution. tests/test_fast_cpu.py pins the diagnostics for both, including the RoPE fp16 compiler error, so a TileLang upgrade that changes them fails loudly instead of silently changing the claim.

Checks

pytest tests/test_fast_cpu.py -q                                      -> 58 passed
pytest tests/test_fast.py -q                                          -> 13 passed before and after
ruff check laya/ --select=E9,F63,F7,F82,F401,F811 --line-length=120   -> All checks passed!
python -m compileall -q laya/ tests/                                  -> exit 0
python tests/test_router.py / test_criteria.py / test_hooks.py
  / test_hooks_api.py / test_packaging.py                             -> 703 / 198 / 240 / 415 / 131 passed
zensical build --strict --clean                                       -> No issues found, no griffe: lines

Numerics: CPU vs the fp32 reference at most 1.073e-6; CPU vs GPU below 0.05 (0.02 for attention, on the documented fixtures). The GPU path is untouched: 24 captured output tensors bit-identical, 10/10 generated CUDA sources byte-identical, maximum before/after probability delta 0.

This is a portability change, not a speed claim: no CPU speedup and no CPU Agent backend is implied.

@NandhaKishorM
NandhaKishorM merged commit ddf3059 into NandhaKishorM:main Oct 3, 2026
25 checks passed
@NandhaKishorM

Copy link
Copy Markdown
Owner

Merged for 0.3.25.

The evidence page is the reason this went in. Replacing a prose claim ("GPU-only") with the exact diagnostics for each of the five kernels, in both dtypes, and then isolating fragments from bf16 by re-probing with cpu=True so you can say which is the first blocker and which is independent, is a better artifact than the code change. The semantic-check error you quote for C_l being indexed by a T.Parallel variable is the kind of detail nobody reconstructs later.

Two things I want to be straight about, neither a blocker.

Nothing in CI runs this. tests/test_fast_cpu.py is in the packaging test's exemptions (correctly, alongside test_fast.py and test_onnx.py), and no workflow names it, so the 58 checks you report run on your machine and nowhere else. TileLang is not installed here either, so I confirmed only that the suite skips cleanly under pytest rather than erroring. That is the existing precedent for this path, not a new gap you introduced, but it does mean your measurements are the only coverage this code has.

#901 landed in the same pass and changes a default you now also read. It takes gemm_geglu_kernel's stages from 3 to 2, measured on a GPU. Your CPU specialization still goes through T.Pipelined(..., num_stages=stages), so the CPU lowering now runs at two stages, and neither of you measured that combination. I kept both (stages=2 with your cpu=False keyword) because each is sound on its own, and I have said the same on #901. If two stages is wrong for the serial CPU path, that is worth a follow-up with the fp32 numbers.

The GPU-preservation evidence is what let me take it without being able to run it: 24 bit-identical output tensors, 10/10 byte-identical generated CUDA sources, and parity_fast.py comparing exactly equal across 288 questions per dtype.

Verified here: the shared suite list and pytest lane green, test_hooks_api and test_packaging pass with your additions, zero decision flips against 0.3.24 on the real checkpoints.

Thank you.

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

2 participants