fix: 修复BatchToSpaceND 等算子 codecheck 问题-ops-math-AtomGit

Mconversion/batch_to_space_nd/op_host/arch35/batch_to_space_nd_tiling_arch35.cpp+11-8

	@@ -37,6 +37,8 @@ static constexpr std::array VALUE_DATA_TYPE_ALL{
37	ge::DT_UINT32, ge::DT_INT64, ge::DT_UINT64, ge::DT_BF16, ge::DT_FLOAT16,	37	ge::DT_UINT32, ge::DT_INT64, ge::DT_UINT64, ge::DT_BF16, ge::DT_FLOAT16,
38	ge::DT_FLOAT, ge::DT_DOUBLE, ge::DT_BOOL, ge::DT_COMPLEX32, ge::DT_COMPLEX64,	38	ge::DT_FLOAT, ge::DT_DOUBLE, ge::DT_BOOL, ge::DT_COMPLEX32, ge::DT_COMPLEX64,
39	};	39	};
		40	// 固定轴数量：batch轴 + remain轴
		41	static constexpr size_t NUM_FIXED_AXES = 2;
40		42
41	// 大尾轴模板常量	43	// 大尾轴模板常量
42	// BUFFER分割数量	44	// BUFFER分割数量
	@@ -56,6 +58,7 @@ static constexpr uint32_t SMALL_C_MAX_BUFFER_SIZE = 64 * 1024U;
56	static constexpr uint32_t SMALL_C_MIN_BUFFER_SIZE = 4 * 1024U;	58	static constexpr uint32_t SMALL_C_MIN_BUFFER_SIZE = 4 * 1024U;
57	static constexpr double MIN_USED_CORES_RATIO = 0.6;	59	static constexpr double MIN_USED_CORES_RATIO = 0.6;
58	static constexpr int16_t MAX_TILING_TIME = 4;	60	static constexpr int16_t MAX_TILING_TIME = 4;
		61	static constexpr uint32_t UB_TILE_SIZE_REDUCTION_FACTOR = 2; // 迭代缩减UB tile尺寸的减半因子
59	// 每块预留大小	62	// 每块预留大小
60	static constexpr uint32_t SMALL_C_RESERVE_BUFFER_SIZE = 256U;	63	static constexpr uint32_t SMALL_C_RESERVE_BUFFER_SIZE = 256U;
61	// 被压缩的轴数量	64	// 被压缩的轴数量
	@@ -192,15 +195,15 @@ static std::string ArrayToString(const T* v, size_t size)
192		195
193	void BatchToSpaceNDTiling::ShowBaseTilingData()	196	void BatchToSpaceNDTiling::ShowBaseTilingData()
194	{	197	{
195	if (unlikely(mergedInput_.rank <= 2)) {	198	if (unlikely(mergedInput_.rank <= NUM_FIXED_AXES)) {
196	return;	199	return;
197	}	200	}
198	// 输入信息	201	// 输入信息
199	OP_LOGI(	202	OP_LOGI(
200	context_, "input: x_shape %s, block_shape %s, crops %s, y_shape %s, data type size %d",	203	context_, "input: x_shape %s, block_shape %s, crops %s, y_shape %s, data type size %d",
201	ArrayToString(mergedInput_.inShape, mergedInput_.rank).c_str(),	204	ArrayToString(mergedInput_.inShape, mergedInput_.rank).c_str(),
202	ArrayToString(mergedInput_.blockShape, mergedInput_.rank - 2).c_str(),	205	ArrayToString(mergedInput_.blockShape, mergedInput_.rank - NUM_FIXED_AXES).c_str(),
203	ArrayToString(mergedInput_.crops, (mergedInput_.rank - 2) 2).c_str(),	206	ArrayToString(mergedInput_.crops, (mergedInput_.rank - NUM_FIXED_AXES) CROPS_DIM_NUM_1).c_str(),
204	ArrayToString(mergedInput_.outShape, mergedInput_.rank).c_str(), dSize_);	207	ArrayToString(mergedInput_.outShape, mergedInput_.rank).c_str(), dSize_);
205	// soc 信息	208	// soc 信息
206	OP_LOGI(	209	OP_LOGI(
	@@ -228,7 +231,7 @@ void BatchToSpaceNDTiling::ShowSmallCTilingData()
228	context_, "tiling data: oriInShape %s, croppedInShape %s, crops %s",	231	context_, "tiling data: oriInShape %s, croppedInShape %s, crops %s",
229	ArrayToString(tilingData->oriInShape, mergedInput_.rank + blockShapeDimNum_).c_str(),	232	ArrayToString(tilingData->oriInShape, mergedInput_.rank + blockShapeDimNum_).c_str(),
230	ArrayToString(tilingData->croppedInShape, mergedInput_.rank + blockShapeDimNum_).c_str(),	233	ArrayToString(tilingData->croppedInShape, mergedInput_.rank + blockShapeDimNum_).c_str(),
231	ArrayToString(tilingData->crops, blockShapeDimNum_ 2).c_str());	234	ArrayToString(tilingData->crops, blockShapeDimNum_ CROPS_DIM_NUM_1).c_str());
232	OP_LOGI(	235	OP_LOGI(
233	context_, "\t: coreNum %u, inUbAxis %u, outUbAxis %u, inUbFactor %u, outUbFactor %u", tilingData->coreNum,	236	context_, "\t: coreNum %u, inUbAxis %u, outUbAxis %u, inUbFactor %u, outUbFactor %u", tilingData->coreNum,
234	tilingData->inUbAxis, tilingData->outUbAxis, tilingData->inUbFactor, tilingData->outUbFactor);	237	tilingData->inUbAxis, tilingData->outUbAxis, tilingData->inUbFactor, tilingData->outUbFactor);
	@@ -598,7 +601,7 @@ ge::graphStatus BatchToSpaceNDTiling::Tiling4LargeC()
598	// batch 不对齐	601	// batch 不对齐
599	ubFactorAlign[0] = 1;	602	ubFactorAlign[0] = 1;
600	// space 对齐 block shape	603	// space 对齐 block shape
601	std::copy(mergedInput_.blockShape, mergedInput_.blockShape + mergedInput_.rank - 2, ubFactorAlign.begin() + 1);	604	std::copy(mergedInput_.blockShape, mergedInput_.blockShape + mergedInput_.rank - NUM_FIXED_AXES, ubFactorAlign.begin() + 1);
602	// remain shape 对齐 ub block	605	// remain shape 对齐 ub block
603	ubFactorAlign[mergedInput_.rank - 1] = ubBlockElements_;	606	ubFactorAlign[mergedInput_.rank - 1] = ubBlockElements_;
604		607
	@@ -666,7 +669,7 @@ void BatchToSpaceNDTiling::SmallCSetInput(B2SNDSmallCTilingData* tilingData)
666	tilingData->croppedInShape[rank - 1] = mergedInput_.inShape[mergedInput_.rank - 1];	669	tilingData->croppedInShape[rank - 1] = mergedInput_.inShape[mergedInput_.rank - 1];
667		670
668	// crops	671	// crops
669	std::copy(mergedInput_.crops, (mergedInput_.crops) + blockShapeDimNum_ * 2, *(tilingData->crops));	672	std::copy(mergedInput_.crops, (mergedInput_.crops) + blockShapeDimNum_ * CROPS_DIM_NUM_1, *(tilingData->crops));
670	}	673	}
671		674
672	std::array<size_t, MAX_EXPAND_RANK> BatchToSpaceNDTiling::SmallCComputeOutputAxisPerm()	675	std::array<size_t, MAX_EXPAND_RANK> BatchToSpaceNDTiling::SmallCComputeOutputAxisPerm()
	@@ -707,7 +710,7 @@ ge::graphStatus BatchToSpaceNDTiling::Tiling4SmallC()
