flagos-ai / flagos-ai/FlagTree

[DOC][TLE.DSA][ascend] tle.dsa.local_ptr is documented in tle.md but absent from every release

Open
#1,111 1 comment 0 reactions 0 assignees View on GitHub
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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.