[需求]A5 支持TPOW/TPOWS普通精度&高精度版本-pto-isa-AtomGit

Minclude/pto/common/event.hpp+4-0

	@@ -50,6 +50,8 @@ enum class Op : uint16_t
50	TPARTMUL,	50	TPARTMUL,
51	TPARTMAX,	51	TPARTMAX,
52	TPARTMIN,	52	TPARTMIN,
		53	+ TPOW,
		54	+ TPOWS,
53	TCMPS,	55	TCMPS,
54	TMRGSORT,	56	TMRGSORT,
55	TSORT32,	57	TSORT32,
	@@ -143,6 +145,8 @@ constexpr pipe_t opPipeList[] = {
143	PIPE_V /* TPARTMUL */,	145	PIPE_V /* TPARTMUL */,
144	PIPE_V /* TPARTMAX */,	146	PIPE_V /* TPARTMAX */,
145	PIPE_V /* TPARTMIN */,	147	PIPE_V /* TPARTMIN */,
		148	+ PIPE_V /* TPOW */,
		149	+ PIPE_V /* TPOWS */,
146	PIPE_V /* TCMPS */,	150	PIPE_V /* TCMPS */,
147	PIPE_V /* TMRGSORT */,	151	PIPE_V /* TMRGSORT */,
148	PIPE_V /* TSORT32 */,	152	PIPE_V /* TSORT32 */,

Minclude/pto/common/pto_instr.hpp+19-0

	@@ -1451,6 +1451,25 @@ PTO_INST RecordEvent TEXP(TileDataDst &dst, TileDataSrc &src, WaitEvents &... ev
1451	return {};	1451	return {};
1452	}	1452	}
1453		1453
		1454	+template <auto PrecisionType = PowAlgorithm::DEFAULT, typename DstTile, typename BaseTile, typename ExpTile,
		1455	+ typename TmpTile, typename... WaitEvents>
		1456	+PTO_INTERNAL RecordEvent TPOW(DstTile &dst, BaseTile &base, ExpTile &exp, TmpTile &tmp, WaitEvents &... events)
		1457	+{
		1458	+ TSYNC(events...);
		1459	+ TPOW_IMPL<PrecisionType>(dst, base, exp, tmp);
		1460	+ return {};
		1461	+}
		1462	+
		1463	+template <auto PrecisionType = PowAlgorithm::DEFAULT, typename DstTile, typename BaseTile, typename TmpTile,
		1464	+ typename... WaitEvents>
		1465	+PTO_INTERNAL RecordEvent TPOWS(DstTile &dst, BaseTile &base, typename DstTile::DType exp, TmpTile &tmp,
		1466	+ WaitEvents &... events)
		1467	+{
		1468	+ TSYNC(events...);
		1469	+ TPOWS_IMPL<PrecisionType>(dst, base, exp, tmp);
		1470	+ return {};
		1471	+}
		1472	+