707	{	710	{
708	// tiling key	711	// tiling key
709	mode_ = TPL_MODE_SMALL_C;	712	mode_ = TPL_MODE_SMALL_C;
710	blockShapeDimNum_ = mergedInput_.rank - 2;	713	blockShapeDimNum_ = mergedInput_.rank - NUM_FIXED_AXES;
711		714
712	// tiling data	715	// tiling data
713	auto tilingData = context_->GetTilingData<B2SNDSmallCTilingData>();	716	auto tilingData = context_->GetTilingData<B2SNDSmallCTilingData>();
	@@ -741,7 +744,7 @@ ge::graphStatus BatchToSpaceNDTiling::Tiling4SmallC()
741	if (tempCnt >= MAX_TILING_TIME) {	744	if (tempCnt >= MAX_TILING_TIME) {
742	break;	745	break;
743	}	746	}
744	inputElements = inputElements / 2;	747	inputElements = inputElements / UB_TILE_SIZE_REDUCTION_FACTOR;
745	auto tiling1 = DualSideTiling(	748	auto tiling1 = DualSideTiling(
746	context_, ubBlockElements_, tilingData->croppedInShape, yAxisPerm.data(), xNeedAlignAxis,	749	context_, ubBlockElements_, tilingData->croppedInShape, yAxisPerm.data(), xNeedAlignAxis,
747	blockShapeDimNum_ + mergedInput_.rank);	750	blockShapeDimNum_ + mergedInput_.rank);

Mconversion/batch_to_space_nd/op_kernel/arch35/batch_to_space_nd_large_c.h+31-22

	@@ -82,6 +82,15 @@ private:
82	constexpr static uint32_t BUFFER_NUM = 2;	82	constexpr static uint32_t BUFFER_NUM = 2;
83	constexpr static uint32_t UB_BLOCK = Ops::Base::GetUbBlockSize();	83	constexpr static uint32_t UB_BLOCK = Ops::Base::GetUbBlockSize();
84	constexpr static uint32_t BLK_ELEMS = UB_BLOCK / sizeof(T);	84	constexpr static uint32_t BLK_ELEMS = UB_BLOCK / sizeof(T);
		85	// Axis offset constants
		86	constexpr static uint32_t AXIS_OFFSET_W = 2; // W轴相对rank末尾的偏移量
		87	constexpr static uint32_t AXIS_OFFSET_H = 3; // H轴相对rank末尾的偏移量
		88	constexpr static uint32_t BLOCK_AXIS_OFFSET_H = 2; // H轴在blockShape数组中的偏移量
		89	// Threshold/rank constants
		90	constexpr static uint32_t MIN_SPATIAL_DIMS = 2; // 处理H轴所需的最少空间维度数
		91	constexpr static uint32_t MIN_RANK_FOR_H_AXIS = 4; // 处理H轴所需的最小tensor rank
		92	constexpr static uint32_t RANK_3D = 3; // 3维tensor的rank值
		93	constexpr static uint32_t RANK_5D = 5; // 5维tensor的rank值
85		94
86	public:	95	public:
87	__aicore__ inline BatchToSpaceLargeC(TPipe* pipe) { pipe_ = pipe; }	96	__aicore__ inline BatchToSpaceLargeC(TPipe* pipe) { pipe_ = pipe; }
	@@ -108,12 +117,12 @@ public:
108	axisPreProduct_ *= tdPtr_->input.outShape[i];	117	axisPreProduct_ *= tdPtr_->input.outShape[i];
109	}	118	}
110	tileQueryNumsDim_ = totalCount_ / axisPreProduct_;	119	tileQueryNumsDim_ = totalCount_ / axisPreProduct_;
111	if (blockShapeSize_ >= 2) {	120	if (blockShapeSize_ >= MIN_SPATIAL_DIMS) {
112	// H 轴	121	// H 轴
113	cropHtop_ = tdPtr_->input.crops[blockShapeSize_ - 2][0];	122	cropHtop_ = tdPtr_->input.crops[blockShapeSize_ - BLOCK_AXIS_OFFSET_H][0];
114	cropHbottom_ = tdPtr_->input.crops[blockShapeSize_ - 2][1];	123	cropHbottom_ = tdPtr_->input.crops[blockShapeSize_ - BLOCK_AXIS_OFFSET_H][1];
115	BSH_ = tdPtr_->input.blockShape[blockShapeSize_ - 2];	124	BSH_ = tdPtr_->input.blockShape[blockShapeSize_ - BLOCK_AXIS_OFFSET_H];
116	outH_ = tdPtr_->input.outShape[rank_ - 3];	125	outH_ = tdPtr_->input.outShape[rank_ - AXIS_OFFSET_H];
117		126
118	CalcBoundaryBlock(cropHtop_, cropHbottom_, BSH_, outH_, headLenH_, tailLenH_, middleLenH_);	127	CalcBoundaryBlock(cropHtop_, cropHbottom_, BSH_, outH_, headLenH_, tailLenH_, middleLenH_);
119		128
	@@ -137,8 +146,8 @@ public:
137	cropLeft_ = tdPtr_->input.crops[blockShapeSize_ - 1][0];	146	cropLeft_ = tdPtr_->input.crops[blockShapeSize_ - 1][0];
138	cropRight_ = tdPtr_->input.crops[blockShapeSize_ - 1][1];	147	cropRight_ = tdPtr_->input.crops[blockShapeSize_ - 1][1];
139	BSW_ = tdPtr_->input.blockShape[blockShapeSize_ - 1];	148	BSW_ = tdPtr_->input.blockShape[blockShapeSize_ - 1];
140	outW_ = tdPtr_->input.outShape[rank_ - 2];	149	outW_ = tdPtr_->input.outShape[rank_ - AXIS_OFFSET_W];
141	inW_ = tdPtr_->input.inShape[rank_ - 2];	150	inW_ = tdPtr_->input.inShape[rank_ - AXIS_OFFSET_W];
142	inC_ = tdPtr_->input.inShape[rank_ - 1];	151	inC_ = tdPtr_->input.inShape[rank_ - 1];
143		152
144	CalcBoundaryBlock(cropLeft_, cropRight_, BSW_, outW_, leftCopyLen_, rightCopyLen_, middleCopyLen_);	153	CalcBoundaryBlock(cropLeft_, cropRight_, BSW_, outW_, leftCopyLen_, rightCopyLen_, middleCopyLen_);
	@@ -181,7 +190,7 @@ public:
181	uint32_t endIdx = (blockIdx_ + 1L) * perCoreCount_;	190	uint32_t endIdx = (blockIdx_ + 1L) * perCoreCount_;
182	endIdx = endIdx < totalCount_ ? endIdx : totalCount_;	191	endIdx = endIdx < totalCount_ ? endIdx : totalCount_;
183		192
184	if ((rank_ == 4 && ubAxis_ == 1) \|\| (rank_ == 5 && ubAxis_ == 2)) {	193	if (ubAxis_ == rank_ - AXIS_OFFSET_H && (rank_ == MIN_RANK_FOR_H_AXIS \|\| rank_ == RANK_5D)) {
185	int sum = getTileNumPrefixSum(startIdx);	194	int sum = getTileNumPrefixSum(startIdx);
186		195
187	uint64_t curOutIndex[MAX_INPUT_RANK] = {0, 0, 0, 0, 0, 0, 0, 0, 0};	196	uint64_t curOutIndex[MAX_INPUT_RANK] = {0, 0, 0, 0, 0, 0, 0, 0, 0};
	@@ -274,11 +283,11 @@ private:
274		283
275	if (ubAxis_ == rank_ - 1) {	284	if (ubAxis_ == rank_ - 1) {
276	DoCopyInAxisC(src, outIndex, inIndex, ubAxisInCopyNum, idx);	285	DoCopyInAxisC(src, outIndex, inIndex, ubAxisInCopyNum, idx);
277	} else if (ubAxis_ == rank_ - 2) {	286	} else if (ubAxis_ == rank_ - AXIS_OFFSET_W) {
278	DoCopyInAxisW(src, outIndex, inIndex, ubAxisInCopyNum, idx);	287	DoCopyInAxisW(src, outIndex, inIndex, ubAxisInCopyNum, idx);
279	} else if (ubAxis_ == rank_ - 3 && rank_ == 3) {	288	} else if (ubAxis_ == rank_ - AXIS_OFFSET_H && rank_ == RANK_3D) {
280	DoCopyInAxisN3(src, outIndex, inIndex, ubAxisInCopyNum, idx);	289	DoCopyInAxisN3(src, outIndex, inIndex, ubAxisInCopyNum, idx);
281	} else if (ubAxis_ == rank_ - 3 && rank_ >= 4) { // H轴	290	} else if (ubAxis_ == rank_ - AXIS_OFFSET_H && rank_ >= MIN_RANK_FOR_H_AXIS) { // H轴
282	DoCopyInAxisH(src, outIndex, inIndex, ubAxisInCopyNum, idx);	291	DoCopyInAxisH(src, outIndex, inIndex, ubAxisInCopyNum, idx);
283	}	292	}
284	}	293	}
	@@ -449,7 +458,7 @@ private:
449	wSize);	458	wSize);
450	uint64_t ubInOffset = W * alignC_;	459	uint64_t ubInOffset = W * alignC_;
451	W += ubWInCopyNum;	460	W += ubWInCopyNum;
452	Conver2InIndex(loopOutIndex, inIndex);	461	Conver2InIndex(static_cast<uint64_t*>(loopOutIndex), inIndex);
453	uint64_t inAddr = 0;	462	uint64_t inAddr = 0;
454	for (uint8_t i = 0; i < rank_; i++) {	463	for (uint8_t i = 0; i < rank_; i++) {
455	inAddr += inIndex[i] * inStride_[i];	464	inAddr += inIndex[i] * inStride_[i];
	@@ -512,8 +521,8 @@ private:
512	if (curUbFactor == 0)	521	if (curUbFactor == 0)
513	return;	522	return;
514		523
515	uint64_t instrideBSW = instrideBS[rank_ - 2];	524	uint64_t instrideBSW = instrideBS[rank_ - AXIS_OFFSET_W];
516	uint64_t instrideBSH = instrideBS[rank_ - 3];	525	uint64_t instrideBSH = instrideBS[rank_ - AXIS_OFFSET_H];
517		526
518	uint64_t leftCopyLen = leftCopyLen_;	527	uint64_t leftCopyLen = leftCopyLen_;
519	uint64_t rightCopyLen = rightCopyLen_;	528	uint64_t rightCopyLen = rightCopyLen_;
	@@ -542,8 +551,8 @@ private:
542	instrideBSH, true, false);	551	instrideBSH, true, false);
543	}	552	}
544		553
545	outIndex[rank_ - 2] = 0;	554	outIndex[rank_ - AXIS_OFFSET_W] = 0;
546	updateOutIndexByCarry(outIndex, curUbFactor, rank_ - 3);	555	updateOutIndexByCarry(outIndex, curUbFactor, rank_ - AXIS_OFFSET_H);
547	}	556	}
548		557
549	__aicore__ inline void CalcBoundaryBlock(	558	__aicore__ inline void CalcBoundaryBlock(
	@@ -692,7 +701,7 @@ private:
692	outIndex[i] = 0;	701	outIndex[i] = 0;
693	}	702	}
694		703
695	uint32_t axis = rank_ - 3; // 例如 rank=4 -> axis=1	704	uint32_t axis = rank_ - AXIS_OFFSET_H; // 例如 rank=4 -> axis=1
696		705
697	outIndex[axis] = sum;	706	outIndex[axis] = sum;
698		707
	@@ -774,7 +783,7 @@ private:
774	// 更新地址和索引	783	// 更新地址和索引
775	if (updateAddr) {	784	if (updateAddr) {
776	ubAddr += copyLen * alignedCLength_;	785	ubAddr += copyLen * alignedCLength_;
777	outIndex[rank_ - 2] += copyLen;	786	outIndex[rank_ - AXIS_OFFSET_W] += copyLen;
778	}	787	}
779	}	788	}
780		789
	@@ -790,14 +799,14 @@ private:
790	if (ubAxis_ == rank_ - 1) {	799	if (ubAxis_ == rank_ - 1) {
791	copyOutParams.blockCount = 1;	800	copyOutParams.blockCount = 1;
792	copyOutParams.blockLen = ubAxisOutCopyNum * outStride_[ubAxis_] * sizeof(T);	801	copyOutParams.blockLen = ubAxisOutCopyNum * outStride_[ubAxis_] * sizeof(T);
793	} else if (ubAxis_ == rank_ - 2) {	802	} else if (ubAxis_ == rank_ - AXIS_OFFSET_W) {
794	copyOutParams.blockCount = ubAxisOutCopyNum;	803	copyOutParams.blockCount = ubAxisOutCopyNum;
795	copyOutParams.blockLen = originC_ * sizeof(T);	804	copyOutParams.blockLen = originC_ * sizeof(T);
796	} else if (ubAxis_ == rank_ - 3 && rank_ == 3) {	805	} else if (ubAxis_ == rank_ - AXIS_OFFSET_H && rank_ == RANK_3D) {
797	copyOutParams.blockCount = ubAxisOutCopyNum * tdPtr_->input.outShape[ubAxis_ + 1]; // N-Axis*L	806	copyOutParams.blockCount = ubAxisOutCopyNum * tdPtr_->input.outShape[ubAxis_ + 1]; // N-Axis*L
798	copyOutParams.blockLen = originC_ * sizeof(T); // C	807	copyOutParams.blockLen = originC_ * sizeof(T); // C
799	} else if (ubAxis_ == rank_ - 3 && rank_ >= 4) { // H轴	808	} else if (ubAxis_ == rank_ - AXIS_OFFSET_H && rank_ >= MIN_RANK_FOR_H_AXIS) { // H轴
800	copyOutParams.blockCount = ubAxisOutCopyNum * tdPtr_->input.outShape[rank_ - 2];	809	copyOutParams.blockCount = ubAxisOutCopyNum * tdPtr_->input.outShape[rank_ - AXIS_OFFSET_W];
801	copyOutParams.blockLen = tdPtr_->input.outShape[rank_ - 1] * sizeof(T);	810	copyOutParams.blockLen = tdPtr_->input.outShape[rank_ - 1] * sizeof(T);
802	}	811	}
803	DataCopyPad(outputGm_[outAddr], src[0], copyOutParams);	812	DataCopyPad(outputGm_[outAddr], src[0], copyOutParams);

Mconversion/batch_to_space_nd/op_kernel/arch35/batch_to_space_nd_small_c.h+33-27

	@@ -27,6 +27,9 @@ using namespace Ops::Base;
