[KMCompiler] Add optimized index_copy and index_copy_ kernels for Ascend - #5507

Closed

Onisen7 wants to merge 192 commits into

flagos-ai:masterfrom

Onisen7:add-ascend-index-copy

Commits on Aug 2, 2026

Commits on Aug 4, 2026

feat: add index_copy_ operator for Hygon backend
Onisen7
committed
0fc7db7
Update index_copy_.py
Onisen7
authored
3daeba0

Commits on Aug 7, 2026

feat: add index_copy_ operator for Ascend backend
Onisen7
committed
03018f1
Merge branch 'master' of github.com:Onisen7/FlagGems into ascend-index-copy-opt
Onisen7
committed
88e5a6e
feat: add index_copy_ for Ascend backend
Onisen7
committed
8e90dd5

Commits on Aug 14, 2026

Add index_copy_ op for ascend backend
Onisen7
committed
9f7d4f0
Update __init__.py to export index_copy_ op
Onisen7
committed
8c39aed

Commits on Aug 15, 2026

Update torch version to 2.9.1 in backends.yaml

tengqm
authored and
Onisen7
committed
ddd3490
Update numpy dependency to allow versions greater than 2

tengqm
authored and
Onisen7
committed
e1f2ac0
Update upload-artifact action to version 7.0.0

