Conv:kernel ut update-ops-nn-AtomGit

Mconv/common/op_kernel/arch35/conv_config.h+2-4

	@@ -287,10 +287,8 @@ public:
287	LocalTensor<typename ConvDataType::FmapT> wholeAl0Tensor;	287	LocalTensor<typename ConvDataType::FmapT> wholeAl0Tensor;
288	LocalTensor<typename ConvDataType::WeightT> wholeBl0Tensor;	288	LocalTensor<typename ConvDataType::WeightT> wholeBl0Tensor;
289	LocalTensor<typename ConvDataType::L0cT> biasBT;	289	LocalTensor<typename ConvDataType::L0cT> biasBT;
290	- LocalTensor<typename ConvDataType::L0cT> wholeCl0Tensor =	290	+ LocalTensor<typename ConvDataType::L0cT> wholeCl0Tensor;
		L L lileizheng6月10日这个好像不能删，52有个场景依赖这个初始化高维健 6月10日评论： 6月10日评论：
291	- LocalTensor<typename ConvDataType::L0cT>(TPosition::CO1, 0, 0);	291	+ LocalTensor<typename ConvDataType::L0cT> cl0;
292	- LocalTensor<typename ConvDataType::L0cT> cl0 =
293	- LocalTensor<typename ConvDataType::L0cT>(TPosition::CO1, 0, 0);
294		292
295	uint8_t enableBias = false; // 是否有bias	293	uint8_t enableBias = false; // 是否有bias
296	uint8_t enableVectorQuant = false; // 是否有vector类型scale，双输出场景下任一scale为vector类型，即为true	294	uint8_t enableVectorQuant = false; // 是否有vector类型scale，双输出场景下任一scale为vector类型，即为true

Mconv/common/op_kernel/arch35/conv_instr_impl.h+2-0

	@@ -154,7 +154,9 @@ public:
154	{	154	{
155	param_.SetKStartPosition(static_cast<uint32_t>(KStartPosition));	155	param_.SetKStartPosition(static_cast<uint32_t>(KStartPosition));
156	param_.SetKStep(static_cast<uint16_t>(kStep));	156	param_.SetKStep(static_cast<uint16_t>(kStep));
		157	+#ifndef ASCENDC_CPU_DEBUG
157	LoadData<TPosition::B2, TPosition::B1, typename Intf::WeightT>(bl0, self_->ctx.bl1, param_);	158	LoadData<TPosition::B2, TPosition::B1, typename Intf::WeightT>(bl0, self_->ctx.bl1, param_);
		L L lileizheng6月10日那ut测试的时候不少了个指令。高维健 6月10日评论： 6月10日评论：
		159	+#endif
158	}	160	}
159		161
160	__aicore__ inline void FullLoadBL0(const LocalTensor<typename Intf::WeightT> &bl0)	162	__aicore__ inline void FullLoadBL0(const LocalTensor<typename Intf::WeightT> &bl0)

