Skip to content

Commit c1e2431

Browse files
authored
Merge branch 'master' into add-ascend-index-copy
2 parents 49e64b7 + bfeca79 commit c1e2431

159 files changed

Lines changed: 19898 additions & 2009 deletions

File tree

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

.github/CODEOWNERS

Lines changed: 6 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -20,12 +20,11 @@
2020
/src/flag_gems/* @0x45f @huangyiqun @103yiran @douxetpur @bin913 @w1120029931-bit @hellojack163 @Caeruleann @tengqm
2121
/src/flag_gems/runtime/backend/*
2222

23-
/tests/* @0x45f @huangyiqun @103yiran @douxetpur @bin913 @w1120029931-bit @hellojack163 @Caeruleann @tengqm
23+
/benchmark/* @0x45f @huangyiqun @103yiran @douxetpur @bin913 @w1120029931-bit @hellojack163 @Caeruleann
24+
/tests/* @0x45f @huangyiqun @103yiran @douxetpur @bin913 @w1120029931-bit @hellojack163 @Caeruleann
2425
/modules_tests/* @0x45f @huangyiqun @103yiran @douxetpur @bin913 @w1120029931-bit @hellojack163 @Caeruleann
25-
/benchmark/* @0x45f @huangyiqun @103yiran @douxetpur @bin913 @w1120029931-bit @hellojack163 @Caeruleann @tengqm
26-
/docs/* @0x45f @huangyiqun @103yiran @tengqm
27-
/container/* @0x45f @huangyiqun @103yiran @tengqm
26+
/docs/* @0x45f @huangyiqun @103yiran
27+
/tools/* @0x45f @huangyiqun @103yiran
2828

29-
/lib/* @0x45f @huangyiqun @103yiran @douxetpur @bin913 @w1120029931-bit @hellojack163 @Caeruleann
30-
/ctests/* @0x45f @huangyiqun @103yiran @douxetpur @bin913 @w1120029931-bit @hellojack163 @Caeruleann
31-
/include/* @0x45f @huangyiqun @103yiran @douxetpur @bin913 @w1120029931-bit @hellojack163 @Caeruleann
29+
/cpp/* @0x45f @huangyiqun @103yiran @douxetpur @bin913 @w1120029931-bit @hellojack163 @Caeruleann
30+
/triton_src/* @0x45f @huangyiqun @103yiran @douxetpur @bin913 @w1120029931-bit @hellojack163 @Caeruleann

.github/backends.json

Lines changed: 8 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -48,9 +48,16 @@
4848
"gpu_check": "tools/gpu_check_kunlunxin.sh",
4949
"enabled": true
5050
},
51+
{
52+
"backend": "metax-maca3720",
53+
"runner_label": "maca372",
54+
"label": "vendor/MetaX",
55+
"gpu_check": "tools/gpu_check_metax.sh",
56+
"enabled": true
57+
},
5158
{
5259
"backend": "metax-maca3810",
53-
"runner_label": "metax",
60+
"runner_label": "maca381",
5461
"label": "vendor/MetaX",
5562
"gpu_check": "tools/gpu_check_metax.sh",
5663
"enabled": true

.github/workflows/command.yaml

Lines changed: 50 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -44,16 +44,19 @@ jobs:
4444
skip_reviews: "true"
4545
fork_review_bypass: "true"
4646
permissions: "write,admin"
47+
param_separator: " "
4748

4849
- id: parse-command
4950
if: ${{ steps.check.outputs.continue == 'true' }}
5051
env:
5152
PULL_NUMBER: ${{ github.event.issue.number }}
53+
COMMENT_BODY: ${{ steps.check.outputs.comment_body }}
5254
run: |
5355
# Parse command arguments
54-
params=${{ steps.check.outputs.params }}
55-
op=$(echo $params | cut -d ':' -f 1)
56-
runner=$(echo $params | cut -d ':' -f 2)
56+
# Comment format: "/test <op>:<runner>", e.g. "/test constant_pad_nd:mthreads"
57+
params="${COMMENT_BODY#/test }"
58+
op="${params%%:*}"
59+
runner="${params##*:}"
5760
echo "OP_ID=${op}" >> $GITHUB_OUTPUT
5861
echo "RUNNER=${runner}" >> $GITHUB_OUTPUT
5962
echo "PR_NUMBER=${PULL_NUMBER}" >> $GITHUB_OUTPUT
@@ -96,6 +99,26 @@ jobs:
9699
BACKEND="nvidia-cuda128"
97100
echo "COMPILER=triton" >> $GITHUB_ENV
98101
;;
102+
ascend)
103+
VENDOR="ascend"
104+
# The ascend runner pool contains nodes with two different CANN
105+
# toolkits (8.5.x and 9.0.x). The job may land on either node, so
106+
# detect the CANN version on the node and pick the matching
107+
# backend: 8.5.x -> ascend-cann850, 9.0.x -> ascend-cann900
108+
CANN_VERSION=$(ls -d /usr/local/Ascend/cann-* 2>/dev/null | sed 's|.*/cann-||' | sort -V | tail -1)
109+
case "${CANN_VERSION}" in
110+
8.5*)
111+
BACKEND="ascend-cann850"
112+
;;
113+
9.0*)
114+
BACKEND="ascend-cann900"
115+
;;
116+
*)
117+
echo "::error::Cannot detect CANN version on ascend runner (got '${CANN_VERSION}')" >&2
118+
exit 1
119+
;;
120+
esac
121+
;;
99122
cann850)
100123
VENDOR="ascend"
101124
BACKEND="ascend-cann850"
@@ -109,6 +132,30 @@ jobs:
109132
BACKEND="mthreads-musa520"
110133
;;
111134
metax)
135+
VENDOR="metax"
136+
# The metax runner pool contains nodes with two different MACA
137+
# drivers (3.7.x and 3.8.x). The job may land on either node, so
138+
# detect the driver on the node and pick the matching backend:
139+
# 3.7.x -> metax-maca3720, 3.8.x -> metax-maca3810
140+
MACA_VERSION=$(mx-smi 2>/dev/null | awk -F': ' '/MACA Version/{split($2,a," "); print a[1]}')
141+
case "${MACA_VERSION}" in
142+
3.7.*)
143+
BACKEND="metax-maca3720"
144+
;;
145+
3.8.*)
146+
BACKEND="metax-maca3810"
147+
;;
148+
*)
149+
echo "::error::Cannot detect MACA version on metax runner (got '${MACA_VERSION}')" >&2
150+
exit 1
151+
;;
152+
esac
153+
;;
154+
maca372)
155+
VENDOR="metax"
156+
BACKEND="metax-maca3720"
157+
;;
158+
maca381)
112159
VENDOR="metax"
113160
BACKEND="metax-maca3810"
114161
;;

