|
1 | | -"""Mthreads backend spec module. |
2 | | -
|
3 | | -This module is loaded by flagtree_spec.py via |
4 | | -importlib.import_module("triton.backends.mthreads.spec") when |
5 | | -FLAGTREE_BACKEND=mthreads. Functions defined here are callable |
6 | | -via flagtree_spec.spec("function_name", ...). |
7 | | -""" |
8 | | - |
9 | | - |
10 | | -def init_language(): |
11 | | - """Add mthreads-specific symbols to triton.language. |
12 | | -
|
13 | | - Called explicitly from the main triton/__init__.py after |
14 | | - triton.language is fully initialized, via: |
15 | | - spec("init_language") |
16 | | - """ |
17 | | - # Delay importing language.core until Triton's language module is ready. |
18 | | - # Keep constexpr global so the JIT can resolve the annotation by name. |
19 | | - global constexpr |
20 | | - from triton.flagtree_spec import bind_language_extension_symbols_to_tl |
21 | | - from triton.runtime.jit import jit as _jit |
22 | | - from triton.language.core import ( |
23 | | - constexpr, |
24 | | - builtin as _builtin, |
25 | | - static_assert as _static_assert, |
26 | | - _unwrap_if_constexpr, |
27 | | - ) |
28 | | - |
29 | | - class _Ext: |
30 | | - __all__ = [ |
31 | | - "squeeze", |
32 | | - "unsqueeze", |
33 | | - "_experimental_descriptor_load", |
34 | | - "_experimental_descriptor_store", |
35 | | - ] |
36 | | - |
37 | | - _ext = _Ext() |
38 | | - |
39 | | - @_jit |
40 | | - def squeeze(x, dim: constexpr): |
41 | | - _static_assert(x.shape[dim] == 1) |
42 | | - return x.reshape(x.shape[:dim] + x.shape[dim + 1:]) |
43 | | - |
44 | | - @_jit |
45 | | - def unsqueeze(x, dim: constexpr): |
46 | | - return x.reshape(x.shape[:dim] + (1, ) + x.shape[dim:]) |
47 | | - |
48 | | - @_builtin |
49 | | - def _experimental_descriptor_load(desc_pointer, offsets, shape, dtype, _semantic=None): |
50 | | - """Legacy compatibility API for descriptor load. |
51 | | -
|
52 | | - New code should prefer ``load_tensor_descriptor``. We keep this symbol |
53 | | - so migrated tests can exercise the same backend descriptor path without |
54 | | - monkeypatching triton.language internals in conftest. |
55 | | - """ |
56 | | - _ = shape |
57 | | - dtype = _unwrap_if_constexpr(dtype) |
58 | | - value = desc_pointer.load(offsets, _semantic=_semantic) |
59 | | - if value.dtype == dtype: |
60 | | - return value |
61 | | - if value.dtype.primitive_bitwidth == dtype.primitive_bitwidth: |
62 | | - return value.to(dtype, bitcast=True, _semantic=_semantic) |
63 | | - return value.to(dtype, _semantic=_semantic) |
64 | | - |
65 | | - @_builtin |
66 | | - def _experimental_descriptor_store(desc_pointer, value, offsets, _semantic=None): |
67 | | - """Legacy compatibility API for descriptor store. |
68 | | -
|
69 | | - New code should prefer ``store_tensor_descriptor``. |
70 | | - """ |
71 | | - value = _semantic.to_tensor(value) |
72 | | - desc_dtype = desc_pointer.dtype |
73 | | - if value.dtype != desc_dtype and value.dtype.primitive_bitwidth == desc_dtype.primitive_bitwidth: |
74 | | - value = value.to(desc_dtype, bitcast=True, _semantic=_semantic) |
75 | | - return desc_pointer.store(offsets, value, _semantic=_semantic) |
76 | | - |
77 | | - _ext.squeeze = squeeze |
78 | | - _ext.unsqueeze = unsqueeze |
79 | | - _ext._experimental_descriptor_load = _experimental_descriptor_load |
80 | | - _ext._experimental_descriptor_store = _experimental_descriptor_store |
81 | | - |
82 | | - bind_language_extension_symbols_to_tl(_ext) |
83 | | - |
84 | | - |
85 | 1 | from ._filecheck import spec_get_stub_target |
86 | 2 | from ._utils import apply_with_path, _tuple_create |
| 3 | +from .language import language_extend_globals |
0 commit comments