已合并
feat(libdevice): support index_put op #838
candyhong创建于 2025年12月1日
feat(libdevice): support index_put op #838
已合并
Pull Request已成功合入, 合并人@ascend-robot
(感谢 candyhong 的贡献)ascend-robot
2025年12月1日 评论:
2025年12月1日 评论:
openLiBingCI
2025年12月1日 评论:
2025年12月1日 评论:
ascend-robot
2025年12月1日 评论:
2025年12月1日 评论:
以下是根据您提交的修改文件推荐的Reviewer和Committer序列,需各模块评审通过后方可合入
| Module List | Reviewers | Committers |
|---|---|---|
| repo-Ascend/triton-ascend | ccdedreams, cxtjjcz, wcleungaj, wangtao489, zhang-chunli01 | tarvlad, cxtjjcz, wcleungaj, shijingchang, weizhan4 |


2025年12月1日 添加了label:ascend-cla/yes
ascend-robot
2025年12月1日 评论:
2025年12月1日 评论:
CLA Signature Pass
@candyhong, thanks for your pull request. All authors of the commits have signed the CLA. 👍


2025年12月1日 添加了label:ci-pipeline-running
CANN-robot
2025年12月1日 评论:
2025年12月1日 评论:
compile#openlibing


openLiBingCI
2025年12月1日 评论:
2025年12月1日 评论:
libing-pipeline创建中,请等待……


2025年12月1日 添加了label:SC-START
2025年12月1日 删除了label:SC-START
2025年12月1日 添加了label:SC-RUNNING
2025年12月1日 删除了label:SC-RUNNING
2025年12月1日 添加了label:SC-RUNNING
openLiBingCI
2025年12月1日 评论:
2025年12月1日 评论:
CANN-robot
2025年12月1日 评论:
2025年12月1日 评论:
流水线 PR-pipeline_Triton-Ascend#1551运行失败
| 阶段 | 任务名 | 状态 | 详情 |
|---|---|---|---|
| 编译构建 | 构建triton-ascend_py39_arm | ✅ | >>> |
| 构建triton-ascend_py310_arm | ✅ | >>> | |
| 构建triton-ascend_py311_arm | ✅ | >>> | |
| 构建triton-ascend_py39_x86 | ✅ | >>> | |
| 构建triton-ascend_py310_x86 | ✅ | >>> | |
| 构建triton-ascend_py311_x86 | ✅ | >>> | |
| 恶意代码检查 | codecheck_anti | ✅ | >>> |
| 编码安全与规范检查 | code_check | ✅ | >>> |
| 开源片段检查 | codecheck_sca | ✅ | >>> |
| 开发者测试 | 验证triton-ascend_py39_ut | ✅ | >>> |
| 验证triton-ascend_py310_ut | ❌ | >>> | |
| 验证triton-ascend_py311_ut | ❌ | >>> | |
| 流水线 | PR-pipeline_Triton-Ascend | ❌ | >>> |


2025年12月1日 删除了label:ci-pipeline-running
2025年12月1日 添加了label:ci-pipeline-passed
2025年12月2日 删除了label:SC-RUNNING
2025年12月2日 添加了label:SC-SUCC
triton_patch/python/triton_patch/language/core.py
@@ -21,7 +21,7 @@
2121# THE SOFTWARE.
2222
2323import os
24-from typing import List, Sequence, Optional, Union
24+from typing import List, Sequence, Optional, Union, Tuple
2525
2626from triton._C.libtriton import ir
2727from triton.language import semantic as real_semantic
@@ -645,6 +645,81 @@ def index_select(src: tensor, idx: tensor, bound, lstdim_blksiz, offsets, numels
645645 return semantic.embedding_gather(src, idx, bound, lstdim_blksiz, offsets, numels, _builder)
646646
647647
648+
之前不是有结论说,非标准的OP,放在libdevice里,也就是 libdevice.index_put 调用。你写在core里,不是变成 tl.index_put 了。


candyhong
2025年12月2日 评论:
2025年12月2日 评论:
因为是一整套方言,现在的方案是沿用core,semantic 进行创建,外层接口通过libdevice引用:
language.extra.ascend.libdevice.index_put = index_put
HaiLijuan
2025年12月2日 评论:
2025年12月2日 评论:
黄线方案不是不改triton代码,由TA的pass分析pattern match到扩展op吗?现在这样加编程接口,怎么推广使用?怎么兼容A5 SIMT?
triton_patch/python/triton_patch/testing.py
@@ -651,6 +651,8 @@ from .triton_patch.language.core import (
651651 insert_slice,
652652 index_select_simd,
653653 index_select,
654+ index_put,
655+ gather_out_to_ub,
654656 gather_out_to_ub,
655657 extract_slice,
656658 trans,
@@ -854,4 +856,5 @@ language.extra.ascend.libdevice.fdiv = language.math.fdiv
854856language.extra.ascend.libdevice.fma = language.math.fma
855857language.extra.ascend.libdevice.abs = language.math.abs
856858language.extra.ascend.libdevice.index_select_simd = index_select_simd
859+language.extra.ascend.libdevice.index_put = index_put
没有加ut,能否增加ut测试


2025年12月2日 解决了最后一个问题
wangzhanpeng5
2025年12月2日 评论:
2025年12月2日 评论:
/lgtm


KanuaK
2025年12月2日 评论:
2025年12月2日 评论:
/lgtm


2025年12月2日 添加了label:lgtm
zhang-chunli01
2025年12月2日 评论:
2025年12月2日 评论:
/approve


2025年12月2日 添加了label:approved
ascend-robot
2025年12月2日 评论:
2025年12月2日 评论:
Review Guide
This Pull-Request Passes Review.
Committers who writed a comment of /approve are: zhang-chunli01.
Reviewers who writed a comment of /lgtm are: zhang-chunli01, wangzhanpeng5, KanuaK.


2025年12月2日 合入了pull request
描述
为支持从统一缓冲区(Unified Buffer, UB)到全局内存(Global Memory, GM)的高效数据散射(
Index Put)操作,实现按指定索引将数据精准写入目标张量,新增libdevice.index_put接口。TA 层在代码生成阶段会构造func.call "triton_index_put"调用,底层对接 SIMT 架构的模板实现,适配硬件级别的高效数据搬移需求,满足 2~5 维张量的索引写入场景。接口说明
index_put接口用于将 UB 中的value张量,按照 UB 中的index索引,沿目标张量的指定维度dim,写入到 GM 中的ptr目标张量。参数说明
ptrtl.tensor(指针类型)indextl.tensorvaluetl.tensordimint0 ≤ dim < value.rank-1dst_shapeTuple[int]dst_offsetTuple[int]约束与限制
ptr(目标张量)与value(待写入张量)的秩(rank)必须相同。ptr和value的数据类型仅支持 float16、bfloat16、float32 三种浮点类型。index必须为整数类型张量,且index固定为 1D 维度。value仅支持 2~5 维张量;dim需为有效维度(满足0 ≤ dim < rank(value)-1),确保沿合法维度进行散射。checklist