.github/workflows/random-test.yaml

Lines changed: 44 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -77,6 +77,26 @@ jobs:
7777
VENDOR="nvidia"
7878
BACKEND="nvidia-cuda133"
7979
;;
80+
ascend)
81+
VENDOR="ascend"
82+
# The ascend runner pool contains nodes with two different CANN
83+
# toolkits (8.5.x and 9.0.x). The job may land on either node, so
84+
# detect the CANN version on the node and pick the matching
85+
# backend: 8.5.x -> ascend-cann850, 9.0.x -> ascend-cann900
86+
CANN_VERSION=$(ls -d /usr/local/Ascend/cann-* 2>/dev/null | sed 's|.*/cann-||' | sort -V | tail -1)
87+
case "${CANN_VERSION}" in
88+
8.5*)
89+
BACKEND="ascend-cann850"
90+
;;
91+
9.0*)
92+
BACKEND="ascend-cann900"
93+
;;
94+
*)
95+
echo "::error::Cannot detect CANN version on ascend runner (got '${CANN_VERSION}')" >&2
96+
exit 1
97+
;;
98+
esac
99+
;;
80100
cann850)
81101
VENDOR="ascend"
82102
BACKEND="ascend-cann850"
@@ -86,6 +106,30 @@ jobs:
86106
BACKEND="ascend-cann900"
87107
;;
88108
metax)
109+
VENDOR="metax"
110+
# The metax runner pool contains nodes with two different MACA
111+
# drivers (3.7.x and 3.8.x). The job may land on either node, so
112+
# detect the driver on the node and pick the matching backend:
113+
# 3.7.x -> metax-maca3720, 3.8.x -> metax-maca3810
114+
MACA_VERSION=$(mx-smi 2>/dev/null | awk -F': ' '/MACA Version/{split($2,a," "); print a[1]}')
115+
case "${MACA_VERSION}" in
116+
3.7.*)
117+
BACKEND="metax-maca3720"
118+
;;
119+
3.8.*)
120+
BACKEND="metax-maca3810"
121+
;;
122+
*)
123+
echo "::error::Cannot detect MACA version on metax runner (got '${MACA_VERSION}')" >&2
124+
exit 1
125+
;;
126+
esac
127+
;;
128+
maca372)
129+
VENDOR="metax"
130+
BACKEND="metax-maca3720"
131+
;;
132+
maca381)
89133
VENDOR="metax"
90134
BACKEND="metax-maca3810"
91135
;;