27	template <typename T, uint8_t BLOCK_DIM_NUM>	27	template <typename T, uint8_t BLOCK_DIM_NUM>
28	class BatchToSpaceSmallC {	28	class BatchToSpaceSmallC {
29	private:	29	private:
		30	constexpr static uint32_t LOG_OFFSET[]{1, 2, 4, 8, 16};
		31	constexpr static uint16_t LOG_OFFSET_NUM = static_cast<uint16_t>(sizeof(LOG_OFFSET) / sizeof(LOG_OFFSET[0]));
		32
30	constexpr static uint32_t BUFFER_NUM = 2;	33	constexpr static uint32_t BUFFER_NUM = 2;
31	constexpr static uint32_t BLK_ELEMS = Ops::Base::GetUbBlockSize() / sizeof(T);	34	constexpr static uint32_t BLK_ELEMS = Ops::Base::GetUbBlockSize() / sizeof(T);
32	constexpr static uint32_t SUB_BASE = uint32_t(4294967296);	35	constexpr static uint32_t SUB_BASE = uint32_t(4294967296);
	@@ -44,8 +47,10 @@ private:
44	constexpr static uint8_t RIGHT_CROP = 1;	47	constexpr static uint8_t RIGHT_CROP = 1;
45	constexpr static uint8_t TWO_DIMENSION = 2;	48	constexpr static uint8_t TWO_DIMENSION = 2;
46	constexpr static uint8_t THIRD_DIMENSION = 3;	49	constexpr static uint8_t THIRD_DIMENSION = 3;
		50	constexpr static uint8_t OUTER_STRIDE_AXIS_OFFSET = 2; // 外层stride的轴偏移量(跨2级轴获取stride)
47	constexpr static int8_t BS_PIXEL_MAP[4][3] = {{-1, -1, -1}, {2, -1, -1}, {3, 4, -1}, {4, 5, 6}};	51	constexpr static int8_t BS_PIXEL_MAP[4][3] = {{-1, -1, -1}, {2, -1, -1}, {3, 4, -1}, {4, 5, 6}};
48	constexpr static uint8_t SHAPE_DIM_NUM = MAX_CROP_DIM_NUM * BLOCK_DIM_NUM + N_C_NUM;	52	constexpr static uint8_t SHAPE_DIM_NUM = MAX_CROP_DIM_NUM * BLOCK_DIM_NUM + N_C_NUM;
		53	constexpr static uint32_t UINT64_PROCESS_FACTOR = 2; // uint64类型数据处理时元素计数翻倍因子
49	const B2SNDSmallCTilingData* td_ = nullptr;	54	const B2SNDSmallCTilingData* td_ = nullptr;
50	GlobalTensor<T> inputGM_;	55	GlobalTensor<T> inputGM_;
51	GlobalTensor<T> outputGM_;	56	GlobalTensor<T> outputGM_;
	@@ -304,7 +309,7 @@ public:
304	{	309	{
305	__ubuf__ T* outputAddrTmp = outputAddr;	310	__ubuf__ T* outputAddrTmp = outputAddr;
306	if constexpr (sizeof(T) == sizeof(uint64_t)) {	311	if constexpr (sizeof(T) == sizeof(uint64_t)) {
307	tiledProcessSize *= 2;	312	tiledProcessSize *= UINT64_PROCESS_FACTOR;
308	}	313	}
309	uint32_t vlSize = vlSize_;	314	uint32_t vlSize = vlSize_;
310	uint64_t offset0 = cropOffset_[0][0];	315	uint64_t offset0 = cropOffset_[0][0];
	@@ -843,11 +848,12 @@ public:
843	tiledInShape_[indexMap_[axis2]] > cropOffset_[indexMap_[axis2]][1]) ?	848	tiledInShape_[indexMap_[axis2]] > cropOffset_[indexMap_[axis2]][1]) ?
844	1 :	849	1 :
845	0;	850	0;
846	loopParams.loop1SrcStride = ubOutStride_[axis2 - 2] * sizeof(T);	851	loopParams.loop1SrcStride = ubOutStride_[axis2 - OUTER_STRIDE_AXIS_OFFSET] * sizeof(T);
847	loopParams.loop1DstStride =	852	loopParams.loop1DstStride =
848	tempAxis > 0 ? outStride_[tempAxis - 1] * sizeof(T) : outStride_[axis2 - 2] * sizeof(T);	853	tempAxis > 0 ? outStride_[tempAxis - 1] * sizeof(T) : outStride_[axis2 - OUTER_STRIDE_AXIS_OFFSET] * sizeof(T);
849	uint64_t tempBlockCount = ubOutStride_[axis2] / ubOutStride_[axis1 - 1];	854	uint64_t tempBlockCount = ubOutStride_[axis2] / ubOutStride_[axis1 - 1];
850	uint64_t tempLoopSize = ubOutStride_[0] * tiledInShape_[indexMap_[0]] / ubOutStride_[axis2 - 2];	855	uint64_t tempLoopSize =
		856	ubOutStride_[0] * tiledInShape_[indexMap_[0]] / ubOutStride_[axis2 - OUTER_STRIDE_AXIS_OFFSET];
851	if (hasFirst) {	857	if (hasFirst) {
852	copyOutParams.blockCount =	858	copyOutParams.blockCount =
853	(tiledInShape_[indexMap_[axis2]] - cropOffset_[indexMap_[axis2]][0]) * tempBlockCount;	859	(tiledInShape_[indexMap_[axis2]] - cropOffset_[indexMap_[axis2]][0]) * tempBlockCount;
	@@ -862,7 +868,7 @@ public:
862	(td_->croppedInShape[indexMap_[axis2]] - cropOffset_[indexMap_[axis2]][0]) * outStride_[axis2] :	868	(td_->croppedInShape[indexMap_[axis2]] - cropOffset_[indexMap_[axis2]][0]) * outStride_[axis2] :
863	0;	869	0;
864	if (hasLast) {	870	if (hasLast) {
865	uint64_t inOffsetL = ubOutStride_[axis2 - 2] - ubOutStride_[axis2 - 1];	871	uint64_t inOffsetL = ubOutStride_[axis2 - OUTER_STRIDE_AXIS_OFFSET] - ubOutStride_[axis2 - 1];
866	uint64_t outOffsetL = outOffset + loopParams.loop1Size / tempLoopSize * outStride_[axis2 - 1];	872	uint64_t outOffsetL = outOffset + loopParams.loop1Size / tempLoopSize * outStride_[axis2 - 1];
867	copyOutParams.blockCount =	873	copyOutParams.blockCount =
868	(tiledInShape_[indexMap_[axis2]] - cropOffset_[indexMap_[axis2]][1]) * tempBlockCount;	874	(tiledInShape_[indexMap_[axis2]] - cropOffset_[indexMap_[axis2]][1]) * tempBlockCount;
	@@ -877,9 +883,9 @@ public:
877	loopParams.loop1Size = tiledInShape_[indexMap_[axis2 - 1]] -	883	loopParams.loop1Size = tiledInShape_[indexMap_[axis2 - 1]] -
878	(cropOffset_[indexMap_[axis2]][0] > 0 ? 1 : 0) -	884	(cropOffset_[indexMap_[axis2]][0] > 0 ? 1 : 0) -
879	(cropOffset_[indexMap_[axis2]][1] > 0 ? 1 : 0);	885	(cropOffset_[indexMap_[axis2]][1] > 0 ? 1 : 0);
880	loopParams.loop2SrcStride = ubOutStride_[axis2 - 2] * sizeof(T);	886	loopParams.loop2SrcStride = ubOutStride_[axis2 - OUTER_STRIDE_AXIS_OFFSET] * sizeof(T);
881	loopParams.loop2DstStride =	887	loopParams.loop2DstStride =
882	tempAxis > 0 ? outStride_[tempAxis - 1] * sizeof(T) : outStride_[axis2 - 2] * sizeof(T);	888	tempAxis > 0 ? outStride_[tempAxis - 1] * sizeof(T) : outStride_[axis2 - OUTER_STRIDE_AXIS_OFFSET] * sizeof(T);
883	loopParams.loop2Size = tempLoopSize;	889	loopParams.loop2Size = tempLoopSize;
884	SetLoopModePara(loopParams, DataCopyMVType::UB_TO_OUT);	890	SetLoopModePara(loopParams, DataCopyMVType::UB_TO_OUT);
885	DataCopyPad(outputGM_[outAddr + outOffset], dst[inOffset], copyOutParams);	891	DataCopyPad(outputGM_[outAddr + outOffset], dst[inOffset], copyOutParams);
	@@ -894,12 +900,12 @@ public:
894	copyOutParams.blockCount = tiledInShape_[indexMap_[axis1 - 1]] - hasFirst - hasLast;	900	copyOutParams.blockCount = tiledInShape_[indexMap_[axis1 - 1]] - hasFirst - hasLast;
895	copyOutParamsF.blockLen =	901	copyOutParamsF.blockLen =
896	ubOutStride_[axis1] * (tiledInShape_[indexMap_[axis1]] - cropOffset_[indexMap_[axis1]][0]) * sizeof(T);	902	ubOutStride_[axis1] * (tiledInShape_[indexMap_[axis1]] - cropOffset_[indexMap_[axis1]][0]) * sizeof(T);
897	copyOutParamsF.srcStride = (ubOutStride_[axis1 - 2] - copyOutParamsF.blockLen / sizeof(T)) / BLK_ELEMS;	903	copyOutParamsF.srcStride = (ubOutStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET] - copyOutParamsF.blockLen / sizeof(T)) / BLK_ELEMS;
898	copyOutParamsF.dstStride = (outStride_[axis1 - 2] - copyOutParamsF.blockLen / sizeof(T)) * sizeof(T);	904	copyOutParamsF.dstStride = (outStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET] - copyOutParamsF.blockLen / sizeof(T)) * sizeof(T);
899	copyOutParamsL.blockLen =	905	copyOutParamsL.blockLen =
900	ubOutStride_[axis1] * (tiledInShape_[indexMap_[axis1]] - cropOffset_[indexMap_[axis1]][1]) * sizeof(T);	906	ubOutStride_[axis1] * (tiledInShape_[indexMap_[axis1]] - cropOffset_[indexMap_[axis1]][1]) * sizeof(T);
901	copyOutParamsL.srcStride = (ubOutStride_[axis1 - 2] - copyOutParamsL.blockLen / sizeof(T)) / BLK_ELEMS;	907	copyOutParamsL.srcStride = (ubOutStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET] - copyOutParamsL.blockLen / sizeof(T)) / BLK_ELEMS;
902	copyOutParamsL.dstStride = (outStride_[axis1 - 2] - copyOutParamsL.blockLen / sizeof(T)) * sizeof(T);	908	copyOutParamsL.dstStride = (outStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET] - copyOutParamsL.blockLen / sizeof(T)) * sizeof(T);
903	tempAxis = FindOuterIndex(axis2, axis1 - 1);	909	tempAxis = FindOuterIndex(axis2, axis1 - 1);
904	if (tempAxis > 0) {	910	if (tempAxis > 0) {
905	copyOutParamsF.dstStride = (outStride_[tempAxis - 1] - copyOutParamsF.blockLen / sizeof(T)) * sizeof(T);	911	copyOutParamsF.dstStride = (outStride_[tempAxis - 1] - copyOutParamsF.blockLen / sizeof(T)) * sizeof(T);
	@@ -941,9 +947,9 @@ public:
941	loopParams.loop2SrcStride = loopParams.loop1SrcStride;	947	loopParams.loop2SrcStride = loopParams.loop1SrcStride;
942	loopParams.loop2DstStride = loopParams.loop1DstStride;	948	loopParams.loop2DstStride = loopParams.loop1DstStride;
943	loopParams.loop1Size = tempFactor;	949	loopParams.loop1Size = tempFactor;
944	loopParams.loop1SrcStride = ubOutStride_[axis1 - 2] * sizeof(T);	950	loopParams.loop1SrcStride = ubOutStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET] * sizeof(T);
945	loopParams.loop1DstStride =	951	loopParams.loop1DstStride =
946	tempAxis > 0 ? outStride_[tempAxis - 1] * sizeof(T) : outStride_[axis1 - 2] * sizeof(T);	952	tempAxis > 0 ? outStride_[tempAxis - 1] * sizeof(T) : outStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET] * sizeof(T);
947	for (auto a = 0; a < outSize; a++) {	953	for (auto a = 0; a < outSize; a++) {
948	SetLoopModePara(loopParams, DataCopyMVType::UB_TO_OUT);	954	SetLoopModePara(loopParams, DataCopyMVType::UB_TO_OUT);
949	DataCopyPad(	955	DataCopyPad(
	@@ -952,14 +958,15 @@ public:
952	ResetLoopModePara(DataCopyMVType::UB_TO_OUT);	958	ResetLoopModePara(DataCopyMVType::UB_TO_OUT);
953	}	959	}
954	} else if (axis2 == 0) {	960	} else if (axis2 == 0) {
955	uint64_t tempFactor = ubOutStride_[0] * tiledInShape_[indexMap_[0]] / ubOutStride_[axis1 - 2];	961	uint64_t tempFactor =
		962	ubOutStride_[0] * tiledInShape_[indexMap_[0]] / ubOutStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET];
956	if (hasFirst) {	963	if (hasFirst) {
957	copyOutParamsF.blockCount = copyOutParamsF.blockCount * tempFactor;	964	copyOutParamsF.blockCount = copyOutParamsF.blockCount * tempFactor;
958	DataCopyPad(outputGM_[outAddr], dst, copyOutParamsF);	965	DataCopyPad(outputGM_[outAddr], dst, copyOutParamsF);
959	}	966	}
960	loopParams.loop1SrcStride = ubOutStride_[axis1 - 2] * sizeof(T);	967	loopParams.loop1SrcStride = ubOutStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET] * sizeof(T);
961	loopParams.loop1DstStride =	968	loopParams.loop1DstStride = tempAxis > 0 ? outStride_[tempAxis - 1] * sizeof(T) :
962	tempAxis > 0 ? outStride_[tempAxis - 1] * sizeof(T) : outStride_[axis1 - 2] * sizeof(T);	969	outStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET] * sizeof(T);
963	loopParams.loop1Size = tempFactor;	970	loopParams.loop1Size = tempFactor;
964	SetLoopModePara(loopParams, DataCopyMVType::UB_TO_OUT);	971	SetLoopModePara(loopParams, DataCopyMVType::UB_TO_OUT);
965	DataCopyPad(outputGM_[outAddr + outOffset], dst[inOffset], copyOutParams);	972	DataCopyPad(outputGM_[outAddr + outOffset], dst[inOffset], copyOutParams);
	@@ -970,7 +977,7 @@ public:
970	}	977	}
971	} else {	978	} else {
972	int8_t tempAxis2 = FindOuterIndex(0, axis2);	979	int8_t tempAxis2 = FindOuterIndex(0, axis2);
973	uint64_t tempFactor = ubOutStride_[axis2 - 1] / ubOutStride_[axis1 - 2];	980	uint64_t tempFactor = ubOutStride_[axis2 - 1] / ubOutStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET];
974	loopParams.loop1SrcStride = ubOutStride_[axis2 - 1] * sizeof(T);	981	loopParams.loop1SrcStride = ubOutStride_[axis2 - 1] * sizeof(T);
975	loopParams.loop1DstStride =	982	loopParams.loop1DstStride =
976	(tempAxis2 > 0) ? outStride_[tempAxis2 - 1] * sizeof(T) : outStride_[axis2 - 1] * sizeof(T);	983	(tempAxis2 > 0) ? outStride_[tempAxis2 - 1] * sizeof(T) : outStride_[axis2 - 1] * sizeof(T);
	@@ -985,9 +992,9 @@ public:
985	DataCopyPad(outputGM_[outAddr + outOffsetL], dst[inOffsetL], copyOutParamsL);	992	DataCopyPad(outputGM_[outAddr + outOffsetL], dst[inOffsetL], copyOutParamsL);
986	}	993	}
987	ResetLoopModePara(DataCopyMVType::UB_TO_OUT);	994	ResetLoopModePara(DataCopyMVType::UB_TO_OUT);
988	loopParams.loop1SrcStride = ubOutStride_[axis1 - 2] * sizeof(T);	995	loopParams.loop1SrcStride = ubOutStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET] * sizeof(T);
989	loopParams.loop1DstStride =	996	loopParams.loop1DstStride = tempAxis > 0 ? outStride_[tempAxis - 1] * sizeof(T) :
990	tempAxis > 0 ? outStride_[tempAxis - 1] * sizeof(T) : outStride_[axis1 - 2] * sizeof(T);	997	outStride_[axis1 - OUTER_STRIDE_AXIS_OFFSET] * sizeof(T);
991	loopParams.loop1Size = tempFactor;	998	loopParams.loop1Size = tempFactor;
992	loopParams.loop2SrcStride = ubOutStride_[axis2 - 1] * sizeof(T);	999	loopParams.loop2SrcStride = ubOutStride_[axis2 - 1] * sizeof(T);
993	loopParams.loop2DstStride =	1000	loopParams.loop2DstStride =
	@@ -1093,7 +1100,7 @@ public:
