fix(comm): run collectives on their operand's device - #3
Closed
lvyufeng wants to merge 20 commits into
Closed
Conversation
…flagos-ai#34) * fix(dist): add _allgather_base/_reduce_scatter_base to ProcessGroupFlagOS The functional all_gather_into_tensor / reduce_scatter_tensor APIs dispatch to group._allgather_base / group._reduce_scatter_base (single flat tensor in and out), which ProcessGroupFlagOS did not override -> 'No backend type associated with device type flagos'. Wrap both, converting the flagos tensors to the inner backend's device view like the other collectives. Add tests/manual/metax/test_flagos_dist_live_metax.py: the MetaX counterpart of the nvidia live test. Live-verified on 2xMetaX C550 via the NCCL(mccl) fallback path (GEMS_VENDOR=metax, flagcx absent): all_reduce / broadcast / all_gather / reduce_scatter_tensor and DDP grad sync ALL PASS. * fix(metax): route torch.cuda.set_device to maca runtime for FlagCX Under metax boxing, stock torch.cuda.set_device only bumps torch's own device counter (torch._C._cuda_setDevice on the CPU wheel) and never reaches maca's mcSetDevice, so torch.cuda.set_device(rank) leaves the maca runtime on device 0 while torch.cuda.current_device() reports rank. ProcessGroupNCCL(mccl) hides this by wrapping every op in a CUDAGuard(tensor.device()); FlagCX does not -- it creates its collective stream (mcStreamCreateWithFlags) on the *current* maca device before binding the tensor's device, so on rank!=0 the stream lands on device 0 while the communicator/tensor are on device r, giving 'CUDA error: invalid resource handle'. Patch torch.cuda.set_device to also move the maca runtime (via _flagos.set_device) and point current_device at the maca-backed reader, so the FlagCX metax path binds a consistent device. Verified 2xC550: all_reduce/broadcast/all_gather/reduce_scatter + DDP grad sync ALL PASS through the FlagCX (flagcx=yes) path. * fix(metax): drop unused ProcessGroupFlagOS import in dist live test The _describe_inner_backend helper imported ProcessGroupFlagOS but never used it, tripping ruff F401 in CI. Remove the import. * style(metax): ruff format dist live test print lines
* ci: use the PyTorch CUDA runner set * test: expect CUDA fallback for recursive mul --------- Co-authored-by: 马函廷 <mahanting@mahantingdeMacBook-Air.local>
…outing (flagos-ai#36) Two independent fixes. The first is NOT DCU-specific -- reviewers on CUDA/MetaX/Ascend should look at it, since it affects every backend that runs the FlagGems Python path on more than one device. 1. FlagGems factory ops ignored their `device` argument (general multi-device bug) CallPythonOp_Factory / CallPythonOp_LikeFactory injected a hardcoded `device=flagos:0` kwarg, dropping the `device` the generated kernels already had in scope. Two consequences: - torch.ones(4, device="flagos:1").device reported flagos:0. - the output was allocated on device 0 while gems' Triton kernel launched on the *current* device -- a cross-device write, which faults the GPU rather than raising. On DCU this surfaced as "Invalid address access, Error code: 3" / KERNEL VMFault, and it made every rank>0 worker crash under FlagGems since DDP workers build their tensors via factories. Both callers now take an optional at::Device and resolve the index once: the argument's index when it carries one, else c10::flagos::CurrentDevice(), matching aten factory semantics (an index-less device means "current" -- the same rule csrc/aten/empty.cc applies via c10::device_or_default). The Python call runs under a c10::DeviceGuard on the resolved device so gems' internal torch.empty and the Triton launch agree. For *_like, an absent device falls back to the source tensor's device ("same as self") before the current-device default. codegen_ops.py forwards the aten `device` arg at both emit sites; the 18 regenerated call sites (13 Factory + 5 LikeFactory) are the only change to flaggems_python_kernels.cc. TensorToPython's CPU-scalar hop moves to the current device too, so a scalar operand on device N no longer drags the computation back to device 0. 2. comm layer did not know GEMS_VENDOR=hygon _VENDOR_PROFILES had no hygon row, so DCU tripped the unknown-vendor warning and fell back to guessing nvidia. Worse, with GEMS_VENDOR unset _patch_flaggems_codegen_config() reached the ascend fallback on DCU (is_nvidia_cuda_available() is False -- DTK has no libcuda.so, only libgalaxyhip), selecting the HCCL profile and failing outright with "no suitable inner backend for GEMS_VENDOR='ascend'". Adds the hygon row (CUDA-ABI: zero-copy view + NCCL, which on DTK is RCCL) and a DCU branch that sets GEMS_VENDOR=hygon, keyed on the existing _build_accelerator() helper and placed before both the generic-NVIDIA branch and the ascend fallback. setdefault, so an explicit GEMS_VENDOR still wins. The transport itself needed no work. Verified on Hygon DCU (DTK 26.04-rc4, 8x BW), GEMS_VENDOR left unset: - ones/zeros/randn/empty/rand/full/arange/eye/linspace and the *_like variants all land on flagos:1 with correct values; dispatch log confirms `ones -> flagos_python`, i.e. the gems path. - tests/manual/test_flagos_dist_live.py with FlagGems ON: all five collectives + DDP fwd/bwd pass at --world-size 2 and 8 (previously VMFaulted). - 88 passed across tests/unit + test_allocator + test_factory_ops + the new test_factory_device_index, with FlagGems both on and off. - full per-file ops sweep: 37/39 clean with FlagGems on, 38/39 with it off (test_full_cuda_coverage 46/46, test_ops 58 passed). Pre-existing upstream failures, unchanged by this work and reproducible on unmodified main: test_mul_dispatch::test_dispatch_log_flaggems_runtime and the orphan/count checks in test_flaggems_conf_consistency (flagos-ai#28 routed mul.Tensor to cuda but left its kFlagOsPython kernel registered).
Event: flagos.Event was a host-timestamp stand-in whose wait() was a no-op, so cross-stream ordering silently raced. Wrap torch.cuda.Event (a real maca event under boxing) instead. Boxing patches torch.cuda.current_stream into a non-Stream shim (triton launch only), which breaks Event.record()/stream() -- so resolve the real current stream straight from torch._C._cuda_getCurrentStream and reimplement the stream() context manager via _cuda_setStream, bypassing the shim. pin_memory: _to_copy hard-rejected pin_memory=True. The host allocator (cudaMallocHost), isPinnedPtr and _pin_memory dispatcher already work, so allow pin_memory=True for a CPU destination (pin the result via at::_pin_memory) and keep rejecting it for device destinations. Verified on 2xC550 (tests/manual/metax/test_event_pin_metax.py, 11/11).
CI runs `lint` first and every other job declares `needs: lint`, so a formatting slip skips the build and integration jobs entirely -- a red lint tells you nothing about whether the change actually works. Cheap to avoid by running the gate locally first. Records the contract from .github/workflows/lint.yml: ruff==0.15.12 (pinned, since formatting rules shift between releases), `ruff check .` plus `ruff format --check .`. Also notes two things that cost time in practice: `ruff format --check` reports only a count in the CI summary so you need `--diff` locally to see what it wants, and on vendor boxes where the torch install is in the system interpreter (DCU/DTK, where bare `python` is conda with no torch) the interpreter has to be spelled out or ruff lands in the wrong environment. Lives in .claude/skills/ rather than CLAUDE.md because CLAUDE.md is gitignored here -- a skill is the mechanism that actually travels with the repo to other machines.
…njection (flagos-ai#39) Unifies the native-CUDA RNG half with the FlagGems RNG half so torch.manual_seed makes ALL flagos RNG reproducible. Prior work unified only the FlagGems Python hot path via torch.cuda.default_generators; native ops routed `= cuda` (normal_/bernoulli/randint/random_/log_normal_/randperm...) still boxed to CUDA and drew from ATen's getDefaultCUDAGenerator(), which is unreachable from torch.manual_seed on the CPU-torch + external libtorch_cuda.so wheel (torch._C has no _cuda_manualSeed binding). Result was two disjoint RNG worlds. Approach -- C++ generator injection: - GetFlagosDefaultCudaGenerator(int64_t idx): new helper in python_op_caller.cc. Under the GIL, fetches torch.cuda.default_generators[idx] -- the same per-device torch.Generator(device="cuda") the vendor compat shim installs and FlagGems reads -- and returns it as at::Generator (process-lifetime cached by index). - Codegen (scripts/codegen_ops.py): for every native kernel whose schema carries Generator?, emit `if (!generator.has_value()) generator = GetFlagosDefaultCudaGenerator(<dev>)` before the at::<op>() call. Covers 80 kernels across functional/inplace/out/tuple/factory templates. No-op when the caller passes a generator (backward compatible). - Generator-less factory overloads (randint/randint.low/randperm): torch.randint(...) / torch.randperm(...) dispatch to overloads whose schema has no Generator? arg, so the rule above did not reach them. These bases each expose a sibling overload with Generator? right after the size arg (per ATen/ops/{randint,randperm,rand,randn}.h), so inject the shared generator there. randperm falls out for free (its randomness is an internal torch.randint). - CMakeLists: build python_op_caller.cc whenever CUDA_KERNEL=ON (cuda_kernels.cc now calls the helper regardless of the FlagGems flags). No routing/config changes -- native RNG stays `= cuda`; the unification is internal to those kernels. Verified on A100 (FLAGOS_USE_FLAGGEMS=1): rand/randn/rand_like/uniform_/ exponential_/multinomial (FlagGems) AND normal_/bernoulli/random_/log_normal_/ cauchy_/geometric_ (native) AND the former gaps randint/randint.low/randperm are all reproducible and seed-sensitive under torch.manual_seed. Distributions correct; explicit generator= still honored. No regression: rng_dispatch 10 passed, flaggems_python 37 passed, native 335 passed/58 skipped/3 xpassed (pre-existing conv1d segfault excluded). Includes design + implementation-plan docs under docs/superpowers/.
…agos-ai#40) * feat: add Moore Threads MUSA backend (musart runtime + mudnn kernels) Adds ACCELERATOR=musa, taking the native operator route like Ascend and GCU: generated kernels over mudnn, the vendor's torch-independent kernel library. There is no libcudart shim on this platform, so the runtime layer is a native port onto the musa* API and no CUDA boxing kernels are built. mudnn (not at::musa::*) is the integration point on purpose. torch_musa also exposes a flat at::musa::* API, which is cheaper to wrap but lives in libmusa_python.so -- that links against torch and embeds its C++ object layout, pinning the plugin to one exact torch build. sizeof(c10::MessageLogger) changed 408 -> 400 between 2.9.1 and 2.10, which corrupts the vendor binary's stack (tensor.sum() segfaults inside empty_musa). libmudnn.so has no torch symbols, so nothing here embeds torch's object layout, and the build needs only the MUSA toolkit -- no torch_musa package, no extracted vendor .so. Coverage is 64 generated ops (scripts/codegen_mudnn.py, category-driven like codegen_gcu.py) plus 2 hand-written convolution kernels; everything else reaches the cpu_fallback. Convolution is hand-written because ATen's default for convolution_overrideable raises rather than being boxable to CPU. Kernel design follows what mudnn was measured to do, not what the headers suggest: - mudnn Tensors honour strides including 0-strides, so broadcasting is expand() alone and strided inputs are read in place -- no .contiguous() as GCU needs. MudnnCopy uses this for _copy_from/contiguous via Unary IDENTITY/CAST. - int64 works across Unary/Binary/Reduce/MatMul, so there is no int64 fallback branch, but bool is rejected for arithmetic modes and needs its own predicate. - TF32 defaults on in mudnn while torch defaults matmul TF32 off; the handle now follows at::globalContext().allowTF32CuBLAS(), without which a 64x64 float mm drifts ~2e-2 from CPU. - mudnn's Convolution is 2D-only, so conv1d runs as 2D with a unit H dim and conv3d falls back. The forward algorithm is probed and cached per shape, since GetRecommendForwardAlgorithm can name one the Run then rejects. - Reduce raises SIGFPE -- uncatchable, not a status -- on a multi-dim reduce of a fully 0-strided input, which conv bias gradients hit via ones.expand(); such inputs are materialized first. Verified on 8x MTT S5000 against the supported target, torch 2.10.0+cpu: clean build and 478 passed / 209 skipped / 3 xpassed, no failures (the 9 errors are tests/integration/test_qwen3_* failing setup because `transformers` is not installed on this box). MLP and conv training loops match CPU to 1.2e-07. The same tree was also run clean against 2.9.1, but 2.10.0+cpu is the version that is supported and tested. Also moves the PrivateUse1-already-claimed check ahead of the _C import: when torch_musa happens to be installed, its autograd fallback registration turns a late check into an uncatchable std::terminate instead of an actionable error. * fix(allocator): zero-byte allocations must report the current device CachingDeviceAllocator::allocate hardcoded index 0 for nbytes == 0, so torch.empty(0, device="flagos:1") came back as flagos:0. That leaks into the composite factories: arange and eye are implemented as `at::empty({0}, options)` followed by an `_out` variant, so both dropped their device index. linspace allocates `steps` elements up front, which is why it was unaffected and the bug went unnoticed. Use the current device instead -- empty_memory_format has already installed a DeviceGuard for the requested device, so it is the right index. A zero-byte request should not require a live runtime, so a failed index query degrades to 0 rather than raising, and the TORCH_CHECK stays on the path that actually allocates. Fixes test_factory_device_index.py::test_{arange,eye}_honors_device_index. Not backend-specific: the passthrough DeviceAllocator already did this correctly, so only the caching path was wrong.
…ily) (flagos-ai#43) Grows GCU coverage from 67 to 101 topsaten kernels, all via scripts/codegen_gcu.py. Ops without a working vendor kernel stay unregistered on PrivateUse1 and keep reaching cpu_fallback, so nothing regresses. Tier-1 compute (16 ops, 14 new categories): amax/amin, tril/triu, flip, clamp, addmm + addmm.out, cat, zeros_like/ones_like, native_layer_norm, _softmax_backward_data, silu_backward, mse_loss/mse_loss_backward. foreach (18 ops, 11 new categories): _foreach_add_/sub_/mul_/div_ across the Scalar, List, Tensor and ScalarList overloads, _foreach_neg_, _foreach_sqrt_, _foreach_sqrt, _foreach_lerp_.Scalar, and _foreach_addcmul_/addcdiv_ in both Scalar and ScalarList forms. The foreach set was chosen by tracing a real workload rather than by guessing: a TorchDispatchMode over two AdamW steps on a 2-block transformer showed the foreach ops were the only remaining *compute* on the fallback path, with every optimizer step copying all parameters to the host and back. That trace now reports 19 falling-back ops instead of 27, and the remainder are view/metadata ops plus three broken vendor kernels. Two hardware findings shaped the implementation, both probed before any codegen: - topsaten's foreach ops accept an output list that aliases an input, so the in-place _foreach_*_ kernels update parameters directly with no temporary. They also broadcast a 1-element rhs across the list, which is what _foreach_add_.Tensor needs (aten requires that operand 0-dim; the wrapper presents rank-0 as shape {1}). - topsatenNativeLayerNormBackward rejects every output_mask ("LNB Output mask is not supported now!") yet still returns TOPSATEN_STATUS_SUCCESS while writing nothing, so a caller would see uninitialized gradients rather than an error. It is deliberately absent from OPS, with a comment; the forward is correct and is registered. A CPU fallback must never call back into the same aten op on a flagos tensor -- it re-enters the kernel and recurses until the stack overflows. zeros_like hit this and core-dumped; the foreach fallbacks loop the equivalent per-tensor method (t.add_, t.addcmul_) instead of at::_foreach_*. Also fixes ones_like under memory_format=Preserve, which autograd passes for every seed gradient and which at::empty rejects outright -- that broke .backward() entirely. New in topsaten_common.h: TopsatenTensorList (owns one wrapper per list element so every sizes/strides array outlives the call), IsForeachEligible (uniform device, contiguous, zero offset, non-empty, representable dtype), and DtypeLowest/DtypeHighest for filling in an absent clamp bound. Also adds the missing "gcu" row to the test suite's _PLATFORM_SKIP_MARKERS. Since flagos-ai#40 the ops conftest resolves the platform from lib/flagos_platform, which reads "gcu" here -- a key absent from the table, so GCU got an empty skip list and ran every other backend's routing tests. Those were passing only because the ops they exercise had no GCU kernel and raised the expected "backend not registered"; adding cat/ones_like/silu_backward made three of them fail for the right reason. GCU skips the same markers MUSA does (plus musa), for the same reason: the tops stack ships no CUDA runtime, so no boxing kernels are compiled and FlagGems is not built. Verified on an S60: every kernel compared against CPU including non-contiguous (transposed) inputs, broadcast, keepdim, negative dims, fp16/bf16, and int64 fallback; in-place foreach ops additionally checked for mutating the original tensors rather than a temporary. An 8-step AdamW run matches the CPU reference to 1.2e-07 with monotonic loss. flagos:1 and flagos:3 agree with CPU. Against a pristine build of the parent commit this branch adds no test failures and removes 95: 6 failed / 387 passed / 294 skipped, versus 101 failed / 526 passed / 60 skipped. The 6 remaining are the pre-existing conv1d cases needing convolution_overrideable; the 9 errors are the absent transformers package.
…er CI (flagos-ai#49) The unified-RNG contract -- one `torch.manual_seed` reaching every RNG op on the flagos device -- had no CI protection at all, and two codegen templates still leaked. Injection gaps. Generator-less RNG overloads carry no `Generator?` in their schema, so `_generator_inject_line` never fires and they fall back to ATen's own default CUDA generator, unreachable from `torch.manual_seed` on a CPU-torch wheel. `torch.randint_like(...)` and every RNG out-variant dispatch to exactly such overloads. The fix has to be made once per template because the ATen insertion point differs each time: `*_like` takes the generator before `dtype` (it carries `self`, so it lands on gen_functional_pure rather than gen_factory), while out-variants take it before the first of names/memory_format/out. That second one covered 11 kernels -- rand/randn `.out` and `.names_out`, rand_like/randn_like `.out`, randint `.out`/`.low_out`, randint_like `.out`/`.low_dtype_out`, randperm `.out`. The three base sets are now module-level constants, one per template, so all three positions read together. `_is_in_bad_fork`. `torch.random._seed_custom_device` needs both it and `manual_seed_all` to seed a custom device; only the latter existed, so every `torch.manual_seed()` warned "Set seed for `flagos` device does not take effect" -- false, since seeding really goes through the patched `torch.cuda.manual_seed_all`. Adding it silences the warning and makes `flagos.initial_seed()` track the global seed. Tests. The old file was marked `flaggems_python` with no `main_ops`, while CI selects only `-m main_ops` and `-m "flaggems and main_ops"` -- so it never ran. It also skipped itself entirely without FLAGOS_USE_FLAGGEMS, on the premise that the native path could not be reproducible; C++ generator injection made that premise obsolete. Rewritten to run under both configs (the native-path injection is only exercised by the vendor run) and to cover 39 ops grouped by the mechanism each stresses, plus the seed plumbing itself, multi-device generators, and distribution correctness. CI now collects 105 RNG tests in the vendor job and 2 in the FlagGems job, up from 0; no workflow change was needed, the markers are the wiring. Two gaps injection cannot reach are recorded as non-strict xfails rather than left implicit: `native_dropout` has no `Generator?` anywhere in its ATen schema, so the vendor path has no argument to inject into; and FlagGems `multinomial` launches its Triton kernel against the wrong device for any index != 0, a device-context bug on that path rather than an RNG one. Verified on 8xA100 under both configs: 102 passed/2 skipped/1 xfailed (vendor) and 104 passed/1 xfailed (FlagGems). Regression clean -- main_ops 115 passed, flaggems main_ops 14 passed, flaggems_python 27 passed, native 445 passed/3 xpassed (pre-existing conv1d segfault excluded). Both ruff gates pass.
…ified broken) (flagos-ai#48) Re-ran codegen discovery against flag_gems 5.4.0dev and executed every newly routed op against a CPU reference on Hygon DCU. discover_flaggems_ops() gates only on static arity/type inspection of the gems function, so it cannot tell whether the call actually works: of 98 new candidate routes, 26 (27%) fail at runtime. All 26 pass on the cuda route and fail as flagos_python, so each is a genuine route regression rather than pre-existing breakage. Net +72 routes: 318->390 (nvidia), 316->388 (dcu), 305->377 (metax). New flaggems_runtime_broken set, grouped by cause: - gems declares `out` as a required keyword-only arg (14) -- the out_variant caller passes only non-out args and copy_'s the result, so these need a kwarg-out category in the C++ caller to be recovered. - gems asserts device.type == "cuda" (8) -- a PrivateUse1 tensor trips it. - DTK triton (hcu) cannot compile the kernel (2): gcd, gcd.out. - wrong numerics on the gems path (2): leaky_relu_, adaptive_max_pool3d_backward. Kernels stay generated (417) and reachable via FLAGOS_OP_<op>=flagos_python; only the default route goes to cuda, same treatment as flaggems_recursive_fallback. Two codegen fixes found along the way: - gen_inplace() did not box optional<Tensor> args, so the guard saw a mix of boxed and unboxed tensors. cuda_kernels.cc had a hand-applied fix for ClampInplaceTensor that any regeneration silently reverted; the generator now emits it. This is the only cuda_kernels.cc change. - conf generation stripped the Apache header off the three checked-in confs on every run; it is now emitted. torch_fl/codegen_skip_ops.txt gains the 64 vendor-only native_fuse_* ops the DTK wheel's bundled torchgen yaml declares, which would otherwise put DTK-only includes into the shared cuda_kernels.cc. Codegen output is now platform-neutral: FLAGOS_CODEGEN_ALL=1 produces no diff on either wheel. test_flaggems_conf_consistency: generalize the AST reader to both forced-cuda sets, and anchor _wrapper_to_dispatcher's regex to `Wrapper` -- the old pattern matched the file's own header comment ("m.impl() lines.") and ate the first real wrapper, which surfaced now that _adaptive_avg_pool2d is routed. Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…flagos-ai#52) The three `default_generators` shims (metax compat, cuda compat, and the flagos module) stand in for a value upstream types as a *tuple*, so callers are entitled to iterate it, slice it or wrap it in `list()`. All three were list-like proxies defining only `__getitem__`/`__len__`, and `__getitem__` was unbounded. With no `__iter__`, Python falls back to the legacy iteration protocol: call `__getitem__(0, 1, 2, ...)` until IndexError. An unbounded `__getitem__` never raises, so `for g in torch.cuda.default_generators` became an infinite loop that allocated a fresh CUDA generator on every step. It presented as a hang rather than an error, which is why nothing surfaced it -- found while walking the generator list on an 8-card C550, where the process sat at 200% CPU indefinitely. The flagos variant did not hang but was equally wrong: out-of-range indices surface as a RuntimeError from C++, which aborts iteration instead of ending it. All three now define `__iter__` and bounds-check `__getitem__`, raising IndexError past the device count and wrapping negative indices the way the tuple they replace does. Slices are supported for the same reason. Regression test is parametrized over both the `torch.cuda` and `flagos` shims and asserts all three properties; verified to hang on the unfixed tree and pass on the fixed one. Verified on 8xC550 (metax boxing): RNG suite 104 passed/2 skipped/1 xfailed (vendor) and 106 passed/1 xfailed (FlagGems). Wider probe of 91 RNG ops covering every overload family, dtype and nn.init entry point: 89/91 clean on the vendor path (the 2 are the known native_dropout schema gap) and 91/91 under FlagGems. Distribution moments checked against theory on all 8 cards. Both ruff gates pass. Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
…lagos-ai#47) The skill covered lint but said nothing about the base the lint runs on, and lint-before-rebase is wasted work: a rebase can reintroduce lint errors or pull in regenerated files. Reorder into an explicit 4-step sequence with rebase first, and document the three ways the rebase actually goes wrong here: - A stale local flagos/main ref. This repo moves fast (ten commits landed in two days recently) and the network to github is slow enough that fetches get skipped, so the cached ref silently drifts. Check the ref's *date*, not just its sha. - A commit that already landed upstream. Squash-merges change the sha, so git cannot dedupe it and the rebase applies the same change twice -- the whole PR comes back conflicted. Drop the local commit; upstream's copy wins. - Generated-file conflicts. Take upstream's side, port only the generator change, re-run codegen, and verify a second run leaves no diff, rather than hand-merging conflict markers in csrc/aten/generated or the backend confs. Also warn that `git checkout <commit> -- <file>` on a hand-written file discards upstream's edits to it silently; `git apply --3way` surfaces the conflict instead. Checklist gains the rebase, commit-range and codegen idempotency steps. Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…lagos-ai#42) * ci: add configured MetaX build and integration workflows * ci: fix MetaX workflow input expressions * ci: persist isolated MetaX Python path * ci: preserve platform environment in integration tests * ci: isolate CUDA wheel from vendor torch * ci: install CUDA integration test dependencies * ci: tolerate CUDA images without standalone libc10_cuda * ci: pin CUDA CI image digest * ci: build MetaX integration wheel locally * ci: diagnose CUDA library layout * ci: fix native library staging * ci: load isolated native runtimes * ci: build CUDA integration wheel locally * ci: stage CUDA runtime dependencies * ci: diagnose MetaX runtime symbols * ci: retain MetaX CUDA runtime dependency * fix: provide MetaX generator ABI shim * fix: use vendor CUDA generator for MetaX boxing * ci: capture first MetaX operator failure * fix(rng): close the last two generator-injection gaps and put RNG under CI (flagos-ai#49) The unified-RNG contract -- one `torch.manual_seed` reaching every RNG op on the flagos device -- had no CI protection at all, and two codegen templates still leaked. Injection gaps. Generator-less RNG overloads carry no `Generator?` in their schema, so `_generator_inject_line` never fires and they fall back to ATen's own default CUDA generator, unreachable from `torch.manual_seed` on a CPU-torch wheel. `torch.randint_like(...)` and every RNG out-variant dispatch to exactly such overloads. The fix has to be made once per template because the ATen insertion point differs each time: `*_like` takes the generator before `dtype` (it carries `self`, so it lands on gen_functional_pure rather than gen_factory), while out-variants take it before the first of names/memory_format/out. That second one covered 11 kernels -- rand/randn `.out` and `.names_out`, rand_like/randn_like `.out`, randint `.out`/`.low_out`, randint_like `.out`/`.low_dtype_out`, randperm `.out`. The three base sets are now module-level constants, one per template, so all three positions read together. `_is_in_bad_fork`. `torch.random._seed_custom_device` needs both it and `manual_seed_all` to seed a custom device; only the latter existed, so every `torch.manual_seed()` warned "Set seed for `flagos` device does not take effect" -- false, since seeding really goes through the patched `torch.cuda.manual_seed_all`. Adding it silences the warning and makes `flagos.initial_seed()` track the global seed. Tests. The old file was marked `flaggems_python` with no `main_ops`, while CI selects only `-m main_ops` and `-m "flaggems and main_ops"` -- so it never ran. It also skipped itself entirely without FLAGOS_USE_FLAGGEMS, on the premise that the native path could not be reproducible; C++ generator injection made that premise obsolete. Rewritten to run under both configs (the native-path injection is only exercised by the vendor run) and to cover 39 ops grouped by the mechanism each stresses, plus the seed plumbing itself, multi-device generators, and distribution correctness. CI now collects 105 RNG tests in the vendor job and 2 in the FlagGems job, up from 0; no workflow change was needed, the markers are the wiring. Two gaps injection cannot reach are recorded as non-strict xfails rather than left implicit: `native_dropout` has no `Generator?` anywhere in its ATen schema, so the vendor path has no argument to inject into; and FlagGems `multinomial` launches its Triton kernel against the wrong device for any index != 0, a device-context bug on that path rather than an RNG one. Verified on 8xA100 under both configs: 102 passed/2 skipped/1 xfailed (vendor) and 104 passed/1 xfailed (FlagGems). Regression clean -- main_ops 115 passed, flaggems main_ops 14 passed, flaggems_python 27 passed, native 445 passed/3 xpassed (pre-existing conv1d segfault excluded). Both ruff gates pass. * fix(rng): make default_generators iterable instead of looping forever (flagos-ai#52) The three `default_generators` shims (metax compat, cuda compat, and the flagos module) stand in for a value upstream types as a *tuple*, so callers are entitled to iterate it, slice it or wrap it in `list()`. All three were list-like proxies defining only `__getitem__`/`__len__`, and `__getitem__` was unbounded. With no `__iter__`, Python falls back to the legacy iteration protocol: call `__getitem__(0, 1, 2, ...)` until IndexError. An unbounded `__getitem__` never raises, so `for g in torch.cuda.default_generators` became an infinite loop that allocated a fresh CUDA generator on every step. It presented as a hang rather than an error, which is why nothing surfaced it -- found while walking the generator list on an 8-card C550, where the process sat at 200% CPU indefinitely. The flagos variant did not hang but was equally wrong: out-of-range indices surface as a RuntimeError from C++, which aborts iteration instead of ending it. All three now define `__iter__` and bounds-check `__getitem__`, raising IndexError past the device count and wrapping negative indices the way the tuple they replace does. Slices are supported for the same reason. Regression test is parametrized over both the `torch.cuda` and `flagos` shims and asserts all three properties; verified to hang on the unfixed tree and pass on the fixed one. Verified on 8xC550 (metax boxing): RNG suite 104 passed/2 skipped/1 xfailed (vendor) and 106 passed/1 xfailed (FlagGems). Wider probe of 91 RNG ops covering every overload family, dtype and nn.init entry point: 89/91 clean on the vendor path (the 2 are the known native_dropout schema gap) and 91/91 under FlagGems. Distribution moments checked against theory on all 8 cards. Both ruff gates pass. Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: nate.river <lvyufeng@cqu.edu.cn> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
…lagos-ai#41) Wires torch.compile into the flagos (PrivateUse1) backend by registering flagos with TorchInductor as a real GPU device, so the traced graph is handed to compile_fx unchanged and inductor emits Triton kernels that operate on flagos tensors directly. The earlier approach rewrote the graph and its example inputs to cuda before compiling. That is not just a copy per call: at::getAccelerator() is PrivateUse1/flagos in this build, and torch::autograd::Node::stream() only yields a stream when a node's input device type equals the accelerator. A cuda-rewritten graph therefore produces stream-less autograd nodes, and AOT autograd's backward trace inside compile_fx trips opt_ready_stream && opt_parent_stream (engine.cpp:1085) -- the cause of 8 of the 11 test failures. Verified by differential test: eager backward on plain cuda fails identically with torch.compile never involved. Registration surface (device_interface.py, inductor_codegen.py): * GPU_TYPES gains "flagos" in place -- is_gpu() is a membership test on that list object, and without it inductor takes the C++/CPU codegen path and never emits Triton. get_gpu_type()'s functools cache is primed while the list is narrowed, since it asserts at most one GPU type is available and the torch.cuda shim reports available too. * DeviceInterface subclass: device state from torch.flagos, hardware properties from torch.cuda (same physical GPU, same allocator). * DeviceProperties.create reports flagos as cuda at the Triton boundary. Triton's NVIDIA backend hard-checks target.backend == "cuda", so a literal "flagos" finds 0 compatible backends. Inductor already does this rewrite in the opposite direction for ROCm (hints.py:149). * Device op overrides + scheduling/wrapper codegen: the stock CUDA/Triton pipeline under the "flagos" key, also published on torch.flagos for inductor's official PrivateUse1 hook. Two generated-kernel bugs that only surface under compilation: * detach re-dispatched into itself. The kernel called at::detach(self), also registered on PrivateUse1, so it dispatched straight back. Eager hid the recursion because DeviceBoxingGuard rewrites self's device metadata first; under FakeTensor it cannot, since the Python dispatch key sits above the backend key. Dynamo traces every nn.Linear through detach, so this was a stack-overflow segfault at trace time. Now emits at::native::detach (NATIVE_DIRECT_VIEW_OPS). * gen_inplace passed only plain at::Tensor args to DeviceBoxingGuard, so clamp_.Tensor handed unboxed flagos min/max to a CUDA self and crashed. Optionals are now materialized into holders, matching gen_functional_pure. Both regressions are covered by tests that were confirmed to fail (segfault at the exact asserting line) against a build with the fixes reverted. CPU-torch wheel accommodations, since torch.cuda's Python layer was frozen without CUDA: re-attach CudaInterface.get_raw_stream (binding exists, the import-time _is_compiled() probe left it None), route torch.cuda.memory_* to the flagos allocator that backs the same pool, hand out flagos Event/Stream in place of the dummy base classes, force triton.cudagraphs off (torch.cuda.CUDAGraph raises on construction) and use_static_cuda_launcher off (not built). flagos_compile_backend now accepts the mode/options/dynamic kwargs dynamo forwards to named backends and expands them into compile_fx config_patches, rather than mutating inductor's global config. Tests: test_compile.py 12 passed / 1 skipped (was 2 passed / 8 failed); test_clamp_dispatch.py 15 passed; ops dispatch sweep 358 passed; test_ops.py 58 passed; allocator/factory/fallback/unit 73 passed. Docs updated to drop the device-aliasing description and the unmeasured performance-parity figures; benchmarking remains open work. Co-authored-by: Claude <claude@anthropic.com> Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…s-ai#54) Two independent multi-device bugs, both reached through csrc/aten/backends/flagos/python_op_caller.cc. 1. Missing DeviceGuard in the non-factory callers flagos-ai#36 gave the two factory callers a DeviceGuard. The other 16 never got one, so the same hazard remained reachable from the other side: gems allocates its intermediates with `device=input.device` but launches Triton on the *current* device. Call an op on a tensor living on device 1 while device 0 is current and the kernel reads and writes across devices. That does not raise -- it faults the GPU (segfault / KERNEL VMFault / "Invalid address access"), so it presents as a process kill with no Python traceback. Found via `torch.multinomial(torch.ones(10, device="flagos:1"), 5)` under FLAGOS_USE_FLAGGEMS=1, which dies while device 0 works and while the boxing path works on both. Every caller now takes a c10::OptionalDeviceGuard resolved from its first tensor operand that names a device (DeviceOfTensor / DeviceOfArgs / DeviceOfFirst). Non-tensor leading args are skipped, so `topk(values, k)` still resolves correctly. nullopt leaves the current device alone, so single-device callers are unaffected. This also fixes TensorToPython's CPU-scalar hop, which resolves to the current device: under the guard that is now the operand's device rather than whatever happened to be current at entry. Not DCU-specific -- any backend running the FlagGems Python path on more than one device hits it. 2. dcu linked with an undefined GetFlagosDefaultCudaGenerator Pre-existing, and the reason the fix above could not be built and tested on dcu at all. The flagos backend was excluded on `NOT FLAGGEMS_KERNEL AND NOT FLAGGEMS_PYTHON AND NOT CUDA_KERNEL`, but generated/cuda_kernels.cc is not governed by CUDA_KERNEL -- it lives under aten/generated/ and is dropped per-accelerator (gcu, musa) instead. Since the unified-RNG generator injection its native RNG kernels call GetFlagosDefaultCudaGenerator() unconditionally, which python_op_caller.cc defines. dcu is the one combination that referenced it without defining it: CUDA_KERNEL=OFF and both FLAGGEMS_* OFF, yet cuda_kernels.cc still compiled, because the DTK torch wheel is hipified and registers its HIP kernels under the CUDA dispatch key. musa sets the same three OFF but also drops cuda_kernels.cc, so it escaped. Result: ImportError: libtorch_fl.so: undefined symbol: _ZN2at6native6flagos29GetFlagosDefaultCudaGeneratorEl The exclusion now runs after the per-accelerator rules and tests the actual source list for cuda_kernels.cc rather than the CUDA_KERNEL option. Verified by configuring both ways: dcu goes BROKEN -> ok, tsingmicro stays ok, gcu/musa still exclude the file. flagos-ai#42's MetaX case is the same shape of exception, so it becomes a third named consumer of the rewritten condition rather than an append inside the removed block. (Its cuda_kernels.cc is in the source list anyway, so the check would already keep the file; the clause stays for intent.) flagos-ai#42's libtorch_cuda.so link block is untouched. Testing (DCU, 8 devices) tests/integration/test_compute_device_index.py -- new, the non-factory companion to test_factory_device_index.py. Pins the current device to 0 and drives one op per caller shape on flagos:1, asserting values and not just placement, since a cross-device read returns whatever the other device's memory holds. Confirmed to catch the bug: with the guard removed from CallPythonOp_Generic alone, the multinomial case dumps core. compute + factory + rng, FlagGems 140 passed, 1 xpassed compute + factory + rng, boxing 138 passed, 2 skipped, 1 xpassed ops sweep, FlagGems, per file 455 passed, 0 failures, 0 crashes conf consistency 6 passed default dcu build (FLAGGEMS_PYTHON=OFF) builds and imports ruff 0.15.12 check + format clean Co-authored-by: Claude Opus 5 (1M context) <noreply@anthropic.com>
…gos-ai#53) * feat: adapt flagos CUDA backend to torch 2.11 via schema codegen Port the codegen approach from 2.13 onto this branch: generate all 71 CUDA boxing ops from native_functions.yaml via torchgen, replacing hand-written per-op .cu/.cc + structured_ops. - Bring over scripts/codegen_ops.py + generated/, device_boxing.h, dispatcher.h, refactored register.cc, external-libtorch scripts/docs, skill - Delete 123 hand-written kernel files superseded by codegen - Fix test_dispatch_log_bmm_out_flagos_default marker (cuda -> flaggems) Generated products are byte-identical to 2.13 (same native_functions.yaml schema for these 71 ops; ARRAYREF_OPS needs no change). codegen validated on torch 2.11 torchgen. End-to-end CPU-only + external libtorch_cuda.so verification deferred until download.pytorch.org (the +cpu wheel source) is reachable again. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * feat: register full CUDA op set into flagos backend via codegen enumeration Expand codegen from the hand-listed 71-op conf to the full CUDA dispatch set with automatic filtering. backends_cuda.conf now routes 1824 ops to the boxing CUDA kernel (zero hand-written kernels). Enumeration predicate is a strict superset of the original 71-op conf (0 missing): direct CUDA kernel union structured_delegate-with-CUDA-target union composite_explicit_autograd. Ops the templates cannot express are auto-skipped and fall back to cpu_fallback, so coverage only grows. Key changes: - codegen_ops.py: FLAGOS_CODEGEN_ALL full-enumeration mode; authoritative torchgen predicates (part_of_structured_group -> IListRef vs ArrayRef, use_const_ref_for_mutable_tensors -> mutable out param signature), root_name for ATen/ops headers, leading/trailing-underscore name disambiguation, compute-factory CUDA-device redirect (fixes randn 0-dim garbage), try/except auto-skip in all mode. - codegen_skip_ops.txt: ~211 ops the templates cannot express (multi-out, exotic signatures, dunder shifts, const-ref out variants). - test_full_cuda_coverage.py: 46 sampling tests across unary/binary/ reduction/shape/factory/foreach with CPU cross-check + dispatch routing assertions + randn 0-dim regression guard. - Fix silu_backward / nll_loss_backward tests to build inputs on CPU then move to device, instead of relying on cross-device same-seed randn (which no longer matches now that randn correctly uses the CUDA RNG). Full op suite: 311 passed, 0 failed, 64 skipped, 3 xpassed. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * feat(cuda): delegate flagos allocation to CUDACachingAllocator The flagos (PrivateUse1) CachingDeviceAllocator maintained its own block pool, so compute-factory ops (torch.randn/mm/etc.) — generated as CUDA boxing kernels that allocate straight on the CUDA device — bypassed it entirely. flagos memory_stats saw only `at::empty` allocations (~3% of the real footprint) and empty_cache could not release the boxed-kernel memory. Since flagos and CUDA share the same physical GPU memory (boxing only relabels the device, no copy), route ALL flagos allocation through the same c10::cuda::CUDACachingAllocator that boxed kernels already use. Now empty + boxed outputs live in one pool: memory stats reflect the true footprint, empty_cache works, and OOM-retry is unified. - DeviceMemoryInterface: add opt-in caching-delegation API (provides_caching / caching_alloc / caching_free / caching_empty_cache / caching_record_stream / caching_get_stats / caching_reset_peak_stats), default off so Ascend and other backends keep the self-built block pool. - CudaDeviceMemory: implement the API over CUDACachingAllocator (raw_alloc / raw_delete / emptyCache / getDeviceStats / resetPeakStats), with a lazy once_flag init(cudaGetDeviceCount()) — no Python lazy-init. - CachingDeviceAllocator: allocate / empty_cache / record_stream / get_stats / reset_stats branch to the backend when provides_caching(); add delegated_deleter for the delegation path. - Extract AllocatorStats into allocator_stats.h to break the include cycle between device_memory_interface.h and caching_device_allocator.h. - CMake: define C10_CUDA_NO_CMAKE_CONFIGURE_FILE (non-Ascend) since the CPU torch wheel ships c10/cuda headers but not cuda_cmake_macros.h. Still links torch_cpu only; CUDA symbols resolve at runtime from the preloaded libtorch_cuda.so (CPU-torch + external-libtorch scheme unchanged). Verified: torch.randn now tracked (allocated_bytes 0 -> 1048576), empty_cache releases reserved (86MB -> 20MB). Regression green — ops 426 passed, allocator 11 passed (test_allocator.py reverted to torch.randn), Qwen3 infer/train + fallback_trace 9 passed. * feat(cuda): recover ~200 skipped ops via codegen template generalization Generalize three codegen templates in scripts/codegen_ops.py so operators that previously fell through to cpu_fallback are boxed directly: - gen_out_variant now supports N mutable Tensor& outputs by calling at::<base>_outf(...) in faithful schema order (outs last), matching the generated signature. Recovers ~121 multi-out ops (sort, topk, svd, native_batch_norm.out). - New foreach_out category: materialize + box every TensorList including out, call _outf, void return. Recovers 75 _foreach_*.out optimizer ops. - New vector_return category: split/unbind returning std::vector<Tensor>, unboxed via UnboxTensorVecToFlagos. Remove the multi_out skip in enumerate_all_cuda_ops and the corresponding 80 entries from codegen_skip_ops.txt. Rename local 'result' to '_ret' in the tuple out-variant branch to avoid collision with out-params named 'result'. Fix with_cuda_libtorch.sh to add .libtorch_cuda_assets to LD_LIBRARY_PATH so SVD's lazy dlopen of libtorch_cuda_linalg.so resolves. Generates 2024 ops (up from ~1900 effective). Regression green: ops 311 passed, full CUDA coverage 46, core+Qwen3 infer/train 124 passed. * feat(cuda): recover 19 more ops — function-only inplace activations + TensorList-out gen_inplace now branches on torchgen variants: function-only inplace ops (silu_/gelu_/celu_/leaky_relu_/threshold_/hardtanh_/mish_/hardsigmoid_/ hardswish_/elu_/embedding_renorm_/rrelu_with_noise_) call the free function at::op_(self,...) instead of the nonexistent Tensor method self.op_(). gen_foreach_out now handles a single mutable Tensor& out with return (cat.out/stack.out/_stack.out/block_diag.out/_chunk_cat.out) and boxes every mutable Tensor& (not just the out list), branching return shape on ret_type. Tuple-returning multi-out RNN out-variants (_cudnn_rnn.out/_lstm_mps.out/ miopen_rnn.out) added to skip list — they were latent -Wreturn-type bugs. 2043 ops generated (was 2024). Regression: ops 311 passed, coverage 46, core+Qwen3 124 — no regressions. * feat(ascend): finish torch-2.11 codegen migration for Ascend backend The torch-2.11 migration (285f52c) unified per-op headers into generated/ops.h but only adapted the CUDA backend, leaving the Ascend backend on the deleted per-op-header layout so it no longer compiled. - Point all ascend + flagos python_wrapper kernels at generated/ops.h - Rename _softmax / sum.dim_IntList dispatchers to the codegen names (PrivSoftmaxFn/priv_softmax_dispatcher, SumDimIntlistFn/sum_dim_intlist_dispatcher) - Split mm/bmm into functional + out variants over a shared aclnn helper to match the new MmFn/MmOutFn (BmmFn/BmmOutFn) signatures - Implement AscendDeviceMemory (ACL-backed DeviceMemoryInterface) so the flagos block-pool allocator works on Ascend; previously GetCachingAllocator hit TORCH_CHECK(false) and .to("flagos:0") threw before any kernel ran Verified on Ascend 910 (torch 2.11.0+cpu, CANN): build+install+import OK; mm/bmm/add/mul/cos/sin/neg/abs/silu/rsqrt/softmax/sum.dim/mean.dim match CPU. * feat(ascend): aclnn codegen MVP — unary category Add a category-based aclnn codegen for the Ascend backend. Unlike the CUDA codegen (which emits one-line at::op boxing bodies and lets PyTorch marshal everything), aclnn needs per-op knowledge the aten schema does not carry (API name, arg marshaling, output allocation), so generation is driven by category templates + an aten->aclnn mapping table. - scripts/codegen_ascend.py: reuses codegen_ops.py:schema_to_cpp_name so symbol names match the dispatcher decls already in generated/ops.h; only fills the Backend::kAscend slot. Validates each aclnn symbol via nm on libopapi.so before emitting; skips unmapped/missing ops with a warning. - csrc/aten/backends/ascend/generated/ascend_kernels.cc: 8 unary ops (sqrt/exp/tanh/sigmoid/reciprocal/log/floor/ceil). Path auto-excluded from non-ascend builds by the existing CMake glob. - backends_ascend.conf: route the 8 new ops to ascend. - docs: design (ascend_aclnn_codegen.md), NPU plan (ascend_npu_plan.md), route-rejection record (cpu_torch_external_libtorch_npu.md), and the standalone feasibility prototype. Verified on Ascend 910: all 8 ops match CPU reference, max_err <= 1.2e-7. * feat(ascend): generalize aclnn codegen to 6 categories (51 ops) Extend the category-based aclnn codegen from unary-only to six categories, driven by a CATEGORIES dict (one kernel-body template each) + an OPS dict mapping each op to (category, aclnn-name override): - unary (28): erf/erfc/expm1/log2/log10/log1p/round/trunc/frac/sign/relu/ cosh/sinh/asin/atan/asinh/acosh/atanh/logical_not/bitwise_not + the 8 from P1 - binary (7): div.Tensor/pow.Tensor_Tensor/atan2/maximum/minimum/ bitwise_or.Tensor/bitwise_xor.Tensor - binary_alpha (1): sub.Tensor - binary_cmp (7, bool out): eq/ne/gt/lt/ge.Tensor + logical_and/logical_or - binary_scalar_alpha (2): add.Scalar/sub.Scalar - binary_scalar_cmp (6, bool out): eq/ne/gt/lt/ge/le.Scalar Device-coercion fix in the shared binary prologue: torch.sub(x, 3.0) and similar tensor-op-python-scalar forms lower to aten::<op>.Tensor (not .Scalar) with the scalar packed as a CPU scalar tensor. The prologue now coerces the other operand to self's device (other.to(self.options()) when not already on the flagos device), mirroring the handwritten add.cc; coercing dtype alone left CPU storage to be read as an NPU device address, producing all-nan output. Both operands are expanded+materialized to the broadcast shape since aclnn does not always broadcast. Symbol validation via nm on libopapi.so auto-excludes ops without the aclnn symbol or dispatcher (square/isnan/isfinite). Verified on Ascend 910: all 51 ops match CPU reference (unary max_err <= 4.4e-5, binary <= 4.7e-6, comparisons exact). * feat(ascend): add reduce categories to aclnn codegen (55 ops) Add three reduce-shaped categories to the aclnn codegen, taking the generated Ascend kernel set from 51 to 55 ops: reduce_dims amax/amin -- (Tensor, IntArrayRef dim, bool keepdim), same dtype. Reuses the handwritten sum.cc dim logic: wrap negative dims, empty list = reduce all, drop (or set to 1 with keepdim) each reduced dim high-to-low. reduce_dim_bool any.dim -- single int64 dim, bool out. aclnnAny takes a dim list, so the dim is wrapped into a one-element vec. cumsum (Tensor, int64 dim, optional dtype) -- same-shape scan. Reduce ops are heterogeneous (each aclnn reduce has its own arg layout), so there is no single "reduce" template; each sub-shape is its own category. The long tail (max.dim/min.dim tuple return, var/std/norm with correction/p args, argmax/argmin/prod/logsumexp with no aclnn symbol in this CANN) is left for later, bespoke handling. All 4 new ops verified on Ascend 910 vs CPU reference across single-dim, dim-list, negative-dim, all-reduce, and keepdim variants (8/8 subtests). * feat(ascend): expand aclnn codegen to 20 categories (77 ops) Add 11 new categories / 22 ops to the Ascend aclnn generator: - unary_bool (isinf), unary_scalar (leaky_relu/clamp_min/clamp_max/ fmod.Scalar), unary_two_scalar (softplus/threshold), unary_int (tril/triu), unary_dims (flip) - addcmul/addcdiv (3-tensor broadcast + Scalar value) - pow_scalar_tensor (pow.Scalar: Scalar self, Tensor exponent) - reduce_max_dim (max.dim/min.dim: tuple(values, int64 indices)) - cumprod (separate from cumsum: aclnnCumprod takes dim as aclScalar*) - act_backward (tanh_backward/sigmoid_backward), threshold_backward -- first training-oriented backward ops - grow binary (fmod.Tensor/floor_divide) and binary_cmp (logical_xor) All 22 verified on Ascend 910 vs CPU reference. Candidates were probed against both the dispatcher decls in generated/ops.h and the aclnn symbols in libopapi.so before templating; ops with no symbol or bespoke args (var/std/norm, argmax, gelu string_view, matmul family) are left long-tail and auto-skipped by the generator's nm validation. * feat(ascend): main-line aclnn codegen batch — activations/loss/gemm (91 ops) Extend the category codegen toward the training/inference main line: +14 ops across 9 new/grown categories, all verified on Ascend 910 vs CPU reference. New categories: - elu (alpha/scale/input_scale) - loss (mse_loss; reduction None=elementwise, Mean/Sum=scalar) - cummax_cummin (tuple values+int64 indices, same shape) - aminmax (tuple min+max, optional dim) - prod (scalar out, optional dtype) - gemm_addmm / gemm_baddbmm (beta/alpha + cubeMathType) - mv / dot Grown: unary_scalar += celu/softshrink/hardshrink; unary_two_scalar += hardtanh. smooth_l1_loss is intentionally left long-tail: its aclnn signature takes a by-value float beta, and EXEC_ASCEND_CMD marshals args through a fully variadic function-pointer typedef, which is unsafe for a bare float on aarch64 (beta arrived as 0 -> pure L1 output). Scalars wrapped as aclScalar* or int64 are varargs-safe; a raw float is not. * feat: codegen-ize FlagGems Python path — auto-discover 235 ops (5 → 235) Replace the hardcoded 5-entry FLAGGEMS_PYTHON_MAP with automatic discovery from flag_gems._FULL_CONFIG (433 ops) plus safety filtering, and generate per-category kFlagOsPython kernels through a schema-driven generic IValue caller. Discovery + safety gates (scripts/codegen_ops.py): - arity gate (hard red line): exclude ops where the flag_gems positional count differs from the aten arg count (gems silently drops trailing scalar args like add.Tensor's alpha / addcmul's value → wrong results). - type gate: every arg type must be in the generic caller's supported set; ScalarType excluded (IValue stores it as a plain int, indistinguishable). - categories: functional_pure / inplace / tuple_return / out_variant. - FLAGGEMS_PYTHON_SKIP holds 8 convergence holdouts: mm.out (required out kwarg the generic caller can't supply) and 7 ops with an unconditional device.type=="cuda" assert flagos PrivateUse1 tensors can't satisfy (maximum, minimum, 5 upsample variants). Result: 150 functional_pure + 57 inplace + 21 tuple_return + 7 out_variant = 235 ops routed to flagos_python. backends_flaggems.conf now auto-generated. Generic caller (python_op_caller.{h,cc}): IValueToPython covers Tensor/int/ double/bool/None/str/Scalar/IntList/DoubleList/BoolList/TensorList; CallPythonOp_Generic uses BuildPyArgs; add CallPythonOp_GenericTuple for tuple-returning ops; GetFunc resolves dotted module.func qualnames. Cleanup: delete stale hand-written flagos wrappers superseded by codegen. Verification: numerical spot-check 23/23 (err <= 3e-5); flaggems_python correctness 27 passed; CUDA-direct path 330 passed / 45 skipped / 3 xpassed (matches baseline, no degradation). * feat(ascend): add layer_norm/group_norm to aclnn codegen (93 ops) * feat(flaggems): recover 18 out-variants where gems takes out positionally (235 -> 253) The out_variant safety gate assumed flag_gems never accepts the aten `out` tensor, so ops whose gems function signature is (…non_out, out) failed the arity check (npos == non_out+out, not == non_out) and were excluded. Add an out_variant_gemsout category: when gems npos == #non_out + #out args, pass the aten out tensor(s) positionally to gems (which writes into them in place) and return the aten out arg. Distinguished from mm.out, whose gems out is a required keyword-only arg the positional caller can't supply (stays in FLAGGEMS_PYTHON_SKIP). Recovers: atan2/bmm/cosh/div/exp/expm1/fmin/hardsigmoid/i0/log10/logaddexp/ pixel_unshuffle/reflection_pad1d/reflection_pad2d/replication_pad1d/softshrink/ special_i0e/where.self .out variants. Verified: 10/10 numerical spot-checks via dispatcher, flaggems_python 27 passed, CUDA-direct 330 passed / 45 skipped / 3 xpassed (no degradation). * feat(ascend): add gelu/log_softmax/softmax-backward to aclnn codegen (98 ops) Add 4 backbone categories (6 ops), all verified on Ascend 910 vs CPU: - gelu / gelu_backward: use aclnnGeluV2 (int64 approximate) + aclnnGeluBackwardV2 (char* approximate). v1 aclnnGelu hardcodes the tanh approximation, but PyTorch's default is approximate="none" (erf form, used by qwen3 et al); V2 selects 0="none"/1="tanh" so both modes match (err 1.8e-07 / 1.2e-07). - _log_softmax: mirrors handwritten softmax.cc (aclnnLogSoftmax(self,dim,out)). - _softmax_backward_data / _log_softmax_backward_data: aclnnSoftmaxBackward / aclnnLogSoftmaxBackward (grad_output, output, dim, grad_input), out dtype = input_dtype. * feat(flaggems): recover 4 ops with trailing-default gems params (253 -> 257) gems funcs addcdiv/round/scatter/scatter_ have extra positional params with defaults beyond the aten args (out=None, decimals=0, reduce=None). Relax the arity gate to admit npos > ncall when every extra gems param is strictly trailing with a default, guarded against the reordering trap (gems gather inserts out=None mid-signature -> aten sparse_grad would be misrouted into the out slot; correctly excluded). * feat(ascend): add addmv/addr + BCE loss family to aclnn codegen (103 ops) Add 5 ops across 5 categories, all verified on Ascend 910 vs CPU: - addmv / addr: gemm-family completion. aclnnAddmv arg order is (self,mat,vec, alpha,beta) -- alpha before beta, opposite of addmm. addr has no cubeMathType. - binary_cross_entropy (+optional weight), _backward, and _with_logits (+optional pos_weight). Optional tensors marshal via value_or(Tensor()) -> AclTensorWrapper nullptr, which aclnn treats as absent. Left out (probed but not shipped): addbmm (hf32 cube accumulation over the batch dim inflates rel-err to ~1e-2 vs ~1e-4 for a single addmm) and native_batch_norm (aclnnBatchNorm returns ACLNN_ERR_INNER_NULLPTR on 4D NCHW input; 2D N,C works). Both deferred to the conv/pool bespoke batch. * refactor(ascend): migrate 19 handwritten seed kernels to aclnn codegen Pre-codegen "seed" kernels (abs/cos/add.Tensor/mul.Scalar/where/softmax/sum/ mean etc.) had bodies expressible by codegen categories -- several byte-identical to the templates. Migrate them into scripts/codegen_ascend.py and delete the handwritten .cc files, per the rule: anything a codegen category can express goes through codegen; handwrite only genuine bespoke ops. - Reuse existing categories: abs/acos/cos/sin/neg/rsqrt/silu (unary), mul.Tensor/bitwise_and.Tensor (binary), add.Tensor (binary_alpha), pow.Tensor_Scalar (unary_scalar). - Add 7 new categories: binary_scalar (mul.Scalar->aclnnMuls, div.Scalar-> aclnnDivs -- headers absent, marshaling from handwritten refs), act_backward_self (silu_backward: grad+self), where (aclnnSWhere 3-tensor), softmax_fwd (_softmax, half_to_float), reduce_all (all), reduce_sum_dtype (sum.dim_IntList), reduce_mean_dtype (mean.dim via aclnnMeanV2). Kept handwritten (SKIP={le.Tensor,mm,bmm}): le (aclnnLe absent, needs runtime version probe), mm/bmm (also register out-variants codegen doesn't emit), factories, TensorList/SymInt ops, embedding, nll_loss. Handwritten kAscend regs 36->17, codegen 103->122, total unchanged. Verified all 19 migrated ops vs CPU on Ascend 910 incl. CPU-scalar coercion, half_to_float, dtype promotion, broadcast -- zero numeric regression. * feat(metax): reuse CUDA boxing kernels on MetaX via FLAGOS_METAX_BOXING Add a boxing build mode for the MetaX backend that reuses the generated CUDA boxing kernels (csrc/aten/generated/cuda_kernels.cc, host g++, no mxcc/nvcc) which dispatch PrivateUse1 -> CUDA into maca's libtorch_cuda.so, instead of the hand-written mxcc .cu kernels under backends/metax/. - CMakeLists.txt / setup.py: FLAGOS_METAX_BOXING=1 sets METAX_KERNEL OFF while keeping the MetaX SDK runtime + cu-bridge headers path. - csrc/CMakeLists.txt, torch_fl/csrc/CMakeLists.txt: define USE_MACA=1 for metax, required because maca's torch headers are a hard fork gated on it (C10_WARP_SIZE, Context.h allow_tf32_cudnn hit static_assert(0) otherwise). - Regenerate generated/* + backends_cuda.conf against maca 2.10 torch (2035 ops; +cudnn_convolution_bias_fused[.out], -upstream-only ldexp/ _foreach_powsum/miopen_ctc_loss/_flash_attention_forward.quantized). Verified on 8x MetaX C550: import + factory ops + add/mul/relu/softmax/ sum/exp/mm/bmm/copy all match CPU reference (matmul matches native maca-cuda exactly; ~1e-3 vs CPU is MetaX TF32 precision). * feat(ascend): aclnn codegen for conv/pool family (5 categories) Add convolution (fwd+bwd) and 2D pooling to the aclnn codegen. These are the first ops needing an explicit output-shape formula since aclnn requires the output pre-allocated -- each template carries a small shape helper. Categories: - adaptive_avg_pool2d, avg_pool2d, max_pool2d_with_indices (tuple w/ int64 idx) - convolution (non-transposed), convolution_backward (3-tuple + output_mask) Two infra additions in op_api_common.h: - AclTensorWrapper gains an optional aclFormat param (default ND). avg_pool2d/ adaptive_avg_pool2d/convolution reject ND 4-D tensors (GetWorkspaceSize ret 161002); they need NCHW/NCL/NCDHW per rank. max_pool2d does NOT care -- the requirement is per-aclnn, so the param defaults to ND and is opt-in. - AclBoolArrayWrapper for convolution_backward's output_mask[3]. conv cubeMathType=0 (KEEP_DTYPE); type 1 loses ~2.5e-3 on the cube unit. Verified vs CPU on Ascend 910: conv fwd/bwd err~1e-6 incl stride/padding/ dilation/grouped; pooling exact-to-1e-7 incl ceil_mode and stride defaults; prior migrated ops regression-clean. codegen 122->127 ops, 52 categories. * feat(flaggems): forward keyword-only gems args, recover 31 ops (257 -> 288) Many gems funcs declare aten's trailing args keyword-only (sum(inp,*,dtype), add(A,B,*,alpha), gelu(self,*,approximate), var(x,dim,*,correction,keepdim)). The positional-only generic caller couldn't pass them, so non-default values were silently dropped -> these were wrongly excluded as arity_short. Add CallPythonOp_GenericKw/GenericKwTuple: forward the trailing aten args by NAME via a PyKwarg vector. dtype (ScalarType) is tagged is_dtype so the caller converts the int payload to a torch.dtype (IValue can't distinguish ScalarType from int); absent optionals set is_none -> Python None. Discovery matches each trailing aten arg to a gems keyword-only param by name (reject on mismatch) and gates types via _FLAGGEMS_KWARG_OK. Recovers 31 ops incl add/sub/addmm/addmv/addr/addcmul (alpha/beta/value), sum/mean/prod/cumsum (dtype), gelu (approximate), var/std/var_mean (correction), isin (invert), sort.stable. Spot-checked all kwarg types numerically (err<=1e-5, non-default alpha/dtype/correction now correctly applied). Blocked: Generator?/ Layout?/Device? args and name-mismatches (multinomial/var.dim/_grouped_mm). * feat(ascend): aclnn codegen for max_pool bwd + batch_norm (CNN train loop) Complete the CNN training closure on top of conv/pool: add max_pool2d backward, native_batch_norm forward and backward. conv + pool + bn now cover a full CNN forward/backward pass. Categories: - max_pool2d_with_indices_backward -> grad_input - native_batch_norm -> (out, save_mean, save_invstd) - native_batch_norm_backward -> (grad_input, grad_weight, grad_bias) Two per-op quirks (each read from the aclnn header @param notes): - max_pool2d FWD ignores format and emits int64 indices, but the BWD kernel requires NCHW format AND int32 indices -- fwd/bwd are not symmetric. The bwd template casts indices to int32 and tags NCHW. - batch_norm's save_invstd uses a different definition than PyTorch CPU (~0.18 apart), but this does not affect training: the backward consumes the same NPU save_invstd and all three grads match CPU to <=4e-6. running_mean/ var are passed non-const (updated in-place). This overturns the earlier "native_batch_norm left out (561103)" note -- that was marshaling, not a real limitation. Verified vs CPU on Ascend 910: max_pool bwd 1e-7; bn fwd out 2e-7, bn bwd grads <=4e-6, bn eval exact; conv/pool/elementwise regression-clean. codegen 127->130 ops, 55 categories. * feat(ascend): aclnn codegen for pool/norm backward (fwd+bwd complete) Round out the backward coverage so conv, avg/adaptive pool, layer_norm and group_norm all have both forward and backward -- covers the remaining norm/ pool grads for CNN and transformer training. Categories: - avg_pool2d_backward, _adaptive_avg_pool2d_backward -> grad_input - native_layer_norm_backward -> (grad_input, grad_weight, grad_bias) - native_group_norm_backward -> (grad_input, grad_gamma, grad_beta) Notes: - avg_pool2d/adaptive_avg_pool2d backward also require NCHW format, matching their forwards. - norm backwards feed the forward's mean/rstd straight through and are self- consistent (no save_invstd-style mismatch: layer/group norm expose rstd, not invstd). - aclnn names: native_layer_norm_backward maps to aclnnLayerNormBackward (aclnnNativeLayerNormBackward does not exist); group_norm likewise. Verified vs CPU on Ascend 910, all <=1e-6 (avg_pool bwd 3e-8, adaptive exact, ln/gn grads <=1e-6); conv/max_pool/elementwise regression-clean. codegen 130->134 ops, 59 categories. * feat(flaggems): promote ScalarType positional args to dtype kwargs, +3 ops (288 -> 291) The positional caller can't carry a ScalarType (IValue stores it as a plain int), so ops taking dtype/input_dtype as a positional aten arg were excluded by the type gate (type_gate_dtype: 4 ops). Now that the kwarg path tags is_dtype and converts to torch.dtype by name, promote any ScalarType positional arg into a by-name kwarg -- safe when gems accepts it by name (not positional-only) and the ScalarType args form a strict suffix (guard rejects middle-dtype reordering). Recovers _softmax_backward_data, _log_softmax_backward_data (input_dtype), linalg_vector_norm (dtype). Spot-checked: backward err<=2e-7, vector_norm f32 ord=1/2 correct. _safe_softmax -> SKIP (gems device assert rejects PrivateUse1, same class as maximum/minimum). vector_norm dtype=f64 raises gems' own NotImplementedError (gems limitation, not a forwarding bug -- f32 path correct). Regressions green: cuda-direct 330 passed, flaggems_python 27 passed. * feat(ascend): aclnn codegen for masked_fill/gather/index_select Fill the remaining Transformer op gaps. A coverage probe showed embedding, gelu, silu, bmm, baddbmm and softmax were already handled (handwritten or earlier codegen), so the real gaps were masking and indexing. Categories: - masked_fill.Scalar, masked_fill.Tensor -> broadcast(self, mask) - gather -> index shape; index_select -> self shape w/ dim -> index.numel() Note on out-of-place via inplace aclnn: aclnn only ships inplace masked_fill (aclnnInplaceMaskedFillScalar/Tensor). Implementing the out-of-place aten op means copy-then-fill, but self.clone() routes through empty_like which is NOT registered for the ascend backend. Allocate via apply_tensor_without_format + out.copy_(self.expand(...)) instead -- applies to any copy-then-mutate kernel. SDPA/flash-attention is fused with a bespoke multi-tensor signature and is left to a dedicated batch. Verified vs CPU on Ascend 910: all exact (err=0) incl broadcast mask, 3D gather, negative dims, vocab-size index_select; core ops regression-clean. codegen 134->138 ops, 63 categories. * feat(flaggems): codegen factory ops (arange/eye/full/ones/zeros/linspace/logspace), +10 (291 -> 301) Factory ops don't take input tensors -- gems generates the tensor itself. New CallPythonOp_Factory injects device=flagos (so gems' internal torch.empty hits OUR allocator -> PrivateUse1 tensor, no CUDA round-trip, no recursion), layout=strided (gems eye/randperm reject layout=None), pin_memory=None, and forwards the aten dtype. discover_flaggems_ops gains a factory branch that strips the TensorOptions fields and passes only shape/scalar positionals; requires gems to accept dtype/layout/device by name (kwonly on every factory). Recovers 10 ops: arange (+.start/.start_step), eye (+.m), full, linspace, logspace, ones, zeros. Spot-checked all exact (err=0) vs CPU incl f64 dtype. Two correctness guards: - arange: gems defaults dtype=None to int64 unconditionally, but aten infers float when any of start/end/step is floating -> arange(0.,3.,.5) was silently wrong ([0,0,1,1,2,2]). Kernel now replicates aten's rule and passes an explicit dtype. Verified arange float now exact. - rand/randn/randperm SKIP: gems reaches default_generators[device], but the PrivateUse1 device has none (IndexError); randperm asserts an int dtype. Same root cause as the Generator? blocked group -- can't express per-device gen. Regressions green: cuda-direct 330 passed, flaggems_python 27 passed. * docs(ascend): record SDPA/flash-attention research findings - Investigated aclnnFlashAttentionScore/Grad API semantics - Confirmed: inputLayout="BNSD", softmaxMax/Sum shape [B,N,S,8] float32 - Confirmed: causal (sparseMode=3, preTokens=INT32_MAX, nextTokens=0) vs full (sparseMode=0, preTokens=65536, nextTokens=65536) - Key blocker: PyTorch logsumexp [B,N,S] vs aclnn (softmaxMax, softmaxSum) [B,N,S,8] shape mismatch; logsumexp=log(Sum)+Max but 8-way tiling collapse unclear - Conclusion: SDPA is bespoke multi-day task (autograd Function wrapper, ctx save/restore, causal/dropout/mask testing), deferred to dedicated pass Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * feat(ascend): implement SDPA forward via aclnnFlashAttentionScore Add handwritten kernel for _scaled_dot_product_efficient_attention (forward only). Wraps aclnnFlashAttentionScore with BNSD layout, handles logsumexp mapping from aclnn's softmaxMax+softmaxSum [B,N,S,8] to PyTorch's [B,N,S] via [:,:,:,0] indexing. Key findings: - attenMask semantics: true=MASK_OUT (opposite of docs), false=KEEP - Causal attention: triu(ones, diagonal=1) masks future positions - Verified: non-causal err=3.34e-06, causal err=7.15e-07 vs CPU Backward NOT implemented: aclnnFlashAttentionScoreGrad needs separate softmaxMax and softmaxSum, but PyTorch's autograd only saves single logsumexp (log addition not invertible). Forward-only covers inference; training needs architectural work. Files: - csrc/aten/backends/ascend/scaled_dot_product_attention.cc (new, 120 lines) - csrc/CMakeLists.txt (add to ascend sources) - test_sdpa_ascend.py (verification script) - docs/ascend_aclnn_codegen.md (document SDPA + attenMask/logsumexp pitfalls) Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * feat(ascend): implement SDPA forward and backward kernels - Add _scaled_dot_product_efficient_attention (forward) - Add _scaled_dot_product_efficient_attention_backward (backward) - Wrap aclnnFlashAttentionScore and aclnnFlashAttentionScoreGrad - Implement activation checkpointing for backward (recompute to get softmaxMax/Sum) - Support causal and non-causal attention - Add dropout=0 constraint (aclnn dropout needs explicit mask handling) - Tests pass: forward+backward, causal mask scenarios Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * feat: single self-contained wheel with runtime flaggems switch Move the flaggems-vs-cuda choice from compile time + external LD_PRELOAD to a pure runtime env switch, so one wheel serves both paths. - Default FLAGGEMS_PYTHON option ON: compile both CUDA boxing kernels and flaggems_python kernels into libtorch_fl.so unconditionally (MetaX keeps it opt-in). No more FLAGGEMS_PYTHON=ON at build time. - Bundle .libtorch_cuda_assets/*.so* into torch_fl/lib and ctypes-preload them (nvidia deps -> torch cpu libs -> cuda libs) before import torch, to satisfy the CUDAHooks hard constraint. FLAGOS_DISABLE_CUDA_ASSETS skips. - FLAGOS_USE_FLAGGEMS selects backends_flaggems.conf vs backends_cuda.conf at import; FLAGOS_OP_<name> per-op overrides still apply. Both confs and the nvidia-*-cu12 runtime deps (CUDA builds) are now packaged. Verified on a clean shell (no wrapper, auto-preload): cuda-path ops 330 passed/45 skipped/3 xpassed; FLAGOS_USE_FLAGGEMS=1 flaggems_python 27 passed; all three switch levels confirmed. * feat(flaggems): route random in-place + *_like factory ops, +6 (301 -> 307) Add two codegen paths to the FlagGems Python bridge: - *_like factory (zeros_like/ones_like/full_like): new CallPythonOp_LikeFactory injects device=flagos/layout=strided/memory_format=None/pin_memory=None and forwards dtype; gems reads shape/device from the source tensor positional. - random in-place (uniform_/exponential_/bernoulli_.float): new CallPythonOp_RandomInplace injects a module-level CUDA torch.Generator as the generator kwarg. gems only reads philox seed+offset from it (randoms computed in the Triton kernel writing into the flagos tensor), so the generator device need not match. A CUDA generator is required: a flagos/CPU generator's 5056-byte MT19937 state fails gems' 16-byte CUDA-philox state unpack. Serialized by the GIL; offset advances per call via gems' set_state. discover_flaggems_ops gains like_factory / random_inplace branches; random ops are whitelisted (_FLAGGEMS_RANDOM_INPLACE) since normal_ is also inplace+Generator? but hardcodes generator=None internally and can't be routed. Explicitly skipped with reasons: normal_/normal.* (gems drops the generator), rand/randn/randperm/rand_like/randn_like (no generator param -> empty PrivateUse1 default_generators), multinomial (name-mismatch + generator). Also add docs/flaggems_no_dispatcher_analysis.md: the 88 no-dispatcher ops all already execute correctly (composite_implicit decomposition / fallback / manual registration), so they are not a functional gap and bulk routing would drop autograd. Validated: like ops err=0 incl dtype override; random ops distribution-correct with distinct streams across calls. Regression unchanged (cuda 330/45/3, flaggems_python 27). * feat(metax): auto-symlink stock torch/lib to maca C++ runtime at import Reuse PyTorch's CUDA boxing kernels on MetaX with a stock torch==X.Y.Z+cpu wheel by symlinking the active wheel's torch/lib core .so to the MetaX torch wheel's copies, so the process loads the MetaX C++ runtime (at::maca::* fork) instead of the upstream one. A runtime hook in torch_fl/__init__.py runs ensure_maca_libtorch_links() BEFORE import torch (afterwards libc10 is already mapped and relinking is too late). Pure ctypes preloading does not work: the official _C.so / libtorch_python.so carry an $ORIGIN RUNPATH that pulls upstream libc10 back in by full path, double-loading it and crashing on duplicate caffe2 static init. Symlinking makes the physical files RUNPATH resolves to be the MetaX ones. Gated on FLAGOS_METAX_BOXING=1; idempotent; backs up originals to torch/lib/_orig_backup/ (restore_original_libtorch() reverts); no-op when the active torch already IS the MetaX wheel. MetaX torch/lib discovered via FLAGOS_MACA_TORCH_LIB env or conda-env scan. Verified on torch 2.10.0+cpu: full_cuda_coverage 45/46 (the 1 fail is addmm TF32 rounding), per-op suite 310 passed / 45 skipped / 3 xpassed (8 fails are all flaggems-backend or test-case issues, not op wiring). * chore(codegen): regen generated ops for torch 2.10 schema Drop cudnn_convolution_bias_fused (2.11-only op absent in 2.10). Regen matches the 2.10 aten schema; flaggems_python and CUDA paths verified against the 2.11 baseline (27 passed; 330 passed/45 skipped/3 xpassed). * feat(metax): bundle forked libtorch into self-contained wheel Package the MetaX-forked libtorch C++ .so (~1.1G) inside the wheel under torch_fl/lib_maca/ so target machines need only the official torch+cpu wheel plus this wheel plus the /opt/maca driver runtime -- no separate MetaX torch wheel required. - setup.py: package_data bundles lib_maca/*.so* and all backends*.conf (boxing modes select backends_cuda.conf via FLAGOS_BACKEND_CONFIG); version gains a +metax local segment (FLAGOS_WHEEL_LOCAL overridable). - pyproject.toml: mark version dynamic so setup.py's computed +metax tag is not overridden by a static [project] version. - _metax_libtorch_link.py: _discover_maca_torch_lib prefers the bundled lib_maca/ over env var / conda scan; symlinks the stock wheel's torch/lib to the bundled forked libtorch at import. - scripts/bundle_maca_libtorch.sh: copy the 8 forked libtorch .so and patchelf their RPATH to $ORIGIN + /opt/maca for the target runtime. - .gitignore: ignore torch_fl/lib_maca/ (build artifact). Verified end-to-end from a fresh wheel install in a clean env (official torch 2.10.0+cpu, no MetaX torch wheel, no LD_LIBRARY_PATH): flagos compute on MetaX GPU works, torch/lib auto-symlinks to the installed lib_maca, libmcblas loads from /opt/maca; op coverage 45/46 (only addmm TF32 rounding fails). * docs(metax): document self-contained boxing wheel packaging and usage Add a MetaX Self-Contained Wheel (CUDA boxing) section covering the FLAGOS_METAX_BOXING=1 path: how to build the wheel (bdist_wheel + bundle_maca_libtorch.sh + repackage), the ~1.1G size / distribution tradeoff, and how to install and run on a clean target (official torch+cpu + this wheel + /opt/maca, no torch+metax wheel, no manual LD_LIBRARY_PATH). Note the two MetaX build modes in the runtime notes and warn that FLAGOS_USE_FLAGGEMS=1 must not be used with the boxing wheel (flagos_python backend is not compiled -> backend not registered). * docs(metax): add MetaX developer portal link for SDK and torch+metax wheel Point readers to https://developer.metax-tech.com/softnova (SoftNova) for the MACA SDK (driver + cu-bridge + mxcc/cucc) and the torch+metax wheel, noting login is required and versions must match the card/driver/Python. Referenced from the top-level prerequisites, the MetaX source-build prerequisites, and the boxing-wheel build step. * docs(metax): drop mxcc source-build path, keep only the boxing wheel We no longer use the hand-written mxcc/cucc kernel build (METAX_KERNEL=ON + torch+metax + Triton). Remove that 'Build from Source (MetaX Platform)' section and promote the self-contained CUDA boxing wheel to be the single MetaX build path (renamed to 'Build from Source (MetaX Platform)'). - Fold the MetaX developer-portal (SoftNova) SDK / torch+metax wheel download note into the boxing build steps (needed to build, not run). - Rewrite the 'Two build modes' runtime note to describe only the boxing wheel and fix its now-stale section anchor. * feat(flaggems): route varargs unary in-place + rng ops, +13 (307 -> 320) varargs unary in-place (7): asinh_/sinh_/log1p_/digamma_/sgn_/hardswish_/ logit_. gems wraps these as (*args, **kwargs) so inspect.signature can't recover arity; add _FLAGGEMS_ARITY_OVERRIDE so the aten schema supplies the authoritative positional count. Whitelist only holds simple elementwise ops verified to run + match CPU (maxdiff <= 1e-6) and drop no kwarg. rng (6): rand/randn (factory), rand_like/randn_like (like_factory), randperm (factory; no generator arg in this torch schema), multinomial (new rng_dropgen category dropping the trailing Generator?). Unblocked by _patch_flaggems_philox() in torch_fl/__init__.py, which monkeypatches gems' philox_backend_seed_offset to fall back to a held CUDA generator when torch.cuda.default_generators is empty (CPU-torch + cuda shim). One patch covers all 6 rng ops; no caller C++ change needed. Excluded: i0_/zero/zero.out hit a hardcoded tensor.is_cuda assert in the gems kernel (flagos is PrivateUse1, never true) so they stay in FLAGGEMS_PYTHON_SKIP; normal_/normal.* hardcode generator=None upstream (can't thread our generator). Verified: 15/15 numeric spot-checks pass; regressions unchanged from baseline (cuda-path 330 passed/45 skipped/3 xpassed, flaggems_python 27 passed). Fresh regen also drops the now-stale cudnn_convolution_bias_fused[.out] from backends_cuda.conf (absent from the torch 2.10 schema). * feat(ascend): enable high-level F.scaled_dot_product_attention Route the high-level SDPA API to the aclnnFlashAttentionScore kernel instead of the math decomposition path, and register the view ops its pre/post-processing needs. - Register _fused_sdp_choice_stub DispatchStub for PrivateUse1 returning efficient_attention (2). PyTorch's scaled_dot_product_attention selects its fused backend via this C++ stub (gated by is_device_supported), not the aten-op-level _fused_sdp_choice. Macro must live in namespace at::native. - Add view/metadata ops for kAscend: transpose.int, permute, select.int, slice.Tensor, squeeze, squeeze.dim, unsqueeze, _unsafe_view, detach. Implemented via at::native::<fn> (not tensor member methods, which re-dispatch through PrivateUse1 and recurse -> segfault). select.int uses select_symint; explicit _native.h includes avoid int64->Dimname overload mis-resolution. Verified end-to-end on Ascend 910 NPU: forward, causal, autograd backward all pass; CPU-reference relative error ~0.0004 (fp16). * feat(ascend): codegen in-place zero_/fill_ via aclnn, +3 (139 -> 142) Add aclnn codegen for the in-place fill primitives: - zero_ -> aclnnInplaceZero - fill_.Scalar -> aclnnInplaceFillScalar - fill_.Tensor -> aclnnInplaceFillTensor Before this, zero_/fill_ had no device implementation, so the handwritten factory ops (zeros/ones_like/new_ones/scalar_tensor) fell back to flaggems/CPU for their internal .zero_()/.fill_() calls -- a hidden h2d path. Now the whole factory chain runs device-side aclnn. The factory ops themselves stay handwritten (their device/dtype inference is not expressible by codegen), but the fill work is pushed down to aclnn. Also move the SDPA + view-op conf entries above the codegen marker: they were appended after the "# --- generated by codegen_ascend.py ---" marker, which codegen truncates on every run, so a regen would silently drop them. Handwritten conf entries must live before the marker. Verified on Ascend 910: zero_/fill_.Scalar/fill_.Tensor and all four factory ops match CPU exactly (diff 0). SDPA + view-op tests still pass. * feat(ascend): codegen embedding/embedding_backward/constant_pad_nd, +3 (142 -> 145) Migrate three single-aclnn-call kernels from handwritten to codegen: - embedding -> aclnnEmbedding - embedding_dense_backward -> aclnnEmbeddingDenseBackward - constant_pad_nd -> aclnnConstantPadNd Each was a straight aclnn call with deterministic output-shape logic, so the handwritten .cc bodies map verbatim into codegen templates. Delete the three .cc files (globbed by CMake, no explicit list to update) and move their conf entries from the handwritten section into the generated block. Handwritten kAscend registrations drop 30 -> 27; codegen 142 -> 145. Verified on Ascend 910: all three match CPU exactly (diff 0), including embedding backward via autograd and multi-dim constant padding. Prior in-place/factory and view-op tests still pass. * feat(ascend): migrate mm/bmm/cat + factory ops from handwritten to codegen Extend codegen_ascend.py with three new structural capabilities so more handwritten kAscend kernels can be expressed as codegen templates: - .out variants: T_MATMUL / T_MATMUL_OUT generate functional + .out pairs (mm/mm.out via aclnnMm, bmm/bmm.out via aclnnBatchMatMul). Reusable pattern for any op whose .out kernel writes into a caller-shaped out&. - TensorList: T_CAT marshals at::ITensorListRef via aclCreateTensorList (aclnnCat), filtering numel==0 tensors; does not aclDestroyTensorList. - factory ops: T_ZEROS/T_SCALAR_TENSOR/T_ONES_LIKE/T_NEW_ONES build TensorOptions on-host + at::empty then fill via device-side zero_/fill_. New NO_ACLNN_CATEGORIES set skips the libopapi symbol guard for kernels that issue no direct aclnn call. Deletes 7 handwritten .cc (mm/bmm/cat/zeros/scalar_tensor/ones_like/new_ones). Handwritten kAscend regs 16 -> 7, codegen 145 -> 154, total unchanged. Verified vs CPU on NPU (test_matmul_codegen.py, test_cat_codegen.py, test_inplace_fill.py): all pass, matmul f32 within hf32-cube tolerance. * feat(metax): enable FlagGems on boxing wheel as runtime switch (FLAGOS_USE_FLAGGEMS) Match the CUDA single-wheel model on MetaX: compile the FlagGems Python path (flagos_python backend) alongside the CUDA boxing kernels and select between them purely at runtime via FLAGOS_USE_FLAGGEMS, instead of forcing it off. - setup.py: drop the metax FLAGGEMS_PYTHON=OFF force; default ON like CUDA (only the C++ FlagGems path FLAGGEMS_KERNEL stays off). FLAGGEMS_PYTHON=0 still available for a slim pure-boxing build. - torch_fl/__init__.py: add explicit MetaX branch in _patch_flaggems_codegen_config (GEMS_VENDOR=metax + patch_torch_cuda_for_metax) before the ascend fallback, fixing metax wrongly getting GEMS_VENDOR=ascend; _select_backend_config picks backends_metax_flaggems.conf when FLAGOS_USE_FLAGGEMS=1 + FLAGOS_METAX_BOXING=1. - _metax_compat.py: add stream/availability/manual_seed shims and _patch_triton_do_bench (wall-clock) so FlagGems Triton kernels run on the CPU-frozen torch wheel against maca libtorch_cuda.so. - scripts/codegen_ops.py: generate backends_metax_flaggems.conf, routing the ops triton-metax cannot run (mm/bmm/mean.dim) and flag_gems device-guarded ops (mul, embedding_dense_backward, etc.) back to cuda boxing. - tests/integration/conftest.py: skip forcing backends_metax.conf in boxing mode (mxcc backend not compiled) so torch_fl's own config selection applies. - tests/integration/ops/conftest.py: skip @mark.metax tests in boxing mode (no metax backend to dispatch to). Verified: 255 passed / 147 skipped; only non-skip failure is the pre-existing out-of-scope cat empty-1d-tensor boxing bug. * feat(ascend): enable Qwen3-0.6B generate() end-to-end on real 910 Walks the transformers generate() path op-by-op until inference runs coherently (~1.8 tok/s) on a pure-aclnn C++ backend, and takes the ops parity suite from 269 failed/31 passed to 6 failed/294 passed. Bespoke handwritten kernels (generation path): - topk/sort/scatter/multinomial via aclnnTopk/Sort/Scatter/Multinomial - arange/argmax/isin/lift_fresh - rng.cc: device randn/rand/randint/randint.low via aclnnInplace Normal/Uniform/Random, seeded from the default CPU generator's random64() so the ops suite's on-device torch.randn(device=DEVICE) inputs work. Codegen (scripts/codegen_ascend.py): - rsub.Scalar -> binary_scalar_alpha (aclnnRsubs) - sum/max/min full-reduce -> reduce_sum_all / reduce_minmax_all - any -> reduce_all; ones/empty_like/full/full_like factory ops SDPA GQA fix (scaled_dot_product_attention.cc): expand kv heads using kv's own S_kv, not query's S -- during decode query S==1 but kv S==full context, so using query S gave an expand size mismatch. View ops (strided_ops): t, unbind.int (pure metadata, via at::native::). Runtime config: torch_fl auto-selects the ascend conf on a /dev/davinci* box; FLAGOS_USE_FLAGGEMS=1 opts into the FlagGems Triton path. Also patch triton-ascend npu_utils.cpp for the CANN 9.0.0 rtLimitType_t enum name. * perf(ascend): reach 0.89x torch_npu on Qwen3-0.6B inference Host-side dispatch optimizations for eager decode on real 910, keeping the op set identical to torch_npu (no fusion in the measured path). - aten::empty fast path: skip the DeviceGuard registry round-trip when the device is unchanged; drop the ptr_to_block_ side map + its mutex from the caching allocator by stashing Block* in the DataPtr context (5.25 -> 2.04 us/call, the single biggest win) - repeatable aclOpExecutor cache in op_api_common.h (ExecAscendCached): owns its aclTensors, rebinds addresses on hit, reuses the workspace tensor - codegen: cached categories for elementwise/unary-scalar/reductions/softmax plus a CPU-scalar fast path routing T+float through aclnnAdds/Muls/... to avoid a per-call H2D copy - register aten::matmul on AutogradPrivateUse1 -> aclnnMatmul, collapsing mm/bmm/view churn and matching torch_npu's op counts exactly - isin: compute on device instead of a triple CPU round-trip (this was the hottest op in the generate() loop) - _to_copy: cache the aclnnCast executor Inference 13.4 -> 24.82 tok/s (0.89x torch_npu); training 507.9 tok/s (0.69x). Adds tests/perf/e2e_qwen3_{infer,train}_ascend.py as the comparison harness. * fix(ascend): do not declare a triton runtime dep on ACCELERATOR=ascend Upstream #22 made flag_gems + triton>=3.5.1 hard runtime deps for every accelerator except dcu. On Ascend the `triton` module is supplied by triton-ascend, which is installed out of band and has no PyPI release satisfying triton>=3.5.1, so `pip install -e .` pulls stock triton over it. Every Triton entry point then fails with "0 active drivers" and the torch_npu-shim patch from scripts/patch_triton_ascend.py is gone. Ascend needs the same carve-out as dcu, for the same reason: the accelerator provides its own triton and PyPI's NVIDIA-targeted wheel is the wrong artifact. flag_gems imports in the Python layer are ImportError-guarded, so omitting the dep is safe. * feat(ascend): _foreach_* kernels for AdamW foreach=True, chunked under aclnn's 50-entry cap Adds the 7 _foreach_* Ascend kernels the AdamW foreach path needs (_foreach_mul_/add_.Scalar, _foreach_lerp_.Scalar, _foreach_addcmul_.Scalar, _foreach_sqrt, _foreach_div_/addcdiv_.ScalarList), plus stack, mean, clamp, clamp.Tensor and bitwise_and_/or_/xor_.Tensor. CANN's aclnnForeach* kernels only process the FIRST 50 entries of an aclTensorList. The ScalarList variants at least error past that (561002/161002); Mul/Add/Addcmul/Lerp/Sqrt instead return success and leave entries >= 50 UNTOUCHED, so the bug is silent -- _foreach_lerp_ over 200 tensors "succeeded" with 150 of them never written. The cap is on the entry count alone: measured identical for numel 8..65536 and fp16/fp32/bf16. AdamW passes 310 tensors for Qwen3-0.6B, so each foreach template is split into a <Kernel>Chunk doing the aclnn call and a wrapper that slices the lists into sub-50 chunks. TensorList and ArrayRef<Scalar> are both ArrayRefs, so the slicing is free, and elementwise semantics make it exact. Note for anyone probing other aclnn list limits: an exception-based search reports "no limit up to 4096" for the silent ops. The values have to be compared against CPU entry by entry. aclnnForeachAddcdivScalarList's scalars param is a device aclTensor whose dtype must match self (float32 scalars against fp16 inputs return 161002), so fp16 addcdiv lands a few ulp off CPU, which keeps the divisor at full precision. AdamW holds optimizer state in fp32, so this does not reach the training path. tests/integration/ops/test_foreach_dispatch.py covers all 7 ops at lengths that straddle the boundary (51/60/128/310) and asserts on every entry, plus AdamW(foreach=True) against foreach=False. Verified it fails (85 cases) when the chunk size is raised to 128. e2e_qwen3_train_ascend.py no longer hard-pins foreach=False; it takes --foreach/--no-foreach, applied to both backends. Measured on 910 (Qwen3-0.6B, seq 128, batch 1, eager, card 10): training 621.0 tok/s vs torch_npu 827.1 (0.75x, up from 0.69x at foreach=False's 507.9), loss curve unchanged at 2.937 -> 0.324. Inference unaffected at 24.5 tok/s. 169 foreach cases pass; ops/ suite and the qwen3 infer/train integration tests show no regressions. * feat(ascend): fused matmul in training via AutogradPrivateUse1 codegen aten::matmul is CompositeImplicitAutograd, so claiming a fused aclnnMatmul kernel on PrivateUse1 stops the mm/bmm/view decomposition -- and with it the sub-op graph autograd was relying on. The op then binds its real derivative, aten::matmul_backward, which the backend must supply. Until now that was sidestepped by taking the fused path only when !requires_grad: inference got one aclnnMatmul, training kept the decomposition. Close the gap the way torch_npu does, by generating the missing autograd layer rather than hand-rolling it. scripts/codegen_autograd.py drives torchgen's own emit_body() -- the generator behind PyTorch's in-tree VariableType_N.cpp -- to produce a VariableType::matmul on AutogradPrivateUse1 that builds MatmulBackward0 and redispatches to the fused kernel. Only that thin layer is generated: the backward node classes already ship in libtorch, so unlike torch_npu we do not regenerate Functions.h/ADInplaceOrView/python bindings. Adding an op is one entry in AUTOGRAD_OPS. matmul_backward itself is implemented with two cached aclnnMatmul calls. op-plugin's MatmulBackwardKernelNpuOpApi.cpp was the starting reference and has two real bugs, deliberately not reproduced: * 2-D x N-D reshapes grad to {M, -1}. mat2^T flattens to (B*N, K), so its row index is the pair (b, n); grad must carry the same pair as its column index, which needs M permuted to the front first. The plain reshape pairs (m, n) against (b, n) and silently mixes batches. * Only *leading* singleton batch dims are stripped, so any interior or trailing broadcast returns a wrong-shaped gradient -- (2,1,3,4) x (2,5,4,6) yields a (2,5,3,4) grad for a (2,1,3,4) input. Replaced with at::sum_to onto the promoted operand shape, a no-op when nothing was broadcast. Guarded by USE_ASCEND throughout: other backends have no fused kernel, keep PyTorch's decomposition, and never bind a matmul_backward they cannot service. Qwen3-0.6B training, batch 1 x seq 128, on a real 910: backward dispatches 5466 -> 4144 (torch_npu 4002) matmul_backward n/a -> 253 (torch_npu 253) backward self-CPU 117.1 -> 109.9 ms/step throughput 566 -> 674 tok/s, 0.69x -> 0.82x torch_npu Verified against a float64 CPU reference -- the kernel runs hf32 cube math, so absolute fp32 comparison misreads ~1e-4 precision as a correctness bug. The new test carries 30 cases and every shape in it catches a real defect in the op-plugin rules on one side or the other. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * fix: guard Ascend-only symbols so non-Ascend builds link and load CUDA CI failed at `import torch_fl` with ImportError: libtorch_fl.so: undefined symbol: _ZN2at6native6flagos18MatmulKernelAscendERKNS_6TensorES4_ Two separate instances of the same mistake, both introduced by this branch. 1. register.cc referenced MatmulKernelAscend behind a *runtime* check (GetBackendForOp("matmul") == kAscend) with only the forward declaration guarded asymmetrically -- the backward decl was inside #if defined(USE_ASCEND), the forward one was not. A runtime branch does not remove a link-time reference, and backends/ascend/matmul.cc is not compiled without USE_ASCEND. A shared library links fine with undefined symbols and only fails at dlopen, which is why "Build wheel (CUDA)" passed and the failure surfaced at import. Fix: guard both the declarations and the call site at compile time. 2. copy_ops.cc / contiguous_ops.cc call ascend::StridedCopy and ascend::DtypeCast from #else branches that cover TsingMicro, GCU and MUSA-without-mudnn as well as Ascend, but included ascend_copy.h only under #ifdef USE_ASCEND. Those platforms failed to *compile* ("'ascend' has not been declared"); no CI builds them, so it stayed hidden. Fix: ascend_copy.h now supplies inline no-op fallbacks for non-Ascend builds (defined, not just declared, so nothing is left undefined at load), and is included unconditionally. The no-ops report "unavailable" and callers take the CPU round-trip they already implement. Verified: register.cc compiled without USE_ASCEND has no undefined MatmulKernelAscend reference (nm -u); copy_ops.cc and contiguous_ops.cc compile clean under each of USE_TSINGMICRO / USE_GCU / USE_MUSA and with no macro at all. On real 910: 30/30 matmul-backward, 175/175 foreach + conf-consistency, 3/3 Qwen3 training. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> --------- Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
…ai#56) The C++ FlagGems path calls flag_gems' C++ entry points in liboperators.so, JIT-compiling and launching Triton kernels without touching Python or the GIL. It was CUDA-only on the torch_fl side because CMakeLists forced FLAGGEMS_KERNEL=OFF for ACCELERATOR=metax. FlagGems itself already supports MetaX (cpp/ -DFLAGGEMS_BACKEND=MACA), and its kernels reach the device through the same DeviceBoxingGuard the metax boxing path uses, so no kernel code needed changing -- only the build gate and the routing table. - CMakeLists.txt: metax no longer force-disables FLAGGEMS_KERNEL. It stays OFF by default (a plain metax build needs no FlagGems) and turns on only in boxing mode when the caller explicitly asks, via FLAGGEMS_KERNEL=1 FLAGGEMS_DIR=<MACA FlagGems build>. - csrc/CMakeLists.txt: fix an RPATH bug. The metax branch set INSTALL_RPATH wholesale, discarding the liboperators.so dir the top level appended to CMAKE_INSTALL_RPATH, so ldd reported "liboperators.so => not found". Re-add it when FLAGGEMS_KERNEL is on. - setup.py: document the opt-in; the default -DFLAGGEMS_KERNEL=OFF is unchanged and the generic env pass-through emits the overriding -D. - torch_fl/__init__.py: CPP + boxing selects backends_metax_flaggems_cpp.conf. - new backends_metax_flaggems_cpp.conf: 17 flagos (C++) + 362 flagos_python + 1654 cuda. C++ ops keep flagos; non-C++ ops inherit the metax Python routing, so the existing metax fallbacks (bmm-via-SPLIT_K, mean.dim, slice_backward) still apply. mm and mm.out route to cuda boxing (mcblas) instead of the C++ kernel: this is a hardware limit, not an integration gap. flag_gems' mm_kernel_general requests 98304 bytes of shared memory and C550 provides 65536, so mcModuleLaunchKernel returns mcErrorInvalidValue. The other 17 C++ ops pass. Verified on real hardware (C550, metax 3.8.1, torch 2.10.0): tests/manual/metax/test_flaggems_cpp_metax.py -- 24/24 ALL PASS. Each op is checked for both routing (FLAGOS_LOG_DISPATCH shows "-> flagos") and numerics against CPU; routing alone does not prove the kernel ran, since the log is emitted before launch. Default path (CPP off) still selects backends_cuda.conf.
… layout fixes (flagos-ai#51) * perf(copy): fast-path _to_copy via direct CUDA redispatch on CUDA/MetaX _to_copy is one of the hottest ops in Qwen decode (FP32<->FP16 casts in RMSNorm and attention). The flagos(PrivateUse1)->CUDA branch allocated an intermediate contiguous tensor, memcpy'd device-to-device, then ran a separate .to(dtype) cast -- an extra allocation and an extra pass. Since flagos and CUDA share the same GPU memory, box self to CUDA and redispatch straight to the native CUDA _to_copy with an explicit CUDA DispatchKeySet (at::_ops::_to_copy::redispatch). That skips re-entering the dispatcher from the top (no chance of routing back to PrivateUse1) and lets the native kernel read self's strides + cast dtype in one on-device pass, allocating the result directly on CUDA. self is unboxed back to PrivateUse1 on DeviceBoxingGuard teardown, matching the boxing pattern used by the generated CUDA kernels (e.g. PrivToCopyOutKernelCuda). Guarded so only platforms with a CUDA runtime take the fast path; USE_ASCEND/TSINGMICRO/GCU/MUSA keep the original contiguous + Memcpy path. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * perf(boxing): track boxed tensors by raw TensorImpl pointer DeviceBoxingGuard and TensorListBoxingGuard recorded each boxed tensor in a std::vector<at::Tensor>, so every box/unbox paid an intrusive refcount atomic (fetch_add on push, fetch_sub on teardown) per tensor. Qwen decode boxes/unboxes tens of tensors per token, so this stacks up on the hot path. Record raw c10::TensorImpl* instead and unbox by calling the new SetTensorImplDevice directly (extracted from SetTensorDevice), skipping the owning-Tensor round trip entirely -- zero refcount atomics on box or unbox. - SmallVector<TensorImpl*, 4>: the common case (<=4 tensor operands) stays on the stack, no heap allocation. - The guard now takes a forwarding pack with a static_assert that every argument is an lvalue reference, rejecting temporary (rvalue) tensor handles at compile time so a recorded impl can never dangle. - CPU-scalar / genuine-CUDA / undefined-tensor skipping and exception-safe restore semantics are unchanged; unboxing still routes through PyTorch's _change_backend_component_keys rather than poking dispatch-key bits directly. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * fix(flagos): register unfold view op to fix repeat data corruption Tensor.repeat() calls aten::unfold internally to build a strided view before copy_-ing into it. flagos had no PrivateUse1 impl for the plain unfold view op (only unfold_backward / unfold_copy.out existed). Because view ops cannot fall back to CPU -- storage can't be shared across devices -- PyTorch emitted a warning and returned an uninitialized shell tensor. repeat() then read garbage. GPTJ's rotary embedding uses repeat() to build gather indices; corrupted indices led to out-of-bounds gather -> illegal memory access -> crash. unfold is pure stride computation (no GPU kernel): delegate to at::native::unfold, which computes the new size/stride and calls as_strided (already registered for flagos). Registered via direct m.impl in the always-on block so it applies to every backend. Verified on MetaX (t210-box): unfold/repeat now numerically correct, the view-fallback warning is gone, and a tiny GPTJ forward runs cleanly (cosine 1.0 vs CPU, rel err 0.05% -- fp32 drift, not corruption). Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> * fix(flagos): normalize group norm input before boxing Diffusers UNet2DModel feeds GroupNorm channels-last (NHWC) or strided (strides>1) inputs. The backend native_group_norm CUDA kernel requires standard-contiguous layout (is_contiguous() == true); the native CUDA path normalizes the layout before the kernel, but the flagos boxing wrapper passed `input` straight through, raising: Expected X.is_contiguous(memory_format) to be true Add a generic _CONTIGUOUS_TENSOR_ARGS_BY_OP map in codegen naming, per op, which Tensor args must be forced contiguous before boxing. The generated kernel now emits: at::Tensor input_contiguous = input.is_contiguous() ? input : input.contiguous(); DeviceBoxingGuard guard(input_contiguous, weight_t, bias_t); auto result = at::native_group_norm(input_contiguous, weight, bias, ...); Design points: - The contiguous() must run before DeviceBoxingGuard: boxing rewrites the tensor's device metadata, so the copy has to happen while it still lives on flagos. - is_contiguous() short-circuits to a no-op when already contiguous, so the only cost is on the exact non-contiguous inputs that would crash. - Wired into gen_functional_pure / gen_out_variant / gen_tuple_return but currently applied only to native_group_norm.input -- precise, scoped. Verified on MetaX (t210-box): channels_last / strided GroupNorm forward + backward now numerically match CPU (max diff ~2e-7), a Diffusers-style channels_last Conv->GroupNorm->SiLU ResnetBlock runs cleanly (cosine 0.99999988), and sort/layernorm/topk (same generators) are unaffected. codegen unit checks pass, including the contiguous-before-guard ordering. Co-Authored-By: Claude Opus 4.8 <noreply@anthropic.com> --------- Co-authored-by: zhangliheng <zhangliheng@chitu.ai> Co-authored-by: Claude Opus 4.8 <noreply@anthropic.com>
…e, kernel metadata (flagos-ai#55) * docs: PrivateUse1 profiler support design (Stage A ProfilerStubs + Stage B CUPTI child profiler) * docs: PrivateUse1 profiler implementation plan (Stage A + Stage B, 8 tasks) * docs: fix backend config path (torch_fl/configs/backends_cuda.conf) * feat(profiler): guard.h carries real CUDA stream for event attribution * feat(profiler): Stage A FlagosProfilerStubs implementation (blocked by torch 2.11 CPU wheel) Implements torch::profiler::impl::ProfilerStubs subclass for flagos backend: - record(): creates CUDA event with timing flag, records on current stream - elapsed(): synchronizes events and returns elapsed time in microseconds - onEachDevice(): iterates all flagos devices - synchronize(): calls flagos DeviceSynchronize - Static initializer registers stubs via registerPrivateUse1Methods() Test added: test_stage_a_privateuse1_device_time checks for non-zero device time. BLOCKER: torch 2.11.0+cpu wheel's Kineto library does not invoke the PRIVATEUSE1_FALLBACK path at runtime. Stubs register successfully (enabled=true) but record()/elapsed() are never called during profiling. Device times remain 0. Evidence: - ProfilerState::KINETO_PRIVATEUSE1_FALLBACK enum exists - registerPrivateUse1Methods() succeeds, privateuse1Stubs() returns valid ptr - Instrumentation confirms record() never invoked during profile() context - CPU wheel likely omits PrivateUse1 fallback implementation (compile-time gate) Next: try torch CUDA wheel, file upstream issue, or pivot to Stage B (NVTX). See task-2-report.md for full diagnosis and unblock options. * feat(profiler): Stage B CUPTI dlopen shim + cmake include path - csrc/profiler/cupti_shim.h: header-only dlopen wrapper for CUPTI Activity API * Forward-declares CUPTI types (CUptiResult, CUpti_ActivityKind, callbacks) to avoid including cupti headers * Loads 8 Activity API functions at runtime (Enable/Disable/RegisterCallbacks/FlushAll/GetNextRecord/GetNumDroppedRecords/Push/Pop) * Tries libcupti.so in priority order: system CUDA 13.0 → .so.13 → .so.12 → generic * Exposes CuptiShim::get().available() + function pointers for Task 5 consumer - csrc/CMakeLists.txt: add CUPTI include path discovery (find_path cupti_activity.h) + FLAGOS_HAVE_CUPTI macro - tests/unit/test_profiler_privateuse1.py: add test_cupti_library_locatable() to verify runtime CUPTI availability Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * feat(profiler): add CUPTI Activity API profiler (Stage B Task 5) Implement libkineto::IActivityProfiler subclass using CUPTI Activity API to collect GPU kernel traces. Register with kineto so torch.profiler automatically captures CUDA kernel timelines for flagos backend. Classes: - FlagosCuptiProfilerSession: manages CUPTI lifecycle (enable/disable/flush) - FlagosCuptiProfiler: factory registered with libkineto::api() CUPTI buffer callbacks parse kernel/memcpy activity records and convert to libkineto::GenericTraceActivity format. LIMITATION: CUPTI Activity API requires initialization before CUDA driver is loaded. In CPU wheel + dlopen'd libtorch_cuda.so architecture, CUDA initializes before our module loads, so CUPTI silently no-ops (callbacks never invoked). Implementation is correct for standard CUDA builds but cannot capture activities in current environment. GPU profiling still works via existing FlagosProfilerStubs (event-based timing) and NVTX (external tools). Files: - csrc/profiler/flagos_cupti_profiler.h (98 lines): class declarations - csrc/profiler/flagos_cupti_profiler.cc (307 lines): implementation - tests/unit/test_profiler_privateuse1.py: updated test Task: torch-fl profiler support - Stage B Task 5 * feat(profiler): Stage B correlation bridge - wire CUPTI push/pop for CPU→GPU linking Implements Task 6: correlation ID bridge for linking CPU ops to GPU kernels. Changes: - Add instrumentation counters (g_correlation_push_count/pop_count) to track push/pop invocations in FlagosCuptiProfilerSession methods - Export C API (flagos_cupti_get_correlation_push_count, _pop_count, _reset_correlation_counters) with visibility("default") for Python testing - Add test_stage_b_correlation_or_degrade that verifies push/pop are called during profiling (3/3 times for single matmul op) KNOWN LIMITATION (degraded acceptance): In the CPU-wheel environment, CUPTI buffer callbacks are never invoked (CUDA initializes before CUPTI registers), so GPU activities are not captured and correlation IDs don't actually link anything. This commit proves CODE CORRECTNESS (push/pop methods call CUPTI functions correctly) rather than functional correlation (which is impossible without a torch+cuda wheel or custom torch build with working CUPTI Activity API). Test result: PASS (degraded mode - push/pop called, 0 GPU activities captured) Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * chore: Task 7 validation + add DeviceGuardImpl::getDeviceCapability Validation results: - qwen3_infer: 4/4 PASS - ops regression: 323 passed, 20 failed (matches Task 1 baseline exactly) - qwen3_train: 3 errors (pre-existing autograd issue, not a regression) Also includes guard.h change from Task 6 that wasn't committed: - Add getDeviceCapability() override to FlagosGuardImpl - Returns default DeviceCapability struct (all scalar types enabled) - Required by torch 2.11's autograd engine Known limitation: qwen3_train fails with 'Backend doesn't support getting device capabilities' error. This is a pre-existing issue documented in Task 1 baseline (20 ops tests fail with same error). The getDeviceCapability implementation is correct but not being invoked by autograd, likely due to vtable/initialization timing issues in the CPU-wheel + external-CUDA-lib architecture. Co-Authored-By: Claude Opus 4.8 (1M context) <noreply@anthropic.com> * fix(profiler): make CUPTI actually capture the GPU kernel timeline Stage B previously registered a CUPTI child profiler that captured 0 activities. Root-caused (see standalone probes) to three real bugs, all fixable within the CPU-torch + external libtorch_cuda.so architecture (no CUDA wheel needed): 1. Version mismatch: CuptiShim dlopen'd the system CUDA-13.0 libcupti first, but the process runs the cu12.8 runtime (libtorch_cuda -> NEEDED libcudart.so.12; _preload_cuda_assets loads the pip cu12 libcupti.so.12). Arming the wrong libcupti no-ops. Now bind the already-loaded copy via dlopen(NULL)+dlsym, then prefer libcupti.so.12. 2. Wrong enum value: CUPTI_ACTIVITY_KIND_CONCURRENT_KERNEL is 10 in cu12, not 9. The wrong value enabled/matched the wrong activity kind. 3. Wrong record layout + self-deadlock: - The activity-record structs were hand-guessed (natural alignment, name at a fixed offset) and decoded garbage (empty names, dur=0). Replaced with byte-accurate __packed__ mirrors of the cu12 CUpti_ActivityKernel9 / CUpti_ActivityMemcpy layouts; kernel name is a const char* field. CMake now prefers the version-matched cu12 cupti_activity.h. - start() held g_session_mutex across cuptiActivityFlushAll(), which synchronously re-enters bufferCompleted (same non-recursive mutex) -> hang in torch.profiler start_trace. Scope the lock to publishing the session only; flush after release. Also arm CUPTI (RegisterCallbacks + ActivityEnable) at module-load time (before the first CUDA context) instead of only in session start(), and gate all diagnostic logging behind FLAGOS_CUPTI_SHIM_DEBUG. Result: torch.profiler now captures the real CUDA kernel timeline with correct names and microsecond durations. Unit test asserts ~N sgemm kernels for N matmuls; qwen3-0.6B inference captures 2261 named GPU kernels. Ops regression unchanged (323 passed / 20 pre-existing fails). * profiler: decouple CUPTI binding from any one CUDA version The shim previously reached for a hardcoded /usr/local/cuda-13.0 path and a cu12-biased soname ordering, and the record decoder trusted a hand-mirrored CUpti_ActivityKernel9 layout unconditionally. Both tie the profiler to one CUPTI release, and the second fails silently: a layout mismatch yields a trace of empty names and zero durations rather than an error. Binding: an explicit FLAGOS_CUPTI_LIBRARY wins first, then whatever libcupti is already loaded in the process (by construction the copy the running CUDA runtime pulled in), then a soname fallback ordering. The override has to outrank the already-loaded copy -- the case that motivates setting it is a preloaded CUPTI we cannot decode, so an override checked only when nothing was loaded would be dead in the one case it is advertised for. Decoding: validate before trusting the mirror. end >= start, non-zero start, a duration under an hour, and a readable non-empty name are all true of any sane kernel record on any CUPTI version. Records failing these are dropped, and one diagnostic naming the bound library and API version replaces the mystery empty timeline. Verified on A100 under torch-fl-211: 21/21 kernels with real names and non-zero durations against both the pip cu12 CUPTI (API 26) and the system CUDA-13 one (API 130000) via the override, confirming the checks do not false-positive across versions; and with the threshold temporarily forced to 1ns, every record is rejected with exactly one diagnostic and no crash. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * docs: design for torch-cuda profiler parity Stage B captures the GPU kernel timeline correctly, but a measured diff against a real torch+cu128 env shows what is still missing. Same code, both stacks: flow arrows 59 vs 0, cuda_runtime events 34 vs 0, kernel arg fields 13 vs 0, kernel names demangled vs mangled, and device time landing on aten::mm vs only on kernel names. Five of those six rows share one root cause: we never collect RUNTIME activities. cuda_runtime is the hub that holds both the External id linking back to cpu_op and the correlation linking forward to the kernel, so without it the whole chain is severed. Two mechanisms found while measuring: processTrace has a four-arg overload whose getLinkedActivity callback is the official channel for filling linked/flow (we only override the one-arg version, hence zero flows), and the Stage A ProfilerStubs sits on an unreachable else-branch -- its recorded skip reason was wrong, though the conclusion that the path is dead was right. Design keeps the external IActivityProfiler and splits it into a vendor-neutral DeviceTracer interface plus a CUPTI implementation, so adding ascend later means one new tracer and no changes to the kineto layer. Acceptance is a structural diff against a baseline snapshot committed to the repo, which keeps CI free of any CUDA torch dependency. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com> * docs: add profiler CUDA parity implementation plan 7 tasks: verification gate, DeviceTracer interface, CUPTI extraction, kineto adaptor rewrite, 13-field metadata, baseline snapshot + CI, cleanup. Refs #spec docs/superpowers/specs/2026-08-03-profiler-cuda-parity-design.md * feat(profiler): verify correlation foundation — RUNTIME activity + flow/device-time metrics Proves 3 design claims before refactoring: 1. Flow arrows (ac2g) appear in chrome trace — 28 entries = 14 paired s->f 2. key_averages aten::mm gains self_device_time_total — ~816us, exactly the sum of the 5 ampere_sgemm kernels in the same trace 3. Runtime events materialize with category privateuse1_runtime — 258 events ActivityType: PRIVATEUSE1_RUNTIME (CUDA_RUNTIME fallback never needed) Metrics: 28 flows (14 paired), 815.97us device time on aten::mm Two spec corrections measured by ablation, see task-1-report.md: - "device time attribution is free" is FALSE: suppressing activity.linked with everything else identical drops aten::mm from ~816us to 0us. Free only in the sense that no Python-side change is needed. - EXTERNAL_CORRELATION (kind 39) is a required step the plan omits: getLinkedActivity() keys on torch correlation ids, CUPTI records carry CUPTI ids, and only these records bridge the two. Also: flow arrows need BOTH halves with a shared CUPTI correlation id, gated on the correlation having a device-side activity. Keying on the torch id emitted 203 unpaired 'f' halves that still passed a bare count>0 assertion. Refs #spec 2026-08-03-profiler-cuda-parity-design.md §3.2 * fix(profiler): review round 1 — pairing assertion, window filter, deferred RUNTIME arming Addresses 6 Important review findings on the Task 1 verification gate. 1. metric 1 asserted a bare ac2g count, reproducing the false-pass documented in the task report. verify_flow_arrows now counts *paired* s<->f ids and requires count(s)==count(f) and equal id sets. 28 entries -> 12 renderable arrows. The old 203-unpaired-'f' trace would now FAIL. 2. metric 2's ground-truth cross-check was a manual observation. Now automated: sums dur over sgemm kernel events and requires |attributed - sum| within max(1us, 1%). Measured delta 0.000us. 3. EXTERNAL_CORRELATION parse branch had no validation — the one branch whose silent misdecode is invisible (garbage externalIds -> nullptr links -> device time silently back to 0). Now checks externalKind == CUSTOM0 via reportLayoutMismatch, consistent with the kernel/memcpy branches. 4. lookupExternalCorrelation returned 0 for both "not found" and a real torch correlation id (id 0 is a live value: autograd/profiler.py treats it as a frontend event). Now returns std::optional; unmapped activities never link. 5. startTime/endTime were accepted and dropped, shipping out-of-window events (66/258 runtime events before the first cpu_op; a 250ms "cudaLaunchKernel" against a 157ms span). Now filtered. Units verified empirically, not assumed: cuptiGetTimestamp is 2.7us from CLOCK_REALTIME and kineto's window is the same epoch-ns domain, so they compare directly — a MONOTONIC reading would have been off ~1.78e18ns and filtered everything. Filter applied both when building device_correlations and at emit, so a dropped kernel cannot leave a dangling flow half. Result: 0 events outside the window. 6. RUNTIME+EXTERNAL_CORRELATION armed at import cost every non-profiling user. MEASURED with a launch-bound workload (a GPU-bound one is insensitive and would have wrongly cleared it): +22%, 10.56 -> 12.88us/op over 3 A/B pairs, with ~15x more variance, far outside the ~0.1ms spread. Moved both kinds to session start(); ActivityRegisterCallbacks still runs at import, which is what cupti-must-arm-before-cuda-context actually requires. Overhead gone (10.48us/op) and all metrics still pass. Re-measured: 12 paired flow arrows, 815.13us aten::mm device time (cross-check delta 0.000us), 192 privateuse1_runtime events, window containment PASS. tests/unit/test_profiler_privateuse1.py: 4 passed, 1 skipped. * feat(profiler): add vendor-agnostic DeviceTracer interface Pure abstract class + DeviceEvent struct with open metadata map. DeviceEvent carries both correlation_id (CUPTI) and external_correlation_id (torch) to support kineto's getLinkedActivity() callback in Task 4. Stub factory in flagos_cupti_profiler.cc (moved to cupti_device_tracer.cc in Task 3). Refs #spec 2026-08-03 §2.1 * refactor(profiler): extract CUPTI logic to cupti_device_tracer.cc CuptiDeviceTracer implements the DeviceTracer interface: it owns every CUPTI detail (buffer callbacks, activity-record layout mirrors, the layout self-check, external-correlation resolution) and hands the layer above a flat vector of vendor-neutral DeviceEvents. flagos_cupti_profiler.{h,cc} is now a pure kineto adaptor with zero CUPTI types or calls -- Task 4 only has to rename it. Behaviour changes beyond the move: * MEMSET is now collected (design §4.1) and mapped to GPU_MEMSET/gpu_memset. cuBLAS zeroes a 512B workspace per gemm, so aten::mm legitimately gains ~9us of memset device time on top of its sgemm kernels. The Task 1 gate's ground truth was sgemm-only and had to learn about it; it now sums every device event linked to aten::mm (stricter -- it also catches a memset attributed to the wrong op) plus a separate assertion that kernels remain >90% of the total, so losing kernel collection cannot pass by shrinking both sides. Confirmed by same-binary A/B: disabling only the MEMSET ActivityEnable returns attribution to the sgemm sum exactly (delta 0.000us). * Runtime events are named from their cbid instead of being hardcoded to "cudaLaunchKernel", which previously reported blocking synchronizes and every malloc/free as kernel launches. The 22-entry table is cross-checked three ways -- pip cu12 header, system CUDA-13.0 header, and a live cuptiGetCallbackName(CUPTI_CB_DOMAIN_RUNTIME_API, cbid) query -- all agreeing. `_ptsz` variants map to their base name so a per-thread-default-stream build does not fall through to the generic label. * deviceCount() resolves cudaGetDeviceCount via RTLD_DEFAULT (returns a real 8 here) rather than hardcoding a guess; libcudart is in the process via the external libtorch_cuda.so but not on our link line. * The adaptor builds fmt header-only: kineto's addMetadata template calls fmt::format and libtorch_cpu.so exports no fmt::vformat, so emitting kernel metadata otherwise fails to link at import time. External-correlation resolution moved with the rest and now runs in drain() rather than per-record, because CUPTI does not order an activity record against the EXTERNAL_CORRELATION record that maps it. Left std::nullopt on a miss, never 0 -- 0 is a valid torch correlation id, and collapsing the two silently mis-attributes device time. Verified: correlation gate ALL METRICS PASSED (18 paired flow arrows, aten::mm 824.125us vs device-event sum 824.125us delta 0.000us); unit tests 4 passed, 1 skipped. Refs #spec 2026-08-03 §2 * fix(profiler): diagnose empty getLinkedActivity, guard buffer alloc, pin cbid mapping Review round 1 follow-ups for Task 3. 1. An empty getLinkedActivity callback silently produced linked == nullptr for every activity -- the exact ablation state Task 1 measured returning self_device_time_total to 0 -- with no log at any verbosity. Warn once, ungated (same rationale as reportLayoutMismatch: rare by construction, catastrophic when it happens). The summary log now reports linked/candidates plus resolver=present|EMPTY, so "0 linked because nothing ran", "0 linked because every id was rejected" and "0 linked because there is no resolver" are three distinguishable states rather than one ambiguous "0/N". Verified by temporary A/B forcing the empty branch: diagnostic fires and aten::mm device time drops to 0, confirming the warning marks the real failure mode; log reads "linked 0/209 ... resolver=EMPTY" vs "209/209 ... resolver=present" normally. 2. bufferRequested ignored aligned_alloc's return value, handing CUPTI a null pointer while claiming size = 8MB. Pre-existing code moved verbatim in the previous commit, but this file now owns buffer management. On null, decline the buffer properly (*size = 0, *maxNumRecords = 0) and warn once, degrading to dropped records instead of a write through address 0 under memory pressure. 3. Interim guard for the cbid->name table until Task 6's integration test lands: the gate now asserts at least one runtime event carries a name that is neither "cudaRuntime" (the default: fallback) nor "cudaLaunchKernel". That is precisely the invariant the old hardcode violated. Negative-tested by rewriting every runtime name to "cudaLaunchKernel" in a captured trace: the check returns False and fails the script. Gate: ALL METRICS PASSED (18 paired arrows; aten::mm delta 0.000us across 3 consecutive runs; new metric3b PASS). Unit tests: 4 passed, 1 skipped. * refactor(profiler): rename flagos_cupti_profiler → flagos_kineto_profiler The profiler adaptor is now vendor-agnostic (CUPTI logic moved to cupti_device_tracer.cc in Task 3). This rename reflects that role. Changes: - File rename: flagos_cupti_profiler.{h,cc} → flagos_kineto_profiler.{h,cc} - Class rename: FlagosCuptiProfiler{,Session} → FlagosKinetoProfiler{,Session} - Function rename: registerFlagosCuptiProfiler → registerFlagosKinetoProfiler - C API rename: flagos_cupti_*_correlation_* → flagos_kineto_*_correlation_* (exported symbols; test updated to match) - Profiler name() string: "flagos_cupti" → "flagos" (no consumers found) - Local debug helpers: flagos_cupti_debug/FLAGOS_CUPTI_LOG → flagos_kineto_* (adaptor copy only; cupti_device_tracer.cc retains its own) Verification: - verify_correlation_foundation.py: all 5 metrics PASS (18 flow arrows, aten::mm device time 824.123µs, cross-check delta 0.000µs) - test_profiler_privateuse1.py: 4 passed, 1 skipped (correlation test now passes) * fix(profiler): quote non-numeric metadata values in kineto trace JSON kineto's GenericTraceActivity::addMetadata stores every value with quoted=false, and metadataJson() emits it as `"key": <raw>`. That is safe only while every value happens to be a JSON literal -- which is true today purely by accident (all current values are numbers or bracketed lists). The blast radius is the reason to fix this BEFORE adding fields: kineto concatenates all activities into one document, so a single bare identifier ("N/A", a memory-kind label, a name) makes json.load() fail on the ENTIRE trace file, not just the offending event. DeviceEvent::metadata is map<string,string>, so values are already text by the time the adaptor sees them and cannot be classified by static type. metadataValueIsJsonLiteral() classifies by textual form instead, and is deliberately sound only in the "may I skip quoting?" direction: a spuriously quoted number is cosmetic, an unquoted non-literal is fatal. The quoted path escapes " and \ because addMetadataQuoted does no escaping of its own. No behaviour change today (every current value classifies as a literal); this is the safety net for the metadata fields that follow. * feat(profiler): add kernel/memcpy/memset metadata parity + occupancy Brings every device event's `args` dict up to what torch-cuda emits, so the two traces can be diffed structurally instead of eyeballed. kernel 5 -> 13 keys (+queued, device, context, stream, correlation, blocks per SM, warps per SM, est. achieved occupancy %) memcpy/set 2 -> 7 keys (+device, context, stream, correlation, memory bandwidth (GB/s)) runtime 1 -> 3 keys (+cbid, correlation) `External id` is NOT added here -- kineto's own logger emits it from linked/correlationId, and adding ours would produce a duplicate JSON key. Kernel9 mirror extended past `name` (reserved0/queued/submitted) to reach the `queued` timestamp. Offsets cross-checked against the installed CUPTI header: name=104, queued=120, submitted=128. The pip cu12 header (the copy actually bound at runtime) and the system CUDA-13 header are byte-identical here, so this is no more version-fragile than the prefix already mirrored. OCCUPANCY. Derived from CUDA's header-only occupancy calculator (no linking, no context); device attributes come via dlsym'd cudaDeviceGetAttribute, mirroring how deviceCount() already resolves cudaGetDeviceCount, and are cached per device because this runs in the buffer-processing hot path. Two things measurement settled, both of which a plausible-looking formula gets wrong: * Only occupancy clamps. "blocks per SM"/"warps per SM" are reported UNCAPPED -- a live 2048-block kernel reports 151.703705 warps/SM, far above the SM's 64-warp capacity. * "warps per SM" uses PLAIN division by warpSize, not a ceiling. A 16-thread kernel reports 0.004630; a ceil() form gives 0.009259, i.e. double. The two agree for every block size that is a multiple of 32, which is why this only shows up on a kernel like that one. Values AND their text formatting are reproduced: occupancy fields as fixed 6-decimal floats (float, not double -- 640/108*4 prints 23.703703 as a float and 23.703704 as a double, and torch-cuda reports the former), bandwidth as shortest-roundtrip. Verified against live torch-cuda traces on identical workloads: 18/18 kernel configurations match exactly on all three derived fields, including the bs=16 and 151.703705 cases above. Two judgment calls, both toward "say nothing rather than something false": * queued: emit 0 when the field is not populated (matches torch-cuda, whose traces show 0). CUPTI was MEASURED writing a literal 0 here, not the documented CUPTI_TIMESTAMP_UNKNOWN, so both spellings are rejected -- checking only the documented sentinel would pass a bogus 0 through as real. * device attributes unavailable: omit the three occupancy keys rather than emit zeros. A 0.0 occupancy is indistinguishable from a real measurement of a badly-occupied kernel; an absent key is honest and a parity test says so. Also widens the metadata JSON-literal classifier to accept exponent notation: %g emits "8e-05" for small bandwidths, which is valid JSON, and quoting it would silently break parity on exactly the values torch-cuda emits bare. Verified: Task 1 gate all metrics (18 paired flow arrows, aten::mm 825.5us, cross-check delta 0.000us); trace json.load()s; arg-key sets identical to torch-cuda for all four categories on two workloads; 4 passed 1 skipped. * test(profiler): add CI parity test + torch-cuda baseline snapshot Baseline: 13-key kernel args, 7-key memcpy/memset, 3-key runtime, ac2g flows. Test asserts: category coverage, PAIRED flow arrows, arg-key supersets, device-time attribution with cross-check, demangled kernel names, and non-hardcoded runtime names (the cbid-table regression guard). CI entry added to .github/configs/cuda.yml integration_tests list. Refs #spec 2026-08-03 §5 * feat(profiler): achieve torch-cuda parity — flows, metadata, device time attribution Removes the scaffolding: unreachable flagos_profiler_stubs.cc and the two tests/scratch verification scripts, whose coverage now lives in the CI parity test (window containment ported in this commit). Adds docs/profiler.md: three-layer architecture, the two correlation-id schemes, the CUPTI arming constraint, debug env vars, and known gaps. Refs #spec 2026-08-03 §6 * chore: untrack machine-local .libtorch_cuda_assets symlink It was committed as an absolute symlink into /nfs/lvyufeng/, which resolves nowhere on any other checkout. setup.py and scripts/with_cuda_libtorch.sh already reference the path without requiring it to be tracked, so ignore it instead. * style(profiler): satisfy pinned ruff on the profiler test files Consolidate scattered mid-file imports to the top, drop unused bindings, and apply ruff format. CI runs lint first with every other job declaring needs: lint, so these would have blocked the build and integration jobs. * docs(profiler): write docs/profiler.md in English Matches the language of the surrounding docs/ tree. --------- Co-authored-by: Claude <claude@anthropic.com> Co-authored-by: Claude Opus 4.8 (1M context) <noreply@anthropic.com>
ProcessGroupFlagOS delegated to the inner backend without pinning the current
device. The inner backend resolves the device it works on from whatever is
current, not from the tensor, so a caller that binds only with
torch.cuda.set_device(rank) -- the natural thing to do -- had its collectives
enqueued on the wrong device.
The current device does not stay where the caller put it. Measured on DTK:
torch.cuda.set_device(1) -> cuda=1 flagos=0
torch.ones(4, device='flagos:1') -> cuda=0 flagos=0
Allocating a flagos tensor resets the *cuda* current device to flagos's. flagos
and cuda track separate currents even though flagos aliases the same physical
GPU, so after the first allocation every rank had device 0 current.
FlagCX turns that into a GPU fault rather than a wrong answer: getStreamByIndex
lazily streamCreate()s one stream on the current device and caches it by index
alone, so the first collective binds the comm to device 0 permanently. Rank 1
then enqueued device-1 buffers onto a device-0 stream and RCCL faulted
(VMFault / "invalid resource handle") on the first all_reduce. A process kill,
not an exception.
Guard in the wrapper rather than asking callers to set both devices: c10d hands
us the tensors, so this is the one layer that always knows the right device.
Same reasoning as the DeviceGuard added to the CallPythonOp_* callers in flagos-ai#54,
one level up. _comm_device_guard moves and restores both currents, and is
applied to all 18 tensor-carrying collectives.
barrier carries no tensor to infer a device from, so it replays the index an
earlier collective bound the comm to; FlagCX rejects a barrier issued from
anywhere else with "flagcx communicator was initialized with different device".
Also merges two pairs of duplicate _allgather_base / _reduce_scatter_base
definitions, where the second silently shadowed the first.
test_flagos_dist_live.py passed with this bug present -- it happened to leave
the two currents agreeing -- so the new test deliberately calls only
torch.cuda.set_device(rank). Validated by neutering the guard: the test
reproduces the VMFault (exit 1), and passes (exit 0) with the guard restored.
Measured on 8x Hygon DCU (DTK), FlagCX inner backend:
2 cards 10/10 collectives OK DDP grads identical exit 0
4 cards 20/20 collectives OK DDP grads identical exit 0
8 cards 40/40 collectives OK DDP grads identical exit 0
matching the RCCL baseline. RCCL path re-checked after the change (4 cards,
20/20, exit 0): no regression.
Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
|
Bugbot is not enabled for your account, so this pull request was not reviewed. Enable Bugbot in the Cursor dashboard to get automatic reviews on future PRs. |
lvyufeng
force-pushed
the
fix-comm-device-index
branch
from
August 6, 2026 05:27
0a2c22c to
a74a9cc
Compare
Owner
Author
|
Opened by mistake against my own fork; the real PR is flagos-ai#57. |
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
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Fixes the device-index bug that causes collectives to enqueue device-N buffers onto a device-0 stream, resulting in VMFault on FlagCX/DCU and silent wrong results on other backends.
Root Cause
csrc/runtime/accelerator/cuda/device.ccmaintained athread_local int gCurrentDeviceshadow variable thatGetDevicereturned instead of asking the driver. Since flagos aliases the same physical GPUs as torch.cuda, andcudaSetDevice's current device is already thread-local, this shadow copy could only diverge:torch.cuda.set_device(1)moves the driver's index without going through ourSetDevice, leaving the shadow at 0DeviceGuardthen reads the stale 0 as "previous" and restores the process to a device that was never currentflagos:1tensor resetstorch.cuda.current_device()back to 0ProcessGroupFlagOSdelegates to an inner backend (FlagCX or RCCL) that resolves its device from whatever is current, so rank 1's device-1 collective gets enqueued on the device-0 stream FlagCX cached. On DCU this triggers VMFault; on other hardware it silently computes wrong results or hangs.Changes
GetDevicecallcudaGetDevicedirectly_allgather_baseand_reduce_scatter_basedefinitions that were silently shadowing each othertests/manual/test_comm_device_index.pydeliberately calls onlytorch.cuda.set_device(rank)to catch this class of bugVerification
test_comm_device_index.py+test_flagos_dist_live.py(20/20 collectives + DDP), all passtest_compute_device_index.py15 passed (this fix also resolves the twin bug that PR fix(flagos): run FlagGems Python ops on their operand's device flagos-ai/PyTorch-Plugin-FL#54 patched in Python)Impact
This also fixes the device-index hazard that flagos-ai#54 addressed on the compute side (
CallPythonOp_*). Both were symptoms of the same C++ accounting bug.