已合并
feat(libdevice): support index_put op #838
candyhong创建于 2025年12月1日
feat(libdevice): support index_put op #838
已合并
candyhong创建于 2025年12月1日
candyhong成员
2025年12月1日

描述

为支持从统一缓冲区(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 目标张量。

参数说明

参数名 类型 是否必填 说明
ptr tl.tensor(指针类型) 全局内存(GM)中的目标张量指针,数据将写入此张量
index tl.tensor 统一缓冲区(UB)中的索引张量,必须为 1D 整数张量
value tl.tensor 统一缓冲区(UB)中的待写入数据张量,支持 2~5 维
dim int 散射操作沿目标张量的维度,需满足0 ≤ dim < value.rank-1
dst_shape Tuple[int] 目标张量的完整形状,用于计算写入位置的全局偏移
dst_offset Tuple[int] 目标张量各维度的偏移量,用于定位写入区域在目标张量中的起始位置

约束与限制

  • 维度一致性:ptr(目标张量)与 value(待写入张量)的秩(rank)必须相同。
  • 数据类型限制:ptrvalue 的数据类型仅支持 float16、bfloat16、float32 三种浮点类型。
  • 索引张量约束:index 必须为整数类型张量,且index 固定为 1D 维度。
  • 维度范围约束:value 仅支持 2~5 维张量;dim 需为有效维度(满足 0 ≤ dim < rank(value)-1),确保沿合法维度进行散射。

checklist

likedislike
Pull Request已成功合入, 合并人@ascend-robot
(感谢 candyhong 的贡献)
Ccandyhong成员
2025年12月1日 创建了 pull request,commit d83b0c8c
ascend-robot
ascend-robot成员
2025年12月1日 评论:

Thank your for your pull-request.

The full list of commands accepted by me can be found at here.

You can get sig-info at here

likedislike
openLiBingCI成员
2025年12月1日 评论:
ascend-robot
ascend-robot成员
2025年12月1日 评论:

以下是根据您提交的修改文件推荐的Reviewer和Committer序列,需各模块评审通过后方可合入

Module List Reviewers Committers
repo-Ascend/triton-ascend ccdedreams, cxtjjcz, wcleungaj, wangtao489, zhang-chunli01 tarvlad, cxtjjcz, wcleungaj, shijingchang, weizhan4
likedislike
ascend-robotascend-robot成员
2025年12月1日 添加了label:ascend-cla/yes
ascend-robot
ascend-robot成员
2025年12月1日 评论:

CLA Signature Pass

@candyhong, thanks for your pull request. All authors of the commits have signed the CLA. 👍

likedislike
candyhong成员
2025年12月1日 评论:

/compile

likedislike
CANN-robotCANN-robot成员
2025年12月1日 添加了label:ci-pipeline-running
CANN-robot
CANN-robot成员
2025年12月1日 评论:

compile#openlibing

likedislike
openLiBingCI成员
2025年12月1日 评论:

libing-pipeline创建中,请等待……

likedislike
OopenLiBingCI成员
2025年12月1日 添加了label:SC-START
OopenLiBingCI成员
2025年12月1日 删除了label:SC-START
OopenLiBingCI成员
2025年12月1日 添加了label:SC-RUNNING
OopenLiBingCI成员
2025年12月1日 删除了label:SC-RUNNING
OopenLiBingCI成员
2025年12月1日 添加了label:SC-RUNNING
openLiBingCI成员
2025年12月1日 评论:
CANN-robot
CANN-robot成员
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 >>>
likedislike
Ccandyhong成员
2025年12月1日 修改了pull request 的描述
CANN-robotCANN-robot成员
2025年12月1日 删除了label:ci-pipeline-running
CANN-robotCANN-robot成员
2025年12月1日 添加了label:ci-pipeline-passed
OopenLiBingCI成员
2025年12月2日 删除了label:SC-RUNNING
OopenLiBingCI成员
2025年12月2日 添加了label:SC-SUCC
wangzhanpeng52025年12月2日进行代码检视5
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+@builtin
wangzhanpeng52025年12月2日评论:

之前不是有结论说,非标准的OP,放在libdevice里,也就是 libdevice.index_put 调用。你写在core里,不是变成 tl.index_put 了。

likedislike
candyhong成员
2025年12月2日 评论:

因为是一整套方言,现在的方案是沿用core,semantic 进行创建,外层接口通过libdevice引用:
language.extra.ascend.libdevice.index_put = index_put

HaiLijuan成员
2025年12月2日 评论:

黄线方案不是不改triton代码,由TA的pass分析pattern match到扩展op吗?现在这样加编程接口,怎么推广使用?怎么兼容A5 SIMT?

candyhong成员
2025年12月2日 评论:

这个不是黄线方案,可以看最新的架构图,现在这个libdevice 接口对接inductor,inductor 会自动生成对应的相关代码。用户手写的话可以参考一下接口描述中的示例案例,后续补op文档进行推广

KanuaK
KanuaK成员
2025年12月2日 评论:

黄线方案在PR806,不是同一个方案

wangzhanpeng52025年12月2日进行代码检视2
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
wangzhanpeng52025年12月2日评论:

没有加ut,能否增加ut测试

likedislike
candyhong成员
2025年12月2日 评论:

只在A5 能跑,且当前NPUIR没上库,暂时不上ut测试,后续流水支持再补

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

/lgtm

likedislike
KanuaK
KanuaK成员2025年12月2日进行代码检视2
triton_patch/python/triton_patch/testing.py
@@ -652,2 +652,4 @@
652652 index_select_simd,
653653 index_select,
654+ index_put,
655+ gather_out_to_ub,
KanuaK
KanuaK2025年12月2日评论:

gather_out_to_ub与L656有重复

likedislike
candyhong成员
2025年12月2日 评论:

我另外提pr修复吧。这个PR今天得合入,下午出主线包转测

KanuaK
KanuaK成员
2025年12月2日 评论:

/lgtm

likedislike
ascend-robotascend-robot成员
2025年12月2日 添加了label:lgtm
zhang-chunli01成员
2025年12月2日 评论:

/approve

likedislike
ascend-robotascend-robot成员
2025年12月2日 添加了label:approved
ascend-robot
ascend-robot成员
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.

likedislike
ascend-robotascend-robot成员
2025年12月2日 合入了pull request