Star1736
1647
Code
Code
Issues475
Pull Requests549
Actions
Discussion
Wiki
Analysis
Star1736
1647
Merged
docs: 修复注释与文档不一致（concat、diag_v2、dynamic_stitch 等 文件注释） #6598
TangPCAdd on 18 days ago
docs: 修复注释与文档不一致（concat、diag_v2、dynamic_stitch 等 文件注释） #6598
Merged
TangPCAdd on 18 days ago
Merged frompingchuantang/ops-math-tpcdocs/comment-cleanup-20261022-03intocann/ops-mathmaster
Total: 5 files changed+29-34
Mconversion/concat/tests/ut/op_host/test_concat_infershape.cpp+1-1
@@ -8,7 +8,7 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10/*!10/*!
11- * \file test_concat_infershpe.cpp11+ * \file test_concat_infershape.cpp
12 * \brief12 * \brief
13 */13 */
14 14 
Mconversion/diag_v2/op_kernel/diag_v2_b32.h+6-7
@@ -21,8 +21,7 @@ namespace DiagV2 {
21using namespace AscendC;21using namespace AscendC;
22 22 
23template <typename T>23template <typename T>
24-class DiagV2B32 : public DiagV2Base<T>24+class DiagV2B32 : public DiagV2Base<T> {
25-{
26public:25public:
27 __aicore__ inline DiagV2B32(){};26 __aicore__ inline DiagV2B32(){};
28 __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, const DiagV2TilingData* tilingData);27 __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, const DiagV2TilingData* tilingData);
@@ -102,8 +101,8 @@ __aicore__ inline void DiagV2B32<T>::Process()
102}101}
103 102 
104template <typename T>103template <typename T>
105-__aicore__ inline void DiagV2B32<T>::ProcessPerBlock(104+__aicore__ inline void DiagV2B32<T>::ProcessPerBlock(const LocalTensor<int16_t>& ubAssist, int64_t loopNum,
106- const LocalTensor<int16_t>& ubAssist, int64_t loopNum, int64_t loopTail)105+ int64_t loopTail)
107{106{
108 for (int64_t idx = 0; idx < loopNum; idx++) {107 for (int64_t idx = 0; idx < loopNum; idx++) {
109 gmOutOffset = blockIdx * m_tilingData.numPerCore + idx * m_tilingData.matrixRowLength;108 gmOutOffset = blockIdx * m_tilingData.numPerCore + idx * m_tilingData.matrixRowLength;
@@ -120,8 +119,8 @@ __aicore__ inline void DiagV2B32<T>::ProcessPerBlock(
120}119}
121 120 
122template <typename T>121template <typename T>
123-__aicore__ inline void DiagV2B32<T>::ProcessLastBlock(122+__aicore__ inline void DiagV2B32<T>::ProcessLastBlock(const LocalTensor<int16_t>& ubAssist, int64_t loopNum,
124- const LocalTensor<int16_t>& ubAssist, int64_t loopNum, int64_t loopTail)123+ int64_t loopTail)
125{124{
126 for (int64_t idx = 0; idx < loopNum; idx++) {125 for (int64_t idx = 0; idx < loopNum; idx++) {
127 gmOutOffset = blockIdx * m_tilingData.numPerCore + idx * m_tilingData.matrixRowLength;126 gmOutOffset = blockIdx * m_tilingData.numPerCore + idx * m_tilingData.matrixRowLength;
@@ -212,4 +211,4 @@ __aicore__ inline void DiagV2B32<T>::CopyInB32WithPad(int64_t index, int64_t dat
212}211}
213 212 
214} // namespace DiagV2213} // namespace DiagV2
215-#endif // DIAG_V2_B64_H214+#endif // DIAG_V2_B32_H
Mconversion/dynamic_stitch/op_kernel/arch35/dynamic_stitch_scatter_simt.h+12-15
@@ -15,7 +15,7 @@
15 */15 */
16 16 
17/*!17/*!
18- * \file dynamic_stitch_simt.h18+ * \file dynamic_stitch_scatter_simt.h
19 * \brief19 * \brief
20 */20 */
21#ifndef DYNAMIC_STITCH_SIMT_H21#ifndef DYNAMIC_STITCH_SIMT_H
@@ -29,8 +29,7 @@ namespace DynamicStitch {
29constexpr uint32_t FIRST_DIM_THREAD_NUM = 2;29constexpr uint32_t FIRST_DIM_THREAD_NUM = 2;
30 30 
31template <typename T>31template <typename T>
32-class DynamicStitchScatterSimt32+class DynamicStitchScatterSimt {
33-{
34public:33public:
35 __aicore__ inline DynamicStitchScatterSimt(TPipe* pipe, const DynamicStitchTilingData* tiling)34 __aicore__ inline DynamicStitchScatterSimt(TPipe* pipe, const DynamicStitchTilingData* tiling)
36 : pipe_(pipe), tilingData_(tiling){};35 : pipe_(pipe), tilingData_(tiling){};
@@ -39,7 +38,7 @@ public:
39 38 
40private:39private:
41 __aicore__ inline void ParseTilingData();40 __aicore__ inline void ParseTilingData();
42- 41+ 
43 TPipe* pipe_;42 TPipe* pipe_;
44 const DynamicStitchTilingData* tilingData_;43 const DynamicStitchTilingData* tilingData_;
45 TBuf<TPosition::VECCALC> assistBuffer_;44 TBuf<TPosition::VECCALC> assistBuffer_;
@@ -64,8 +63,8 @@ __simt_vf__ __aicore__ LAUNCH_BOUND(THREAD_NUM) inline void SimtDataCopy(
64 GM_ADDR inTensorsPtr, int64_t startTensorIndex, int64_t endTensorIndex, int64_t startOffset, int64_t endOffset,63 GM_ADDR inTensorsPtr, int64_t startTensorIndex, int64_t endTensorIndex, int64_t startOffset, int64_t endOffset,
65 __ubuf__ int64_t* tensorCumsum, __gm__ int32_t* workspaceGmAddr, __gm__ T* yGmAddr, const int64_t sliceSize)64 __ubuf__ int64_t* tensorCumsum, __gm__ int32_t* workspaceGmAddr, __gm__ T* yGmAddr, const int64_t sliceSize)
66{65{
67- for (int32_t tensorIndex = startTensorIndex + static_cast<int32_t>(threadIdx.x);66+ for (int32_t tensorIndex = startTensorIndex + static_cast<int32_t>(threadIdx.x); tensorIndex <= endTensorIndex;
68- tensorIndex <= endTensorIndex; tensorIndex += static_cast<int32_t>(blockDim.x)) {67+ tensorIndex += static_cast<int32_t>(blockDim.x)) {
69 __gm__ T* inputTensor = GetTensorSimtAddr<T>(tensorIndex, inTensorsPtr);68 __gm__ T* inputTensor = GetTensorSimtAddr<T>(tensorIndex, inTensorsPtr);
70 int64_t curElemCount = tensorCumsum[tensorIndex + 1] - tensorCumsum[tensorIndex]; // 获取当前tensor的个数69 int64_t curElemCount = tensorCumsum[tensorIndex + 1] - tensorCumsum[tensorIndex]; // 获取当前tensor的个数
71 int64_t curStartIndex = 0;70 int64_t curStartIndex = 0;
@@ -105,8 +104,7 @@ __aicore__ inline void DynamicStitchScatterSimt<T>::Init(GM_ADDR x, GM_ADDR y, G
105 xGmAddr_ = x;104 xGmAddr_ = x;
106 yGm_.SetGlobalBuffer((__gm__ T*)y);105 yGm_.SetGlobalBuffer((__gm__ T*)y);
107 workspaceGm_.SetGlobalBuffer((__gm__ int32_t*)workspace);106 workspaceGm_.SetGlobalBuffer((__gm__ int32_t*)workspace);
108- pipe_->InitBuffer(107+ pipe_->InitBuffer(assistBuffer_, AlignUp((MAX_LIST_TENSOR_CNT + 1) * sizeof(int64_t), Ops::Base::GetUbBlockSize()));
109- assistBuffer_, AlignUp((MAX_LIST_TENSOR_CNT + 1) * sizeof(int64_t), Ops::Base::GetUbBlockSize()));
110}108}
111 109 
112template <typename T>110template <typename T>
@@ -135,15 +133,14 @@ __aicore__ inline void DynamicStitchScatterSimt<T>::Process()
135 __ubuf__ int64_t* tensorCumsumAddr = (__ubuf__ int64_t*)tensorCumsumListLocalTensor.GetPhyAddr();133 __ubuf__ int64_t* tensorCumsumAddr = (__ubuf__ int64_t*)tensorCumsumListLocalTensor.GetPhyAddr();
136 134 
137 if (THREAD_NUM / (curTensorNum_ * sliceSize_) > 0) {135 if (THREAD_NUM / (curTensorNum_ * sliceSize_) > 0) {
138- asc_vf_call<SimtDataCopy<T>>(136+ asc_vf_call<SimtDataCopy<T>>(dim3{curTensorNum_, THREAD_NUM / (curTensorNum_ * sliceSize_), sliceSize_},
139- dim3{curTensorNum_, THREAD_NUM / (curTensorNum_ * sliceSize_), sliceSize_}, xGmAddr_,137+ xGmAddr_, startTensorIndex_, endTensorIndex_, startOffset_, endOffset_,
140- startTensorIndex_, endTensorIndex_, startOffset_, endOffset_, tensorCumsumAddr, workspaceGmAddr,138+ tensorCumsumAddr, workspaceGmAddr, yGmAddr, sliceSize_);
141- yGmAddr, sliceSize_);
142 } else {139 } else {
143 asc_vf_call<SimtDataCopy<T>>(140 asc_vf_call<SimtDataCopy<T>>(
144- dim3{FIRST_DIM_THREAD_NUM, THREAD_NUM / (FIRST_DIM_THREAD_NUM * sliceSize_), sliceSize_},141+ dim3{FIRST_DIM_THREAD_NUM, THREAD_NUM / (FIRST_DIM_THREAD_NUM * sliceSize_), sliceSize_}, xGmAddr_,
145- xGmAddr_, startTensorIndex_, endTensorIndex_, startOffset_, endOffset_, tensorCumsumAddr,142+ startTensorIndex_, endTensorIndex_, startOffset_, endOffset_, tensorCumsumAddr, workspaceGmAddr,
146- workspaceGmAddr, yGmAddr, sliceSize_);143+ yGmAddr, sliceSize_);
147 }144 }
148 }145 }
149}146}
Mconversion/fill/op_api/aclnn_fill_scalar.h+5-5
@@ -33,8 +33,8 @@ extern "C" {
33 * @param [out] executor: 返回op执行器，包含算子计算流程。33 * @param [out] executor: 返回op执行器，包含算子计算流程。
34 * @return aclnnStatus: 返回状态码。34 * @return aclnnStatus: 返回状态码。
35 */35 */
36-ACLNN_API aclnnStatus aclnnInplaceFillScalarGetWorkspaceSize(36+ACLNN_API aclnnStatus aclnnInplaceFillScalarGetWorkspaceSize(aclTensor* selfRef, const aclScalar* value,
37- aclTensor* selfRef, const aclScalar* value, uint64_t* workspaceSize, aclOpExecutor** executor);37+ uint64_t* workspaceSize, aclOpExecutor** executor);
38/**38/**
39 * @brief aclnnInplaceFillScalar的第二段接口，用于执行计算。39 * @brief aclnnInplaceFillScalar的第二段接口，用于执行计算。
40 *40 *
@@ -48,11 +48,11 @@ ACLNN_API aclnnStatus aclnnInplaceFillScalarGetWorkspaceSize(
48 * @param [in] executor: op执行器，包含了算子计算流程。48 * @param [in] executor: op执行器，包含了算子计算流程。
49 * @return aclnnStatus: 返回状态码。49 * @return aclnnStatus: 返回状态码。
50 */50 */
51-ACLNN_API aclnnStatus51+ACLNN_API aclnnStatus aclnnInplaceFillScalar(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
52-aclnnInplaceFillScalar(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);52+ aclrtStream stream);
53 53 
54#ifdef __cplusplus54#ifdef __cplusplus
55}55}
56#endif56#endif
57 57 
58-#endif // OP_API_INC_FILL_H_58+#endif // OP_API_ACLNN_FILL_SCALAR_H_
Mconversion/fill/op_api/aclnn_fill_tensor.h+5-6
@@ -8,7 +8,6 @@
8 * See LICENSE in the root of the software repository for the full text of the License.8 * See LICENSE in the root of the software repository for the full text of the License.
9 */9 */
10 10 
11- 
12#ifndef OP_API_ACLNN_FILL_TENSOR_H_11#ifndef OP_API_ACLNN_FILL_TENSOR_H_
13#define OP_API_ACLNN_FILL_TENSOR_H_12#define OP_API_ACLNN_FILL_TENSOR_H_
14 13 
@@ -45,8 +44,8 @@ extern "C" {
45 * @param [out] executor: 返回op执行器，包含算子计算流程。44 * @param [out] executor: 返回op执行器，包含算子计算流程。
46 * @return aclnnStatus: 返回状态码。45 * @return aclnnStatus: 返回状态码。
47 */46 */
48-ACLNN_API aclnnStatus aclnnInplaceFillTensorGetWorkspaceSize(47+ACLNN_API aclnnStatus aclnnInplaceFillTensorGetWorkspaceSize(aclTensor* selfRef, const aclTensor* value,
49- aclTensor* selfRef, const aclTensor* value, uint64_t* workspaceSize, aclOpExecutor** executor);48+ uint64_t* workspaceSize, aclOpExecutor** executor);
50 49 
51/**50/**
52 * @brief aclnnInplaceFillTensor的第二段接口，用于执行计算。51 * @brief aclnnInplaceFillTensor的第二段接口，用于执行计算。
@@ -56,11 +55,11 @@ ACLNN_API aclnnStatus aclnnInplaceFillTensorGetWorkspaceSize(
56 * @param [in] stream: acl stream流。55 * @param [in] stream: acl stream流。
57 * @return aclnnStatus: 返回状态码。56 * @return aclnnStatus: 返回状态码。
58 */57 */
59-ACLNN_API aclnnStatus58+ACLNN_API aclnnStatus aclnnInplaceFillTensor(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
60-aclnnInplaceFillTensor(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor, aclrtStream stream);59+ aclrtStream stream);
61 60 
62#ifdef __cplusplus61#ifdef __cplusplus
63}62}
64#endif63#endif
65 64 
66-#endif // OP_API_INC_LEVEL2_ACLNN_FILL_TENSOR_H_65+#endif // OP_API_ACLNN_FILL_TENSOR_H_