回退polar算子到simt实现-ops-math-AtomGit

Mmath/polar/op_host/arch35/polar_tiling.cpp+199-73

	@@ -9,87 +9,213 @@
9	*/	9	*/
10		10
11	/*!	11	/*!
12	- * \file polar_tiling.cpp	12	+ * \file polar_tiling.cpp
13	- * \brief polar broadcast tiling implementation	13	+ * \brief polar tiling
14	*/	14	*/
15	-#include <graph/utils/type_utils.h>
16	-#include "op_host/math_tiling_templates_registry.h"
17	-#include "log/log.h"
18	-#include "atvoss/broadcast/broadcast_tiling.h"
19	-#include "../../op_kernel/arch35/polar_dag.h"
20	-#include "../../op_kernel/arch35/polar_struct.h"
21	-#include "polar_tiling.h"
22		15
23	-using namespace Ops::Base;	16	+#include "polar_tiling.h"
24	-using namespace AscendC;	17	+#include <graph/utils/type_utils.h>
		18	+
25	using namespace ge;	19	using namespace ge;
26		20
27	namespace optiling {	21	namespace optiling {
		22	+static constexpr uint64_t INPUT_ABS = 0;
		23	+static constexpr uint64_t INPUT_ANGLE = 1;
		24	+static constexpr uint64_t OUTPUT_Y = 0;
28		25
29	-constexpr static uint64_t POLAR_COMMON_TILING_PRIORITY = 0;	26	+ge::graphStatus PolarTiling::GetPlatformInfo()
30	-
31	-ge::graphStatus PolarTiling::GetShapeAttrsInfo() { return ge::GRAPH_SUCCESS; }
32	-
33	-bool PolarTiling::IsCapable() { return true; }
34	-
35	-ge::graphStatus PolarTiling::DoOpTiling()
36	{	27	{
37	- auto outputDesc = context_->GetOutputDesc(0);	28	+ OP_LOGD(context_, "PolarTiling GetPlatformInfo.");
38	- OP_CHECK_NULL_WITH_CONTEXT(context_, outputDesc);	29	+ compileInfo_ = static_cast<const PolarCompileInfo*>(context_->GetCompileInfo());
39	- ge::DataType outputDtype = outputDesc->GetDataType();	30	+ OP_CHECK_NULL_WITH_CONTEXT(context_, compileInfo_);
40	-
41	- ge::graphStatus ret = ge::GRAPH_SUCCESS;
42	- if (outputDtype == ge::DT_COMPLEX64) {
43	- BroadcastBaseTiling<PolarOp::PolarBrcDag<complex64, float>::OpDag> brcBaseTiling(context_);
44	- ret = brcBaseTiling.DoTiling();
45	- tilingKey = GET_TPL_TILING_KEY(brcBaseTiling.GetSchMode());
46	- } else {
47	- OP_LOGE_FOR_INVALID_DTYPE(context_->GetNodeName(), "output", ge::TypeUtils::DataTypeToSerialString(outputDtype),
48	- "COMPLEX64");
49	- return ge::GRAPH_FAILED;
50	- }
51	-
52	- return ret;
53	-}
54	-
55	-ge::graphStatus PolarTiling::DoLibApiTiling() { return ge::GRAPH_SUCCESS; }
56	-
57	-uint64_t PolarTiling::GetTilingKey() const { return tilingKey; }
58	-
59	-ge::graphStatus PolarTiling::GetWorkspaceSize() { return ge::GRAPH_SUCCESS; }
60	-
61	-ge::graphStatus PolarTiling::PostTiling() { return ge::GRAPH_SUCCESS; }
62	-
63	-ge::graphStatus PolarTiling::GetPlatformInfo() { return ge::GRAPH_SUCCESS; }
64	-
65	-ge::graphStatus TilingForPolar(gert::TilingContext* context)
66	-{
67	- OP_LOGD("PolarTiling", "Enter TilingForPolar");
68	- if (context == nullptr) {
69	- OP_LOGE("PolarTiling", "Tiling context is nullptr");
70	- return ge::GRAPH_FAILED;
71	- }
72	-
73	- auto compileInfo = reinterpret_cast<const PolarCompileInfo*>(context->GetCompileInfo());
74	- OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
75	-
76	- OP_LOGD(context, "Enter ascendc PolarTiling");
77	- return Ops::Math::OpTiling::TilingRegistry::GetInstance().DoTilingImpl(context);
78	-}
79	-
80	-ge::graphStatus TilingPrepareForPolar(gert::TilingParseContext* context)
81	-{
82	- auto compileInfoPtr = context->GetCompiledInfo<PolarCompileInfo>();
83	- OP_CHECK_NULL_WITH_CONTEXT(context, compileInfoPtr);
84	- fe::PlatFormInfos* platformInfoPtr = context->GetPlatformInfo();
85	- OP_CHECK_NULL_WITH_CONTEXT(context, platformInfoPtr);
86	- auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfoPtr);
87	- compileInfoPtr->coreNum = ascendcPlatform.GetCoreNumAiv();
88	- ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, compileInfoPtr->ubSize);
89	return ge::GRAPH_SUCCESS;	31	return ge::GRAPH_SUCCESS;
90	}	32	}
91		33
92	-IMPL_OP_OPTILING(Polar).Tiling(TilingForPolar).TilingParse<PolarCompileInfo>(TilingPrepareForPolar);	34	+ge::graphStatus PolarTiling::CheckDtype()
		35	+{
		36	+ OP_LOGD(context_, "PolarTiling CheckDtype.");
		37	+ auto input0Desc = context_->GetInputDesc(INPUT_ABS);
		38	+ OP_CHECK_NULL_WITH_CONTEXT(context_, input0Desc);
		39	+ ge::DataType input0Dtype = input0Desc->GetDataType();
		40	+ auto input1Desc = context_->GetInputDesc(INPUT_ANGLE);
		41	+ OP_CHECK_NULL_WITH_CONTEXT(context_, input1Desc);
		42	+ ge::DataType input1Dtype = input1Desc->GetDataType();
		43	+ auto outputDesc = context_->GetOutputDesc(OUTPUT_Y);
		44	+ OP_CHECK_NULL_WITH_CONTEXT(context_, outputDesc);
		45	+ ge::DataType outputDtype = outputDesc->GetDataType();
		46	+ if (input0Dtype != ge::DT_FLOAT \|\| input1Dtype != ge::DT_FLOAT \|\| outputDtype != ge::DT_COMPLEX64) {
		47	+ std::string dtypesStr = ge::TypeUtils::DataTypeToSerialString(input0Dtype) + ", " +
		48	+ ge::TypeUtils::DataTypeToSerialString(input1Dtype) + " and " +
		49	+ ge::TypeUtils::DataTypeToSerialString(outputDtype);
		50	+ OP_LOGE_FOR_INVALID_DTYPES_WITH_REASON(
		51	+ context_->GetNodeName(), "abs, angle and y", dtypesStr.c_str(),
		52	+ "The dtypes of abs and angle must be float, and the dtype of y must be complex64");
		53	+ return ge::GRAPH_FAILED;
		54	+ }
		55	+ return ge::GRAPH_SUCCESS;
		56	+}