1454	template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents>	1473	template <typename TileDataDst, typename TileDataSrc, typename... WaitEvents>
1455	PTO_INST RecordEvent TNOT(TileDataDst &dst, TileDataSrc &src, WaitEvents &... events)	1474	PTO_INST RecordEvent TNOT(TileDataDst &dst, TileDataSrc &src, WaitEvents &... events)
1456	{	1475	{

Minclude/pto/common/pto_instr_impl.hpp+2-0

Minclude/pto/common/type.hpp+6-0

	@@ -333,6 +333,12 @@ using FloatUnion = FloatIntUnion<float>;
333	using HalfUnion = FloatIntUnion<half>;	333	using HalfUnion = FloatIntUnion<half>;
334	#endif	334	#endif
335		335
		336	+enum class PowAlgorithm : uint8_t
		337	+{
		338	+ DEFAULT,
		339	+ HIGH_PRECISION
		340	+};
		341	+
336	enum class DivAlgorithm : uint8_t	342	enum class DivAlgorithm : uint8_t
337	{	343	{
338	DEFAULT,	344	DEFAULT,

Minclude/pto/common/utils.hpp+5-0

	@@ -11,10 +11,15 @@ See LICENSE in the root of the software repository for the full text of the Lice
11	#ifndef __UIILS_HPP__	11	#ifndef __UIILS_HPP__
12	#define __UIILS_HPP__	12	#define __UIILS_HPP__
13		13
		14	+#include <type_traits>
14	#include <pto/common/constants.hpp>	15	#include <pto/common/constants.hpp>
15	#pragma once	16	#pragma once
16		17
17	namespace pto {	18	namespace pto {
		19	+template <typename T, typename... Types>
		20	+using isSupportTypeImpl = std::disjunction<std::is_same<T, Types>...>;
		21	+template <typename T, typename... Types>
		22	+inline constexpr bool isSupportType = isSupportTypeImpl<T, Types...>::value;
18	template <typename T>	23	template <typename T>
19	struct LoadTypeBySize {	24	struct LoadTypeBySize {
20	using type = std::conditional_t<sizeof(T) == sizeof(uint8_t), uint8_t,	25	using type = std::conditional_t<sizeof(T) == sizeof(uint8_t), uint8_t,

Ainclude/pto/npu/a2a3/TPow.hpp+25-0

@@ -0,0 +1,25 @@
	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 TPOW_HPP
	12	+#define TPOW_HPP
	13	+
	14	+namespace pto {
	15	+template <PowAlgorithm algo, typename DstTile, typename BaseTile, typename ExpTile, typename TmpTile>
	16	+PTO_INTERNAL void TPOW_IMPL(DstTile &dst, BaseTile &base, ExpTile &exp, TmpTile &tmp)
	17	+{}
	18	+
	19	+template <PowAlgorithm algo, typename DstTile, typename BaseTile, typename TmpTile>
	20	+PTO_INTERNAL void TPOWS_IMPL(DstTile &dst, BaseTile &base, typename DstTile::DType exp, TmpTile &tmp)
	21	+{}
	22	+
	23	+} // namespace pto
	24	+
	25	+#endif

Minclude/pto/npu/a2a3/TPrint.hpp+7-1

	@@ -55,12 +55,18 @@ PTO_INTERNAL void PrintValue(T &val, int col)
55	} else if constexpr (Format == PrintFormat::Width10_Precision6) {	55	} else if constexpr (Format == PrintFormat::Width10_Precision6) {
56	cce::printf("%10.6f", static_cast<float>(val));	56	cce::printf("%10.6f", static_cast<float>(val));
57	}	57	}
58	- } else if constexpr (std::is_integral_v<T>) {	58	+ } else if constexpr (std::is_signed_v<T>) {
59	if constexpr (Format == PrintFormat::Width10_Precision6) {	59	if constexpr (Format == PrintFormat::Width10_Precision6) {
60	cce::printf("%10d", static_cast<int>(val));	60	cce::printf("%10d", static_cast<int>(val));
61	} else {	61	} else {
62	cce::printf("%8d", static_cast<int>(val));	62	cce::printf("%8d", static_cast<int>(val));
63	}	63	}
		64	+ } else if constexpr (std::is_unsigned_v<T>) {
		65	+ if constexpr (Format == PrintFormat::Width10_Precision6) {
		66	+ cce::printf("%10u", static_cast<unsigned int>(val));
		67	+ } else {
		68	+ cce::printf("%8u", static_cast<unsigned int>(val));
		69	+ }
64	} else {	70	} else {
65	static_assert(sizeof(T) == 0, "Unsupported data type for Print.");	71	static_assert(sizeof(T) == 0, "Unsupported data type for Print.");
66	}	72	}

Minclude/pto/npu/a5/TPartBinOps.hpp+41-85

	@@ -8,8 +8,8 @@ INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A
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	-#ifndef TPATIALBINOPS_HPP	11	+#ifndef TPARTIALBINOPS_HPP
12	-#define TPATIALBINOPS_HPP	12	+#define TPARTIALBINOPS_HPP
13		13
14	#include <pto/common/constants.hpp>	14	#include <pto/common/constants.hpp>
15	#include <pto/common/utils.hpp>	15	#include <pto/common/utils.hpp>
	@@ -18,94 +18,50 @@ namespace pto {
18		18
19	template <typename T>	19	template <typename T>
20	struct Padding {	20	struct Padding {
21	- using Type = std::make_unsigned_t<T>;	21	+ using Type = std::conditional_t<sizeof(T) == 4, uint32_t, std::conditional_t<sizeof(T) == 2, uint16_t, uint8_t>>;
22	- static constexpr Type Null = (Type)0;
23	- static constexpr Type Zero = (Type)0;
24	- static constexpr Type Min = (Type)0;
25	- static constexpr Type Max = (Type)0xffffffffffffffffUL;
26	-};
27		22
28	-template <>	23	+ PTO_INTERNAL static constexpr Type GetPaddingMin()
29	-struct Padding<float> {	24	+ {
30	- using Type = uint32_t;	25	+ if constexpr (std::is_same_v<T, float>) {
31	- static constexpr Type Null = (Type)0;	26	+ return (Type)0xff800000UL;
32	- static constexpr Type Zero = (Type)0;	27	+ } else if constexpr (std::is_same_v<T, half>) {
33	- static constexpr Type Min = (Type)0xff800000UL;	28	+ return (Type)0xfc00UL;
34	- static constexpr Type Max = (Type)0x7f800000UL;	29	+ } else if constexpr (std::is_same_v<T, bfloat16_t>) {
35	-};	30	+ return (Type)0xff80UL;
		31	+ } else if constexpr (std::is_same_v<T, int32_t>) {
		32	+ return (Type)0x80000000UL;
		33	+ } else if constexpr (std::is_same_v<T, int16_t>) {
		34	+ return (Type)0x8000UL;
		35	+ } else if constexpr (std::is_same_v<T, int8_t>) {
		36	+ return (Type)0x80UL;
		37	+ } else {
		38	+ return (Type)0;
		39	+ }
		40	+ }
36		41
37	-template <>	42	+ PTO_INTERNAL static constexpr Type GetPaddingMax()
38	-struct Padding<int32_t> {	43	+ {
39	- using Type = uint32_t;	44	+ if constexpr (std::is_same_v<T, float>) {
40	- static constexpr Type Null = (Type)0;	45	+ return (Type)0x7f800000UL;
41	- static constexpr Type Zero = (Type)0;	46	+ } else if constexpr (std::is_same_v<T, half>) {
42	- static constexpr Type Min = (Type)0x80000000UL;	47	+ return (Type)0x7c00UL;
43	- static constexpr Type Max = (Type)0x7fffffffUL;	48	+ } else if constexpr (std::is_same_v<T, bfloat16_t>) {
44	-};	49	+ return (Type)0x7f80UL;
		50	+ } else if constexpr (std::is_same_v<T, int32_t>) {
		51	+ return (Type)0x7fffffffUL;
		52	+ } else if constexpr (std::is_same_v<T, int16_t>) {
		53	+ return (Type)0x7fffUL;
		54	+ } else if constexpr (std::is_same_v<T, int8_t>) {
		55	+ return (Type)0x7fUL;
		56	+ } else {
		57	+ return (Type)(~(Type)0);
		58	+ }
		59	+ }
45		60
46	-template <>
47	-struct Padding<uint32_t> {
48	- using Type = uint32_t;
49	static constexpr Type Null = (Type)0;	61	static constexpr Type Null = (Type)0;
50	static constexpr Type Zero = (Type)0;	62	static constexpr Type Zero = (Type)0;
51	- static constexpr Type Min = (Type)0;	63	+ static constexpr Type Min = GetPaddingMin();
52	- static constexpr Type Max = (Type)0xffffffffUL;	64	+ static constexpr Type Max = GetPaddingMax();
53	-};
54	-
55	-#ifndef PTO_NPU_ARCH_KIRIN9030
56	-template <>
57	-struct Padding<bfloat16_t> {
58	- using Type = uint16_t;
59	- static constexpr Type Null = (Type)0;
60	- static constexpr Type Zero = (Type)0;
61	- static constexpr Type Min = (Type)0xff80UL;
62	- static constexpr Type Max = (Type)0x7f80UL;
63	-};
64	-#endif
65	-
66	-template <>
67	-struct Padding<half> {
68	- using Type = uint16_t;
69	- static constexpr Type Null = (Type)0;
70	- static constexpr Type Zero = (Type)0;
71	- static constexpr Type Min = (Type)0xfc00UL;
72	- static constexpr Type Max = (Type)0x7c00UL;
73	-};
74	-
75	-template <>
76	-struct Padding<int16_t> {
77	- using Type = uint16_t;
78	- static constexpr Type Null = (Type)0;
79	- static constexpr Type Zero = (Type)0;
80	- static constexpr Type Min = (Type)0x8000UL;
81	- static constexpr Type Max = (Type)0x7fffUL;
82	-};
83	-
84	-template <>
85	-struct Padding<uint16_t> {
86	- using Type = uint16_t;
87	- static constexpr Type Null = (Type)0;
88	- static constexpr Type Zero = (Type)0;
89	- static constexpr Type Min = (Type)0;
90	- static constexpr Type Max = (Type)0xffffUL;
91	-};
92	-
93	-template <>
94	-struct Padding<int8_t> {
95	- using Type = uint8_t;
96	- static constexpr Type Null = (Type)0;
97	- static constexpr Type Zero = (Type)0;
98	- static constexpr Type Min = (Type)0x80UL;
99	- static constexpr Type Max = (Type)0x7fUL;
100	-};
101	-
102	-template <>
103	-struct Padding<uint8_t> {
104	- using Type = uint8_t;
105	- static constexpr Type Null = (Type)0;
106	- static constexpr Type Zero = (Type)0;
107	- static constexpr Type Min = (Type)0;
108	- static constexpr Type Max = (Type)0xffUL;
109	};	65	};
110		66
111	template <typename Op, typename T, unsigned elementsPerRepeat, unsigned dstStride>	67	template <typename Op, typename T, unsigned elementsPerRepeat, unsigned dstStride>

Ainclude/pto/npu/a5/TPow.hpp+676-0

@@ -0,0 +1,676 @@
	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 TPOW_HPP
	12	+#define TPOW_HPP
	13	+
	14	+#include <pto/common/constants.hpp>
	15	+#include <pto/common/utils.hpp>
	16	+#include <pto/common/type.hpp>
	17	+
	18	+#include "utils.hpp"
	19	+
	20	+namespace pto {
	21	+namespace PowF {
	22	+constexpr float LOG2_LOWEST_VALUE = 1.175494351e-38f;
	23	+constexpr float LOG2_LOWEST_VALUE_MULS = 8388608.0f;
	24	+constexpr float LOG2_REDUCE_COEFF1 = 0.70710678f;
	25	+constexpr int32_t LOG2_REDUCE_COEFF2 = 0xff800000;
	26	+constexpr float LOG2_REDUCE_FMAF_COEFF1 = 1.19209290e-7f;
	27	+constexpr float LOG2_BEST_FMAF_COEFF1 = 0.129394531f;
	28	+constexpr float LOG2_BEST_FMAF_COEFF2 = 0.141957462f;
	29	+constexpr float LOG2_BEST_FMAF_COEFF3 = 0.200015724f;
	30	+constexpr float LOG2_BEST_FMAF_COEFF4 = 0.333333254f;
	31	+constexpr float LOG2_HI1 = 6.93147182e-1f;
	32	+constexpr float LOG2_HI2 = -6.93147182e-1f;
	33	+constexpr float LOG2_LO = -1.90465421e-9f;
	34	+constexpr float EXP_OVFL_UNFL_F = -104.0f;
	35	+constexpr float EXP_MIN_F = 88.7228390f;
	36	+constexpr int32_t INF = 0x7F800000;
	37	+constexpr int32_t NEG_INF = 0xff800000;
	38	+constexpr int32_t F32_NAN = 0x7fc00000;
	39	+constexpr int32_t R10_COEFF = 0x7F800000;
	40	+constexpr int32_t R12_COEFF = 0x7FFFFFFF;
	41	+constexpr int16_t COMPARE_ZERO_OFFSET = 31;
	42	+constexpr float F32_FRACTIONS = -23.0f;
	43	+
	44	+PTO_INTERNAL void IsInfNum(MaskReg &infMask, RegTensor<float> &srcReg, RegTensor<int32_t> &tmpR12Reg, MaskReg &mask)
	45	+{
	46	+ RegTensor<float> tmpFloatReg;
	47	+ vand((RegTensor<int32_t> &)tmpFloatReg, (RegTensor<int32_t> &)srcReg, tmpR12Reg, mask);
	48	+ vcmps_eq(infMask, (RegTensor<int32_t> &)tmpFloatReg, INF, mask);
	49	+}
	50	+
	51	+PTO_INTERNAL void IsNanNum(MaskReg &nanMask, RegTensor<float> &srcReg, MaskReg &mask)
	52	+{
	53	+ vcmp_ne(nanMask, srcReg, srcReg, mask);
	54	+}
	55	+
	56	+PTO_INTERNAL void RFloor(RegTensor<float> &dstReg, RegTensor<float> &srcReg, MaskReg &mask)
	57	+{
	58	+ vtrc(dstReg, srcReg, ROUND_F, mask, MODE_ZEROING);
	59	+}
	60	+
	61	+PTO_INTERNAL void ComputeExpoOddInt(MaskReg &oddMask, RegTensor<float> &expReg, RegTensor<float> &twoReg, MaskReg &mask)
	62	+{
	63	+ // calculate exp is odd or not: expo_odd_int = fmaf (-2.0f, floorf (0.5f * b), b) == 1.0f;
	64	+ RegTensor<float> tmpFloatReg;
	65	+ vmuls(tmpFloatReg, expReg, 0.5f, mask);
	66	+ RFloor(tmpFloatReg, tmpFloatReg, mask);
	67	+ vmadd(tmpFloatReg, twoReg, expReg, mask);
	68	+ vcmps_eq(oddMask, tmpFloatReg, 1.0f, mask);
	69	+}
	70	+
	71	+PTO_INTERNAL void ProcessFloatSpecialCase(RegTensor<float> &dstReg, RegTensor<float> &baseReg, RegTensor<float> &expReg,
	72	+ RegTensor<int32_t> &tmpR10Reg, RegTensor<int32_t> &tmpR12Reg,
	73	+ RegTensor<float> &twoReg, MaskReg &mask)
	74	+{
	75	+ RegTensor<float> tmpFloatReg, tmpFloatReg2;
	76	+ MaskReg cmpMask1, cmpMask2, curMask;
	77	+
	78	+ // 1. 基本情况
	79	+ // if (exp == 0.0f \|\| base == 1.0f)
	80	+ // return 1.0f;
	81	+ vcmps_eq(cmpMask1, expReg, 0.0f, mask);
	82	+ vcmps_eq(cmpMask2, baseReg, 1.0f, mask);
	83	+ por(cmpMask2, cmpMask1, cmpMask2, mask);
	84	+ vdup(dstReg, 1.0f, cmpMask2, MODE_MERGING);
	85	+
	86	+ // 2. NaN处理
	87	+ // if (isnan(base) \|\| isnan(exp))
	88	+ // return NAN;
	89	+ pnot(curMask, cmpMask2, mask);
	90	+ IsNanNum(cmpMask1, baseReg, mask);
	91	+ IsNanNum(cmpMask2, expReg, mask);
	92	+ por(cmpMask2, cmpMask1, cmpMask2, curMask);
	93	+ vdup((RegTensor<int32_t> &)dstReg, F32_NAN, cmpMask2, MODE_MERGING);
	94	+
	95	+ // 3. 无穷大和零处理
	96	+ // if (isinf(base) \|\| base == 0.0f) {
	97	+ // if (exp < 0.0f)
	98	+ // return base ^ 0x7F800000; // 反转指数位
	99	+ // return base & 0x7FFFFFFF; // 取绝对值
	100	+ // }
	101	+ pxor(curMask, cmpMask2, curMask, mask);
	102	+ IsInfNum(cmpMask1, baseReg, tmpR12Reg, curMask);
	103	+ vcmps_eq(cmpMask2, baseReg, 0.0f, curMask);
	104	+ por(cmpMask1, cmpMask1, cmpMask2, mask);
	105	+ vcmps_lt(cmpMask2, expReg, 0.0f, cmpMask1);
	106	+ vxor((RegTensor<int32_t> &)tmpFloatReg, (RegTensor<int32_t> &)baseReg, tmpR10Reg, curMask);
	107	+ vsel(tmpFloatReg, tmpFloatReg, baseReg, cmpMask2);
	108	+ vand((RegTensor<int32_t> &)tmpFloatReg2, (RegTensor<int32_t> &)tmpFloatReg, tmpR12Reg, curMask);
	109	+ ComputeExpoOddInt(cmpMask2, expReg, twoReg, mask);
	110	+ vsel(tmpFloatReg, tmpFloatReg, tmpFloatReg2, cmpMask2);
	111	+ vsel(dstReg, tmpFloatReg, dstReg, cmpMask1);
	112	+
	113	+ // 4. 负数底数处理
	114	+ // if (base < 0.0f) {
	115	+ // if (exp != floor(exp))
	116	+ // return NAN; // 负数的非整数幂为NaN
	117	+ // if (is_odd(exp))
	118	+ // return -result;
	119	+ // }
	120	+ pxor(curMask, cmpMask1, curMask, mask);
	121	+ vneg(tmpFloatReg, dstReg, curMask);
	122	+ vsel(tmpFloatReg, tmpFloatReg, dstReg, cmpMask2);
	123	+ RFloor(tmpFloatReg2, expReg, curMask);
	124	+ vcmp_eq(cmpMask1, expReg, tmpFloatReg2, curMask);
	125	+ vdup((RegTensor<int32_t> &)tmpFloatReg, F32_NAN, cmpMask1, MODE_MERGING);
	126	+ vcmps_lt(cmpMask2, baseReg, 0.0f, curMask);
	127	+ vsel(dstReg, tmpFloatReg, dstReg, cmpMask2);
	128	+
	129	+ // 5. 特殊组合处理
	130	+ // if (base == -1.0f && isinf(exp))
	131	+ // return 1.0f;
	132	+ vcmps_eq(cmpMask1, expReg, INF, curMask);
	133	+ vcmps_eq(cmpMask2, expReg, NEG_INF, curMask);
	134	+ por(cmpMask1, cmpMask1, cmpMask2, mask);
	135	+ vcmps_eq(cmpMask2, baseReg, -1.0f, cmpMask1);
	136	+ vdup(dstReg, 1.0f, cmpMask2, MODE_MERGING);
	137	+}
	138	+
	139	+template <typename T>
	140	+PTO_INTERNAL void LoadSrcData(RegTensor<float> &srcReg, __ubuf__ T *src, uint16_t offset, MaskReg &mask)
	141	+{
	142	+ if constexpr (isSupportType<T, half, bfloat16_t>) {
	143	+ RegTensor<T> tmpReg;
	144	+ vlds(tmpReg, src, offset, UNPK_B16);
	145	+ vcvt(srcReg, tmpReg, mask, PART_EVEN);
	146	+ } else {
	147	+ vlds(srcReg, src, offset, NORM);
	148	+ }
	149	+}
	150	+
	151	+template <typename T>
	152	+PTO_INTERNAL void StoreDstData(RegTensor<float> &dstReg, __ubuf__ T *dst, uint16_t offset, MaskReg &mask)
	153	+{
	154	+ constexpr auto distValue =
	155	+ std::integral_constant<::DistVST, static_cast<::DistVST>(GetDistVst<T, DistVST::DIST_NORM>())>();
	156	+ if constexpr (isSupportType<T, half, bfloat16_t>) {
	157	+ RegTensor<T> tmpReg;
	158	+ MaskReg tmpMask;
	159	+ vcvt(tmpReg, dstReg, mask, ROUND_R, RS_DISABLE, PART_EVEN);
	160	+ vpack((RegTensor<uint16_t> &)tmpReg, (RegTensor<uint32_t> &)tmpReg, LOWER);
	161	+ ppack(tmpMask, mask, LOWER);
	162	+ vsts(tmpReg, dst, offset, distValue, tmpMask);
	163	+ } else {
	164	+ vsts(dstReg, dst, offset, distValue, mask);
	165	+ }
	166	+}
	167	+
	168	+PTO_INTERNAL void GetTPowFloatCore(RegTensor<float> &dstReg, RegTensor<float> &baseReg, RegTensor<float> &expReg,
	169	+ MaskReg &mask)
	170	+{
	171	+ RegTensor<float> tmpReg;
	172	+ vabs(tmpReg, baseReg, mask);
	173	+ vln(tmpReg, tmpReg, mask, MODE_ZEROING);
	174	+ vmul(dstReg, expReg, tmpReg, mask, MODE_ZEROING);
	175	+ vexp(dstReg, dstReg, mask, MODE_ZEROING);
	176	+}
	177	+
	178	+template <typename T, uint32_t DstStride, uint32_t BaseStride, uint32_t ExpStride>
	179	+PTO_INTERNAL void TPowFloat(__ubuf__ T dst, __ubuf__ T base, __ubuf__ T *exp, unsigned validRow, unsigned validCol)
	180	+{
	181	+ constexpr uint16_t nElemPerRpt = CCE_VL / sizeof(float);
	182	+ uint16_t repeatTime = CeilDivision(validCol, nElemPerRpt);
	183	+
	184	+ __VEC_SCOPE__
	185	+ {
	186	+ unsigned sReg = validCol;
	187	+ MaskReg mask = CreatePredicate<float>(sReg);
	188	+ MaskReg tmpMask;
	189	+ RegTensor<float> baseReg, expReg, dstReg, twoReg;
	190	+ RegTensor<int32_t> tmpR10Reg, tmpR12Reg;
	191	+ vdup(tmpR10Reg, R10_COEFF, mask, MODE_ZEROING);
	192	+ vdup(tmpR12Reg, R12_COEFF, mask, MODE_ZEROING);
	193	+ vdup(twoReg, -2.0f, mask, MODE_ZEROING);
	194	+
	195	+ for (uint16_t i = 0; i < (uint16_t)validRow; i++) {
	196	+ sReg = validCol;
	197	+ for (uint16_t j = 0; j < repeatTime; j++) {
	198	+ mask = CreatePredicate<float>(sReg);
	199	+ tmpMask = mask;
	200	+ LoadSrcData(baseReg, base, i * BaseStride + j * nElemPerRpt, mask);
	201	+ LoadSrcData(expReg, exp, i * ExpStride + j * nElemPerRpt, mask);
	202	+ GetTPowFloatCore(dstReg, baseReg, expReg, mask);
	203	+ ProcessFloatSpecialCase(dstReg, baseReg, expReg, tmpR10Reg, tmpR12Reg, twoReg, tmpMask);
	204	+ StoreDstData(dstReg, dst, i * DstStride + j * nElemPerRpt, mask);
	205	+ }
	206	+ }
	207	+ }
	208	+}
	209	+
	210	+template <typename T, uint32_t DstStride, uint32_t BaseStride>
	211	+PTO_INTERNAL void TPowFloat(__ubuf__ T dst, __ubuf__ T base, T exp, unsigned validRow, unsigned validCol)
	212	+{
	213	+ constexpr uint16_t nElemPerRpt = CCE_VL / sizeof(float);
	214	+ uint16_t repeatTime = CeilDivision(validCol, nElemPerRpt);
	215	+ __VEC_SCOPE__
	216	+ {
	217	+ unsigned sReg = validCol;
	218	+ MaskReg mask = CreatePredicate<float>(sReg);
	219	+ MaskReg tmpMask;
	220	+ RegTensor<float> baseReg, expReg, dstReg, twoReg;
	221	+ RegTensor<int32_t> tmpR10Reg, tmpR12Reg;
	222	+ vdup(expReg, exp, mask, MODE_ZEROING);
	223	+ vdup(tmpR10Reg, R10_COEFF, mask, MODE_ZEROING);
	224	+ vdup(tmpR12Reg, R12_COEFF, mask, MODE_ZEROING);
	225	+ vdup(twoReg, -2.0f, mask, MODE_ZEROING);
	226	+ for (uint16_t i = 0; i < (uint16_t)validRow; i++) {
	227	+ sReg = validCol;
	228	+ for (uint16_t j = 0; j < repeatTime; j++) {
	229	+ mask = CreatePredicate<float>(sReg);
	230	+ tmpMask = mask;
	231	+ LoadSrcData(baseReg, base, i * BaseStride + j * nElemPerRpt, mask);
	232	+ GetTPowFloatCore(dstReg, baseReg, expReg, mask);
	233	+ ProcessFloatSpecialCase(dstReg, baseReg, expReg, tmpR10Reg, tmpR12Reg, twoReg, tmpMask);
	234	+ StoreDstData(dstReg, dst, i * DstStride + j * nElemPerRpt, mask);
	235	+ }
	236	+ }
	237	+ }
	238	+}
	239	+
	240	+struct PowerLogParams {
	241	+ RegTensor<float> zeroReg;
	242	+ RegTensor<float> oneReg;
	243	+ RegTensor<float> fractionReg;
	244	+ RegTensor<float> subReg;
	245	+ RegTensor<int32_t> intReg;
	246	+ RegTensor<float> rReg;
	247	+ RegTensor<float> addReg1;
	248	+ RegTensor<float> addReg2;
	249	+ RegTensor<float> twoReg;
	250	+ RegTensor<int32_t> tmpR10Reg;
	251	+ RegTensor<int32_t> tmpR12Reg;
	252	+};
	253	+
	254	+PTO_INTERNAL void PowerLogParamsInit(PowerLogParams &params, MaskReg &mask)
	255	+{
	256	+ vdup(params.zeroReg, 0.0f, mask, MODE_ZEROING);
	257	+ vdup(params.oneReg, 1.0f, mask, MODE_ZEROING);
	258	+ vdup(params.fractionReg, F32_FRACTIONS, mask, MODE_ZEROING);
	259	+ vdup(params.subReg, LOG2_REDUCE_COEFF1, mask, MODE_ZEROING);
	260	+ vdup(params.intReg, LOG2_REDUCE_COEFF2, mask, MODE_ZEROING);
	261	+ vdup(params.rReg, LOG2_BEST_FMAF_COEFF2, mask, MODE_ZEROING);
	262	+ vdup(params.addReg1, LOG2_BEST_FMAF_COEFF3, mask, MODE_ZEROING);
	263	+ vdup(params.addReg2, LOG2_BEST_FMAF_COEFF4, mask, MODE_ZEROING);
	264	+ vdup(params.tmpR10Reg, R10_COEFF, mask, MODE_ZEROING);
	265	+ vdup(params.tmpR12Reg, R12_COEFF, mask, MODE_ZEROING);
	266	+ vdup(params.twoReg, -2.0f, mask, MODE_ZEROING);
	267	+}
	268	+
	269	+PTO_INTERNAL void GetLogFExtStepOne(RegTensor<float> &logHighReg, RegTensor<float> &logLowReg,
	270	+ RegTensor<float> &tmpResultReg, RegTensor<float> &baseReg, PowerLogParams &params,
	271	+ MaskReg &mask)
	272	+{
	273	+ RegTensor<float> tmpAReg, tmpFloatReg, absReg;
	274	+ RegTensor<int32_t> tmpEReg;
	275	+ MaskReg cmpMask;
	276	+
	277	+ vabs(absReg, baseReg, mask);
	278	+ vcmps_lt(cmpMask, absReg, LOG2_LOWEST_VALUE, mask);
	279	+ vmuls(tmpAReg, absReg, LOG2_LOWEST_VALUE_MULS, mask);
	280	+ vsel(logHighReg, params.fractionReg, params.zeroReg, cmpMask);
	281	+
	282	+ tmpFloatReg = params.subReg;
	283	+ vsub(tmpEReg, (RegTensor<int32_t> &)absReg, (RegTensor<int32_t> &)tmpFloatReg, mask);
	284	+ vand(tmpEReg, tmpEReg, params.intReg, mask);
	285	+ vsub((RegTensor<int32_t> &)logLowReg, (RegTensor<int32_t> &)absReg, tmpEReg, mask);
	286	+ vcvt(tmpFloatReg, tmpEReg, mask, ROUND_A);
	287	+ vaxpy(logHighReg, tmpFloatReg, LOG2_REDUCE_FMAF_COEFF1, mask);
	288	+ RegTensor<float> tmpPReg;
	289	+ vadds(tmpPReg, logLowReg, 1.0f, mask);
	290	+ vadds(logLowReg, logLowReg, -1.0f, mask);
	291	+
	292	+ vdiv(tmpResultReg, params.oneReg, tmpPReg, mask);
	293	+}
	294	+
	295	+PTO_INTERNAL void GetLogFExtStepTwo(RegTensor<float> &logHigh, RegTensor<float> &logLow, RegTensor<float> &tmpRReg,
	296	+ PowerLogParams &params, MaskReg &mask)
	297	+{
	298	+ RegTensor<float> tmpQHIReg, tmpQLOReg;
	299	+ RegTensor<float> tmpFloatReg, tmpFloatReg2;
	300	+ vmul(tmpQHIReg, logLow, tmpRReg, mask, MODE_ZEROING);
	301	+ vmuls(tmpFloatReg, tmpQHIReg, -2.0f, mask);
	302	+ vadd(tmpFloatReg, tmpFloatReg, logLow, mask);
	303	+ vneg(tmpFloatReg2, logLow, mask);
	304	+ vmadd(tmpFloatReg2, tmpQHIReg, tmpFloatReg, mask, MODE_ZEROING);
	305	+ vmul(tmpQLOReg, tmpRReg, tmpFloatReg2, mask, MODE_ZEROING);
	306	+ RegTensor<float> tmpSReg;
	307	+ vmul(tmpSReg, tmpQHIReg, tmpQHIReg, mask, MODE_ZEROING);
	308	+ tmpRReg = params.rReg;
	309	+ vaxpy(tmpRReg, tmpSReg, LOG2_BEST_FMAF_COEFF1, mask, MODE_ZEROING);
	310	+ vmadd(tmpRReg, tmpSReg, params.addReg1, mask, MODE_ZEROING);
	311	+ vmadd(tmpRReg, tmpSReg, params.addReg2, mask, MODE_ZEROING);
	312	+ vmul(tmpRReg, tmpRReg, tmpSReg, mask, MODE_ZEROING);
	313	+ vadd(tmpQHIReg, tmpQHIReg, tmpQHIReg, mask);
	314	+ vadd(tmpQLOReg, tmpQLOReg, tmpQLOReg, mask);
	315	+ RegTensor<float> tmpFHIReg, tmpFLOReg;
	316	+ vmuls(tmpFHIReg, logHigh, LOG2_HI1, mask);
	317	+ vadd(tmpFHIReg, tmpFHIReg, tmpQHIReg, mask);
	318	+ tmpFloatReg2 = tmpFHIReg;
	319	+ vaxpy(tmpFloatReg2, logHigh, LOG2_HI2, mask);
	320	+ vsub(tmpFLOReg, tmpQHIReg, tmpFloatReg2, mask);
	321	+ vmadd(tmpQHIReg, tmpRReg, tmpFLOReg, mask, MODE_ZEROING);
	322	+ vmuls(tmpFloatReg, tmpQLOReg, 3.0f, mask);
	323	+ vmula(tmpQLOReg, tmpFloatReg, tmpRReg, mask);
	324	+ vaxpy(tmpQLOReg, logHigh, LOG2_LO, mask);
	325	+ vadd(tmpQLOReg, tmpQLOReg, tmpQHIReg, mask);
	326	+ vadd(logHigh, tmpFHIReg, tmpQLOReg, mask);
	327	+ vsub(tmpFloatReg, tmpFHIReg, logHigh, mask);
	328	+ vadd(logLow, tmpFloatReg, tmpQLOReg, mask);
	329	+}
	330	+
	331	+PTO_INTERNAL void GetExpCore(RegTensor<float> &dstReg, RegTensor<float> &logHighReg, RegTensor<float> &logLowReg,
	332	+ RegTensor<float> &expReg, MaskReg &mask)
	333	+{
	334	+ RegTensor<float> tmPHIReg, tmPLOReg, tmpRReg;
	335	+ vmul(tmPHIReg, logHighReg, expReg, mask, MODE_ZEROING);
	336	+ RegTensor<float> tmpFloatReg, tmpFloatReg2;
	337	+ vneg(tmPLOReg, tmPHIReg, mask);
	338	+ vmula(tmPLOReg, logHighReg, expReg, mask);
	339	+ vmula(tmPLOReg, logLowReg, expReg, mask);
	340	+ vexp(tmpRReg, tmPHIReg, mask, MODE_ZEROING);
	341	+ vmula(tmpRReg, tmPLOReg, tmpRReg, mask);
	342	+ MaskReg cmpMask1, cmpMask2;
	343	+ vcmps_ge(cmpMask1, tmPHIReg, 0.0f, mask);
	344	+ vdup((RegTensor<int32_t> &)tmpFloatReg, INF, cmpMask1, MODE_ZEROING);
	345	+ vcmps_ge(cmpMask2, tmPHIReg, EXP_MIN_F, mask);
	346	+ vcmps_lt(cmpMask1, tmPHIReg, EXP_OVFL_UNFL_F, mask);
	347	+ por(cmpMask2, cmpMask1, cmpMask2, mask);
	348	+ vsel(dstReg, tmpFloatReg, tmpRReg, cmpMask2);
	349	+}
	350	+
	351	+template <typename T, uint32_t DstStride, uint32_t BaseStride, uint32_t ExpStride>
	352	+PTO_INTERNAL void TPowFloatHighPrecisionImpl(__ubuf__ T dst, __ubuf__ T base, __ubuf__ T *exp, unsigned validRow,
	353	+ unsigned validCol)
	354	+{
	355	+ constexpr uint16_t nElemPerRpt = CCE_VL / sizeof(float);
	356	+ uint16_t repeatTime = CeilDivision(validCol, nElemPerRpt);
	357	+
	358	+ __VEC_SCOPE__
	359	+ {
	360	+ unsigned sReg = validCol;
	361	+ MaskReg mask = CreatePredicate<float>(sReg);
	362	+ PowerLogParams params;
	363	+ PowerLogParamsInit(params, mask);
	364	+
	365	+ RegTensor<float> baseReg, expReg, dstReg;
	366	+ RegTensor<float> logHighReg, logLowReg, tmpResultReg;
	367	+ for (uint16_t i = 0; i < (uint16_t)validRow; i++) {
	368	+ sReg = validCol;
	369	+ for (uint16_t j = 0; j < repeatTime; j++) {
	370	+ mask = CreatePredicate<float>(sReg);
	371	+
	372	+ LoadSrcData(expReg, exp, i * ExpStride + j * nElemPerRpt, mask);
	373	+ LoadSrcData(baseReg, base, i * BaseStride + j * nElemPerRpt, mask);
	374	+
	375	+ GetLogFExtStepOne(logHighReg, logLowReg, tmpResultReg, baseReg, params, mask);
	376	+ GetLogFExtStepTwo(logHighReg, logLowReg, tmpResultReg, params, mask);
	377	+ GetExpCore(dstReg, logHighReg, logLowReg, expReg, mask);
	378	+ ProcessFloatSpecialCase(dstReg, baseReg, expReg, params.tmpR10Reg, params.tmpR12Reg, params.twoReg,
	379	+ mask);
	380	+
	381	+ StoreDstData(dstReg, dst, i * DstStride + j * nElemPerRpt, mask);
	382	+ }
	383	+ }
	384	+ }
	385	+}
	386	+
	387	+template <typename T, uint32_t DstStride, uint32_t BaseStride>
	388	+PTO_INTERNAL void TPowFloatHighPrecisionImpl(__ubuf__ T dst, __ubuf__ T base, T exp, unsigned validRow,
	389	+ unsigned validCol)
	390	+{
	391	+ constexpr uint16_t nElemPerRpt = CCE_VL / sizeof(float);
	392	+ uint16_t repeatTime = CeilDivision(validCol, nElemPerRpt);
	393	+
	394	+ __VEC_SCOPE__
	395	+ {
	396	+ unsigned sReg = validCol;
	397	+ MaskReg mask = CreatePredicate<float>(sReg);
	398	+ PowerLogParams param;
	399	+ PowerLogParamsInit(param, mask);
	400	+
	401	+ RegTensor<float> baseReg, expReg, dstReg;
	402	+ RegTensor<float> logLowReg, logHighReg, tmpResultReg;
	403	+ vdup(expReg, exp, mask, MODE_ZEROING);
	404	+ for (uint16_t i = 0; i < (uint16_t)validRow; i++) {
	405	+ sReg = validCol;
	406	+ for (uint16_t j = 0; j < repeatTime; j++) {
	407	+ mask = CreatePredicate<float>(sReg);
	408	+ LoadSrcData(baseReg, base, i * BaseStride + j * nElemPerRpt, mask);
	409	+
	410	+ GetLogFExtStepOne(logHighReg, logLowReg, tmpResultReg, baseReg, param, mask);
	411	+ GetLogFExtStepTwo(logHighReg, logLowReg, tmpResultReg, param, mask);
	412	+ GetExpCore(dstReg, logHighReg, logLowReg, expReg, mask);
	413	+ ProcessFloatSpecialCase(dstReg, baseReg, expReg, param.tmpR10Reg, param.tmpR12Reg, param.twoReg, mask);
	414	+
	415	+ StoreDstData(dstReg, dst, i * DstStride + j * nElemPerRpt, mask);
	416	+ }
	417	+ }
	418	+ }
	419	+}
	420	+
	421	+} // namespace PowF
	422	+
	423	+namespace PowI {
	424	+constexpr int16_t SHIFT_ONE_BIT = 1;
	425	+constexpr int16_t BITS_PER_BYTE = 8;
	426	+
	427	+template <typename T, typename ConvType>
	428	+PTO_INTERNAL void LoadSrcData(RegTensor<ConvType> &srcReg, __ubuf__ T *src, uint32_t offset, MaskReg &mask)
	429	+{
	430	+ if constexpr (sizeof(T) == 1) {
	431	+ RegTensor<T> tmpReg;
	432	+ vlds(tmpReg, src, offset, UNPK_B8);
	433	+ vcvt(srcReg, tmpReg, mask, PART_EVEN);
	434	+ } else {
	435	+ vlds(srcReg, src, offset, NORM);
	436	+ }
	437	+}
	438	+
	439	+template <typename T, typename ConvType>
	440	+PTO_INTERNAL void StoreDstData(RegTensor<ConvType> &dstReg, __ubuf__ T *dst, uint32_t offset, MaskReg &mask)
	441	+{
	442	+ constexpr auto distValue =
	443	+ std::integral_constant<::DistVST, static_cast<::DistVST>(GetDistVst<T, DistVST::DIST_NORM>())>();
	444	+ if constexpr (sizeof(T) == 1) {
	445	+ RegTensor<T> tmpReg;
	446	+ MaskReg tmpMask;
	447	+ vpack((RegTensor<uint8_t> &)tmpReg, (RegTensor<uint16_t> &)dstReg, LOWER);
	448	+ ppack(tmpMask, mask, LOWER);
	449	+ vsts(tmpReg, dst, offset, distValue, tmpMask);
	450	+ } else {
	451	+ vsts(dstReg, dst, offset, distValue, mask);
	452	+ }
	453	+}
	454	+
	455	+template <typename T>
	456	+PTO_INTERNAL void GetPowI(T &dstReg, T &baseReg, T &expReg, MaskReg &mask)
	457	+{
	458	+ T selReg;
	459	+ MaskReg selMask;
	460	+ vdup(selReg, 1, mask, MODE_ZEROING);
	461	+ vand(selReg, expReg, selReg, mask);
	462	+ vcmps_eq(selMask, selReg, 1, mask);
	463	+
	464	+ T tmpReg;
	465	+ vmul(tmpReg, dstReg, baseReg, mask, MODE_ZEROING);
	466	+ vsel(dstReg, tmpReg, dstReg, selMask);
	467	+
	468	+ vshrs(expReg, expReg, SHIFT_ONE_BIT, mask);
	469	+ vmul(baseReg, baseReg, baseReg, mask, MODE_ZEROING);
	470	+}
	471	+
	472	+template <typename T>
	473	+PTO_INTERNAL void ProcessSpecialCaseForPowI(T &dstReg, T &baseReg, T &expReg, MaskReg &mask)
	474	+{
	475	+ T tmpReg;
	476	+ vdup(tmpReg, 1, mask, MODE_ZEROING);
	477	+
	478	+ MaskReg cmpMask1, cmpMask2, condMask;
	479	+ vcmps_eq(cmpMask1, expReg, 0, mask);
	480	+ vcmps_eq(cmpMask2, baseReg, 1, mask);
	481	+ por(condMask, cmpMask1, cmpMask2, mask);
	482	+
	483	+ vsel(dstReg, tmpReg, dstReg, condMask);
	484	+ pxor(mask, mask, condMask, mask);
	485	+}
	486	+
	487	+template <typename ConvType, typename T>
	488	+PTO_INTERNAL void GetPowICompute(T &dstReg, T &baseReg, T &expReg, MaskReg &mask)
	489	+{
	490	+ // TODO: vcmax(dst, exp); maxLoop = __buildin_clz(dst[0])
	491	+ constexpr uint16_t maxLoop = sizeof(ConvType) * BITS_PER_BYTE;
	492	+ T tmpBaseReg = baseReg;
	493	+ T tmpExpReg = expReg;
	494	+ MaskReg tmpMask = mask;
	495	+ for (uint16_t j = 0; j < maxLoop; j++) {
	496	+ GetPowI(dstReg, tmpBaseReg, tmpExpReg, mask);
	497	+ }
	498	+ ProcessSpecialCaseForPowI(dstReg, baseReg, expReg, tmpMask);
	499	+}
	500	+
	501	+template <typename T, uint32_t DstStride, uint32_t BaseStride, uint32_t ExpStride>
	502	+PTO_INTERNAL void PowIComputeImpl(__ubuf__ T dst, __ubuf__ T base, __ubuf__ T *exp, unsigned validRow,
	503	+ unsigned validCol)
	504	+{
	505	+ using ConvType = std::conditional_t<std::is_same_v<T, int8_t>, int16_t,
	506	+ std::conditional_t<std::is_same_v<T, uint8_t>, uint16_t, T>>;
	507	+ constexpr uint16_t nElemPerRpt = CCE_VL / sizeof(ConvType);
	508	+ uint16_t repeatTime = CeilDivision(validCol, nElemPerRpt);
	509	+
	510	+ __VEC_SCOPE__
	511	+ {
	512	+ unsigned sReg = validCol;
	513	+ RegTensor<ConvType> baseReg, expReg;
	514	+ RegTensor<ConvType> initRetReg, dstReg;
	515	+
	516	+ MaskReg mask = CreatePredicate<ConvType>(sReg);
	517	+ vdup(initRetReg, 1, mask, MODE_ZEROING);
	518	+ for (uint16_t i = 0; i < (uint16_t)validRow; i++) {
	519	+ sReg = validCol;
	520	+ for (uint16_t j = 0; j < repeatTime; j++) {
	521	+ mask = CreatePredicate<ConvType>(sReg);
	522	+ dstReg = initRetReg;
	523	+
	524	+ LoadSrcData(baseReg, base, i * BaseStride + j * nElemPerRpt, mask);
	525	+ LoadSrcData(expReg, exp, i * ExpStride + j * nElemPerRpt, mask);
	526	+ GetPowICompute<ConvType>(dstReg, baseReg, expReg, mask);
	527	+ StoreDstData(dstReg, dst, i * DstStride + j * nElemPerRpt, mask);
	528	+ }
	529	+ }
	530	+ }
	531	+}
	532	+
	533	+template <typename T, uint32_t DstStride, uint32_t BaseStride>
	534	+PTO_INTERNAL void PowIComputeImpl(__ubuf__ T dst, __ubuf__ T base, T exp, unsigned validRow, unsigned validCol)
	535	+{
	536	+ using ConvType = std::conditional_t<std::is_same_v<T, int8_t>, int16_t,
	537	+ std::conditional_t<std::is_same_v<T, uint8_t>, uint16_t, T>>;
	538	+ constexpr uint16_t nElemPerRpt = CCE_VL / sizeof(ConvType);
	539	+ uint16_t repeatTime = CeilDivision(validCol, nElemPerRpt);
	540	+
	541	+ __VEC_SCOPE__
	542	+ {
	543	+ unsigned sReg = validCol;
	544	+ RegTensor<ConvType> initRetReg, dstReg;
	545	+ RegTensor<ConvType> baseReg, expReg;
	546	+
	547	+ MaskReg mask = CreatePredicate<ConvType>(sReg);
	548	+ vdup(initRetReg, 1, mask, MODE_ZEROING);
	549	+ vdup(expReg, exp, mask, MODE_ZEROING);
	550	+ for (uint16_t i = 0; i < (uint16_t)validRow; i++) {
	551	+ sReg = validCol;
	552	+ for (uint16_t j = 0; j < repeatTime; j++) {
	553	+ mask = CreatePredicate<ConvType>(sReg);
	554	+ dstReg = initRetReg;
	555	+
	556	+ LoadSrcData(baseReg, base, i * BaseStride + j * nElemPerRpt, mask);
	557	+ GetPowICompute<ConvType>(dstReg, baseReg, expReg, mask);
	558	+ StoreDstData(dstReg, dst, i * DstStride + j * nElemPerRpt, mask);
	559	+ }
	560	+ }
	561	+ }
	562	+}
	563	+
	564	+} // namespace PowI
	565	+
	566	+template <typename T>
	567	+inline constexpr bool IsFloatNum = isSupportType<T, float, half, bfloat16_t>;
	568	+template <typename T>
	569	+inline constexpr bool IsIntegerNum = isSupportType<T, uint8_t, int8_t, uint16_t, int16_t, uint32_t, int32_t>;
	570	+
	571	+template <PowAlgorithm algo, typename DstTile, typename BaseTile, typename ExpTile>
	572	+__tf__ PTO_INTERNAL void TPowImpl(typename DstTile::TileDType __out__ dstData,
	573	+ typename BaseTile::TileDType __in__ baseData,
	574	+ typename ExpTile::TileDType __in__ expData, unsigned validRow, unsigned validCol,
	575	+ unsigned version = VFImplKind::VFIMPL_DEFAULT)
	576	+{
	577	+ using T = typename DstTile::DType;
	578	+ __ubuf__ T dst = (__ubuf__ T )__cce_get_tile_ptr(dstData);
	579	+ __ubuf__ T base = (__ubuf__ T )__cce_get_tile_ptr(baseData);
	580	+ __ubuf__ T exp = (__ubuf__ T )__cce_get_tile_ptr(expData);
	581	+
	582	+ if constexpr (IsFloatNum<T>) {
	583	+ if constexpr (algo == PowAlgorithm::DEFAULT) {
	584	+ PowF::TPowFloat<T, DstTile::RowStride, BaseTile::RowStride, ExpTile::RowStride>(dst, base, exp, validRow,
	585	+ validCol);
	586	+ } else if (algo == PowAlgorithm::HIGH_PRECISION) {
	587	+ PowF::TPowFloatHighPrecisionImpl<T, DstTile::RowStride, BaseTile::RowStride, ExpTile::RowStride>(
	588	+ dst, base, exp, validRow, validCol);
	589	+ }
	590	+ } else if constexpr (IsIntegerNum<T>) {
	591	+ PowI::PowIComputeImpl<T, DstTile::RowStride, BaseTile::RowStride, ExpTile::RowStride>(dst, base, exp, validRow,
	592	+ validCol);
	593	+ }
	594	+}
	595	+
	596	+template <PowAlgorithm algo, typename DstTile, typename BaseTile, typename ExpTile>
	597	+PTO_INTERNAL void PowCheckType()
	598	+{
	599	+ static_assert(DstTile::isRowMajor && BaseTile::isRowMajor && ExpTile::isRowMajor,
	600	+ "TPOW: Not supported Layout type");
	601	+ static_assert(DstTile::Loc == TileType::Vec && BaseTile::Loc == TileType::Vec && ExpTile::Loc == TileType::Vec,
	602	+ "TPOW: TileType of dst, base and exp tiles must be TileType::Vec.");
	603	+ static_assert(DstTile::ValidCol <= DstTile::Cols,
	604	+ "TPOW: Number of dst's valid columns must not be greater than number of tile columns.");
	605	+ static_assert(DstTile::ValidRow <= DstTile::Rows,
	606	+ "TPOW: Number of dst's valid rows must not be greater than number of tile rows.");
	607	+ static_assert(BaseTile::ValidCol <= BaseTile::Cols,
	608	+ "TPOW: Number of base's valid columns must not be greater than number of tile columns.");
	609	+ static_assert(BaseTile::ValidRow <= BaseTile::Rows,
	610	+ "TPOW: Number of base's valid rows must not be greater than number of tile rows.");
	611	+ static_assert(ExpTile::ValidCol <= ExpTile::Cols,
	612	+ "TPOW: Number of exp's valid columns must not be greater than number of tile columns.");
	613	+ static_assert(ExpTile::ValidRow <= ExpTile::Rows,
	614	+ "TPOW: Number of exp's valid rows must not be greater than number of tile rows.");
	615	+
	616	+ using T = typename DstTile::DType;
	617	+
	618	+ if constexpr (algo == PowAlgorithm::HIGH_PRECISION) {
	619	+ static_assert(isSupportType<T, float, half, bfloat16_t>,
	620	+ "Type must be half/float/bfloat16 in high precision algorithm.");
	621	+ } else {
	622	+ static_assert(isSupportType<T, float, half, int32_t, uint32_t, int16_t, uint16_t, int8_t, uint8_t>,
	623	+ "Type must be uint8/int8/uint16/int16/uint32/int32/half/float in default algorithm.");
	624	+ }
	625	+ static_assert(std::is_same_v<T, typename BaseTile::DType> && std::is_same_v<T, typename ExpTile::DType>,
	626	+ "TPOW: The data type of dst, base and exp must be consistent");
	627	+}
	628	+
	629	+template <PowAlgorithm algo, typename DstTile, typename BaseTile, typename ExpTile, typename TmpTile>
	630	+PTO_INTERNAL void TPOW_IMPL(DstTile &dst, BaseTile &base, ExpTile &exp, TmpTile &tmp)
	631	+{
	632	+ PowCheckType<algo, DstTile, BaseTile, ExpTile>();
	633	+ unsigned validRow = dst.GetValidRow();
	634	+ unsigned validCol = dst.GetValidCol();
	635	+ PTO_ASSERT(validCol == base.GetValidCol(), "TPOW: Number of columns of base and dst must be same.");
	636	+ PTO_ASSERT(validRow == base.GetValidRow(), "TPOW: Number of rows of base and dst must be same.");
	637	+ PTO_ASSERT(validCol == exp.GetValidCol(), "TPOW: Number of columns of exp and dst must be same.");
	638	+ PTO_ASSERT(validRow == exp.GetValidRow(), "TPOW: Number of rows of exp and dst must be same.");
	639	+
	640	+ TPowImpl<algo, DstTile, BaseTile, ExpTile>(dst.data(), base.data(), exp.data(), validRow, validCol);
	641	+}
	642	+
	643	+template <PowAlgorithm algo, typename DstTile, typename BaseTile>
	644	+__tf__ PTO_INTERNAL void TPowSImpl(typename DstTile::TileDType __out__ dstData,
	645	+ typename BaseTile::TileDType __in__ baseData, typename DstTile::DType exp,
	646	+ unsigned validRow, unsigned validCol, unsigned version = VFImplKind::VFIMPL_DEFAULT)
	647	+{
	648	+ using T = typename DstTile::DType;
	649	+ __ubuf__ T dst = (__ubuf__ T )__cce_get_tile_ptr(dstData);
	650	+ __ubuf__ T base = (__ubuf__ T )__cce_get_tile_ptr(baseData);
	651	+
	652	+ if constexpr (IsFloatNum<T>) {
	653	+ if constexpr (algo == PowAlgorithm::DEFAULT) {
	654	+ PowF::TPowFloat<T, DstTile::RowStride, BaseTile::RowStride>(dst, base, exp, validRow, validCol);
	655	+ } else if (algo == PowAlgorithm::HIGH_PRECISION) {
	656	+ PowF::TPowFloatHighPrecisionImpl<T, DstTile::RowStride, BaseTile::RowStride>(dst, base, exp, validRow,
	657	+ validCol);
	658	+ }
	659	+ } else if constexpr (IsIntegerNum<T>) {
	660	+ PowI::PowIComputeImpl<T, DstTile::RowStride, BaseTile::RowStride>(dst, base, exp, validRow, validCol);
	661	+ }
	662	+}
	663	+
	664	+template <PowAlgorithm algo, typename DstTile, typename BaseTile, typename TmpTile>
	665	+PTO_INTERNAL void TPOWS_IMPL(DstTile &dst, BaseTile &base, typename DstTile::DType exp, TmpTile &tmp)
	666	+{
	667	+ PowCheckType<algo, DstTile, BaseTile, DstTile>();
	668	+ unsigned validRow = dst.GetValidRow();
	669	+ unsigned validCol = dst.GetValidCol();
	670	+ PTO_ASSERT(validCol == base.GetValidCol(), "TPOW: Number of columns of base and dst must be same.");
	671	+ PTO_ASSERT(validRow == base.GetValidRow(), "TPOW: Number of rows of base and dst must be same.");
	672	+
	673	+ TPowSImpl<algo, DstTile, BaseTile>(dst.data(), base.data(), exp, validRow, validCol);
	674	+}
	675	+} // namespace pto
	676	+#endif

Minclude/pto/npu/kirinX90/TMov.hpp+1-1

	@@ -11,7 +11,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
11	#ifndef TMOV_HPP	11	#ifndef TMOV_HPP
12	#define TMOV_HPP	12	#define TMOV_HPP
13		13
14	-#include "pto/npu/kirinX90/TExtract.hpp"	14	+#include "pto/npu/a2a3/TExtract.hpp"
15	#include "pto/npu/a2a3/TCopy.hpp"	15	#include "pto/npu/a2a3/TCopy.hpp"
16		16
17	namespace pto {	17	namespace pto {

Minclude/pto/npu/kirinX90/header.hpp+18-4

	@@ -10,20 +10,27 @@ See LICENSE in the root of the software repository for the full text of the Lice
10		10
11	#ifndef HEADER_HPP	11	#ifndef HEADER_HPP
12	#define HEADER_HPP	12	#define HEADER_HPP
		13	+#define bfloat16_t half
		14	+#define hifloat8_t int8_t
		15	+#define float8_e4m3_t int8_t
		16	+#define float8_e5m2_t int8_t
		17	+#define float8_e8m0_t int8_t
		18	+#define float4_e2m1x2_t int64_t
		19	+#define float4_e1m2x2_t int64_t
13	#include "pto/common/utils.hpp"	20	#include "pto/common/utils.hpp"
14	#include "pto/common/constants.hpp"	21	#include "pto/common/constants.hpp"
15	-#include "pto/npu/a2a3/TAssign.hpp"
16	-#include "pto/npu/a2a3/TExtract.hpp"
17	#include "pto/npu/kirinX90/datatype.hpp"	22	#include "pto/npu/kirinX90/datatype.hpp"
18	#include "pto/npu/kirinX90/common.hpp"	23	#include "pto/npu/kirinX90/common.hpp"
19	-#include "pto/npu/kirin9030/utils.hpp"
20	-#include "pto/npu/kirin9030/TSync.hpp"
21	#include "pto/npu/kirinX90/TLoad.hpp"	24	#include "pto/npu/kirinX90/TLoad.hpp"
22	#include "pto/npu/kirinX90/TStore.hpp"	25	#include "pto/npu/kirinX90/TStore.hpp"
23	#include "pto/npu/kirinX90/TMov.hpp"	26	#include "pto/npu/kirinX90/TMov.hpp"
24	#ifdef __DAV_VEC__	27	#ifdef __DAV_VEC__
25	#include "pto/npu/kirinX90/TCvt.hpp"	28	#include "pto/npu/kirinX90/TCvt.hpp"
26	#endif	29	#endif
		30	+#include "pto/npu/a2a3/TAssign.hpp"
		31	+#include "pto/npu/a2a3/TExtract.hpp"
		32	+#include "pto/npu/kirin9030/utils.hpp"
		33	+#include "pto/npu/kirin9030/TSync.hpp"
27	#include "pto/npu/kirin9030/TAdd.hpp"	34	#include "pto/npu/kirin9030/TAdd.hpp"
28	#include "pto/npu/kirin9030/TAddS.hpp"	35	#include "pto/npu/kirin9030/TAddS.hpp"
29	#include "pto/npu/kirin9030/TDivS.hpp"	36	#include "pto/npu/kirin9030/TDivS.hpp"
	@@ -53,4 +60,11 @@ See LICENSE in the root of the software repository for the full text of the Lice
53	#include "pto/npu/kirin9030/TBinSOp.hpp"	60	#include "pto/npu/kirin9030/TBinSOp.hpp"
54	#include "pto/npu/kirin9030/TDiv.hpp"	61	#include "pto/npu/kirin9030/TDiv.hpp"
55	#include "pto/npu/kirin9030/TMul.hpp"	62	#include "pto/npu/kirin9030/TMul.hpp"
		63	+#undef bfloat16_t
		64	+#undef hifloat8_t
		65	+#undef float8_e4m3_t
		66	+#undef float8_e5m2_t
		67	+#undef float8_e8m0_t
		68	+#undef float4_e2m1x2_t
		69	+#undef float4_e1m2x2_t
56	#endif	70	#endif

Mtests/npu/a5/src/st/testcase/CMakeLists.txt+2-1

	@@ -224,7 +224,8 @@ set(ALL_TESTCASES
224	mgather	224	mgather
225	mscatter	225	mscatter
226	t_dhrystone	226	t_dhrystone
227	-	227	+ tpow
		228	+ tpows
228	tconcatidx	229	tconcatidx
229	tconcat	230	tconcat
230	thistogram	231	thistogram

Atests/npu/a5/src/st/testcase/tpow/CMakeLists.txt+11-0

@@ -0,0 +1,11 @@
	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	+pto_vec_st(tpow)

Atests/npu/a5/src/st/testcase/tpow/gen_data.py+79-0

@@ -0,0 +1,79 @@
	1	+#!/usr/bin/python3
	2	+# coding=utf-8
	3	+# --------------------------------------------------------------------------------
	4	+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
	5	+# This program is free software, you can redistribute it and/or modify it under terms and conditions of
	6	+# CANN Open Software License Agreement Version 2.0 (the "License").
	7	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	8	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	9	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	10	+# See LICENSE in the root of the software repository for the full text of the License.
	11	+# --------------------------------------------------------------------------------
	12	+
	13	+import os
	14	+import numpy as np
	15	+
	16	+np.random.seed(19)
	17	+
	18	+def gen_golden_data(param):
	19	+ dtype = param.dtype
	20	+ row, col = param.row, param.col
	21	+ valid_row, valid_col = param.valid_row, param.valid_col
	22	+
	23	+ if dtype in (np.int8, np.uint8, np.int16, np.uint16, np.int32, np.uint32):
	24	+ dtype_info = np.iinfo(dtype)
	25	+ base_arr = np.random.randint(1, 10, size=[row, col]).astype(dtype)
	26	+ exp_arr = np.random.randint(0, 5, size=[row, col]).astype(dtype)
	27	+ else:
	28	+ dtype_info = np.finfo(dtype)
	29	+ base_arr = np.random.uniform(0.1, 5.0, size=[row, col]).astype(dtype)
	30	+ exp_arr = np.random.uniform(0, 3.0, size=[row, col]).astype(dtype)
	31	+
	32	+ golden = np.zeros((row, col), dtype=dtype)
	33	+ golden[0:valid_row, 0:valid_col] = np.power(base_arr[0:valid_row, 0:valid_col], exp_arr[0:valid_row, 0:valid_col])
	34	+
	35	+ base_arr.tofile("base.bin")
	36	+ exp_arr.tofile("exp.bin")
	37	+ golden.tofile("golden.bin")
	38	+
	39	+
	40	+class TPowParams:
	41	+ def __init__(self, name, dtype, row, col, valid_row, valid_col):
	42	+ self.name = name
	43	+ self.dtype = dtype
	44	+ self.row = row
	45	+ self.col = col
	46	+ self.valid_row = valid_row
	47	+ self.valid_col = valid_col
	48	+
	49	+
	50	+if __name__ == "__main__":
	51	+ # Get the absolute path of the script
	52	+ script_dir = os.path.dirname(os.path.abspath(__file__))
	53	+ testcases_dir = os.path.join(script_dir, "testcases")
	54	+
	55	+ # Ensure the testcases directory exists
	56	+ if not os.path.exists(testcases_dir):
	57	+ os.makedirs(testcases_dir)
	58	+
	59	+ case_params_list = [
	60	+ TPowParams("TPOWTest.case1", np.float32, 64, 64, 63, 63),
	61	+ TPowParams("TPOWTest.case2", np.float16, 64, 64, 63, 63),
	62	+ TPowParams("TPOWTest.case3", np.int32, 64, 64, 63, 63),
	63	+ TPowParams("TPOWTest.case4", np.int16, 64, 64, 63, 63),
	64	+ TPowParams("TPOWTest.case5", np.int8, 64, 64, 63, 63),
	65	+ TPowParams("TPOWTest.case6", np.uint32, 64, 64, 63, 63),
	66	+ TPowParams("TPOWTest.case7", np.uint8, 64, 64, 63, 63),
	67	+ TPowParams("TPOWTest.case8", np.float32, 64, 64, 63, 63),
	68	+ TPowParams("TPOWTest.case9", np.float16, 64, 64, 63, 63),
	69	+ TPowParams("TPOWTest.case10", np.float32, 16, 256, 15, 231),
	70	+ TPowParams("TPOWTest.case11", np.float16, 16, 512, 16, 400),
	71	+ ]
	72	+
	73	+ for param in case_params_list:
	74	+ if not os.path.exists(param.name):
	75	+ os.makedirs(param.name)
	76	+ original_dir = os.getcwd()
	77	+ os.chdir(param.name)
	78	+ gen_golden_data(param)
	79	+ os.chdir(original_dir)

Atests/npu/a5/src/st/testcase/tpow/main.cpp+134-0

@@ -0,0 +1,134 @@
	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	+#include "test_common.h"
	12	+#include "acl/acl.h"
	13	+#include <gtest/gtest.h>
	14	+
	15	+using namespace std;
	16	+using namespace PtoTestCommon;
	17	+
	18	+class TPOWTest : public testing::Test {
	19	+protected:
	20	+ void SetUp() override
	21	+ {}
	22	+ void TearDown() override
	23	+ {}
	24	+};
	25	+
	26	+std::string GetGoldenDir()
	27	+{
	28	+ const testing::TestInfo *testInfo = testing::UnitTest::GetInstance()->current_test_info();
	29	+ const std::string caseName = testInfo->name();
	30	+ std::string suiteName = testInfo->test_suite_name();
	31	+ std::string fullPath = "../" + suiteName + "." + caseName;
	32	+ return fullPath;
	33	+}
	34	+
	35	+template <typename T, int TRow, int TCol, int validRow, int validCol, bool isHighPrecision>
	36	+void LaunchTPow(T out, T base, T exp, void stream);
	37	+
	38	+template <typename T, int Row, int Col, int validRow, int validCol, bool isHighPrecision = false>
	39	+void test_tpow()
	40	+{
	41	+ size_t fileSize = Row * Col * sizeof(T);
	42	+
	43	+ aclInit(nullptr);
	44	+ aclrtSetDevice(0);
	45	+ aclrtStream stream;
	46	+ aclrtCreateStream(&stream);
	47	+
	48	+ T dstHost, baseHost, *expHost;
	49	+ T dstDevice, baseDevice, *expDevice;
	50	+
	51	+ aclrtMallocHost((void **)(&dstHost), fileSize);
	52	+ aclrtMallocHost((void **)(&baseHost), fileSize);
	53	+ aclrtMallocHost((void **)(&expHost), fileSize);
	54	+ aclrtMalloc((void **)&dstDevice, fileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	55	+ aclrtMalloc((void **)&baseDevice, fileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	56	+ aclrtMalloc((void **)&expDevice, fileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	57	+
	58	+ ReadFile(GetGoldenDir() + "/base.bin", fileSize, baseHost, fileSize);
	59	+ ReadFile(GetGoldenDir() + "/exp.bin", fileSize, expHost, fileSize);
	60	+ aclrtMemcpy(baseDevice, fileSize, baseHost, fileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	61	+ aclrtMemcpy(expDevice, fileSize, expHost, fileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	62	+ LaunchTPow<T, Row, Col, validRow, validCol, isHighPrecision>(dstDevice, baseDevice, expDevice, stream);
	63	+
	64	+ aclrtSynchronizeStream(stream);
	65	+ aclrtMemcpy(dstHost, fileSize, dstDevice, fileSize, ACL_MEMCPY_DEVICE_TO_HOST);
	66	+
	67	+ WriteFile(GetGoldenDir() + "/output.bin", dstHost, fileSize);
	68	+
	69	+ aclrtFree(dstDevice);
	70	+ aclrtFree(baseDevice);
	71	+ aclrtFree(expDevice);
	72	+
	73	+ aclrtFreeHost(dstHost);
	74	+ aclrtFreeHost(baseHost);
	75	+ aclrtFreeHost(expHost);
	76	+ aclrtDestroyStream(stream);
	77	+ aclrtResetDevice(0);
	78	+ aclFinalize();
	79	+
	80	+ std::vector<T> golden(fileSize / sizeof(T));
	81	+ std::vector<T> devFinal(fileSize / sizeof(T));
	82	+ ReadFile(GetGoldenDir() + "/golden.bin", fileSize, golden.data(), fileSize);
	83	+ ReadFile(GetGoldenDir() + "/output.bin", fileSize, devFinal.data(), fileSize);
	84	+
	85	+ constexpr float eps = std::is_same_v<T, float> ? 0.0005f : 0.00005f;
	86	+ bool ret = ResultCmp<T>(golden, devFinal, eps);
	87	+
	88	+ EXPECT_TRUE(ret);
	89	+}
	90	+
	91	+TEST_F(TPOWTest, case1)
	92	+{
	93	+ test_tpow<float, 64, 64, 63, 63>();
	94	+}
	95	+TEST_F(TPOWTest, case2)
	96	+{
	97	+ test_tpow<aclFloat16, 64, 64, 63, 63>(); // typedef uint16_t aclFloat16
	98	+}
	99	+TEST_F(TPOWTest, case3)
	100	+{
	101	+ test_tpow<int32_t, 64, 64, 63, 63>();
	102	+}
	103	+TEST_F(TPOWTest, case4)
	104	+{
	105	+ test_tpow<int16_t, 64, 64, 63, 63>();
	106	+}
	107	+TEST_F(TPOWTest, case5)
	108	+{
	109	+ test_tpow<int8_t, 64, 64, 63, 63>();
	110	+}
	111	+TEST_F(TPOWTest, case6)
	112	+{
	113	+ test_tpow<uint32_t, 64, 64, 63, 63>();
	114	+}
	115	+TEST_F(TPOWTest, case7)
	116	+{
	117	+ test_tpow<uint8_t, 64, 64, 63, 63>();
	118	+}
	119	+TEST_F(TPOWTest, case8)
	120	+{
	121	+ test_tpow<float, 64, 64, 63, 63, true>();
	122	+}
	123	+TEST_F(TPOWTest, case9)
	124	+{
	125	+ test_tpow<aclFloat16, 64, 64, 63, 63, true>();
	126	+}
	127	+TEST_F(TPOWTest, case10)
	128	+{
	129	+ test_tpow<float, 16, 256, 15, 231>();
	130	+}
	131	+TEST_F(TPOWTest, case11)
	132	+{
	133	+ test_tpow<aclFloat16, 16, 512, 16, 400, true>();
	134	+}

Atests/npu/a5/src/st/testcase/tpow/tpow_kernel.cpp+67-0

@@ -0,0 +1,67 @@
	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	+#include <pto/pto-inst.hpp>
	12	+#include <pto/common/constants.hpp>
	13	+
	14	+using namespace pto;
	15	+
	16	+template <typename T, int TRow, int TCol, int validRow, int validCol, bool isHighPrecision>
	17	+__global__ AICORE void runTPow(__gm__ T __out__ out, __gm__ T __in__ base, __gm__ T __in__ *exp)
	18	+{
	19	+ using ShapeDim5 = Shape<1, 1, 1, validRow, validCol>;
	20	+ using StrideDim5 = pto::Stride<TRow * TCol, TRow * TCol, TRow * TCol, TCol, 1>;
	21	+ using GlobalData = GlobalTensor<T, ShapeDim5, StrideDim5>;
	22	+
	23	+ GlobalData baseGlobal(base);
	24	+ GlobalData expGlobal(exp);
	25	+ GlobalData dstGlobal(out);
	26	+
	27	+ using TileData = Tile<TileType::Vec, T, TRow, TCol, BLayout::RowMajor, validRow, validCol>;
	28	+
	29	+ TileData baseTile;
	30	+ TileData expTile;
	31	+ TileData dstTile;
	32	+ TileData tmpTile;
	33	+
	34	+ TASSIGN(baseTile, 0x0);
	35	+ TASSIGN(expTile, 1 * TileData::Numel * sizeof(T));
	36	+ TASSIGN(dstTile, 2 * TileData::Numel * sizeof(T));
	37	+ TASSIGN(tmpTile, 3 * TileData::Numel * sizeof(T));
	38	+
	39	+ constexpr PowAlgorithm algo = isHighPrecision ? PowAlgorithm::HIGH_PRECISION : PowAlgorithm::DEFAULT;
	40	+ Event<Op::TLOAD, Op::TPOW> evt0 = TLOAD(baseTile, baseGlobal);
	41	+ Event<Op::TLOAD, Op::TPOW> evt1 = TLOAD(expTile, expGlobal);
	42	+ Event<Op::TPOW, Op::TSTORE_VEC> evt2 = TPOW<algo>(dstTile, baseTile, expTile, tmpTile, evt0, evt1);
	43	+ TSTORE(dstGlobal, dstTile, evt2);
	44	+}
	45	+
	46	+template <typename T, int TRow, int TCol, int validRow, int validCol, bool isHighPrecision>
	47	+void LaunchTPow(T out, T base, T exp, void stream)
	48	+{
	49	+ if constexpr (std::is_same_v<T, uint16_t>) {
	50	+ runTPow<half, TRow, TCol, validRow, validCol, isHighPrecision>
	51	+ <<<1, nullptr, stream>>>((half )(out), (half )(base), (half *)(exp));
	52	+ } else {
	53	+ runTPow<T, TRow, TCol, validRow, validCol, isHighPrecision><<<1, nullptr, stream>>>(out, base, exp);
	54	+ }
	55	+}
	56	+
	57	+template void LaunchTPow<float, 64, 64, 63, 63, false>(float out, float base, float exp, void stream);
	58	+template void LaunchTPow<uint16_t, 64, 64, 63, 63, false>(uint16_t out, uint16_t base, uint16_t exp, void stream);
	59	+template void LaunchTPow<int32_t, 64, 64, 63, 63, false>(int32_t out, int32_t base, int32_t exp, void stream);
	60	+template void LaunchTPow<int16_t, 64, 64, 63, 63, false>(int16_t out, int16_t base, int16_t exp, void stream);
	61	+template void LaunchTPow<int8_t, 64, 64, 63, 63, false>(int8_t out, int8_t base, int8_t exp, void stream);
	62	+template void LaunchTPow<uint32_t, 64, 64, 63, 63, false>(uint32_t out, uint32_t base, uint32_t exp, void stream);
	63	+template void LaunchTPow<uint8_t, 64, 64, 63, 63, false>(uint8_t out, uint8_t base, uint8_t exp, void stream);
	64	+template void LaunchTPow<float, 64, 64, 63, 63, true>(float out, float base, float exp, void stream);
	65	+template void LaunchTPow<uint16_t, 64, 64, 63, 63, true>(uint16_t out, uint16_t base, uint16_t exp, void stream);
	66	+template void LaunchTPow<float, 16, 256, 15, 231, false>(float out, float base, float exp, void stream);
	67	+template void LaunchTPow<uint16_t, 16, 512, 16, 400, true>(uint16_t out, uint16_t base, uint16_t exp, void stream);

Atests/npu/a5/src/st/testcase/tpows/CMakeLists.txt+11-0

@@ -0,0 +1,11 @@
	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	+pto_vec_st(tpows)

Atests/npu/a5/src/st/testcase/tpows/gen_data.py+79-0

@@ -0,0 +1,79 @@
	1	+#!/usr/bin/python3
	2	+# coding=utf-8
	3	+# --------------------------------------------------------------------------------
	4	+# Copyright (c) 2025 Huawei Technologies Co., Ltd.
	5	+# This program is free software, you can redistribute it and/or modify it under terms and conditions of
	6	+# CANN Open Software License Agreement Version 2.0 (the "License").
	7	+# Please refer to the License for details. You may not use this file except in compliance with the License.
	8	+# THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	9	+# INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	10	+# See LICENSE in the root of the software repository for the full text of the License.
	11	+# --------------------------------------------------------------------------------
	12	+
	13	+import os
	14	+import numpy as np
	15	+
	16	+np.random.seed(19)
	17	+
	18	+def gen_golden_data(param):
	19	+ dtype = param.dtype
	20	+ row, col = param.row, param.col
	21	+ valid_row, valid_col = param.valid_row, param.valid_col
	22	+
	23	+ if dtype in (np.int8, np.uint8, np.int16, np.uint16, np.int32, np.uint32):
	24	+ dtype_info = np.iinfo(dtype)
	25	+ base_arr = np.random.randint(1, 10, size=[row, col]).astype(dtype)
	26	+ exp_arr = np.random.randint(0, 5, size=1).astype(dtype)
	27	+ else:
	28	+ dtype_info = np.finfo(dtype)
	29	+ base_arr = np.random.uniform(0.1, 5.0, size=[row, col]).astype(dtype)
	30	+ exp_arr = np.random.uniform(0, 3.0, size=1).astype(dtype)
	31	+
	32	+ golden = np.zeros((row, col), dtype=dtype)
	33	+ golden[0:valid_row, 0:valid_col] = np.power(base_arr[0:valid_row, 0:valid_col], exp_arr[0])
	34	+
	35	+ base_arr.tofile("base.bin")
	36	+ exp_arr.tofile("exp.bin")
	37	+ golden.tofile("golden.bin")
	38	+
	39	+
	40	+class TPOWSParams:
	41	+ def __init__(self, name, dtype, row, col, valid_row, valid_col):
	42	+ self.name = name
	43	+ self.dtype = dtype
	44	+ self.row = row
	45	+ self.col = col
	46	+ self.valid_row = valid_row
	47	+ self.valid_col = valid_col
	48	+
	49	+
	50	+if __name__ == "__main__":
	51	+ # Get the absolute path of the script
	52	+ script_dir = os.path.dirname(os.path.abspath(__file__))
	53	+ testcases_dir = os.path.join(script_dir, "testcases")
	54	+
	55	+ # Ensure the testcases directory exists
	56	+ if not os.path.exists(testcases_dir):
	57	+ os.makedirs(testcases_dir)
	58	+
	59	+ case_params_list = [
	60	+ TPOWSParams("TPOWSTest.case1", np.float32, 64, 64, 63, 63),
	61	+ TPOWSParams("TPOWSTest.case2", np.float16, 64, 64, 63, 63),
	62	+ TPOWSParams("TPOWSTest.case3", np.int32, 64, 64, 63, 63),
	63	+ TPOWSParams("TPOWSTest.case4", np.int16, 64, 64, 63, 63),
	64	+ TPOWSParams("TPOWSTest.case5", np.int8, 64, 64, 63, 63),
	65	+ TPOWSParams("TPOWSTest.case6", np.uint32, 64, 64, 63, 63),
	66	+ TPOWSParams("TPOWSTest.case7", np.uint8, 64, 64, 63, 63),
	67	+ TPOWSParams("TPOWSTest.case8", np.float32, 64, 64, 63, 63),
	68	+ TPOWSParams("TPOWSTest.case9", np.float16, 64, 64, 63, 63),
	69	+ TPOWSParams("TPOWSTest.case10", np.float32, 16, 256, 15, 231),
	70	+ TPOWSParams("TPOWSTest.case11", np.float16, 16, 512, 16, 400),
	71	+ ]
	72	+
	73	+ for param in case_params_list:
	74	+ if not os.path.exists(param.name):
	75	+ os.makedirs(param.name)
	76	+ original_dir = os.getcwd()
	77	+ os.chdir(param.name)
	78	+ gen_golden_data(param)
	79	+ os.chdir(original_dir)

Atests/npu/a5/src/st/testcase/tpows/main.cpp+135-0

@@ -0,0 +1,135 @@
	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	+#include "test_common.h"
	12	+#include "acl/acl.h"
	13	+#include <gtest/gtest.h>
	14	+
	15	+using namespace std;
	16	+using namespace PtoTestCommon;
	17	+
	18	+class TPOWSTest : public testing::Test {
	19	+protected:
	20	+ void SetUp() override
	21	+ {}
	22	+ void TearDown() override
	23	+ {}
	24	+};
	25	+
	26	+std::string GetGoldenDir()
	27	+{
	28	+ const testing::TestInfo *testInfo = testing::UnitTest::GetInstance()->current_test_info();
	29	+ const std::string caseName = testInfo->name();
	30	+ std::string suiteName = testInfo->test_suite_name();
	31	+ std::string fullPath = "../" + suiteName + "." + caseName;
	32	+ return fullPath;
	33	+}
	34	+
	35	+template <typename T, int TRow, int TCol, int validRow, int validCol, bool isHighPrecision>
	36	+void LaunchTPows(T out, T base, T exp, void stream);
	37	+
	38	+template <typename T, int Row, int Col, int validRow, int validCol, bool isHighPrecision = false>
	39	+void test_tpows()
	40	+{
	41	+ size_t fileSize = Row * Col * sizeof(T);
	42	+ size_t expSize = sizeof(T);
	43	+
	44	+ aclInit(nullptr);
	45	+ aclrtSetDevice(0);
	46	+ aclrtStream stream;
	47	+ aclrtCreateStream(&stream);
	48	+
	49	+ T dstHost, baseHost, *expHost;
	50	+ T dstDevice, baseDevice, *expDevice;
	51	+
	52	+ aclrtMallocHost((void **)(&dstHost), fileSize);
	53	+ aclrtMallocHost((void **)(&baseHost), fileSize);
	54	+ aclrtMallocHost((void **)(&expHost), expSize);
	55	+ aclrtMalloc((void **)&dstDevice, fileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	56	+ aclrtMalloc((void **)&baseDevice, fileSize, ACL_MEM_MALLOC_HUGE_FIRST);
	57	+ aclrtMalloc((void **)&expDevice, expSize, ACL_MEM_MALLOC_HUGE_FIRST);
	58	+
	59	+ ReadFile(GetGoldenDir() + "/base.bin", fileSize, baseHost, fileSize);
	60	+ ReadFile(GetGoldenDir() + "/exp.bin", expSize, expHost, expSize);
	61	+ aclrtMemcpy(baseDevice, fileSize, baseHost, fileSize, ACL_MEMCPY_HOST_TO_DEVICE);
	62	+ aclrtMemcpy(expDevice, expSize, expHost, expSize, ACL_MEMCPY_HOST_TO_DEVICE);
	63	+ LaunchTPows<T, Row, Col, validRow, validCol, isHighPrecision>(dstDevice, baseDevice, expDevice, stream);
	64	+
	65	+ aclrtSynchronizeStream(stream);
	66	+ aclrtMemcpy(dstHost, fileSize, dstDevice, fileSize, ACL_MEMCPY_DEVICE_TO_HOST);
	67	+
	68	+ WriteFile(GetGoldenDir() + "/output.bin", dstHost, fileSize);
	69	+
	70	+ aclrtFree(dstDevice);
	71	+ aclrtFree(baseDevice);
	72	+ aclrtFree(expDevice);
	73	+
	74	+ aclrtFreeHost(dstHost);
	75	+ aclrtFreeHost(baseHost);
	76	+ aclrtFreeHost(expHost);
	77	+ aclrtDestroyStream(stream);
	78	+ aclrtResetDevice(0);
	79	+ aclFinalize();
	80	+
	81	+ std::vector<T> golden(fileSize / sizeof(T));
	82	+ std::vector<T> devFinal(fileSize / sizeof(T));
	83	+ ReadFile(GetGoldenDir() + "/golden.bin", fileSize, golden.data(), fileSize);
	84	+ ReadFile(GetGoldenDir() + "/output.bin", fileSize, devFinal.data(), fileSize);
	85	+
	86	+ constexpr float eps = std::is_same_v<T, float> ? 0.0005f : 0.00005f;
	87	+ bool ret = ResultCmp<T>(golden, devFinal, eps);
	88	+
	89	+ EXPECT_TRUE(ret);
	90	+}
	91	+
	92	+TEST_F(TPOWSTest, case1)
	93	+{
	94	+ test_tpows<float, 64, 64, 63, 63>();
	95	+}
	96	+TEST_F(TPOWSTest, case2)
	97	+{
	98	+ test_tpows<aclFloat16, 64, 64, 63, 63>(); // typedef uint16_t aclFloat16
	99	+}
	100	+TEST_F(TPOWSTest, case3)
	101	+{
	102	+ test_tpows<int32_t, 64, 64, 63, 63>();
	103	+}
	104	+TEST_F(TPOWSTest, case4)
	105	+{
	106	+ test_tpows<int16_t, 64, 64, 63, 63>();
	107	+}
	108	+TEST_F(TPOWSTest, case5)
	109	+{
	110	+ test_tpows<int8_t, 64, 64, 63, 63>();
	111	+}
	112	+TEST_F(TPOWSTest, case6)
	113	+{
	114	+ test_tpows<uint32_t, 64, 64, 63, 63>();
	115	+}
	116	+TEST_F(TPOWSTest, case7)
	117	+{
	118	+ test_tpows<uint8_t, 64, 64, 63, 63>();
	119	+}
	120	+TEST_F(TPOWSTest, case8)
	121	+{
	122	+ test_tpows<float, 64, 64, 63, 63, true>();
	123	+}
	124	+TEST_F(TPOWSTest, case9)
	125	+{
	126	+ test_tpows<aclFloat16, 64, 64, 63, 63, true>();
	127	+}
	128	+TEST_F(TPOWSTest, case10)
	129	+{
	130	+ test_tpows<float, 16, 256, 15, 231>();
	131	+}
	132	+TEST_F(TPOWSTest, case11)
	133	+{
	134	+ test_tpows<aclFloat16, 16, 512, 16, 400, true>();
	135	+}

Atests/npu/a5/src/st/testcase/tpows/tpows_kernel.cpp+65-0

@@ -0,0 +1,65 @@
	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	+#include <pto/pto-inst.hpp>
	12	+#include <pto/common/constants.hpp>
	13	+
	14	+using namespace pto;
	15	+
	16	+template <typename T, int TRow, int TCol, int validRow, int validCol, bool isHighPrecision>
	17	+__global__ AICORE void runTPows(__gm__ T out, __gm__ T base, __gm__ T *exp)
	18	+{
	19	+ T expVal = *exp;
	20	+ constexpr PowAlgorithm algo = isHighPrecision ? PowAlgorithm::HIGH_PRECISION : PowAlgorithm::DEFAULT;
	21	+
	22	+ using ShapeDim5 = Shape<1, 1, 1, validRow, validCol>;
	23	+ using StrideDim5 = pto::Stride<TRow * TCol, TRow * TCol, TRow * TCol, TCol, 1>;
	24	+ using GlobalData = GlobalTensor<T, ShapeDim5, StrideDim5>;
	25	+
	26	+ GlobalData baseGlobal(base);
	27	+ GlobalData dstGlobal(out);
	28	+
	29	+ using TileData = Tile<TileType::Vec, T, TRow, TCol, BLayout::RowMajor, validRow, validCol>;
	30	+
	31	+ TileData baseTile;
	32	+ TileData dstTile;
	33	+ TileData tmpTile;
	34	+
	35	+ TASSIGN<0x0>(baseTile);
	36	+ TASSIGN<1 * TileData::Numel * sizeof(T)>(dstTile);
	37	+ TASSIGN<2 * TileData::Numel * sizeof(T)>(tmpTile);
	38	+
	39	+ Event<Op::TLOAD, Op::TPOW> evt0 = TLOAD(baseTile, baseGlobal);
	40	+ Event<Op::TPOW, Op::TSTORE_VEC> evt1 = TPOWS<algo>(dstTile, baseTile, expVal, tmpTile, evt0);
	41	+ TSTORE(dstGlobal, dstTile, evt1);
	42	+}
	43	+
	44	+template <typename T, int TRow, int TCol, int validRow, int validCol, bool isHighPrecision>
	45	+void LaunchTPows(T out, T base, T exp, void stream)
	46	+{
	47	+ if constexpr (std::is_same_v<T, uint16_t>) {
	48	+ runTPows<half, TRow, TCol, validRow, validCol, isHighPrecision>
	49	+ <<<1, nullptr, stream>>>((half )(out), (half )(base), (half *)(exp));
	50	+ } else {
	51	+ runTPows<T, TRow, TCol, validRow, validCol, isHighPrecision><<<1, nullptr, stream>>>(out, base, exp);
	52	+ }
	53	+}
	54	+
	55	+template void LaunchTPows<float, 64, 64, 63, 63, false>(float out, float base, float exp, void stream);
	56	+template void LaunchTPows<uint16_t, 64, 64, 63, 63, false>(uint16_t out, uint16_t base, uint16_t exp, void stream);
	57	+template void LaunchTPows<int32_t, 64, 64, 63, 63, false>(int32_t out, int32_t base, int32_t exp, void stream);
	58	+template void LaunchTPows<int16_t, 64, 64, 63, 63, false>(int16_t out, int16_t base, int16_t exp, void stream);
	59	+template void LaunchTPows<int8_t, 64, 64, 63, 63, false>(int8_t out, int8_t base, int8_t exp, void stream);
	60	+template void LaunchTPows<uint32_t, 64, 64, 63, 63, false>(uint32_t out, uint32_t base, uint32_t exp, void stream);
	61	+template void LaunchTPows<uint8_t, 64, 64, 63, 63, false>(uint8_t out, uint8_t base, uint8_t exp, void stream);
	62	+template void LaunchTPows<float, 64, 64, 63, 63, true>(float out, float base, float exp, void stream);
	63	+template void LaunchTPows<uint16_t, 64, 64, 63, 63, true>(uint16_t out, uint16_t base, uint16_t exp, void stream);
	64	+template void LaunchTPows<float, 16, 256, 15, 231, false>(float out, float base, float exp, void stream);
	65	+template void LaunchTPows<uint16_t, 16, 512, 16, 400, true>(uint16_t out, uint16_t base, uint16_t exp, void stream);

Mtests/npu/kirin9030/src/st/testcase/tfillpad/gen_data.py+79-18

	@@ -12,31 +12,92 @@
12		12
13	import os	13	import os
14	import numpy as np	14	import numpy as np
		15	+np.random.seed(42)
		16	+
		17	+def get_pad_value(dtype, pad_val):
		18	+ if pad_val == 'PADMAX':
		19	+ if dtype == np.float32 or dtype == np.float16:
		20	+ return np.inf
		21	+ else:
		22	+ return np.iinfo(dtype).max
		23	+ elif pad_val == 'PADMIN':
		24	+ if dtype == np.float32 or dtype == np.float16:
		25	+ return -np.inf
		26	+ else:
		27	+ return np.iinfo(dtype).min
		28	+ else:
		29	+ return 0
		30	+
		31	+
		32	+def gen_golden_data(param):
		33	+ dtype = param.dtype
		34	+ src_rows, src_cols = param.src_rows, param.src_cols
		35	+ dst_rows, dst_cols = param.dst_rows, param.dst_cols
		36	+ pad_val = param.pad_val
		37	+
		38	+ pad_value = get_pad_value(dtype, pad_val)
		39	+
		40	+ in_arr = np.random.uniform(low=-8, high=8, size=(src_rows, src_cols)).astype(dtype)
		41	+ gold_arr = np.full((dst_rows, dst_cols), pad_value, dtype=dtype)
		42	+ gold_arr[:src_rows, :src_cols] = in_arr
		43	+
		44	+ in_arr.tofile("input.bin")
		45	+ gold_arr.tofile("golden.bin")
		46	+
		47	+
		48	+class TFILLPADParams:
		49	+ def __init__(self, dtype, src_rows, src_cols, dst_rows, dst_cols, pad_val):
		50	+ self.dtype = dtype
		51	+ self.src_rows = src_rows
		52	+ self.src_cols = src_cols
		53	+ self.dst_rows = dst_rows
		54	+ self.dst_cols = dst_cols
		55	+ self.pad_val = pad_val
		56	+
		57	+
		58	+def get_params_for_test_key(test_key):
		59	+ params_map = {
		60	+ 1: TFILLPADParams(np.float32, 64, 127, 64, 128, 'PADMAX'),
		61	+ 2: TFILLPADParams(np.float32, 64, 127, 64, 144, 'PADMAX'),
		62	+ 3: TFILLPADParams(np.float32, 64, 127, 64, 160, 'PADMAX'),
		63	+ 4: TFILLPADParams(np.float32, 260, 7, 260, 16, 'PADMAX'),
		64	+ 5: TFILLPADParams(np.float32, 260, 7, 260, 16, 'PADMAX'),
		65	+ 6: TFILLPADParams(np.uint16, 260, 7, 260, 32, 'PADMAX'),
		66	+ 7: TFILLPADParams(np.int8, 260, 7, 260, 64, 'PADMAX'),
		67	+ 8: TFILLPADParams(np.uint16, 259, 7, 260, 32, 'PADMAX'),
		68	+ 9: TFILLPADParams(np.int8, 259, 7, 260, 64, 'PADMAX'),
		69	+ 10: TFILLPADParams(np.int16, 260, 7, 260, 32, 'PADMIN'),
		70	+ 11: TFILLPADParams(np.int32, 260, 7, 260, 32, 'PADMIN'),
		71	+ }
		72	+ return params_map.get(test_key)
		73	+
15		74
16	if __name__ == "__main__":	75	if __name__ == "__main__":
17	- # 用例名称	76	+ script_dir = os.path.dirname(os.path.abspath(__file__))
		77	+ testcases_dir = os.path.join(script_dir, "testcases")
		78	+
		79	+ if not os.path.exists(testcases_dir):
		80	+ os.makedirs(testcases_dir)
		81	+
18	case_name_list = [	82	case_name_list = [
19	- "TFILLPADTest.case_float_GT_128_127_VT_128_128_BLK1_PADMAX_PADMAX",	83	+ "TFILLPADTest.case_float_GT_64_127_VT_64_128_BLK1_PADMAX",
20	- "TFILLPADTest.case_float_GT_128_127_VT_128_160_BLK1_PADMAX_PADMAX",	84	+ "TFILLPADTest.case_float_GT_64_127_VT_64_144_BLK1_PADMAX",
21	- "TFILLPADTest.case_float_GT_128_127_VT_128_160_BLK1_PADMIN_PADMAX",	85	+ "TFILLPADTest.case_float_GT_64_127_VT_64_160_BLK1_PADMAX",
22	- "TFILLPADTest.case_float_GT_260_7_VT_260_16_BLK1_PADMIN_PADMAX",	86	+ "TFILLPADTest.case_float_GT_260_7_VT_260_16_BLK1_PADMAX",
23	- "TFILLPADTest.case_float_GT_260_7_VT_260_16_BLK1_PADMIN_PADMAX_INPLACE",	87	+ "TFILLPADTest.case_float_GT_260_7_VT_260_16_BLK1_PADMAX_INPLACE",
24	- "TFILLPADTest.case_u16_GT_260_7_VT_260_32_BLK1_PADMIN_PADMAX",	88	+ "TFILLPADTest.case_u16_GT_260_7_VT_260_32_BLK1_PADMAX",
25	- "TFILLPADTest.case_s8_GT_260_7_VT_260_64_BLK1_PADMIN_PADMAX",	89	+ "TFILLPADTest.case_s8_GT_260_7_VT_260_64_BLK1_PADMAX",
26	- "TFILLPADTest.case_u16_GT_259_7_VT_260_32_BLK1_PADMIN_PADMAX_EXPAND",	90	+ "TFILLPADTest.case_u16_GT_259_7_VT_260_32_BLK1_PADMAX_EXPAND",
27	- "TFILLPADTest.case_s8_GT_259_7_VT_260_64_BLK1_PADMIN_PADMAX_EXPAND",	91	+ "TFILLPADTest.case_s8_GT_259_7_VT_260_64_BLK1_PADMAX_EXPAND",
28	- "TFILLPADTest.case_s16_GT_260_7_VT_260_32_BLK1_PADMIN_PADMIN",	92	+ "TFILLPADTest.case_s16_GT_260_7_VT_260_32_BLK1_PADMIN",
29	- "TFILLPADTest.case_s32_GT_260_7_VT_260_32_BLK1_PADMIN_PADMIN",	93	+ "TFILLPADTest.case_s32_GT_260_7_VT_260_32_BLK1_PADMIN",
30	- "TFILLPADTest.case_float_GT_128_64_VT_128_128_PADCUSTOM_NEG1",
31	- "TFILLPADTest.case_float_GT_128_127_VT_128_160_BLK1_PADCUSTOM_NEG1_PADCUSTOM_NEG1",
32	]	94	]
33		95
34	- for i, case_name in enumerate(case_name_list):	96	+ for test_key, case_name in enumerate(case_name_list, start=1):
35	if not os.path.exists(case_name):	97	if not os.path.exists(case_name):
36	os.makedirs(case_name)	98	os.makedirs(case_name)
37	original_dir = os.getcwd()	99	original_dir = os.getcwd()
38	os.chdir(case_name)	100	os.chdir(case_name)
39	- pass	101	+ param = get_params_for_test_key(test_key)
		102	+ gen_golden_data(param)
40	os.chdir(original_dir)	103	os.chdir(original_dir)
41	-
42	- pass

Mtests/npu/kirin9030/src/st/testcase/tfillpad/main.cpp+45-83

	@@ -8,19 +8,14 @@ INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A
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	+#include <type_traits>
		12	+#include <gtest/gtest.h>
11	#include "test_common.h"	13	#include "test_common.h"
12	#include "acl/acl.h"	14	#include "acl/acl.h"
13	-#include <gtest/gtest.h>
14		15
15	using namespace std;	16	using namespace std;
16	using namespace PtoTestCommon;	17	using namespace PtoTestCommon;
17		18
18	-template <int32_t testKey>
19	-void launchTFILLPAD(uint8_t out, uint8_t src, void *stream);
20	-
21	-template <int32_t testKey>
22	-int get_input_golden(uint8_t input, uint8_t golden);
23	-
24	class TFILLPADTest : public testing::Test {	19	class TFILLPADTest : public testing::Test {
25	protected:	20	protected:
26	void SetUp() override	21	void SetUp() override
	@@ -38,145 +33,112 @@ std::string GetGoldenDir()
38	return fullPath;	33	return fullPath;
39	}	34	}
40		35
41	-template <typename T>	36	+template <int32_t testKey>
42	-constexpr T getGoldenZero()	37	+void launchTFILLPAD(uint8_t out, uint8_t src, void *stream);
43	-{
44	- return T{0};
45	-}
46		38
47	-template <int32_t testKey, typename T, int32_t kBlock>	39	+template <typename T, int32_t srcRows, int32_t srcCols, int32_t dstRows, int32_t dstCols, int32_t testKey>
48	-void tfillpad_test()	40	+void test_tfillpad()
49	{	41	{
50	- uint32_t M = 1024;	42	+ size_t fileSizeSrc = srcRows * srcCols * sizeof(T);
51	- uint32_t N = 1024;	43	+ size_t fileSizeDst = dstRows * dstCols * sizeof(T);
52		44
53	aclInit(nullptr);	45	aclInit(nullptr);
54	aclrtSetDevice(0);	46	aclrtSetDevice(0);
55	aclrtStream stream;	47	aclrtStream stream;
56	aclrtCreateStream(&stream);	48	aclrtCreateStream(&stream);
57		49
58	- int in_byteSize = M * N * sizeof(float);	50	+ void dstHost, srcHost;
59	- int out_byteSize = M * N * sizeof(float);
60	-
61	- void dstHost, srcHost, *goldHost;
62	void dstDevice, srcDevice;	51	void dstDevice, srcDevice;
63		52
64	- aclrtMallocHost((void **)(&srcHost), in_byteSize);	53	+ aclrtMallocHost((void **)(&srcHost), fileSizeSrc);
65	- aclrtMallocHost((void **)(&dstHost), out_byteSize);	54	+ aclrtMallocHost((void **)(&dstHost), fileSizeDst);
66	- aclrtMallocHost((void **)(&goldHost), out_byteSize);
67		55
68	- aclrtMalloc((void **)&dstDevice, in_byteSize, ACL_MEM_MALLOC_HUGE_FIRST);	56	+ aclrtMalloc((void **)&dstDevice, fileSizeDst, ACL_MEM_MALLOC_HUGE_FIRST);
69	- aclrtMalloc((void **)&srcDevice, out_byteSize, ACL_MEM_MALLOC_HUGE_FIRST);	57	+ aclrtMalloc((void **)&srcDevice, fileSizeSrc, ACL_MEM_MALLOC_HUGE_FIRST);
70		58
71	- int actual_out_byteSize = 0;	59	+ ReadFile(GetGoldenDir() + "/input.bin", fileSizeSrc, srcHost, fileSizeSrc);
72	- actual_out_byteSize = get_input_golden<testKey>((uint8_t )srcHost, (uint8_t )goldHost);	60	+ aclrtMemset(dstHost, fileSizeDst, 0, fileSizeDst);
73	- cout << "Golden size:" << actual_out_byteSize << " B" << endl;
74	- std::fill((uint8_t )dstHost, ((uint8_t )(dstHost)) + out_byteSize, 0);
75		61
76	- aclrtMemcpy(srcDevice, in_byteSize, srcHost, in_byteSize, ACL_MEMCPY_HOST_TO_DEVICE);	62	+ aclrtMemcpy(dstDevice, fileSizeDst, dstHost, fileSizeDst, ACL_MEMCPY_HOST_TO_DEVICE);
77	- aclrtMemcpy(dstDevice, out_byteSize, dstHost, out_byteSize, ACL_MEMCPY_HOST_TO_DEVICE);	63	+ aclrtMemcpy(srcDevice, fileSizeSrc, srcHost, fileSizeSrc, ACL_MEMCPY_HOST_TO_DEVICE);
78		64
79	launchTFILLPAD<testKey>((uint8_t )dstDevice, (uint8_t )srcDevice, stream);	65	launchTFILLPAD<testKey>((uint8_t )dstDevice, (uint8_t )srcDevice, stream);
80		66
81	aclrtSynchronizeStream(stream);	67	aclrtSynchronizeStream(stream);
82	- aclrtMemcpy(dstHost, out_byteSize, dstDevice, out_byteSize, ACL_MEMCPY_DEVICE_TO_HOST);	68	+ aclrtMemcpy(dstHost, fileSizeDst, dstDevice, fileSizeDst, ACL_MEMCPY_DEVICE_TO_HOST);
83		69
84	- std::ofstream inFile(GetGoldenDir() + "/input.bin", std::ios::binary \| std::ios::out);	70	+ WriteFile(GetGoldenDir() + "/output.bin", dstHost, fileSizeDst);
85	- std::ofstream outFile(GetGoldenDir() + "/output.bin", std::ios::binary \| std::ios::out);
86	- std::ofstream goldFile(GetGoldenDir() + "/golden.bin", std::ios::binary \| std::ios::out);
87	- inFile.write((const char *)srcHost, actual_out_byteSize);
88	- outFile.write((const char *)dstHost, actual_out_byteSize);
89	- goldFile.write((const char *)goldHost, actual_out_byteSize);
90	- inFile.close();
91	- outFile.close();
92	- goldFile.close();
93		71
94	aclrtFree(dstDevice);	72	aclrtFree(dstDevice);
95	aclrtFree(srcDevice);	73	aclrtFree(srcDevice);
96		74
97	aclrtFreeHost(dstHost);	75	aclrtFreeHost(dstHost);
98	aclrtFreeHost(srcHost);	76	aclrtFreeHost(srcHost);
99	- aclrtFreeHost(goldHost);
100	-
101	aclrtDestroyStream(stream);	77	aclrtDestroyStream(stream);
102	aclrtResetDevice(0);	78	aclrtResetDevice(0);
103	aclFinalize();	79	aclFinalize();
104		80
105	- int elements = actual_out_byteSize / sizeof(T);	81	+ std::vector<T> golden(dstRows * dstCols);
106	-	82	+ std::vector<T> devFinal(dstRows * dstCols);
107	- auto zero = getGoldenZero<T>();	83	+ ReadFile(GetGoldenDir() + "/golden.bin", fileSizeDst, golden.data(), fileSizeDst);
108	- using CT = decltype(zero);	84	+ ReadFile(GetGoldenDir() + "/output.bin", fileSizeDst, devFinal.data(), fileSizeDst);
109	- std::vector<CT> golden(elements);
110	- std::vector<CT> devFinal(elements);
111	- size_t oFileSize = actual_out_byteSize;
112	- ReadFile(GetGoldenDir() + "/golden.bin", oFileSize, golden.data(), oFileSize);
113	- ReadFile(GetGoldenDir() + "/output.bin", oFileSize, devFinal.data(), oFileSize);
114		85
115	bool ret = ResultCmp(golden, devFinal, 0);	86	bool ret = ResultCmp(golden, devFinal, 0);
116		87
117	EXPECT_TRUE(ret);	88	EXPECT_TRUE(ret);
118	}	89	}
119		90
120	-TEST_F(TFILLPADTest, case_float_GT_128_127_VT_128_128_BLK1_PADMAX_PADMAX)	91	+TEST_F(TFILLPADTest, case_float_GT_64_127_VT_64_128_BLK1_PADMAX)
121	{	92	{
122	- tfillpad_test<1, float, 1>();	93	+ test_tfillpad<float, 64, 127, 64, 128, 1>();
123	}	94	}
124		95
125	-TEST_F(TFILLPADTest, case_float_GT_128_127_VT_128_160_BLK1_PADMAX_PADMAX)	96	+TEST_F(TFILLPADTest, case_float_GT_64_127_VT_64_144_BLK1_PADMAX)
126	{	97	{
127	- tfillpad_test<2, float, 1>();	98	+ test_tfillpad<float, 64, 127, 64, 144, 2>();
128	}	99	}
129		100
130	-TEST_F(TFILLPADTest, case_float_GT_128_127_VT_128_160_BLK1_PADMIN_PADMAX)	101	+TEST_F(TFILLPADTest, case_float_GT_64_127_VT_64_160_BLK1_PADMAX)
131	{	102	{
132	- tfillpad_test<3, float, 1>();	103	+ test_tfillpad<float, 64, 127, 64, 160, 3>();
133	}	104	}
134		105
135	-TEST_F(TFILLPADTest, case_float_GT_260_7_VT_260_16_BLK1_PADMIN_PADMAX)	106	+TEST_F(TFILLPADTest, case_float_GT_260_7_VT_260_16_BLK1_PADMAX)
136	{	107	{
137	- tfillpad_test<4, float, 1>();	108	+ test_tfillpad<float, 260, 7, 260, 16, 4>();
138	}	109	}
139		110
140	-TEST_F(TFILLPADTest, case_float_GT_260_7_VT_260_16_BLK1_PADMIN_PADMAX_INPLACE)	111	+TEST_F(TFILLPADTest, case_float_GT_260_7_VT_260_16_BLK1_PADMAX_INPLACE)
141	{	112	{
142	- tfillpad_test<5, float, 1>();	113	+ test_tfillpad<float, 260, 7, 260, 16, 5>();
143	}	114	}
144		115
145	-TEST_F(TFILLPADTest, case_u16_GT_260_7_VT_260_32_BLK1_PADMIN_PADMAX)	116	+TEST_F(TFILLPADTest, case_u16_GT_260_7_VT_260_32_BLK1_PADMAX)
146	{	117	{
147	- tfillpad_test<6, uint16_t, 1>();	118	+ test_tfillpad<uint16_t, 260, 7, 260, 32, 6>();
148	}	119	}
149		120
150	-TEST_F(TFILLPADTest, case_s8_GT_260_7_VT_260_64_BLK1_PADMIN_PADMAX)	121	+TEST_F(TFILLPADTest, case_s8_GT_260_7_VT_260_64_BLK1_PADMAX)
151	{	122	{
152	- tfillpad_test<7, int8_t, 1>();	123	+ test_tfillpad<int8_t, 260, 7, 260, 64, 7>();
153	}	124	}
154		125
155	-TEST_F(TFILLPADTest, case_u16_GT_259_7_VT_260_32_BLK1_PADMIN_PADMAX_EXPAND)	126	+TEST_F(TFILLPADTest, case_u16_GT_259_7_VT_260_32_BLK1_PADMAX_EXPAND)
156	{	127	{
157	- tfillpad_test<8, uint16_t, 1>();	128	+ test_tfillpad<uint16_t, 259, 7, 260, 32, 8>();
158	}	129	}
159		130
160	-TEST_F(TFILLPADTest, case_s8_GT_259_7_VT_260_64_BLK1_PADMIN_PADMAX_EXPAND)	131	+TEST_F(TFILLPADTest, case_s8_GT_259_7_VT_260_64_BLK1_PADMAX_EXPAND)
161	{	132	{
162	- tfillpad_test<9, int8_t, 1>();	133	+ test_tfillpad<int8_t, 259, 7, 260, 64, 9>();
163	}	134	}
164		135
165	-TEST_F(TFILLPADTest, case_s16_GT_260_7_VT_260_32_BLK1_PADMIN_PADMIN)	136	+TEST_F(TFILLPADTest, case_s16_GT_260_7_VT_260_32_BLK1_PADMIN)
166	{	137	{
167	- tfillpad_test<10, int16_t, 1>();	138	+ test_tfillpad<int16_t, 260, 7, 260, 32, 10>();
168	}	139	}
169		140
170	-TEST_F(TFILLPADTest, case_s32_GT_260_7_VT_260_32_BLK1_PADMIN_PADMIN)	141	+TEST_F(TFILLPADTest, case_s32_GT_260_7_VT_260_32_BLK1_PADMIN)
171	{	142	{
172	- tfillpad_test<11, int32_t, 1>();	143	+ test_tfillpad<int32_t, 260, 7, 260, 32, 11>();
173	-}
174	-TEST_F(TFILLPADTest, case_float_GT_128_64_VT_128_128_PADCUSTOM_NEG1)
175	-{
176	- tfillpad_test<12, float, 1>();
177	-}
178	-
179	-TEST_F(TFILLPADTest, case_float_GT_128_127_VT_128_160_BLK1_PADCUSTOM_NEG1_PADCUSTOM_NEG1)
180	-{
181	- tfillpad_test<13, float, 1>();
182	}	144	}

Mtests/npu/kirin9030/src/st/testcase/tfillpad/tfillpad_kernel.cpp+75-381

	@@ -16,453 +16,147 @@ See LICENSE in the root of the software repository for the full text of the Lice
16	using namespace std;	16	using namespace std;
17	using namespace pto;	17	using namespace pto;
18		18
19	-// Custom pad value for test case 12
20	-// -1.0f has bit pattern 0xBF800000
21	-constexpr PadValue PadCustomNeg1 = PadValueCustom(-1.0f);
22	-
23	-// case shape is static, but testing would do dynamic or static test
24	-template <int shape0, int shape1, int shape2, int shape3, int shape4>
25	-AICORE __inline__ auto getOptDynShape(int gShape0, int gShape1, int gShape2, int gShape3, int gShape4)
26	-{
27	- if constexpr (shape0 == 1) {
28	- using DynShapeDim5 = Shape<1, -1, -1, -1, -1>;
29	- DynShapeDim5 dynShape(gShape1, gShape2, gShape3, gShape4);
30	- return dynShape;
31	- } else if constexpr (shape0 == 1 && shape1 == 1) {
32	- using DynShapeDim5 = Shape<1, 1, -1, -1, -1>;
33	- DynShapeDim5 dynShape(gShape2, gShape3, gShape4);
34	- return dynShape;
35	- } else if constexpr (shape0 == 1 && shape1 == 1 && shape2 == 1) {
36	- using DynShapeDim5 = Shape<1, 1, 1, -1, -1>;
37	- DynShapeDim5 dynShape(gShape3, gShape4);
38	- return dynShape;
39	- } else if constexpr (shape0 == 1 && shape1 == 1 && shape2 == 1 && shape3 == 1) {
40	- using DynShapeDim5 = Shape<1, 1, 1, 1, -1>;
41	- DynShapeDim5 dynShape(gShape4);
42	- return dynShape;
43	- } else {
44	- using DynShapeDim5 = Shape<-1, -1, -1, -1, -1>;
45	- DynShapeDim5 dynShape(gShape0, gShape1, gShape2, gShape3, gShape4);
46	- return dynShape;
47	- }
48	-}
49	-
50	-// case shape is static, but testing would do dynamic or static test
51	-template <typename T, int shape0, int shape1, int shape2, int shape3, int shape4, int tRows, int tCols, BLayout major,
52	- int dyn>
53	-AICORE __inline__ auto getGlobalTensor(__gm__ T *addr, int gShape0, int gShape1, int gShape2, int gShape3, int gShape4)
54	-{
55	- if constexpr (dyn) {
56	- int stride0 = gShape1 * gShape2 * shape3 * shape4;
57	- int stride1 = gShape2 * shape3 * shape4;
58	- int stride2 = shape3 * shape4;
59	-
60	- using DynStrideDim5 = pto::Stride<-1, -1, -1, -1, -1>;
61	- auto dynShape =
62	- getOptDynShape<shape0, shape1, shape2, shape3, shape4>(gShape0, gShape1, gShape2, gShape3, gShape4);
63	- using GlobalData = GlobalTensor<T, decltype(dynShape), DynStrideDim5>;
64	-
65	- if constexpr (major == BLayout::RowMajor) {
66	- GlobalData srcGlobal(addr, dynShape, DynStrideDim5(stride0, stride1, stride2, shape4, 1));
67	- return srcGlobal;
68	- } else {
69	- GlobalData srcGlobal(addr, dynShape, DynStrideDim5(stride0, stride1, stride2, 1, shape4));
70	- return srcGlobal;
71	- }
72	- } else // static
73	- {
74	- constexpr int stride0 = shape1 * shape2 * shape3 * shape4;
75	- constexpr int stride1 = shape2 * shape3 * shape4;
76	- constexpr int stride2 = shape3 * shape4;
77	- using StaticShapeDim5 = Shape<shape0, shape1, shape2, tRows, tCols>;
78	-
79	- if constexpr (major == BLayout::RowMajor) {
80	- using StaticStrideDim5 = pto::Stride<stride0, stride1, stride2, shape4, 1>;
81	- using GlobalData = GlobalTensor<T, StaticShapeDim5, StaticStrideDim5>;
82	- GlobalData srcGlobal(addr);
83	- return srcGlobal;
84	- } else {
85	- using StaticStrideDim5 = pto::Stride<stride0, stride1, stride2, 1, shape4>;
86	- using GlobalData = GlobalTensor<T, StaticShapeDim5, StaticStrideDim5>;
87	- GlobalData srcGlobal(addr);
88	- return srcGlobal;
89	- }
90	- }
91	-}
92	-
93	-inline AICORE uint64_t get_syscnt() // dont use get_sys_cnt(), need volatile for profiling
94	-{
95	- uint64_t syscnt;
96	- asm volatile("MOV %0, SYS_CNT\n" : "+l"(syscnt));
97	- return syscnt;
98	-}
99	-
100	#define type_32_aligned(T) (32 / sizeof(T))	19	#define type_32_aligned(T) (32 / sizeof(T))
101	-#define align_to_32B(x, T) ((((x) + type_32_aligned(T) - 1) / type_32_aligned(T)) * (type_32_aligned(T)));	20	+#define align_to_32B(x, T) ((((x) + type_32_aligned(T) - 1) / type_32_aligned(T)) * (type_32_aligned(T)))
102		21
103	-template <typename T, int shape0, int shape1, int shape2, int shape3, int shape4, int kTRows_, int kTCols_, int dyn_,	22	+template <typename T, int srcRows, int srcCols, int dstRows, int dstCols, PadValue LoadPadVal_ = PadValue::Null,
104	- PadValue LoadPadVal_ = PadValue::Null, PadValue FillPadVal_ = PadValue::Null, bool inplace = false,	23	+ PadValue FillPadVal_ = PadValue::Null, bool inplace = false, bool expand = false>
105	- bool expand = false>	24	+AICORE void runTFILLPAD(__gm__ T out, __gm__ T src)
106	-AICORE void runTFILLPAD(__gm__ T out, __gm__ T src, int gShape0, int gShape1, int gShape2, int gRows, int gCols)
107	{	25	{
108	-#ifndef __PTO_AUTO__	26	+ constexpr int srcTileCols = expand ? align_to_32B(srcCols, T) : dstCols;
109	- // Avoid stack dcache miss
110	- {
111	-#define INIT_STACK 8192
112	- uint64_t stack[INIT_STACK / sizeof(uint64_t)]; // 8KB
113	- volatile uint64_t *pStack = stack;
114	- for (int i = 0; i < INIT_STACK; i += 64 / sizeof(uint64_t)) // cacheline is 64B
115	- {
116	- *(pStack++) = 0;
117	- }
118	- dsb(DSB_ALL);
119	- }
120		27
121	- // Avoid icache miss in profiling: preload 4KB icache and wait	28	+ using SrcShape = Shape<1, 1, 1, srcRows, srcCols>;
122	- uint64_t pc;	29	+ using DstShape = Shape<1, 1, 1, dstRows, dstCols>;
123	- asm volatile("MOV %0, PC\n" : "+l"(pc));	30	+ using SrcStride = pto::Stride<srcRows * srcCols, srcRows * srcCols, srcRows * srcCols, srcCols, 1>;
124	- preload((void *)pc, 2);	31	+ using DstStride = pto::Stride<dstRows * dstCols, dstRows * dstCols, dstRows * dstCols, dstCols, 1>;
125	- while (get_icache_prl_st()) {	32	+ using SrcGlobal = GlobalTensor<T, SrcShape, SrcStride>;
126	-#if defined(__DAV_C220_CUBE__) \|\| defined(__DAV_C220_VEC__)	33	+ using DstGlobal = GlobalTensor<T, DstShape, DstStride>;
127	- // seems to compile for a2a3; will crash in HiIPUJumpOpt pass for A5
128	- asm("nop");
129	-#endif
130	- }
131	-#endif
132		34
133	- __ubuf__ T *ubaddr0 = 0x0;	35	+ SrcGlobal srcGlobal(src);
134	- __ubuf__ T ubaddr1 = (__ubuf__ T )0x18000;	36	+ DstGlobal dstGlobal(out);
135	- if (inplace)
136	- ubaddr1 = ubaddr0;
137		37
138	- constexpr int shape4_aligned = align_to_32B(shape4, T);	38	+ using SrcTile = Tile<TileType::Vec, T, srcRows, srcTileCols, BLayout::RowMajor, srcRows, srcCols, SLayout::NoneBox,
139	- constexpr int kGTRows = kTRows_ / shape0 / shape1 / shape2; // Dst Tile Rows, merged all shape0shape1shape2 row	39	+ 512, LoadPadVal_>;
140	- int srcOffset = (block_idx) * (shape3 / block_num) * shape4;	40	+ using DstTile = Tile<TileType::Vec, T, dstRows, dstCols, BLayout::RowMajor, dstRows, dstCols, SLayout::NoneBox, 512,
141	- auto srcGlobal =	41	+ FillPadVal_>;
142	- getGlobalTensor<T, shape0, shape1, shape2, shape3, shape4, kGTRows, shape4, BLayout::RowMajor, dyn_>(	42	+ SrcTile srcTile;
143	- src + srcOffset, gShape0, gShape1, gShape2, kGTRows, shape4);	43	+ DstTile dstTile;
144	- int dstOffset = (block_idx) * (shape3 / block_num) * kTCols_;
145	- auto dstGlobal =
146	- getGlobalTensor<T, shape0, shape1, shape2, shape3, kTCols_, kGTRows, kTCols_, BLayout::RowMajor, 0>(
147	- out + dstOffset, gShape0, gShape1, gShape2, kGTRows, kTCols_); // dst TStore GlobalTensor just use static
148		44
149	- volatile uint64_t t0, t1, t2;	45	+ TASSIGN<0x0>(srcTile);
150	- constexpr PadValue PadCustomNeg1_Test = PadValueCustom(-1.0f); // Test device usage	46	+ TASSIGN<inplace ? 0x0 : SrcTile::Numel * sizeof(T)>(dstTile);
151	- static_assert(PadCustomNeg1_Test == static_cast<PadValue>(0x00000001BF800000ULL),
152	- "PadValueCustom float device test");
153	- constexpr PadValue PadCustomNeg1_Half_Test = PadValueCustom((half)-1.0); // fp16 using half type
154	- static_assert(PadCustomNeg1_Half_Test == static_cast<PadValue>(0x000000010000BC00ULL),
155	- "PadValueCustom16 fp16 device test");
156	- constexpr PadValue PadCustomNeg1_Bf16_Test = PadValueCustom((bfloat16_t)-1.0); // bf16 using bfloat16_t type
157	- static_assert(PadCustomNeg1_Bf16_Test == static_cast<PadValue>(0x000000010000BF80ULL),
158	- "PadValueCustom bf16 encoding test");
159	- // Verify decoding: getCustomPadBits should return 0xBF80 (bf16 -1.0), NOT 0 from bits >> 16
160	- static_assert(getCustomPadBits(PadCustomNeg1_Bf16_Test) == 0xBF80U, "PadValueCustom bf16 decoding test");
161		47
162	- // Test custom pad bits extraction for each type (catches decode bugs!)	48	+ TLOAD(srcTile, srcGlobal);
163	- // For 16-bit types, bits & 0xFFFF must return the correct fp16/bf16 bits
164	- constexpr uint32_t float_bits = getCustomPadBits(PadCustomNeg1_Test);
165	- constexpr uint32_t half_bits = getCustomPadBits(PadCustomNeg1_Half_Test) & 0xFFFF;
166	- constexpr uint32_t bf16_bits = getCustomPadBits(PadCustomNeg1_Bf16_Test) & 0xFFFF;
167	- static_assert(float_bits == 0xBF800000U, "Custom pad float: expected -1.0f bits");
168	- static_assert(half_bits == 0xBC00U, "Custom pad half: expected fp16 -1.0 bits (0xBC00)");
169	- static_assert(bf16_bits == 0xBF80U, "Custom pad bf16: expected bf16 -1.0 bits (0xBF80)");
170		49
171	- using TileDataP =	50	+ set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
172	- Tile<TileType::Vec, T, kTRows_, kTCols_, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, FillPadVal_>;	51	+ wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
173	- TileDataP vecTileP(kTRows_, kTCols_);
174	- TASSIGN(vecTileP, (uint64_t)ubaddr1);
175		52
176	if constexpr (expand) {	53	if constexpr (expand) {
177	- using TileData = Tile<TileType::Vec, T, kTRows_, shape4_aligned, BLayout::RowMajor, -1, -1, SLayout::NoneBox,	54	+ TFILLPAD_EXPAND(dstTile, srcTile);
178	- 512, LoadPadVal_>;	55	+ } else if (inplace) {
179	- // using TileData = Tile<TileType::Vec, T, kTRows_, kTCols_, BLayout::RowMajor, -1, -1>;	56	+ TFILLPAD_INPLACE(dstTile, srcTile);
180	-
181	- TileData vecTile(shape3, shape4);
182	- TASSIGN(vecTile, (uint64_t)ubaddr0);
183	-
184	- // TLOAD(vecTile, srcGlobal); //warm up...
185	- TLOAD(vecTile, srcGlobal);
186	-#ifndef __PTO_AUTO__
187	- set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
188	- wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
189	-#endif
190	- TFILLPAD_EXPAND(vecTileP, vecTile);
191	} else {	57	} else {
192	- using TileData =	58	+ TFILLPAD(dstTile, srcTile);
193	- Tile<TileType::Vec, T, kTRows_, kTCols_, BLayout::RowMajor, -1, -1, SLayout::NoneBox, 512, LoadPadVal_>;
194	- // using TileData = Tile<TileType::Vec, T, kTRows_, kTCols_, BLayout::RowMajor, -1, -1>;
195	-
196	- TileData vecTile(shape3, shape4);
197	- TASSIGN(vecTile, (uint64_t)ubaddr0);
198	-
199	- // TLOAD(vecTile, srcGlobal); //warm up...
200	- TLOAD(vecTile, srcGlobal);
201	-#ifndef __PTO_AUTO__
202	- set_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
203	- wait_flag(PIPE_MTE2, PIPE_V, EVENT_ID0);
204	-#endif
205	- if constexpr (inplace) {
206	-#ifdef __PTO_AUTO__
207	- TRESHAPE(vecTileP, vecTile);
208	-#endif
209	- TFILLPAD_INPLACE(vecTileP, vecTile);
210	- } else
211	- TFILLPAD(vecTileP, vecTile);
212	}	59	}
213		60
214	-#ifndef __PTO_AUTO__
215	set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);	61	set_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
216	wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);	62	wait_flag(PIPE_V, PIPE_MTE3, EVENT_ID0);
217	-#endif	63	+ TSTORE(dstGlobal, dstTile);
218	- TSTORE(dstGlobal, vecTileP);
219	-#ifndef __PTO_AUTO__
220	- set_flag(PIPE_MTE2, PIPE_S, EVENT_ID0);
221	- wait_flag(PIPE_MTE2, PIPE_S, EVENT_ID0);
222	-#endif
223	}	64	}
224		65
225	-extern "C" __global__ AICORE void launchTFILLPAD_1(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,	66	+extern "C" __global__ AICORE void launchTFILLPAD_1(__gm__ uint8_t out, __gm__ uint8_t src)
226	- int gShape2, int gRows, int gCols)
227	{	67	{
228	- runTFILLPAD<float, 1, 1, 1, 128, 127, 128, 128, 1, PadValue::Max, PadValue::Max>(	68	+ runTFILLPAD<float, 64, 127, 64, 128, PadValue::Max, PadValue::Max>((__gm__ float )out, (__gm__ float )src);
229	- (__gm__ float )out, (__gm__ float )src, gShape0, gShape1, gShape2, gRows, gCols);
230	}	69	}
231		70
232	-extern "C" __global__ AICORE void launchTFILLPAD_2(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,	71	+extern "C" __global__ AICORE void launchTFILLPAD_2(__gm__ uint8_t out, __gm__ uint8_t src)
233	- int gShape2, int gRows, int gCols)
234	{	72	{
235	- runTFILLPAD<float, 1, 1, 1, 128, 127, 128, 160, 1, PadValue::Max, PadValue::Max>(	73	+ runTFILLPAD<float, 64, 127, 64, 144, PadValue::Max, PadValue::Max>((__gm__ float )out, (__gm__ float )src);
236	- (__gm__ float )out, (__gm__ float )src, gShape0, gShape1, gShape2, gRows, gCols);
237	}	74	}
238		75
239	-extern "C" __global__ AICORE void launchTFILLPAD_3(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,	76	+extern "C" __global__ AICORE void launchTFILLPAD_3(__gm__ uint8_t out, __gm__ uint8_t src)
240	- int gShape2, int gRows, int gCols)
241	{	77	{
242	- runTFILLPAD<float, 1, 1, 1, 128, 127, 128, 160, 1, PadValue::Min, PadValue::Max>(	78	+ runTFILLPAD<float, 64, 127, 64, 160, PadValue::Min, PadValue::Max>((__gm__ float )out, (__gm__ float )src);
243	- (__gm__ float )out, (__gm__ float )src, gShape0, gShape1, gShape2, gRows, gCols);
244	}	79	}
245		80
246	-extern "C" __global__ AICORE void launchTFILLPAD_4(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,	81	+extern "C" __global__ AICORE void launchTFILLPAD_4(__gm__ uint8_t out, __gm__ uint8_t src)
247	- int gShape2, int gRows, int gCols)
248	{	82	{
249	- runTFILLPAD<float, 1, 1, 1, 260, 7, 260, 16, 1, PadValue::Min, PadValue::Max>(	83	+ runTFILLPAD<float, 260, 7, 260, 16, PadValue::Min, PadValue::Max>((__gm__ float )out, (__gm__ float )src);
250	- (__gm__ float )out, (__gm__ float )src, gShape0, gShape1, gShape2, gRows, gCols);
251	}	84	}
252		85
253	-extern "C" __global__ AICORE void launchTFILLPAD_5(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,	86	+extern "C" __global__ AICORE void launchTFILLPAD_5(__gm__ uint8_t out, __gm__ uint8_t src)
254	- int gShape2, int gRows, int gCols)
255	{	87	{
256	- runTFILLPAD<float, 1, 1, 1, 260, 7, 260, 16, 1, PadValue::Min, PadValue::Max, true>(	88	+ runTFILLPAD<float, 260, 7, 260, 16, PadValue::Min, PadValue::Max, true>((__gm__ float )out, (__gm__ float )src);
257	- (__gm__ float )out, (__gm__ float )src, gShape0, gShape1, gShape2, gRows, gCols);
258	}	89	}
259		90
260	-extern "C" __global__ AICORE void launchTFILLPAD_6(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,	91	+extern "C" __global__ AICORE void launchTFILLPAD_6(__gm__ uint8_t out, __gm__ uint8_t src)
261	- int gShape2, int gRows, int gCols)
262	{	92	{
263	- runTFILLPAD<uint16_t, 1, 1, 1, 260, 7, 260, 32, 1, PadValue::Min, PadValue::Max>(	93	+ runTFILLPAD<uint16_t, 260, 7, 260, 32, PadValue::Min, PadValue::Max>((__gm__ uint16_t *)out,
264	- (__gm__ uint16_t )out, (__gm__ uint16_t )src, gShape0, gShape1, gShape2, gRows, gCols);	94	+ (__gm__ uint16_t *)src);
265	}	95	}
266		96
267	-extern "C" __global__ AICORE void launchTFILLPAD_7(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,	97	+extern "C" __global__ AICORE void launchTFILLPAD_7(__gm__ uint8_t out, __gm__ uint8_t src)
268	- int gShape2, int gRows, int gCols)
269	{	98	{
270	- runTFILLPAD<int8_t, 1, 1, 1, 260, 7, 260, 64, 1, PadValue::Min, PadValue::Max>(	99	+ runTFILLPAD<int8_t, 260, 7, 260, 64, PadValue::Min, PadValue::Max>((__gm__ int8_t )out, (__gm__ int8_t )src);
271	- (__gm__ int8_t )out, (__gm__ int8_t )src, gShape0, gShape1, gShape2, gRows, gCols);
272	}	100	}
273		101
274	-extern "C" __global__ AICORE void launchTFILLPAD_8(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,	102	+extern "C" __global__ AICORE void launchTFILLPAD_8(__gm__ uint8_t out, __gm__ uint8_t src)
275	- int gShape2, int gRows, int gCols)
276	{	103	{
277	- runTFILLPAD<uint16_t, 1, 1, 1, 259, 7, 260, 32, 1, PadValue::Min, PadValue::Max, false, true>(	104	+ runTFILLPAD<uint16_t, 259, 7, 260, 32, PadValue::Min, PadValue::Max, false, true>((__gm__ uint16_t *)out,
278	- (__gm__ uint16_t )out, (__gm__ uint16_t )src, gShape0, gShape1, gShape2, gRows, gCols);	105	+ (__gm__ uint16_t *)src);
279	}	106	}
280		107
281	-extern "C" __global__ AICORE void launchTFILLPAD_9(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,	108	+extern "C" __global__ AICORE void launchTFILLPAD_9(__gm__ uint8_t out, __gm__ uint8_t src)
282	- int gShape2, int gRows, int gCols)
283	{	109	{
284	- runTFILLPAD<int8_t, 1, 1, 1, 259, 7, 260, 64, 1, PadValue::Min, PadValue::Max, false, true>(	110	+ runTFILLPAD<int8_t, 259, 7, 260, 64, PadValue::Min, PadValue::Max, false, true>((__gm__ int8_t *)out,
285	- (__gm__ int8_t )out, (__gm__ int8_t )src, gShape0, gShape1, gShape2, gRows, gCols);	111	+ (__gm__ int8_t *)src);
286	}	112	}
287		113
288	-extern "C" __global__ AICORE void launchTFILLPAD_10(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,	114	+extern "C" __global__ AICORE void launchTFILLPAD_10(__gm__ uint8_t out, __gm__ uint8_t src)
289	- int gShape2, int gRows, int gCols)
290	{	115	{
291	- runTFILLPAD<int16_t, 1, 1, 1, 260, 7, 260, 32, 1, PadValue::Min, PadValue::Min>(	116	+ runTFILLPAD<int16_t, 260, 7, 260, 32, PadValue::Min, PadValue::Min>((__gm__ int16_t )out, (__gm__ int16_t )src);
292	- (__gm__ int16_t )out, (__gm__ int16_t )src, gShape0, gShape1, gShape2, gRows, gCols);
293	}	117	}
294		118
295	-extern "C" __global__ AICORE void launchTFILLPAD_11(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,	119	+extern "C" __global__ AICORE void launchTFILLPAD_11(__gm__ uint8_t out, __gm__ uint8_t src)
296	- int gShape2, int gRows, int gCols)
297	{	120	{
298	- runTFILLPAD<int32_t, 1, 1, 1, 260, 7, 260, 32, 1, PadValue::Min, PadValue::Min>(	121	+ runTFILLPAD<int32_t, 260, 7, 260, 32, PadValue::Min, PadValue::Min>((__gm__ int32_t )out, (__gm__ int32_t )src);
299	- (__gm__ int32_t )out, (__gm__ int32_t )src, gShape0, gShape1, gShape2, gRows, gCols);
300	-}
301	-
302	-// Case 12: Custom pad value (-1.0f)
303	-extern "C" __global__ AICORE void launchTFILLPAD_12(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,
304	- int gShape2, int gRows, int gCols)
305	-{
306	- runTFILLPAD<float, 1, 1, 1, 128, 64, 128, 128, 1, PadValue::Null, PadCustomNeg1>(
307	- (__gm__ float )out, (__gm__ float )src, gShape0, gShape1, gShape2, gRows, gCols);
308	-}
309	-
310	-// Case 13: Custom pad value for both TLOAD and TFILLPAD (32B unaligned: 127 cols)
311	-extern "C" __global__ AICORE void launchTFILLPAD_13(__gm__ uint8_t out, __gm__ uint8_t src, int gShape0, int gShape1,
312	- int gShape2, int gRows, int gCols)
313	-{
314	- runTFILLPAD<float, 1, 1, 1, 128, 127, 128, 160, 1, PadCustomNeg1, PadCustomNeg1>(
315	- (__gm__ float )out, (__gm__ float )src, gShape0, gShape1, gShape2, gRows, gCols);
316	}	122	}
317		123
318	template <int32_t testKey>	124	template <int32_t testKey>
319	void launchTFILLPAD(uint8_t out, uint8_t src, void *stream)	125	void launchTFILLPAD(uint8_t out, uint8_t src, void *stream)
320	{	126	{
321	if constexpr (testKey == 1) {	127	if constexpr (testKey == 1) {
322	- launchTFILLPAD_1<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 128, 127);	128	+ launchTFILLPAD_1<<<1, nullptr, stream>>>(out, src);
323	} else if constexpr (testKey == 2) {	129	} else if constexpr (testKey == 2) {
324	- launchTFILLPAD_2<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 128, 160);	130	+ launchTFILLPAD_2<<<1, nullptr, stream>>>(out, src);
325	} else if constexpr (testKey == 3) {	131	} else if constexpr (testKey == 3) {
326	- launchTFILLPAD_3<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 128, 160);	132	+ launchTFILLPAD_3<<<1, nullptr, stream>>>(out, src);
327	} else if constexpr (testKey == 4) {	133	} else if constexpr (testKey == 4) {
328	- launchTFILLPAD_4<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 260, 7);	134	+ launchTFILLPAD_4<<<1, nullptr, stream>>>(out, src);
329	} else if constexpr (testKey == 5) {	135	} else if constexpr (testKey == 5) {
330	- launchTFILLPAD_5<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 260, 7);	136	+ launchTFILLPAD_5<<<1, nullptr, stream>>>(out, src);
331	} else if constexpr (testKey == 6) {	137	} else if constexpr (testKey == 6) {
332	- launchTFILLPAD_6<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 260, 7);	138	+ launchTFILLPAD_6<<<1, nullptr, stream>>>(out, src);
333	} else if constexpr (testKey == 7) {	139	} else if constexpr (testKey == 7) {
334	- launchTFILLPAD_7<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 260, 7);	140	+ launchTFILLPAD_7<<<1, nullptr, stream>>>(out, src);
335	} else if constexpr (testKey == 8) {	141	} else if constexpr (testKey == 8) {
336	- launchTFILLPAD_8<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 260, 7);	142	+ launchTFILLPAD_8<<<1, nullptr, stream>>>(out, src);
337	} else if constexpr (testKey == 9) {	143	} else if constexpr (testKey == 9) {
338	- launchTFILLPAD_9<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 260, 7);	144	+ launchTFILLPAD_9<<<1, nullptr, stream>>>(out, src);
339	} else if constexpr (testKey == 10) {	145	} else if constexpr (testKey == 10) {
340	- launchTFILLPAD_10<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 260, 7);	146	+ launchTFILLPAD_10<<<1, nullptr, stream>>>(out, src);
341	} else if constexpr (testKey == 11) {	147	} else if constexpr (testKey == 11) {
342	- launchTFILLPAD_11<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 260, 7);	148	+ launchTFILLPAD_11<<<1, nullptr, stream>>>(out, src);
343	- } else if constexpr (testKey == 12) {
344	- launchTFILLPAD_12<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 128, 64);
345	- } else if constexpr (testKey == 13) {
346	- launchTFILLPAD_13<<<1, nullptr, stream>>>(out, src, 1, 1, 1, 128, 127);
347	}	149	}
348	}	150	}
349		151
350	-template <typename T>	152	+template void launchTFILLPAD<1>(uint8_t out, uint8_t src, void *stream);
351	-constexpr T getGoldenZero()	153	+template void launchTFILLPAD<2>(uint8_t out, uint8_t src, void *stream);
352	-{	154	+template void launchTFILLPAD<3>(uint8_t out, uint8_t src, void *stream);
353	- return T{0};	155	+template void launchTFILLPAD<4>(uint8_t out, uint8_t src, void *stream);
354	-}	156	+template void launchTFILLPAD<5>(uint8_t out, uint8_t src, void *stream);
355	-	157	+template void launchTFILLPAD<6>(uint8_t out, uint8_t src, void *stream);
356	-template <typename U, int Shape0, int Shape1, int Shape2, int Shape3, int Shape4, int kTRows_, int kTCols_,	158	+template void launchTFILLPAD<7>(uint8_t out, uint8_t src, void *stream);
357	- auto PadVal_ = PadValue::Null>	159	+template void launchTFILLPAD<8>(uint8_t out, uint8_t src, void *stream);
358	-int get_input_golden_case(uint8_t input, uint8_t golden)	160	+template void launchTFILLPAD<9>(uint8_t out, uint8_t src, void *stream);
359	-{	161	+template void launchTFILLPAD<10>(uint8_t out, uint8_t src, void *stream);
360	- auto arr = getGoldenZero<U>();
361	- using T = decltype(arr);
362	-
363	- constexpr int shape4_aligned = align_to_32B(Shape4, T);
364	- int in_shape[5] = {Shape0, Shape1, Shape2, Shape3, Shape4};
365	- int out_shape[5] = {Shape0, Shape1, Shape2, kTRows_, kTCols_};
366	- int in_capacity = in_shape[0] * in_shape[1] * in_shape[2] * in_shape[3] * in_shape[4];
367	- int out_capacity = out_shape[0] * out_shape[1] * out_shape[2] * out_shape[3] * out_shape[4];
368	- int in_byteSize = in_capacity * sizeof(T);
369	- int out_byteSize = out_capacity * sizeof(T);
370	-
371	- U u_padVal[1] = {0};
372	- if constexpr (static_cast<uint64_t>(PadVal_) >= static_cast<uint64_t>(PadValue::CustomBase)) {
373	- // Custom pad value - extract float bits
374	- uint32_t bits = static_cast<uint32_t>(static_cast<uint64_t>(PadVal_) & 0xFFFFFFFFULL);
375	- u_padVal[0] = reinterpret_cast<const U >(&bits);
376	- } else if (std::numeric_limits<U>::has_infinity) {
377	- if (PadVal_ == PadValue::Max)
378	- u_padVal[0] = std::numeric_limits<U>::infinity();
379	- else if (PadVal_ == PadValue::Min)
380	- u_padVal[0] = -std::numeric_limits<U>::infinity();
381	- } else {
382	- if (PadVal_ == PadValue::Max)
383	- u_padVal[0] = std::numeric_limits<U>::max();
384	- else if (PadVal_ == PadValue::Min)
385	- u_padVal[0] = std::numeric_limits<U>::min();
386	- }
387	- T t_padVal = (T )(u_padVal);
388	-
389	- T in_arr[Shape0][Shape1][Shape2][Shape3][Shape4] = {};
390	- T gold_arr[Shape0][Shape1][Shape2][kTRows_][kTCols_] = {};
391	- for (int x0 = 0; x0 < Shape0; x0++)
392	- for (int x1 = 0; x1 < Shape1; x1++)
393	- for (int x2 = 0; x2 < Shape2; x2++)
394	- for (int i = 0; i < kTRows_; i++) {
395	- for (int j = 0; j < kTCols_; j++) {
396	- if (i < Shape3 && j < Shape4) {
397	- in_arr[x0][x1][x2][i][j] = x0 * Shape1 * Shape2 * Shape3 * Shape4 +
398	- x1 * Shape2 * Shape3 * Shape4 + x2 * Shape3 * Shape4 +
399	- i * Shape4 + j;
400	- gold_arr[x0][x1][x2][i][j] = in_arr[x0][x1][x2][i][j];
401	- } else {
402	- gold_arr[x0][x1][x2][i][j] = t_padVal;
403	- }
404	- } // j
405	- } // i
406	-
407	- std::copy((uint8_t )in_arr, ((uint8_t )(in_arr)) + in_byteSize, input);
408	- std::copy((uint8_t )gold_arr, ((uint8_t )(gold_arr)) + out_byteSize, golden);
409	- return sizeof(gold_arr);
410	-}
411	-
412	-template <int32_t testKey>
413	-int get_input_golden(uint8_t input, uint8_t golden)
414	-{
415	- if constexpr (testKey == 1) {
416	- return get_input_golden_case<float, 1, 1, 1, 128, 127, 128, 128, PadValue::Max>(input, golden);
417	- } else if constexpr (testKey == 2 \|\| testKey == 3) {
418	- return get_input_golden_case<float, 1, 1, 1, 128, 127, 128, 160, PadValue::Max>(input, golden);
419	- } else if constexpr (testKey == 4 \|\| testKey == 5) {
420	- return get_input_golden_case<float, 1, 1, 1, 260, 7, 260, 16, PadValue::Max>(input, golden);
421	- } else if constexpr (testKey == 6) {
422	- return get_input_golden_case<uint16_t, 1, 1, 1, 260, 7, 260, 32, PadValue::Max>(input, golden);
423	- } else if constexpr (testKey == 7) {
424	- return get_input_golden_case<int8_t, 1, 1, 1, 260, 7, 260, 64, PadValue::Max>(input, golden);
425	- } else if constexpr (testKey == 8) {
426	- return get_input_golden_case<uint16_t, 1, 1, 1, 259, 7, 260, 32, PadValue::Max>(input, golden);
427	- } else if constexpr (testKey == 9) {
428	- return get_input_golden_case<int8_t, 1, 1, 1, 259, 7, 260, 64, PadValue::Max>(input, golden);
429	- } else if constexpr (testKey == 10) {
430	- return get_input_golden_case<int16_t, 1, 1, 1, 260, 7, 260, 32, PadValue::Min>(input, golden);
431	- } else if constexpr (testKey == 11) {
432	- return get_input_golden_case<int32_t, 1, 1, 1, 260, 7, 260, 32, PadValue::Min>(input, golden);
433	- } else if constexpr (testKey == 12) {
434	- return get_input_golden_case<float, 1, 1, 1, 128, 64, 128, 128, PadCustomNeg1>(input, golden);
435	- } else if constexpr (testKey == 13) {
436	- return get_input_golden_case<float, 1, 1, 1, 128, 127, 128, 160, PadCustomNeg1>(input, golden);
437	- }
438	-
439	- return 0;
440	-}
441	-
442	-template void launchTFILLPAD<1>(uint8_t out, uint8_t src, void *stream); // 实例化 Key=0 的版本
443	-template void launchTFILLPAD<2>(uint8_t out, uint8_t src, void *stream); // 实例化 Key=0 的版本
444	-template void launchTFILLPAD<3>(uint8_t out, uint8_t src, void *stream); // 实例化 Key=0 的版本
445	-template void launchTFILLPAD<4>(uint8_t out, uint8_t src, void *stream); // 实例化 Key=0 的版本
446	-template void launchTFILLPAD<5>(uint8_t out, uint8_t src, void *stream); // 实例化 Key=0 的版本
447	-template void launchTFILLPAD<6>(uint8_t out, uint8_t src, void *stream); // 实例化 Key=0 的版本
448	-template void launchTFILLPAD<7>(uint8_t out, uint8_t src, void *stream); // 实例化 Key=0 的版本
449	-template void launchTFILLPAD<8>(uint8_t out, uint8_t src, void *stream); // 实例化 Key=0 的版本
450	-template void launchTFILLPAD<9>(uint8_t out, uint8_t src, void *stream); // 实例化 Key=0 的版本
451	-template void launchTFILLPAD<10>(uint8_t out, uint8_t src, void *stream); // 实例化 Key=0 的版本
452	template void launchTFILLPAD<11>(uint8_t out, uint8_t src, void *stream);	162	template void launchTFILLPAD<11>(uint8_t out, uint8_t src, void *stream);
453	-template void launchTFILLPAD<12>(uint8_t out, uint8_t src, void *stream); // 实例化 Key=0 的版本
454	-template void launchTFILLPAD<13>(uint8_t out, uint8_t src, void *stream);
455	-
456	-template int get_input_golden<1>(uint8_t input, uint8_t golden);
457	-template int get_input_golden<2>(uint8_t input, uint8_t golden);
458	-template int get_input_golden<3>(uint8_t input, uint8_t golden);
459	-template int get_input_golden<4>(uint8_t input, uint8_t golden);
460	-template int get_input_golden<5>(uint8_t input, uint8_t golden);
461	-template int get_input_golden<6>(uint8_t input, uint8_t golden);
462	-template int get_input_golden<7>(uint8_t input, uint8_t golden);
463	-template int get_input_golden<8>(uint8_t input, uint8_t golden);
464	-template int get_input_golden<9>(uint8_t input, uint8_t golden);
465	-template int get_input_golden<10>(uint8_t input, uint8_t golden);
466	-template int get_input_golden<11>(uint8_t input, uint8_t golden);
467	-template int get_input_golden<12>(uint8_t input, uint8_t golden);
468	-template int get_input_golden<13>(uint8_t input, uint8_t golden);

