Skip to content

Commit ad95157

Browse files
authored
Merge branch 'flagos-ai:master' into moe_align_debug
2 parents ea0c087 + db14b4b commit ad95157

54 files changed

Lines changed: 2479 additions & 73 deletions

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.
Lines changed: 57 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,57 @@
1+
name: gems-experimental-test
2+
3+
on:
4+
push:
5+
branches: [ "master" ]
6+
paths:
7+
- 'src/flag_gems/experimental_ops/**'
8+
pull_request:
9+
branches: [ "master" ]
10+
paths:
11+
- 'src/flag_gems/experimental_ops/**'
12+
13+
env:
14+
CUDA_VISIBLE_DEVICES: 7
15+
http_proxy: ${{ secrets.HTTP_PROXY }}
16+
https_proxy: ${{ secrets.HTTPS_PROXY }}
17+
18+
jobs:
19+
experimental-test-on-hopper:
20+
runs-on: hopper
21+
concurrency:
22+
group: experimental-test-on-hopper-${{ github.event.pull_request.number || github.ref }}
23+
cancel-in-progress: true
24+
steps:
25+
- name: Checkout code
26+
uses: actions/checkout@v4
27+
with:
28+
fetch-depth: 0
29+
30+
- name: FlagGems experimental tests on hopper
31+
shell: bash
32+
run: |
33+
source "/home/zhangzhihui/miniconda3/etc/profile.d/conda.sh"
34+
conda activate flag_gems
35+
source tools/run_command.sh
36+
BASE_SHA=${{ github.event.pull_request.base.sha }}
37+
HEAD_SHA=${{ github.sha }}
38+
39+
echo "Diffing $BASE_SHA...$HEAD_SHA"
40+
41+
changed_ops=$(git diff --name-only $BASE_SHA...$HEAD_SHA | grep '^src/flag_gems/experimental_ops/.*\.py$' || true)
42+
43+
tests=""
44+
for f in $changed_ops; do
45+
base=$(basename "$f" .py)
46+
test_file="src/flag_gems/experimental_ops/exp_tests/${base}_test.py"
47+
if [ -f "$test_file" ]; then
48+
tests="$tests $test_file"
49+
fi
50+
done
51+
52+
if [ -n "$tests" ]; then
53+
echo "Running tests:$tests"
54+
run_command pytest -s $tests
55+
else
56+
echo "No relevant ops changes, skipping tests"
57+
fi

benchmark/test_attention_perf.py

100644100755
Lines changed: 6 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -23,11 +23,14 @@ def set_more_shapes(self):
2323

2424

2525
@pytest.mark.skipif(vendor_name == "kunlunxin", reason="RESULT TODOFIX")
26-
@pytest.mark.skipif(vendor_name == "hygon", reason="RuntimeError")
26+
# @pytest.mark.skipif(vendor_name == "hygon", reason="RuntimeError")
2727
@pytest.mark.scaled_dot_product_attention
2828
@pytest.mark.parametrize("dropout_p", [0.0])
2929
@pytest.mark.parametrize("is_causal", [True, False])
3030
def test_perf_scaled_dot_product_attention(dropout_p, is_causal):
31+
if flag_gems.vendor_name == "hygon":
32+
os.environ["TRITON_HIP_USE_NEW_STREAM_PIPELINE"] = "0"
33+
3134
def scaled_dot_product_attention_kwargs(shape, dtype, device):
3235
query = torch.randn(shape, device=device, dtype=dtype)
3336
key = torch.randn(shape, device=device, dtype=dtype)
@@ -61,6 +64,8 @@ def sdpa_flash(
6164
)
6265
bench.set_gems(flag_gems.scaled_dot_product_attention)
6366
bench.run()
67+
if flag_gems.vendor_name == "hygon":
68+
del os.environ["TRITON_HIP_USE_NEW_STREAM_PIPELINE"]
6469

6570

6671
class FlashMLABenchmark(GenericBenchmark):