93		57
94	-REGISTER_OPS_TILING_TEMPLATE(Polar, PolarTiling, POLAR_COMMON_TILING_PRIORITY);	58	+ge::graphStatus PolarTiling::CheckBroadcastAndMergeShape()
95	-} // namespace optiling	59	+{
		60	+ OP_LOGD(context_, "PolarTiling CheckBroadcastAndMergeShape.");
		61	+ const gert::StorageShape* absStorageShape = context_->GetInputShape(INPUT_ABS);
		62	+ OP_CHECK_NULL_WITH_CONTEXT(context_, absStorageShape);
		63	+ const gert::StorageShape* angleStorageShape = context_->GetInputShape(INPUT_ANGLE);
		64	+ OP_CHECK_NULL_WITH_CONTEXT(context_, angleStorageShape);
		65	+
		66	+ auto absShape = absStorageShape->GetStorageShape();
		67	+ auto angleShape = angleStorageShape->GetStorageShape();
		68	+
		69	+ int64_t absDimNum = static_cast<int64_t>(absShape.GetDimNum());
		70	+ int64_t angleDimNum = static_cast<int64_t>(angleShape.GetDimNum());
		71	+ dimNum_ = std::max(absDimNum, angleDimNum);
		72	+ OP_CHECK_IF(dimNum_ > POLAR_MAX_DIM,
		73	+ OP_LOGE(context_, "dimNum %ld exceeds POLAR_MAX_DIM %ld", dimNum_, POLAR_MAX_DIM),
		74	+ return ge::GRAPH_FAILED);
		75	+
		76	+ for (int64_t i = 0; i < dimNum_; i++) {
		77	+ int64_t absOffset = i - (dimNum_ - absDimNum);
		78	+ int64_t angleOffset = i - (dimNum_ - angleDimNum);
		79	+ int64_t absDim = (absOffset >= 0) ? absShape.GetDim(absOffset) : 1;
		80	+ int64_t angleDim = (angleOffset >= 0) ? angleShape.GetDim(angleOffset) : 1;
		81	+ absDims_[i] = absDim;
		82	+ angleDims_[i] = angleDim;
		83	+ OP_CHECK_IF(absDim != angleDim && absDim != 1 && angleDim != 1,
		84	+ OP_LOGE(context_, "Shapes not broadcastable at dim %ld: %ld vs %ld", i, absDim, angleDim),
		85	+ return ge::GRAPH_FAILED);
		86	+ mergedShape_[i] = std::max(absDim, angleDim);
		87	+ }
		88	+
		89	+ totalElements_ = 1;
		90	+ for (int64_t i = 0; i < dimNum_; i++) {
		91	+ totalElements_ *= mergedShape_[i];
		92	+ }
		93	+ return ge::GRAPH_SUCCESS;
		94	+}
		95	+
		96	+ge::graphStatus PolarTiling::CalcStride()
		97	+{
		98	+ OP_LOGD(context_, "PolarTiling CalcStride.");
		99	+ int64_t strideAbs = 1;
		100	+ int64_t strideAngle = 1;
		101	+ int64_t strideMerged = 1;
		102	+ int64_t strideY = 1;
		103	+ for (int64_t i = dimNum_ - 1; i >= 0; i--) {
		104	+ absStride_[i] = (absDims_[i] == 1) ? 0 : strideAbs;
		105	+ angleStride_[i] = (angleDims_[i] == 1) ? 0 : strideAngle;
		106	+ mergedStride_[i] = strideMerged;
		107	+ yStride_[i] = strideY;
		108	+ strideAbs *= absDims_[i];
		109	+ strideAngle *= angleDims_[i];
		110	+ strideMerged *= mergedShape_[i];
		111	+ strideY *= mergedShape_[i];
		112	+ }
		113	+ return ge::GRAPH_SUCCESS;
		114	+}
		115	+
		116	+ge::graphStatus PolarTiling::GetShapeAttrsInfo()
		117	+{
		118	+ if (CheckDtype() != ge::GRAPH_SUCCESS)
		119	+ return ge::GRAPH_FAILED;
		120	+ if (CheckBroadcastAndMergeShape() != ge::GRAPH_SUCCESS)
		121	+ return ge::GRAPH_FAILED;
		122	+ if (CalcStride() != ge::GRAPH_SUCCESS)
		123	+ return ge::GRAPH_FAILED;
		124	+ return ge::GRAPH_SUCCESS;
		125	+}
		126	+
		127	+ge::graphStatus PolarTiling::DoOpTiling()
		128	+{
		129	+ OP_LOGD(context_, "PolarTiling DoOpTiling.");
		130	+
		131	+ int64_t coreNum = compileInfo_->coreNum;
		132	+ int64_t elementsPerCore = totalElements_ / coreNum;
		133	+ int64_t formerCore = totalElements_ % coreNum;
		134	+
		135	+ tilingData_.totalElements = totalElements_;
		136	+ tilingData_.elementsPerCore = elementsPerCore;
		137	+ tilingData_.coreNum = coreNum;
		138	+ tilingData_.formerCore = formerCore;
		139	+ tilingData_.dimNum = dimNum_;
		140	+
		141	+ for (int64_t i = 0; i < POLAR_MAX_DIM; i++) {
		142	+ if (i < dimNum_) {
		143	+ tilingData_.mergedStride[i] = mergedStride_[i];
		144	+ tilingData_.absStride[i] = absStride_[i];
		145	+ tilingData_.angleStride[i] = angleStride_[i];
		146	+ tilingData_.yStride[i] = yStride_[i];
		147	+ } else {
		148	+ tilingData_.mergedStride[i] = 1;
		149	+ tilingData_.absStride[i] = 0;
		150	+ tilingData_.angleStride[i] = 0;
		151	+ tilingData_.yStride[i] = 0;
		152	+ }
		153	+ }
		154	+
		155	+ blockDim_ = (totalElements_ < coreNum) ? totalElements_ : coreNum;
		156	+
		157	+ return ge::GRAPH_SUCCESS;
		158	+}
		159	+
		160	+ge::graphStatus PolarTiling::PostTiling()
		161	+{
		162	+ OP_LOGD(context_, "PolarTiling PostTiling.");
		163	+
		164	+ auto workspaces = context_->GetWorkspaceSizes(1);
		165	+ OP_CHECK_NULL_WITH_CONTEXT(context_, workspaces);
		166	+ workspaces[0] = 0;
		167	+
		168	+ auto res = context_->SetBlockDim(static_cast<uint32_t>(blockDim_));
		169	+ OP_CHECK_IF((res != ge::GRAPH_SUCCESS), OP_LOGE(context_, "SetBlockDim failed."), return ge::GRAPH_FAILED);
		170	+
		171	+ errno_t ret = memcpy_s(context_->GetRawTilingData()->GetData(), context_->GetRawTilingData()->GetCapacity(),
		172	+ &tilingData_, sizeof(PolarTilingData));
		173	+ if (ret != EOK) {
		174	+ OP_LOGE(context_->GetNodeName(), "memcpy_s failed, ret=%d", ret);
		175	+ return ge::GRAPH_FAILED;
		176	+ }
		177	+ context_->GetRawTilingData()->SetDataSize(sizeof(PolarTilingData));
		178	+
		179	+ return ge::GRAPH_SUCCESS;
		180	+}
		181	+
		182	+static ge::graphStatus Tiling4Polar(gert::TilingContext* context)
		183	+{
		184	+ OP_LOGD(context, "Tiling4Polar start.");
		185	+
		186	+ PolarTiling polarTiling(context);
		187	+ auto ret = polarTiling.DoTiling();
		188	+ OP_CHECK_IF((ret == ge::GRAPH_FAILED), OP_LOGD(context, "Tiling4Polar failed!"), return ge::GRAPH_FAILED);
		189	+ OP_LOGD(context, "Tiling4Polar end.");
		190	+ return ge::GRAPH_SUCCESS;
		191	+}
		192	+
		193	+static ge::graphStatus TilingPrepare4PolarAscendc(gert::TilingParseContext* context)
		194	+{
		195	+ OP_LOGD(context->GetNodeName(), "Enter TilingPrepare4PolarAscendc.");
		196	+
		197	+ auto compileInfo = context->GetCompiledInfo<PolarCompileInfo>();
		198	+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
		199	+ auto platformInfo = context->GetPlatformInfo();
		200	+ OP_CHECK_NULL_WITH_CONTEXT(context, platformInfo);
		201	+ auto ascendcPlatform = platform_ascendc::PlatformAscendC(platformInfo);
		202	+
		203	+ compileInfo->coreNum = ascendcPlatform.GetCoreNumAiv();
		204	+ OP_CHECK_IF((compileInfo->coreNum <= 0), OP_LOGE(context->GetNodeName(), "core num is negative."),
		205	+ return ge::GRAPH_FAILED);
		206	+
		207	+ OP_LOGD(context->GetNodeName(), "Exit TilingPrepare4PolarAscendc.");
		208	+ return ge::GRAPH_SUCCESS;
		209	+}
		210	+
		211	+static ge::graphStatus TilingPrepare4Polar(gert::TilingParseContext* context)
		212	+{
		213	+ auto compileInfo = context->GetCompiledInfo<PolarCompileInfo>();
		214	+ OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
		215	+ OP_LOGD("TilingPrepare4Polar", "Ascend C TilingPrepare4Polar success.");
		216	+ return TilingPrepare4PolarAscendc(context);
		217	+}
		218	+
		219	+IMPL_OP_OPTILING(Polar).Tiling(Tiling4Polar).TilingParse<PolarCompileInfo>(TilingPrepare4Polar);
		220	+
		221	+} // namespace optiling