Mtests/run_st.sh+5-0

	@@ -111,6 +111,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
111	#include "pto/npu/a2a3/TPartMul.hpp"	111	#include "pto/npu/a2a3/TPartMul.hpp"
112	#include "pto/npu/a2a3/TPartMax.hpp"	112	#include "pto/npu/a2a3/TPartMax.hpp"
113	#include "pto/npu/a2a3/TPartMin.hpp"	113	#include "pto/npu/a2a3/TPartMin.hpp"
		114	+#include "pto/npu/a2a3/TPow.hpp"
114	#include "pto/npu/a2a3/TImg2col.hpp"	115	#include "pto/npu/a2a3/TImg2col.hpp"
115	#include "pto/npu/a2a3/TSetFmatrix.hpp"	116	#include "pto/npu/a2a3/TSetFmatrix.hpp"
116	#include "pto/npu/a2a3/TSetImg2colRpt.hpp"	117	#include "pto/npu/a2a3/TSetImg2colRpt.hpp"
	@@ -229,6 +230,7 @@ See LICENSE in the root of the software repository for the full text of the Lice
229	#include "pto/npu/a5/TPartMul.hpp"	230	#include "pto/npu/a5/TPartMul.hpp"
230	#include "pto/npu/a5/TPartMax.hpp"	231	#include "pto/npu/a5/TPartMax.hpp"
231	#include "pto/npu/a5/TPartMin.hpp"	232	#include "pto/npu/a5/TPartMin.hpp"
		233	+#include "pto/npu/a5/TPow.hpp"