benchmark/test_convolution_perf.py

100644100755
Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1,3 +1,5 @@
1+
import os
2+
13
import pytest
24
import torch
35

@@ -103,6 +105,8 @@ def conv2d_input_fn(shape, dtype, device):
103105
"padding": padding,
104106
},
105107

108+
if flag_gems.vendor_name == "hygon":
109+
os.environ["TRITON_HIP_USE_NEW_STREAM_PIPELINE"] = "0"
106110
torch.backends.cudnn.allow_tf32 = False
107111
bench = Conv2DBenchmark(
108112
input_fn=conv2d_input_fn,
@@ -112,6 +116,8 @@ def conv2d_input_fn(shape, dtype, device):
112116
)
113117
bench.set_gems(flag_gems.conv2d)
114118
bench.run()
119+
if flag_gems.vendor_name == "hygon":
120+
del os.environ["TRITON_HIP_USE_NEW_STREAM_PIPELINE"]
115121

116122

117123
class Conv3DBenchmark(GenericBenchmark):

benchmark/test_reduction_perf.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -221,7 +221,7 @@ def test_generic_reduction_benchmark(op_name, torch_op, input_fn, dtypes):
221221
bench.run()
222222

223223

224-
@pytest.mark.skipif(vendor_name == "hygon", reason="RESULT TODOFIX")
224+
# @pytest.mark.skipif(vendor_name == "hygon", reason="RESULT TODOFIX")
225225
@pytest.mark.count_nonzero
226226
def test_perf_count_nonzero():
227227
def count_nonzero_input_fn(shape, dtype, device):

benchmark/test_select_and_slice_perf.py

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -179,6 +179,27 @@ def test_perf_scatter_add():
179179
bench.run()
180180

181181

