[libdevice](fix) Remove duplicated deprecated function

third_party/ascend/language/cann/extension/_utils.py

@@ -1,35 +1,6 @@
-import inspect
-from functools import wraps
-from warnings import warn
-
 import triton.language.core as tl
 from triton._C.libtriton import ir
 
-_DEPRECATED_MESSAGE_ATTR = "_deprecated_message"
-
-
-def _deprecated(fn_name=None, replacement=None):
-
-    def decorator(fn):
-        name = fn_name or f"triton.language.extra.cann.extension.{fn.__name__}"
-        message = (f"{name} is deprecated and will be removed in the next release"
-                   f"{f'; use {replacement} instead.' if replacement else '.'}")
-
-        if inspect.isclass(fn):
-            setattr(fn, _DEPRECATED_MESSAGE_ATTR, message)
-            fn.__doc__ = f"{fn.__doc__ or ''}\n\n.. warning::\n   {message}"
-            return fn
-
-        @wraps(fn)
-        def wrapper(*args, **kwargs):
-            warn(message, FutureWarning, stacklevel=2)
-            return fn(*args, **kwargs)
-
-        wrapper.__doc__ = f"{fn.__doc__ or ''}\n\n.. warning::\n   {message}"
-        return wrapper
-
-    return decorator
-
 
 def custom_op(builder: ir.builder, op_name: str, **kwargs):
     if op_name == "sync_block_all":

third_party/ascend/language/cann/extension/builtin_custom_ops.py

@@ -23,7 +23,8 @@
 import triton.language.core as tl
 from .custom_op import register_custom_op
 from .core import CORE, PIPE, MODE
-from ._utils import _is_int_like_elem, _assert_int_like_tuple, _deprecated
+from ..utils import _deprecated
+from ._utils import _is_int_like_elem, _assert_int_like_tuple
 
 
 @_deprecated(fn_name="triton.language.extra.cann.extension.custom('__builtin_index_select', ...)")

third_party/ascend/language/cann/extension/mem_ops.py

@@ -15,7 +15,8 @@
 from typing import Optional, Tuple, List, overload, Union
 from triton._C.libtriton import ir
 
-from ._utils import _convert_elem_to_ir_value, _deprecated
+from ..utils import _deprecated
+from ._utils import _convert_elem_to_ir_value
 
 
 @_deprecated()

third_party/ascend/language/cann/libdevice.py

@@ -18,14 +18,14 @@
 # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
 # THE SOFTWARE.
 
-from functools import wraps
 from math import pi as math_pi
-from warnings import warn
 from triton.language import core, math, semantic, standard
 from triton._C.libtriton import ir
 from triton.runtime.jit import jit
 from triton.backends.ascend.utils import is_compile_on_910_95, triton_enable_libdevice_simt
 
+from .utils import _deprecated
+
 
 def _is_libdevice_simt_enabled(_semantic) -> bool:
     return triton_enable_libdevice_simt(_semantic.builder.options.arch)
@@ -35,23 +35,6 @@ def _is_a5_target(_semantic) -> bool:
     return is_compile_on_910_95(_semantic.builder.options.arch)
 
 
-def _deprecated(replacement):
-
-    def decorator(fn):
-        message = (f"cann.libdevice.{fn.__name__} is deprecated and will be removed in the next release; "
-                   f"use cann.libdevice.{replacement} instead.")
-
-        @wraps(fn)
-        def wrapper(*args, **kwargs):
-            warn(message, FutureWarning, stacklevel=2)
-            return fn(*args, **kwargs)
-
-        wrapper.__doc__ = f"{fn.__doc__ or ''}\n\n.. warning::\n   {message}"
-        return wrapper
-
-    return decorator
-
-
 class _FlipStaticRange:
 
     def __init__(self, arg1, arg2=None, step=None):