Mmath/polar/op_host/arch35/polar_tiling.h+29-12

	@@ -10,39 +10,56 @@
10		10
11	/*!	11	/*!
12	* \file polar_tiling.h	12	* \file polar_tiling.h
13	- * \brief polar broadcast tiling header	13	+ * \brief polar tiling header
14	*/	14	*/
15	#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_POLAR_TILING_H	15	#ifndef OPS_BUILD_IN_OP_TILING_RUNTIME_POLAR_TILING_H
16	#define OPS_BUILD_IN_OP_TILING_RUNTIME_POLAR_TILING_H	16	#define OPS_BUILD_IN_OP_TILING_RUNTIME_POLAR_TILING_H
17		17
18	-#include "register/op_def_registry.h"	18	+#include <cstdint>
19	-#include "tiling/tiling_api.h"	19	+#include "register/tilingdata_base.h"
20	#include "op_host/tiling_base_class.h"	20	#include "op_host/tiling_base_class.h"
		21	+#include "register/op_impl_registry.h"
		22	+#include "platform/platform_ascendc.h"
		23	+#include "log/log.h"
		24	+#include "../../op_kernel/arch35/polar_struct.h"
21		25
22	namespace optiling {	26	namespace optiling {
23		27
24	struct PolarCompileInfo {	28	struct PolarCompileInfo {
25	- uint64_t coreNum = 0;	29	+ int64_t coreNum = 0;
26	- uint64_t ubSize = 0;
27	};	30	};
28		31
29	class PolarTiling : public Ops::Base::TilingBaseClass {	32	class PolarTiling : public Ops::Base::TilingBaseClass {
30	public:	33	public:
31	- explicit PolarTiling(gert::TilingContext* context) : Ops::Base::TilingBaseClass(context) {}	34	+ explicit PolarTiling(gert::TilingContext* context) : TilingBaseClass(context) {}
32		35
33	protected:	36	protected:
34	- bool IsCapable() override;	37	+ bool IsCapable() override { return true; }
35	ge::graphStatus GetPlatformInfo() override;	38	ge::graphStatus GetPlatformInfo() override;
36	ge::graphStatus GetShapeAttrsInfo() override;	39	ge::graphStatus GetShapeAttrsInfo() override;
37	ge::graphStatus DoOpTiling() override;	40	ge::graphStatus DoOpTiling() override;
38	- ge::graphStatus DoLibApiTiling() override;	41	+ ge::graphStatus DoLibApiTiling() override { return ge::GRAPH_SUCCESS; }
39	- uint64_t GetTilingKey() const override;	42	+ uint64_t GetTilingKey() const override { return 0; }
40	- ge::graphStatus GetWorkspaceSize() override;	43	+ ge::graphStatus GetWorkspaceSize() override { return ge::GRAPH_SUCCESS; }
41	ge::graphStatus PostTiling() override;	44	ge::graphStatus PostTiling() override;
		45	+ ge::graphStatus CheckDtype();
		46	+ ge::graphStatus CheckBroadcastAndMergeShape();
		47	+ ge::graphStatus CalcStride();
42		48
43	private:	49	private:
44	- uint64_t tilingKey = 0;	50	+ const PolarCompileInfo* compileInfo_;
		51	+ PolarTilingData tilingData_{};
		52	+ uint32_t blockDim_{1};
		53	+ int64_t totalElements_ = 0;
		54	+ int64_t dimNum_ = 0;
		55	+ int64_t absDims_[POLAR_MAX_DIM] = {0};
		56	+ int64_t angleDims_[POLAR_MAX_DIM] = {0};
		57	+ int64_t mergedShape_[POLAR_MAX_DIM] = {0};
		58	+ int64_t mergedStride_[POLAR_MAX_DIM] = {0};
		59	+ int64_t absStride_[POLAR_MAX_DIM] = {0};
		60	+ int64_t angleStride_[POLAR_MAX_DIM] = {0};
		61	+ int64_t yStride_[POLAR_MAX_DIM] = {0};
45	};	62	};
46		63
47	} // namespace optiling	64	} // namespace optiling
48	-#endif // OPS_BUILD_IN_OP_TILING_RUNTIME_POLAR_TILING_H	65	+#endif // OPS_BUILD_IN_OP_TILING_RUNTIME_POLAR_TILING_H

Dmath/polar/op_kernel/arch35/polar.h+0-33

	@@ -1,33 +0,0 @@
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 polar_simd.h
13	- * \brief polar operator SIMD kernel entry for arch35 (ascend950)
14	- */
15	-#ifndef POLAR_SIMD_H_
16	-#define POLAR_SIMD_H_
17	-
18	-#include "kernel_operator.h"
19	-#include "atvoss/broadcast/broadcast_sch.h"
20	-#include "polar_dag.h"
21	-#include "polar_struct.h"
22	-
23	-using namespace Ops::Base;
24	-
25	-template <uint64_t schMode>
26	-__global__ __aicore__ void polar(GM_ADDR abs, GM_ADDR angle, GM_ADDR out, GM_ADDR workspace, GM_ADDR tiling)
27	-{
28	- using OpDag = PolarOp::PolarBrcDag<complex64, float>::OpDag;
29	- BroadcastSch<schMode, OpDag> sch(tiling);
30	- sch.Process(abs, angle, out);
31	-}
32	-
33	-#endif // POLAR_SIMD_H_

Dmath/polar/op_kernel/arch35/polar_dag.h+0-105

	@@ -1,105 +0,0 @@
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 polar_dag.h
13	- * \brief polar operator BRC DAG definition for arch35 (ascend950)
14	- *
15	- * Polar operator: abs(float) and angle(float) -> out(complex64)
16	- * out.real = abs * cos(angle)
17	- * out.imag = abs * sin(angle)
18	- * Supports broadcast between abs and angle inputs via CopyInBrc.
19	- * Uses AscendC SIMD vector operations (Cos, Sin, Mul) and Interleave
20	- * with double buffer pipeline.
21	- */
22	-#ifndef POLAR_DAG_H_
23	-#define POLAR_DAG_H_
24	-
25	-#include "atvoss/util/dag.h"
26	-#include "atvoss/util/vec.h"
27	-#include "atvoss/util/placeholder.h"
28	-#include "atvoss/util/elems.h"
29	-
30	-#ifndef __CCE_AICORE__
31	-struct complex64 {
32	- float real;
33	- float imag;
34	-};
35	-#else
36	-namespace PolarOp {
37	-using namespace AscendC;
38	-constexpr uint32_t POLAR_TAIL_THREAD_NUM = 1024;
39	-
40	-template <typename T>
41	-__simt_vf__ __aicore__ inline void PolarTailCompute(__ubuf__ T* dst, __ubuf__ T* abs, __ubuf__ T* angle,
42	- uint64_t alignedCount, uint64_t count)
43	-{
44	- for (uint64_t i = alignedCount + threadIdx.x; i < count; i += blockDim.x) {
45	- T cosVal = Simt::Cos(angle[i]);
46	- T sinVal = Simt::Sin(angle[i]);
47	- dst[i * 2] = abs[i] * cosVal;
48	- dst[i * 2 + 1] = abs[i] * sinVal;
49	- }
50	-}
51	-} // namespace PolarOp
52	-#endif
53	-
54	-namespace PolarOp {
55	-using namespace Ops::Base;
56	-
57	-template <class C, class T>
58	-struct PolarMerge : public Vec::ElemwiseBinaryOP<C, T, T> {
59	- __aicore__ inline PolarMerge(LocalTensor<C>& dst, LocalTensor<T>& abs, LocalTensor<T>& angle, uint64_t count)
60	- {
61	-#ifdef __CCE_AICORE__
62	- using namespace AscendC;
63	- constexpr uint32_t ALIGN_ELEMS = 32 / sizeof(T); // 32B的元素数量
64	- LocalTensor<T> dstT = dst.template ReinterpretCast<T>();
65	-
66	- uint64_t alignedCount = (count / ALIGN_ELEMS) * ALIGN_ELEMS;
67	- if (alignedCount > 0) {
68	- LocalTensor<T> tmpReal = dstT;
69	- LocalTensor<T> tmpImag = dstT[alignedCount];
70	-
71	- AscendC::Cos(tmpReal, angle, alignedCount); // tmpReal = cos(angle)
72	- AscendC::Sin(tmpImag, angle, alignedCount); // tmpImag = sin(angle)
73	- AscendC::Mul(tmpReal, abs, tmpReal, alignedCount); // tmpReal = abs * cos(angle)
74	- AscendC::Mul(tmpImag, abs, tmpImag, alignedCount); // tmpImag = abs * sin(angle)
75	-
76	- Interleave(dstT, dstT[alignedCount], tmpReal, tmpImag, alignedCount);
77	- }
78	-
79	- // 非32B对齐的尾部数据（<32B，VEC无法处理，用Simt::VF_CALL计算）
80	- if (count > alignedCount) {
81	- __ubuf__ T* dstAddr = (__ubuf__ T*)dstT.GetPhyAddr();
82	- __ubuf__ T* absAddr = (__ubuf__ T*)abs.GetPhyAddr();
83	- __ubuf__ T* angleAddr = (__ubuf__ T*)angle.GetPhyAddr();
84	- Simt::VF_CALL<PolarTailCompute<T>>(Simt::Dim3(POLAR_TAIL_THREAD_NUM), dstAddr, absAddr, angleAddr,
85	- alignedCount, count);
86	- }
87	-
88	-#endif
89	- }
90	-};
91	-
92	-template <typename C, typename T>
93	-struct PolarBrcDag {
94	- using OpCopyInAbs = Bind<Vec::CopyInBrc<T>, Placeholder::In0<T>>;
95	- using OpCopyInAngle = Bind<Vec::CopyInBrc<T>, Placeholder::In1<T>>;
96	- using OpMerge = Bind<PolarMerge<C, T>, OpCopyInAbs, OpCopyInAngle>;
97	- using OpCopyOut = Bind<Vec::CopyOut<C>, Placeholder::Out0<C>, OpMerge>;
98	-
99	- using Outputs = Elems<OpCopyOut>;
100	- using MemCfg = MemOptCfg<MemLevel::LEVEL_2>;
101	- using OpDag = DAGSch<Outputs, void, MemCfg>;
102	-};
103	-
104	-} // namespace PolarOp
105	-#endif // POLAR_DAG_H_

Amath/polar/op_kernel/arch35/polar_simt.h+122-0

@@ -0,0 +1,122 @@
	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 polar_simt.h
	13	+ * \brief polar simt kernel
	14	+ */
	15	+#ifndef POLAR_SIMT_H
	16	+#define POLAR_SIMT_H
	17	+
	18	+#include "kernel_operator.h"
	19	+#include "polar_struct.h"
	20	+#ifdef __CCE_AICORE__
	21	+#include "simt_api/asc_simt.h"
	22	+#endif
	23	+
	24	+namespace PolarOp {
	25	+using namespace AscendC;
	26	+
	27	+constexpr int32_t POLAR_THREAD_DIM = 1024;
	28	+
	29	+struct PolarStridePara {
	30	+ int64_t ms[POLAR_MAX_DIM];
	31	+ int64_t absS[POLAR_MAX_DIM];
	32	+ int64_t angleS[POLAR_MAX_DIM];
	33	+ int64_t yS[POLAR_MAX_DIM];
	34	+};
	35	+
	36	+template <typename T>
	37	+class PolarSimt {
	38	+public:
	39	+ __aicore__ inline PolarSimt() {}
	40	+ __aicore__ inline ~PolarSimt() {}
	41	+
	42	+ __aicore__ inline void Init(GM_ADDR abs, GM_ADDR angle, GM_ADDR y, const PolarTilingData& tilingData)
	43	+ {
	44	+ totalElements_ = tilingData.totalElements;
	45	+ elementsPerCore_ = tilingData.elementsPerCore;
	46	+ coreNum_ = tilingData.coreNum;
	47	+ formerCore_ = tilingData.formerCore;
	48	+ dimNum_ = tilingData.dimNum;
	49	+
	50	+ for (int i = 0; i < POLAR_MAX_DIM; i++) {
	51	+ para_.ms[i] = tilingData.mergedStride[i];
	52	+ para_.absS[i] = tilingData.absStride[i];
	53	+ para_.angleS[i] = tilingData.angleStride[i];
	54	+ para_.yS[i] = tilingData.yStride[i];
	55	+ }
	56	+
	57	+ absGm_.SetGlobalBuffer((__gm__ T*)abs);
	58	+ angleGm_.SetGlobalBuffer((__gm__ T*)angle);
	59	+ yGm_.SetGlobalBuffer((__gm__ T*)y);
	60	+ }
	61	+
	62	+ __aicore__ inline void Process()
	63	+ {
	64	+ int32_t blockIdx = static_cast<int32_t>(GetBlockIdx());
	65	+ int64_t startIdx = (blockIdx < formerCore_) ? (elementsPerCore_ + 1) * blockIdx :
	66	+ formerCore_ + elementsPerCore_ * blockIdx;
	67	+ int64_t count = (blockIdx < formerCore_) ? (elementsPerCore_ + 1) : elementsPerCore_;
	68	+
	69	+ if (count <= 0)
	70	+ return;
	71	+
	72	+ asc_vf_call<SimtPolarCompute<T>>(dim3(POLAR_THREAD_DIM), (__gm__ T*)(absGm_.GetPhyAddr()),
	73	+ (__gm__ T)(angleGm_.GetPhyAddr()), (__gm__ T)(yGm_.GetPhyAddr()), startIdx,
	74	+ count, dimNum_, para_);
	75	+ }
	76	+
	77	+private:
	78	+ template <typename U>
	79	+ __simt_vf__ LAUNCH_BOUND(POLAR_THREAD_DIM) static void SimtPolarCompute(__gm__ U* absGm, __gm__ U* angleGm,
	80	+ __gm__ U* yGm, int64_t startIdx,
	81	+ int64_t count, int64_t dimNum,
	82	+ PolarStridePara para)
	83	+ {
	84	+ const int64_t idx = threadIdx.x;
	85	+ const int64_t step = blockDim.x;
	86	+ int64_t i = idx;
	87	+ while (i < count) {
	88	+ int64_t v = startIdx + i;
	89	+ int64_t absBase = 0;
	90	+ int64_t angleBase = 0;
	91	+ int64_t yBase = 0;
	92	+ int64_t vv = v;
	93	+ for (int64_t d = 0; d < dimNum; d++) {
	94	+ int64_t c = vv / para.ms[d];
	95	+ vv -= c * para.ms[d];
	96	+ absBase += c * para.absS[d];
	97	+ angleBase += c * para.angleS[d];
	98	+ yBase += c * para.yS[d];
	99	+ }
	100	+ U absVal = absGm[absBase];
	101	+ U cosVal = Simt::Cos(angleGm[angleBase]);
	102	+ U sinVal = Simt::Sin(angleGm[angleBase]);
	103	+ yGm[2 * yBase] = absVal * cosVal;
	104	+ yGm[2 * yBase + 1] = absVal * sinVal;
	105	+ i += step;
	106	+ }
	107	+ }
	108	+
	109	+ GlobalTensor<T> absGm_;
	110	+ GlobalTensor<T> angleGm_;
	111	+ GlobalTensor<T> yGm_;
	112	+ int64_t totalElements_{0};
	113	+ int64_t elementsPerCore_{0};
	114	+ int64_t coreNum_{0};
	115	+ int64_t formerCore_{0};
	116	+ int64_t dimNum_{0};
	117	+ PolarStridePara para_;
	118	+};
	119	+
	120	+} // namespace PolarOp
	121	+
	122	+#endif // POLAR_SIMT_H

Mmath/polar/op_kernel/arch35/polar_struct.h+16-5

	@@ -15,11 +15,22 @@
15	#ifndef POLAR_STRUCT_H_	15	#ifndef POLAR_STRUCT_H_
16	#define POLAR_STRUCT_H_	16	#define POLAR_STRUCT_H_
17		17
18	-#include "atvoss/broadcast/broadcast_base_struct.h"	18	+#include <cstdint>
19		19
20	-using namespace Ops::Base;	20	+constexpr int64_t POLAR_MAX_DIM = 8;
21		21
22	-ASCENDC_TPL_ARGS_DECL(Polar, BRC_TEMP_SCH_MODE_KEY_DECL(schMode));	22	+#pragma pack(push, 8)
23	-ASCENDC_TPL_SEL(ASCENDC_TPL_ARGS_SEL(BRC_TEMP_SCH_MODE_KEY_SEL(schMode)));	23	+struct PolarTilingData {
		24	+ int64_t totalElements;
		25	+ int64_t elementsPerCore;
		26	+ int64_t coreNum;
		27	+ int64_t formerCore;
		28	+ int64_t dimNum;
		29	+ int64_t mergedStride[POLAR_MAX_DIM];
		30	+ int64_t absStride[POLAR_MAX_DIM];
		31	+ int64_t angleStride[POLAR_MAX_DIM];
		32	+ int64_t yStride[POLAR_MAX_DIM];
		33	+};
		34	+#pragma pack(pop)
24		35
25	-#endif // POLAR_STRUCT_H_	36	+#endif // POLAR_STRUCT_H_

Mmath/polar/op_kernel/polar_apt.cpp+14-2

	@@ -10,8 +10,20 @@
10		10
11	/* !	11	/* !
12	* \file polar_apt.cpp	12	* \file polar_apt.cpp
13	- * \brief polar kernel entry - dispatch to arch35 SIMD implementation	13	+ * \brief polar kernel
14	*/	14	*/
15		15
16	#include "kernel_operator.h"	16	#include "kernel_operator.h"
17	-#include "arch35/polar.h"	17	+#include "arch35/polar_struct.h"
		18	+#include "arch35/polar_simt.h"
		19	+
		20	+using namespace AscendC;
		21	+
		22	+__global__ __aicore__ void polar(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
		23	+{
		24	+ REGISTER_TILING_DEFAULT(PolarTilingData);
		25	+ GET_TILING_DATA(tilingData, tiling);
		26	+ PolarOp::PolarSimt<float> polarOp;
		27	+ polarOp.Init(x1, x2, y, tilingData);
		28	+ polarOp.Process();
		29	+}

Mmath/polar/tests/ut/op_host/arch35/test_polar_tiling_arch35.cpp+41-7

	@@ -35,8 +35,8 @@ TEST_F(PolarTilingTest, polar_test_fp32_same_shape)
35	{{{16, 7, 14}, {16, 7, 14}}, ge::DT_COMPLEX64, ge::FORMAT_ND},	35	{{{16, 7, 14}, {16, 7, 14}}, ge::DT_COMPLEX64, ge::FORMAT_ND},
36	},	36	},
37	&compileInfo);	37	&compileInfo);
38	- uint64_t expectTilingKey = 8;	38	+ uint64_t expectTilingKey = 0;
39	- std::vector<size_t> expectWorkspaces = {16777216};	39	+ std::vector<size_t> expectWorkspaces = {0};
40	ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);	40	ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
41	}	41	}
42		42
	@@ -53,7 +53,7 @@ TEST_F(PolarTilingTest, polar_test_fp32_broadcast)
