Star1296
1410
代码
代码
Issues108
Pull Requests213
Actions
项目讨论
Wiki
分析
Star1296
1410
已合并
添加 ascend950 算子 cross #2232
sakuraqqz创建于 4月15日
添加 ascend950 算子 cross #2232
已合并
sakuraqqz创建于 4月15日
从sakuraqqz/ops-mathmaster合入到cann/ops-mathmaster
共 14 个文件变更+814-95
Mmath/cross/CMakeLists.txt+9-7
@@ -1,11 +1,13 @@
1-# ---------------------------------------------------------------------------------------------------------1+# ----------------------------------------------------------------------------
2-# Copyright (c) 2025 Huawei Technologies Co., Ltd.2+# Copyright (c) 2026 Huawei Technologies Co., Ltd.
3-# This program is free software, you can redistribute it and/or modify it under the terms and conditions of 3+# This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4# CANN Open Software License Agreement Version 2.0 (the "License").4# CANN Open Software License Agreement Version 2.0 (the "License").
5# Please refer to the License for details. You may not use this file except in compliance with the License.5# Please refer to the License for details. You may not use this file except in compliance with the License.
6-# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED, 6+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7-# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE. 7+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
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-add_all_modules_sources(OPTYPE cross ACLNNTYPE aclnn_exclude)11+set(SUPPORT_COMPUTE_UNIT "ascend950")
12+set(SUPPORT_TILING_DIR "arch35")
13+add_all_modules_sources(OPTYPE cross ACLNNTYPE aclnn_exclude COMPUTE_UNIT ${SUPPORT_COMPUTE_UNIT} TILING_DIR ${SUPPORT_TILING_DIR} DISABLE_IN_OPP TRUE)
Mmath/cross/op_api/aclnn_linalg_cross.cpp+8-4
@@ -16,6 +16,7 @@
16#include "opdev/op_dfx.h"16#include "opdev/op_dfx.h"
17#include "conversion/broadcast_to/op_api/broadcast_to.h"17#include "conversion/broadcast_to/op_api/broadcast_to.h"
18#include "op_api/op_api_def.h"18#include "op_api/op_api_def.h"
19+#include "op_api/aclnn_check.h"
19 20 
20using namespace op;21using namespace op;
21#ifdef __cplusplus22#ifdef __cplusplus
@@ -43,7 +44,8 @@ inline static bool CheckNotNull(const aclTensor *self, const aclTensor *other, c
43static bool CheckDtypeValid(const aclTensor *self, const aclTensor *other, const aclTensor *out)44static bool CheckDtypeValid(const aclTensor *self, const aclTensor *other, const aclTensor *out)
44{45{
45 // 检查self的数据类型是否在linalg cross算子的支持列表内46 // 检查self的数据类型是否在linalg cross算子的支持列表内
46- if (op::GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201) {47+ if (op::GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201 ||
48+ op::GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_3510) {
47 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST_910B, return false);49 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST_910B, return false);
48 } else {50 } else {
49 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST, return false);51 OP_CHECK_DTYPE_NOT_SUPPORT(self, DTYPE_SUPPORT_LIST, return false);
@@ -161,7 +163,7 @@ static aclnnStatus ExecLinalgCrossGetWorkspaceSize(const aclTensor *self, const
161 auto dimInner = dim;163 auto dimInner = dim;
162 164 
163 // 若形状不同，则做broadcast165 // 若形状不同，则做broadcast
164- if (selfContiguous->GetViewShape() != otherContiguous->GetViewShape()) {166+ if (selfContiguous->GetViewShape() != otherContiguous->GetViewShape() && !IsRegBase()) {
165 selfBroadCast = BroadcastTensor(out->GetViewShape(), selfContiguous, uniqueExecutor.get());167 selfBroadCast = BroadcastTensor(out->GetViewShape(), selfContiguous, uniqueExecutor.get());
166 CHECK_RET(selfBroadCast != nullptr, ACLNN_ERR_INNER_NULLPTR);168 CHECK_RET(selfBroadCast != nullptr, ACLNN_ERR_INNER_NULLPTR);
167 otherBroadCast = BroadcastTensor(out->GetViewShape(), otherContiguous, uniqueExecutor.get());169 otherBroadCast = BroadcastTensor(out->GetViewShape(), otherContiguous, uniqueExecutor.get());
@@ -189,7 +191,8 @@ static aclnnStatus ExecLinalgCrossGetWorkspaceSize(const aclTensor *self, const
189 }191 }
190 192 
191 if (self->GetDataType() == op::DataType::DT_BF16 &&193 if (self->GetDataType() == op::DataType::DT_BF16 &&
192- op::GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_2201) {194+ op::GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_2201 &&
195+ op::GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_3510) {
193 selfBroadCast = l0op::Cast(selfBroadCast, op::DataType::DT_FLOAT, uniqueExecutor.get());196 selfBroadCast = l0op::Cast(selfBroadCast, op::DataType::DT_FLOAT, uniqueExecutor.get());
194 CHECK_RET(selfBroadCast != nullptr, ACLNN_ERR_INNER_NULLPTR);197 CHECK_RET(selfBroadCast != nullptr, ACLNN_ERR_INNER_NULLPTR);
195 otherBroadCast = l0op::Cast(otherBroadCast, op::DataType::DT_FLOAT, uniqueExecutor.get());198 otherBroadCast = l0op::Cast(otherBroadCast, op::DataType::DT_FLOAT, uniqueExecutor.get());
@@ -201,7 +204,8 @@ static aclnnStatus ExecLinalgCrossGetWorkspaceSize(const aclTensor *self, const
201 CHECK_RET(crossOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);204 CHECK_RET(crossOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
202 205 
203 if (self->GetDataType() == op::DataType::DT_BF16 &&206 if (self->GetDataType() == op::DataType::DT_BF16 &&
204- op::GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_2201) {207+ op::GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_2201 &&
208+ op::GetCurrentPlatformInfo().GetCurNpuArch() != NpuArch::DAV_3510) {
205 crossOpOut = l0op::Cast(crossOpOut, op::DataType::DT_BF16, uniqueExecutor.get());209 crossOpOut = l0op::Cast(crossOpOut, op::DataType::DT_BF16, uniqueExecutor.get());
206 CHECK_RET(crossOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);210 CHECK_RET(crossOpOut != nullptr, ACLNN_ERR_INNER_NULLPTR);
207 }211 }
Mmath/cross/op_api/cross.cpp+27-5
@@ -12,14 +12,21 @@
12#include "opdev/aicpu/aicpu_task.h"12#include "opdev/aicpu/aicpu_task.h"
13#include "opdev/make_op_executor.h"13#include "opdev/make_op_executor.h"
14#include "opdev/op_dfx.h"14#include "opdev/op_dfx.h"
15+#include "op_api/aclnn_check.h"
15 16 
16using namespace op;17using namespace op;
17namespace l0op {18namespace l0op {
18OP_TYPE_REGISTER(Cross);19OP_TYPE_REGISTER(Cross);
19OP_TYPE_REGISTER(CrossV2);20OP_TYPE_REGISTER(CrossV2);
20 21 
21-static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {22+static const int64_t FORMER_LLENGTH_BOUND = 10240;
22- op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT8,23+static const int64_t TAIL_LLENGTH_BOUND = 64;
24+ 
25+static const std::initializer_list<op::DataType> REGBASE_AICORE_DTYPE_SUPPORT_LIST = {
26+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_BF16};
27+ 
28+static const std::initializer_list<op::DataType> AICORE_DTYPE_SUPPORT_LIST = {
29+ op::DataType::DT_FLOAT, op::DataType::DT_FLOAT16, op::DataType::DT_INT8,
23 op::DataType::DT_INT16, op::DataType::DT_INT32, op::DataType::DT_UINT8};30 op::DataType::DT_INT16, op::DataType::DT_INT32, op::DataType::DT_UINT8};
24 31 
25static const std::initializer_list<op::DataType> V2_AICORE_DTYPE_SUPPORT_LIST = {32static const std::initializer_list<op::DataType> V2_AICORE_DTYPE_SUPPORT_LIST = {
@@ -30,14 +37,22 @@ static const std::initializer_list<op::DataType> V2_AICORE_DTYPE_SUPPORT_LIST =
30// 根据芯片类型、dtype判断算子是否支持走aicore37// 根据芯片类型、dtype判断算子是否支持走aicore
31inline static bool IsAiCoreSupport(const aclTensor *self)38inline static bool IsAiCoreSupport(const aclTensor *self)
32{39{
33- // Cross只需要判断dtype40+ // Cross只需要判断dtype是否支持即可
41+ auto curArch = GetCurrentPlatformInfo().GetCurNpuArch();
42+ if (IsRegBase(curArch)) {
43+ return CheckType(self->GetDataType(), REGBASE_AICORE_DTYPE_SUPPORT_LIST);
44+ }
34 return CheckType(self->GetDataType(), AICORE_DTYPE_SUPPORT_LIST);45 return CheckType(self->GetDataType(), AICORE_DTYPE_SUPPORT_LIST);
35}46}
36 47 
37// 根据芯片类型、dtype判断算子是否支持走aicore48// 根据芯片类型、dtype判断算子是否支持走aicore
38static inline bool IsV2AiCoreSupport(const aclTensor *self) {49static inline bool IsV2AiCoreSupport(const aclTensor *self) {
50+ if (IsRegBase()){
51+ return false;
52+ }
39 // 获取芯片类型53 // 获取芯片类型
40- if (op::GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201) {54+ if (op::GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_2201 ||
55+ op::GetCurrentPlatformInfo().GetCurNpuArch() == NpuArch::DAV_3510) {
41 return CheckType(self->GetDataType(), V2_AICORE_DTYPE_SUPPORT_LIST);56 return CheckType(self->GetDataType(), V2_AICORE_DTYPE_SUPPORT_LIST);
42 }57 }
43 return false;58 return false;
@@ -85,7 +100,14 @@ inline static const aclTensor *CrossAiCpu(const aclTensor *self, const aclTensor
85 100 
86const aclTensor *Cross(const aclTensor *self, const aclTensor *other, int64_t dim, aclOpExecutor *executor)101const aclTensor *Cross(const aclTensor *self, const aclTensor *other, int64_t dim, aclOpExecutor *executor)
87{102{
88- auto crossOut = executor->AllocTensor(self->GetViewShape(), self->GetDataType(), op::Format::FORMAT_ND);103+ Shape broadcastShape;
104+ if (!BroadcastInferShape(self->GetViewShape(), other->GetViewShape(), broadcastShape)) {
105+ OP_LOGE(
106+ ACLNN_ERR_PARAM_INVALID, "Broadcast %s and %s failed.", op::ToString(self->GetViewShape()).GetString(),
107+ op::ToString(other->GetViewShape()).GetString());
108+ return nullptr;
109+ }
110+ auto crossOut = executor->AllocTensor(broadcastShape, self->GetDataType(), op::Format::FORMAT_ND);
89 if (IsV2AiCoreSupport(self)) {111 if (IsV2AiCoreSupport(self)) {
90 return CrossV2AiCore(self, other, dim, crossOut, executor);112 return CrossV2AiCore(self, other, dim, crossOut, executor);
91 } else if (IsAiCoreSupport(self)) {113 } else if (IsAiCoreSupport(self)) {
Amath/cross/op_host/arch35/cross_tiling.cpp+242-0
@@ -0,0 +1,242 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file cross_tiling.cpp
13+ * \brief cross tiling
14+ */
15+ 
16+#include "cross_tiling.h"
17+ 
18+using namespace ge;
19+ 
20+namespace optiling {
21+static constexpr uint64_t INPUT_X1 = 0;
22+static constexpr uint64_t INPUT_X2 = 1;
23+static constexpr uint64_t DIM = 0;
24+static constexpr int64_t INT_MAX = 2147483647;
25+ 
26+ge::graphStatus CrossTiling::GetPlatformInfo()
27+{
28+ OP_LOGD(context_, "CrossTiling GetPlatformInfo.");
29+ compileInfo_ = static_cast<const CrossCompileInfo*>(context_->GetCompileInfo());
30+ OP_CHECK_NULL_WITH_CONTEXT(context_, compileInfo_);
31+ return ge::GRAPH_SUCCESS;
32+}
33+ 
34+// 1. 基础参数校验（输入形状、维度合法性、尺寸检查）
35+ge::graphStatus CrossTiling::CheckBaseShapeAndAttrs() {
36+ OP_LOGD(context_, "CrossTiling CheckBaseShapeAndAttrs.");
37+ const gert::StorageShape* shape1 = context_->GetInputShape(INPUT_X1);
38+ const gert::StorageShape* shape2 = context_->GetInputShape(INPUT_X2);
39+ OP_CHECK_NULL_WITH_CONTEXT(context_, shape1);
40+ OP_CHECK_NULL_WITH_CONTEXT(context_, shape2);
41+ 
42+ auto s1 = shape1->GetStorageShape();
43+ auto s2 = shape2->GetStorageShape();
44+ auto attrs = context_->GetAttrs();
45+ OP_CHECK_IF((attrs == nullptr), OP_LOGE(context_, "Get attrs Failed."), return ge::GRAPH_FAILED);
46+ 
47+ int64_t dim = *(attrs->GetAttrPointer<int64_t>(DIM));
48+ dimNum1_ = s1.GetDimNum();
49+ dimNum2_ = s2.GetDimNum();
50+ 
51+ for (int64_t i = 0; i < dimNum1_; i++) {
52+ x1Dims_[i] = s1.GetDim(i);
53+ }
54+ for (int64_t i = 0; i < dimNum2_; i++) {
55+ x2Dims_[i] = s2.GetDim(i);
56+ }
57+ 
58+ OP_CHECK_IF((dim < -dimNum1_ || dim >= dimNum1_),
59+ OP_LOGE(context_, "dim must be in [%ld, %ld], dim: [%ld].", -dimNum1_, dimNum1_ - 1, dim),
60+ return ge::GRAPH_FAILED);
61+ OP_CHECK_IF((dimNum1_ != dimNum2_),
62+ OP_LOGE(context_, "x1 and x2 dim count mismatch: %ld vs %ld.", dimNum1_, dimNum2_),
63+ return ge::GRAPH_FAILED);
64+ 
65+ // 标准化维度
66+ int64_t tempDimNum = dimNum1_ == 0 ? 1 : dimNum1_;
67+ normalizedDim_ = (dim + tempDimNum) % tempDimNum;
68+ 
69+ // 校验指定维度必须为3
70+ int64_t dimSize1 = x1Dims_[normalizedDim_];
71+ int64_t dimSize2 = x2Dims_[normalizedDim_];
72+ OP_CHECK_IF((dimSize1 != 3), OP_LOGE(context_, "x1 dim[%ld] must be 3, got %ld.", normalizedDim_, dimSize1), return ge::GRAPH_FAILED);
73+ OP_CHECK_IF((dimSize2 != 3), OP_LOGE(context_, "x2 dim[%ld] must be 3, got %ld.", normalizedDim_, dimSize2), return ge::GRAPH_FAILED);
74+ 
75+ return ge::GRAPH_SUCCESS;
76+}
77+ 
78+// 2. 广播兼容性校验 + 计算合并形状
79+ge::graphStatus CrossTiling::CheckBroadcastAndMergeShape() {
80+ OP_LOGD(context_, "CrossTiling CheckBroadcastAndMergeShape.");
81+ for (int64_t i = 0; i < dimNum1_; i++) {
82+ if (i == normalizedDim_) continue;
83+ int64_t size1 = x1Dims_[i];
84+ int64_t size2 = x2Dims_[i];
85+ OP_CHECK_IF((size1 != size2 && size1 != 1 && size2 != 1),
86+ OP_LOGE(context_, "Shapes not broadcastable at dim %ld: %ld vs %ld.", i, size1, size2),
87+ return ge::GRAPH_FAILED);
88+ }
89+ 
90+ // 计算合并形状 & 输出总大小
91+ dim_ = normalizedDim_;
92+ dimNum_ = dimNum1_;
93+ ySize_ = 1;
94+ for (int64_t i = 0; i < dimNum1_; i++) {
95+ if (i == normalizedDim_) {
96+ mergedShape_[i] = 1;
97+ ySize_ *= x1Dims_[i];
98+ } else {
99+ mergedShape_[i] = std::max(x1Dims_[i], x2Dims_[i]);
100+ ySize_ *= mergedShape_[i];
101+ }
102+ }
103+ tilingData_.usedInt64 = ySize_ > INT_MAX;
104+ return ge::GRAPH_SUCCESS;
105+}
106+ 
107+// 3. 计算步长stride + 最终向量参数
108+ge::graphStatus CrossTiling::CalcStrideAndVectors() {
109+ OP_LOGD(context_, "CrossTiling CalcStrideAndVectors.");
110+ int64_t stride[4] = {1, 1, 1, 1};
111+ for (int64_t i = dimNum1_ - 1; i >= 0; i--) {
112+ x1Stride_[i] = (x1Dims_[i] == 1) ? 0 : stride[0];
113+ x2Stride_[i] = (x2Dims_[i] == 1) ? 0 : stride[1];
114+ mergedStride_[i] = stride[2];
115+ yStride_[i] = stride[3];
116+ 
117+ stride[0] *= x1Dims_[i];
118+ stride[1] *= x2Dims_[i];
119+ stride[2] *= mergedShape_[i];
120+ stride[3] *= (i == dim_ ? 3 : mergedShape_[i]);
121+ }
122+ 
123+ // 计算维度步长
124+ dimStride_ = 1;
125+ for (int64_t i = normalizedDim_ + 1; i < dimNum1_; i++) {
126+ dimStride_ *= mergedShape_[i];
127+ }
128+ 
129+ // 计算总向量数
130+ totalVectors_ = 1;
131+ for (int64_t i = 0; i < dimNum1_; i++) {
132+ if (i != normalizedDim_) totalVectors_ *= mergedShape_[i];
133+ }
134+ return ge::GRAPH_SUCCESS;
135+}
136+ 
137+ge::graphStatus CrossTiling::GetShapeAttrsInfo() {
138+ if (CheckBaseShapeAndAttrs() != ge::GRAPH_SUCCESS) return ge::GRAPH_FAILED;
139+ if (CheckBroadcastAndMergeShape() != ge::GRAPH_SUCCESS) return ge::GRAPH_FAILED;
140+ if (CalcStrideAndVectors() != ge::GRAPH_SUCCESS) return ge::GRAPH_FAILED;
141+ return ge::GRAPH_SUCCESS;
142+}
143+ 
144+ge::graphStatus CrossTiling::DoOpTiling()
145+{
146+ OP_LOGD(context_, "CrossTiling DoOpTiling.");
147+
148+ int64_t coreNum = compileInfo_->coreNum;
149+ // coreNum = 1;
150+ int64_t vectorsPerCore = totalVectors_ / coreNum;
151+ int64_t formerCore = totalVectors_ % coreNum;
152+ 
153+ tilingData_.totalVectors = totalVectors_;
154+ tilingData_.vectorsPerCore = vectorsPerCore;
155+ tilingData_.coreNum = coreNum;
156+ tilingData_.dim = dim_;
157+ tilingData_.dimNum = dimNum_;
158+ tilingData_.dimStride = dimStride_;
159+ tilingData_.formerCore = formerCore;
160+
161+ for (int64_t i = 0; i < MAX_DIM; i++) {
162+ if (i < dimNum_) {
163+ tilingData_.mergedStride[i] = mergedStride_[i];
164+ tilingData_.x1Stride[i] = x1Stride_[i];
165+ tilingData_.x2Stride[i] = x2Stride_[i];
166+ tilingData_.yStride[i] = yStride_[i];
167+ } else {
168+ tilingData_.mergedStride[i] = 1;
169+ tilingData_.x1Stride[i] = 0;
170+ tilingData_.x2Stride[i] = 0;
171+ tilingData_.yStride[i] = 0;
172+ }
173+ }
174+ 
175+ blockDim_ = (totalVectors_ < coreNum) ? totalVectors_ : coreNum;
176+ 
177+ return ge::GRAPH_SUCCESS;
178+}
179+ 
180+ge::graphStatus CrossTiling::PostTiling()
181+{
182+ OP_LOGD(context_, "CrossTiling PostTiling.");
183+ 
184+ auto workspaces = context_->GetWorkspaceSizes(1);
185+ OP_CHECK_NULL_WITH_CONTEXT(context_, workspaces);
186+ workspaces[0] = 0;
187+ 
188+ auto res = context_->SetBlockDim(static_cast<uint32_t>(blockDim_));
189+ OP_CHECK_IF((res != ge::GRAPH_SUCCESS), OP_LOGE(context_, "SetBlockDim failed."), return ge::GRAPH_FAILED);
190+ 
191+ errno_t ret = memcpy_s(
192+ context_->GetRawTilingData()->GetData(), context_->GetRawTilingData()->GetCapacity(), &tilingData_,
193+ sizeof(CrossRegbaseTilingData));
194+ if (ret != EOK) {
195+ OP_LOGE(context_->GetNodeName(), "memcpy_s failed, ret=%d", ret);
196+ return ge::GRAPH_FAILED;
197+ }
198+ context_->GetRawTilingData()->SetDataSize(sizeof(CrossRegbaseTilingData));
199+ 
200+ return ge::GRAPH_SUCCESS;
201+}
202+ 
203+static ge::graphStatus Tiling4Cross(gert::TilingContext* context)
204+{
205+ OP_LOGD(context, "Tiling4Cross start.");
206+ 
207+ CrossTiling crossTiling(context);
208+ auto ret = crossTiling.DoTiling();
209+ OP_CHECK_IF((ret == ge::GRAPH_FAILED), OP_LOGD(context, "Tiling4Cross failed!"), return ge::GRAPH_FAILED);
210+ OP_LOGD(context, "Tiling4Cross end.");
211+ return ge::GRAPH_SUCCESS;
212+}
213+ 
214+static ge::graphStatus TilingPrepare4CrossAscendc(gert::TilingParseContext* context)
215+{
216+ OP_LOGD(context->GetNodeName(), "Enter TilingPrepare4CrossAscendc.");
217+ 
218+ auto compileInfo = context->GetCompiledInfo<CrossCompileInfo>();
219+ 
220+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
221+ auto platformInfo = context->GetPlatformInfo();
222+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
223+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
224+ 
225+ compileInfo->coreNum = ascendcPlatform.GetCoreNumAiv();
226+ OP_CHECK_IF(
227+ (compileInfo->coreNum <= 0), OP_LOGE(context->GetNodeName(), "core num is negative."), return ge::GRAPH_FAILED);
228+ 
229+ OP_LOGD(context->GetNodeName(), "Exit TilingPrepare4CrossAscendc.");
230+ return ge::GRAPH_SUCCESS;
231+}
232+ 
233+static ge::graphStatus TilingPrepare4Cross(gert::TilingParseContext* context)
234+{
235+ auto compile_info = context->GetCompiledInfo<CrossCompileInfo>();
236+ OP_CHECK_NULL_WITH_CONTEXT(context, compile_info);
237+ OP_LOGD("TilingPrepare4Cross", "Ascend C TilingPrepare4Cross success.");
238+ return TilingPrepare4CrossAscendc(context);
239+}
240+ 
241+IMPL_OP_OPTILING(Cross).Tiling(Tiling4Cross).TilingParse<CrossCompileInfo>(TilingPrepare4Cross);
242+} // namespace optiling
Amath/cross/op_host/arch35/cross_tiling.h+72-0
@@ -0,0 +1,72 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file cross_tiling.h
13+ * \brief cross tiling header
14+ */
15+#ifndef OPS_BUILT_IN_OP_TILING_RUNTIME_CROSS_H_
16+#define OPS_BUILT_IN_OP_TILING_RUNTIME_CROSS_H_
17+ 
18+#include <cstdint>
19+#include "register/tilingdata_base.h"
20+#include "op_host/tiling_base.h"
21+#include "register/op_impl_registry.h"
22+#include "platform/platform_ascendc.h"
23+#include "log/log.h"
24+#include "../../op_kernel/arch35/cross_struct.h"
25+ 
26+namespace optiling {
27+ 
28+struct CrossCompileInfo {
29+ int64_t coreNum = 0;
30+};
31+ 
32+class CrossTiling : public Ops::Math::OpTiling::TilingBaseClass {
33+public:
34+ explicit CrossTiling(gert::TilingContext* context) : TilingBaseClass(context)
35+ {}
36+ 
37+protected:
38+ bool IsCapable() override { return true; }
39+ ge::graphStatus GetPlatformInfo() override;
40+ ge::graphStatus GetShapeAttrsInfo() override;
41+ ge::graphStatus DoOpTiling() override;
42+ ge::graphStatus DoLibApiTiling() override { return ge::GRAPH_SUCCESS; }
43+ uint64_t GetTilingKey() const override { return 0; }
44+ ge::graphStatus GetWorkspaceSize() override { return ge::GRAPH_SUCCESS; }
45+ ge::graphStatus PostTiling() override;
46+ ge::graphStatus CheckBaseShapeAndAttrs();
47+ ge::graphStatus CheckBroadcastAndMergeShape();
48+ ge::graphStatus CalcStrideAndVectors();
49+ 
50+private:
51+ const CrossCompileInfo* compileInfo_;
52+ CrossRegbaseTilingData tilingData_{};
53+ uint32_t blockDim_{1};
54+ int64_t totalVectors_{0};
55+ int64_t dimNum_{0};
56+ int64_t dim_{0};
57+ int64_t dimNum1_{0};
58+ int64_t dimNum2_{0};
59+ int64_t dimStride_{1};
60+ int64_t normalizedDim_{0};
61+ int64_t ySize_{0};
62+ int64_t x1Dims_[MAX_DIM] = {0};
63+ int64_t x2Dims_[MAX_DIM] = {0};
64+ int64_t x1Stride_[MAX_DIM] = {0};
65+ int64_t x2Stride_[MAX_DIM] = {0};
66+ int64_t mergedStride_[MAX_DIM] = {1};
67+ int64_t mergedShape_[MAX_DIM] = {1};
68+ int64_t yStride_[MAX_DIM] = {0};
69+};
70+ 
71+} // namespace optiling
72+#endif // OPS_BUILT_IN_OP_TILING_RUNTIME_CROSS_H_
Amath/cross/op_host/config/ascend950/cross_binary.json+119-0
@@ -0,0 +1,119 @@
1+{
2+ "op_type": "Cross",
3+ "op_list": [
4+ {
5+ "bin_filename": "Cross_bf16_high_performance",
6+ "inputs": [
7+ {
8+ "name": "x1",
9+ "index": 0,
10+ "dtype": "bfloat16",
11+ "format": "ND",
12+ "paramType": "required",
13+ "shape": [-2]
14+ },
15+ {
16+ "name": "x2",
17+ "index": 1,
18+ "dtype": "bfloat16",
19+ "format": "ND",
20+ "paramType": "required",
21+ "shape": [-2]
22+ }
23+ ],
24+ "outputs": [
25+ {
26+ "name": "y",
27+ "index": 0,
28+ "dtype": "bfloat16",
29+ "format": "ND",
30+ "paramType": "required",
31+ "shape": [-2]
32+ }
33+ ],
34+ "attrs": [
35+ {
36+ "name": "dim",
37+ "dtype": "int",
38+ "value": null
39+ }
40+ ]
41+ },
42+ {
43+ "bin_filename": "Cross_fp16_high_performance",
44+ "inputs": [
45+ {
46+ "name": "x1",
47+ "index": 0,
48+ "dtype": "float16",
49+ "format": "ND",
50+ "paramType": "required",
51+ "shape": [-2]
52+ },
53+ {
54+ "name": "x2",
55+ "index": 1,
56+ "dtype": "float16",
57+ "format": "ND",
58+ "paramType": "required",
59+ "shape": [-2]
60+ }
61+ ],
62+ "outputs": [
63+ {
64+ "name": "y",
65+ "index": 0,
66+ "dtype": "float16",
67+ "format": "ND",
68+ "paramType": "required",
69+ "shape": [-2]
70+ }
71+ ],
72+ "attrs": [
73+ {
74+ "name": "dim",
75+ "dtype": "int",
76+ "value": null
77+ }
78+ ]
79+ },
80+ {
81+ "bin_filename": "Cross_fp32_high_performance",
82+ "inputs": [
83+ {
84+ "name": "x1",
85+ "index": 0,
86+ "dtype": "float32",
87+ "format": "ND",
88+ "paramType": "required",
89+ "shape": [-2]
90+ },
91+ {
92+ "name": "x2",
93+ "index": 1,
94+ "dtype": "float32",
95+ "format": "ND",
96+ "paramType": "required",
97+ "shape": [-2]
98+ }
99+ ],
100+ "outputs": [
101+ {
102+ "name": "y",
103+ "index": 0,
104+ "dtype": "float32",
105+ "format": "ND",
106+ "paramType": "required",
107+ "shape": [-2]
108+ }
109+ ],
110+ "attrs": [
111+ {
112+ "name": "dim",
113+ "dtype": "int",
114+ "value": null
115+ }
116+ ]
117+ }
118+ ]
119+}
Amath/cross/op_host/cross_def.cpp+53-0
@@ -0,0 +1,53 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file cross_def.cpp
13+ * \brief cross_def infomation
14+ */
15+ 
16+#include <cstdint>
17+#include "register/op_def_registry.h"
18+ 
19+namespace ops {
20+class Cross : public OpDef {
21+public:
22+ explicit Cross(const char* name) : OpDef(name)
23+ {
24+ this->Input("x1")
25+ .ParamType(REQUIRED)
26+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
27+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
28+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
29+ this->Input("x2")
30+ .ParamType(REQUIRED)
31+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
32+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
33+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
34+ this->Output("y")
35+ .ParamType(REQUIRED)
36+ .DataType({ge::DT_FLOAT16, ge::DT_FLOAT, ge::DT_BF16})
37+ .Format({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND})
38+ .UnknownShapeFormat({ge::FORMAT_ND, ge::FORMAT_ND, ge::FORMAT_ND});
39+ this->Attr("dim").AttrType(OPTIONAL).Int(-1);
40+ 
41+ OpAICoreConfig aicoreConfig;
42+ aicoreConfig.DynamicCompileStaticFlag(true)
43+ .DynamicFormatFlag(false)
44+ .DynamicRankSupportFlag(true)
45+ .DynamicShapeSupportFlag(true)
46+ .NeedCheckSupportFlag(false)
47+ .ExtendCfgInfo("opFile.value", "cross_apt");
48+ this->AICore().AddConfig("ascend950", aicoreConfig);
49+ }
50+};
51+ 
52+OP_ADD(Cross);
53+} // namespace ops
Dmath/cross/op_host/op_api/aclnn_linalg_cross.h+0-58
@@ -1,58 +0,0 @@
1-/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9- */
10- 
11-#ifndef OP_API_INC_LINALG_CROSS_H_
12-#define OP_API_INC_LINALG_CROSS_H_
13- 
14-#include "aclnn/aclnn_base.h"
15-#include "aclnn_util.h"
16- 
17-#ifdef __cplusplus
18-extern "C" {
19-#endif
20- 
21-/**
22- * @brief aclnnLinalgCross的第一段接口，根据具体的计算流程，计算workspace大小。
23- * @domain aclnn_math
24- *
25- * 算子功能： 对输入Tensor完成linalg_cross操作
26- * @param [in] self: npu device侧的aclTensor, 数据类型支持INT8、INT16、INT32、UINT8、FLOAT16、FLOAT、FLOAT64、
27- * COMPLEX64、COMPLEX128、BFLOAT16, shape为非空，支持非连续的Tensor，与other有broadcast关系，数据格式支持ND。
28- * @param [in] other: npu device侧的aclTensor, 数据类型支持INT8、INT16、INT32、UINT8、FLOAT16、FLOAT、FLOAT64、
29- * COMPLEX64、COMPLEX128、BFLOAT16, shape为非空，支持非连续的Tensor，与self有broadcast关系，数据格式支持ND。
30- * @param [in] dim: 输入INT，默认值为-1。
31- * @param [in] out: npu device侧的aclTensor, 数据类型支持INT8、INT16、INT32、UINT8、FLOAT16、FLOAT、FLOAT64、
32- * COMPLEX64、COMPLEX128、BFLOAT16, shape与self和other的broadcast相同，数据格式支持ND。
33- * @param [out] workspaceSize: 返回用户需要在npu device侧申请的workspace大小。
34- * @param [out] executor: 返回op执行器，包含算子计算流程。
35- * @return aclnnStatus: 返回状态码
36- */
37-ACLNN_API aclnnStatus aclnnLinalgCrossGetWorkspaceSize(const aclTensor* self, const aclTensor* other, int64_t dim,
38- aclTensor* out, uint64_t* workspaceSize,
39- aclOpExecutor** executor);
40- 
41-/**
42- * @brief: aclnnLinalgCross的第二段接口，用于执行计算
43- *
44- * 算子功能： 对输入Tensor完成linalg_cross操作
45- * @param [in] workspace: 在npu device侧申请的workspace内存起址。
46- * @param [in] workspaceSize: 在npu device侧申请的workspace大小，由第一段接口aclnnLinalgCrossGetWorkspaceSize获取。
47- * @param [in] stream: acl stream流。
48- * @param [in] executor: op执行器，包含了算子计算流程。
49- * @return aclnnStatus: 返回状态码。
50- */
51-ACLNN_API aclnnStatus aclnnLinalgCross(void* workspace, uint64_t workspaceSize, aclOpExecutor* executor,
52- aclrtStream stream);
53- 
54-#ifdef __cplusplus
55-}
56-#endif
57- 
58-#endif // OP_API_INC_LINALG_CROSS_H_
Dmath/cross/op_host/op_api/cross.h+0-21
@@ -1,21 +0,0 @@
1-/**
2- * Copyright (c) 2025 Huawei Technologies Co., Ltd.
3- * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4- * CANN Open Software License Agreement Version 2.0 (the "License").
5- * Please refer to the License for details. You may not use this file except in compliance with the License.
6- * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7- * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8- * See LICENSE in the root of the software repository for the full text of the License.
9- */
10- 
11-#ifndef PTA_NPU_OP_API_INC_LEVEL0_OP_CROSS_OP_H_
12-#define PTA_NPU_OP_API_INC_LEVEL0_OP_CROSS_OP_H_
13- 
14-#include "opdev/op_executor.h"
15- 
16-namespace l0op {
17-const aclTensor *Cross(const aclTensor *self, const aclTensor *other, int64_t dim,
18- aclOpExecutor *executor);
19-}
20- 
21-#endif // PTA_NPU_OP_API_INC_LEVEL0_OP_CROSS_OP_H_
Amath/cross/op_kernel/arch35/cross.h+199-0
@@ -0,0 +1,199 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#ifndef CROSS_H
12+#define CROSS_H
13+ 
14+#include "kernel_operator.h"
15+#include "simt_api/asc_simt.h"
16+#include "cross_struct.h"
17+#include "cross_tiling_key.h"
18+ 
19+using namespace AscendC;
20+ 
21+constexpr int32_t CROSS_THREAD_DIM = 1024;
22+constexpr int32_t CROSS_THREAD_DIM_FOR_INT64 = 512;
23+ 
24+struct CrossStridePara {
25+ int32_t ms[8];
26+ int32_t x1s[8];
27+ int32_t x2s[8];
28+ int32_t ys[8];
29+};
30+ 
31+namespace CrossKernel {
32+ 
33+template <typename T>
34+class Cross {
35+public:
36+ __aicore__ inline Cross() {}
37+ __aicore__ inline ~Cross() {}
38+ 
39+ __aicore__ inline void Init(GM_ADDR x1, GM_ADDR x2, GM_ADDR y,
40+ const CrossRegbaseTilingData& tilingData)
41+ {
42+ totalVectors_ = tilingData.totalVectors;
43+ vectorsPerCore_ = tilingData.vectorsPerCore;
44+ coreNum_ = tilingData.coreNum;
45+ dim_ = tilingData.dim;
46+ dimNum_ = tilingData.dimNum;
47+ dimStride_ = tilingData.dimStride;
48+ formerCore_ = tilingData.formerCore;
49+ usedInt64_ = tilingData.usedInt64;
50+ 
51+ for (int i = 0; i < 8; i++) {
52+ para_.ms[i] = tilingData.mergedStride[i];
53+ para_.x1s[i] = tilingData.x1Stride[i];
54+ para_.x2s[i] = tilingData.x2Stride[i];
55+ para_.ys[i] = tilingData.yStride[i];
56+ }
57+ 
58+ x1Gm_.SetGlobalBuffer((__gm__ T*)x1);
59+ x2Gm_.SetGlobalBuffer((__gm__ T*)x2);
60+ yGm_.SetGlobalBuffer((__gm__ T*)y);
61+ }
62+ 
63+ __aicore__ inline void Process()
64+ {
65+ int32_t block_idx = GetBlockIdx();
66+ int32_t startIdx = (block_idx < formerCore_) ? (vectorsPerCore_+1)*block_idx : formerCore_ + vectorsPerCore_*block_idx;
67+ int32_t count = (block_idx < formerCore_) ? (vectorsPerCore_+1) : vectorsPerCore_;
68+ 
69+ if (count <= 0) return;
70+ 
71+ if (!usedInt64_) {
72+ asc_vf_call<SimtCrossCompute<T>>(
73+ dim3(CROSS_THREAD_DIM),
74+ (__gm__ T*)x1Gm_.GetPhyAddr(),
75+ (__gm__ T*)x2Gm_.GetPhyAddr(),
76+ (__gm__ T*)yGm_.GetPhyAddr(),
77+ startIdx, count, dim_, dimNum_, dimStride_, para_
78+ );
79+ } else {
80+ asc_vf_call<SimtCrossComputeInt64<T>>(
81+ dim3(CROSS_THREAD_DIM_FOR_INT64),
82+ (__gm__ T*)x1Gm_.GetPhyAddr(),
83+ (__gm__ T*)x2Gm_.GetPhyAddr(),
84+ (__gm__ T*)yGm_.GetPhyAddr(),
85+ startIdx, count, dim_, dimNum_, dimStride_, para_
86+ );
87+ }
88+ }
89+ 
90+private:
91+ // 普通 int32 偏移版本
92+ template <typename U>
93+ __simt_vf__ LAUNCH_BOUND(CROSS_THREAD_DIM)
94+ static void SimtCrossCompute(
95+ __gm__ U* x1Gm, __gm__ U* x2Gm, __gm__ U* yGm,
96+ int32_t startIdx, int32_t count, int32_t dim, int32_t dimNum, int32_t dimStride,
97+ CrossStridePara para)
98+ {
99+ const int32_t idx = threadIdx.x;
100+ const int32_t step = blockDim.x;
101+ const int32_t sx1 = para.x1s[dim];
102+ const int32_t sx2 = para.x2s[dim];
103+ const int32_t sy = para.ys[dim];
104+ 
105+ int32_t i = idx;
106+ while (i < count) {
107+ int32_t v = startIdx + i;
108+ int32_t x1b = 0, x2b = 0, yb = 0;
109+ 
110+ for (int d = 0; d < dimNum; d++) {
111+ int32_t c = v / para.ms[d];
112+ v -= c * para.ms[d];
113+ x1b += c * para.x1s[d];
114+ x2b += c * para.x2s[d];
115+ yb += c * para.ys[d];
116+ }
117+ 
118+ float a0 = (float)x1Gm[x1b];
119+ float a1 = (float)x1Gm[x1b + sx1];
120+ float a2 = (float)x1Gm[x1b + 2*sx1];
121+ float b0 = (float)x2Gm[x2b];
122+ float b1 = (float)x2Gm[x2b + sx2];
123+ float b2 = (float)x2Gm[x2b + 2*sx2];
124+ 
125+ yGm[yb] = (U)(a1*b2 - a2*b1);
126+ yGm[yb + sy] = (U)(a2*b0 - a0*b2);
127+ yGm[yb + 2*sy] = (U)(a0*b1 - a1*b0);
128+ 
129+ i += step;
130+ }
131+ }
132+ 
133+ template <typename U>
134+ __simt_vf__ LAUNCH_BOUND(CROSS_THREAD_DIM_FOR_INT64)
135+ static void SimtCrossComputeInt64(
136+ __gm__ U* x1Gm, __gm__ U* x2Gm, __gm__ U* yGm,
137+ int32_t startIdx, int32_t count, int32_t dim, int32_t dimNum, int32_t dimStride,
138+ CrossStridePara para)
139+ {
140+ const int32_t idx = threadIdx.x;
141+ const int32_t step = blockDim.x;
142+ 
143+ const int64_t sx1 = (int64_t)para.x1s[dim];
144+ const int64_t sx2 = (int64_t)para.x2s[dim];
145+ const int64_t sy = (int64_t)para.ys[dim];
146+ 
147+ int32_t i = idx;
148+ while (i < count) {
149+ int32_t v = startIdx + i;
150+ int64_t x1b = 0;
151+ int64_t x2b = 0;
152+ int64_t yb = 0;
153+ 
154+ for (int d = 0; d < dimNum; d++) {
155+ int64_t m = (int64_t)para.ms[d];
156+ int64_t x1s = (int64_t)para.x1s[d];
157+ int64_t x2s = (int64_t)para.x2s[d];
158+ int64_t ys = (int64_t)para.ys[d];
159+ 
160+ int64_t c = (int64_t)v / m;
161+ v = (int32_t)((int64_t)v - c * m);
162+ 
163+ x1b += c * x1s;
164+ x2b += c * x2s;
165+ yb += c * ys;
166+ }
167+ 
168+ float a0 = (float)x1Gm[x1b];
169+ float a1 = (float)x1Gm[x1b + sx1];
170+ float a2 = (float)x1Gm[x1b + 2 * sx1];
171+ 
172+ float b0 = (float)x2Gm[x2b];
173+ float b1 = (float)x2Gm[x2b + sx2];
174+ float b2 = (float)x2Gm[x2b + 2 * sx2];
175+ 
176+ yGm[yb] = (U)(a1 * b2 - a2 * b1);
177+ yGm[yb + sy] = (U)(a2 * b0 - a0 * b2);
178+ yGm[yb + 2 * sy] = (U)(a0 * b1 - a1 * b0);
179+ 
180+ i += step;
181+ }
182+ }
183+ 
184+ GlobalTensor<T> x1Gm_;
185+ GlobalTensor<T> x2Gm_;
186+ GlobalTensor<T> yGm_;
187+ int32_t totalVectors_{0};
188+ int32_t vectorsPerCore_{0};
189+ int32_t coreNum_{0};
190+ int32_t dim_{0};
191+ int32_t dimNum_{0};
192+ int32_t dimStride_{1};
193+ int32_t formerCore_{0};
194+ bool usedInt64_{false};
195+ CrossStridePara para_;
196+};
197+ 
198+} // namespace CrossKernel
199+#endif
Amath/cross/op_kernel/arch35/cross_struct.h+40-0
@@ -0,0 +1,40 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file cross_struct.h
13+ * \brief cross tiling data
14+ */
15+ 
16+#ifndef CROSS_STRUCT_H
17+#define CROSS_STRUCT_H
18+ 
19+#include <cstdint>
20+ 
21+constexpr int64_t MAX_DIM = 8;
22+ 
23+#pragma pack(push, 8)
24+struct CrossRegbaseTilingData {
25+ int64_t totalVectors;
26+ int64_t vectorsPerCore;
27+ int64_t coreNum;
28+ int64_t dim;
29+ int64_t dimNum;
30+ int64_t mergedStride[MAX_DIM];
31+ int64_t x1Stride[MAX_DIM];
32+ int64_t x2Stride[MAX_DIM];
33+ int64_t yStride[MAX_DIM];
34+ int64_t dimStride;
35+ int64_t formerCore;
36+ int64_t usedInt64;
37+};
38+#pragma pack(pop)
39+ 
40+#endif
Amath/cross/op_kernel/arch35/cross_tiling_key.h+21-0
@@ -0,0 +1,21 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+/*!
12+ * \file cross_tiling_key.h
13+ * \brief cross tiling key definitions
14+ */
15+ 
16+#ifndef CROSS_TILING_KEY_H
17+#define CROSS_TILING_KEY_H
18+ 
19+#define CROSS_DEFAULT 0
20+ 
21+#endif
Amath/cross/op_kernel/cross_apt.cpp+23-0
@@ -0,0 +1,23 @@
1+/**
2+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
3+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
4+ * CANN Open Software License Agreement Version 2.0 (the "License").
5+ * Please refer to the License for details. You may not use this file except in compliance with the License.
6+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
7+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
8+ * See LICENSE in the root of the software repository for the full text of the License.
9+ */
10+ 
11+#include "kernel_operator.h"
12+#include "arch35/cross.h"
13+ 
14+__global__ __aicore__ void cross(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
15+{
16+ REGISTER_TILING_DEFAULT(CrossRegbaseTilingData);
17+ GET_TILING_DATA(tilingData, tiling);
18+ if (TILING_KEY_IS(CROSS_DEFAULT)) {
19+ CrossKernel::Cross<DTYPE_X1> op;
20+ op.Init(x1, x2, y, tilingData);
21+ op.Process();
22+ }
23+}
Mscripts/ci/ascend950/ops_math_operator_list.yaml+1-0
@@ -42,6 +42,7 @@ operator_group_1:
42 - cholesky42 - cholesky
43 - cos43 - cos
44 - cosh44 - cosh
45+ - cross
45 - cumsum46 - cumsum
46 - cummin47 - cummin
47 - diag48 - diag