Measured on an eight-device Enflame S60 host. Installation and configuration are covered in flaggems-setup.md.
The flagos device reuses the u = 2 cut of the eight GCU dies, so
torch.flagos.device_count() reports 8 while the vendor runtime addresses 2
devices. Every measurement below was taken on flagos:0.
- Environment: venv (
gcu-flagtree), Python 3.12 - Location:
/public-flash/lvyufeng/venvs/gcu-flagtree
| Package | Version |
|---|---|
| torch | 2.10.0+cpu |
| torch_fl | 0.1.0 (in-place build_ext against the working tree) |
| flagtree | 0.6.1+enflame3.6 |
| triton | 3.6.0 (bundled by flagtree, backends {'enflame'}) |
| flag_gems | 0.0+g3c6f7537d, vendor_name == 'enflame' |
FlagGems is an editable install of master revision
3c6f7537d2d5d3aa680c55bbee5c70f2100c5b85. No triton distribution is
registered: flagtree installs its own triton package and the NVIDIA wheel has
to be uninstalled first (see flaggems-setup.md).
numpy is absent from this venv, and deliberately so. It is not a dependency of
the CPU torch wheel, and nothing on the GCU route imports it. pip warns at
install time that flag_gems declares a numpy requirement it cannot satisfy;
that warning is expected, because flag_gems is installed with --no-deps and
the GCU route never reaches a numpy import.
The measurement box runs Ubuntu 22.04 / glibc 2.35. The flagtree wheel is built
for glibc >= 2.38 and takes exactly one symbol newer than that: the C23
__isoc23_strtol. On this box import triton therefore fails with
ImportError: /lib/x86_64-linux-gnu/libc.so.6: version `GLIBC_2.38' not found
(required by .../site-packages/triton/_C/libtriton.so)
To take the measurements below, that one dynamic-symbol reference in the
measurement venv's libtriton.so was retargeted in place to glibc's existing
strtol (identical behaviour here; the C23 variant only adds 0b prefix
parsing). The edit is width-preserving and touches no code, only the dynsym
name and the version-need entry, and it is not committed anywhere.
CI needs none of this: .github/configs/gcu.yml pins the
flagtree-enflame3.6-py312-torch2.10.0-ubuntu24.04 image, whose glibc 2.39
satisfies the wheel as shipped. This has not been run through CI yet — the
change is still a local branch — so the claim rests only on that glibc version
comparison.
python3.12 -m venv /public-flash/lvyufeng/venvs/gcu-flagtree
source /public-flash/lvyufeng/venvs/gcu-flagtree/bin/activate
pip install torch==2.10.0 --index-url https://download.pytorch.org/whl/cpu
python3 -m pip uninstall -y triton # repeat until "not installed"
RES="--index-url=https://resource.flagos.net/repository/flagos-pypi-hosted/simple"
python3 -m pip install "flagtree===0.6.1+enflame3.6" $RES
python3 -m pip uninstall -y torch_gcu triton_gcu
cd /tmp && git clone https://github.com/FlagOpen/FlagGems.git && cd FlagGems
git checkout 3c6f7537d2d5d3aa680c55bbee5c70f2100c5b85
pip install --no-deps -e .
pip install packaging 'PyYAML==6.0.1' 'sqlalchemy==2.0.48'The shell used for every run below:
export TOPS_HOME=/opt/tops FLAGOS_ACCELERATOR=gcu GCU_KERNEL=1 FLAGGEMS_PYTHON=1 FLAGOS_BUILD_FLAGGEMS=0
export PYTHONPATH=/public-flash/lvyufeng/PyTorch-Plugin-FL
source /public-flash/lvyufeng/venvs/gcu-flagtree/bin/activateRebuild the in-place extension after any change under csrc/ or scripts/, and
re-run both generators after a change to the gap set:
python3 scripts/codegen/codegen_gcu_flaggems.py
python3 scripts/codegen/gen_vendor_confs.py
python3 setup.py build_ext --inplacetorch_fl/configs/backends_gcu.conf is generated by
scripts/codegen/gen_vendor_confs.py and is FlagGems-first by construction: for every
one of the 2037 routable overloads the generator picks the first available
backend in the order flaggems_cpp > flaggems > tileops > gcu > none, then
re-routes everything in NATIVE_TRITON_GAPS["gcu"] to the native topsaten
kernel (gcu), or to none where topsaten has no kernel either and the call
reaches ATen's cpu_fallback.
| Route | Overloads | Share |
|---|---|---|
flaggems |
255 | 12.5% |
gcu (native topsaten) |
176 | 8.6% |
none (cpu_fallback) |
1606 | 78.8% |
| accelerated | 431 | 21.2% |
backends_gcu.conf, sha256
366d986f0c7aa292eb398e9de8b3dd9a1533240fc4f026472299a3560245e40d for the
change this file records, and
4e8265afbf12b472097588ea1f7ae67f627b37711f2e3c17200f190588e0ea40 as committed
at the tip of this branch; the difference is the banner the flagos/main merge
rewrote, described below, plus the one route line and the coverage banner
upstream's 2f6b10c adds to the op list this branch was rebased onto.
The banner line is the only part of either generated file that the merge with
flagos/main changed: the generator paths inside them moved from scripts/ to
scripts/codegen/ in 1f62f74, so the hashes recorded before that merge
(10b8ab4b… for the conf, 39cd03e3… for the .inc) differ from these without
any route or registration changing.
The previous revision routed 152 overloads to gcu and 1884 to none, with no
FlagGems route at all, so this is 152 -> 431 accelerated routes. "Previous"
means the conf committed at the branch point -- 251c9c8, and unchanged in
flagos/main at bc39a83 -- which is what the wheel ships and what CI reads.
The two most recent revisions that touched that file (this one and d0e2d1a)
both ship zero flaggems routes. The generator is a different story and worth
separating out: it was already FlagGems-first for GCU at the branch point and
still is at bc39a83, emitting 88 flaggems routes, but no
gcu_flaggems_register.inc existed to back them, and the committed conf was
never regenerated -- scripts/codegen/gen_vendor_confs.py --check reports
backends_gcu.conf stale at both revisions. Regenerating alone would therefore
have shipped 88 routes pointing at an empty dispatcher slot. What this change
adds is the registration and the measured gap set that make the FlagGems-first
generation safe to ship.
The two counts reconcile with the generated registration:
csrc/aten/backends/gcu/generated/gcu_flaggems_register.incholds 248m.impllines (FLAGGEMS_PYTHON_OPSminus the 227-op gap set, minus the opsgcu_register.incalready registers).- 7 further
flaggemsroutes are claimed natively first bycsrc/aten/backends/gcu/generated/gcu_register.inc(clamp,fmod.Tensor,gelu,mean,mean.dim,remainder.Tensor,silu), which is where the 255 comes from.codegen_gcu_flaggems.pycomputes its op list as(FLAGGEMS_PYTHON_OPS - gaps) - native - pendingand omits those ops on purpose -- a secondm.impl()for the same op and dispatch key would only override the first and warn at import -- so the conf still namesflaggemsas their route while the native kernel serves them. The routing table and the two registration files reconcile at255 = 248 + 7. gcu_register.incregisters 183 native ops; 7 of them are the ops just listed, and the other 176 are thegcuroutes.
Every accelerated route is present in the registration its name implies: no op
is routed to gcu without a gcu_register.inc entry, so no route reaches a
registered-but-empty dispatcher slot.
A second, later change to the same conf moved eleven overloads from none to
gcu -- arange, arange.start, arange.start_step, fill_.Scalar,
fill_.Tensor, masked_fill_.Scalar, masked_fill_.Tensor, zero_,
view_as_real, view_as_complex, linalg_vector_norm -- which is the whole of
the +11 in the gcu column above and the whole of the -11 in none. It was
selected by census rather than by survey: a 50-step Qwen-Image-2512 run at
1024x1024 with FLAGOS_LOG=fallback recorded 80,536 cpu_fallback calls in
exactly six operators (fill_ 23,570, zero_ 19,221, view_as_complex
18,794, view_as_real 18,794, arange 79, linalg_vector_norm 78) against
468,227 device dispatches, with no seventh operator anywhere in the log.
view_as_real and view_as_complex are the two handwritten metadata kernels in
csrc/aten/strided_ops.cc; everything else was generated. Within this conf
gcu_register.inc moves 152 -> 163 m.impl lines and
gcu_flaggems_register.inc keeps its 248 (only its banner's count of FlagGems
ops already claimed natively moves, 88 -> 97), so at that point the
reconciliations read 255 = 248 + 7 and 156 = 163 - 7. No flaggems route was
touched.
A third change to the same conf moved three more overloads from none to gcu
-- add.out, sub.out and mul.out -- which is the whole of the remaining
156 -> 159 in the gcu column and the 1625 -> 1622 in none. They were
selected the same way, by census rather than by survey: two Adam steps over a
five-parameter model with FLAGOS_LOG=fallback recorded 45 cpu_fallback
calls, 25 of them aten::add and 10 aten::mul, all from the optimizer's own
state updates. Those two names were misleading -- csrc/aten/fallback.cc:21
logs op.schema().name(), so add_.Tensor, add.out and add.Tensor all
print as aten::add -- and ATen registers no PrivateUse1 kernel for
add_.Tensor / sub_.Tensor / mul_.Tensor at all: they are structured
in-place ops whose composite redispatches onto add.out / sub.out / mul.out
with out=self. The conf's add_.Tensor = none, add_.Scalar = none and the
mul_.* / sub_.* lines are therefore inert, and one kernel per out= spelling
serves both forms. add.out and sub.out are generated from the
binary_alpha_out template and mul.out from binary_out, all three through
topsaten, and the kernel sizes its output because the out= contract lets a
caller pass a zero-size or differently-shaped out. A cast guard
(c10::canCast) and an output device guard were added to these and to the three
pre-existing out= kernels bmm.out, mm.out and addmm.out. Two defects on
that path were fixed with it: at::native::flagos::resize_ used to delegate to
the CPU at::native::resize_, whose resize_bytes_cpu copies the old bytes
with a host memcpy and therefore faulted on an HBM source, so growing an out
storage segfaulted; and _copy_from took its byte count from the source while
writing into the destination's dtype in both directions. Within this conf
gcu_register.inc moves 163 -> 166 m.impl lines
(c6526ed3… -> f961aff0…) and gcu_flaggems_register.inc keeps its 248 again,
so the reconciliations above are what the tree now reads: 255 = 248 + 7 and
159 = 166 - 7. Two of the moves this conf makes come off flaggems rather than
none: _softmax and linalg_vector_norm gained entries in the generator's
NATIVE_TRITON_GAPS["gcu"] set, so the route generated for both changed from
flaggems to gcu. That is the 257 -> 255 in the table, and
linalg_vector_norm leaving this file's own m.impl list is the 249 -> 248.
A fourth change to the same conf moved twelve more overloads from none to
gcu -- _upsample_nearest_exact2d[.out], embedding_dense_backward[.out],
index_select[.out], index_fill.int_Scalar[_out], index_fill.int_Tensor[_out],
index_fill_.int_Scalar and index_fill_.int_Tensor -- which is the whole of the
159 -> 171 in the gcu column and the 1622 -> 1610 in none. Two censuses
picked them. A training step (two Adam steps over a five-parameter model with an
embedding lookup and an index_select path) with FLAGOS_LOG=fallback recorded
11 cpu_fallback calls in four operators -- aten::index_select 4,
aten::embedding_dense_backward 4, aten::index_fill_ 2, aten::nonzero_static
- The first three names cover all eight
index_*/embedding_dense_backwardspellings above:csrc/aten/fallback.cc:21logsop.schema().name(), and an in-placeindex_fill_reports asaten::index_fill_whichever value form it took. A Qwen-Image-2512 VAE decode at 1024x1024 supplied the fourth name: it made three_upsample_nearest_exact2dcalls, all of themcpu_fallbackon the pre-change tree. After the change the same training step records exactly onecpu_fallbackcall,aten::nonzero_static, with both step losses unchanged (-32.7952,-51.3676), and the VAE decode reportscpu_fallback ops: 0and runs to completion.
aten::nonzero_static is deliberately left on the fallback and is the only
remaining one on either path. Its result is int64, which
TopsatenSupportsDtype declines everywhere, so no gcu kernel can be written
for it; flag_gems ships no nonzero_static module, so there is no flaggems
kernel either; and the schema has no out= overload, so the FL layer's out=
lane cannot serve it. There is nothing to route it to, and the conf says none
because that is true rather than because the op was overlooked.
The kernels were added to the generator rather than by hand --
scripts/codegen/codegen_gcu.py gains the twelve overloads, twelve templates and
five helpers in csrc/aten/backends/gcu/topsaten_common.h -- and running it a
second time leaves all three generated artifacts byte-identical. Within this conf
gcu_register.inc moves 166 -> 178 m.impl lines
(f961aff0… -> 806d7ce7…) and gcu_kernels.cc 35c8a37e… -> 7f34288a…,
while gcu_flaggems_register.inc keeps its 248 (9c9c99d0…), so the
reconciliations above are what the tree now reads: 255 = 248 + 7 and
171 = 178 - 7. The fourth change moves nothing off flaggems; all twelve of
its routes come from none.
Because registration and routing are checked against each other, the twelve ops
cannot be measured against the pre-change conf: with the new gcu_register.inc
in place and none still in the conf, the first _upsample_nearest_exact2d call
raises routed to 'none' (no accelerated impl on this platform) but the op is registered on PrivateUse1 instead of reaching cpu_fallback. The before-numbers
above therefore come from a build that predates the registration, and the after
numbers from a build that has both halves of the change.
A fifth change to the same conf moved five more overloads from none to gcu --
all, index.Tensor, nonzero, nonzero_static and where.self -- which is the
whole of the 171 -> 176 in the gcu column, the 1611 -> 1606 in none and the
426 -> 431 accelerated. It is the one change here selected by a census of the
pipeline under test rather than by a probe standing in for it. Both halves of the
paired
Qwen-Image-2512 flagos cohort on the S60 (torch_fl + diffusers,
FLAGOS_LOG=dispatch,fallback) log cpu_fallback before their first prompt, and
the two logs agree operator for operator: 62 cpu_fallback calls in exactly four
operators -- aten::where 56, aten::nonzero 2, aten::index 2 and aten::all
2 -- against 489,770 dispatch records in the first half and 429,692 in the
second. No fifth operator appears in either log. Within this conf
gcu_register.inc moves 178 -> 183 m.impl lines, and gcu_flaggems_register.inc
keeps its 248 while its banner moves 97 -> 100: three of the five (all,
nonzero, where.self) were on the list of FlagGems ops gcu_register.inc
already claims, so claiming them natively is what takes them off it. The
reconciliations above are what the tree now reads: 255 = 248 + 7 and
176 = 183 - 7. The siblings stay where they were: where.self_out, all.dim
and all.dims are already on flaggems, and all.out, all.all_out,
nonzero.out and nonzero_static.out stay on none because the pipeline does not
call them. This change neither adds nor removes a FlagGems route, but neither was
the survey re-run against the cohort the four earlier changes left: its kept set is
the 257 routes that predate the two the third change moves off flaggems, so the
per-operator results below describe a cohort two operators wider than the one the
conf ships.
flaggems-- FlagGems Triton kernel through the Python bridgegcu-- native topsaten kernelnone-- not registered on PrivateUse1; ATencpu_fallback
FLAGOS_LOG=dispatch # print "[flagos dispatch] <op> -> <backend>"
FLAGOS_OP_abs=flaggems # force one op onto a backend (dots -> __)
FLAGOS_BACKEND_CONFIG=/path/to.conf # force a different conf (testing only)The eleven none -> gcu moves recorded above came from a census of the host
round trips a real model makes, not from the survey. FLAGOS_LOG=fallback
over a 50-step Qwen-Image-2512 run at 1024x1024 lists six operators and no
seventh, which is the set scripts/codegen/codegen_gcu.py was taught to serve:
arange (three overloads, from at::native::compute_arange_size on the host
plus topsatenArange), the in-place writes zero_ -> topsatenZero,
fill_.Scalar / fill_.Tensor -> topsatenFill_ and masked_fill_.Scalar /
masked_fill_.Tensor -> topsatenMasked_fill, linalg_vector_norm ->
topsatenLinalgVectorNorm, and the two complex views. The in-place writes are
what torch.zeros / ones / full decompose into on a device, so they are
worth more than their own call count; the complex views matter because a view op
on the cpu_fallback path is copied, so its result stops aliasing its input.
view_as_real and view_as_complex are the two handwritten registrations in
csrc/aten/strided_ops.cc: every category template the generator emits ends in
a topsaten:: call, so a pure metadata alias is not expressible in it, and the
generator lists both in METADATA_OPS. The TopsatenArangeDtype predicate
accepts Float / Int / Short / Char / Byte / Half / BFloat16 only, so int64 and
float64 arange stay on the host path. The full measurement, including the
1-ULP bound on a fractional arange step and the probe coverage, is in
docs/reference/operator-support.md; the NATIVE_TRITON_GAPS["gcu"] entries
for arange, zero_ and the two fill_ overloads stay, because a gap entry
only takes an op off FlagGems and for the factory family it is what selects
the vendor.
The last three none -> gcu moves came from a second census, this one over a
training step rather than an inference pipeline. Two Adam steps over a
five-parameter model, run once with foreach=True and once with foreach=False,
logged 45 cpu_fallback calls before the change and 11 after, with the 25
aten::add and 10 aten::mul entries gone and matching 25 add.out -> gcu and
10 mul.out -> gcu dispatch lines in their place; both losses are byte-identical
across the two runs (-32.7952, -51.3676), so the reroute changed the route
and not the numbers. The remaining 11 are index_select (4),
embedding_dense_backward (4), index_fill_ (2) and nonzero_static (1), none
of which the optimizer itself produces.
The ops are worth more than three routes suggest, because one kernel serves both
the explicit out= spelling and the in-place one: ATen gives add_.Tensor,
sub_.Tensor and mul_.Tensor no PrivateUse1 registration of their own, and
their composite redispatches onto add.out / sub.out / mul.out with
out=self. The conf's add_.Tensor, add_.Scalar and mul_.* / sub_.* lines
predate this and are inert -- the dispatcher never sees those spellings -- which
is why the fallbacks looked invisible in a log that prints the schema name rather
than the overload. On a 4096x4096 in-place add_ the device path measures
0.507 ms/call against 114.185 ms/call for an int64 operand, which takes
the host path by construction and is the in-run control: 225.3x. The full
measurement -- the nine in-place spellings, both broadcasts, five dtypes, a
strided-view write, the explicit out= forms including a zero-size and a
differently-shaped out, the narrowing rejections, and the resize_ and
copy_ defects this path reaches -- is in docs/reference/operator-support.md.
Nothing here is GCU-specific in the shared code: the resize_ and copy_ fixes
land on Ascend, MUSA, DCU, MetaX, PPU and Tsingmicro too, and those platforms are
not revalidated.
tests/manual/flaggems_overload_survey.py (harness v4) was run against the GCU
conf as it stood before the measured gap set was applied: 374 flaggems
routes, seven profiles per overload, 2,618 cases total.
That draft is a transient working-tree version of
torch_fl/configs/backends_gcu.conf and is not byte-recoverable now. The survey
records both the path and the hash it read --
"conf": ".../torch_fl/configs/backends_gcu.conf",
"conf_sha256": "82f801778c392eafa6aba7f581cdbcaa4f633ae5766b7796efc2c9cd25bfb8d8"
-- and 82f801778c matches neither the base commit's conf (039fb323…) nor the
one this change ships (7b87541e…, later 74aab449…). It was overwritten when
the final conf was
regenerated. What survives is the per-op JSON, which stays the auditable
evidence.
Its route count is the un-gapped candidate set, and it reconciles with the
shipped conf: 374 - 117 - 2 = 255, where 117 = 81 ops wrong on a
float16/float32 profile plus 36 ops wrong only for int64 with a topsaten
kernel to fall back to -- precisely the ops the measurement sent back to the
vendor, as broken out below -- and the further 2 are the Group C gaps added
after the survey (see "Group C" below).
Those two are invisible to this harness by construction: every one of its seven
profiles is at most 32 rows wide, and the defect is a launch-grid axis limit
that only a real workload reaches. The 76 int64-only failures left on FlagGems
and the 60 unexercised routes were candidates at measurement time and are
candidates in the shipped conf; they were never part of that 117.
| Profile | dtype | shape | layout |
|---|---|---|---|
2d-f32 |
float32 | 32x32 | contiguous |
4d-f32 |
float32 | 2x3x16x16 | contiguous |
1d-f32 |
float32 | 32 | contiguous |
2d-f16 |
float16 | 32x32 | contiguous |
2d-i64 |
int64 | 32x32 | contiguous |
2d-bool |
bool | 32x32 | contiguous |
2d-f32-strided |
float32 | 32x32 | strided |
Survey summary: 374 registered, 314 tested, 239 basic-executable, 121 strict.
| Verdict | Ops | Meaning |
|---|---|---|
STRICT |
121 | every valid profile matches the CPU reference within tolerance |
BASIC_ONLY |
118 | runs, but at least one valid profile is outside tolerance |
FAILED |
75 | no valid profile produced a correct result |
UNTESTED |
60 | the harness could not construct a valid case for any profile |
Verdicts decide whether a route works; they do not decide whether to gap it, because a route that is wrong only for int64 still carries a working float32 path. The routing decision therefore used a second, profile-aware pass:
| Class | Ops | Route taken |
|---|---|---|
| clean on every exercised profile | 121 | stays on flaggems |
| wrong on a float16/float32 profile | 81 | reverted to the vendor route |
| wrong only for int64/bool, native topsaten kernel exists | 36 | reverted to gcu |
| wrong only for int64/bool, no topsaten kernel | 76 | stays on flaggems (see below) |
The 60 UNTESTED ops are disjoint from all four rows: none of them produced a
valid case on any profile, so they are counted nowhere above. They are all
shape/metadata ops (avg_pool2d(_backward), col2im, reflection_pad*,
native_group_norm*, scatter.reduce, _thnn_fused_lstm_cell,
_scaled_dot_product_*_backward, ...). They are not covered by this
measurement and are not individually gapped; they stay on FlagGems and are
listed under "Not covered" below.
121 + 81 + 112 + 60 = 374, and 81 + 36 = 117 of the 227 gap entries come from this measurement; 2 more come from the Group C scale failures measured on Qwen-Image-2512 (below), and the remaining 108 are the pre-existing entries.
The seven-profile survey above cannot see this family: its widest profile is 32
rows, and the defect is in the launch grid the enflame backend builds from a
data dimension. GCU's grid limits are far below CUDA's -- grid.x 65535
against 2**31-1, grid.y 255 against 65535 -- so a kernel launched once per
row (or per 8-row tile) passes the survey and raises on a real workload.
Measured on the S60 with torch._softmax(x, -1, False) on bf16:
| Input | Leading dim product M |
Result |
|---|---|---|
(1, 24, 2729, 2729) |
65496 | PASS |
(1, 24, 2731, 2731) |
65544 | FAIL grid.x Required 65544 |
(1, 24, 4114, 4114) |
98736 | FAIL grid.x Required 98736 |
The non-inner kernel has the same defect on the other axis: a softmax that is
not over the last dimension launches with grid = (M, cdiv(K, TILE_K), 1), so
(24, 4114, 4114) over dim 0 fails with grid.y Required 2067.
| Op | Kernel and launch | Failure | Route |
|---|---|---|---|
_softmax |
softmax_kernel_inner, grid = (M, 1, 1) |
grid.x above 65535 |
flaggems -> gcu |
linalg_vector_norm |
l2_norm_kernel, grid = (cdiv(M, BLOCK_M),), BLOCK_M <= 8 |
grid.x above 524280 |
flaggems -> gcu |
Both were reached from Qwen-Image-2512 on the S60: _softmax on
(1, 24, 4114, 4114) from F.scaled_dot_product_attention in every one of the
60 transformer blocks, and linalg_vector_norm via F.normalize(x, dim=1) in
diffusers' QwenImageRMS_norm on the VAE decode at (1, 128, 1, 1024, 1024),
M = 1048576, failing with grid.x Required 131072.
topsatenSoftmaxForward serves every shape above -- it has no grid to size --
so _softmax goes back to the vendor kernel. That was originally the only
native destination for the pair: linalg_vector_norm had no topsaten kernel at
the time, so its gap landed on none and the call reached cpu_fallback, and
the vendor API itself was not the missing piece -- topsaten carries
topsteL2Norm, but topste is not wired into the generator. It is now served
by topsatenLinalgVectorNorm instead (see "Native kernels for the six measured
CPU fallbacks" under Routing Configuration), so this gap no longer costs a host
round trip.
FLAGOS_LOG=dispatch python -c "
import torch, torch.nn.functional as F, torch_fl
x = torch.randn(32, 32, device='flagos:0')
y = F.silu(x); y = torch.abs(x); y = torch.relu_(x)
print(torch.sum(x, dim=1).sum().item())
"[flagos] loading backend config from .../torch_fl/configs/backends_gcu.conf
[flagos dispatch] randn -> gcu
[flagos dispatch] silu -> flagos_python
[flagos dispatch] abs -> gcu
[flagos dispatch] relu_ -> flagos_python
[flagos dispatch] sum.dim_IntList -> gcu
[flagos dispatch] sum -> gcu
silu and relu_ take FlagGems; abs and sum.dim_IntList take the native
kernel because the gap set reverted them. On the base commit's conf both were
gcu # gcu, so silu -> flagos_python is the reroute this change is about and
abs -> gcu is the one the measurement sent back. randn is native in both
configs and is printed only as the buffer the probe starts from.
[flagos] loading backend config from ... is printed once per process and names
the conf that decided the routes.
The two sum lines are the same call: this PyTorch build reaches
sum.dim_IntList first and its native kernel then dispatches the inner
sum.
Every entry added to NATIVE_TRITON_GAPS["gcu"] came from one of the families
below. Reproduce any of them by forcing the op back onto FlagGems:
FLAGOS_OP_asin=flaggems FLAGOS_LOG=dispatch python -c "
import torch, torch_fl
torch.asin(torch.randn(32, 32, device='flagos:0'))
"The compiler emits error: 64-bit data type not supported on GCU300! and the
runtime surfaces it to Python as RuntimeError: Pipeline run failed: PassManager execution failed. The rejected type is a pointer parameter, so the operand
dtype is what has to be avoided; the op itself is fine for f16/f32/i32/bool.
Reproduced by every reduction, scan and index overload whose kernel carries an
index accumulator: nonzero, argmax, argmin, count_nonzero, _unique2,
unique_dim, unique_consecutive, topk, median, nanmedian, max.dim,
min.dim, mode, kthvalue, range, repeat_interleave.Tensor,
unfold_backward, index_copy(_), sum.out, var.correction,
var_mean.correction, _weight_norm_interface(_backward), norm.Scalar,
norm.ScalarOpt_dim, native_layer_norm, mse_loss, and the scatter/renorm
family. Measured on all seven profiles for nonzero; on the float profiles only
for the rest.
add logic to support op arith.maxsi
UNREACHABLE executed at .../third_party/enflame/triton_gcu/triton_gcu300/lib/Analysis/PtrAnalysis.cpp:1711!
An abort, not an exception: the process exits with SIGABRT, so the op cannot be
caught at the Python level. Reproduced on every profile, all dtypes, by
_adaptive_avg_pool2d. Nine further ops are in the same family but abort in the
elementwise lowering instead (family 3).
unsupported extern elementwise: __nv_asinf
UNREACHABLE executed at .../TritonToGCU/ElementwiseFusionOpToGCU.cpp:874!
The enflame backend does not lower the libm externs flag_gems emits for the
inverse-trig and special-function families. Reproduced by asin, asin_,
erf-adjacent specials and the special_chebyshev_* /
special_shifted_chebyshev_* / special_hermite_polynomial_h group.
EE: [tid:498312] efdrv/src/driver/dtu_context_obj.cc:693:submit_sip_assertion_task
##abort as Detected SIP assert###
terminate called after throwing an instance of 'std::runtime_error'
what(): Detected SIP assert!!!
Reproduced by addr on the 1d profile (SIGABRT), and by
_batch_norm_no_update and native_batch_norm (SIGSEGV, exit status -11).
The harness truncates captured stderr at 300 bytes, so for the two
batch-norm ops the recorded evidence is the exit status plus the process's last
stderr output (FlagGems' active Triton backend does not provide a replay benchmarker warning); the faulting frame is not in the record. They are
gapped because they kill the process, but the signal is the whole of the
evidence.
These run to completion and return the wrong tensor, which is the reason they must not be inferred from a "did it raise" test:
| Op | Measured difference from CPU |
|---|---|
elu, elu_, elu_backward |
max_diff 0.38 -- 0.89 on f32, 0.70 on f16 |
histc |
max_diff 32 -- 1536 (bin counts) |
_softmax_backward_data, _log_softmax_backward_data |
returns torch.int8 where torch.float32 is expected |
sum.out |
returns shape (32, 32) where the scalar output shape () is expected |
_pdist_backward, _weight_norm_interface(_backward) |
outside tolerance on all float profiles |
addmm.out, addmm_, silu_backward |
outside tolerance on all float profiles |
tril.out, tril_, triu, triu_ |
outside tolerance on all float profiles |
Not device defects -- the FlagGems Python wrapper is stricter than the ATen schema, so a caller that works on CPU raises on GCU:
topk:AssertionError: Currently only support topk in last dimensionspecial_modified_bessel_k0(.out):AssertionError: Tensors must be CUDA tensorsdequantize.self:NotImplementedError: Could not run 'aten::int_repr' with arguments from the 'CPU' backend_linalg_eigvals:RuntimeError: PyTorch was compiled without CUDA support
cumsum, cumsum.out, cumprod, cumprod_, cummax, cummin,
logcumsumexp, logcumsumexp.out were already in the gap set before this
measurement and remain there.
76 overloads are correct for float16, float32, int32 and bool and fail only
for an int64 operand, and have no topsaten kernel to fall back to. Gapping them
would send their float path to cpu_fallback as well, so they stay on FlagGems.
The consequence is explicit: these ops raise Pipeline run failed: PassManager execution failed for an int64 operand, where the previous
configuration served the same call through cpu_fallback. The set is dominated
by ops a float model never hands an int64 tensor:
abs_ acosh addcmul addcmul_ angle asinh atan2 atanh ceil.out ceil_
clamp_max clamp_min digamma exp2 floor.out floor_ floor_divide.Scalar fmax
fmax.out fmin fmin.out gcd_ gt_.Scalar gt_.Tensor isin.Scalar_Tensor isinf
isnan isneginf lcm lcm_ logical_and_ logical_not logical_not_ logical_or
logical_xor logical_xor_ logit logsumexp lt_.Scalar lt_.Tensor masked_scatter
masked_scatter_ max min nan_to_num neg_ new_ones pow.Scalar pow_.Tensor prod
prod.dim_int rad2deg relu_ remainder.Scalar_Tensor rsub.Scalar sgn_ signbit
sinc slice_backward special_airy_ai special_chebyshev_polynomial_u.n_scalar
special_i0e special_i1 special_modified_bessel_k1
special_shifted_chebyshev_polynomial_v special_xlog1py sum.IntList_out tan
threshold threshold_ threshold_backward trace trunc_ xlogy.Scalar_Other
xlogy.Scalar_Self xlogy.Tensor
The 36 overloads in the same class that do have a topsaten kernel are gapped
and route to gcu, where the generated kernel's TopsatenSupportsDtype guard
round-trips int64 through the CPU. That is slower than FlagGems but correct, and
it keeps int64 working for the three highest-traffic names in the group
(abs, neg, sum.dim_IntList), which the previous configuration also served:
abs acos addmm amax amin atan ceil cos cosh erf exp expm1 flip floor
fmod.Scalar log log10 log1p log2 logical_and logical_or mse_loss_backward neg
pow.Tensor_Scalar pow.Tensor_Tensor reciprocal relu remainder.Scalar sigmoid sin
sinh sqrt sum.dim_IntList tanh tril trunc
Root cause: the GCU300 front end in triton_gcu300 rejects 64-bit types in
kernel IR. It is a property of the compiler, not of flag_gems: the same
flag_gems kernel compiles on CPU/CUDA Triton backends, and the equivalent
hand-written Triton kernel fails identically on this device.
Reproduced locally:
FLAGOS_OP_abs=flaggems python -c "
import torch, torch_fl
torch.abs(torch.tensor([-1, -2, 3], device='flagos:0'))
"RuntimeError: Pipeline run failed: PassManager execution failed
Impact: any FlagGems route whose op is handed an int64 tensor raises. Where topsaten has a kernel the op is gapped and the call succeeds on the vendor path; where it does not, the op raises rather than falling back (76 overloads, listed above).
Resolution: not fixable in this repository. Routing is per operator and the rejected type is the operand dtype, so there is no per-dtype fallback to hang this on without a dispatcher change. The gap set is the mitigation; the remainder is documented rather than hidden.
Root cause: the lowered elu kernel produces values that disagree with the
CPU reference on all four float profiles. Whether this is the flagtree lowering
or the flag_gems formulation was not isolated; the observable is what the
routing decision was made on.
Reproduced by the survey's 2d-f32 profile, dtype float32:
max_diff=0.8942340614211499 # 2d-f32, torch.elu
max_diff=0.7372526526451111 # 2d-f32, torch.elu
Impact: silent wrong answers on every float profile, including elu_backward.
Resolution: elu, elu_, elu_backward moved to
NATIVE_TRITON_GAPS["gcu"]. A wrong answer is worse than a slow one, so this
family was treated as more urgent than the crashes.
Impact: the measurement above does not cover these ops -- shapes and
metadata-driven overloads the generic harness cannot construct a valid call for:
_conv_depthwise2d, _efficient_attention_backward,
_embedding_bag_dense_backward, _flash_attention_backward, _masked_scale,
_scaled_dot_product_{cudnn,efficient,flash}_attention_backward,
_sparse_semi_structured_mm, _thnn_fused_lstm_cell(_backward_impl),
_upsample_bicubic2d_aa_backward, addmm.dtype(_out), addmv,
avg_pool2d/3d(_backward), bincount, clamp, col2im, cudnn_convolution,
mv, native_batch_norm_backward, native_group_norm(_backward),
nll_loss2d_*, reflection_pad1d/2d/3d(_backward), replication_pad1d/3d,
rot90, scatter.reduce, scatter_reduce.two(_out), scatter_reduce_.two,
upsample_linear1d_backward, fractional_max_pool2d(_backward),
max_pool2d/3d_with_indices, max_unpool2d, diagonal_backward,
glu_backward, index_reduce_, affine_grid_generator,
rrelu_with_noise_backward, softplus_backward, lift, polar, and the
special_scaled_modified_bessel_k1 pair.
Resolution: not gapped. They stay on FlagGems, and this document records that they are not measured rather than claiming they pass. A targeted harness is the follow-up.
RuntimeWarning: active Triton backend does not provide a replay benchmarker; falling back to event timing
The enflame Triton backend does not implement the replay benchmarking API
FlagGems' libentry prefers. Timing falls back to event measurement. Correctness
is unaffected and no route was changed for it.
- Other GCU hosts: the S60 is the only Enflame machine available to this
change. No GCU400 or multi-card claim is made, and the device-1 (
flagos:1) paths were exercised only indirectly through the unified RNG suite. - The 60 unexercised routes above.
- torch.compile on FlagGems-routed ops:
test_compile.pypasses, but its cases are not selected by the FlagGems route set, so it does not establish that a FlagGems-routed op survivestorch.compile. - Non-GCU platforms: no route was measured for Ascend, MUSA, DCU, MetaX or PPU, and no route was changed for them.
export PYTHONPATH=/public-flash/lvyufeng/PyTorch-Plugin-FL
source /public-flash/lvyufeng/venvs/gcu-flagtree/bin/activate
# FlagGems runtime path
FLAGOS_USE_FLAGGEMS=1 pytest tests/integration/ops/ -m "flaggems and main_ops" -v
# Vendor/generic operator cohort
pytest tests/integration/ops/ -m "(anyplatform or main_ops) and not flaggems_python and not flaggems_cpp" -v
# Generators and conf consistency
python3 scripts/codegen/codegen_gcu_flaggems.py --check
python3 scripts/codegen/gen_vendor_confs.py --check
pytest tests/unit/test_gen_vendor_confs.py tests/integration/ops/test_flaggems_conf_consistency.py -vEvery group of .github/configs/gcu.yml was run locally on the S60 against the
final tree, in one uninterrupted pass, in manifest order. All eight exit 0.
| # | Group | Command | Result |
|---|---|---|---|
| 1 | Check isolated GCU environment | bash .github/scripts/set_env_gcu.sh |
see below |
| 2 | Vendor backend, main ops | pytest tests/integration/ops/ -m "(anyplatform or main_ops) and not flaggems_python and not flaggems_cpp" |
595 passed, 32 skipped, 499 deselected, 2 xfailed, 2 xpassed |
| 3 | FlagGems runtime path, main ops | FLAGOS_USE_FLAGGEMS=1 pytest tests/integration/ops/ -m "flaggems and main_ops" |
9 passed, 4 skipped, 1116 deselected, 1 xpassed |
| 4 | Unified RNG | pytest tests/integration/ops/test_rng_dispatch.py -m main_ops |
111 passed, 4 skipped, 1 deselected, 1 xpassed |
| 5 | General (factory ops) | pytest tests/integration/test_factory_ops.py |
46 passed |
| 6 | Unified AMP contract | pytest tests/integration/test_amp_contract.py -m amp |
27 passed |
| 7 | Unified math-bits contract | pytest tests/integration/test_math_bits_contract.py -m math_bits |
12 passed |
| 8 | torch.compile | pytest tests/integration/test_compile.py |
29 passed, 18 skipped |
Two of these groups failed on the first attempt, for reasons that are not defects in the shipped routes and that a reader should be able to check against this record:
- Group 3, 3 of 14.
test_abs_dispatch.py,test_neg_dispatch.pyandtest_sum_dispatch.pyasserted-> flagos_pythonunconditionally, but this change movesabs,negandsum.dim_IntListto the vendor route. They were rewritten to assert the route the platform configuration actually lists (routed_backend(...)), the patterntest_add_dispatch.pyalready used for the MUSAadd.Tensorcase. This is the positive control for the rerouting: the tests fail if the gap set is reverted, and the group passes after the rewrite. - Group 8, 1 of 29.
test_torch_backends_entry_point_is_registeredfailed withAssertionError: torch_fl must register a torch.backends entry point; found {}against a localtorch_fl.egg-infobaked withoutentry_points.txt.setup.py egg_inforegenerated it and the group passes. A stale local build artifact, not a property of the change.
pytest tests/unit/ is not part of the GCU CI manifest, which runs integration
groups only. Run on the same tree it is 2 failed, 317 passed, 97 skipped, and
both failures pre-date this change:
test_gen_vendor_confs.py::test_shipped_confs_are_up_to_datereportsstale: ['backends_ascend.conf']. Run on a pristineflagos/mainworktree (bc39a83), the same test reportsstale: ['backends_gcu.conf', 'backends_ascend.conf'], so both halves are pre-existing and this change repairs the GCU one. The repair is not just the regeneration: that worktree's generator wants to emit 88 GCUflaggemsroutes and not one of them had a registration behind it, so regenerating there would have produced a conf whoseflaggemsroutes dispatch to an empty slot. The change supplies the missinggcu_flaggems_register.incand the measured gap set first, and only then regenerates. Ascend is left alone: it is a different platform and out of scope.test_profiler_privateuse1.py::test_guard_stream_is_real_not_syntheticfails withassert 0 != 0(csrc/runtime/guard.hreturns stream id 0); unrelated to routing and untouched here.
.github/scripts/set_env_gcu.sh was run end to end against a scratch venv
(TORCH_FL_VENV_ROOT pointing at an empty directory, CI_STAGE=integration), so
the whole provisioning path ran for real rather than being read. Every step
before the final import check succeeded:
| Step | Observed |
|---|---|
| TopsRider SDK discovery | TOPS_HOME=/opt/tops, lib/libtopsrt.so + include/gcu/topsaten/topsaten.h found |
| topsaten library | /usr/lib/libtopsaten.so |
| Device node | /dev/gcu0 present |
| venv bootstrap | provision-venv/bin/python, Python 3.12 |
| CPU torch | 2.10.0+cpu, from site-packages/torch/__init__.py |
torch_gcu excluded |
find_spec("torch_gcu") is None asserted clean |
| NVIDIA triton removed | pip show triton empty before the flagtree install |
| flagtree | flagtree-0.6.1+enflame3.6 installed from the FlagOS index |
| FlagGems | flag_gems-5.4.0rc2.post1+g3c6f7537d built from git and installed |
| FlagGems deps | packaging, PyYAML==6.0.1, sqlalchemy==2.0.48 |
| Triton import check | Triton: 3.6.0 (backends: ['enflame']), flag_gems: installed |
The last row is the one step the box could not reach unassisted: the check
aborted with the glibc 2.38 error described under
Host toolchain caveat. After
applying that same local-only retarget to the scratch venv's libtriton.so, the
script's own verification snippet — the exact heredoc it runs, asserting
"enflame" in triton.backends.backends and that flag_gems is importable —
passes verbatim, which is the output quoted above. Everything else in the run is
unmodified, and the failure is a property of the 22.04 host, not of the script.
The remaining statement to be exact about: this is a local reproduction of the
provisioning path, not a CI run. .github/configs/gcu.yml has not executed this
change yet.
Both generators were run twice in a row and the second run produced no change:
scripts/codegen/gen_vendor_confs.py--backends_gcu.confsha256 identical across the two runs (74aab449...), and--checkexits 0.scripts/codegen/codegen_gcu_flaggems.py--gcu_flaggems_register.incsha256 identical across the two runs (9c9c99d0...), and--checkexits 0.scripts/codegen/codegen_gcu.pyhas no--checkmode, so its idempotency was proven by re-running it three times and comparing the three generated files: byte-identical each time (74aab449...conf,35c8a37e...kernels,f961aff0...register).
backends_gcu.confwent from 152 to 431 accelerated routes out of 2037 (7.5% -> 21.2%): 255flaggems, 176gcu, 1606none.- 374 FlagGems routes were measured across seven profiles each: 121 are clean on every exercised profile, 81 are wrong at float16/float32 and were reverted to the vendor route, and 112 are wrong only for int64/bool (36 of them reverted because topsaten has a kernel, 76 kept on FlagGems because it does not).
- A census of a real workload replaced the survey for the last nineteen moves:
a 50-step Qwen-Image-2512 run at 1024x1024 made 80,536
cpu_fallbackcalls in exactly six operators, two Adam steps over a five-parameter model made 45, of which 35 were the in-place arithmetic the optimizer runs every step, and the pipeline's own warm-up on the S60 made 62 in four operators. The six reach a kernel on the device; of the 35, 25aten::addand 10aten::mulnow do too, and the 11 the optimizer does not produce are what remains. - FlagGems is not patched or forked. Every route that could not work on the
enflame stack was moved into
NATIVE_TRITON_GAPS["gcu"]in the generator. - All eight
.github/configs/gcu.ymlgroups pass locally on the same tree.