tengqm
authored and
Onisen7
committed
6e13fbf
Update upload-artifact action to version 7.0.0 (#5002)

tengqm
authored and
Onisen7
committed
1a02e73
Revise numpy version requirement comment

tengqm
authored and
Onisen7
committed
fb157e0
Fix mul fallback error (#4999)

0x45f
authored and
Onisen7
committed
510fb24
Add te rmsnorm_fwd and rmsnorm_bwd op (#4640)

0x45f
authored and
Onisen7
committed
e2f9594
Fix rms_norm implementation (#5009)

CLYtou
authored and
Onisen7
committed
29f8c82
fix issue 4055 (#4743)

huangyiqun
authored and
Onisen7
committed
533965b
fix issue 4057 (#4742)

huangyiqun
authored and
Onisen7
committed
2b44d40
fix issue 4058 (#4741)

huangyiqun
authored and
Onisen7
committed
693fbe1
[_sunrise] update sunrise ops (#5024)

Dayrker
authored and
Onisen7
committed
6a04eff
Make upload-artifact step optional in on-demand test workflow (#5007)

tengqm
authored and
Onisen7
committed
d0ad5e8
Fix hygon (#5016)

huangyiqun
authored and
Onisen7
committed
c4d5e4a
[KernelGen][Nvidia] Add float_power_ operator with Triton kernel (#4141)

authored and
Onisen7
committed
98be909
[KernelGen][Nvidia] Add max_unpool3d operator with Triton kernel (#4341)

authored and
Onisen7
committed
ff48da8
Pick MXFP4 MoE block_m by padding cost, not an absolute M cutoff (#5027)

zeroherolin
authored and
Onisen7
committed
2863040
[KernelGen][Nvidia] Add arccosh_ operator with Triton kernel (#4002)

authored and
Onisen7
committed
597690b
[KernelGen][Nvidia] Add _native_batch_norm_legit_functional operator with Triton kernel (#4003)

authored and
Onisen7
committed
9f1dce7
Add wheel to build tools (#5047)

tengqm
authored and
Onisen7
committed
6311c73
fix(config): lazy import c_operators — detect .so without loading (#4959)

tengqm
authored and
Onisen7
committed
7bc7d61
Add wheel to setup.sh for some backends (#5048)

tengqm
authored and
Onisen7
committed
b443219
[ENFLAME]enflame_to_FlagOS_20260729 (#5044)

chongzhouyang
authored and
Onisen7
committed
1915247
fix(utils): add fallback_j1 (#5038)

103yiran
authored and
Onisen7
committed
a93c514
[QC] Optimize flash_mla_with_kvcache for model1 (#5010)

monellz
authored and
Onisen7
committed
1198017
[Runtime][AMD] Add RDNA4 autotune configs and match arch by exact GPU target (#4991)

authored and
Onisen7
committed
5624142
fix: replace triton_src kernel refs with ops/ paths for contiguous and zeros (#4931)

yysheng26
authored and
Onisen7
committed
159c46b
fix(hygon): add fp64 path to mm (#4799)

douxetpur
authored and
Onisen7
committed
1ddf2d1
[KernelGen][Nvidia] Add arctan2 operator with Triton kernel (#4006)

authored and
Onisen7
committed
f3d8340
[KernelGen][Nvidia] Add special_legendre_polynomial_p operator with Triton kernel (#4042)

authored and
Onisen7
committed
b8bd04a
[KernelGen][Nvidia] Add erfc and special_erfc operator with Triton kernel (#4041)

authored and
Onisen7
committed
a4544a7
Add _pdist_forward operator with Triton kernel (#3948)

authored and
Onisen7
committed
bdf878c
[KernelGen][Nvidia] Add log_sigmoid_forward operator with Triton kernel (#3874)

authored and
Onisen7
committed
c027c92
build(deps): bump the actions-minor group across 1 directory with 3 updates (#5071)

dependabot[bot]
authored and
Onisen7
committed
65a422e
Add logging.debug in common ops (#5057)

0x45f
authored and
Onisen7
committed
4a85be2
Fix CI job to runner mapping (#5049)

tengqm
authored and
Onisen7
committed
cb660ae
Fix import error in ascend (#5101)

0x45f
authored and
Onisen7
committed
d75eca0
[Ascend] Fix ascend ops when running Qwen3.6 vLLM with flagtree ascend3.5 (#5113)

zhzhcookie
authored and
Onisen7
committed
4aac8eb
[KMCompiler][Nvidia] add operator linalg_lu_factor (#4945)

YangLong114514
authored and
Onisen7
committed
e81f89c
[KernelGen][Nvidia] Add subtract operator with Triton kernel (#4512)

authored and
Onisen7
committed
16768a6
[KernelGen][Nvidia] Move sign operator to ops (#5030)

authored and
Onisen7
committed
980e273
[KernelGen][Nvidia] Move sgn operator to ops (#4990)

authored and
Onisen7
committed
5cb4d07
[KernelGen][Nvidia] Add mish_backward operator with Triton kernel (#4362)

authored and
Onisen7
committed
a78e857
[KernelGen][Nvidia] Add lstm operator with Triton kernel (#4288)

authored and
Onisen7
committed
b41ca5e
[KernelGen][Nvidia] Add miopen_batch_norm_backward operator with Triton kernel (#4353)

authored and
Onisen7
committed
4ff3b69
[KernelGen][Nvidia] Add igamma_ operator with Triton kernel (#4145)

authored and
Onisen7
committed
e8857a0
Optimize index_add for contiguous suffix layouts (#4670)

ZYX223
authored and
Onisen7
committed
b32b27d
add Devcontainer support (#5137)

authored and
Onisen7
committed
2d4e5af
[KMCompiler][Nvidia] opt for log_normal_ operator (#4621)

littlekid-zxt
authored and
Onisen7
committed
c4c12c6
[KMCompiler] Add Triton nonzero_static operator (#4661)

chenkc2025
authored and
Onisen7
committed
e4e3d7d
[KMCompiler][NVIDIA][Hygon] Add nansum operator with triton kernel (#4804)

Willing2329
authored and
Onisen7
committed
09ade37
[KMCompiler][NVIDIA] Add NVIDIA cholesky solve operator (#5011)

YangLong114514
authored and
Onisen7
committed
671bc3d
[KernelGen][Nvidia] Add _cdist_forward operator with Triton kernel (#4022)

authored and
Onisen7
committed
6b32024
[KernelGen][Nvidia] Add _convert_weight_to_int4pack operator with Triton kernel (#4025)

authored and
Onisen7
committed
92e8bfb
[KernelGen][Nvidia] Add binary_cross_entropy_backward operator with Triton kernel (#5090)

authored and
Onisen7
committed
4b67150
[KernelGen][Nvidia] Add expand_as operator (#5041)

authored and
Onisen7
committed
748d192
[KernelGen][Nvidia] Add flatten operator (#5039)

authored and
Onisen7
committed
da70a36
[KernelGen][Nvidia] Add unfold operator (#5037)

authored and
Onisen7
committed
6c8da19
[KernelGen][Nvidia] Add unsafe_chunk operator with Triton kernel (#5028)

authored and
Onisen7
committed
d9ceb2b
[KernelGen][Nvidia] Add eq_ operator with Triton kernel (#5108)

ShawnsYing
authored and
Onisen7
committed
98c68c3
Optimize wna16 MoE main loop: one mma per K tile instead of eight (#5140)

zeroherolin
authored and
Onisen7
committed
8d6c4a6
[KernelGen][Nvidia] Add special_erfcx operator with Triton kernel (#3965)

authored and
Onisen7
committed
7dd1ce1
[KernelGen][Nvidia] Add alias operator (#5043)

authored and
Onisen7
committed
1425b7f
fix adaptive_avg_pool2d (#4952)

fpzh2011
authored and
Onisen7
committed
e4cdfad
TEST CASE FIX -- iluvatar (#5131)

authored and
Onisen7
committed
da5202b
[TSINGMICRO] update tsingmicro backend to 0804 (#5181)

tsingmicro-public-e
authored and
Onisen7
committed
2912c62
fix(mul): gate optimized path on active backend device, not hardcoded "cuda" (#5130)

tengqm
authored and
Onisen7
committed
64520d3
fix(apply_rotary_pos_emb): fix cos/sin error (#5073)

103yiran
authored and
Onisen7
committed
501ece7
fix: resolve_conj performance test failure on mthreads backend. (#5109)

lyujheng
authored and
Onisen7
committed
0581a6c
fix: resolve copy op accuracy failure on mthreads backend (#5107)

lyujheng
authored and
Onisen7
committed
eab9cd2
fix: vector_norm accuracy test failure on mthreads backend (#5106)

lyujheng
authored and
Onisen7
committed
08e124b
fix: mul accuracy and performance on mthreads (#5111)

lyujheng
authored and
Onisen7
committed
92c2d15
[KernelGen][Nvidia] Add chunk operator (#5033)

authored and
Onisen7
committed
96a644e
[KernelGen][Nvidia] Add isposinf operator with Triton kernel (#3923)

authored and
Onisen7
committed
ba143f6
[KernelGen][Nvidia] Add _make_dep_token operator with Triton kernel (#4026)

authored and
Onisen7
committed
9285f40
[KernelGen][Nvidia] Add view_as_complex operator with Triton kernel (#5031)

LoserCheems
authored and
Onisen7
committed
ae5223e
[KernelGen][Nvidia] Move hardtanh_ operator from experimental to ops (#5115)

authored and
Onisen7
committed
94f7b67
[KernelGen][Nvidia] Add _reshape_alias operator (#5121)

authored and
Onisen7
committed
e17b13d
[KernelGen][Nvidia] Add less_ operator with Triton kernel (#5126)

authored and
Onisen7
committed
b35e06a
[KernelGen][Nvidia] Add less_equal_ operator with Triton kernel (#5124)

authored and
Onisen7
committed
834623b
[KernelGen][Nvidia] Add le_ operator with Triton kernel (#5122)

authored and
Onisen7
committed
03bffd6
fix: register empty as empty.memory_format so torch.empty uses FlagGems (#4929)

authored and
Onisen7
committed
bb785ed
fix: recognize vLLM stable ABI extensions (#5150)

CherryLemon
authored and
Onisen7
committed
f370b85
[_sunrise] update sunrise ops (#5238)

Dayrker
authored and
Onisen7
committed
02583cd
Fix IX gate by pinning numpy version < 2 (#5244)

tengqm
authored and
Onisen7
committed
093ff90
fix: index_add accuracy failure on mthreads (#5246)

lyujheng
authored and
Onisen7
committed
7d388b9
[KernelGen][metax] Fix num_warps exceeding hardware thread limit for softmax, logsumexp, dropout, exponential_, rand, randn, and uniform (#5258)

yzw1128
authored and
Onisen7
committed
7c5f531
Revert "[KernelGen][Nvidia] Add view_as_complex operator with Triton kernel (…" (#5262)

103yiran
authored and
Onisen7
committed
860d923
[KernelGen][Nvidia] Move take operator to ops (#5070)

authored and
Onisen7
committed
95ee7e3
[KernelGen][Nvidia] Move log_ operator to ops (#5045)

authored and
Onisen7
committed
1901aff
Ship flaggems-setup as a top-level module and auto-install a compiler (#5207)

tengqm
authored and
Onisen7
committed
7d1d43c
fix(cpp): disable aten dispatch + add PrivateUse1 guard in fallback (#4962)

tengqm
authored and
Onisen7
committed
530cc20
fix(cmake): add global include_directories for NPU backend (#4960)

tengqm
authored and
Onisen7
committed
6921cec
fix: add tsingmicro/txda to device detection (#5056)

authored and
Onisen7
committed
fe01a9d
docs(release): place dev tag on an empty commit after release, not on release tag (#4912)

tengqm
authored and
Onisen7
committed
28e68a2
Upgrade triton for Sunrise (#5268)

tengqm
authored and
Onisen7
committed
dd8c4b4
#19382 tang -> ptpu (#5267)

Dayrker
authored and
Onisen7
committed
d2f7f99
Bump wheel version for CVE (#5271)

tengqm
authored and
Onisen7
committed
5dc305e
Revert "#19382 tang -> ptpu (#5267)" (#5278)

Dayrker
authored and
Onisen7
committed
83fa63b
[KernelGen][Nvidia] Add native_layer_norm operator with Triton kernel (#5180)

authored and
Onisen7
committed
91fc7df
[KernelGen][Nvidia] Add _flash_attention_forward operator with Triton kernel (#5175)

authored and
Onisen7
committed
ac44f07
[KernelGen][Nvidia] Add _scaled_dot_product_flash_attention operator with Triton kernel (#5164)

authored and
Onisen7
committed
0a92a62
[KernelGen][Nvidia] Add lift_fresh operator with Triton kernel (#5148)

authored and
Onisen7
committed
0ffc01d
[KernelGen][Nvidia] Move fix_ operator to ops (#5158)

authored and
Onisen7
committed
3232b2f
[KernelGen][Nvidia] Move hypot_ operator to ops (#5157)

authored and
Onisen7
committed
12decb9
[KernelGen][Nvidia] Move heaviside_ operator to ops (#5156)

authored and
Onisen7
committed
bf9c414
[KernelGen][Nvidia] Move huber_loss operator to ops (#5154)

authored and
Onisen7
committed
d81e24d
[KernelGen][Nvidia] Move arccosh operator to ops (#5163)

authored and
Onisen7
committed
f6dbaa7
[KernelGe][Nvidia] Move hardsigmoid operator to ops (#5149)

authored and
Onisen7
committed
e49a5f5
[KernelGen][Nvidia] Add _cudnn_attention_forward operator with Triton kernel (#5160)

authored and
Onisen7
committed
a69d0df
[KernelGen][Nvidia] Add _adaptive_avg_pool2d_backward operator with Triton kernel (#5084)

Yukun-Cui
authored and
Onisen7
committed
dd31f12
[KernelGen][Nvidia] Add replication_pad3d_backward operator with Triton kernel (#5078)

Yukun-Cui
authored and
Onisen7
committed
5c9bfdd
[KernelGen][Nvidia] Add _weight_norm operator with Triton kernel (#5080)

Yukun-Cui
authored and
Onisen7
committed
d4ba514
[KernelGen][Nvidia] Add replication_pad2d operator with Triton kernel (#5172)

authored and
Onisen7
committed
2bf1b7f
[KernelGen][Nvidia] Add reflection_pad2d_backward operator with Triton kernel (#5165)

authored and
Onisen7
committed
b4bb015
[KernelGen][Nvidia] Add _fused_rms_norm operator with Triton kernel (#5152)

authored and
Onisen7
committed
ddcfba1
[KernelGen][Nvidia] Add cholesky_inverse operator with Triton kernel (#5161)

authored and
Onisen7
committed
0a72e92
[KernelGen][Nvidia] Move hardshrink operator to ops (#5155)

authored and
Onisen7
committed
af679cf
[KernelGen][Nvidia] Add max_pool2d_with_indices_backward operator (#5159)

authored and
Onisen7
committed
59ae4f5
fix(utils): add pure-triton nextafter fallback for non-CUDA backends (#5088)

tengqm
authored and
Onisen7
committed
e00a819
[Fix] fix enflame backend operator bug. (#5302)

chaaa-a
authored and
Onisen7
committed
95d14b0
[KMCompiler][Nvidia]Add Polygamma Operator (#4500)

cheersluvs
authored and
Onisen7
committed
9449e68
[KMCompiler][Nvidia] Add pairwise_distance operator with triton kernel. (#4949)

Caeruleann
authored and
Onisen7
committed
519d7e9
[KMCompiler][Ascend]Add linalg_lstsq NPU backend support (#5173)

cheersluvs
authored and
Onisen7
committed
044bc5f
[KMCompiler] [Nvidia] add linalg_lu_factor_ex operator with Triton kernel (#5209)

YangLong114514
authored and
Onisen7
committed
819d4ee
Contribution doc update (#5307)

103yiran
authored and
Onisen7
committed
9a2469b
Drop experimental operator sinc (#4810)

authored and
Onisen7
committed
f40ff8b
[KernelGen][Nvidia] Add special_bessel_j0 operator with Triton kernel (#4039)

authored and
Onisen7
committed
c9425b2
fix sinc (#4892)

authored and
Onisen7
committed
ff9a6be
Fix cumsum ZeroDivisionError on empty tensor (#4541)

DannyP0
authored and
Onisen7
committed
35154c6
fix add op test (#5325)

huangyiqun
authored and
Onisen7
committed
fd6e811
[mthreads] Optimize flip for multi-dim tensors with trailing non-flipped dims (#5261)

Oslomayor
authored and
Onisen7
committed
e818d44
add logger debug (#5210)

douxetpur
authored and
Onisen7
committed
27638a8
cambricon: update and fix kernel (#4435)

authored and
Onisen7
committed
f87c075
Fix code style (#5331)

0x45f
authored and
Onisen7
committed
d0adba8
[KernelGen][Nvidia] Add lu_unpack operator with Triton kernel (#5072)

LoserCheems
authored and
Onisen7
committed
cea9f5a
[KernelGen][Nvidia] Add block_diag operator with Triton kernel (#5068)

LoserCheems
authored and
Onisen7
committed
3880202
[KernelGen][Nvidia] Add addbmm operator with Triton kernel (#5055)

LoserCheems
authored and
Onisen7
committed
0085515
fix(ops): skip triton kernels for complex dtypes in empty/lift_fresh/randn (#5332)

CLYtou
authored and
Onisen7
committed
9821f07
[KernelGen][Nvidia] Add scatter_add operator with Triton kernel (#4502)

authored and
Onisen7
committed
8bf14ec
fix(conf): add missing info for special_log_softmax (#5333)

103yiran
authored and
Onisen7
committed
0feaee0
fused: downcast slot_mapping to int32 for enflame on reshape_and_cache_flash (#5345)

authored and
Onisen7
committed
01effe8
Append commit id to __version__ for dev builds (#5346)

tengqm
authored and
Onisen7
committed
e551554
fix(utils): add pure-triton j0 and log2 fallbacks for non-CUDA backends (#5348)

tengqm
authored and
Onisen7
committed
f2e0546
Split multi-version vendors into versioned backend keys (#5350)

tengqm
authored and
Onisen7
committed
97c599c
fix(cambricon): tolerate triton forks lacking TRITON_MAX_TENSOR_NUMEL (#5351)

tengqm
authored and
Onisen7
committed
902c0ec
Drop experimental _unsafe_view operator (#4723)

tengqm
authored and
Onisen7
committed
f21a6b2
[KernelGen][Nvidia] Add _has_compatible_shallow_copy_type operator (#5280)

authored and
Onisen7
committed
84e47e3
[KernelGen][Nvidia] Add sym_storage_offset operator with Triton kernel (#5279)

authored and
Onisen7
committed
82acdbf
fix: div accuracy test failures on mthreads backend (#5295)

lyujheng
authored and
Onisen7
committed
b40b641
fix: add mthreads mul specialization for complex dtype operations (#5294)

lyujheng
authored and
Onisen7
committed
a38cc4f
[AUTOTUNER] Add (XGB+GA) FlagTune support (via FlagTree) (#4740)

HenryRenYz
authored and
Onisen7
committed
95fc69a
Metax fix shared ops (#5334)

CLYtou
authored and
Onisen7
committed
5d86d41
Metax c550 operator fixes (#5337)

CLYtou
authored and
Onisen7
committed
0a013a7
[MTHREADS] optimize conv1d_padding autotune configs (#5343)

Oslomayor
authored and
Onisen7
committed
30142d6
[MTHREADS] optimize conv3d_padding autotune configs (#5353)

Oslomayor
authored and
Onisen7
committed
6941300
[KernelGen][Nvidia] Add special_i1e operator with Triton kernel (#5276)

authored and
Onisen7
committed
a2af686
Fix special_chebyshev_polynomial_w ut (#5352)

0x45f
authored and
Onisen7
committed
bb8907d
[MTHREADS] Fix flip op: correct multi-dim index decoding and per-dim flip logic (#5374)

lyujheng
authored and
Onisen7
committed
efbb669
[KernelGen][Nvidia] Add special_exp2 operator with Triton kernel (#5275)

authored and
Onisen7
committed
0a0e318
[KernelGen][Nvidia] Add special_erf operator with Triton kernel (#5274)

authored and
Onisen7
committed
bc28753
fix(ascend): align full_like with updated full_kernel API (#5308)

yzw1128
authored and
Onisen7
committed
0a9d125
fix: prevent run_tests.py deadlock when worker crashes (#5358)

authored and
Onisen7
committed
5d0f16c
[KMCompiler][Ascend] Add polygamma NPU backend support (#4809)

cheersluvs
authored and
Onisen7
committed
64fa801
[KMCompiler] Add igammac operator with Triton kernel (#4860)

LittleShun1214
authored and
Onisen7
committed
d4db72b
[KMCompiler][Ascend] nonzero static ascend (#4921)

chenkc2025
authored and
Onisen7
committed
736fdb8
[KMCompiler][NVIDIA][Hygon]Add replication_pad2d_backward operator with triton kernel (#5146)

Willing2329
authored and
Onisen7
committed
13a87b4
[KMCompiler] Add Iluvatar cholesky solve backend (#5242)

YangLong114514
authored and
Onisen7
committed
2422a2e
[KMCompiler] Add Ascend cholesky solve backend (#5299)

YangLong114514
authored and
Onisen7
committed
228a826
[KMCompiler] Add MetaX cholesky solve backend (#5247)

YangLong114514
authored and
Onisen7
committed
c1bc66c
fix: increase op test job timeout (#5391)

wuwentao
authored and
Onisen7
committed
e79b393
Use Triton for metax CI (#5396)

tengqm
authored and
Onisen7
committed
3bcb766
fix testcase iluvatar (#5375)

authored and
Onisen7
committed
1c63f26
Restore flagtree backend for metax CI (#5405)

tengqm
authored and
Onisen7
committed
ac5cca1
restore: use stable_two_path version of index_copy_
Onisen7
committed
412cfba
Update index_copy_.py
Onisen7
authored
22c7c04
Update index_copy_.py
Onisen7
authored
2c6783e
Merge branch 'master' into add-ascend-index-copy
Onisen7
committed
a80ea2b
Delete src/flag_gems/runtime/backend/_nvidia/ops/index_copy_.py
Onisen7
authored
da72fcc
Delete src/flag_gems/runtime/backend/_hygon/ops/index_copy_.py
Onisen7
authored
49e64b7
Merge branch 'master' into add-ascend-index-copy
Onisen7
authored
c1e2431