benchmark/base.py

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -310,8 +310,9 @@ def get_latency(self, op, *args, **kwargs):
310310
do_bench = triton.backends.ascend.testing.do_bench_npu
311311
latency = do_bench(
312312
fn,
313-
warmup=Config.warm_up,
314-
active=Config.repetition,
313+
# do_bench_npu requires iterations, rather than duration
314+
# warmup=Config.warm_up,
315+
# active=Config.repetition,
315316
)
316317
else:
317318
do_bench = triton.testing.do_bench

benchmark/core_shapes.yaml

Lines changed: 32 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -65,42 +65,6 @@ svd:
6565
- [256, 256]
6666
shape_desc: "(*B), M, N"
6767

68-
linalg_lu_factor:
69-
shapes:
70-
- [16, 16]
71-
- [32, 32]
72-
- [64, 64]
73-
- [128, 128]
74-
- [256, 256]
75-
- [1024, 512]
76-
- [32, 16]
77-
- [16, 32]
78-
- [128, 64]
79-
- [64, 128]
80-
- [4, 32, 32]
81-
- [128, 16, 16]
82-
- [1024, 512, 512]
83-
- [4096, 512, 512]
84-
shape_desc: "input shape"
85-
86-
linalg_lu_factor_out:
87-
shapes:
88-
- [16, 16]
89-
- [32, 32]
90-
- [64, 64]
91-
- [128, 128]
92-
- [256, 256]
93-
- [1024, 512]
94-
- [32, 16]
95-
- [16, 32]
96-
- [128, 64]
97-
- [64, 128]
98-
- [4, 32, 32]
99-
- [128, 16, 16]
100-
- [1024, 512, 512]
101-
- [4096, 512, 512]
102-
shape_desc: "input shape"
103-
10468
segment_reduce: &segment_reduce_shapes
10569
shapes:
10670
- [1048576]
@@ -308,6 +272,38 @@ grouped_mm:
308272
- [64, 2048, 128]
309273
shape_desc: "Groups, N, K" # shapes are defined as (Groups, N, K)
310274

275+
sparse_sampled_addmm:
276+
shapes:
277+
- [2, 128, 128, 128]
278+
- [4, 256, 256, 256]
279+
- [8, 512, 512, 512]
280+
- [16, 512, 512, 512]
281+
- [4, 1024, 1024, 1024]
282+
- [16, 1024, 1024, 1024]
283+
- [8, 2048, 2048, 64]
284+
- [2, 2048, 2048, 128]
285+
- [4, 4096, 4096, 128]
286+
- [1, 4096, 4096, 64]
287+
- [2, 4096, 1024, 512]
288+
- [2, 1024, 1000, 384]
289+
shape_desc: "B, M, N, K"
290+
291+
sparse_sampled_addmm_out:
292+
shapes:
293+
- [2, 128, 128, 128]
294+
- [4, 256, 256, 256]
295+
- [8, 512, 512, 512]
296+
- [16, 512, 512, 512]
297+
- [4, 1024, 1024, 1024]
298+
- [16, 1024, 1024, 1024]
299+
- [8, 2048, 2048, 64]
300+
- [2, 2048, 2048, 128]
301+
- [4, 4096, 4096, 128]
302+
- [1, 4096, 4096, 64]
303+
- [2, 4096, 1024, 512]
304+
- [2, 1024, 1000, 384]
305+
shape_desc: "B, M, N, K"
306+
311307
MvAndOuterBenchmark:
312308
shapes:
313309
- [384, 384]
Lines changed: 63 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,63 @@
1+
# Copyright 2026 FlagOS Contributors
2+
#
3+
# Licensed under the Apache License, Version 2.0 (the "License");
4+
# you may not use this file except in compliance with the License.
5+
# You may obtain a copy of the License at
6+
#
7+
# http://www.apache.org/licenses/LICENSE-2.0
8+
#
9+
# Unless required by applicable law or agreed to in writing, software
10+
# distributed under the License is distributed on an "AS IS" BASIS,
11+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+
# See the License for the specific language governing permissions and
13+
# limitations under the License.
14+
15+
import pytest
16+
import torch
17+
import torch.nn.functional as F
18+
19+
from . import base, consts
20+
21+
22+
class NormBenchmark(base.GenericBenchmark):
23+
def set_more_shapes(self):
24+
return [
25+
# 3D shapes represented as [batch_size, channels, hidden_size]
26+
(16, 16, 64),
27+
(16, 16, 1024),
28+
(16, 16, 4098),
29+
# 4D shapes represented as [batch_size, channels, H, W]
30+
(1, 8, 4, 4),
31+
(16, 8, 128, 128),
32+
]
33+
34+
35+
def native_batch_norm_legit_no_training_input_fn(shape, dtype, device):
36+
C = shape[1]
37+
inp = torch.randn(shape, dtype=dtype, device=device)
38+
weight = torch.randn((C,), dtype=dtype, device=device)
39+
bias = torch.randn((C,), dtype=dtype, device=device)
40+
running_mean = torch.randn((C,), dtype=dtype, device=device)
41+
running_var = torch.abs(torch.randn((C,), dtype=dtype, device=device)) + 0.1
42+
momentum = 0.1
43+
eps = 1e-5
44+
yield inp, weight, bias, running_mean, running_var, momentum, eps
45+
46+
47+
def torch_native_batch_norm_legit_no_training(
48+
inp, weight, bias, running_mean, running_var, momentum, eps
49+
):
50+
return F.batch_norm(
51+
inp, running_mean, running_var, weight, bias, training=False, eps=eps
52+
)
53+
54+
55+
@pytest.mark.native_batch_norm_legit_no_training
56+
def test_native_batch_norm_legit_no_training():
57+
bench = NormBenchmark(
58+
input_fn=native_batch_norm_legit_no_training_input_fn,
59+
op_name="native_batch_norm_legit_no_training",
60+
torch_op=torch_native_batch_norm_legit_no_training,
61+
dtypes=consts.FLOAT_DTYPES,
62+
)
63+
bench.run()