53	},	53	},
54	&compileInfo);	54	&compileInfo);
55	uint64_t expectTilingKey = 0;	55	uint64_t expectTilingKey = 0;
56	- std::vector<size_t> expectWorkspaces = {16777216};	56	+ std::vector<size_t> expectWorkspaces = {0};
57	ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);	57	ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
58	}	58	}
59		59
	@@ -71,7 +71,7 @@ TEST_F(PolarTilingTest, polar_test_fp32_broadcast_multidim)
71	},	71	},
72	&compileInfo);	72	&compileInfo);
73	uint64_t expectTilingKey = 0;	73	uint64_t expectTilingKey = 0;
74	- std::vector<size_t> expectWorkspaces = {16777216};	74	+ std::vector<size_t> expectWorkspaces = {0};
75	ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);	75	ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
76	}	76	}
77		77
	@@ -88,8 +88,8 @@ TEST_F(PolarTilingTest, polar_test_fp32_scalar_broadcast)
88	{{{16, 1, 4, 4, 8}, {16, 1, 4, 4, 8}}, ge::DT_COMPLEX64, ge::FORMAT_ND},	88	{{{16, 1, 4, 4, 8}, {16, 1, 4, 4, 8}}, ge::DT_COMPLEX64, ge::FORMAT_ND},
89	},	89	},
90	&compileInfo);	90	&compileInfo);
91	- uint64_t expectTilingKey = 8;	91	+ uint64_t expectTilingKey = 0;
92	- std::vector<size_t> expectWorkspaces = {16777216};	92	+ std::vector<size_t> expectWorkspaces = {0};
93	ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);	93	ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
94	}	94	}
95		95
	@@ -106,10 +106,44 @@ TEST_F(PolarTilingTest, polar_test_fp32_broadcast_diff_shape)
