fix(indexer_quant_cache): 修复GEIR内存泄漏并补充数值golden-ops-transformer-AtomGit

已合并

fix(indexer_quant_cache): 修复GEIR内存泄漏并补充数值golden #11302

chendunyang创建于 9月4日从qq_37913898/ops-transformer： fix/indexer-quant-cache-geir-leak-golden 合入到 cann/ops-transformer： master

fix(indexer_quant_cache): 修复GEIR内存泄漏并补充数值golden #11302

已合并

chendunyang创建于 9月4日

从qq_37913898/ops-transformer fix/indexer-quant-cache-geir-leak-golden合入到cann/ops-transformer master

共 2 个文件变更+157-1

Mattention/indexer_quant_cache/examples/test_geir_indexer_quant_cache.cpp+4-1

	@@ -42,6 +42,7 @@ using std::map;
42	using std::string;	42	using std::string;
43	using std::vector;	43	using std::vector;
44		44
		45	+// clang-format off
45	#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \	46	#define ADD_INPUT(intputIndex, intputName, intputDtype, inputShape) \
46	vector<int64_t> placeholder##intputIndex##_shape = inputShape; \	47	vector<int64_t> placeholder##intputIndex##_shape = inputShape; \
47	auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \	48	auto placeholder##intputIndex = op::Data("placeholder" + intputIndex).set_attr_index(0); \
	@@ -72,6 +73,7 @@ using std::vector;