232	#include "pto/npu/a5/TQuant.hpp"	234	#include "pto/npu/a5/TQuant.hpp"
233	#include "pto/npu/a5/TDeQuant.hpp"	235	#include "pto/npu/a5/TDeQuant.hpp"
234	#include "pto/npu/a5/TImg2col.hpp"	236	#include "pto/npu/a5/TImg2col.hpp"

	@@ -369,6 +369,8 @@ if [ "$ENABLE_A5" = "true" ]; then
369	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmul -g TPARTMULTest.case_float_64x64_64x64_64x64	369	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmul -g TPARTMULTest.case_float_64x64_64x64_64x64
370	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmax -g TPARTMAXTest.case_fp32_64x64_64x64_64x64	370	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmax -g TPARTMAXTest.case_fp32_64x64_64x64_64x64
371	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmin -g TPARTMINTest.case_fp32_64x64_64x64_64x64	371	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmin -g TPARTMINTest.case_fp32_64x64_64x64_64x64
		372	+ python3 tests/script/run_st.py $ARGS -w -v a5 -t tpow -g TPOWTest.case1
		373	+ python3 tests/script/run_st.py $ARGS -w -v a5 -t tpows -g TPOWSTest.case1
372	python3 tests/script/run_st.py $ARGS -w -v a5 -t tprelu -g TPRELUTest.case1	374	python3 tests/script/run_st.py $ARGS -w -v a5 -t tprelu -g TPRELUTest.case1
373	python3 tests/script/run_st.py $ARGS -w -v a5 -t trem -g TREMTest.case1	375	python3 tests/script/run_st.py $ARGS -w -v a5 -t trem -g TREMTest.case1
374	python3 tests/script/run_st.py $ARGS -w -v a5 -t trowexpand -g TROWEXPANDTest.case5_float_16_8_16_127	376	python3 tests/script/run_st.py $ARGS -w -v a5 -t trowexpand -g TROWEXPANDTest.case5_float_16_8_16_127
	@@ -511,6 +513,8 @@ if [ "$ENABLE_A5" = "true" ]; then