182+
@pytest.mark.scatter_add_
183+
def test_perf_scatter_add_():
184+
def scatter_input_fn(shape, dtype, device):
185+
input_gen = gather_input_fn(shape, dtype, device)
186+
inp, dim, index = next(input_gen)
187+
src_shape = list(size + 16 for size in index.shape)
188+
src = torch.randn(src_shape, dtype=dtype, device=device)
189+
190+
yield inp, dim, index, src
191+
192+
bench = TensorSelectBenchmark(
193+
op_name="scatter_add_",
194+
torch_op=torch.Tensor.scatter_add_,
195+
input_fn=scatter_input_fn,
196+
get_gbps=gather_scatter_gbps,
197+
dtypes=FLOAT_DTYPES,
198+
)
199+
bench.run()
200+
201+
202+
@pytest.mark.scatter_multiply
182203
@pytest.mark.scatter
183204
def test_perf_scatter_multiply():
184205
bench = TensorSelectBenchmark(

benchmark/test_special_perf.py

100644100755
Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -70,7 +70,7 @@ def test_special_operations_benchmark(op_name, torch_op, dtypes, input_fn):
7070
bench.run()
7171

7272

73-
@pytest.mark.skipif(flag_gems.vendor_name == "hygon", reason="RuntimeError")
73+
# @pytest.mark.skipif(flag_gems.vendor_name == "hygon", reason="RuntimeError")
7474
@pytest.mark.isin
7575
def test_isin_perf():
7676
def isin_input_fn(shape, dtype, device):
@@ -94,7 +94,7 @@ def isin_input_fn(shape, dtype, device):
9494
bench.run()
9595

9696

97-
@pytest.mark.skipif(flag_gems.vendor_name == "hygon", reason="RuntimeError")
97+
# @pytest.mark.skipif(flag_gems.vendor_name == "hygon", reason="RuntimeError")
9898
@pytest.mark.unique
9999
def test_perf_unique():
100100
def unique_input_fn(shape, dtype, device):
@@ -110,7 +110,7 @@ def unique_input_fn(shape, dtype, device):
110110
bench.run()
111111

112112

113-
@pytest.mark.skipif(flag_gems.vendor_name == "hygon", reason="RuntimeError")
113+
# @pytest.mark.skipif(flag_gems.vendor_name == "hygon", reason="RuntimeError")
114114
@pytest.mark.sort
115115
def test_perf_sort():
116116
class SortBenchmark(GenericBenchmark2DOnly):

benchmark/test_tensor_constructor_perf.py

100644100755
Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -184,7 +184,7 @@ def test_tensor_constructor_inplace_benchmark(op_name, torch_op, input_fn):
184184
bench.run()
185185

186186

187-
@pytest.mark.skipif(vendor_name == "hygon", reason="RESULT TODOFIX")
187+
# @pytest.mark.skipif(vendor_name == "hygon", reason="RESULT TODOFIX")
188188
@pytest.mark.randperm
189189
def test_perf_randperm():
190190
if flag_gems.vendor_name == "mthreads":

docs/get_start_with_flaggems.md

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -61,3 +61,18 @@ B = torch.randn((K, N), dtype=torch.float16, device=flag_gems.device)
6161
with flag_gems.use_gems():
6262
C = torch.mm(A, B)
6363
```
64+
65+
66+
## How To Use Experimental Gems
67+
The `experimental_ops` module provides a space for new operators that are not yet ready for production release. Operators in this module are accessible via `flag_gems.experimental_ops.*` and follow the same development patterns as core operators.
68+
```
69+
import flag_gems
70+
71+
# Global enablement
72+
flag_gems.enable()
73+
result = flag_gems.experimental_ops.rmsnorm(*args)
74+
75+
# Or scoped usage
76+
with flag_gems.use_gems():
77+
result = flag_gems.experimental_ops.rmsnorm(*args)
78+
```

src/flag_gems/__init__.py

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,7 @@
66
from flag_gems import testing # noqa: F401
77
from flag_gems import runtime
88
from flag_gems.config import aten_patch_list
9+
from flag_gems.experimental_ops import * # noqa: F403
910
from flag_gems.fused import * # noqa: F403
1011
from flag_gems.logging_utils import setup_flaggems_logging
1112
from flag_gems.modules import * # noqa: F403
@@ -344,6 +345,7 @@ def enable(
344345
("where.self_out", where_self_out),
345346
("zeros", zeros),
346347
("zeros_like", zeros_like),
348+
("scatter_add_", scatter_add_),
347349
("dreglu", dreglu),
348350
("reglu", reglu),
349351
("scaled_softmax_forward", scaled_softmax_forward),
@@ -394,6 +396,12 @@ def __exit__(self, exc_type, exc_val, exc_tb):
394396
logging.root.removeHandler(handler)
395397
logging.basicConfig(level=logging.INFO)
396398

399+
@property
400+
def experimental_ops(self):
401+
import flag_gems.experimental_ops
402+
403+
return flag_gems.experimental_ops
404+
397405

398406
def all_ops():
399407
return current_work_registrar.get_all_ops()
Lines changed: 141 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,141 @@
1+
# Overview
2+
The `experimental_ops` module provides a space for new operators that are not yet ready for production release. Operators in this module are accessible via `flag_gems.experimental_ops.*` and follow the same development patterns as core operators.
3+
4+
# Usage Example
5+
Users can access operators as:
6+
```
7+
import flag_gems
8+
9+
# Global enablement
10+
flag_gems.enable()
11+
result = flag_gems.experimental_ops.your_operator(*args)
12+
13+
# Or scoped usage
14+
with flag_gems.use_gems():
15+
result = flag_gems.experimental_ops.your_operator(*args)
16+
```
17+
18+
19+
# File Structure
20+
```
21+
src/flag_gems/experimental_ops/
22+
├── __init__.py # Module initialization
23+
├── rmsnorm.py # Example operator implementation
24+
├── [other_operators].py # Additional operators
25+
├── exp_tests/ # Accuracy test and performance test
26+
├── __init__.py
27+
├── rmsnorm_test.py
28+
├── [other_operators]_test.py
29+
```
30+
31+
# Adding New Operators
32+
## 1. Create Operator Implementation
33+
Create your operator file in `src/flag_gems/experimental_ops/`:
34+
```
35+
# src/flag_gems/experimental_ops/your_operator.py
36+
from flag_gems.utils import libentry
37+
38+
@libentry()
39+
@triton.autotune(
40+
configs=[...],
41+
key=[...]
42+
)
43+
def your_operator_kernel(...):
44+
# Triton kernel implementation
45+
pass
46+
47+
def your_operator(*args, **kwargs):
48+
# Python wrapper
49+
return your_operator_kernel(*args, **kwargs)
50+
```
51+
52+
## 2. Update Module Exports
53+
Add your operator to `src/flag_gems/experimental_ops/__init__.py` :
54+
```
55+
from .your_operator import your_operator
56+
__all__ = ["rmsnorm", "your_operator"]
57+
```
58+
59+
## 3. Update Main Module
60+
The experimental_ops module is already integrated in the main `__init__.py` . No changes needed there.
61+
62+
63+
# Testing
64+
## Accuracy Tests
65+
Add accuracy test in `exp_tests/your_ops_test.py`:
66+
```
67+
import pytest
68+
import torch
69+
import flag_gems
70+
from tests.accuracy_utils import (
71+
FLOAT_DTYPES,
72+
gems_assert_close,
73+
to_reference,
74+
)
75+
76+
@pytest.mark.your_operator
77+
@pytest.mark.parametrize("shape", [...])
78+
@pytest.mark.parametrize("dtype", FLOAT_DTYPES)
79+
def test_accuracy_your_operator(shape, dtype):
80+
# Test implementation
81+
inp = torch.randn(shape, dtype=dtype, device=flag_gems.device)
82+
ref_inp = to_reference(inp, True)
83+
84+
# Reference implementation
85+
ref_out = torch.your_operator(ref_inp, ...)
86+
87+
# FlagGems implementation
88+
with flag_gems.use_gems():
89+
res_out = flag_gems.experimental_ops.your_operator(inp, ...)
90+
91+
gems_assert_close(res_out, ref_out, dtype)
92+
```
93+
94+
## Performance Tests
95+
Add performance test in `exp_tests/your_ops_test.py`:
96+
```
97+
import pytest
98+
import torch
99+
import time
100+
import flag_gems
101+
102+
class TestYourOperatorPerf:
103+
def setup_method(self):
104+
flag_gems.enable()
105+
106+
def teardown_method(self):
107+
flag_gems.disable()
108+
109+
@pytest.mark.your_operator
110+
@pytest.mark.parametrize("shape", [...])
111+
def test_perf_your_operator(self, shape):
112+
inp = torch.randn(shape, device=flag_gems.device)
113+
114+
# Warmup
115+
for _ in range(10):
116+
_ = flag_gems.experimental_ops.your_operator(inp)
117+
118+
torch.cuda.synchronize()
119+
120+
# Benchmark FlagGems
121+
start_time = time.time()
122+
for _ in range(100):
123+
out = flag_gems.experimental_ops.your_operator(inp)
124+
torch.cuda.synchronize()
125+
gems_time = (time.time() - start_time) / 100
126+
127+
# Benchmark PyTorch
128+
start_time = time.time()
129+
for _ in range(100):
130+
ref_out = torch.your_operator(inp)
131+
torch.cuda.synchronize()
132+
torch_time = (time.time() - start_time) / 100
133+
134+
speedup = torch_time / gems_time
135+
print(f"YourOperator {shape}: Speedup {speedup:.2f}x")
136+
137+
assert speedup > 1.0, "Should be faster than PyTorch"
138+
```
139+
140+
# CI Integration
141+
Add tests ad performace tests to the CI workflow `.github/workflows/gems-experimental-test.yaml` .

0 commit comments

Comments
 (0)