72	do { \	73	do { \
73	printf(message, ##__VA_ARGS__); \	74	printf(message, ##__VA_ARGS__); \
74	} while (0)	75	} while (0)
		76	+// clang-format on
75		77
76	string GetTime()	78	string GetTime()
77	{	79	{
	@@ -182,7 +184,7 @@ int CreateOppInGraph(DataType inDtype, std::vector<ge::Tensor> &input, std::vect
182	const int64_t tokenNum = 1024;	184	const int64_t tokenNum = 1024;
183	const int64_t d = 128;	185	const int64_t d = 128;
184	const int64_t blockNum = 128;	186	const int64_t blockNum = 128;
185	- const int64_t blockSizeDim = 16; // num_slots = 128*16 = 2048	187	+ const int64_t blockSizeDim = 16; // num_slots = 128*16 = 2048
186	std::vector<int64_t> cacheShape = {blockNum, blockSizeDim, 1, d};	188	std::vector<int64_t> cacheShape = {blockNum, blockSizeDim, 1, d};
187	std::vector<int64_t> cacheScaleShape = {blockNum, blockSizeDim, 1, 1};	189	std::vector<int64_t> cacheScaleShape = {blockNum, blockSizeDim, 1, 1};
188	std::vector<int64_t> xShape = {tokenNum, d};	190	std::vector<int64_t> xShape = {tokenNum, d};
	@@ -304,6 +306,7 @@ int main(int argc, char *argv[])
304	std::string warning_str(warning_msg.GetString());	306	std::string warning_str(warning_msg.GetString());
305	std::cout << "Warning message: " << warning_str << std::endl;	307	std::cout << "Warning message: " << warning_str << std::endl;
306	printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());	308	printf("%s - INFO - [XIR]: Start to finalize ir graph session\n", GetTime().c_str());
		309	+ delete session;
307	ret = ge::GEFinalize();	310	ret = ge::GEFinalize();
308	if (ret != SUCCESS) {	311	if (ret != SUCCESS) {
309	printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());	312	printf("%s - INFO - [XIR]: Finalize ir graph session failed\n", GetTime().c_str());

Aattention/indexer_quant_cache/tests/ut/op_kernel/arch35/test_indexer_quant_cache_golden.cpp+153-0

@@ -0,0 +1,153 @@
	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 test_indexer_quant_cache_golden.cpp
	13	+ * \brief Numeric golden tests for the IndexerQuantCache Normal quantization branch.
	14	+ */
	15	+#include <array>
	16	+#include <cassert>
	17	+#include <cstdint>
	18	+#include <cstring>
	19	+#include <gtest/gtest.h>
	20	+
	21	+#ifdef __CCE_KT_TEST__
	22	+#include "data_utils.h"
	23	+#include "tikicpulib.h"
	24	+#endif
	25	+
	26	+extern "C" __global__ __aicore__ void indexer_quant_cache(GM_ADDR cache, GM_ADDR cache_scale, GM_ADDR x,
	27	+ GM_ADDR slot_mapping, GM_ADDR cache_out,
	28	+ GM_ADDR cache_scale_out, GM_ADDR workspace, GM_ADDR tiling);
	29	+
	30	+namespace {
	31	+constexpr int64_t BS = 4;
	32	+constexpr int64_t D = 64;
	33	+constexpr int64_t CACHE_ROWS = BS;
	34	+constexpr int64_t NORMAL_QUANT_MODE = 1;
	35	+constexpr int64_t TILING_KEY = 10001;
	36	+constexpr size_t SYS_WORKSPACE = 16 * 1024 * 1024;
	37	+constexpr std::array<int32_t, BS> SLOT_MAPPING = {0, -1, 2, 3};
	38	+
	39	+struct NormalGolden {
	40	+ int64_t roundScale;
	41	+ std::array<uint8_t, 2> cache;
	42	+ uint32_t scale;
	43	+};
	44	+
	45	+template <typename T>
	46	+T *GmAlloc(size_t size)
	47	+{
	48	+ T ptr = reinterpret_cast<T >(AscendC::GmAlloc(size));
	49	+ assert(ptr != nullptr && "GM allocation failed");
	50	+ return ptr;
	51	+}
	52	+
	53	+void FillInput(uint8_t *input)
	54	+{
	55	+ uint16_t halfInput = reinterpret_cast<uint16_t >(input);
	56	+ for (int64_t i = 0; i < BS * D; ++i) {
	57	+ halfInput[i] = (i & 1) ? 0xB800 : 0x3C00; // -0.5, 1.0
	58	+ }
	59	+}
	60	+
	61	+IndexerQuantCacheTilingData MakeTiling(const NormalGolden &golden)
	62	+{
	63	+ IndexerQuantCacheTilingData tiling;
	64	+ std::memset(&tiling, 0, sizeof(tiling));
	65	+ tiling.bs = BS;
	66	+ tiling.d = D;
	67	+ tiling.scaleCol = 1;
	68	+ tiling.rowOfFormerBlock = BS;
	69	+ tiling.rowOfTailBlock = BS;
	70	+ tiling.rowLoopOfFormerBlock = BS;
	71	+ tiling.rowLoopOfTailBlock = BS;
	72	+ tiling.rowFactor = 1;
	73	+ tiling.tailRowFactorOfFormerBlock = 1;
	74	+ tiling.tailRowFactorOfTailBlock = 1;
	75	+ tiling.quantMode = NORMAL_QUANT_MODE;
	76	+ tiling.roundScale = golden.roundScale;
	77	+ tiling.scalesAttr = 1.0f;
	78	+ tiling.blockSize = 1;
	79	+ tiling.cacheRowStride = D;
	80	+ tiling.cacheBlockStride = D;
	81	+ tiling.scaleRowStride = 1;
	82	+ tiling.scaleBlockStride = 1;
	83	+ return tiling;
	84	+}
	85	+
	86	+void RunAndCheckNormalGolden(const NormalGolden &golden)
	87	+{
	88	+ AscendC::SetKernelMode(KernelMode::AIV_MODE);
	89	+
	90	+ constexpr size_t X_SIZE = BS * D * sizeof(uint16_t);
	91	+ constexpr size_t SLOT_SIZE = BS * sizeof(int32_t);
	92	+ constexpr size_t CACHE_SIZE = CACHE_ROWS * D * sizeof(uint8_t);
	93	+ constexpr size_t SCALE_SIZE = CACHE_ROWS * sizeof(uint32_t);
	94	+
	95	+ uint8_t *cache = GmAlloc<uint8_t>(CACHE_SIZE);
	96	+ uint8_t *cacheScale = GmAlloc<uint8_t>(SCALE_SIZE);
	97	+ uint8_t *x = GmAlloc<uint8_t>(X_SIZE);
	98	+ uint8_t *slotMapping = GmAlloc<uint8_t>(SLOT_SIZE);
	99	+ uint8_t *cacheOut = GmAlloc<uint8_t>(CACHE_SIZE);
	100	+ uint8_t *cacheScaleOut = GmAlloc<uint8_t>(SCALE_SIZE);
	101	+ uint8_t *workspace = GmAlloc<uint8_t>(SYS_WORKSPACE);
	102	+ uint8_t *tiling = GmAlloc<uint8_t>(sizeof(IndexerQuantCacheTilingData));
	103	+
	104	+ std::memset(cache, 0, CACHE_SIZE);
	105	+ std::memset(cacheScale, 0, SCALE_SIZE);
	106	+ FillInput(x);
	107	+ std::memcpy(slotMapping, SLOT_MAPPING.data(), SLOT_SIZE);
	108	+ std::memset(cacheOut, 0, CACHE_SIZE);
	109	+ std::memset(cacheScaleOut, 0, SCALE_SIZE);
	110	+ std::memset(workspace, 0, SYS_WORKSPACE);
	111	+ const IndexerQuantCacheTilingData tilingData = MakeTiling(golden);
	112	+ std::memcpy(tiling, &tilingData, sizeof(tilingData));
	113	+
	114	+ ICPU_SET_TILING_KEY(TILING_KEY);
	115	+ ICPU_RUN_KF(indexer_quant_cache, 1, cache, cacheScale, x, slotMapping, cacheOut, cacheScaleOut, workspace, tiling);
	116	+
	117	+ int64_t cacheMismatch = 0;
	118	+ int64_t scaleMismatch = 0;
	119	+ for (int64_t row = 0; row < BS; ++row) {
	120	+ const bool skipped = SLOT_MAPPING[row] < 0;
	121	+ for (int64_t col = 0; col < D; ++col) {
	122	+ const uint8_t expected = skipped ? 0 : golden.cache[col & 1];
	123	+ cacheMismatch += cache[row * D + col] != expected;
	124	+ }
	125	+ uint32_t actualScale = 0;
	126	+ std::memcpy(&actualScale, cacheScale + row * sizeof(uint32_t), sizeof(actualScale));
	127	+ scaleMismatch += actualScale != (skipped ? 0 : golden.scale);
	128	+ }
	129	+ EXPECT_EQ(cacheMismatch, 0) << "cache differs from the literal float8_e4m3fn golden";
	130	+ EXPECT_EQ(scaleMismatch, 0) << "cache scale differs from the literal float32 golden";
	131	+
	132	+ AscendC::GmFree(cache);
	133	+ AscendC::GmFree(cacheScale);
	134	+ AscendC::GmFree(x);
	135	+ AscendC::GmFree(slotMapping);
	136	+ AscendC::GmFree(cacheOut);
	137	+ AscendC::GmFree(cacheScaleOut);
	138	+ AscendC::GmFree(workspace);
	139	+ AscendC::GmFree(tiling);
	140	+}
	141	+} // namespace
	142	+
	143	+TEST(IndexerQuantCacheNormalGolden, RoundScale)
	144	+{
	145	+ // Input alternates 1.0/-0.5. 2^-8 scale produces E4M3FN bytes 0x78/0xF0.
	146	+ RunAndCheckNormalGolden(NormalGolden{1, {0x78, 0xF0}, 0x3B800000U});
	147	+}
	148	+
	149	+TEST(IndexerQuantCacheNormalGolden, NoRoundScale)
	150	+{
	151	+ // Input alternates 1.0/-0.5. 1/448 scale produces E4M3FN bytes 0x7E/0xF6.
	152	+ RunAndCheckNormalGolden(NormalGolden{0, {0x7E, 0xF6}, 0x3B124925U});
	153	+}