106	},	106	},
107	&compileInfo);	107	&compileInfo);
108	uint64_t expectTilingKey = 0;	108	uint64_t expectTilingKey = 0;
109	- std::vector<size_t> expectWorkspaces = {16777216};	109	+ std::vector<size_t> expectWorkspaces = {0};
110	ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);	110	ExecuteTestCase(tilingContextPara, ge::GRAPH_SUCCESS, expectTilingKey, expectWorkspaces);
111	}	111	}
112		112
		113	+TEST_F(PolarTilingTest, polar_test_failed_dtype_mismatch_input)
		114	+{
		115	+ optiling::PolarCompileInfo compileInfo = {64};
		116	+ gert::TilingContextPara tilingContextPara("Polar",
		117	+ {
		118	+ {{{16, 7, 14}, {16, 7, 14}}, ge::DT_BF16, ge::FORMAT_ND},
		119	+ {{{16, 7, 14}, {16, 7, 14}}, ge::DT_BF16, ge::FORMAT_ND},
		120	+ },
		121	+ {
		122	+ {{{16, 7, 14}, {16, 7, 14}}, ge::DT_COMPLEX64, ge::FORMAT_ND},
		123	+ },
		124	+ &compileInfo);
		125	+ uint64_t expectTilingKey = 0;
		126	+ std::vector<size_t> expectWorkspaces = {0};
		127	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectWorkspaces);
		128	+}
		129	+
		130	+TEST_F(PolarTilingTest, polar_test_failed_dtype_mismatch_output)
		131	+{
		132	+ optiling::PolarCompileInfo compileInfo = {64};
		133	+ gert::TilingContextPara tilingContextPara("Polar",
		134	+ {
		135	+ {{{16, 7, 14}, {16, 7, 14}}, ge::DT_FLOAT, ge::FORMAT_ND},
		136	+ {{{16, 7, 14}, {16, 7, 14}}, ge::DT_FLOAT, ge::FORMAT_ND},
		137	+ },
		138	+ {
		139	+ {{{16, 7, 14}, {16, 7, 14}}, ge::DT_FLOAT, ge::FORMAT_ND},
		140	+ },
		141	+ &compileInfo);
		142	+ uint64_t expectTilingKey = 0;
		143	+ std::vector<size_t> expectWorkspaces = {0};
		144	+ ExecuteTestCase(tilingContextPara, ge::GRAPH_FAILED, expectTilingKey, expectWorkspaces);
		145	+}
		146	+
113	TEST_F(PolarTilingTest, polar_test_failed_not_broadcastable)	147	TEST_F(PolarTilingTest, polar_test_failed_not_broadcastable)
114	{	148	{
115	optiling::PolarCompileInfo compileInfo = {64};	149	optiling::PolarCompileInfo compileInfo = {64};