Mconv/conv2d_v2/op_kernel/conv2d_v2.cpp+1-1

	@@ -113,7 +113,7 @@ __global__ __aicore__ void conv2dv2(GM_ADDR x, GM_ADDR filter, GM_ADDR bias, GM_
113	#endif	113	#endif
114	using scaleType = ConvType<TPosition::GM, scaleFormat, uint64_t>;	114	using scaleType = ConvType<TPosition::GM, scaleFormat, uint64_t>;
115		115
116	- if (SmallKernel == 1 && OutputOrder == static_cast<int8_t>(ConvOutputOrder::M_MODE) &&	116	+ if constexpr (SmallKernel == 1 && OutputOrder == static_cast<int8_t>(ConvOutputOrder::M_MODE) &&
117	fmapFormat == ConvFormat::NCHW && outputFormat == ConvFormat::NCHW &&	117	fmapFormat == ConvFormat::NCHW && outputFormat == ConvFormat::NCHW &&
118	AscendC::IsSameType<DTYPE_X, half>::value) {	118	AscendC::IsSameType<DTYPE_X, half>::value) {
119	Conv2dSmallKernel<DTYPE_X, DTYPE_FILTER, biasType::T, DTYPE_Y> op;	119	Conv2dSmallKernel<DTYPE_X, DTYPE_FILTER, biasType::T, DTYPE_Y> op;

Mconv/conv2d_v2/tests/ut/op_kernel/test_conv2d_v2.cpp+45-23

	@@ -10,20 +10,26 @@
10		10
11		11
12	#include <iostream>	12	#include <iostream>
		13	+#include <vector>
		14	+#include <numeric>
		15	+#include <functional>
13	#include "gtest/gtest.h"	16	#include "gtest/gtest.h"
14	#include "tikicpulib.h"	17	#include "tikicpulib.h"
		18	+#include "version/asc_devkit_version.h"
15	#include "conv2d_v2_tiling_def.h"	19	#include "conv2d_v2_tiling_def.h"
16		20
		21	+
		22	+#ifndef CONV_KERNEL
		23	+#include "conv2d_v2/conv2d_v2.cpp"
		24	+#define CONV_KERNEL
		25	+#endif
		26	+
17	constexpr uint8_t FP16_SIZE = 2;	27	constexpr uint8_t FP16_SIZE = 2;
18	-constexpr uint8_t N0 = 16;
19	constexpr uint32_t DIM2 = 2;	28	constexpr uint32_t DIM2 = 2;
20	constexpr uint32_t DIM3 = 3;	29	constexpr uint32_t DIM3 = 3;
21	constexpr uint32_t SIZE_1K = 1024;	30	constexpr uint32_t SIZE_1K = 1024;
22	constexpr uint32_t NUM_16 = 16;	31	constexpr uint32_t NUM_16 = 16;
23		32
24	-extern "C" __global__ __aicore__ void conv2dv2(GM_ADDR x, GM_ADDR filter, GM_ADDR bias, GM_ADDR offset_w,
25	- GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling);
26	-
27	class Conv2DV2KernelTest : public testing::Test {	33	class Conv2DV2KernelTest : public testing::Test {
28	protected:	34	protected:
29	static void SetUpTestCase() {	35	static void SetUpTestCase() {
	@@ -110,7 +116,8 @@ void SetConv2dApiPartOne(Conv2DTilingData* tiling, const TilingInput& tilingInpu
110	tiling->orgHo = outputShape[DIM2];	116	tiling->orgHo = outputShape[DIM2];
111	tiling->orgWo = outputShape[DIM3];	117	tiling->orgWo = outputShape[DIM3];
112	tiling->singleCoreBatch = inputShape[0];	118	tiling->singleCoreBatch = inputShape[0];
113	- tiling->singleCoreHo = outputShape[DIM2];	119	+ // M_MODE: singleCoreHo stores flattened M (Ho * Wo) per core
		120	+ tiling->singleCoreHo = outputShape[DIM2] * outputShape[DIM3];
114	tiling->singleCoreWo = outputShape[DIM3];	121	tiling->singleCoreWo = outputShape[DIM3];
115	tiling->orgCi = inputShape[1];	122	tiling->orgCi = inputShape[1];
116	tiling->orgCo = weightShape[0];	123	tiling->orgCo = weightShape[0];
	@@ -118,8 +125,12 @@ void SetConv2dApiPartOne(Conv2DTilingData* tiling, const TilingInput& tilingInpu
118	tiling->singleCoreCo = weightShape[0];	125	tiling->singleCoreCo = weightShape[0];
119	tiling->hoL1 = NUM_16;	126	tiling->hoL1 = NUM_16;
120	tiling->woL1 = 0;	127	tiling->woL1 = 0;
121	- tiling->kAL1 = NUM_16;	128	+ tiling->orgHixWi = inputShape[DIM2] * inputShape[DIM3];
122	- tiling->kBL1 = NUM_16;	129	+ tiling->kernelHxkernelW = weightShape[DIM2] * weightShape[DIM3];
		130	+ tiling->kernelHxkernelWxkernelD = weightShape[DIM2] * weightShape[DIM3];
		131	+ uint64_t kL1 = tiling->singleCoreCi * tiling->kernelHxkernelW;
		132	+ tiling->kAL1 = kL1;
		133	+ tiling->kBL1 = kL1;
123	tiling->nBL1 = NUM_16;	134	tiling->nBL1 = NUM_16;
124	tiling->hoL0 = NUM_16;	135	tiling->hoL0 = NUM_16;
125	tiling->woL0 = 0;	136	tiling->woL0 = 0;
	@@ -132,15 +143,12 @@ void SetConv2dApiPartOne(Conv2DTilingData* tiling, const TilingInput& tilingInpu
132	tiling->singleCoreGroupOpt = 0;	143	tiling->singleCoreGroupOpt = 0;
133	tiling->bUbNStep = 0;	144	tiling->bUbNStep = 0;
134	tiling->bUbKStep = 0;	145	tiling->bUbKStep = 0;
135	- tiling->orgHixWi = inputShape[DIM2] * inputShape[DIM3];
136	- tiling->kernelHxkernelW = weightShape[DIM2] * weightShape[DIM3];
137	- tiling->kernelHxkernelWxkernelD = weightShape[DIM2] * weightShape[DIM3];
138	- tiling->aL1SpaceSize = SIZE_1K;
139	tiling->multiNBL1 = 1;	146	tiling->multiNBL1 = 1;
140	tiling->cinAInCore = tiling->kAL1 / tiling->kernelHxkernelW;	147	tiling->cinAInCore = tiling->kAL1 / tiling->kernelHxkernelW;
141	tiling->cinATailInCore = tiling->cinAInCore;	148	tiling->cinATailInCore = tiling->cinAInCore;
142	tiling->cinBInCore = tiling->kBL1 / tiling->kernelHxkernelW;	149	tiling->cinBInCore = tiling->kBL1 / tiling->kernelHxkernelW;
143	tiling->cinBTailInCore = tiling->cinBInCore;	150	tiling->cinBTailInCore = tiling->cinBInCore;
		151	+ tiling->aL1SpaceSize = tiling->cinAInCore * tiling->orgHixWi * FP16_SIZE;
144	}	152	}
145		153
146	void SetConv2dApiPartTwo(Conv2DTilingData* tiling, const TilingInput& tilingInput)	154	void SetConv2dApiPartTwo(Conv2DTilingData* tiling, const TilingInput& tilingInput)
	@@ -187,18 +195,32 @@ void SetConv2dApiPartTwo(Conv2DTilingData* tiling, const TilingInput& tilingInpu
187	tiling->roundMode = 0;	195	tiling->roundMode = 0;
188	}	196	}
189		197
		198	+void SetUnionDataXt(Conv2DTilingData* tiling)
		199	+{
		200	+ uint64_t xt = 0;
		201	+ xt \|= static_cast<uint64_t>(tiling->strideW & 0x3f);
		202	+ xt \|= (static_cast<uint64_t>(tiling->strideH & 0x3f) << 6);
		203	+ xt \|= (static_cast<uint64_t>(tiling->kw & 0xff) << 12);
		204	+ xt \|= (static_cast<uint64_t>(tiling->kh & 0xff) << 20);
		205	+ xt \|= (static_cast<uint64_t>(tiling->dilationW & 0xff) << 28);
		206	+ xt \|= (static_cast<uint64_t>(tiling->dilationH & 0xff) << 36);
		207	+ xt \|= ((static_cast<uint64_t>(tiling->kw) & 0x100) >> 8) << 44;
		208	+ xt \|= ((static_cast<uint64_t>(tiling->kh) & 0x100) >> 8) << 45;
		209	+ tiling->unionDataXt = xt;
		210	+}
		211	+
190	void SetTilingData(Conv2DTilingData* tiling, const TilingInput& tilingInput)	212	void SetTilingData(Conv2DTilingData* tiling, const TilingInput& tilingInput)
191	{	213	{
192	SetConv2dRunInfo(tiling, tilingInput);	214	SetConv2dRunInfo(tiling, tilingInput);
193	SetConv2dApiPartOne(tiling, tilingInput);	215	SetConv2dApiPartOne(tiling, tilingInput);
194	SetConv2dApiPartTwo(tiling, tilingInput);	216	SetConv2dApiPartTwo(tiling, tilingInput);
		217	+ SetUnionDataXt(tiling);
195	}	218	}
196		219
197	void TestSimpleKernel(const std::vector<uint64_t>& inputShape, const std::vector<uint64_t>& weightShape)	220	void TestSimpleKernel(const std::vector<uint64_t>& inputShape, const std::vector<uint64_t>& weightShape)
198	{	221	{
199	AscendC::SetKernelMode(KernelMode::MIX_MODE);	222	AscendC::SetKernelMode(KernelMode::MIX_MODE);
200	const uint64_t numBlocks = 1;	223	const uint64_t numBlocks = 1;
201	- const uint64_t groups = 1;
202	std::vector<uint64_t> pads = {0, 0, 0, 0};	224	std::vector<uint64_t> pads = {0, 0, 0, 0};
203	std::vector<uint64_t> strides = {1, 1};	225	std::vector<uint64_t> strides = {1, 1};
204	std::vector<uint64_t> dilations = {1, 1};	226	std::vector<uint64_t> dilations = {1, 1};
	@@ -214,12 +236,10 @@ void TestSimpleKernel(const std::vector<uint64_t>& inputShape, const std::vector
214	size_t weightBytes = VectorReduceMul(weightShape) * FP16_SIZE;	236	size_t weightBytes = VectorReduceMul(weightShape) * FP16_SIZE;
215	size_t outputBytes = VectorReduceMul(outputShape) * FP16_SIZE;	237	size_t outputBytes = VectorReduceMul(outputShape) * FP16_SIZE;
216	size_t workspaceSize = SIZE_1K * SIZE_1K * NUM_16;	238	size_t workspaceSize = SIZE_1K * SIZE_1K * NUM_16;
217	- size_t tilingDataSize = sizeof(Conv2DTilingData);
218	-
219	uint8_t* input = (uint8_t*)AscendC::GmAlloc(inputBtyes);	239	uint8_t* input = (uint8_t*)AscendC::GmAlloc(inputBtyes);
220	uint8_t* weight = (uint8_t*)AscendC::GmAlloc(weightBytes);	240	uint8_t* weight = (uint8_t*)AscendC::GmAlloc(weightBytes);
221	uint8_t* output = (uint8_t*)AscendC::GmAlloc(outputBytes);	241	uint8_t* output = (uint8_t*)AscendC::GmAlloc(outputBytes);
222	- uint8_t* workspace = (uint8_t)AscendC::GmAlloc(SIZE_1K SIZE_1K * 16);	242	+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceSize);
223	uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(Conv2DTilingData));	243	uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(Conv2DTilingData));
224		244
225	memset(workspace, 0, workspaceSize);	245	memset(workspace, 0, workspaceSize);
	@@ -228,11 +248,13 @@ void TestSimpleKernel(const std::vector<uint64_t>& inputShape, const std::vector
228	TilingInput tilingInput = {inputShape, weightShape, outputShape, pads, strides, dilations};	248	TilingInput tilingInput = {inputShape, weightShape, outputShape, pads, strides, dilations};
229	SetTilingData(tilingData, tilingInput);	249	SetTilingData(tilingData, tilingInput);
230		250
231	- // auto conv2dv2_func = [](GM_ADDR x, GM_ADDR filter, GM_ADDR bias, GM_ADDR offset_w,	251	+ auto conv2dv2_func = [](GM_ADDR x, GM_ADDR filter, GM_ADDR bias, GM_ADDR offset_w,
232	- // GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {	252	+ GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
233	- // conv2dv2<0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0>(x, filter, bias, offset_w, y, workspace, tiling);	253	+ // OutputOrder=1 (M_MODE): tiling uses hoL1>0 && woL1==0
234	- // };	254	+ ::conv2dv2<0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0>(
235	- // ICPU_RUN_KF(conv2dv2_func, numBlocks, input, weight, nullptr, nullptr, output, workspace, tiling);	255	+ x, filter, bias, offset_w, y, workspace, tiling);
		256	+ };
		257	+ ICPU_RUN_KF(conv2dv2_func, numBlocks, input, weight, nullptr, nullptr, output, workspace, tiling);
236		258
237	AscendC::GmFree(input);	259	AscendC::GmFree(input);
238	AscendC::GmFree(weight);	260	AscendC::GmFree(weight);
	@@ -243,8 +265,8 @@ void TestSimpleKernel(const std::vector<uint64_t>& inputShape, const std::vector
243	} // namespace	265	} // namespace
244	TEST_F(Conv2DV2KernelTest, conv2dv2_kernel_test_base)	266	TEST_F(Conv2DV2KernelTest, conv2dv2_kernel_test_base)
245	{	267	{
246	- std::vector<uint64_t> inputShape = {1, 1, 1, 1};	268	+ std::vector<uint64_t> inputShape = {1, 16, 16, 16};
247	- std::vector<uint64_t> weightShape = {1, 1, 1, 1};	269	+ std::vector<uint64_t> weightShape = {16, 16, 1, 1};
248		270
249	- //TestSimpleKernel(inputShape, weightShape);	271	+ TestSimpleKernel(inputShape, weightShape);
250	}	272	}

Aconv/extend_conv2d/tests/ut/op_kernel/CMakeLists.txt+13-0

@@ -0,0 +1,13 @@
	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	+if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
	12	+ AddOpTestCase(extend_conv2d "ascend950pr_9599" "-DDTYPE_X=half -DDTYPE_FILTER=half -DDTYPE_Y0=half" conv2d_v2)
	13	+endif()

Aconv/extend_conv2d/tests/ut/op_kernel/extend_conv2d_tiling_def.h+244-0

@@ -0,0 +1,244 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+
	12	+/*!
	13	+ * \file extend_conv2d_tiling_def.h
	14	+ * \brief
	15	+ */
	16	+
	17	+#ifndef _GE_EXTEND_CONV2D_TILING_H_
	18	+#define _GE_EXTEND_CONV2D_TILING_H_
	19	+
	20	+#include "securec.h"
	21	+#include "kernel_operator.h"
	22	+#include "kernel_tiling/kernel_tiling.h"
	23	+
	24	+#define __forceinline__
	25	+#define FORMAT_X FORMAT_NCHW
	26	+#define FORMAT_FILTER FORMAT_NCHW
	27	+#define FORMAT_Y0 FORMAT_NCHW
	28	+
	29	+#pragma pack(1)
	30	+
	31	+struct Conv2DRunInfo {
	32	+ uint64_t hin = 0;
	33	+ uint64_t win = 0;
	34	+ uint64_t hout = 0;
	35	+ uint64_t wout = 0;
	36	+ uint32_t batch = 0;
	37	+ uint32_t cin = 0;
	38	+ uint32_t cout = 0;
	39	+ uint32_t kh = 0;
	40	+ uint32_t kw = 0;
	41	+ uint32_t batchDim = 0;
	42	+ uint32_t groupDim = 0;
	43	+ uint32_t nDim = 0;
	44	+ uint32_t hoDim = 0;
	45	+ uint32_t woDim = 0;
	46	+ uint32_t strideH = 0;
	47	+ uint32_t strideW = 0;
	48	+ uint32_t dilationH = 0;
	49	+ uint32_t dilationW = 0;
	50	+ uint32_t padTop = 0;
	51	+ uint32_t padLeft = 0;
	52	+ uint32_t groups = 0;
	53	+ uint32_t enlarge = 0;
	54	+ uint32_t cinOpt = 0;
	55	+ uint32_t coutOpt = 0;
	56	+ uint32_t groupOpt = 0;
	57	+ uint8_t hasBias = 0;
	58	+};
	59	+
	60	+struct TConv2DTiling {
	61	+ uint64_t orgHi = 0;
	62	+ uint64_t orgWi = 0;
	63	+ uint64_t orgHo = 0;
	64	+ uint64_t orgWo = 0;
	65	+ uint64_t singleCoreBatch = 0;
	66	+ uint64_t singleCoreHo = 0;
	67	+ uint64_t singleCoreWo = 0;
	68	+ uint32_t orgCi = 0;
	69	+ uint32_t orgCo = 0;
	70	+ uint32_t singleCoreCi = 0;
	71	+ uint32_t singleCoreCo = 0;
	72	+ uint32_t hoL1 = 0;
	73	+ uint32_t woL1 = 0;
	74	+ uint32_t kAL1 = 0;
	75	+ uint32_t kBL1 = 0;
	76	+ uint32_t nBL1 = 0;
	77	+ uint32_t hoL0 = 0;
	78	+ uint32_t woL0 = 0;
	79	+ uint32_t kL0 = 0;
	80	+ uint32_t nL0 = 0;
	81	+ uint32_t pBufferFlag = 0;
	82	+ uint32_t groups = 0;
	83	+ uint32_t enlarge = 0;
	84	+ uint32_t singleCoreGroups = 0;
	85	+ uint32_t singleCoreGroupOpt = 0;
	86	+ uint32_t bUbNStep = 0;
	87	+ uint32_t bUbKStep = 0;
	88	+ uint32_t orgHixWi = 0;
	89	+ uint32_t kernelHxkernelW = 0;
	90	+ uint32_t kernelHxkernelWxkernelD = 0;
	91	+ uint32_t aL1SpaceSize = 0;
	92	+ uint32_t multiNBL1 = 0;
	93	+ uint32_t cinAInCore = 0;
	94	+ uint32_t cinATailInCore = 0;
	95	+ uint32_t cinBInCore = 0;
	96	+ uint32_t cinBTailInCore = 0;
	97	+ uint32_t mStep = 0;
	98	+ uint32_t kStep = 0;
	99	+ uint32_t nStep = 0;
	100	+ uint32_t fmapKStride = 0;
	101	+ uint32_t weightKStride = 0;
	102	+ uint32_t cinOffsetBlockInGM = 0;
	103	+ uint32_t coutOffsetBlock = 0;
	104	+ uint32_t nL1DivBlockSize = 0;
	105	+ uint32_t kernelH = 0;
	106	+ uint32_t kernelW = 0;
	107	+ uint32_t strideH = 0;
	108	+ uint32_t strideW = 0;
	109	+ uint32_t dilationH = 0;
	110	+ uint32_t dilationW = 0;
	111	+ uint32_t padTop = 0;
	112	+ uint32_t padBottom = 0;
	113	+ uint32_t padLeft = 0;
	114	+ uint32_t padRight = 0;
	115	+ uint32_t innerBatch = 1;
	116	+ uint8_t iterateMNOrder = 0;
	117	+ uint8_t biasFullLoadFlag = 0;
	118	+ uint8_t fixpParamsFullLoadFlag = 0;
	119	+ uint8_t hf32Enable = 0;
	120	+ uint8_t hf32TransMode = 0;
	121	+ uint8_t hasBias = 0;
	122	+ uint8_t hasScale = 0;
	123	+ uint8_t dualOutput = 0;
	124	+ uint8_t quantMode0 = 0;
	125	+ uint8_t reluMode0 = 0;
	126	+ uint8_t clipMode0 = 0;
	127	+ uint8_t quantMode1 = 0;
	128	+ uint8_t reluMode1 = 0;
	129	+ uint8_t clipMode1 = 0;
	130	+ int8_t offsetx = 0;
	131	+ int8_t roundMode = 0;
	132	+};
	133	+
	134	+struct Conv2DTilingData
	135	+{
	136	+ uint64_t orgHi = 0;
	137	+ uint64_t orgWi = 0;
	138	+ uint64_t orgHo = 0;
	139	+ uint64_t orgWo = 0;
	140	+ uint64_t orgHixWi = 0;
	141	+ uint64_t singleCoreBatch = 0;
	142	+ uint64_t singleCoreHo = 0;
	143	+ uint64_t singleCoreWo = 0;
	144	+ uint64_t unionDataXt = 0;
	145	+ uint64_t hin = 0;
	146	+ uint64_t win = 0;
	147	+ uint64_t hout = 0;
	148	+ uint64_t wout = 0;
	149	+
	150	+ uint32_t orgCi = 0;
	151	+ uint32_t orgCo = 0;
	152	+ uint32_t singleCoreCi = 0;
	153	+ uint32_t singleCoreCo = 0;
	154	+ uint32_t hoL1 = 0;
	155	+ uint32_t woL1 = 0;
	156	+ uint32_t kAL1 = 0;
	157	+ uint32_t kBL1 = 0;
	158	+ uint32_t khL1 = 0;
	159	+ uint32_t kwL1 = 0;
	160	+ uint32_t nBL1 = 0;
	161	+ uint32_t hoL0 = 0;
	162	+ uint32_t woL0 = 0;
	163	+ uint32_t kL0 = 0;
	164	+ uint32_t nL0 = 0;
	165	+ uint32_t pBufferFlag = 0;
	166	+ uint32_t groups = 0;
	167	+ uint32_t enlarge = 0;
	168	+ uint32_t singleCoreGroups = 0;
	169	+ uint32_t singleCoreGroupOpt = 0;
	170	+ uint32_t bUbNStep = 0;
	171	+ uint32_t bUbKStep = 0;
	172	+ uint32_t khUb = 0;
	173	+ uint32_t kwUb = 0;
	174	+ uint32_t kernelHxkernelW = 0;
	175	+ uint32_t kernelHxkernelWxkernelD = 0;
	176	+ uint32_t aL1SpaceSize = 0;
	177	+ uint32_t multiNBL1 = 0;
	178	+ uint32_t cinAInCore = 0;
	179	+ uint32_t cinATailInCore = 0;
	180	+ uint32_t cinBInCore = 0;
	181	+ uint32_t cinBTailInCore = 0;
	182	+ uint32_t mStep = 0;
	183	+ uint32_t kStep = 0;
	184	+ uint32_t nStep = 0;
	185	+ uint32_t fmapKStride = 0;
	186	+ uint32_t weightKStride = 0;
	187	+ uint32_t cinOffsetBlockInGM = 0;
	188	+ uint32_t coutOffsetBlock = 0;
	189	+ uint32_t nL1DivBlockSize = 0;
	190	+ uint32_t kernelH = 0;
	191	+ uint32_t kernelW = 0;
	192	+ uint32_t strideH = 0;
	193	+ uint32_t strideW = 0;
	194	+ uint32_t dilationH = 0;
	195	+ uint32_t dilationW = 0;
	196	+ uint32_t padTop = 0;
	197	+ uint32_t padBottom = 0;
	198	+ uint32_t padLeft = 0;
	199	+ uint32_t padRight = 0;
	200	+ uint32_t innerBatch = 1;
	201	+ uint32_t batch = 0;
	202	+ uint32_t cin = 0;
	203	+ uint32_t cout = 0;
	204	+ uint32_t kh = 0;
	205	+ uint32_t kw = 0;
	206	+ uint32_t batchDim = 0;
	207	+ uint32_t groupDim = 0;
	208	+ uint32_t nDim = 0;
	209	+ uint32_t hoDim = 0;
	210	+ uint32_t woDim = 0;
	211	+ uint32_t cinOpt = 0;
	212	+ uint32_t coutOpt = 0;
	213	+ uint32_t groupOpt = 0;
	214	+
	215	+ uint8_t iterateMNOrder = 0;
	216	+ uint8_t biasFullLoadFlag = 0;
	217	+ uint8_t fixpParamsFullLoadFlag = 0;
	218	+ uint8_t hf32Enable = 0;
	219	+ uint8_t hf32TransMode = 0;
	220	+ uint8_t hasBias = 0;
	221	+ uint8_t hasScale = 0;
	222	+ uint8_t dualOutput = 0;
	223	+ uint8_t quantMode0 = 0;
	224	+ uint8_t reluMode0 = 0;
	225	+ uint8_t clipMode0 = 0;
	226	+ uint8_t quantMode1 = 0;
	227	+ uint8_t reluMode1 = 0;
	228	+ uint8_t clipMode1 = 0;
	229	+ int8_t offsetx = 0;
	230	+ int8_t roundMode = 0;
	231	+};
	232	+
	233	+#pragma pack()
	234	+
	235	+inline void InitTilingData(uint8_t* tiling, Conv2DTilingData* constData) {
	236	+ if (constData != nullptr && tiling != nullptr) {
	237	+ memcpy_s(constData, sizeof(Conv2DTilingData), tiling, sizeof(Conv2DTilingData));
	238	+ }
	239	+}
	240	+
	241	+#define GET_TILING_DATA(tilingData, tilingArg) \
	242	+ Conv2DTilingData tilingData; \
	243	+ InitTilingData(tilingArg, &tilingData)
	244	+#endif

Aconv/extend_conv2d/tests/ut/op_kernel/test_extend_conv2d.cpp+277-0

@@ -0,0 +1,277 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+
	12	+#include <iostream>
	13	+#include <vector>
	14	+#include <numeric>
	15	+#include <functional>
	16	+#include "gtest/gtest.h"
	17	+#include "tikicpulib.h"
	18	+#include "version/asc_devkit_version.h"
	19	+#include "extend_conv2d_tiling_def.h"
	20	+
	21	+
	22	+#ifndef CONV_KERNEL
	23	+#include "extend_conv2d/extend_conv2d.cpp"
	24	+#define CONV_KERNEL
	25	+#endif
	26	+
	27	+constexpr uint8_t FP16_SIZE = 2;
	28	+constexpr uint32_t DIM2 = 2;
	29	+constexpr uint32_t DIM3 = 3;
	30	+constexpr uint32_t SIZE_1K = 1024;
	31	+constexpr uint32_t NUM_16 = 16;
	32	+
	33	+class ExtendConv2DKernelTest : public testing::Test {
	34	+protected:
	35	+ static void SetUpTestCase() {
	36	+ std::cout << "ExtendConv2DKernelTest SetUp." << std::endl;
	37	+ }
	38	+
	39	+ static void TearDownTestCase() {
	40	+ std::cout << "ExtendConv2DKernelTest TearDown." << std::endl;
	41	+ }
	42	+};
	43	+namespace {
	44	+struct ConvShape {
	45	+ uint64_t shape;
	46	+ uint64_t pad1;
	47	+ uint64_t pad2;
	48	+ uint64_t dilation;
	49	+ uint64_t stride;
	50	+ uint64_t kernelSize;
	51	+};
	52	+
	53	+struct TilingInput {
	54	+ std::vector<uint64_t> inputShape;
	55	+ std::vector<uint64_t> weightShape;
	56	+ std::vector<uint64_t> outputShape;
	57	+ std::vector<uint64_t> pads;
	58	+ std::vector<uint64_t> strides;
	59	+ std::vector<uint64_t> dilations;
	60	+};
	61	+
	62	+uint64_t CalcOutputShape(const ConvShape& convShape)
	63	+{
	64	+ return ((convShape.shape + convShape.pad1 + convShape.pad2 - convShape.dilation * (convShape.kernelSize - 1) - 1) /
	65	+ convShape.stride + 1);
	66	+}
	67	+
	68	+size_t VectorReduceMul(const std::vector<uint64_t>& vec)
	69	+{
	70	+ return std::accumulate(std::begin(vec), std::end(vec), 1, std::multiplies<size_t>());
	71	+}
	72	+
	73	+void SetConv2dRunInfo(Conv2DTilingData* tiling, const TilingInput& tilingInput)
	74	+{
	75	+ std::vector<uint64_t> inputShape = tilingInput.inputShape;
	76	+ std::vector<uint64_t> weightShape = tilingInput.weightShape;
	77	+ std::vector<uint64_t> outputShape = tilingInput.outputShape;
	78	+ std::vector<uint64_t> pads = tilingInput.pads;
	79	+ std::vector<uint64_t> strides = tilingInput.strides;
	80	+ std::vector<uint64_t> dilations = tilingInput.dilations;
	81	+ tiling->hin = inputShape[DIM2];
	82	+ tiling->win = inputShape[DIM3];
	83	+ tiling->hout = outputShape[DIM2];
	84	+ tiling->wout = outputShape[DIM3];
	85	+ tiling->batch = inputShape[0];
	86	+ tiling->cin = inputShape[1];
	87	+ tiling->cout = weightShape[0];
	88	+ tiling->kh = weightShape[DIM2];
	89	+ tiling->kw = weightShape[DIM3];
	90	+ tiling->batchDim = 1;
	91	+ tiling->groupDim = 1;
	92	+ tiling->nDim = 1;
	93	+ tiling->hoDim = 1;
	94	+ tiling->woDim = 1;
	95	+ tiling->strideH = strides[0];
	96	+ tiling->strideW = strides[1];
	97	+ tiling->dilationH = dilations[0];
	98	+ tiling->dilationW = dilations[1];
	99	+ tiling->padTop = pads[0];
	100	+ tiling->padLeft = pads[DIM2];
	101	+ tiling->groups = 1;
	102	+ tiling->enlarge = 0;
	103	+ tiling->cinOpt = 0;
	104	+ tiling->coutOpt = 0;
	105	+ tiling->groupOpt = 0;
	106	+ tiling->hasBias = 0;
	107	+}
	108	+
	109	+void SetConv2dApiPartOne(Conv2DTilingData* tiling, const TilingInput& tilingInput)
	110	+{
	111	+ std::vector<uint64_t> inputShape = tilingInput.inputShape;
	112	+ std::vector<uint64_t> weightShape = tilingInput.weightShape;
	113	+ std::vector<uint64_t> outputShape = tilingInput.outputShape;
	114	+ tiling->orgHi = inputShape[DIM2];
	115	+ tiling->orgWi = inputShape[DIM3];
	116	+ tiling->orgHo = outputShape[DIM2];
	117	+ tiling->orgWo = outputShape[DIM3];
	118	+ tiling->singleCoreBatch = inputShape[0];
	119	+ tiling->singleCoreHo = outputShape[DIM2] * outputShape[DIM3];
	120	+ tiling->singleCoreWo = outputShape[DIM3];
	121	+ tiling->orgCi = inputShape[1];
	122	+ tiling->orgCo = weightShape[0];
	123	+ tiling->singleCoreCi = inputShape[1];
	124	+ tiling->singleCoreCo = weightShape[0];
	125	+ tiling->hoL1 = NUM_16;
	126	+ tiling->woL1 = 0;
	127	+ tiling->orgHixWi = inputShape[DIM2] * inputShape[DIM3];
	128	+ tiling->kernelHxkernelW = weightShape[DIM2] * weightShape[DIM3];
	129	+ tiling->kernelHxkernelWxkernelD = tiling->kernelHxkernelW;
	130	+ uint64_t kL1 = tiling->singleCoreCi * tiling->kernelHxkernelW;
	131	+ tiling->kAL1 = kL1;
	132	+ tiling->kBL1 = kL1;
	133	+ tiling->nBL1 = NUM_16;
	134	+ tiling->hoL0 = NUM_16;
	135	+ tiling->woL0 = 0;
	136	+ tiling->kL0 = NUM_16;
	137	+ tiling->nL0 = NUM_16;
	138	+ tiling->pBufferFlag = 0;
	139	+ tiling->groups = 1;
	140	+ tiling->enlarge = 0;
	141	+ tiling->singleCoreGroups = 0;
	142	+ tiling->singleCoreGroupOpt = 0;
	143	+ tiling->bUbNStep = 0;
	144	+ tiling->bUbKStep = 0;
	145	+ tiling->multiNBL1 = 1;
	146	+ tiling->cinAInCore = tiling->kAL1 / tiling->kernelHxkernelW;
	147	+ tiling->aL1SpaceSize = tiling->cinAInCore * tiling->orgHixWi * FP16_SIZE;
	148	+ tiling->cinATailInCore = tiling->cinAInCore;
	149	+ tiling->cinBInCore = tiling->kBL1 / tiling->kernelHxkernelW;
	150	+ tiling->cinBTailInCore = tiling->cinBInCore;
	151	+}
	152	+
	153	+void SetConv2dApiPartTwo(Conv2DTilingData* tiling, const TilingInput& tilingInput)
	154	+{
	155	+ std::vector<uint64_t> weightShape = tilingInput.weightShape;
	156	+ std::vector<uint64_t> pads = tilingInput.pads;
	157	+ std::vector<uint64_t> strides = tilingInput.strides;
	158	+ std::vector<uint64_t> dilations = tilingInput.dilations;
	159	+ tiling->mStep = NUM_16;
	160	+ tiling->kStep = 1;
	161	+ tiling->nStep = 1;
	162	+ tiling->fmapKStride = 1;
	163	+ tiling->weightKStride = 1;
	164	+ tiling->cinOffsetBlockInGM = tiling->kAL1 /
	165	+ tiling->kernelHxkernelW * tiling->orgHixWi;
	166	+ tiling->coutOffsetBlock = (tiling->orgCi /
	167	+ tiling->groups) * tiling->kernelHxkernelW;
	168	+ tiling->nL1DivBlockSize = tiling->nBL1 / NUM_16;
	169	+ tiling->kernelH = weightShape[DIM2];
	170	+ tiling->kernelW = weightShape[DIM3];
	171	+ tiling->strideH = strides[0];
	172	+ tiling->strideW = strides[1];
	173	+ tiling->dilationH = dilations[0];
	174	+ tiling->dilationW = dilations[1];
	175	+ tiling->padTop = pads[0];
	176	+ tiling->padBottom = pads[1];
	177	+ tiling->padLeft = pads[DIM2];
	178	+ tiling->padRight = pads[DIM3];
	179	+ tiling->iterateMNOrder = 0;
	180	+ tiling->biasFullLoadFlag = 1;
	181	+ tiling->fixpParamsFullLoadFlag = 1;
	182	+ tiling->hf32Enable = 0;
	183	+ tiling->hf32TransMode = 0;
	184	+ tiling->hasBias = 0;
	185	+ tiling->hasScale = 0;
	186	+ tiling->dualOutput = 0;
	187	+ tiling->quantMode0 = 0;
	188	+ tiling->reluMode0 = 0;
	189	+ tiling->clipMode0 = 0;
	190	+ tiling->quantMode1 = 0;
	191	+ tiling->reluMode1 = 0;
	192	+ tiling->clipMode1 = 0;
	193	+ tiling->offsetx = 0;
	194	+ tiling->roundMode = 0;
	195	+}
	196	+
	197	+void SetUnionDataXt(Conv2DTilingData* tiling)
	198	+{
	199	+ uint64_t xt = 0;
	200	+ xt \|= static_cast<uint64_t>(tiling->strideW & 0x3f);
	201	+ xt \|= (static_cast<uint64_t>(tiling->strideH & 0x3f) << 6);
	202	+ xt \|= (static_cast<uint64_t>(tiling->kw & 0xff) << 12);
	203	+ xt \|= (static_cast<uint64_t>(tiling->kh & 0xff) << 20);
	204	+ xt \|= (static_cast<uint64_t>(tiling->dilationW & 0xff) << 28);
	205	+ xt \|= (static_cast<uint64_t>(tiling->dilationH & 0xff) << 36);
	206	+ xt \|= ((static_cast<uint64_t>(tiling->kw) & 0x100) >> 8) << 44;
	207	+ xt \|= ((static_cast<uint64_t>(tiling->kh) & 0x100) >> 8) << 45;
	208	+ tiling->unionDataXt = xt;
	209	+}
	210	+
	211	+void SetTilingData(Conv2DTilingData* tiling, const TilingInput& tilingInput)
	212	+{
	213	+ SetConv2dRunInfo(tiling, tilingInput);
	214	+ SetConv2dApiPartOne(tiling, tilingInput);
	215	+ SetConv2dApiPartTwo(tiling, tilingInput);
	216	+ SetUnionDataXt(tiling);
	217	+}
	218	+
	219	+void TestSimpleKernel(const std::vector<uint64_t>& inputShape, const std::vector<uint64_t>& weightShape)
	220	+{
	221	+ AscendC::SetKernelMode(KernelMode::MIX_MODE);
	222	+ const uint64_t numBlocks = 1;
	223	+ std::vector<uint64_t> pads = {0, 0, 0, 0};
	224	+ std::vector<uint64_t> strides = {1, 1};
	225	+ std::vector<uint64_t> dilations = {1, 1};
	226	+ ConvShape convShapeH =
	227	+ {inputShape[DIM2], pads[0], pads[1], dilations[0], strides[0], weightShape[DIM2]};
	228	+ ConvShape convShapeW =
	229	+ {inputShape[DIM3], pads[DIM2], pads[DIM3], dilations[1], strides[1], weightShape[DIM3]};
	230	+ uint64_t ho = CalcOutputShape(convShapeH);
	231	+ uint64_t wo = CalcOutputShape(convShapeW);
	232	+ std::vector<uint64_t> outputShape = {inputShape[0], weightShape[0], ho, wo};
	233	+
	234	+ size_t inputBtyes = VectorReduceMul(inputShape) * FP16_SIZE;
	235	+ size_t weightBytes = VectorReduceMul(weightShape) * FP16_SIZE;
	236	+ size_t outputBytes = VectorReduceMul(outputShape) * FP16_SIZE;
	237	+ size_t workspaceSize = SIZE_1K * SIZE_1K * NUM_16;
	238	+ uint8_t* input = (uint8_t*)AscendC::GmAlloc(inputBtyes);
	239	+ uint8_t* weight = (uint8_t*)AscendC::GmAlloc(weightBytes);
	240	+ uint8_t* output = (uint8_t*)AscendC::GmAlloc(outputBytes);
	241	+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceSize);
	242	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(Conv2DTilingData));
	243	+
	244	+ memset(workspace, 0, workspaceSize);
	245	+
	246	+ Conv2DTilingData* tilingData = reinterpret_cast<Conv2DTilingData*>(tiling);
	247	+ TilingInput tilingInput = {inputShape, weightShape, outputShape, pads, strides, dilations};
	248	+ SetTilingData(tilingData, tilingInput);
	249	+
	250	+ auto extend_conv2d_func = [](GM_ADDR x, GM_ADDR filter, GM_ADDR bias, GM_ADDR offset_w,
	251	+ GM_ADDR scale0, GM_ADDR relu_weight0, GM_ADDR clip_value0, GM_ADDR scale1, GM_ADDR relu_weight1,
	252	+ GM_ADDR clip_value1, GM_ADDR y0, GM_ADDR y1, GM_ADDR workspace, GM_ADDR tiling) {
	253	+ ::extend_conv2d<0, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0>(
	254	+ x, filter, bias, offset_w, scale0, relu_weight0, clip_value0, scale1, relu_weight1,
	255	+ clip_value1, y0, y1, workspace, tiling);
	256	+ };
	257	+ ICPU_RUN_KF(extend_conv2d_func, numBlocks,
	258	+ input, weight, nullptr, nullptr,
	259	+ nullptr, nullptr, nullptr,
	260	+ nullptr, nullptr, nullptr,
	261	+ output, nullptr,
	262	+ workspace, tiling);
	263	+
	264	+ AscendC::GmFree(input);
	265	+ AscendC::GmFree(weight);
	266	+ AscendC::GmFree(output);
	267	+ AscendC::GmFree(workspace);
	268	+ AscendC::GmFree(tiling);
	269	+}
	270	+} // namespace
	271	+TEST_F(ExtendConv2DKernelTest, extend_conv2d_kernel_test_base)
	272	+{
	273	+ std::vector<uint64_t> inputShape = {1, 16, 16, 16};
	274	+ std::vector<uint64_t> weightShape = {16, 16, 1, 1};
	275	+
	276	+ TestSimpleKernel(inputShape, weightShape);
	277	+}

Aconv/quant_conv3d/tests/ut/op_kernel/CMakeLists.txt+13-0

@@ -0,0 +1,13 @@
	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	+if ((UT_TEST_ALL OR OP_KERNEL_UT) AND NOT UT_DONE)
	12	+ AddOpTestCase(quant_conv3d "ascend950pr_9599" "-DDTYPE_X=half -DDTYPE_FILTER=half -DDTYPE_Y=half" conv3d_v2)
	13	+endif()

Aconv/quant_conv3d/tests/ut/op_kernel/quant_conv3d_tiling_def.h+42-0

@@ -0,0 +1,42 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+/*!
	12	+ * \file quant_conv3d_tiling_def.h
	13	+ * \brief Kernel UT tiling header for quant_conv3d (uses Conv3DV2TilingDataV2).
	14	+ */
	15	+
	16	+#ifndef QUANT_CONV3D_TILING_DEF_H
	17	+#define QUANT_CONV3D_TILING_DEF_H
	18	+
	19	+#include <cstring>
	20	+#include "kernel_tiling/kernel_tiling.h"
	21	+#include "conv3d_v2/conv3d_v2_tiling_data.h"
	22	+
	23	+#define __forceinline__
	24	+
	25	+namespace Ops {
	26	+namespace NN {
	27	+namespace Conv3dV2 {
	28	+
	29	+inline void InitQuantConv3dTilingData(uint8_t* tiling, Conv3DV2TilingDataV2* constData)
	30	+{
	31	+ memcpy(constData, tiling, sizeof(Conv3DV2TilingDataV2));
	32	+}
	33	+
	34	+} // namespace Conv3dV2
	35	+} // namespace NN
	36	+} // namespace Ops
	37	+
	38	+#define GET_TILING_DATA(tilingData, tilingArg) \
	39	+ Ops::NN::Conv3dV2::Conv3DV2TilingDataV2 tilingData; \
	40	+ Ops::NN::Conv3dV2::InitQuantConv3dTilingData(tilingArg, &tilingData)
	41	+
	42	+#endif // QUANT_CONV3D_TILING_DEF_H

Aconv/quant_conv3d/tests/ut/op_kernel/test_quant_conv3d.cpp+270-0

@@ -0,0 +1,270 @@
	1	+/**
	2	+ * Copyright (c) 2026 Huawei Technologies Co., Ltd.
	3	+ * This program is free software, you can redistribute it and/or modify it under the terms and conditions of
	4	+ * CANN Open Software License Agreement Version 2.0 (the "License").
	5	+ * Please refer to the License for details. You may not use this file except in compliance with the License.
	6	+ * THIS SOFTWARE IS PROVIDED ON AN "AS IS" BASIS, WITHOUT WARRANTIES OF ANY KIND, EITHER EXPRESS OR IMPLIED,
	7	+ * INCLUDING BUT NOT LIMITED TO NON-INFRINGEMENT, MERCHANTABILITY, OR FITNESS FOR A PARTICULAR PURPOSE.
	8	+ * See LICENSE in the root of the software repository for the full text of the License.
	9	+ */
	10	+
	11	+
	12	+#include <iostream>
	13	+#include <vector>
	14	+#include <numeric>
	15	+#include <functional>
	16	+#include "gtest/gtest.h"
	17	+#include "tikicpulib.h"
	18	+#include "version/asc_devkit_version.h"
	19	+#include "quant_conv3d_tiling_def.h"
	20	+
	21	+#ifndef CONV_KERNEL
	22	+#include "quant_conv3d/quant_conv3d.cpp"
	23	+#define CONV_KERNEL
	24	+#endif
	25	+
	26	+constexpr uint8_t FP16_SIZE = 2;
	27	+constexpr uint32_t DIM1 = 1;
	28	+constexpr uint32_t DIM2 = 2;
	29	+constexpr uint32_t DIM3 = 3;
	30	+constexpr uint32_t DIM4 = 4;
	31	+constexpr uint32_t SIZE_1K = 1024;
	32	+constexpr uint32_t NUM_16 = 16;
	33	+
	34	+using Ops::NN::Conv3dV2::Conv3DV2TilingDataV2;
	35	+
	36	+class QuantConv3DKernelTest : public testing::Test {
	37	+protected:
	38	+ static void SetUpTestCase() {
	39	+ std::cout << "QuantConv3DKernelTest SetUp." << std::endl;
	40	+ }
	41	+
	42	+ static void TearDownTestCase() {
	43	+ std::cout << "QuantConv3DKernelTest TearDown." << std::endl;
	44	+ }
	45	+};
	46	+namespace {
	47	+struct Conv3DShape {
	48	+ uint64_t shape;
	49	+ uint64_t pad1;
	50	+ uint64_t pad2;
	51	+ uint64_t dilation;
	52	+ uint64_t stride;
	53	+ uint64_t kernelSize;
	54	+};
	55	+
	56	+uint64_t CalcOutputShape(const Conv3DShape& convShape)
	57	+{
	58	+ return ((convShape.shape + convShape.pad1 + convShape.pad2 - convShape.dilation * (convShape.kernelSize - 1) - 1) /
	59	+ convShape.stride + 1);
	60	+}
	61	+
	62	+size_t VectorReduceMul(const std::vector<uint64_t>& vec)
	63	+{
	64	+ return std::accumulate(std::begin(vec), std::end(vec), 1, std::multiplies<size_t>());
	65	+}
	66	+
	67	+void SetQuantConv3dTiling(Conv3DV2TilingDataV2* tiling, const std::vector<uint64_t>& inputShape,
	68	+ const std::vector<uint64_t>& weightShape, const std::vector<uint64_t>& outputShape)
	69	+{
	70	+ tiling->din = inputShape[DIM2];
	71	+ tiling->hin = inputShape[DIM3];
	72	+ tiling->win = inputShape[DIM4];
	73	+ tiling->dout = outputShape[DIM2];
	74	+ tiling->hout = outputShape[DIM3];
	75	+ tiling->wout = outputShape[DIM4];
	76	+ tiling->batch = inputShape[0];
	77	+ tiling->cin = inputShape[1];
	78	+ tiling->cout = weightShape[0];
	79	+ tiling->kd = weightShape[DIM2];
	80	+ tiling->kh = weightShape[DIM3];
	81	+ tiling->kw = weightShape[DIM4];
	82	+
	83	+ tiling->orgDi = inputShape[DIM2];
	84	+ tiling->orgHi = inputShape[DIM3];
	85	+ tiling->orgWi = inputShape[DIM4];
	86	+ tiling->orgDo = outputShape[DIM2];
	87	+ tiling->orgHo = outputShape[DIM3];
	88	+ tiling->orgWo = outputShape[DIM4];
	89	+ tiling->orgHixWi = inputShape[DIM3] * inputShape[DIM4];
	90	+ tiling->orgHoxWo = outputShape[DIM3] * outputShape[DIM4];
	91	+
	92	+ tiling->singleCoreBatch = inputShape[0];
	93	+ tiling->singleCoreDo = outputShape[DIM2];
	94	+ tiling->singleCoreM = outputShape[DIM2] * outputShape[DIM3] * outputShape[DIM4];
	95	+ tiling->singleCoreHo = outputShape[DIM3];
	96	+ tiling->singleCoreWo = 0;
	97	+
	98	+ tiling->orgCi = inputShape[1];
	99	+ tiling->orgCo = weightShape[0];
	100	+ tiling->singleCoreCi = inputShape[1];
	101	+ tiling->singleCoreCo = weightShape[0];
	102	+ tiling->kernelD = weightShape[DIM2];
	103	+ tiling->kernelH = weightShape[DIM3];
	104	+ tiling->kernelW = weightShape[DIM4];
	105	+ tiling->kernelHxkernelW = weightShape[DIM3] * weightShape[DIM4];
	106	+ tiling->kernelHxkernelWxkernelD = weightShape[DIM2] * weightShape[DIM3] * weightShape[DIM4];
	107	+
	108	+ tiling->batchDim = 1;
	109	+ tiling->doDim = 1;
	110	+ tiling->mDim = 1;
	111	+ tiling->wDim = 1;
	112	+ tiling->nDim = 1;
	113	+ tiling->groupDim = 1;
	114	+ tiling->hoDim = 1;
	115	+
	116	+ tiling->strideD = 1;
	117	+ tiling->strideH = 1;
	118	+ tiling->strideW = 1;
	119	+ tiling->dilationD = 1;
	120	+ tiling->dilationH = 1;
	121	+ tiling->dilationW = 1;
	122	+ tiling->padHead = 0;
	123	+ tiling->padTail = 0;
	124	+ tiling->padTop = 0;
	125	+ tiling->padBottom = 0;
	126	+ tiling->padLeft = 0;
	127	+ tiling->padRight = 0;
	128	+
	129	+ tiling->groups = 1;
	130	+ tiling->enlarge = 0;
	131	+ tiling->singleCoreGroups = 0;
	132	+ tiling->singleCoreGroupOpt = 0;
	133	+ tiling->groupOpt = 0;
	134	+ tiling->cinOpt = 0;
	135	+ tiling->coutOpt = 0;
	136	+
	137	+ uint32_t ci = inputShape[1];
	138	+ uint32_t kernelVol = weightShape[DIM2] * weightShape[DIM3] * weightShape[DIM4];
	139	+ uint32_t kL1 = ci * kernelVol;
	140	+ tiling->kAL1 = kL1;
	141	+ tiling->kBL1 = kL1;
	142	+ tiling->kAL1Tail = 0;
	143	+ tiling->kBL1Tail = 0;
	144	+ tiling->nBL1 = NUM_16;
	145	+
	146	+ tiling->hoL1 = NUM_16;
	147	+ tiling->hoL0 = NUM_16;
	148	+ tiling->woL1 = 0;
	149	+ tiling->woL0 = 0;
	150	+ tiling->kL0 = NUM_16;
	151	+ tiling->nL0 = NUM_16;
	152	+ tiling->mL0 = NUM_16;
	153	+
	154	+ tiling->mAL1 = NUM_16;
	155	+ tiling->pBufferFlag = 0;
	156	+ tiling->fmapKStride = 1;
	157	+ tiling->weightKStride = 1;
	158	+ tiling->cinOffsetBlockInGM = tiling->kAL1 / kernelVol * inputShape[DIM3] * inputShape[DIM4];
	159	+ tiling->coutOffsetBlock = ci * kernelVol;
	160	+ tiling->nL1DivBlockSize = tiling->nBL1 / NUM_16;
	161	+ tiling->mStep = NUM_16;
	162	+ tiling->kStep = 1;
	163	+ tiling->nStep = 1;
	164	+
	165	+ tiling->cinAInCore = tiling->kAL1 / kernelVol;
	166	+ tiling->cinATailInCore = tiling->cinAInCore;
	167	+ tiling->cinBInCore = tiling->kBL1 / kernelVol;
	168	+ tiling->cinBTailInCore = tiling->cinBInCore;
	169	+ tiling->cin1InAL1 = tiling->kAL1 / kernelVol;
	170	+ tiling->cin1InAL1Tail = tiling->cin1InAL1;
	171	+ tiling->aL1SpaceSize = tiling->cinAInCore * tiling->orgHixWi * FP16_SIZE;
	172	+ tiling->multiNBL1 = 1;
	173	+
	174	+ tiling->nL0xk0 = NUM_16;
	175	+ tiling->nBL1DivnL0 = 1;
	176	+ tiling->mAL1DivmL0 = 1;
	177	+ tiling->KBL1Divk0 = tiling->kBL1 / NUM_16;
	178	+ tiling->KBL1TailDivk0 = 0;
	179	+
	180	+ tiling->oriWixk0 = inputShape[DIM4];
	181	+ tiling->oriHixOriWixk0 = inputShape[DIM3] * inputShape[DIM4];
	182	+ tiling->cin1xOriHixOriWixk0 = tiling->cin1InAL1 * inputShape[DIM3] * inputShape[DIM4];
	183	+ tiling->kL0xorgCoAlignN0 = NUM_16;
	184	+
	185	+ tiling->biasFullLoadFlag = 1;
	186	+ tiling->fixpParamsFullLoadFlag = 1;
	187	+ tiling->hasBias = 0;
	188	+ tiling->hasScale = 0;
	189	+ tiling->iterateMNOrder = 0;
	190	+ tiling->hf32Enable = 0;
	191	+ tiling->hf32TransMode = 0;
	192	+ tiling->outputOrder = 1;
	193	+ tiling->offsetx = 0;
	194	+ tiling->roundMode = 0;
	195	+
	196	+ tiling->mUB = 0;
	197	+ tiling->nUB = 0;
	198	+ tiling->scaleAndBiasLoadType = 0;
	199	+ tiling->workspaceSize = 0;
	200	+}
	201	+
	202	+void SetUnionDataXt(Conv3DV2TilingDataV2* tiling)
	203	+{
	204	+ uint64_t xt = 0;
	205	+ xt \|= static_cast<uint64_t>(tiling->strideW & 0x3f);
	206	+ xt \|= (static_cast<uint64_t>(tiling->strideH & 0x3f) << 6);
	207	+ xt \|= (static_cast<uint64_t>(tiling->kw & 0xff) << 12);
	208	+ xt \|= (static_cast<uint64_t>(tiling->kh & 0xff) << 20);
	209	+ xt \|= (static_cast<uint64_t>(tiling->dilationW & 0xff) << 28);
	210	+ xt \|= (static_cast<uint64_t>(tiling->dilationH & 0xff) << 36);
	211	+ xt \|= ((static_cast<uint64_t>(tiling->kw) & 0x100) >> 8) << 44;
	212	+ xt \|= ((static_cast<uint64_t>(tiling->kh) & 0x100) >> 8) << 45;
	213	+ tiling->unionDataXt = xt;
	214	+}
	215	+
	216	+void TestSimpleKernel(const std::vector<uint64_t>& inputShape, const std::vector<uint64_t>& weightShape)
	217	+{
	218	+ AscendC::SetKernelMode(KernelMode::MIX_MODE);
	219	+ const uint64_t numBlocks = 1;
	220	+ Conv3DShape convShapeD =
	221	+ {inputShape[DIM2], static_cast<uint64_t>(0), static_cast<uint64_t>(0), static_cast<uint64_t>(1),
	222	+ static_cast<uint64_t>(1), weightShape[DIM2]};
	223	+ Conv3DShape convShapeH =
	224	+ {inputShape[DIM3], static_cast<uint64_t>(0), static_cast<uint64_t>(0), static_cast<uint64_t>(1),
	225	+ static_cast<uint64_t>(1), weightShape[DIM3]};
	226	+ Conv3DShape convShapeW =
	227	+ {inputShape[DIM4], static_cast<uint64_t>(0), static_cast<uint64_t>(0), static_cast<uint64_t>(1),
	228	+ static_cast<uint64_t>(1), weightShape[DIM4]};
	229	+ uint64_t dout = CalcOutputShape(convShapeD);
	230	+ uint64_t hout = CalcOutputShape(convShapeH);
	231	+ uint64_t wout = CalcOutputShape(convShapeW);
	232	+ std::vector<uint64_t> outputShape = {inputShape[0], weightShape[0], dout, hout, wout};
	233	+
	234	+ size_t inputBtyes = VectorReduceMul(inputShape) * FP16_SIZE;
	235	+ size_t weightBytes = VectorReduceMul(weightShape) * FP16_SIZE;
	236	+ size_t outputBytes = VectorReduceMul(outputShape) * FP16_SIZE;
	237	+ size_t workspaceSize = SIZE_1K * SIZE_1K * NUM_16;
	238	+ uint8_t* input = (uint8_t*)AscendC::GmAlloc(inputBtyes);
	239	+ uint8_t* weight = (uint8_t*)AscendC::GmAlloc(weightBytes);
	240	+ uint8_t* output = (uint8_t*)AscendC::GmAlloc(outputBytes);
	241	+ uint8_t* workspace = (uint8_t*)AscendC::GmAlloc(workspaceSize);
	242	+ uint8_t* tiling = (uint8_t*)AscendC::GmAlloc(sizeof(Conv3DV2TilingDataV2));
	243	+
	244	+ memset(workspace, 0, workspaceSize);
	245	+
	246	+ Conv3DV2TilingDataV2* tilingData = reinterpret_cast<Conv3DV2TilingDataV2*>(tiling);
	247	+ SetQuantConv3dTiling(tilingData, inputShape, weightShape, outputShape);
	248	+ SetUnionDataXt(tilingData);
	249	+
	250	+ auto quant_conv3d_func = [](GM_ADDR x, GM_ADDR filter, GM_ADDR scale, GM_ADDR bias, GM_ADDR offset,
	251	+ GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
	252	+ ::quant_conv3d<0, 0, 0, 0, 1, 0, 0>(
	253	+ x, filter, scale, bias, offset, y, workspace, tiling);
	254	+ };
	255	+ ICPU_RUN_KF(quant_conv3d_func, numBlocks, input, weight, nullptr, nullptr, nullptr, output, workspace, tiling);
	256	+
	257	+ AscendC::GmFree(input);
	258	+ AscendC::GmFree(weight);
	259	+ AscendC::GmFree(output);
	260	+ AscendC::GmFree(workspace);
	261	+ AscendC::GmFree(tiling);
	262	+}
	263	+} // namespace
	264	+TEST_F(QuantConv3DKernelTest, quant_conv3d_kernel_test_base)
	265	+{
	266	+ std::vector<uint64_t> inputShape = {1, 16, 4, 4, 4};
	267	+ std::vector<uint64_t> weightShape = {16, 16, 1, 1, 1};
	268	+
	269	+ TestSimpleKernel(inputShape, weightShape);
	270	+}