flagos-ai / flagos-ai/FlagTree
[DOC][TLE.DSA][ascend] tle.dsa.local_ptr is documented in tle.md but absent from every release
- Dominant language
- Python
- Stars
- 350
- Forks
- 149
- Avg merge
- 2d 4h
- Merged PRs (30d)
- 87
Description
`tle.md` documents `tle.dsa.local_ptr` as a working primitive, with a full
signature, semantics and two runnable examples. It does not exist in any
released version.
## What the documentation says
`tle.md` §3.3.2.1.3:
> - Signature: `tle.dsa.local_ptr(buffer, indices=None) -> tl.tensor | tl.ptr`
> - Purpose: build pointer views over DSA local buffers (for example UB/L1) for
> explicit local-memory access patterns.
> - Semantics: Shape and indexing behavior follow `tle.gpu.local_ptr` (same
> pointer-view model).
>
> ```python
> a_ub = tle.dsa.alloc([BM, BK], dtype=tl.float16, mem_addr_space=tle.dsa.ascend.UB)
> rows = tl.broadcast_to(tl.arange(0, BM)[:, None], (BM, BK))
> cols = tl.broadcast_to(tl.arange(0, BK)[None, :], (BM, BK))
> a_ptr = tle.dsa.local_ptr(a_ub, (rows, cols))
> a_val = tl.load(a_ptr)
> ```
§3.3.2.1.4 documents a second form, `tle.dsa.local_ptr(remote_buffer, indices)`,
for buffers obtained from `tle.remote(...)`. The same text appears in
`tle_cn.md`, in `flagos-ai/docs` (`user_guide/use-tle-struct.md`, both
languages), and in the v0.5.0 release notes.
## What is actually shipped
```python
>>> import triton.experimental.tle.language.dsa as dsa
>>> dsa.local_ptr
AttributeError: module 'triton.experimental.tle.language.dsa' has no attribute 'local_ptr'
```
`dsa` exports 33 public names on `0.6.1+ascend3.5`; the only one containing
"pointer" is `from_buffer_to_tensor_pointer`, which returns a value tensor, not
a pointer — using it with `tl.store` gives
`ValueError: Unsupported ptr type <[128], int32>`. Its 18 call sites are all
internal, converting buffers into operands for `dsa.add`/`sub`/`mul`/`div`/
`max`/`min`.
Checked, all with the same result:
| version | how obtained |
|---|---|
| `0.6.0+ascend3.5` | vendor-built, shipped in a CANN 9.0.0 container |
| `0.6.1+ascend3.5` | built from source |
| `0.6.2a1+ascend3.5` | built from source |
Wider checks:
- `git grep local_ptr -- 'python/triton/experimental/tle/language/dsa/*'`
returns nothing on **any** ascend branch or tag, and on `main`.
- `main` has no `language/dsa/` directory at all — only `gpu/` and `raw/`.
- `main`'s primitive registry, `language/primitives.py`, lists **no `dsa.*`
entry**. `local_ptr` appears there only as `"gpu.local_ptr"`.
- On the C++ side the builder exposes twelve `create_dsa_*` methods
(`alloc`, `copy`, `add`, `sub`, `mul`, `div`, `max`, `min`, `extract_scalar`,
`extract_slice`, `insert_slice`, `subview`) — **none of them returns a
pointer**.
- A GitHub-wide code search for `dsa.local_ptr` returns 7 hits, all of them
documentation (this repo's `tle.md`/`tle_cn.md`, `flagos-ai/docs`,
`flagos-ai/FlagFFT`) and **no code**. The control search, `gpu.local_ptr`,
returns 143 including real operators.
## Why it matters
The section reads exactly like a shipped feature, so it is easy to plan around
it before checking the sources. Without a pointer into UB, `tle.dsa` is
value-semantic only (`alloc` / `copy` / `to_tensor` / `subview` /
`extract_slice` / `insert_slice`), so an operator that needs to scatter into
on-chip memory — a radix histogram, for instance — cannot be expressed through
DSA at all. Knowing that from the documentation would have saved a
considerable amount of work.
## Request
Either implement `tle.dsa.local_ptr` as documented, or mark §3.3.2.1.3 and
§3.3.2.1.4 as planned / not yet implemented, and remove the mention from the
v0.5.0 release notes. The same applies to the copies in `flagos-ai/docs` and
`flagos-ai/FlagFFT`.
## Note on the `gpu` form
`tle.gpu.local_ptr` does exist, but it is not usable on Ascend either: building
its op aborts the process with
```
LLVM ERROR: Building op `tle.local_pointers` but it isn't known in this
MLIRContext: the dialect may not be loaded or this operation hasn't been added
by the dialect.
```
`libtriton.tle.passes` carries five TLE passes while `libtriton.tle.ir` is
empty, on the vendor build and on two of our own — the dialect's IR bindings
appear not to be built for the ascend target. Happy to open that separately if
it is a distinct problem.
Contributor guide
Research direction
Start with tle.md §3.3.2.1.3-§3.3.2.1.4 and compare them with the available DSA sources, language/primitives.py, and the C++ DSA builder methods. Check the listed release and documentation copies; done means either implementing the documented primitive with coverage or marking both sections as unavailable and removing the stale release-note and external references.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- python
- Domain
- compilers, documentation
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Active
- Clarity
- Mostly clear
- Newbie friendliness
- 38/100