fix(nms): 对齐 BatchMultiClassNonMaxSuppression 和 SortedNMS 的旧 InferShape 协议-ops-cv-AtomGit

Mobjdetect/batch_multi_class_non_max_suppression/op_host/arch35/batch_multi_class_non_max_suppression_tiling_arch35.cpp+29-9

	@@ -156,8 +156,12 @@ ge::graphStatus ParseRequiredInputsAndAttrs(gert::TilingContext* context, Parsed
156	OP_CHECK_IF(boxesShape.GetDim(3) != 4, OP_LOGE(context, "boxes last dimension must be 4."),	156	OP_CHECK_IF(boxesShape.GetDim(3) != 4, OP_LOGE(context, "boxes last dimension must be 4."),
157	return ge::GRAPH_FAILED);	157	return ge::GRAPH_FAILED);
158	}	158	}
159	- params.classesNum = scoresShape.GetDim(2);	159	+ constexpr size_t kScoresFirstInnerAxis = 1;
160	- OP_CHECK_IF(scoresShape.GetDim(0) != params.batch \|\| scoresShape.GetDim(1) != params.boxesNum \|\|	160	+ constexpr size_t kScoresSecondInnerAxis = 2;
		161	+ const size_t scoresClassAxis = params.transposeBox ? kScoresFirstInnerAxis : kScoresSecondInnerAxis;
		162	+ const size_t scoresBoxesAxis = params.transposeBox ? kScoresSecondInnerAxis : kScoresFirstInnerAxis;
		163	+ params.classesNum = scoresShape.GetDim(scoresClassAxis);
		164	+ OP_CHECK_IF(scoresShape.GetDim(0) != params.batch \|\| scoresShape.GetDim(scoresBoxesAxis) != params.boxesNum \|\|
161	(params.boxClassesNum != 1 && params.boxClassesNum != params.classesNum),	165	(params.boxClassesNum != 1 && params.boxClassesNum != params.classesNum),
162	OP_LOGE(context, "boxes and scores shapes are incompatible."), return ge::GRAPH_FAILED);	166	OP_LOGE(context, "boxes and scores shapes are incompatible."), return ge::GRAPH_FAILED);
163	OP_CHECK_IF(params.batch > std::numeric_limits<int64_t>::max() / params.classesNum,	167	OP_CHECK_IF(params.batch > std::numeric_limits<int64_t>::max() / params.classesNum,
	@@ -227,6 +231,14 @@ ge::graphStatus ValidateOptionalInputs(gert::TilingContext* context, ParsedTilin
227		231
228	ge::graphStatus ValidateOutputContract(gert::TilingContext* context, const ParsedTilingParams& params)	232	ge::graphStatus ValidateOutputContract(gert::TilingContext* context, const ParsedTilingParams& params)
229	{	233	{
		234	+ constexpr int64_t kBoxCoordinateCount = 4;
		235	+ constexpr int64_t kPaddedCountStride = 8;
		236	+ constexpr size_t kBoxesOutputRank = 3;
		237	+ constexpr size_t kMatrixOutputRank = 2;
		238	+ constexpr size_t kVectorOutputRank = 1;
		239	+ constexpr size_t kBatchAxis = 0;
		240	+ constexpr size_t kFirstInnerAxis = 1;
		241	+ constexpr size_t kSecondInnerAxis = 2;
230	const gert::StorageShape* outputBoxes = context->GetOutputShape(kNmsedBoxesIndex);	242	const gert::StorageShape* outputBoxes = context->GetOutputShape(kNmsedBoxesIndex);
231	const gert::StorageShape* outputScores = context->GetOutputShape(kNmsedScoresIndex);	243	const gert::StorageShape* outputScores = context->GetOutputShape(kNmsedScoresIndex);
232	const gert::StorageShape* outputClasses = context->GetOutputShape(kNmsedClassesIndex);	244	const gert::StorageShape* outputClasses = context->GetOutputShape(kNmsedClassesIndex);
	@@ -239,13 +251,21 @@ ge::graphStatus ValidateOutputContract(gert::TilingContext* context, const Parse
239	const gert::Shape& outputScoresShape = outputScores->GetStorageShape();	251	const gert::Shape& outputScoresShape = outputScores->GetStorageShape();
240	const gert::Shape& outputClassesShape = outputClasses->GetStorageShape();	252	const gert::Shape& outputClassesShape = outputClasses->GetStorageShape();
241	const gert::Shape& outputNumShape = outputNum->GetStorageShape();	253	const gert::Shape& outputNumShape = outputNum->GetStorageShape();
242	- OP_CHECK_IF(outputBoxesShape.GetDimNum() != 3 \|\| outputBoxesShape.GetDim(0) != params.batch \|\|	254	+ OP_CHECK_IF(outputBoxesShape.GetDimNum() != kBoxesOutputRank \|\|
243	- outputBoxesShape.GetDim(1) != params.maxTotalSize \|\| outputBoxesShape.GetDim(2) != 4 \|\|	255	+ outputBoxesShape.GetDim(kBatchAxis) != params.batch \|\|
244	- outputScoresShape.GetDimNum() != 2 \|\| outputScoresShape.GetDim(0) != params.batch \|\|	256	+ outputBoxesShape.GetDim(kFirstInnerAxis) !=
245	- outputScoresShape.GetDim(1) != params.maxTotalSize \|\| outputClassesShape.GetDimNum() != 2 \|\|	257	+ (params.transposeBox ? kBoxCoordinateCount : params.maxTotalSize) \|\|
246	- outputClassesShape.GetDim(0) != params.batch \|\|	258	+ outputBoxesShape.GetDim(kSecondInnerAxis) !=
247	- outputClassesShape.GetDim(1) != params.maxTotalSize \|\| outputNumShape.GetDimNum() != 1 \|\|	259	+ (params.transposeBox ? params.maxTotalSize : kBoxCoordinateCount) \|\|
248	- outputNumShape.GetDim(0) != params.batch,	260	+ outputScoresShape.GetDimNum() != kMatrixOutputRank \|\|
		261	+ outputScoresShape.GetDim(kBatchAxis) != params.batch \|\|
		262	+ outputScoresShape.GetDim(kFirstInnerAxis) != params.maxTotalSize \|\|
		263	+ outputClassesShape.GetDimNum() != kMatrixOutputRank \|\|
		264	+ outputClassesShape.GetDim(kBatchAxis) != params.batch \|\|
		265	+ outputClassesShape.GetDim(kFirstInnerAxis) != params.maxTotalSize \|\|
		266	+ outputNumShape.GetDimNum() != (params.transposeBox ? kMatrixOutputRank : kVectorOutputRank) \|\|
		267	+ outputNumShape.GetDim(kBatchAxis) != params.batch \|\|
		268	+ (params.transposeBox && outputNumShape.GetDim(kFirstInnerAxis) != kPaddedCountStride),
249	OP_LOGE(context, "Output shapes do not match BatchMultiClassNonMaxSuppression contract."),	269	OP_LOGE(context, "Output shapes do not match BatchMultiClassNonMaxSuppression contract."),
250	return ge::GRAPH_FAILED);	270	return ge::GRAPH_FAILED);
251	const auto* outputBoxesDesc = context->GetOutputDesc(kNmsedBoxesIndex);	271	const auto* outputBoxesDesc = context->GetOutputDesc(kNmsedBoxesIndex);

Mobjdetect/batch_multi_class_non_max_suppression/op_host/batch_multi_class_non_max_suppression_infershape.cpp+38-23

	@@ -12,33 +12,37 @@
12	#include "register/op_impl_registry.h"	12	#include "register/op_impl_registry.h"
13		13
14	namespace {	14	namespace {
15	-constexpr int64_t kBoxesIndex = 0;
16	constexpr int64_t kScoresIndex = 1;	15	constexpr int64_t kScoresIndex = 1;
17	constexpr int64_t kNmsedBoxesIndex = 0;	16	constexpr int64_t kNmsedBoxesIndex = 0;
18	constexpr int64_t kNmsedScoresIndex = 1;	17	constexpr int64_t kNmsedScoresIndex = 1;
19	constexpr int64_t kNmsedClassesIndex = 2;	18	constexpr int64_t kNmsedClassesIndex = 2;
20	constexpr int64_t kNmsedNumIndex = 3;	19	constexpr int64_t kNmsedNumIndex = 3;
21	constexpr int64_t kMaxTotalSizeAttrIndex = 3;	20	constexpr int64_t kMaxTotalSizeAttrIndex = 3;
		21	+constexpr int64_t kTransposeBoxAttrIndex = 5;
22	} // namespace	22	} // namespace
23		23
24	namespace ops {	24	namespace ops {
25	static ge::graphStatus InferShapeForBatchMultiClassNonMaxSuppression(gert::InferShapeContext* context)	25	static ge::graphStatus InferShapeForBatchMultiClassNonMaxSuppression(gert::InferShapeContext* context)
26	{	26	{
27	- const gert::Shape* boxesShape = context->GetInputShape(kBoxesIndex);
28	const gert::Shape* scoresShape = context->GetInputShape(kScoresIndex);	27	const gert::Shape* scoresShape = context->GetInputShape(kScoresIndex);
29	const gert::RuntimeAttrs* attrs = context->GetAttrs();	28	const gert::RuntimeAttrs* attrs = context->GetAttrs();
30	- OP_CHECK_NULL_WITH_CONTEXT(context, boxesShape);
31	OP_CHECK_NULL_WITH_CONTEXT(context, scoresShape);	29	OP_CHECK_NULL_WITH_CONTEXT(context, scoresShape);
32	OP_CHECK_NULL_WITH_CONTEXT(context, attrs);	30	OP_CHECK_NULL_WITH_CONTEXT(context, attrs);
33	-	31	+ OP_CHECK_IF(scoresShape->GetDimNum() == 0, OP_LOGE(context, "scores shape is empty."), return ge::GRAPH_FAILED);
34	- OP_CHECK_IF(boxesShape->GetDimNum() != 4 \|\| scoresShape->GetDimNum() != 3,
35	- OP_LOGE(context, "boxes must be rank 4 and scores must be rank 3."), return ge::GRAPH_FAILED);
36	- const int64_t batch = boxesShape->GetDim(0);
37	- const int64_t scoresBatch = scoresShape->GetDim(0);
38	const int64_t* maxTotalSize = attrs->GetAttrPointer<int64_t>(kMaxTotalSizeAttrIndex);	32	const int64_t* maxTotalSize = attrs->GetAttrPointer<int64_t>(kMaxTotalSizeAttrIndex);
39	OP_CHECK_NULL_WITH_CONTEXT(context, maxTotalSize);	33	OP_CHECK_NULL_WITH_CONTEXT(context, maxTotalSize);
40	- OP_CHECK_IF(batch <= 0 \|\| scoresBatch != batch \|\| *maxTotalSize <= 0,	34	+ const bool* transposeBoxAttr = attrs->GetAttrPointer<bool>(kTransposeBoxAttrIndex);
41	- OP_LOGE(context, "Invalid batch dimension or max_total_size."), return ge::GRAPH_FAILED);	35	+ const bool transposeBox = transposeBoxAttr != nullptr && *transposeBoxAttr;
		36	+
		37	+ // Match the legacy GE inference, including its conservative unknown-batch
		38	+ // rule. The fusion pass uses transpose_box for both input and output layouts.
		39	+ int64_t batch = scoresShape->GetDim(0);
		40	+ for (size_t dim = 0; dim < scoresShape->GetDimNum(); ++dim) {
		41	+ if (scoresShape->GetDim(dim) < 0) {
		42	+ batch = ge::UNKNOWN_DIM;
		43	+ break;
		44	+ }
		45	+ }
42		46
43	gert::Shape* nmsedBoxesShape = context->GetOutputShape(kNmsedBoxesIndex);	47	gert::Shape* nmsedBoxesShape = context->GetOutputShape(kNmsedBoxesIndex);
44	gert::Shape* nmsedScoresShape = context->GetOutputShape(kNmsedScoresIndex);	48	gert::Shape* nmsedScoresShape = context->GetOutputShape(kNmsedScoresIndex);
	@@ -49,26 +53,37 @@ static ge::graphStatus InferShapeForBatchMultiClassNonMaxSuppression(gert::Infer
49	OP_CHECK_NULL_WITH_CONTEXT(context, nmsedClassesShape);	53	OP_CHECK_NULL_WITH_CONTEXT(context, nmsedClassesShape);
50	OP_CHECK_NULL_WITH_CONTEXT(context, nmsedNumShape);	54	OP_CHECK_NULL_WITH_CONTEXT(context, nmsedNumShape);
51		55
52	- nmsedBoxesShape->SetDimNum(3);	56	+ constexpr int64_t kBoxCoordinateCount = 4;
53	- nmsedBoxesShape->SetDim(0, batch);	57	+ constexpr int64_t kPaddedCountStride = 8;
54	- nmsedBoxesShape->SetDim(1, *maxTotalSize);	58	+ constexpr size_t kBoxesOutputRank = 3;
55	- nmsedBoxesShape->SetDim(2, 4);	59	+ constexpr size_t kMatrixOutputRank = 2;
		60	+ constexpr size_t kVectorOutputRank = 1;
		61	+ constexpr size_t kBatchAxis = 0;
		62	+ constexpr size_t kFirstInnerAxis = 1;
		63	+ constexpr size_t kSecondInnerAxis = 2;
		64	+ nmsedBoxesShape->SetDimNum(kBoxesOutputRank);
		65	+ nmsedBoxesShape->SetDim(kBatchAxis, batch);
		66	+ nmsedBoxesShape->SetDim(kFirstInnerAxis, transposeBox ? kBoxCoordinateCount : *maxTotalSize);
		67	+ nmsedBoxesShape->SetDim(kSecondInnerAxis, transposeBox ? *maxTotalSize : kBoxCoordinateCount);
56	for (gert::Shape* output : {nmsedScoresShape, nmsedClassesShape}) {	68	for (gert::Shape* output : {nmsedScoresShape, nmsedClassesShape}) {
57	- output->SetDimNum(2);	69	+ output->SetDimNum(kMatrixOutputRank);
58	- output->SetDim(0, batch);	70	+ output->SetDim(kBatchAxis, batch);
59	- output->SetDim(1, *maxTotalSize);	71	+ output->SetDim(kFirstInnerAxis, *maxTotalSize);
		72	+ }
		73	+ nmsedNumShape->SetDimNum(transposeBox ? kMatrixOutputRank : kVectorOutputRank);
		74	+ nmsedNumShape->SetDim(kBatchAxis, batch);
		75	+ if (transposeBox) {
		76	+ nmsedNumShape->SetDim(kFirstInnerAxis, kPaddedCountStride);
60	}	77	}
61	- nmsedNumShape->SetDimNum(1);
62	- nmsedNumShape->SetDim(0, batch);
63	return ge::GRAPH_SUCCESS;	78	return ge::GRAPH_SUCCESS;
64	}	79	}
65		80
66	static ge::graphStatus InferDataTypeForBatchMultiClassNonMaxSuppression(gert::InferDataTypeContext* context)	81	static ge::graphStatus InferDataTypeForBatchMultiClassNonMaxSuppression(gert::InferDataTypeContext* context)
67	{	82	{
68	- const ge::DataType boxesType = context->GetInputDataType(kBoxesIndex);	83	+ const ge::DataType scoresType = context->GetInputDataType(kScoresIndex);
69	- context->SetOutputDataType(kNmsedBoxesIndex, boxesType);	84	+ context->SetOutputDataType(kNmsedBoxesIndex, scoresType);
70	- context->SetOutputDataType(kNmsedScoresIndex, boxesType);	85	+ context->SetOutputDataType(kNmsedScoresIndex, scoresType);
71	- context->SetOutputDataType(kNmsedClassesIndex, boxesType);	86	+ context->SetOutputDataType(kNmsedClassesIndex, scoresType);
72	context->SetOutputDataType(kNmsedNumIndex, ge::DT_INT32);	87	context->SetOutputDataType(kNmsedNumIndex, ge::DT_INT32);
73	return ge::GRAPH_SUCCESS;	88	return ge::GRAPH_SUCCESS;
74	}	89	}

Mobjdetect/batch_multi_class_non_max_suppression/op_kernel/arch35/batch_multi_class_non_max_suppression_kernel.h+43-21

	@@ -20,6 +20,9 @@
20	namespace BatchMultiClassNonMaxSuppressionOp {	20	namespace BatchMultiClassNonMaxSuppressionOp {
21	using namespace AscendC;	21	using namespace AscendC;
22		22
		23	+constexpr int64_t kBoxCoordinateCount = 4;
		24	+constexpr int64_t kPaddedCountStride = 8;
		25	+constexpr int64_t kUnpaddedCountStride = 1;
23	constexpr uint32_t kGatherThreadNum32 = 1024;	26	constexpr uint32_t kGatherThreadNum32 = 1024;
24	constexpr uint32_t kGatherThreadNum64 = 512;	27	constexpr uint32_t kGatherThreadNum64 = 512;
25	constexpr uint32_t kMergeThreadNum = 256;	28	constexpr uint32_t kMergeThreadNum = 256;
	@@ -42,9 +45,8 @@ struct GatherIndexType<true> {
42	template <bool Use32Bit>	45	template <bool Use32Bit>
43	using GatherIndex = typename GatherIndexType<Use32Bit>::type;	46	using GatherIndex = typename GatherIndexType<Use32Bit>::type;
44		47
45	-// Scores are strided by class and boxes can be either [B, N, q, 4] or	48	+// Fused inputs use scores [B, C, N] and boxes [B, q, 4, N].
46	-// [B, q, 4, N]. This is an irregular GM access pattern, so use the 950 SIMT	49	+// Use the 950 SIMT unit to compact one (batch, class) task into five contiguous FP32 arrays.
47	-// unit to compact one (batch, class) task into five contiguous FP32 arrays.
48	// Subsequent score reduction and IoU work is entirely vectorized on UB tiles.	50	// Subsequent score reduction and IoU work is entirely vectorized on UB tiles.
49	template <typename T, bool Use32Bit>	51	template <typename T, bool Use32Bit>
50	__simt_vf__ __aicore__	52	__simt_vf__ __aicore__
	@@ -66,7 +68,7 @@ __launch_bounds__(Use32Bit ? kGatherThreadNum32 : kGatherThreadNum64) inline voi
66	}	68	}
67	for (IndexT boxIndex = static_cast<IndexT>(threadIdx.x); boxIndex < boxesNum;	69	for (IndexT boxIndex = static_cast<IndexT>(threadIdx.x); boxIndex < boxesNum;
68	boxIndex += static_cast<IndexT>(blockDim.x)) {	70	boxIndex += static_cast<IndexT>(blockDim.x)) {
69	- const IndexT scoreOffset = (batchIndex * boxesNum + boxIndex) * classesNum + classIndex;	71	+ const IndexT scoreOffset = (batchIndex * classesNum + classIndex) * boxesNum + boxIndex;
70	stageScores[boxIndex] = boxIndex < validBoxes ? static_cast<float>(scores[scoreOffset]) : kNoCandidate;	72	stageScores[boxIndex] = boxIndex < validBoxes ? static_cast<float>(scores[scoreOffset]) : kNoCandidate;
71	// Each [B, q, 4, N] field is contiguous. Their bases are computed	73	// Each [B, q, 4, N] field is contiguous. Their bases are computed
72	// by AIV before the SIMT launch, so this path only needs a per-box	74	// by AIV before the SIMT launch, so this path only needs a per-box
	@@ -128,9 +130,16 @@ template <typename T>
128	__simt_vf__ __aicore__ __launch_bounds__(kMergeThreadNum) inline void GatherMergedOutput(	130	__simt_vf__ __aicore__ __launch_bounds__(kMergeThreadNum) inline void GatherMergedOutput(
129	const __gm__ float* classBoxes, const __gm__ float* mergeScores, const __gm__ int32_t* mergeIndices,	131	const __gm__ float* classBoxes, const __gm__ float* mergeScores, const __gm__ int32_t* mergeIndices,
130	__gm__ T* nmsedBoxes, __gm__ T* nmsedScores, __gm__ T* nmsedClasses, __gm__ int32_t* nmsedNum, uint64_t batchIndex,	132	__gm__ T* nmsedBoxes, __gm__ T* nmsedScores, __gm__ T* nmsedClasses, __gm__ int32_t* nmsedNum, uint64_t batchIndex,
131	- uint64_t classesNum, uint64_t maxSizePerClass, uint64_t maxTotalSize)	133	+ uint64_t classesNum, uint64_t maxSizePerClass, uint64_t maxTotalSize, bool transposeBox)
132	{	134	{
133	- const uint64_t validOutputCount = static_cast<uint64_t>(nmsedNum[batchIndex]);	135	+ constexpr uint64_t kXMinCoordinate = 1;
		136	+ constexpr uint64_t kYMaxCoordinate = 2;
		137	+ constexpr uint64_t kXMaxCoordinate = 3;
		138	+ const uint64_t countOffset = batchIndex * (transposeBox ? kPaddedCountStride : kUnpaddedCountStride);
		139	+ const uint64_t validOutputCount = static_cast<uint64_t>(nmsedNum[countOffset]);
		140	+ if (transposeBox && threadIdx.x < kPaddedCountStride - kUnpaddedCountStride) {
		141	+ nmsedNum[countOffset + threadIdx.x + 1] = 0;
		142	+ }
134	for (uint64_t outputIndex = static_cast<uint64_t>(threadIdx.x); outputIndex < maxTotalSize;	143	for (uint64_t outputIndex = static_cast<uint64_t>(threadIdx.x); outputIndex < maxTotalSize;
135	outputIndex += static_cast<uint64_t>(blockDim.x)) {	144	outputIndex += static_cast<uint64_t>(blockDim.x)) {
136	const uint64_t outputOffset = batchIndex * maxTotalSize + outputIndex;	145	const uint64_t outputOffset = batchIndex * maxTotalSize + outputIndex;
	@@ -140,20 +149,29 @@ __simt_vf__ __aicore__ __launch_bounds__(kMergeThreadNum) inline void GatherMerg
140	const uint64_t classIndex = flatIndex / maxSizePerClass;	149	const uint64_t classIndex = flatIndex / maxSizePerClass;
141	const uint64_t classPosition = flatIndex % maxSizePerClass;	150	const uint64_t classPosition = flatIndex % maxSizePerClass;
142	const uint64_t candidateOffset = (batchIndex * classesNum + classIndex) * maxSizePerClass + classPosition;	151	const uint64_t candidateOffset = (batchIndex * classesNum + classIndex) * maxSizePerClass + classPosition;
143	- const uint64_t boxOffset = candidateOffset * 4;	152	+ const uint64_t boxOffset = candidateOffset * kBoxCoordinateCount;
144	- const uint64_t outputBoxOffset = outputOffset * 4;	153	+ const uint64_t outputBoxOffset = transposeBox ?
		154	+ batchIndex * maxTotalSize * kBoxCoordinateCount + outputIndex :
		155	+ outputOffset * kBoxCoordinateCount;
		156	+ const uint64_t coordinateStride = transposeBox ? maxTotalSize : 1;
145	nmsedBoxes[outputBoxOffset] = static_cast<T>(classBoxes[boxOffset]);	157	nmsedBoxes[outputBoxOffset] = static_cast<T>(classBoxes[boxOffset]);
146	- nmsedBoxes[outputBoxOffset + 1] = static_cast<T>(classBoxes[boxOffset + 1]);	158	+ nmsedBoxes[outputBoxOffset + kXMinCoordinate * coordinateStride] = static_cast<T>(
147	- nmsedBoxes[outputBoxOffset + 2] = static_cast<T>(classBoxes[boxOffset + 2]);	159	+ classBoxes[boxOffset + kXMinCoordinate]);
148	- nmsedBoxes[outputBoxOffset + 3] = static_cast<T>(classBoxes[boxOffset + 3]);	160	+ nmsedBoxes[outputBoxOffset + kYMaxCoordinate * coordinateStride] = static_cast<T>(
		161	+ classBoxes[boxOffset + kYMaxCoordinate]);
		162	+ nmsedBoxes[outputBoxOffset + kXMaxCoordinate * coordinateStride] = static_cast<T>(
		163	+ classBoxes[boxOffset + kXMaxCoordinate]);
149	nmsedScores[outputOffset] = static_cast<T>(score);	164	nmsedScores[outputOffset] = static_cast<T>(score);
150	nmsedClasses[outputOffset] = static_cast<T>(classIndex);	165	nmsedClasses[outputOffset] = static_cast<T>(classIndex);
151	} else {	166	} else {
152	- const uint64_t outputBoxOffset = outputOffset * 4;	167	+ const uint64_t outputBoxOffset = transposeBox ?
		168	+ batchIndex * maxTotalSize * kBoxCoordinateCount + outputIndex :
		169	+ outputOffset * kBoxCoordinateCount;
		170	+ const uint64_t coordinateStride = transposeBox ? maxTotalSize : 1;
153	nmsedBoxes[outputBoxOffset] = static_cast<T>(0);	171	nmsedBoxes[outputBoxOffset] = static_cast<T>(0);
154	- nmsedBoxes[outputBoxOffset + 1] = static_cast<T>(0);	172	+ nmsedBoxes[outputBoxOffset + kXMinCoordinate * coordinateStride] = static_cast<T>(0);
155	- nmsedBoxes[outputBoxOffset + 2] = static_cast<T>(0);	173	+ nmsedBoxes[outputBoxOffset + kYMaxCoordinate * coordinateStride] = static_cast<T>(0);
156	- nmsedBoxes[outputBoxOffset + 3] = static_cast<T>(0);	174	+ nmsedBoxes[outputBoxOffset + kXMaxCoordinate * coordinateStride] = static_cast<T>(0);
157	nmsedScores[outputOffset] = static_cast<T>(0);	175	nmsedScores[outputOffset] = static_cast<T>(0);
158	nmsedClasses[outputOffset] = static_cast<T>(0);	176	nmsedClasses[outputOffset] = static_cast<T>(0);
159	}	177	}
	@@ -204,7 +222,7 @@ __simt_callee__ __aicore__ __attribute__((always_inline)) inline void SiftClassH
204	__simt_vf__ __aicore__ __launch_bounds__(kMergeThreadNum) inline void MergeClassOutput(	222	__simt_vf__ __aicore__ __launch_bounds__(kMergeThreadNum) inline void MergeClassOutput(
205	const __gm__ float* classScores, __gm__ float* classPositions, __gm__ float* mergeScores,	223	const __gm__ float* classScores, __gm__ float* classPositions, __gm__ float* mergeScores,
206	__gm__ int32_t* mergeIndices, __gm__ int32_t* nmsedNum, __ubuf__ float* heapScores, __ubuf__ int32_t* heapIndices,	224	__gm__ int32_t* mergeIndices, __gm__ int32_t* nmsedNum, __ubuf__ float* heapScores, __ubuf__ int32_t* heapIndices,
207	- uint64_t batchIndex, uint64_t classesNum, uint64_t maxSizePerClass, uint64_t maxTotalSize)	225	+ uint64_t batchIndex, uint64_t classesNum, uint64_t maxSizePerClass, uint64_t maxTotalSize, bool transposeBox)
208	{	226	{
209	const uint64_t classBase = batchIndex * classesNum;	227	const uint64_t classBase = batchIndex * classesNum;
210	const uint64_t outputBase = batchIndex * maxTotalSize;	228	const uint64_t outputBase = batchIndex * maxTotalSize;
	@@ -238,7 +256,8 @@ __simt_vf__ __aicore__ __launch_bounds__(kMergeThreadNum) inline void MergeClass
238	}	256	}
239	SiftClassHeap(heapScores, heapIndices, classesNum, 0);	257	SiftClassHeap(heapScores, heapIndices, classesNum, 0);
240	}	258	}
241	- nmsedNum[batchIndex] = static_cast<int32_t>(validOutputCount);	259	+ nmsedNum[batchIndex * (transposeBox ? kPaddedCountStride : kUnpaddedCountStride)] = static_cast<int32_t>(
		260	+ validOutputCount);
242	return;	261	return;
243	}	262	}
244		263
	@@ -272,7 +291,8 @@ __simt_vf__ __aicore__ __launch_bounds__(kMergeThreadNum) inline void MergeClass
272	static_cast<int32_t>(position);	291	static_cast<int32_t>(position);
273	classPositions[classBase + bestClass] = static_cast<float>(position + 1);	292	classPositions[classBase + bestClass] = static_cast<float>(position + 1);
274	}	293	}
275	- nmsedNum[batchIndex] = static_cast<int32_t>(validOutputCount);	294	+ nmsedNum[batchIndex * (transposeBox ? kPaddedCountStride : kUnpaddedCountStride)] = static_cast<int32_t>(
		295	+ validOutputCount);
276	}	296	}
277	}	297	}
278		298
	@@ -389,7 +409,9 @@ __aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::Init(
389	nmsedBoxesGm_.SetGlobalBuffer((__gm__ T)nmsedBoxes, resultElements 4);	409	nmsedBoxesGm_.SetGlobalBuffer((__gm__ T)nmsedBoxes, resultElements 4);
390	nmsedScoresGm_.SetGlobalBuffer((__gm__ T*)nmsedScores, resultElements);	410	nmsedScoresGm_.SetGlobalBuffer((__gm__ T*)nmsedScores, resultElements);
391	nmsedClassesGm_.SetGlobalBuffer((__gm__ T*)nmsedClasses, resultElements);	411	nmsedClassesGm_.SetGlobalBuffer((__gm__ T*)nmsedClasses, resultElements);
392	- nmsedNumGm_.SetGlobalBuffer((__gm__ int32_t*)nmsedNum, tilingData_->batch);	412	+ nmsedNumGm_.SetGlobalBuffer(
		413	+ (__gm__ int32_t*)nmsedNum,
		414	+ tilingData_->batch * (tilingData_->transposeBox != 0U ? kPaddedCountStride : kUnpaddedCountStride));
393	InitWorkspace();	415	InitWorkspace();
394	InitTileBuffers();	416	InitTileBuffers();
395	}	417	}
	@@ -789,7 +811,7 @@ __aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::MergeBatch(int
789	reinterpret_cast<__ubuf__ float*>(mergeInputScoresLocal_.GetPhyAddr()),	811	reinterpret_cast<__ubuf__ float*>(mergeInputScoresLocal_.GetPhyAddr()),
790	reinterpret_cast<__ubuf__ int32_t*>(mergeInputIndicesLocal_.GetPhyAddr()), static_cast<uint64_t>(batchIndex),	812	reinterpret_cast<__ubuf__ int32_t*>(mergeInputIndicesLocal_.GetPhyAddr()), static_cast<uint64_t>(batchIndex),
791	static_cast<uint64_t>(tilingData_->classesNum), static_cast<uint64_t>(tilingData_->maxSizePerClass),	813	static_cast<uint64_t>(tilingData_->classesNum), static_cast<uint64_t>(tilingData_->maxSizePerClass),
792	- static_cast<uint64_t>(tilingData_->maxTotalSize));	814	+ static_cast<uint64_t>(tilingData_->maxTotalSize), tilingData_->transposeBox != 0U);
793		815
794	asc_vf_call<GatherMergedOutput<T>>(	816	asc_vf_call<GatherMergedOutput<T>>(
795	dim3{kMergeThreadNum}, (__gm__ float)classBoxesGm_.GetPhyAddr(), (__gm__ float)mergeScoresGm_.GetPhyAddr(),	817	dim3{kMergeThreadNum}, (__gm__ float)classBoxesGm_.GetPhyAddr(), (__gm__ float)mergeScoresGm_.GetPhyAddr(),
	@@ -797,7 +819,7 @@ __aicore__ inline void BatchMultiClassNonMaxSuppressionKernel<T>::MergeBatch(int
797	(__gm__ T)nmsedScoresGm_.GetPhyAddr(), (__gm__ T)nmsedClassesGm_.GetPhyAddr(),	819	(__gm__ T)nmsedScoresGm_.GetPhyAddr(), (__gm__ T)nmsedClassesGm_.GetPhyAddr(),
798	(__gm__ int32_t*)nmsedNumGm_.GetPhyAddr(), static_cast<uint64_t>(batchIndex),	820	(__gm__ int32_t*)nmsedNumGm_.GetPhyAddr(), static_cast<uint64_t>(batchIndex),
799	static_cast<uint64_t>(tilingData_->classesNum), static_cast<uint64_t>(tilingData_->maxSizePerClass),	821	static_cast<uint64_t>(tilingData_->classesNum), static_cast<uint64_t>(tilingData_->maxSizePerClass),
800	- static_cast<uint64_t>(tilingData_->maxTotalSize));	822	+ static_cast<uint64_t>(tilingData_->maxTotalSize), tilingData_->transposeBox != 0U);
801	}	823	}
802		824
803	template <typename T>	825	template <typename T>

Mobjdetect/batch_multi_class_non_max_suppression/tests/ut/op_host/test_batch_multi_class_non_max_suppression_infershape.cpp+87-12

	@@ -11,11 +11,13 @@
11	#include <gtest/gtest.h>	11	#include <gtest/gtest.h>
12	#include <vector>	12	#include <vector>
13		13
		14	+#include "base/registry/op_impl_space_registry_v2.h"
		15	+#include "op_infer_datatype_context_builder.h"
14	#include "infershape_case_executor.h"	16	#include "infershape_case_executor.h"
15	#include "infershape_context_faker.h"	17	#include "infershape_context_faker.h"
16		18
17	namespace {	19	namespace {
18	-std::vector<gert::InfershapeContextPara::OpAttr> GetAttrs(int64_t maxTotalSize)	20	+std::vector<gert::InfershapeContextPara::OpAttr> GetAttrs(int64_t maxTotalSize, bool transposeBox = false)
19	{	21	{
20	return {	22	return {
21	gert::InfershapeContextPara::OpAttr("score_threshold", Ops::Cv::AnyValue::CreateFrom<float>(0.3F)),	23	gert::InfershapeContextPara::OpAttr("score_threshold", Ops::Cv::AnyValue::CreateFrom<float>(0.3F)),
	@@ -23,7 +25,7 @@ std::vector<gert::InfershapeContextPara::OpAttr> GetAttrs(int64_t maxTotalSize)
23	gert::InfershapeContextPara::OpAttr("max_size_per_class", Ops::Cv::AnyValue::CreateFrom<int64_t>(3)),	25	gert::InfershapeContextPara::OpAttr("max_size_per_class", Ops::Cv::AnyValue::CreateFrom<int64_t>(3)),
24	gert::InfershapeContextPara::OpAttr("max_total_size", Ops::Cv::AnyValue::CreateFrom<int64_t>(maxTotalSize)),	26	gert::InfershapeContextPara::OpAttr("max_total_size", Ops::Cv::AnyValue::CreateFrom<int64_t>(maxTotalSize)),
25	gert::InfershapeContextPara::OpAttr("change_coordinate_frame", Ops::Cv::AnyValue::CreateFrom<bool>(false)),	27	gert::InfershapeContextPara::OpAttr("change_coordinate_frame", Ops::Cv::AnyValue::CreateFrom<bool>(false)),
26	- gert::InfershapeContextPara::OpAttr("transpose_box", Ops::Cv::AnyValue::CreateFrom<bool>(false)),	28	+ gert::InfershapeContextPara::OpAttr("transpose_box", Ops::Cv::AnyValue::CreateFrom<bool>(transposeBox)),
27	};	29	};
28	}	30	}
29	} // namespace	31	} // namespace
	@@ -41,15 +43,88 @@ TEST(BatchMultiClassNonMaxSuppressionInfershape, derives_public_output_layout)
41	ExecuteTestCase(context, ge::GRAPH_SUCCESS, {{2, 5, 4}, {2, 5}, {2, 5}, {2}});	43	ExecuteTestCase(context, ge::GRAPH_SUCCESS, {{2, 5, 4}, {2, 5}, {2, 5}, {2}});
42	}	44	}
43		45
44	-TEST(BatchMultiClassNonMaxSuppressionInfershape, rejects_non_rank_four_boxes)	46	+namespace {
		47	+void ExpectLegacyShape(const gert::StorageShape& boxes, const gert::StorageShape& scores, bool transposeBox,
		48	+ int64_t expectedBatch)
45	{	49	{
46	- gert::InfershapeContextPara context(	50	+ constexpr int64_t kMaxTotalSize = 5;
47	- "BatchMultiClassNonMaxSuppression",	51	+ constexpr int64_t kBoxCoordinateCount = 4;
48	- {{{{2, 8, 4}, {2, 8, 4}}, ge::DT_FLOAT, ge::FORMAT_ND}, {{{2, 8, 3}, {2, 8, 3}}, ge::DT_FLOAT, ge::FORMAT_ND}},	52	+ constexpr int64_t kPaddedCountStride = 8;
49	- {{{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},	53	+ gert::InfershapeContextPara context("BatchMultiClassNonMaxSuppression",
50	- {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},	54	+ {{boxes, ge::DT_FLOAT, ge::FORMAT_ND}, {scores, ge::DT_FLOAT, ge::FORMAT_ND}},
51	- {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},	55	+ {{{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
52	- {{{}, {}}, ge::DT_INT32, ge::FORMAT_ND}},	56	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
53	- GetAttrs(5));	57	+ {{{}, {}}, ge::DT_FLOAT, ge::FORMAT_ND},
54	- ExecuteTestCase(context, ge::GRAPH_FAILED);	58	+ {{{}, {}}, ge::DT_INT32, ge::FORMAT_ND}},
		59	+ GetAttrs(kMaxTotalSize, transposeBox));
		60	+ const std::vector<int64_t> expectedBoxes = transposeBox ? std::vector<int64_t>{expectedBatch, kBoxCoordinateCount,
		61	+ kMaxTotalSize} :
		62	+ std::vector<int64_t>{expectedBatch, kMaxTotalSize,
		63	+ kBoxCoordinateCount};
		64	+ const std::vector<int64_t> expectedNum = transposeBox ? std::vector<int64_t>{expectedBatch, kPaddedCountStride} :
		65	+ std::vector<int64_t>{expectedBatch};
		66	+ ExecuteTestCase(context, ge::GRAPH_SUCCESS,
		67	+ {expectedBoxes, {expectedBatch, kMaxTotalSize}, {expectedBatch, kMaxTotalSize}, expectedNum});
		68	+}
		69	+} // namespace
		70	+
		71	+TEST(BatchMultiClassNonMaxSuppressionInfershape, derives_legacy_transposed_outputs)
		72	+{
		73	+ // The old fusion path provides scores as [B,C,N], with transposed boxes and padded counts.
		74	+ ExpectLegacyShape({{2, 1, 4, 8}, {2, 1, 4, 8}}, {{2, 3, 8}, {2, 3, 8}}, true, 2);
		75	+}
		76	+
		77	+TEST(BatchMultiClassNonMaxSuppressionInfershape, derives_batch_from_scores_with_unknown_boxes_rank)
		78	+{
		79	+ ExpectLegacyShape({{ge::UNKNOWN_DIM_NUM}, {ge::UNKNOWN_DIM_NUM}}, {{2, 8, 3}, {2, 8, 3}}, false, 2);
		80	+}
		81	+
		82	+TEST(BatchMultiClassNonMaxSuppressionInfershape, preserves_unknown_batch_for_dynamic_scores)
		83	+{
		84	+ const std::vector<gert::StorageShape> dynamicScores = {{{ge::UNKNOWN_DIM_NUM}, {ge::UNKNOWN_DIM_NUM}},
		85	+ {{ge::UNKNOWN_DIM, 8, 3}, {ge::UNKNOWN_DIM, 8, 3}},
		86	+ {{2, ge::UNKNOWN_DIM, 3}, {2, ge::UNKNOWN_DIM, 3}},
		87	+ {{2, 8, ge::UNKNOWN_DIM}, {2, 8, ge::UNKNOWN_DIM}}};
		88	+ for (const bool transposeBox : {false, true}) {
		89	+ for (size_t index = 0; index < dynamicScores.size(); ++index) {
		90	+ SCOPED_TRACE(::testing::Message() << "transpose_box=" << transposeBox << ", case=" << index);
		91	+ ExpectLegacyShape({{2, 8, 1, 4}, {2, 8, 1, 4}}, dynamicScores[index], transposeBox, ge::UNKNOWN_DIM);
		92	+ }
		93	+ }
		94	+}
		95	+
		96	+TEST(BatchMultiClassNonMaxSuppressionInfershape, floating_outputs_follow_scores_dtype)
		97	+{
		98	+ constexpr size_t kInputCount = 2;
		99	+ constexpr size_t kOutputCount = 4;
		100	+ constexpr size_t kFloatingOutputCount = 3;
		101	+ constexpr size_t kBoxesInputIndex = 0;
		102	+ constexpr size_t kScoresInputIndex = 1;
		103	+ constexpr size_t kCountOutputIndex = kFloatingOutputCount;
		104	+ auto registry = gert::DefaultOpImplSpaceRegistryV2::GetInstance().GetSpaceRegistry();
		105	+ ASSERT_NE(registry, nullptr);
		106	+ const auto* opImpl = registry->GetOpImpl("BatchMultiClassNonMaxSuppression");
		107	+ ASSERT_NE(opImpl, nullptr);
		108	+ ASSERT_NE(opImpl->infer_datatype, nullptr);
		109	+ for (const ge::DataType scoresType : {ge::DT_FLOAT, ge::DT_FLOAT16}) {
		110	+ SCOPED_TRACE(scoresType);
		111	+ gert::OpInferDataTypeContextBuilder builder;
		112	+ builder.OpType("BatchMultiClassNonMaxSuppression").OpName("BatchMultiClassNonMaxSuppression");
		113	+ builder.IONum(kInputCount, kOutputCount);
		114	+ // Deliberately differ so this detects accidentally taking the boxes dtype.
		115	+ const ge::DataType boxesType = scoresType == ge::DT_FLOAT ? ge::DT_FLOAT16 : ge::DT_FLOAT;
		116	+ builder.InputTensorDesc(kBoxesInputIndex, boxesType, ge::FORMAT_ND, ge::FORMAT_ND);
		117	+ builder.InputTensorDesc(kScoresInputIndex, scoresType, ge::FORMAT_ND, ge::FORMAT_ND);
		118	+ for (size_t index = 0; index < kOutputCount; ++index) {
		119	+ builder.OutputTensorDesc(index, ge::FORMAT_ND, ge::FORMAT_ND);
		120	+ }
		121	+ auto holder = builder.Build();
		122	+ auto* context = holder.GetContext();
		123	+ ASSERT_NE(context, nullptr);
		124	+ ASSERT_EQ(opImpl->infer_datatype(context), ge::GRAPH_SUCCESS);
		125	+ for (size_t index = 0; index < kFloatingOutputCount; ++index) {
		126	+ EXPECT_EQ(context->GetOutputDataType(index), scoresType);
		127	+ }
		128	+ EXPECT_EQ(context->GetOutputDataType(kCountOutputIndex), ge::DT_INT32);
		129	+ }
55	}	130	}

Mobjdetect/sorted_nms/op_host/sorted_nms_infershape.cpp+4-16

	@@ -14,8 +14,6 @@
14	namespace {	14	namespace {
15	constexpr int64_t INPUT_BOXES = 0;	15	constexpr int64_t INPUT_BOXES = 0;
16	constexpr int64_t OUTPUT_SELECTED_INDICES = 0;	16	constexpr int64_t OUTPUT_SELECTED_INDICES = 0;
17	-constexpr int64_t BOX_RANK = 2;
18	-constexpr int64_t BOX_COORDS = 4;
19	constexpr int64_t OUTPUT_RANK = 1;	17	constexpr int64_t OUTPUT_RANK = 1;
20	} // namespace	18	} // namespace
21		19
	@@ -29,15 +27,8 @@ static ge::graphStatus InferShapeForSortedNMS(gert::InferShapeContext* context)
29	return ge::GRAPH_FAILED;	27	return ge::GRAPH_FAILED;
30	}	28	}
31		29
32	- if (boxesShape->GetDimNum() != BOX_RANK) {	30	+ // Shape inference does not require a known input rank. Concrete input
33	- OP_LOGE(context, "boxes rank must be 2");	31	+ // constraints belong to tiling, as in the legacy runtime callback.
34	- return ge::GRAPH_FAILED;
35	- }
36	- if (boxesShape->GetDim(1) != ge::UNKNOWN_DIM && boxesShape->GetDim(1) != BOX_COORDS) {
37	- OP_LOGE(context, "boxes second dim must be 4");
38	- return ge::GRAPH_FAILED;
39	- }
40	-
41	selectedShape->SetDimNum(OUTPUT_RANK);	32	selectedShape->SetDimNum(OUTPUT_RANK);
42	selectedShape->SetDim(0, ge::UNKNOWN_DIM);	33	selectedShape->SetDim(0, ge::UNKNOWN_DIM);
43	return ge::GRAPH_SUCCESS;	34	return ge::GRAPH_SUCCESS;
	@@ -50,17 +41,14 @@ static ge::graphStatus InferShapeRangeForSortedNMS(gert::InferShapeRangeContext*
50	auto boxesRange = context->GetInputShapeRange(INPUT_BOXES);	41	auto boxesRange = context->GetInputShapeRange(INPUT_BOXES);
51	auto selectedRange = context->GetOutputShapeRange(OUTPUT_SELECTED_INDICES);	42	auto selectedRange = context->GetOutputShapeRange(OUTPUT_SELECTED_INDICES);
52	OP_CHECK_NULL_WITH_CONTEXT(context, boxesRange);	43	OP_CHECK_NULL_WITH_CONTEXT(context, boxesRange);
53	- OP_CHECK_NULL_WITH_CONTEXT(context, boxesRange->GetMin());
54	OP_CHECK_NULL_WITH_CONTEXT(context, boxesRange->GetMax());	44	OP_CHECK_NULL_WITH_CONTEXT(context, boxesRange->GetMax());
55	OP_CHECK_NULL_WITH_CONTEXT(context, selectedRange);	45	OP_CHECK_NULL_WITH_CONTEXT(context, selectedRange);
56	OP_CHECK_NULL_WITH_CONTEXT(context, selectedRange->GetMin());	46	OP_CHECK_NULL_WITH_CONTEXT(context, selectedRange->GetMin());
57	OP_CHECK_NULL_WITH_CONTEXT(context, selectedRange->GetMax());	47	OP_CHECK_NULL_WITH_CONTEXT(context, selectedRange->GetMax());
58		48
59	- OP_CHECK_IF(boxesRange->GetMin()->GetDimNum() != BOX_RANK \|\| boxesRange->GetMax()->GetDimNum() != BOX_RANK,	49	+ OP_CHECK_IF(boxesRange->GetMax()->GetDimNum() == 0,
60	- OP_LOGE(context, "boxes shape range rank must be 2"), return ge::GRAPH_FAILED);	50	+ OP_LOGE(context, "boxes shape range must contain a first dimension"), return ge::GRAPH_FAILED);
61	const int64_t maxBoxesNum = boxesRange->GetMax()->GetDim(0);	51	const int64_t maxBoxesNum = boxesRange->GetMax()->GetDim(0);
62	- OP_CHECK_IF(maxBoxesNum < 0, OP_LOGE(context, "boxes shape range first dim must be known"),
63	- return ge::GRAPH_FAILED);
64		52
65	selectedRange->GetMin()->SetDimNum(OUTPUT_RANK);	53	selectedRange->GetMin()->SetDimNum(OUTPUT_RANK);
66	selectedRange->GetMin()->SetDim(0, 0);	54	selectedRange->GetMin()->SetDim(0, 0);