benchmark/test_absolute.py

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -24,3 +24,14 @@ def test_absolute():
2424
op_name="absolute", torch_op=torch.absolute, dtypes=consts.FLOAT_DTYPES
2525
)
2626
bench.run()
27+
28+
29+
@pytest.mark.absolute_
30+
def test_absolute_inplace():
31+
bench = base.UnaryPointwiseBenchmark(
32+
op_name="absolute_",
33+
torch_op=lambda x: x.absolute_(),
34+
dtypes=consts.FLOAT_DTYPES,
35+
is_inplace=True,
36+
)
37+
bench.run()
Lines changed: 48 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,48 @@
1+
# Copyright 2026 FlagOS Contributors
2+
#
3+
# Licensed under the Apache License, Version 2.0 (the "License");
4+
# you may not use this file except in compliance with the License.
5+
# You may obtain a copy of the License at
6+
#
7+
# http://www.apache.org/licenses/LICENSE-2.0
8+
#
9+
# Unless required by applicable law or agreed to in writing, software
10+
# distributed under the License is distributed on an "AS IS" BASIS,
11+
# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12+
# See the License for the specific language governing permissions and
13+
# limitations under the License.
14+
15+
from typing import Generator
16+
17+
import pytest
18+
import torch
19+
20+
from . import base, consts
21+
22+
23+
class AdaptiveMaxPool2dBackwardBenchmark(base.Benchmark):
24+
def set_shapes(self, shape_file_path=None):
25+
# Common CNN feature-map shapes paired with representative output sizes.
26+
self.shapes = [
27+
(4, 3, 32, 32, 7, 7),
28+
(8, 64, 56, 56, 7, 7),
29+
(4, 128, 112, 112, 14, 14),
30+
]
31+
self.shape_desc = "N, C, H, W, OUT_H, OUT_W"
32+
33+
def get_input_iter(self, cur_dtype) -> Generator:
34+
for n, c, h, w, out_h, out_w in self.shapes:
35+
inp = torch.randn((n, c, h, w), dtype=cur_dtype, device=self.device)
36+
output, indices = torch.ops.aten.adaptive_max_pool2d(inp, (out_h, out_w))
37+
grad_output = torch.randn_like(output)
38+
yield grad_output, inp, indices
39+
40+
41+
@pytest.mark.adaptive_max_pool2d_backward
42+
def test_adaptive_max_pool2d_backward():
43+
bench = AdaptiveMaxPool2dBackwardBenchmark(
44+
op_name="adaptive_max_pool2d_backward",
45+
torch_op=torch.ops.aten.adaptive_max_pool2d_backward,
46+
dtypes=consts.FLOAT_DTYPES,
47+
)
48+
bench.run()

0 commit comments

Comments
 (0)