1093	}	1100	}
1094		1101
1095	// 计算输入shape截取数据后的偏移	1102	// 计算输入shape截取数据后的偏移
1096	__aicore__ inline void CalcCropIndex(uint64_t result[8])	1103	__aicore__ inline void CalcCropIndex(uint64_t result[MAX_DIMS_NUM])
1097	{	1104	{
1098	for (uint8_t dim = 0; dim < BLOCK_DIM_NUM; ++dim) {	1105	for (uint8_t dim = 0; dim < BLOCK_DIM_NUM; ++dim) {
1099	int8_t pixelDim = BS_PIXEL_MAP[BLOCK_DIM_NUM][dim];	1106	int8_t pixelDim = BS_PIXEL_MAP[BLOCK_DIM_NUM][dim];
	@@ -1122,11 +1129,10 @@ public:
1122	__aicore__ inline uint32_t CeilLog2(uint32_t input)	1129	__aicore__ inline uint32_t CeilLog2(uint32_t input)
1123	{	1130	{
1124	input--;	1131	input--;
1125	input \|= input >> 1;	1132	#pragma unroll
1126	input \|= input >> 2;	1133	for (uint16_t i=0; i < LOG_OFFSET_NUM; ++i) {
1127	input \|= input >> 4;	1134	input \|= input >> LOG_OFFSET[i];
1128	input \|= input >> 8;	1135	}
1129	input \|= input >> 16;
1130	input++;	1136	input++;
1131	uint32_t res = 0;	1137	uint32_t res = 0;
1132	while (input >>= 1) {	1138	while (input >>= 1) {

Mconversion/broadcast_to/op_host/arch35/broadcast_to_tiling_base.cpp+13-12

	@@ -21,10 +21,11 @@
21	#include "util/const_util.h"	21	#include "util/const_util.h"
22	#include "util/math_util.h"	22	#include "util/math_util.h"
23		23
24	namespace optiling	24	namespace optiling {
25	{	25	namespace brcto {
26	namespace brcto	26	static constexpr int64_t DIM_NUM_THRESHOLD_FOR_R4_SIZE = 5; // 触发尾轴4D尺寸计算的维度数阈值
27	{	27	static constexpr int64_t TRAILING_DIM_NUM_FOR_R4_SIZE = 4; // 用于4D尺寸计算的尾轴维度数
		28
28	ge::graphStatus GetABFlag(const gert::TilingContext* context, const gert::Shape& inShape, const gert::Shape& outShape,	29	ge::graphStatus GetABFlag(const gert::TilingContext* context, const gert::Shape& inShape, const gert::Shape& outShape,
29	std::array<bool, MAX_DIM_NUM>& abInfo)	30	std::array<bool, MAX_DIM_NUM>& abInfo)
30	{	31	{
	@@ -571,14 +572,14 @@ void BroadcastToTilingAscendC::CalcTensorSize()
571		572
572	int64_t ubGate = maxTensorSize_ / nTwo / nTwo;	573	int64_t ubGate = maxTensorSize_ / nTwo / nTwo;
573	int64_t r4DimSize = minTensorSize_;	574	int64_t r4DimSize = minTensorSize_;
574	if (dimNum > 5) {	575	if (dimNum > DIM_NUM_THRESHOLD_FOR_R4_SIZE) {
575	r4DimSize = CalcDimSize(inShapePtr_, dimNum - 4, dimNum);	576	r4DimSize = CalcDimSize(inShapePtr_, dimNum - TRAILING_DIM_NUM_FOR_R4_SIZE, dimNum);
576	}	577	}
577		578
578	isDMABrcA_ =	579	isDMABrcA_ =
579	(dimNum > 1 && (nTwo * outLastDim <= LAST_DIM_GATE \|\| (outLastDim == LAST_DIM_GATE / nTwo + 1 &&	580	(dimNum > 1 && (nTwo * outLastDim <= LAST_DIM_GATE \|\|
580	outShapePtr_->GetDim(dimNum - nTwo) <= LAST_DIM_GATE)	581	(outLastDim == LAST_DIM_GATE / nTwo + 1 && outShapePtr_->GetDim(dimNum - nTwo) <= LAST_DIM_GATE)
581	\|\| ((outLastDim < 8) && (r4DimSize < minTensorSize_))));	582	\|\| ((outLastDim < LAST_DIM_GATE) && (r4DimSize < minTensorSize_))));
582	if ((!abInfo_[dimNum - 1] && outLastDim <= ubGate && !isDMABrcA_) \|\|	583	if ((!abInfo_[dimNum - 1] && outLastDim <= ubGate && !isDMABrcA_) \|\|
583	(abInfo_[dimNum - 1] && outLastDim >= LAST_DIM_GATE)) { // UB broadcast	584	(abInfo_[dimNum - 1] && outLastDim >= LAST_DIM_GATE)) { // UB broadcast
584	tmpTensorSize = std::min(ubGate, MAX_TENSOR_SIZE);	585	tmpTensorSize = std::min(ubGate, MAX_TENSOR_SIZE);
	@@ -769,7 +770,7 @@ ge::graphStatus BroadcastToTilingAscendC::DoTiling()
769	return WriteTilingData();	770	return WriteTilingData();
770	}	771	}
771		772
772	} // namespace brcto	773	} // namespace brcto
773		774
774	ge::graphStatus Tiling4BroadcastToAscendC(gert::TilingContext* context, const gert::Shape* inShapePtr,	775	ge::graphStatus Tiling4BroadcastToAscendC(gert::TilingContext* context, const gert::Shape* inShapePtr,
775	const gert::Shape* outShapePtr)	776	const gert::Shape* outShapePtr)
	@@ -789,4 +790,4 @@ ge::graphStatus Tiling4BroadcastToAscendC(gert::TilingContext* context, const ge
789	return brcToTiling.DoTiling();	790	return brcToTiling.DoTiling();
790	}	791	}
791		792
792	} // namespace optiling	793	} // namespace optiling

Mconversion/broadcast_to/op_host/arch35/broadcast_to_tiling_base.h+3-1

	@@ -120,7 +120,9 @@ public:
120	return ge::GRAPH_FAILED;	120	return ge::GRAPH_FAILED;
121	}	121	}
122		122
123	auto dtype = context_->GetInputDesc(0)->GetDataType();	123	auto inputDesc = context_->GetInputDesc(0);
		124	OP_CHECK_NULL_WITH_CONTEXT(context_, inputDesc);
		125	auto dtype = inputDesc->GetDataType();
124	dtypeSize_ = GetSizeByDataType(dtype);	126	dtypeSize_ = GetSizeByDataType(dtype);
125	OP_LOGI(context_->GetNodeName(), "The ub size is: %ld", ubSize_);	127	OP_LOGI(context_->GetNodeName(), "The ub size is: %ld", ubSize_);
126		128

Mconversion/im2col/op_host/arch35/im2col_tiling_arch35.cpp+4-4

	@@ -75,7 +75,7 @@ private:
75	int32_t gatherVRegElements_{0};	75	int32_t gatherVRegElements_{0};
76		76
77	// tiling key param	77	// tiling key param
78	ge::Format inputFormat_;	78	ge::Format inputFormat_{ge::FORMAT_ND};
79	bool isSIMT_{false};	79	bool isSIMT_{false};
80	bool isPadding_{false};	80	bool isPadding_{false};
81	bool isBigShape_{false};	81	bool isBigShape_{false};
	@@ -83,7 +83,7 @@ private:
83		83
84	// 输入参数	84	// 输入参数
85	int32_t dSize_{0};	85	int32_t dSize_{0};
86	Im2ColInputInfo input_;	86	Im2ColInputInfo input_{};
87		87
88	// 中间计算结果	88	// 中间计算结果
89	// 卷积核影响HW	89	// 卷积核影响HW
	@@ -427,7 +427,7 @@ std::tuple<int32_t, int32_t> Im2ColTiling::NCHWCalcBufSize(int32_t validBufSize)
427	double ratio = static_cast<double>(tmpOutBufSize) / (tmpInBufSize + tmpOutBufSize);	427	double ratio = static_cast<double>(tmpOutBufSize) / (tmpInBufSize + tmpOutBufSize);
428	OP_LOGD(context_, "The ratio of the output buffer size to total size is %f", ratio);	428	OP_LOGD(context_, "The ratio of the output buffer size to total size is %f", ratio);
429	// 分配 buffsize	429	// 分配 buffsize
430	tmpOutBufSize = validBufSize * ratio;	430	tmpOutBufSize = static_cast<int64_t>(validBufSize * ratio);
431	// 向下对齐 vector length	431	// 向下对齐 vector length
432	tmpOutBufSize = Ops::Base::FloorAlign(tmpOutBufSize, static_cast<int64_t>(vRegSize_));	432	tmpOutBufSize = Ops::Base::FloorAlign(tmpOutBufSize, static_cast<int64_t>(vRegSize_));
433	tmpInBufSize = validBufSize - tmpOutBufSize;	433	tmpInBufSize = validBufSize - tmpOutBufSize;
	@@ -466,7 +466,7 @@ bool Im2ColTiling::NCHWTryFullLoad(int32_t validBufSize)
466		466
467	// 输入输出大小是否满足全载条件	467	// 输入输出大小是否满足全载条件
468	int64_t allNeedSize = inHWNeedSize + outHWNeedSize;	468	int64_t allNeedSize = inHWNeedSize + outHWNeedSize;
469	if (allNeedSize > validBufSize) {	469	if (allNeedSize <= 0 \|\| allNeedSize > validBufSize) {
470	return false;	470	return false;
471	}	471	}
472	// 输入1个补pad后的HW的元素个数不能超出gather索引大小	472	// 输入1个补pad后的HW的元素个数不能超出gather索引大小

Mconversion/matrix_set_diag/op_host/arch35/matrix_set_diag_tiling_arch35.cpp+17-10

	@@ -21,9 +21,13 @@
21	#include "exe_graph/runtime/runtime_attrs.h"	21	#include "exe_graph/runtime/runtime_attrs.h"
22		22
23	namespace optiling {	23	namespace optiling {
24	// NCHW 常量	24	// 公共常量
25	// BUFFER分割数量	25	// BUFFER分割数量
26	static constexpr uint32_t BUFFER_NUM = 2;	26	static constexpr uint32_t BUFFER_NUM = 2;
		27	// UB对齐预留块数
		28	static constexpr uint32_t RESERVED_ALIGN_BLOCK_COUNT = 2;
		29	static constexpr uint32_t COL_DIM_OFFSET = 1; // 列维度（-1轴）距shape末尾的偏移量
		30	static constexpr uint32_t ROW_DIM_OFFSET = 2; // 行维度（-2轴）距shape末尾的偏移量
27		31
28	static constexpr uint8_t MIN_INPUT_DIMNUM = 2;	32	static constexpr uint8_t MIN_INPUT_DIMNUM = 2;
29	static constexpr uint8_t MAX_INPUT_DIMNUM = 8;	33	static constexpr uint8_t MAX_INPUT_DIMNUM = 8;
	@@ -32,6 +36,8 @@ static constexpr double MIN_USED_CORES_RATIO = 0.8;
32	static constexpr int64_t MIN_PER_UB_SIZE = 4096;	36	static constexpr int64_t MIN_PER_UB_SIZE = 4096;
33	// UB内 scatter 操作的最大元素个数	37	// UB内 scatter 操作的最大元素个数
34	static constexpr int32_t MAX_UB_SCATTER_ELEMENT_NUM = std::numeric_limits<uint16_t>::max();	38	static constexpr int32_t MAX_UB_SCATTER_ELEMENT_NUM = std::numeric_limits<uint16_t>::max();
		39	// scatter操作元素数限制的最小数据类型字节数
		40	static constexpr int32_t MIN_DSIZE_FOR_SCATTER_LIMIT = 2;
35		41
36	// SIMT 常量	42	// SIMT 常量
37	static constexpr int64_t MAX_SHAPE_SIZE_FOR_SIMT = 1024;	43	static constexpr int64_t MAX_SHAPE_SIZE_FOR_SIMT = 1024;
	@@ -67,7 +73,7 @@ private:
67		73
68	// 输入参数	74	// 输入参数
69	int32_t dSize_{0};	75	int32_t dSize_{0};
70	MatrixSetDiagTilingData* tilingData_;	76	MatrixSetDiagTilingData* tilingData_{nullptr};
71		77
72	// tiling context	78	// tiling context
73	gert::TilingContext* context_;	79	gert::TilingContext* context_;
	@@ -200,10 +206,10 @@ ge::graphStatus MatrixSetDiagTiling::ParamCheck()
200	"diagonal dim num must equal input dim num minus 1"),	206	"diagonal dim num must equal input dim num minus 1"),
201	return ge::GRAPH_FAILED);	207	return ge::GRAPH_FAILED);
202		208
203	xColNum_ = inputShapeVal.GetDim(dimNum_ - 1);	209	xColNum_ = inputShapeVal.GetDim(dimNum_ - COL_DIM_OFFSET);
204	xRowNum_ = inputShapeVal.GetDim(dimNum_ - 2);	210	xRowNum_ = inputShapeVal.GetDim(dimNum_ - ROW_DIM_OFFSET);
205	tailAxisDataSize_ = xColNum_ * xRowNum_;	211	tailAxisDataSize_ = xColNum_ * xRowNum_;
206	diagLen_ = diagShapeVal.GetDim(diagDimNum_ - 1);	212	diagLen_ = diagShapeVal.GetDim(static_cast<size_t>(diagDimNum_ - 1));
207	OP_CHECK_IF(	213	OP_CHECK_IF(
208	diagLen_ != std::min(xColNum_, xRowNum_),	214	diagLen_ != std::min(xColNum_, xRowNum_),
209	OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(	215	OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(
	@@ -227,7 +233,7 @@ ge::graphStatus MatrixSetDiagTiling::ParamCheck()
227		233
228	void MatrixSetDiagTiling::CalUbFactor()	234	void MatrixSetDiagTiling::CalUbFactor()
229	{	235	{
230	uint64_t validBufSize = bufferSize_ - ubBlockSize_ * 2;	236	uint64_t validBufSize = bufferSize_ - ubBlockSize_ * RESERVED_ALIGN_BLOCK_COUNT;
231	if (isCutW_) {	237	if (isCutW_) {
232	if (xColNum_ * dSize_ >= bufferSize_) {	238	if (xColNum_ * dSize_ >= bufferSize_) {
233	ubFactor_ = validBufSize / dSize_;	239	ubFactor_ = validBufSize / dSize_;
	@@ -241,7 +247,7 @@ void MatrixSetDiagTiling::CalUbFactor()
241	}	247	}
242	} else {	248	} else {
243	uint64_t totalTailSize = (tailAxisDataSize_ + diagLen_) * dSize_;	249	uint64_t totalTailSize = (tailAxisDataSize_ + diagLen_) * dSize_;
244	ubFactor_ = validBufSize >= totalTailSize ? validBufSize / totalTailSize : 1;	250	ubFactor_ = validBufSize >= totalTailSize ? Ops::Base::FloorDiv(validBufSize, totalTailSize) : 1;
245	}	251	}
246	}	252	}
247		253
	@@ -264,7 +270,7 @@ ge::graphStatus MatrixSetDiagTiling::Tiling4CutW()
264	isCutW_ = true;	270	isCutW_ = true;
265	CalUbFactor();	271	CalUbFactor();
266	OP_CHECK_IF((ubFactor_ == 0U), OP_LOGE(context_, "ubFactor is 0"), return ge::GRAPH_FAILED);	272	OP_CHECK_IF((ubFactor_ == 0U), OP_LOGE(context_, "ubFactor is 0"), return ge::GRAPH_FAILED);
267	if (dSize_ <= 2) {	273	if (dSize_ <= MIN_DSIZE_FOR_SCATTER_LIMIT) {
268	ubFactor_ = ubFactor_ < MAX_UB_SCATTER_ELEMENT_NUM ? ubFactor_ : MAX_UB_SCATTER_ELEMENT_NUM;	274	ubFactor_ = ubFactor_ < MAX_UB_SCATTER_ELEMENT_NUM ? ubFactor_ : MAX_UB_SCATTER_ELEMENT_NUM;
269	}	275	}
270	// 设置核数	276	// 设置核数
	@@ -283,7 +289,7 @@ ge::graphStatus MatrixSetDiagTiling::Tiling4NoCutW()
283	{	289	{
284	CalUbFactor();	290	CalUbFactor();
285	OP_CHECK_IF((ubFactor_ == 0U), OP_LOGE(context_, "ubFactor is 0"), return ge::GRAPH_FAILED);	291	OP_CHECK_IF((ubFactor_ == 0U), OP_LOGE(context_, "ubFactor is 0"), return ge::GRAPH_FAILED);
286	if (dSize_ <= 2) {	292	if (dSize_ <= MIN_DSIZE_FOR_SCATTER_LIMIT) {
287	ubFactor_ = ubFactor_ * tailAxisDataSize_ < MAX_UB_SCATTER_ELEMENT_NUM ?	293	ubFactor_ = ubFactor_ * tailAxisDataSize_ < MAX_UB_SCATTER_ELEMENT_NUM ?
288	ubFactor_ :	294	ubFactor_ :
289	MAX_UB_SCATTER_ELEMENT_NUM / tailAxisDataSize_;	295	MAX_UB_SCATTER_ELEMENT_NUM / tailAxisDataSize_;
	@@ -381,7 +387,8 @@ ge::graphStatus MatrixSetDiagTiling::Tiling4MatrixSetDiag()
381	uint64_t totalTailSize = (AlignBlock(tailAxisDataSize_) + AlignBlock(diagLen_)) * dSize_;	387	uint64_t totalTailSize = (AlignBlock(tailAxisDataSize_) + AlignBlock(diagLen_)) * dSize_;
382	bufferSize_ = ubSize_ / BUFFER_NUM - vectorSize_;	388	bufferSize_ = ubSize_ / BUFFER_NUM - vectorSize_;
383	OP_LOGI(context_, "bufferSize_ %lu, totalTailSize %lu", bufferSize_, totalTailSize);	389	OP_LOGI(context_, "bufferSize_ %lu, totalTailSize %lu", bufferSize_, totalTailSize);
384	if (totalTailSize >= bufferSize_ \|\| (dSize_ <= 2 && tailAxisDataSize_ >= MAX_UB_SCATTER_ELEMENT_NUM)) {	390	if (totalTailSize >= bufferSize_ \|\|
		391	(dSize_ <= MIN_DSIZE_FOR_SCATTER_LIMIT && tailAxisDataSize_ >= MAX_UB_SCATTER_ELEMENT_NUM)) {
385	return Tiling4CutW();	392	return Tiling4CutW();
386	} else {	393	} else {
387	return Tiling4NoCutW();	394	return Tiling4NoCutW();

Mconversion/mem_set_v2/op_host/arch35/mem_set_v2_tiling_arch35.cpp+1-1

	@@ -207,7 +207,7 @@ ge::graphStatus DoTiling(gert::TilingContext* context, uint64_t tilingKey, int64
207	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);	207	ascendcPlatform.GetCoreMemSize(platform_ascendc::CoreMemType::UB, ubSize);
208	} else {	208	} else {
209	OP_LOGD(context->GetNodeName(), "Get compileInfo with tilingParse.");	209	OP_LOGD(context->GetNodeName(), "Get compileInfo with tilingParse.");
210	auto compileInfo = reinterpret_cast<const MemSetV2CompileInfo*>(context->GetCompileInfo());	210	auto compileInfo = context->GetCompileInfo<MemSetV2CompileInfo>();
211	OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);	211	OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
212	ubSize = compileInfo->ubSize;	212	ubSize = compileInfo->ubSize;
213	coreNum = compileInfo->coreNum;	213	coreNum = compileInfo->coreNum;

Mconversion/pad_v3/op_api/aclnn_constant_pad_nd.cpp+1-1

	@@ -172,7 +172,7 @@ static bool Checkformat(const aclTensor* self, const aclTensor* out)
172	return true;	172	return true;
173	}	173	}
174		174
175	static bool CheckPadForFp8(const aclTensor* self, uint32_t& signSymbol)	175	static bool CheckPadForFp8(const aclTensor* self, const uint32_t& signSymbol)
176	{	176	{
177	// self的数据类型为fp8时，pad数组中不能有负数，StridedSlice不支持FLOAT8_E8M0类型	177	// self的数据类型为fp8时，pad数组中不能有负数，StridedSlice不支持FLOAT8_E8M0类型
178	if (CheckType(self->GetDataType(), DTYPE_SUPPORT_FP8_FP4_LIST) && (signSymbol & NEGETIVE) == NEGETIVE) {	178	if (CheckType(self->GetDataType(), DTYPE_SUPPORT_FP8_FP4_LIST) && (signSymbol & NEGETIVE) == NEGETIVE) {

Mconversion/pad_v3/op_host/arch35/pad_v3_tiling_arch35.cpp+2-2

Mconversion/pad_v3_grad/op_host/arch35/pad_v3_grad_tiling_arch35.cpp+1-1

	@@ -311,7 +311,7 @@ void PadV3GradACTiling::CalculateTilingKeyMirror()
311	cutMode_ = TPL_SIMD_BIG;	311	cutMode_ = TPL_SIMD_BIG;
312	bufferSize_ =	312	bufferSize_ =
313	GetSizeOfBlockAlign(ubSize_ / (CONST2 * dtypeBytes_ + CONST4 * FP32_SIZE) - alignNum, alignNum);	313	GetSizeOfBlockAlign(ubSize_ / (CONST2 * dtypeBytes_ + CONST4 * FP32_SIZE) - alignNum, alignNum);
314	ubAxis_ = dimNum_ - 1;	314	ubAxis_ = static_cast<uint8_t>(static_cast<int8_t>(dimNum_) - 1);
315	ubFactor_ = tilingData_->inShape[dimNum_ - 1];	315	ubFactor_ = tilingData_->inShape[dimNum_ - 1];
316	outTileSize_ = bufferSize_;	316	outTileSize_ = bufferSize_;
317	return TilingInfoTune();	317	return TilingInfoTune();

Mconversion/pad_v3_grad/op_kernel/arch35/pad_v3_grad_circular_huge_width.h+10-5

	@@ -26,6 +26,11 @@ using namespace AscendC;
26	template <typename T>	26	template <typename T>
27	class KernelPadV3GradCircularHugeWidth {	27	class KernelPadV3GradCircularHugeWidth {
28	private:	28	private:
		29	constexpr static size_t DIM_IDX_4 = 4U;
		30	constexpr static size_t GM_ADDR_IDX_CENTER = 0U;
		31	constexpr static size_t GM_ADDR_IDX_LEFT = 1U;
		32	constexpr static size_t GM_ADDR_IDX_RIGHT = 2U;
		33
29	uint32_t inResStart_;	34	uint32_t inResStart_;
30	uint32_t inSrcStart_;	35	uint32_t inSrcStart_;
31		36
	@@ -199,7 +204,7 @@ private:
199	} else {	204	} else {
200	leftUbStartIdx_ = (outIndex_[ubAxis_] < leftStartOnInner_) ? (leftStartOnInner_ - outIndex_[ubAxis_]) : 0;	205	leftUbStartIdx_ = (outIndex_[ubAxis_] < leftStartOnInner_) ? (leftStartOnInner_ - outIndex_[ubAxis_]) : 0;
201	leftUbAddLen_ = dataLen_ - leftUbStartIdx_;	206	leftUbAddLen_ = dataLen_ - leftUbStartIdx_;
202	inIdxCnt[4].inGmIdx[1] = outIndex_[ubAxis_] + leftUbStartIdx_ - leftStartOnInner_;	207	inIdxCnt[DIM_IDX_4].inGmIdx[GM_ADDR_IDX_LEFT] = outIndex_[ubAxis_] + leftUbStartIdx_ - leftStartOnInner_;
203	}	208	}
204		209
205	if (tilingData_->rightPad[ubAxis_] == 0 \|\| outIndex_[ubAxis_] >= tilingData_->rightPad[ubAxis_]) {	210	if (tilingData_->rightPad[ubAxis_] == 0 \|\| outIndex_[ubAxis_] >= tilingData_->rightPad[ubAxis_]) {
	@@ -208,14 +213,14 @@ private:
208	rightUbAddLen_ = (outIndex_[ubAxis_] + dataLen_ <= tilingData_->rightPad[ubAxis_]) ?	213	rightUbAddLen_ = (outIndex_[ubAxis_] + dataLen_ <= tilingData_->rightPad[ubAxis_]) ?
209	dataLen_ :	214	dataLen_ :
210	tilingData_->rightPad[ubAxis_] - outIndex_[ubAxis_];	215	tilingData_->rightPad[ubAxis_] - outIndex_[ubAxis_];
211	inIdxCnt[4].inGmIdx[2] =	216	inIdxCnt[DIM_IDX_4].inGmIdx[GM_ADDR_IDX_RIGHT] =
212	outIndex_[ubAxis_] + tilingData_->leftPad[ubAxis_] + tilingData_->outShape[ubAxis_];	217	outIndex_[ubAxis_] + tilingData_->leftPad[ubAxis_] + tilingData_->outShape[ubAxis_];
213	// 当前块的右pad在输入中的索引	218	// 当前块的右pad在输入中的索引
214	}	219	}
215		220
216	for (uint8_t i = 0; i < dimNum_ - 1; ++i) {	221	for (uint8_t i = 0; i < dimNum_ - 1; ++i) {
217	// self	222	// self
218	inIdxCnt[i].inGmIdx[0] = inIndex_[i] * tilingData_->inStride[i];	223	inIdxCnt[i].inGmIdx[GM_ADDR_IDX_CENTER] = inIndex_[i] * tilingData_->inStride[i];
219	// left	224	// left
220	if (tilingData_->leftPad[i] != 0 && outIndex_[i] >= tilingData_->outShape[i] - tilingData_->leftPad[i]) {	225	if (tilingData_->leftPad[i] != 0 && outIndex_[i] >= tilingData_->outShape[i] - tilingData_->leftPad[i]) {
221	inIdxCnt[i].inGmIdx[inIdxCnt[i].cnt++] =	226	inIdxCnt[i].inGmIdx[inIdxCnt[i].cnt++] =
	@@ -315,7 +320,7 @@ private:
315	__ubuf__ T* srcAddr, __ubuf__ CalType* resAddr, LocalTensor<T> src, uint32_t idx, uint64_t a3Offset,	320	__ubuf__ T* srcAddr, __ubuf__ CalType* resAddr, LocalTensor<T> src, uint32_t idx, uint64_t a3Offset,
316	IdxAndTimes* inIdxCnt)	321	IdxAndTimes* inIdxCnt)
317	{	322	{
318	uint64_t inLeftAddr = a3Offset + inIdxCnt[4].inGmIdx[1];	323	uint64_t inLeftAddr = a3Offset + inIdxCnt[DIM_IDX_4].inGmIdx[GM_ADDR_IDX_LEFT];
319	copyInParams_.blockLen = leftUbAddLen_ * sizeof(T);	324	copyInParams_.blockLen = leftUbAddLen_ * sizeof(T);
320		325
321	uint16_t leftMainTimes = leftUbAddLen_ / oneRepeatSize_;	326	uint16_t leftMainTimes = leftUbAddLen_ / oneRepeatSize_;
	@@ -390,7 +395,7 @@ private:
390	__ubuf__ T* srcAddr, __ubuf__ CalType* resAddr, LocalTensor<T> src, uint32_t idx, uint64_t a3Offset,	395	__ubuf__ T* srcAddr, __ubuf__ CalType* resAddr, LocalTensor<T> src, uint32_t idx, uint64_t a3Offset,
391	IdxAndTimes* inIdxCnt)	396	IdxAndTimes* inIdxCnt)
392	{	397	{
393	uint64_t inRightAddr = a3Offset + inIdxCnt[4].inGmIdx[2];	398	uint64_t inRightAddr = a3Offset + inIdxCnt[DIM_IDX_4].inGmIdx[GM_ADDR_IDX_RIGHT];
394	copyInParams_.blockLen = rightUbAddLen_ * sizeof(T);	399	copyInParams_.blockLen = rightUbAddLen_ * sizeof(T);
395		400
396	uint16_t rightMainTimes = CeilDiv(rightUbAddLen_, oneRepeatSize_);	401	uint16_t rightMainTimes = CeilDiv(rightUbAddLen_, oneRepeatSize_);

Mconversion/strided_slice/op_host/arch35/strided_slice_tiling_arch35.cpp+1-1

	@@ -211,7 +211,7 @@ void StrideSliceTiling::CalcBlockSplitInfo()
211	// 说明整个shape都很小，blk不切最后一根轴	211	// 说明整个shape都很小，blk不切最后一根轴
212	if (blkIndex_ == static_cast<int64_t>(-1)) {	212	if (blkIndex_ == static_cast<int64_t>(-1)) {
213	blkIndex_ = static_cast<int64_t>(dimNum_ - 1);	213	blkIndex_ = static_cast<int64_t>(dimNum_ - 1);
214	blkFactor_ = outputShape.GetDim(dimNum_ - 1);	214	blkFactor_ = outputShape.GetDim(static_cast<size_t>(dimNum_ - 1));
215	}	215	}
216		216
217	blkTailFactor_ = outputShape.GetDim(blkIndex_) % blkFactor_;	217	blkTailFactor_ = outputShape.GetDim(blkIndex_) % blkFactor_;

Mconversion/strided_slice_v2/op_host/arch35/strided_slice_v2_tiling_arch35.cpp+18-17

	@@ -14,6 +14,7 @@
14	*/	14	*/
15	#include "strided_slice_v2_tiling_arch35.h"	15	#include "strided_slice_v2_tiling_arch35.h"
16	#include <numeric>	16	#include <numeric>
		17	#include "common/inc/op_host/math_log.h"
17		18
18	namespace {	19	namespace {
19	const std::string OP_NAME = "StridedSliceV2";	20	const std::string OP_NAME = "StridedSliceV2";
	@@ -42,7 +43,7 @@ static void ConstructSliceShape(const gert::StorageShape *storage, gert::Shape &
42	}	43	}
43		44
44	static bool CheckStride(ops::QuickVector &stride, const gert::TilingContext *context) {	45	static bool CheckStride(ops::QuickVector &stride, const gert::TilingContext *context) {
45	auto compileInfo = reinterpret_cast<const StridedSliceV2CompileInfo *>(context->GetCompileInfo());	46	auto compileInfo = context->GetCompileInfo<StridedSliceV2CompileInfo>();
46	OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);	47	OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
47		48
48	for (size_t i = 0; i < stride.GetDimNum(); i++) {	49	for (size_t i = 0; i < stride.GetDimNum(); i++) {
	@@ -405,6 +406,15 @@ ge::graphStatus StrideSliceV2TilingForAscendC(gert::TilingContext* context, int6
405	return tilingObject.RunStrideSliceTiling();	406	return tilingObject.RunStrideSliceTiling();
406	}	407	}
407		408
		409	static ge::graphStatus GetMaskValue(
		410	gert::TilingContext* context, const gert::RuntimeAttrs* attrs, size_t attrIndex, uint64_t& maskValue)
		411	{
		412	const int64_t* mask = attrs->GetAttrPointer<int64_t>(attrIndex);
		413	OP_CHECK_NULL_WITH_CONTEXT(context, mask);
		414	maskValue = static_cast<uint64_t>(*mask);
		415	return ge::GRAPH_SUCCESS;
		416	}
		417
408	ge::graphStatus Tiling4StridedSliceV2(gert::TilingContext *context) {	418	ge::graphStatus Tiling4StridedSliceV2(gert::TilingContext *context) {
409	SliceParametersRuntime2 sliceParam;	419	SliceParametersRuntime2 sliceParam;
410	if (ConstructSliceParam(context, sliceParam) != ge::GRAPH_SUCCESS) {	420	if (ConstructSliceParam(context, sliceParam) != ge::GRAPH_SUCCESS) {
	@@ -415,21 +425,12 @@ ge::graphStatus Tiling4StridedSliceV2(gert::TilingContext *context) {
415	auto attrs = context->GetAttrs();	425	auto attrs = context->GetAttrs();
416	OP_CHECK_NULL_WITH_CONTEXT(context, attrs);	426	OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
417		427
418	uint64_t beginValue, endValue, ellipsisValue, newAxisValue, shrinkAxisValue;	428	uint64_t beginValue{0}, endValue{0}, ellipsisValue{0}, newAxisValue{0}, shrinkAxisValue{0};
419		429	CHECK_RET_SUCC(GetMaskValue(context, attrs, IDX_MASK_BEGIN, beginValue));
420	#define GET_MASK_VALUE(index, mask_name) \	430	CHECK_RET_SUCC(GetMaskValue(context, attrs, IDX_MASK_END, endValue));
421	do { \	431	CHECK_RET_SUCC(GetMaskValue(context, attrs, IDX_MASK_ELLIPSIS, ellipsisValue));
422	const int64_t* mask_##mask_name = attrs->GetAttrPointer<int64_t>(index); \	432	CHECK_RET_SUCC(GetMaskValue(context, attrs, IDX_MASK_NEW_AXIS, newAxisValue));
423	OP_CHECK_NULL_WITH_CONTEXT(context, mask_##mask_name); \	433	CHECK_RET_SUCC(GetMaskValue(context, attrs, IDX_MASK_SHRINK_AXIS, shrinkAxisValue));
424	mask_name##Value = static_cast<uint64_t>(*mask_##mask_name); \
425	} while(0)
426
427	GET_MASK_VALUE(IDX_MASK_BEGIN, begin);
428	GET_MASK_VALUE(IDX_MASK_END, end);
429	GET_MASK_VALUE(IDX_MASK_ELLIPSIS, ellipsis);
430	GET_MASK_VALUE(IDX_MASK_NEW_AXIS, newAxis);
431	GET_MASK_VALUE(IDX_MASK_SHRINK_AXIS, shrinkAxis);
432	#undef GET_MASK_VALUE
433		434
434	// Infer shape	435	// Infer shape
435	const gert::StorageShape *xStorage = context->GetInputShape(INDEX_X);	436	const gert::StorageShape *xStorage = context->GetInputShape(INDEX_X);
	@@ -470,7 +471,7 @@ ge::graphStatus Tiling4StridedSliceV2(gert::TilingContext *context) {
470	OP_LOGI(context->GetNodeName(), "perf slice params: %s", sliceParam.to_string().c_str());	471	OP_LOGI(context->GetNodeName(), "perf slice params: %s", sliceParam.to_string().c_str());
471		472
472	// Infer tiling mode	473	// Infer tiling mode
473	auto compileInfo = reinterpret_cast<const StridedSliceV2CompileInfo *>(context->GetCompileInfo());	474	auto compileInfo = context->GetCompileInfo<StridedSliceV2CompileInfo>();
474	OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);	475	OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
475	OP_LOGD(context->GetNodeName(), "compile info: %s.", compileInfo->to_string().c_str());	476	OP_LOGD(context->GetNodeName(), "compile info: %s.", compileInfo->to_string().c_str());
476		477

Mconversion/strided_slice_v3/op_host/arch35/strided_slice_v3_tiling_arch35.cpp+4-5

	@@ -245,10 +245,9 @@ static int64_t GetTilingMode4StridesLargerThanOne(const SliceParameters& paramet
245	int64_t multiTimes = dtypeSize / DTYPE_SIZE_FP16 > 1 ? dtypeSize / DTYPE_SIZE_FP16 : 1;	245	int64_t multiTimes = dtypeSize / DTYPE_SIZE_FP16 > 1 ? dtypeSize / DTYPE_SIZE_FP16 : 1;
246	int64_t vconvUbSize = alignUbSize * multiTimes / 2;	246	int64_t vconvUbSize = alignUbSize * multiTimes / 2;
247	int64_t multiBlockElement = blockElement * multiTimes;	247	int64_t multiBlockElement = blockElement * multiTimes;
248	int64_t vconvColumn = dtypeSize % DTYPE_SIZE_FP16 == 0 ?	248	int64_t vconvColumn = dtypeSize % DTYPE_SIZE_FP16 == 0 ? VNCHWCONV_COLUMNS_FP16 : VNCHWCONV_COLUMNS_INT8;
249	VNCHWCONV_COLUMNS_FP16 : VNCHWCONV_COLUMNS_INT8;	249	int64_t alignColSize = Ops::Base::CeilAlign(inputInner * output32bytesAlignRows, multiBlockElement);
250	int64_t maxRowsInUb = vconvUbSize / vconvColumn / Ops::Base::CeilAlign(inputInner *	250	int64_t maxRowsInUb = Ops::Base::FloorDiv(vconvUbSize / vconvColumn, alignColSize) * output32bytesAlignRows;
251	output32bytesAlignRows, multiBlockElement) * output32bytesAlignRows;
252	bool isValidLen = outputInner - multiTimes > 0;	251	bool isValidLen = outputInner - multiTimes > 0;
253	int64_t validLen = (outputInner - multiTimes) * lastStride + multiTimes;	252	int64_t validLen = (outputInner - multiTimes) * lastStride + multiTimes;
254	int64_t dstStride = (lastStride - 1) * multiTimes;	253	int64_t dstStride = (lastStride - 1) * multiTimes;
	@@ -680,7 +679,7 @@ static ge::graphStatus TilingForStridedSliceV3(gert::TilingContext* context) {
680	ConstructStrideList(context->GetOptionalInputTensor(INDEX_STRIDES), inputDimNum, newAxes,	679	ConstructStrideList(context->GetOptionalInputTensor(INDEX_STRIDES), inputDimNum, newAxes,
681	sliceParam.stride_list);	680	sliceParam.stride_list);
682		681
683	auto compileInfo = reinterpret_cast<const StridedSliceV3CompileInfo*>(context->GetCompileInfo());	682	auto compileInfo = context->GetCompileInfo<StridedSliceV3CompileInfo>();
684	OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);	683	OP_CHECK_NULL_WITH_CONTEXT(context, compileInfo);
685		684
686	OP_LOGD(OP_NAME.c_str(), "before make performance, slice params: %s", sliceParam.to_string().c_str());	685	OP_LOGD(OP_NAME.c_str(), "before make performance, slice params: %s", sliceParam.to_string().c_str());

Mconversion/trans_data/op_host/arch35/transdata_tiling_arch35.cpp+4-4

	@@ -29,7 +29,7 @@ namespace transdata_asc {
29		29
30	ge::graphStatus TransDataTilingAscendC::GetHardwareInfo()	30	ge::graphStatus TransDataTilingAscendC::GetHardwareInfo()
31	{	31	{
32	auto compileInfo = reinterpret_cast<const TransDataCompileInfo*>(context_->GetCompileInfo());	32	auto compileInfo = context_->GetCompileInfo<TransDataCompileInfo>();
33	OP_CHECK_NULL_WITH_CONTEXT(context_, compileInfo);	33	OP_CHECK_NULL_WITH_CONTEXT(context_, compileInfo);
34	coreNum_ = static_cast<uint32_t>(compileInfo->coreNum);	34	coreNum_ = static_cast<uint32_t>(compileInfo->coreNum);
35	ubSize_ = compileInfo->ubSize;	35	ubSize_ = compileInfo->ubSize;
	@@ -141,10 +141,10 @@ bool TransDataTilingAscendC::CalcNzToNdShapeSize()
141	outShape_ = Ops::Base::EnsureNotScalar(yStorage->GetStorageShape());	141	outShape_ = Ops::Base::EnsureNotScalar(yStorage->GetStorageShape());
142		142
143	OP_CHECK_IF(	143	OP_CHECK_IF(
144	inShape_.GetDimNum() < (int64_t)nFour,	144	inShape_.GetDimNum() < nFour, OP_LOGE(context_->GetNodeName(), "The input shape dim is less than 4!"),
145	OP_LOGE(context_->GetNodeName(), "The input shape dim is less than 4!"), return false);	145	return false);
146	OP_CHECK_IF(	146	OP_CHECK_IF(
147	inShape_.GetDimNum() - outShape_.GetDimNum() != (int64_t)nTwo,	147	inShape_.GetDimNum() != outShape_.GetDimNum() + nTwo,
148	OP_LOGE(context_->GetNodeName(), "The input dim must be 2 higher than the output dim."), return false);	148	OP_LOGE(context_->GetNodeName(), "The input dim must be 2 higher than the output dim."), return false);
149		149
150	auto inputDimCnt = inShape_.GetDimNum();	150	auto inputDimCnt = inShape_.GetDimNum();

Mconversion/trans_data/op_host/arch35/transdata_tiling_arch35.h+20-20

	@@ -34,7 +34,7 @@ TILING_DATA_FIELD_DEF(int64_t, c0);
34	TILING_DATA_FIELD_DEF(int64_t, h);	34	TILING_DATA_FIELD_DEF(int64_t, h);
35	TILING_DATA_FIELD_DEF(int64_t, n);	35	TILING_DATA_FIELD_DEF(int64_t, n);
36	TILING_DATA_FIELD_DEF(int64_t, c);	36	TILING_DATA_FIELD_DEF(int64_t, c);
37	TILING_DATA_FIELD_DEF(int64_t, tNum); // thread number	37	TILING_DATA_FIELD_DEF(int64_t, tNum); // thread number
38	END_TILING_DATA_DEF;	38	END_TILING_DATA_DEF;
39		39
40	BEGIN_TILING_DATA_DEF(TransDataNzToNdTilingData)	40	BEGIN_TILING_DATA_DEF(TransDataNzToNdTilingData)
	@@ -74,7 +74,7 @@ constexpr int64_t SIMT_RSV_SIZE = 128 * 1024L;
74		74
75	class TransDataTilingAscendC {	75	class TransDataTilingAscendC {
76	public:	76	public:
77	explicit TransDataTilingAscendC(gert::TilingContext* context) : context_(context){};	77	explicit TransDataTilingAscendC(gert::TilingContext* context) : context_(context) {};
78	ge::graphStatus DoTiling();	78	ge::graphStatus DoTiling();
79	ge::graphStatus DoNz2NdTiling();	79	ge::graphStatus DoNz2NdTiling();
80	ge::graphStatus GetHardwareInfo();	80	ge::graphStatus GetHardwareInfo();
	@@ -102,27 +102,27 @@ private:
102	TransDataNzToNdTilingData tilingNzToNdData_;	102	TransDataNzToNdTilingData tilingNzToNdData_;
103	gert::Shape inShape_;	103	gert::Shape inShape_;
104	gert::Shape outShape_;	104	gert::Shape outShape_;
105	ge::Format dstFormat_;	105	ge::Format dstFormat_{ge::FORMAT_MD};
106	size_t dtypeSize_;	106	size_t dtypeSize_{0};
107	ge::DataType srcDtype_;	107	ge::DataType srcDtype_{ge::DT_UNDEFINED};
108		108
109	uint32_t coreNum_{1};	109	uint32_t coreNum_{1};
110	uint32_t bNum_;	110	uint32_t bNum_{0};
111	int64_t ubSize_;	111	int64_t ubSize_{0};
112	int64_t blockSize_;	112	int64_t blockSize_{0};
113	int64_t tilingKey_{TILING_MODE_SIMT};	113	int64_t tilingKey_{TILING_MODE_SIMT};
114		114
115	int64_t expectC0_;	115	int64_t expectC0_{0};
116	int64_t c0_;	116	int64_t c0_{0};
117	int64_t h_;	117	int64_t h_{0};
118	int64_t n_;	118	int64_t n_{0};
119	int64_t c_;	119	int64_t c_{0};
120	int64_t c1_;	120	int64_t c1_{0};
121	int64_t n0_;	121	int64_t n0_{0};
122	int64_t n1_;	122	int64_t n1_{0};
123	int64_t tNum_;	123	int64_t tNum_{0};
124	};	124	};
125	} // namespace transdata_asc	125	} // namespace transdata_asc
126		126
127	} // namespace optiling	127	} // namespace optiling
128	#endif // OPS_MATH_CONVERSION_TRANSDATA_TILING_ARCH35_H_	128	#endif // OPS_MATH_CONVERSION_TRANSDATA_TILING_ARCH35_H_

Mmath/arg_max_with_value/op_host/arch35/arg_common_base_tiling_arch35.cpp+4-3

Mmath/sinc/op_host/sinc_def.cpp+0-1

	@@ -13,7 +13,6 @@
13	* \brief sinc def	13	* \brief sinc def
14	*/	14	*/
15		15
16	#include <cstdint>
17	#include "register/op_def_registry.h"	16	#include "register/op_def_registry.h"
18		17
19	namespace ops {	18	namespace ops {

Mmath/truncate_div/op_kernel/truncate_div_apt.cpp+133-57

	@@ -8,11 +8,6 @@
8	* See LICENSE in the root of the software repository for the full text of the License.	8	* See LICENSE in the root of the software repository for the full text of the License.
9	*/	9	*/
10		10
11	/*!
12	* \file truncate_div_apt.cpp
13	* \brief truncate_div kernel
14	*/
15
16	#include "kernel_operator.h"	11	#include "kernel_operator.h"
17	#include "arch35/truncate_div_dag.h"	12	#include "arch35/truncate_div_dag.h"
18	#include "arch35/truncate_div_struct.h"	13	#include "arch35/truncate_div_struct.h"
	@@ -22,9 +17,9 @@ using namespace AscendC;
22	using namespace Ops::Base;	17	using namespace Ops::Base;
23		18
24	template <uint64_t schMode, bool canUseMul>	19	template <uint64_t schMode, bool canUseMul>
25	__global__ __aicore__ void truncate_div(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)	20	__aicore__ inline void truncate_div_half(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR tiling)
26	{	21	{
27	if constexpr (std::is_same<DTYPE_X1, half>::value && std::is_same<DTYPE_X2, float>::value) {	22	if constexpr (std::is_same<DTYPE_X2, float>::value) {
28	if constexpr (canUseMul) {	23	if constexpr (canUseMul) {
29	using OpDag = TruncateDivOp::TruncateDivFloatWithCastScalar<half, float, float>::OpDag;	24	using OpDag = TruncateDivOp::TruncateDivFloatWithCastScalar<half, float, float>::OpDag;
30	BroadcastSch<schMode, OpDag> sch(tiling);	25	BroadcastSch<schMode, OpDag> sch(tiling);
	@@ -34,67 +29,148 @@ __global__ __aicore__ void truncate_div(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_AD
34	BroadcastSch<schMode, OpDag> sch(tiling);	29	BroadcastSch<schMode, OpDag> sch(tiling);
35	sch.Process(x1, x2, y);	30	sch.Process(x1, x2, y);
36	}	31	}
37	} else if constexpr (std::is_same<DTYPE_X1, half>::value \|\| std::is_same<DTYPE_X1, bfloat16_t>::value) {	32	} else if constexpr (std::is_same<DTYPE_X2, half>::value) {
38	if constexpr (std::is_same<DTYPE_X2, DTYPE_X1>::value) {	33	if constexpr (canUseMul) {
39	if constexpr (canUseMul) {	34	using OpDag = TruncateDivOp::TruncateDivFloat16Scalar<half, float>::OpDag;
40	using OpDag = TruncateDivOp::TruncateDivFloat16Scalar<DTYPE_X1, float>::OpDag;
41	BroadcastSch<schMode, OpDag> sch(tiling);
42	sch.Process(x1, y);
43	} else {
44	using OpDag = TruncateDivOp::TruncateDivFloat16<DTYPE_X1, float>::OpDag;
45	BroadcastSch<schMode, OpDag> sch(tiling);
46	sch.Process(x1, x2, y);
47	}
48	}
49	} else if constexpr (std::is_same<DTYPE_X1, float>::value) {
50	if constexpr (std::is_same<DTYPE_X2, float>::value) {
51	if constexpr (canUseMul) {
52	using OpDag = TruncateDivOp::TruncateDivFloatScalar<float>::OpDag;
53	BroadcastSch<schMode, OpDag> sch(tiling);
54	sch.Process(x1, y);
55	} else {
56	using OpDag = TruncateDivOp::TruncateDivFloat<float>::OpDag;
57	BroadcastSch<schMode, OpDag> sch(tiling);
58	sch.Process(x1, x2, y);
59	}
60	} else if constexpr (std::is_same<DTYPE_X2, int32_t>::value) {
61	using OpDag = TruncateDivOp::TruncateDivFloatToLowBit<float, int32_t, float>::OpDag;
62	BroadcastSch<schMode, OpDag> sch(tiling);	35	BroadcastSch<schMode, OpDag> sch(tiling);
63	sch.Process(x1, x2, y);	36	sch.Process(x1, y);
64	} else if constexpr (std::is_same<DTYPE_X2, half>::value) {
65	if constexpr (canUseMul) {
66	using OpDag = TruncateDivOp::TruncateDivFloatScalar<float>::OpDag;
67	BroadcastSch<schMode, OpDag> sch(tiling);
68	sch.Process(x1, y);
69	} else {
70	using OpDag = TruncateDivOp::TruncateDivFloatToLowBit<float, half, float>::OpDag;
71	BroadcastSch<schMode, OpDag> sch(tiling);
72	sch.Process(x1, x2, y);
73	}
74	}
75	} else if constexpr (std::is_same<DTYPE_X1, int8_t>::value \|\| std::is_same<DTYPE_X1, uint8_t>::value) {
76	if constexpr (std::is_same<DTYPE_X1, int8_t>::value) {
77	using OpDag = TruncateDivOp::TruncateDivIntS8<int8_t, half>::OpDag;
78	BroadcastSch<schMode, OpDag> sch(tiling);
79	sch.Process(x1, x2, y);
80	} else {	37	} else {
81	using OpDag = TruncateDivOp::TruncateDivIntU8<uint8_t, uint16_t>::OpDag;	38	using OpDag = TruncateDivOp::TruncateDivFloat16<half, float>::OpDag;
82	BroadcastSch<schMode, OpDag> sch(tiling);	39	BroadcastSch<schMode, OpDag> sch(tiling);
83	sch.Process(x1, x2, y);	40	sch.Process(x1, x2, y);
84	}	41	}
85	} else if constexpr (	42	}
86	std::is_same<DTYPE_X2, DTYPE_X1>::value &&	43	}
87	(std::is_same<DTYPE_X1, int16_t>::value \|\| std::is_same<DTYPE_X1, int32_t>::value)) {	44
88	using OpDag = TruncateDivOp::TruncateDivInt<DTYPE_X1>::OpDag;	45	template <uint64_t schMode, bool canUseMul>
		46	__aicore__ inline void truncate_div_bfloat16(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR tiling)
		47	{
		48	if constexpr (std::is_same<DTYPE_X2, bfloat16_t>::value) {
		49	if constexpr (canUseMul) {
		50	using OpDag = TruncateDivOp::TruncateDivFloat16Scalar<bfloat16_t, float>::OpDag;
		51	BroadcastSch<schMode, OpDag> sch(tiling);
		52	sch.Process(x1, y);
		53	} else {
		54	using OpDag = TruncateDivOp::TruncateDivFloat16<bfloat16_t, float>::OpDag;
		55	BroadcastSch<schMode, OpDag> sch(tiling);
		56	sch.Process(x1, x2, y);
		57	}
		58	}
		59	}
		60
		61	template <uint64_t schMode, bool canUseMul>
		62	__aicore__ inline void truncate_div_float(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR tiling)
		63	{
		64	if constexpr (std::is_same<DTYPE_X2, float>::value) {
		65	if constexpr (canUseMul) {
		66	using OpDag = TruncateDivOp::TruncateDivFloatScalar<float>::OpDag;
		67	BroadcastSch<schMode, OpDag> sch(tiling);
		68	sch.Process(x1, y);
		69	} else {
		70	using OpDag = TruncateDivOp::TruncateDivFloat<float>::OpDag;
		71	BroadcastSch<schMode, OpDag> sch(tiling);
		72	sch.Process(x1, x2, y);
		73	}
		74	} else if constexpr (std::is_same<DTYPE_X2, int32_t>::value) {
		75	using OpDag = TruncateDivOp::TruncateDivFloatToLowBit<float, int32_t, float>::OpDag;
89	BroadcastSch<schMode, OpDag> sch(tiling);	76	BroadcastSch<schMode, OpDag> sch(tiling);
90	sch.Process(x1, x2, y);	77	sch.Process(x1, x2, y);
91	} else if constexpr (std::is_same<DTYPE_X1, int64_t>::value) {	78	} else if constexpr (std::is_same<DTYPE_X2, half>::value) {
92	using OpDag = TruncateDivOp::TruncateDivInt64<DTYPE_X1>::OpDag;	79	if constexpr (canUseMul) {
		80	using OpDag = TruncateDivOp::TruncateDivFloatScalar<float>::OpDag;
		81	BroadcastSch<schMode, OpDag> sch(tiling);
		82	sch.Process(x1, y);
		83	} else {
		84	using OpDag = TruncateDivOp::TruncateDivFloatToLowBit<float, half, float>::OpDag;
		85	BroadcastSch<schMode, OpDag> sch(tiling);
		86	sch.Process(x1, x2, y);
		87	}
		88	}
		89	}
		90
		91	template <uint64_t schMode, bool canUseMul>
		92	__aicore__ inline void truncate_div_int8(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR tiling)
		93	{
		94	using OpDag = TruncateDivOp::TruncateDivIntS8<int8_t, half>::OpDag;
		95	BroadcastSch<schMode, OpDag> sch(tiling);
		96	sch.Process(x1, x2, y);
		97	}
		98
		99	template <uint64_t schMode, bool canUseMul>
		100	__aicore__ inline void truncate_div_uint8(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR tiling)
		101	{
		102	using OpDag = TruncateDivOp::TruncateDivIntU8<uint8_t, uint16_t>::OpDag;
		103	BroadcastSch<schMode, OpDag> sch(tiling);
		104	sch.Process(x1, x2, y);
		105	}
		106
		107	template <uint64_t schMode, bool canUseMul>
		108	__aicore__ inline void truncate_div_int16(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR tiling)
		109	{
		110	if constexpr (std::is_same<DTYPE_X2, int16_t>::value) {
		111	using OpDag = TruncateDivOp::TruncateDivInt<int16_t>::OpDag;
93	BroadcastSch<schMode, OpDag> sch(tiling);	112	BroadcastSch<schMode, OpDag> sch(tiling);
94	sch.Process(x1, x2, y);	113	sch.Process(x1, x2, y);
95	} else if constexpr (std::is_same<DTYPE_X1, int32_t>::value && std::is_same<DTYPE_X2, float>::value) {	114	}
		115	}
		116
		117	template <uint64_t schMode, bool canUseMul>
		118	__aicore__ inline void truncate_div_int32(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR tiling)
		119	{
		120	if constexpr (std::is_same<DTYPE_X2, int32_t>::value) {
		121	using OpDag = TruncateDivOp::TruncateDivInt<int32_t>::OpDag;
		122	BroadcastSch<schMode, OpDag> sch(tiling);
		123	sch.Process(x1, x2, y);
		124	} else if constexpr (std::is_same<DTYPE_X2, float>::value) {
96	using OpDag = TruncateDivOp::TruncateDivIntToFloat<int32_t, float, float>::OpDag;	125	using OpDag = TruncateDivOp::TruncateDivIntToFloat<int32_t, float, float>::OpDag;
97	BroadcastSch<schMode, OpDag> sch(tiling);	126	BroadcastSch<schMode, OpDag> sch(tiling);
98	sch.Process(x1, x2, y);	127	sch.Process(x1, x2, y);
99	}	128	}
		129	}
		130
		131	template <uint64_t schMode, bool canUseMul>
		132	__aicore__ inline void truncate_div_int64(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR tiling)
		133	{
		134	using OpDag = TruncateDivOp::TruncateDivInt64<int64_t>::OpDag;
		135	BroadcastSch<schMode, OpDag> sch(tiling);
		136	sch.Process(x1, x2, y);
		137	}
		138
		139	/**
		140	* Supported data type combinations (x1, x2, y):
		141	* \| x1 (DTYPE_X1) \| x2 (DTYPE_X2) \| y (DTYPE_Y) \|
		142	* \|------------------\|-----------------\|-----------------\|
		143	* \| half \| float \| float \|
		144	* \| half \| half \| half \|
		145	* \| bfloat16_t \| bfloat16_t \| bfloat16_t \|
		146	* \| float \| float \| float \|
		147	* \| float \| int32_t \| float \|
		148	* \| float \| half \| float \|
		149	* \| int8_t \| int8_t \| int8_t \|
		150	* \| uint8_t \| uint8_t \| uint8_t \|
		151	* \| int16_t \| int16_t \| int16_t \|
		152	* \| int32_t \| int32_t \| int32_t \|
		153	* \| int32_t \| float \| float \|
		154	* \| int64_t \| int64_t \| int64_t \|
		155	*/
		156	template <uint64_t schMode, bool canUseMul>
		157	__global__ __aicore__ void truncate_div(GM_ADDR x1, GM_ADDR x2, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
		158	{
		159	if constexpr (std::is_same<DTYPE_X1, half>::value) {
		160	truncate_div_half<schMode, canUseMul>(x1, x2, y, tiling);
		161	} else if constexpr (std::is_same<DTYPE_X1, bfloat16_t>::value) {
		162	truncate_div_bfloat16<schMode, canUseMul>(x1, x2, y, tiling);
		163	} else if constexpr (std::is_same<DTYPE_X1, float>::value) {
		164	truncate_div_float<schMode, canUseMul>(x1, x2, y, tiling);
		165	} else if constexpr (std::is_same<DTYPE_X1, int8_t>::value) {
		166	truncate_div_int8<schMode, canUseMul>(x1, x2, y, tiling);
		167	} else if constexpr (std::is_same<DTYPE_X1, uint8_t>::value) {
		168	truncate_div_uint8<schMode, canUseMul>(x1, x2, y, tiling);
		169	} else if constexpr (std::is_same<DTYPE_X1, int16_t>::value) {
		170	truncate_div_int16<schMode, canUseMul>(x1, x2, y, tiling);
		171	} else if constexpr (std::is_same<DTYPE_X1, int32_t>::value) {
		172	truncate_div_int32<schMode, canUseMul>(x1, x2, y, tiling);
		173	} else if constexpr (std::is_same<DTYPE_X1, int64_t>::value) {
		174	truncate_div_int64<schMode, canUseMul>(x1, x2, y, tiling);
		175	}
100	}	176	}

	@@ -1223,7 +1223,7 @@ ge::graphStatus PadACTiling::Fp4ValidatePaddings()
1223	// fp4 输入数据类型，左右pad的最后一维均为偶数	1223	// fp4 输入数据类型，左右pad的最后一维均为偶数
1224	size_t frontDimNum = paddings_.padFront.GetDimNum();	1224	size_t frontDimNum = paddings_.padFront.GetDimNum();
1225	int64_t frontValue = paddings_.padFront.GetDim(frontDimNum - 1);	1225	int64_t frontValue = paddings_.padFront.GetDim(frontDimNum - 1);
1226	if (frontValue % HALF_FACTOR != 0) {	1226	if (frontValue % static_cast<int64_t>(HALF_FACTOR) != 0) {
1227	std::string paramMsg = "the last axis of padFront";	1227	std::string paramMsg = "the last axis of padFront";
1228	std::string reasonMsg = "When the dtype is fp4, the last axis of padFront must be an even number.";	1228	std::string reasonMsg = "When the dtype is fp4, the last axis of padFront must be an even number.";
1229	OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(	1229	OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(
	@@ -1233,7 +1233,7 @@ ge::graphStatus PadACTiling::Fp4ValidatePaddings()
1233		1233
1234	size_t backDimNum = paddings_.padBack.GetDimNum();	1234	size_t backDimNum = paddings_.padBack.GetDimNum();
1235	int64_t backValue = paddings_.padBack.GetDim(backDimNum - 1);	1235	int64_t backValue = paddings_.padBack.GetDim(backDimNum - 1);
1236	if (backValue % HALF_FACTOR != 0) {	1236	if (backValue % static_cast<int64_t>(HALF_FACTOR) != 0) {
1237	std::string paramMsg = "the last axis of padBack";	1237	std::string paramMsg = "the last axis of padBack";
1238	std::string reasonMsg = "When the dtype is fp4, the last axis of padBack must be an even number.";	1238	std::string reasonMsg = "When the dtype is fp4, the last axis of padBack must be an even number.";
1239	OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(	1239	OP_LOGE_FOR_INVALID_SHAPE_WITH_REASON(

	@@ -42,6 +42,7 @@ static const uint64_t MAX_SIZE_USING_SINGLE_CORE = static_cast<uint64_t>(4 * 102
42	static const uint64_t CUTR_MIN_MOV = static_cast<uint64_t>(4 * 1024);	42	static const uint64_t CUTR_MIN_MOV = static_cast<uint64_t>(4 * 1024);
43	static const uint64_t CUTR_MIN_SIZE = 4;	43	static const uint64_t CUTR_MIN_SIZE = 4;
44	static const float AR_BLOCK_NUM_FACTOR = 0.85f;	44	static const float AR_BLOCK_NUM_FACTOR = 0.85f;
		45	static constexpr int BF16_BUFFER_MULTIPLIER = 2; // BF16数据类型因float32中间结果所需的缓冲区倍数
45		46
46	ge::graphStatus ArgCommonBaseTiling::Init(const uint64_t& coreNum, const uint64_t& ubSize, const uint64_t& vRegSize)	47	ge::graphStatus ArgCommonBaseTiling::Init(const uint64_t& coreNum, const uint64_t& ubSize, const uint64_t& vRegSize)
47	{	48	{
	@@ -390,7 +391,7 @@ ge::graphStatus ArgCommonBaseTiling::CalcSplitInfoForAr()
390	blkFactor_ = static_cast<uint64_t>(1);	391	blkFactor_ = static_cast<uint64_t>(1);
391	blkTailFactor_ = static_cast<uint64_t>(0);	392	blkTailFactor_ = static_cast<uint64_t>(0);
392	} else {	393	} else {
393	realCoreNum_ = coreNum_;	394	realCoreNum_ = (coreNum_ <= 0) ? 1 : coreNum_;
394	blkFactor_ = tilingData_.get_aSize() / realCoreNum_;	395	blkFactor_ = tilingData_.get_aSize() / realCoreNum_;
395	blkTailFactor_ = tilingData_.get_aSize() % realCoreNum_;	396	blkTailFactor_ = tilingData_.get_aSize() % realCoreNum_;
396	}	397	}
	@@ -718,7 +719,7 @@ ge::graphStatus ArgCommonBaseTiling::CalcSplitInfoForCopyOnly()
718	blkFactor_ = static_cast<uint64_t>(1);	719	blkFactor_ = static_cast<uint64_t>(1);
719	blkTailFactor_ = static_cast<uint64_t>(0);	720	blkTailFactor_ = static_cast<uint64_t>(0);
720	} else {	721	} else {
721	realCoreNum_ = coreNum_;	722	realCoreNum_ = (coreNum_ <= 0) ? 1 : coreNum_;
722	blkFactor_ = sumSize / realCoreNum_;	723	blkFactor_ = sumSize / realCoreNum_;
723	blkTailFactor_ = sumSize % realCoreNum_;	724	blkTailFactor_ = sumSize % realCoreNum_;
724	}	725	}
	@@ -774,7 +775,7 @@ void ArgCommonBaseTiling::AddExtraBufferNeed(uint64_t& fixedNeed, int& isBfloatN
774	{	775	{
775	if (ge::DT_BF16 == valueDtype_) {	776	if (ge::DT_BF16 == valueDtype_) {
776	fixedNeed += alignSize * DOUBLE_BUFFER_NUM * sizeof(float);	777	fixedNeed += alignSize * DOUBLE_BUFFER_NUM * sizeof(float);
777	isBfloatNum = 2;	778	isBfloatNum = BF16_BUFFER_MULTIPLIER;
778	}	779	}
779	if ((indiceDtypeSize_ != t2Size_) \|\| (indiceDtype_ != ge::DT_INT32 && indiceDtype_ != ge::DT_INT64)) {	780	if ((indiceDtypeSize_ != t2Size_) \|\| (indiceDtype_ != ge::DT_INT32 && indiceDtype_ != ge::DT_INT64)) {
780	fixedNeed += alignSize * indiceDtypeSize_;	781	fixedNeed += alignSize * indiceDtypeSize_;