511	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartadd	513	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartadd
512	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmul	514	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmul
513	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmax	515	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmax
		516	+ python3 tests/script/run_st.py $ARGS -w -v a5 -t tpow
		517	+ python3 tests/script/run_st.py $ARGS -w -v a5 -t tpows
514	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmin	518	python3 tests/script/run_st.py $ARGS -w -v a5 -t tpartmin
515	python3 tests/script/run_st.py $ARGS -w -v a5 -t tprelu	519	python3 tests/script/run_st.py $ARGS -w -v a5 -t tprelu
516	python3 tests/script/run_st.py $ARGS -w -v a5 -t trem	520	python3 tests/script/run_st.py $ARGS -w -v a5 -t trem
	@@ -564,6 +568,7 @@ fi
564	if [ "$ENABLE_KIRIN9030" = "true" ]; then	568	if [ "$ENABLE_KIRIN9030" = "true" ]; then
565	python3 tests/script/build_st.py $ARGS -v kirin9030 -t all	569	python3 tests/script/build_st.py $ARGS -v kirin9030 -t all
566	python3 tests/script/run_st.py $ARGS -w -v kirin9030 -t textract	570	python3 tests/script/run_st.py $ARGS -w -v kirin9030 -t textract
		571	+ python3 tests/script/run_st.py $ARGS -w -v kirin9030 -t tfillpad
567	python3 tests/script/run_st.py $ARGS -w -v kirin9030 -t tmov	572	python3 tests/script/run_st.py $ARGS -w -v kirin9030 -t tmov
568	python3 tests/script/run_st.py $ARGS -w -v kirin9030 -t tadd	573	python3 tests/script/run_st.py $ARGS -w -v kirin9030 -t tadd
569	python3 tests/script/run_st.py $ARGS -w -v kirin9030 -t tcolsum	574	python3 tests/script/run_st.py $ARGS -w -v kirin9030 -t tcolsum