@@ -571,7 +554,7 @@ def hadd(arg0, arg1, _semantic=None):
     return core.extern_elementwise(
         "", "", [arg0, arg1], {
             (core.dtype("int32"), core.dtype("int32")): ("__hmf_hadd_i32", core.dtype("int32")),
-            (core.dtype("uint32"), core.dtype("uint32")): ("__hmf_uhadd_u32_u32", core.dtype("uint32")),
+            (core.dtype("uint32"), core.dtype("uint32")): ("__hmf_uhadd_u32", core.dtype("uint32")),
         }, is_pure=True, _semantic=_semantic)
 
 
@@ -583,7 +566,7 @@ def rhadd(arg0, arg1, _semantic=None):
     return core.extern_elementwise(
         "", "", [arg0, arg1], {
             (core.dtype("int32"), core.dtype("int32")): ("__hmf_rhadd_i32", core.dtype("int32")),
-            (core.dtype("uint32"), core.dtype("uint32")): ("__hmf_urhadd_u32_u32", core.dtype("uint32")),
+            (core.dtype("uint32"), core.dtype("uint32")): ("__hmf_urhadd_u32", core.dtype("uint32")),
         }, is_pure=True, _semantic=_semantic)
 
 
@@ -2913,30 +2896,35 @@ def sqrt(arg0, _semantic=None):
 
 
 @core.extern
-@_deprecated("hadd")
+@_deprecated(fn_name="triton.language.extra.cann.libdevice.uhadd",
+             replacement="triton.language.extra.cann.libdevice.hadd")
 def uhadd(arg0, arg1, _semantic=None):
     return hadd(arg0, arg1, _semantic=_semantic)
 
 
 @core.extern
-@_deprecated("mul24")
+@_deprecated(fn_name="triton.language.extra.cann.libdevice.umul24",
+             replacement="triton.language.extra.cann.libdevice.mul24")
 def umul24(arg0, arg1, _semantic=None):
     return mul24(arg0, arg1, _semantic=_semantic)
 
 
 @core.extern
-@_deprecated("mulhi")
+@_deprecated(fn_name="triton.language.extra.cann.libdevice.umulhi",
+             replacement="triton.language.extra.cann.libdevice.mulhi")
 def umulhi(arg0, arg1, _semantic=None):
     return mulhi(arg0, arg1, _semantic=_semantic)
 
 
 @core.extern
-@_deprecated("rhadd")
+@_deprecated(fn_name="triton.language.extra.cann.libdevice.urhadd",
+             replacement="triton.language.extra.cann.libdevice.rhadd")
 def urhadd(arg0, arg1, _semantic=None):
     return rhadd(arg0, arg1, _semantic=_semantic)
 
 
 @core.extern
-@_deprecated("sad")
+@_deprecated(fn_name="triton.language.extra.cann.libdevice.usad",
+             replacement="triton.language.extra.cann.libdevice.sad")
 def usad(arg0, arg1, arg2, _semantic=None):
     return sad(arg0, arg1, arg2, _semantic=_semantic)

third_party/ascend/language/cann/utils.py

@@ -0,0 +1,28 @@
+import inspect
+from functools import wraps
+from warnings import warn
+
+_DEPRECATED_MESSAGE_ATTR = "_deprecated_message"
+
+
+def _deprecated(fn_name=None, replacement=None):
+
+    def decorator(fn):
+        name = fn_name or f"triton.language.extra.cann.extension.{fn.__name__}"
+        message = (f"{name} is deprecated and will be removed in the next release"
+                   f"{f'; use {replacement} instead.' if replacement else '.'}")
+
+        if inspect.isclass(fn):
+            setattr(fn, _DEPRECATED_MESSAGE_ATTR, message)
+            fn.__doc__ = f"{fn.__doc__ or ''}\n\n.. warning::\n   {message}"
+            return fn
+
+        @wraps(fn)
+        def wrapper(*args, **kwargs):
+            warn(message, FutureWarning, stacklevel=2)
+            return fn(*args, **kwargs)
+
+        wrapper.__doc__ = f"{fn.__doc__ or ''}\n\n.. warning::\n   {message}"
+        return wrapper
+
+    return decorator