Skip to content

Commit 905fe7f

Browse files
committed
refactor(dpa4): triton rotation
1 parent 6359dd8 commit 905fe7f

19 files changed

Lines changed: 2915 additions & 5370 deletions

deepmd/pt/model/descriptor/sezm.py

Lines changed: 0 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,6 @@
3333
)
3434

3535
import math
36-
import os
3736
from contextlib import (
3837
contextmanager,
3938
)
@@ -555,12 +554,6 @@ def __init__(
555554
self.layer_scale = bool(layer_scale)
556555
self.use_amp = bool(use_amp) # and self.training
557556
self.trainable = bool(trainable)
558-
self.use_triton = os.environ.get("DP_TRITON", "0").lower() in (
559-
"1",
560-
"true",
561-
"yes",
562-
"on",
563-
)
564557
self.seed = seed
565558
self.random_gamma = bool(random_gamma)
566559
self.add_chg_spin_ebd = bool(add_chg_spin_ebd)
@@ -899,7 +892,6 @@ def __init__(
899892
ffn_activation_function=self.ffn_activation_function,
900893
ffn_glu_activation=self.ffn_glu_activation,
901894
mlp_bias=self.mlp_bias,
902-
use_triton=self.use_triton,
903895
eps=self.eps,
904896
dtype=self.dtype,
905897
seed=child_seed(seed_blocks, block_idx),
@@ -1128,7 +1120,6 @@ def forward(
11281120
# the model is roll-equivariant, so inference fixes gamma.
11291121
random_gamma=self.random_gamma and self.training,
11301122
wigner_calc=self.wigner_calc,
1131-
use_geometry_rbf_triton=(self.use_triton and not self.training),
11321123
)
11331124

11341125
ebed_dim_0 = self.node_ebed_dims[0] # (node_lmax+1)^2

deepmd/pt/model/descriptor/sezm_nn/block.py

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -214,9 +214,6 @@ class SeZMInteractionBlock(nn.Module):
214214
- SO3Linear: l=0 bias
215215
- SO2Linear: l=0 bias
216216
- GatedActivation: gate linear bias
217-
use_triton
218-
If True, opt into fused Triton SO(2) rotation kernels inside
219-
``SO2Convolution`` when the runtime supports them.
220217
eps
221218
Small epsilon for numerical stability.
222219
dtype
@@ -275,7 +272,6 @@ def __init__(
275272
ffn_activation_function: str,
276273
ffn_glu_activation: bool = True,
277274
mlp_bias: bool = False,
278-
use_triton: bool = False,
279275
eps: float = 1e-7,
280276
dtype: torch.dtype,
281277
seed: int | list[int] | None,
@@ -370,7 +366,6 @@ def __init__(
370366
self.ffn_activation_function = str(ffn_activation_function)
371367
self.ffn_glu_activation = bool(ffn_glu_activation)
372368
self.mlp_bias = bool(mlp_bias)
373-
self.use_triton = bool(use_triton)
374369
self.eps = float(eps)
375370
self.dtype = dtype
376371
self.device = env.DEVICE
@@ -436,7 +431,6 @@ def __init__(
436431
lebedev_quadrature=self.so2_lebedev_quadrature,
437432
activation_function=self.so2_activation_function,
438433
mlp_bias=self.mlp_bias,
439-
use_triton=self.use_triton,
440434
eps=self.eps,
441435
dtype=dtype,
442436
seed=seed_so2_conv,
Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
# SPDX-License-Identifier: LGPL-3.0-or-later
2+
"""
3+
CuTe-DSL accelerated SO(2) rotation operators for SeZM / DPA4.
4+
5+
This package provides a self-contained, ``torch.compile``-friendly implementation
6+
of the two fused gather + batched-GEMM operators used by the SeZM SO(2) edge
7+
convolution:
8+
9+
* ``rotate_to_local`` : ``out[e] = wigner[e][coeff_index] @ x[src[e]]``
10+
* ``rotate_back`` : ``out[e] = wigner[e][:, coeff_index] @ x_local[e]``
11+
12+
The kernels are written with the NVIDIA CuTe DSL (``cutlass.cute``) and fuse the
13+
Wigner-row/column gather and the source-node gather directly into the matmul, so
14+
the large ``D_to_m`` / ``x_src`` intermediates are never materialized. They are
15+
exposed through the modern ``torch.library.custom_op`` API (functional, with
16+
``register_fake`` + ``register_autograd``) so that they compose correctly with
17+
``torch.compile`` and autograd.
18+
19+
The top-level entry points are re-exported here for convenience.
20+
"""
21+
22+
from __future__ import (
23+
annotations,
24+
)
25+
26+
from .so2_rotation import (
27+
SEZM_CUTE_AVAILABLE,
28+
rotate_back_cute,
29+
rotate_to_local_cute,
30+
)
31+
32+
__all__ = [
33+
"SEZM_CUTE_AVAILABLE",
34+
"rotate_back_cute",
35+
"rotate_to_local_cute",
36+
]

0 commit comments

Comments
 (0)