[Fix] Fix static check errors detected by SPACES-pytorch-AtomGit

M.lintrunner.toml+29-38

	@@ -428,44 +428,35 @@ command = [
428	]	428	]
429	is_formatter = true	429	is_formatter = true
430		430
431	-# [[linter]]	431	+[[linter]]
432	-# code = 'SPACES'	432	+code = 'SPACES'
433	-# include_patterns = ['**']	433	+include_patterns = ['**']
434	-# exclude_patterns = [	434	+exclude_patterns = [
435	-# 'test_upstream/**',	435	+ 'test_upstream/**',
436	-# '/contrib/',	436	+ '/contrib/',
437	-# '*/.diff',	437	+ '*/.diff',
438	-# '*/.patch',	438	+ '*/.patch',
439	-# 'third_party/**',	439	+ 'third_party/**',
440	-# 'aten/src/ATen/native/vulkan/api/vk_mem_alloc.h',	440	+ 'aten/src/ATen/native/vulkan/api/vk_mem_alloc.h',
441	-# 'fb/**',	441	+ 'fb/**',
442	-# '/fb/',	442	+ '/fb/',
443	-# 'test/cpp/jit/upgrader_models/*.ptl',	443	+ 'test/cpp/jit/upgrader_models/*.ptl',
444	-# 'test/cpp/jit/upgrader_models/*.ptl.ff',	444	+ 'test/cpp/jit/upgrader_models/*.ptl.ff',
445	-# # NPUGraph logs files	445	+ '*/.md',
446	-# 'torch_npu/_logging/_internal.py',	446	+]
447	-# 'torch_npu/csrc/core/npu/NPUGraph.cpp',	447	+command = [
448	-# 'torch_npu/csrc/core/npu/NPUGraph.h',	448	+ 'python3',
449	-# 'torch_npu/csrc/npu/Graph.cpp',	449	+ 'tools/linter/adapters/grep_linter.py',
450	-# 'torch_npu/csrc/core/npu/NPUCachingAllocator.cpp',	450	+ '--pattern=[[:blank:]]$',
451	-# 'torch_npu/csrc/core/npu/NPUWorkspaceAllocator.cpp',	451	+ '--linter-name=SPACES',
452	-# 'torch_npu/npu/_graph_tree.py',	452	+ '--error-name=trailing spaces',
453	-# 'torch_npu/npu/graphs.py',	453	+ '--replace-pattern=s/[[:blank:]]+$//',
454	-# 'torch_npu/utils/_graph_tree.py',	454	+ """--error-description=\
455	-# ]	455	+ This line has trailing spaces; please remove them.\
456	-# command = [	456	+ """,
457	-# 'python3',	457	+ '--',
458	-# 'tools/linter/adapters/grep_linter.py',	458	+ '@{{PATHSFILE}}'
459	-# '--pattern=[[:blank:]]$',	459	+]
460	-# '--linter-name=SPACES',
461	-# '--error-name=trailing spaces',
462	-# '--replace-pattern=s/[[:blank:]]+$//',
463	-# """--error-description=\
464	-# This line has trailing spaces; please remove them.\
465	-# """,
466	-# '--',
467	-# '@{{PATHSFILE}}'
468	-# ]
469		460
470	# [[linter]]	461	# [[linter]]
471	# code = 'TABS'	462	# code = 'TABS'

Mbenchmarks/torchbench/extract_log.py+26-26

	@@ -8,7 +8,7 @@ import pandas as pd
8	def extract_log_info(log_file, profile_dir='./profile', output_file='log_analysis.xlsx'):	8	def extract_log_info(log_file, profile_dir='./profile', output_file='log_analysis.xlsx'):
9	"""	9	"""
10	提取日志文件中各模型的训练时间信息和profile数据	10	提取日志文件中各模型的训练时间信息和profile数据
11	-	11	+
12	Args:	12	Args:
13	log_file: 日志文件路径	13	log_file: 日志文件路径
14	profile_dir: profile数据目录路径	14	profile_dir: profile数据目录路径
	@@ -19,14 +19,14 @@ def extract_log_info(log_file, profile_dir='./profile', output_file='log_analysi
19	model_pattern = r'(npu\|cuda)\s+train\s+(\S+)'	19	model_pattern = r'(npu\|cuda)\s+train\s+(\S+)'
20	eager_pattern = r'eager.avg step time:\s([\d\.]+)\s*ms'	20	eager_pattern = r'eager.avg step time:\s([\d\.]+)\s*ms'
21	compile_pattern = r'compile.avg step time:\s([\d\.]+)\s*ms'	21	compile_pattern = r'compile.avg step time:\s([\d\.]+)\s*ms'
22	-	22	+
23	# 模式匹配：算子编译时间	23	# 模式匹配：算子编译时间
24	op_compile_time_pattern = r'op_compile_time:\s([\d\.]+)\sms'	24	op_compile_time_pattern = r'op_compile_time:\s([\d\.]+)\sms'
25	-	25	+
26	# 存储结果	26	# 存储结果
27	data = defaultdict(lambda: {	27	data = defaultdict(lambda: {
28	'accuracy': None, # 修改2: 存储完整的精度校验日志	28	'accuracy': None, # 修改2: 存储完整的精度校验日志
29	- 'eager_E2E_avg_time': None,	29	+ 'eager_E2E_avg_time': None,
30	'compile_E2E_avg_time': None,	30	'compile_E2E_avg_time': None,
31	'op_compile_time': None,	31	'op_compile_time': None,
32	'eager_OP_avg_time': None,	32	'eager_OP_avg_time': None,
	@@ -35,10 +35,10 @@ def extract_log_info(log_file, profile_dir='./profile', output_file='log_analysi
35	current_model = None	35	current_model = None
36	in_compile_block = False	36	in_compile_block = False
37	compile_block_lines = []	37	compile_block_lines = []
38	-	38	+
39	with open(log_file, 'r', encoding='utf-8') as f:	39	with open(log_file, 'r', encoding='utf-8') as f:
40	lines = f.readlines()	40	lines = f.readlines()
41	-	41	+
42	for _, line in enumerate(lines):	42	for _, line in enumerate(lines):
43	# 匹配模型名	43	# 匹配模型名
44	model_match = re.search(model_pattern, line)	44	model_match = re.search(model_pattern, line)
	@@ -50,57 +50,57 @@ def extract_log_info(log_file, profile_dir='./profile', output_file='log_analysi
50	if 'pass_accuracy' in log_line:	50	if 'pass_accuracy' in log_line:
51	data[current_model]['accuracy'] = log_line.strip()	51	data[current_model]['accuracy'] = log_line.strip()
52	break	52	break
53	-	53	+
54	# 提取op_compile_time	54	# 提取op_compile_time
55	for log_line in compile_block_lines:	55	for log_line in compile_block_lines:
56	op_compile_match = re.search(op_compile_time_pattern, log_line)	56	op_compile_match = re.search(op_compile_time_pattern, log_line)
57	if op_compile_match:	57	if op_compile_match:
58	data[current_model]['op_compile_time'] = float(op_compile_match.group(1))	58	data[current_model]['op_compile_time'] = float(op_compile_match.group(1))
59	break	59	break
60	-	60	+
61	# 重置状态	61	# 重置状态
62	current_model = model_match.group(2) # 第二个分组是模型名	62	current_model = model_match.group(2) # 第二个分组是模型名
63	in_compile_block = False	63	in_compile_block = False
64	compile_block_lines = []	64	compile_block_lines = []
65	continue	65	continue
66	-	66	+
67	# 匹配eager模式时间	67	# 匹配eager模式时间
68	if current_model:	68	if current_model:
69	eager_match = re.search(eager_pattern, line)	69	eager_match = re.search(eager_pattern, line)
70	if eager_match:	70	if eager_match:
71	data[current_model]['eager_E2E_avg_time'] = float(eager_match.group(1))	71	data[current_model]['eager_E2E_avg_time'] = float(eager_match.group(1))
72	-	72	+
73	# 匹配compile模式时间	73	# 匹配compile模式时间
74	compile_match = re.search(compile_pattern, line)	74	compile_match = re.search(compile_pattern, line)
75	if compile_match:	75	if compile_match:
76	data[current_model]['compile_E2E_avg_time'] = float(compile_match.group(1))	76	data[current_model]['compile_E2E_avg_time'] = float(compile_match.group(1))
77	in_compile_block = True	77	in_compile_block = True
78	compile_block_lines = [] # 开始收集compile块日志	78	compile_block_lines = [] # 开始收集compile块日志
79	-	79	+
80	# 收集compile块的日志	80	# 收集compile块的日志
81	if current_model and in_compile_block:	81	if current_model and in_compile_block:
82	compile_block_lines.append(line)	82	compile_block_lines.append(line)
83	-	83	+
84	# 处理最后一个模型的compile块日志	84	# 处理最后一个模型的compile块日志
85	if current_model and in_compile_block and compile_block_lines:	85	if current_model and in_compile_block and compile_block_lines:
86	for log_line in compile_block_lines:	86	for log_line in compile_block_lines:
87	if 'pass_accuracy' in log_line:	87	if 'pass_accuracy' in log_line:
88	data[current_model]['accuracy'] = log_line.strip()	88	data[current_model]['accuracy'] = log_line.strip()
89	break	89	break
90	-	90	+
91	for log_line in compile_block_lines:	91	for log_line in compile_block_lines:
92	op_compile_match = re.search(op_compile_time_pattern, log_line)	92	op_compile_match = re.search(op_compile_time_pattern, log_line)
93	if op_compile_match:	93	if op_compile_match:
94	data[current_model]['op_compile_time'] = float(op_compile_match.group(1))	94	data[current_model]['op_compile_time'] = float(op_compile_match.group(1))
95	break	95	break
96	-	96	+
97	# 需求3: 读取profile目录中的step_trace_time.csv文件	97	# 需求3: 读取profile目录中的step_trace_time.csv文件
98	profile_path = Path(profile_dir)	98	profile_path = Path(profile_dir)
99	if profile_path.exists():	99	if profile_path.exists():
100	for model_dir in profile_path.iterdir():	100	for model_dir in profile_path.iterdir():
101	if model_dir.is_dir():	101	if model_dir.is_dir():
102	model_name = model_dir.name	102	model_name = model_dir.name
103	-	103	+
104	# 读取eager模式下的step_trace_time.csv	104	# 读取eager模式下的step_trace_time.csv
105	# 修改3: 自动获取下一级目录	105	# 修改3: 自动获取下一级目录
106	eager_dir = model_dir / 'eager'	106	eager_dir = model_dir / 'eager'
	@@ -119,7 +119,7 @@ def extract_log_info(log_file, profile_dir='./profile', output_file='log_analysi
119	print(f"警告: {eager_csv_path} 中没有Computing列")	119	print(f"警告: {eager_csv_path} 中没有Computing列")
120	except Exception as e:	120	except Exception as e:
121	print(f"读取{eager_csv_path}时出错: {e}")	121	print(f"读取{eager_csv_path}时出错: {e}")
122	-	122	+
123	# 读取compile模式下的step_trace_time.csv	123	# 读取compile模式下的step_trace_time.csv
124	compile_dir = model_dir / 'compile'	124	compile_dir = model_dir / 'compile'
125	if compile_dir.exists() and compile_dir.is_dir():	125	if compile_dir.exists() and compile_dir.is_dir():
	@@ -139,7 +139,7 @@ def extract_log_info(log_file, profile_dir='./profile', output_file='log_analysi
139	print(f"读取{compile_csv_path}时出错: {e}")	139	print(f"读取{compile_csv_path}时出错: {e}")
140	else:	140	else:
141	print(f"警告: profile目录不存在: {profile_dir}")	141	print(f"警告: profile目录不存在: {profile_dir}")
142	-	142	+
143	# 转换为DataFrame	143	# 转换为DataFrame
144	df = pd.DataFrame.from_dict(data, orient='index')	144	df = pd.DataFrame.from_dict(data, orient='index')
145	df.index.name = 'model_name'	145	df.index.name = 'model_name'
	@@ -149,17 +149,17 @@ def extract_log_info(log_file, profile_dir='./profile', output_file='log_analysi
149	# 1. 计算E2E_speed_up_rate = eager_E2E_avg_time / compile_E2E_avg_time	149	# 1. 计算E2E_speed_up_rate = eager_E2E_avg_time / compile_E2E_avg_time
150	# 2. 计算OP_speed_up_rate = eager_OP_avg_time / compile_OP_avg_time	150	# 2. 计算OP_speed_up_rate = eager_OP_avg_time / compile_OP_avg_time
151	df['E2E_speed_up_rate'] = df.apply(	151	df['E2E_speed_up_rate'] = df.apply(
152	- lambda row: row['eager_E2E_avg_time'] / row['compile_E2E_avg_time']	152	+ lambda row: row['eager_E2E_avg_time'] / row['compile_E2E_avg_time']
153	- if row['compile_E2E_avg_time'] and row['compile_E2E_avg_time'] != 0 else None,	153	+ if row['compile_E2E_avg_time'] and row['compile_E2E_avg_time'] != 0 else None,
154	axis=1	154	axis=1
155	)	155	)
156	-	156	+
157	df['OP_speed_up_rate'] = df.apply(	157	df['OP_speed_up_rate'] = df.apply(
158	- lambda row: row['eager_OP_avg_time'] / row['compile_OP_avg_time']	158	+ lambda row: row['eager_OP_avg_time'] / row['compile_OP_avg_time']
159	- if row['compile_OP_avg_time'] and row['compile_OP_avg_time'] != 0 else None,	159	+ if row['compile_OP_avg_time'] and row['compile_OP_avg_time'] != 0 else None,
160	axis=1	160	axis=1
161	)	161	)
162	-	162	+
163	# 重排列顺序，使相关列更清晰	163	# 重排列顺序，使相关列更清晰
164	column_order = [	164	column_order = [
165	'model_name', 'accuracy', 'op_compile_time',	165	'model_name', 'accuracy', 'op_compile_time',
	@@ -168,7 +168,7 @@ def extract_log_info(log_file, profile_dir='./profile', output_file='log_analysi
168	]	168	]
169	existing_columns = [col for col in column_order if col in df.columns]	169	existing_columns = [col for col in column_order if col in df.columns]
170	df = df[existing_columns + [col for col in df.columns if col not in existing_columns]]	170	df = df[existing_columns + [col for col in df.columns if col not in existing_columns]]
171	-	171	+
172	# 保存到Excel	172	# 保存到Excel
173	df.to_excel(output_file, index=False)	173	df.to_excel(output_file, index=False)
174	return df	174	return df
	@@ -179,9 +179,9 @@ def main():
179	parser.add_argument('--log_file', required=True, help='日志文件路径')	179	parser.add_argument('--log_file', required=True, help='日志文件路径')
180	parser.add_argument('--profile_dir', default='./profile', help='profile数据目录路径，默认为./profile')	180	parser.add_argument('--profile_dir', default='./profile', help='profile数据目录路径，默认为./profile')
181	parser.add_argument('--output_file', default='analysis.xlsx', help='输出Excel文件路径，默认为log_analysis.xlsx')	181	parser.add_argument('--output_file', default='analysis.xlsx', help='输出Excel文件路径，默认为log_analysis.xlsx')
182	-	182	+
183	args = parser.parse_args()	183	args = parser.parse_args()
184	-	184	+
185	result = extract_log_info(	185	result = extract_log_info(
186	log_file=args.log_file,	186	log_file=args.log_file,
187	profile_dir=args.profile_dir,	187	profile_dir=args.profile_dir,

Mci/build.sh+1-1

	@@ -46,7 +46,7 @@ function parse_script_args() {
46	export PGO_MODE=1	46	export PGO_MODE=1
47	args_num=$((args_num-1))	47	args_num=$((args_num-1))
48	;;	48	;;
49	- 2)	49	+ 2)
50	export PGO_MODE=2	50	export PGO_MODE=2
51	args_num=$((args_num-1))	51	args_num=$((args_num-1))
52	;;	52	;;

Mexamples/libtorch_hccl/CMakeLists.txt+4-4

	@@ -68,9 +68,9 @@ endif()
68	add_executable(example_allreduce_hccl allreduce_hccl.cpp)	68	add_executable(example_allreduce_hccl allreduce_hccl.cpp)
69		69
70	# 链接库	70	# 链接库
71	-target_link_libraries(example_allreduce_hccl	71	+target_link_libraries(example_allreduce_hccl
72	- -ltorch	72	+ -ltorch
73	- -ltorch_cpu	73	+ -ltorch_cpu
74	- -lc10	74	+ -lc10
75	-ltorch_npu	75	-ltorch_npu
76	)	76	)

Mexamples/libtorch_hccl/allreduce_hccl.cpp+17-17

	@@ -24,64 +24,64 @@ int main(int argc, char** argv)
24	{	24	{
25	int rank = g_rank;	25	int rank = g_rank;
26	int size = g_size;	26	int size = g_size;
27	-	27	+
28	std::cout << "启动 HCCL allreduce 示例: rank=" << rank << ", size=" << size << std::endl;	28	std::cout << "启动 HCCL allreduce 示例: rank=" << rank << ", size=" << size << std::endl;
29	-	29	+
30	// 初始化NPU设备 - 使用npu字符串格式	30	// 初始化NPU设备 - 使用npu字符串格式
31	std::string device_str = "npu:" + std::to_string(rank);	31	std::string device_str = "npu:" + std::to_string(rank);
32	torch_npu::init_npu(device_str);	32	torch_npu::init_npu(device_str);
33	std::cout << "NPU设备 " << rank << " 初始化完成" << std::endl;	33	std::cout << "NPU设备 " << rank << " 初始化完成" << std::endl;
34	-	34	+
35	// 创建FileStore用于进程间通信协调	35	// 创建FileStore用于进程间通信协调
36	auto store = c10::make_intrusive<FileStore>("/tmp/c10d_hccl_example", size);	36	auto store = c10::make_intrusive<FileStore>("/tmp/c10d_hccl_example", size);
37	-	37	+
38	// 创建ProcessGroupHCCL选项	38	// 创建ProcessGroupHCCL选项
39	auto options = ProcessGroupHCCL::Options::create();	39	auto options = ProcessGroupHCCL::Options::create();
40	-	40	+
41	// 创建ProcessGroupHCCL实例	41	// 创建ProcessGroupHCCL实例
42	auto pg = c10::make_intrusive<ProcessGroupHCCL>(store, rank, size, options);	42	auto pg = c10::make_intrusive<ProcessGroupHCCL>(store, rank, size, options);
43	-	43	+
44	// 通过传NPU字符串构造NPU设备	44	// 通过传NPU字符串构造NPU设备
45	auto device = at::Device(device_str);	45	auto device = at::Device(device_str);
46	-	46	+
47	// 创建10个张量用于测试	47	// 创建10个张量用于测试
48	const auto ntensors = 10;	48	const auto ntensors = 10;
49	std::vector<at::Tensor> tensors;	49	std::vector<at::Tensor> tensors;
50	-	50	+
51	for (const auto i : c10::irange(ntensors)) {	51	for (const auto i : c10::irange(ntensors)) {
52	// 在NPU设备上创建全1张量	52	// 在NPU设备上创建全1张量
53	auto x = at::ones({1000, 16 * (i + 1)}, at::TensorOptions(device).dtype(at::kFloat));	53	auto x = at::ones({1000, 16 * (i + 1)}, at::TensorOptions(device).dtype(at::kFloat));
54	tensors.push_back(x);	54	tensors.push_back(x);
55	}	55	}
56	-	56	+
57	std::cout << "在NPU设备 " << rank << " 上创建了 " << ntensors << " 个张量" << std::endl;	57	std::cout << "在NPU设备 " << rank << " 上创建了 " << ntensors << " 个张量" << std::endl;
58	-	58	+
59	// 提交所有allreduce操作	59	// 提交所有allreduce操作
60	std::vector<c10::intrusive_ptr<Work>> pending;	60	std::vector<c10::intrusive_ptr<Work>> pending;
61	for (const auto i : c10::irange(ntensors)) {	61	for (const auto i : c10::irange(ntensors)) {
62	std::vector<at::Tensor> tmp = {tensors[i]};	62	std::vector<at::Tensor> tmp = {tensors[i]};
63	pending.push_back(pg->allreduce(tmp));	63	pending.push_back(pg->allreduce(tmp));
64	}	64	}
65	-	65	+
66	std::cout << "已提交 " << ntensors << " 个allreduce操作" << std::endl;	66	std::cout << "已提交 " << ntensors << " 个allreduce操作" << std::endl;
67	-	67	+
68	// 等待所有操作完成	68	// 等待所有操作完成
69	for (auto& work : pending) {	69	for (auto& work : pending) {
70	work->wait();	70	work->wait();
71	}	71	}
72	-	72	+
73	std::cout << "所有操作已完成!" << std::endl;	73	std::cout << "所有操作已完成!" << std::endl;
74	-	74	+
75	// 验证结果 - 打印前3个张量的第一个元素	75	// 验证结果 - 打印前3个张量的第一个元素
76	for (const auto i : c10::irange(std::min(ntensors, 3))) {	76	for (const auto i : c10::irange(std::min(ntensors, 3))) {
77	auto cpu_tensor = tensors[i].to(at::kCPU);	77	auto cpu_tensor = tensors[i].to(at::kCPU);
78	std::cout << "张量 " << i << " 第一个元素: " << cpu_tensor.data_ptr<float>()[0] << std::endl;	78	std::cout << "张量 " << i << " 第一个元素: " << cpu_tensor.data_ptr<float>()[0] << std::endl;
79	}	79	}
80	-	80	+
81	std::cout << "HCCL allreduce示例运行成功!" << std::endl;	81	std::cout << "HCCL allreduce示例运行成功!" << std::endl;
82	-	82	+
83	// 使用NPU设备结束需进行反初始化	83	// 使用NPU设备结束需进行反初始化
84	torch_npu::finalize_npu();	84	torch_npu::finalize_npu();
85	-	85	+
86	return 0;	86	return 0;
87	}	87	}

Msetup.py+7-7

	@@ -355,12 +355,12 @@ class CPPLibBuild(build_clib, object):
355	if DISABLE_RPC == 'FALSE':	355	if DISABLE_RPC == 'FALSE':
356	if check_tensorpipe_valid(BASE_DIR):	356	if check_tensorpipe_valid(BASE_DIR):
357	cmake_args.append('-DBUILD_TENSORPIPE=on')	357	cmake_args.append('-DBUILD_TENSORPIPE=on')
358	-	358	+
359	if ENABLE_LTO == "TRUE":	359	if ENABLE_LTO == "TRUE":
360	cmake_args.append('-DENABLE_LTO=on')	360	cmake_args.append('-DENABLE_LTO=on')
361	if PGO_MODE != 0:	361	if PGO_MODE != 0:
362	cmake_args.append('-DPGO_MODE=' + str(PGO_MODE))	362	cmake_args.append('-DPGO_MODE=' + str(PGO_MODE))
363	-	363	+
364	if USE_CXX11_ABI:	364	if USE_CXX11_ABI:
365	cmake_args.append('-DGLIBCXX_USE_CXX11_ABI=1')	365	cmake_args.append('-DGLIBCXX_USE_CXX11_ABI=1')
366		366
	@@ -426,12 +426,12 @@ def add_ops_python_files(ret_list):
426	opplugin_path = os.path.join(BASE_DIR, 'third_party/op-plugin/op_plugin/python')	426	opplugin_path = os.path.join(BASE_DIR, 'third_party/op-plugin/op_plugin/python')
427		427
428	if os.path.exists(opplugin_path):	428	if os.path.exists(opplugin_path):
429	- ops_python_files = glob.glob(os.path.join(opplugin_path, '*/.py'), recursive=True)	429	+ ops_python_files = glob.glob(os.path.join(opplugin_path, '*/.py'), recursive=True)
430	for src in ops_python_files:	430	for src in ops_python_files:
431	dst = os.path.join(	431	dst = os.path.join(
432	os.path.join(BASE_DIR, "build/packages/torch_npu/op_plugin"),	432	os.path.join(BASE_DIR, "build/packages/torch_npu/op_plugin"),
433	os.path.relpath(src, opplugin_path))	433	os.path.relpath(src, opplugin_path))
434	- os.makedirs(os.path.dirname(dst), exist_ok=True)	434	+ os.makedirs(os.path.dirname(dst), exist_ok=True)
435	ret_list.append((src, dst))	435	ret_list.append((src, dst))
436	return	436	return
437		437
	@@ -544,11 +544,11 @@ def get_src_py_and_dst():
544	# 按原目录结构复制到目标路径	544	# 按原目录结构复制到目标路径
545	for src in codegen_files:	545	for src in codegen_files:
546	# 仅过滤指定目录下的根级__init__.py	546	# 仅过滤指定目录下的根级__init__.py
547	- if (exclude_root_init is not None and	547	+ if (exclude_root_init is not None and
548	- os.path.basename(src) == '__init__.py' and	548	+ os.path.basename(src) == '__init__.py' and
549	os.path.dirname(src) == exclude_root_init):	549	os.path.dirname(src) == exclude_root_init):
550	continue # 跳过op-plugin/codegen根目录的__init__.py	550	continue # 跳过op-plugin/codegen根目录的__init__.py
551	-	551	+
552	# 计算目标路径（保留原目录层级）	552	# 计算目标路径（保留原目录层级）
553	dst = os.path.join(	553	dst = os.path.join(
554	codegen_dst_dir,	554	codegen_dst_dir,

Mtest/_inductor/test_cat.py+1-1

	@@ -44,7 +44,7 @@ class TestCat(TestUtils):
44	output_tensor = torch.cat(slices, self.dim)	44	output_tensor = torch.cat(slices, self.dim)
45		45
46	return output_tensor	46	return output_tensor
47	-	47	+
48	@parametrize('shape', [(128, 50, 128)])	48	@parametrize('shape', [(128, 50, 128)])
49	@parametrize('dim', [2])	49	@parametrize('dim', [2])
50	@parametrize('dtype', ['float32', 'bfloat16'])	50	@parametrize('dtype', ['float32', 'bfloat16'])

Mtest/_inductor/test_check_accuracy.py+3-3

	@@ -15,7 +15,7 @@ class TestCheckAccuracy(TestUtils):
15	def test_check_accuracy_1(self):	15	def test_check_accuracy_1(self):
16	count_data_dump = 0	16	count_data_dump = 0
17	count_check_accuracy = 0	17	count_check_accuracy = 0
18	-	18	+
19	def run(x, y):	19	def run(x, y):
20	return F.relu(x) - y	20	return F.relu(x) - y
21		21
	@@ -50,7 +50,7 @@ class TestCheckAccuracy(TestUtils):
50	patch.object(torch_npu._inductor.npu_triton_heuristics, "check_accuracy_triton", wrap_check_accuracy):	50	patch.object(torch_npu._inductor.npu_triton_heuristics, "check_accuracy_triton", wrap_check_accuracy):
51	self.assertTrue(torch_npu._inductor.config.dump_fx_graph)	51	self.assertTrue(torch_npu._inductor.config.dump_fx_graph)
52	self.assertTrue(torch_npu._inductor.config.check_accuracy)	52	self.assertTrue(torch_npu._inductor.config.check_accuracy)
53	-	53	+
54	# Try run custom path and make sure no data_dump and check_accuracy is invoked.	54	# Try run custom path and make sure no data_dump and check_accuracy is invoked.
55	torch_npu._inductor.config.dump_fx_graph = False	55	torch_npu._inductor.config.dump_fx_graph = False
56	torch_npu._inductor.config.check_accuracy = False	56	torch_npu._inductor.config.check_accuracy = False
	@@ -65,7 +65,7 @@ class TestCheckAccuracy(TestUtils):
65	self.assertEqual(count_data_dump, 1)	65	self.assertEqual(count_data_dump, 1)
66	self.assertEqual(count_check_accuracy, 1)	66	self.assertEqual(count_check_accuracy, 1)
67	self.assertEqual(z, g)	67	self.assertEqual(z, g)
68	-	68	+
69		69
70	if __name__ == "__main__":	70	if __name__ == "__main__":
71	run_tests()	71	run_tests()

Mtest/_inductor/test_clamp.py+1-1

	@@ -68,7 +68,7 @@ class TestClamp(TestUtils):
68	self.assertEqual(std_result, inductor_result)	68	self.assertEqual(std_result, inductor_result)
69		69
70	@parametrize('shape', TestUtils._pointwise_demo_shapes)	70	@parametrize('shape', TestUtils._pointwise_demo_shapes)
71	- @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32', 'int64'])	71	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32', 'int64'])
72	def test_pointwise_cases_min_only(self, shape, dtype):	72	def test_pointwise_cases_min_only(self, shape, dtype):
73	min_numel = 0	73	min_numel = 0
74		74

Mtest/_inductor/test_debug_msg.py+1-1

	@@ -14,7 +14,7 @@ os.environ["INDUCTOR_ASCEND_DUMP_FX_GRAPH"] = "1"
14	os.environ["TORCH_COMPILE_DEBUG"] = "1"	14	os.environ["TORCH_COMPILE_DEBUG"] = "1"
15		15
16		16
17	-class TestDebugMsg(TestUtils):	17	+class TestDebugMsg(TestUtils):
18	@parametrize('shape_x', [(32, 512, 64)])	18	@parametrize('shape_x', [(32, 512, 64)])
19	@parametrize('shape_y', [(32, 1, 64)])	19	@parametrize('shape_y', [(32, 1, 64)])
20	@parametrize('dtype', ['float32'])	20	@parametrize('dtype', ['float32'])

Mtest/_inductor/test_dropout_with_checkpoint_recompute.py+8-8

	@@ -6,14 +6,14 @@ from torch.testing._internal.common_utils import (
6	)	6	)
7	from testutils import TestUtils	7	from testutils import TestUtils
8	import torch_npu	8	import torch_npu
9	-	9	+
10	class TestDropoutWithCheckpointRecompute(TestUtils):	10	class TestDropoutWithCheckpointRecompute(TestUtils):
11	def test_dropout_with_checkpoint_recompute(self):	11	def test_dropout_with_checkpoint_recompute(self):
12	device = "npu"	12	device = "npu"
13	-	13	+
14	def gn(x):	14	def gn(x):
15	return torch.sigmoid(torch.dropout(torch.sigmoid(x), p=0.5, train=True))	15	return torch.sigmoid(torch.dropout(torch.sigmoid(x), p=0.5, train=True))
16	-	16	+
17	def fn(x):	17	def fn(x):
18	return checkpoint(	18	return checkpoint(
19	gn,	19	gn,
	@@ -21,22 +21,22 @@ class TestDropoutWithCheckpointRecompute(TestUtils):
21	use_reentrant=False,	21	use_reentrant=False,
22	preserve_rng_state=True,	22	preserve_rng_state=True,
23	)	23	)
24	-	24	+
25	x = torch.randn(4, 4, requires_grad=True, device=device)	25	x = torch.randn(4, 4, requires_grad=True, device=device)
26	-	26	+
27	torch.manual_seed(42)	27	torch.manual_seed(42)
28	eager_out = fn(x)	28	eager_out = fn(x)
29	eager_out.sum().backward()	29	eager_out.sum().backward()
30	eager_grad = x.grad.clone()	30	eager_grad = x.grad.clone()
31	-	31	+
32	x.grad = None	32	x.grad = None
33	-	33	+
34	torch.manual_seed(42)	34	torch.manual_seed(42)
35	compiled_fn = torch.compile(fn, backend="inductor")	35	compiled_fn = torch.compile(fn, backend="inductor")
36	compiled_out = compiled_fn(x)	36	compiled_out = compiled_fn(x)
37	compiled_out.sum().backward()	37	compiled_out.sum().backward()
38	compiled_grad = x.grad.clone()	38	compiled_grad = x.grad.clone()
39	-	39	+
40	self.assertEqual(eager_out, compiled_out)	40	self.assertEqual(eager_out, compiled_out)
41	self.assertEqual(eager_grad, compiled_grad)	41	self.assertEqual(eager_grad, compiled_grad)
42		42

Mtest/_inductor/test_embedding.py+1-1

	@@ -13,7 +13,7 @@ class TestEmbeddingDense(TestUtils):
13	# UT skip, reason: precision fail	13	# UT skip, reason: precision fail
14	# Added to pytorch-disable-tests.json	14	# Added to pytorch-disable-tests.json
15	def test_pointwise_cases(self):	15	def test_pointwise_cases(self):
16	-	16	+
17	arg0 = torch.tensor([[14, 1, 2, 10, 0, 10, 0],	17	arg0 = torch.tensor([[14, 1, 2, 10, 0, 10, 0],
18	[9, 13, 13, 4, 7, 15, 14],	18	[9, 13, 13, 4, 7, 15, 14],
19	[8, 0, 3, 15, 4, 2, 6],	19	[8, 0, 3, 15, 4, 2, 6],

Mtest/_inductor/test_exceptions.py+5-5

	@@ -29,13 +29,13 @@ import torch_npu
29	size_hints=[16384, 32], tile_hint=TileHint.DEFAULT,	29	size_hints=[16384, 32], tile_hint=TileHint.DEFAULT,
30	filename=__file__,	30	filename=__file__,
31	triton_meta={'signature': {'in_ptr0': 'fp16', 'in_ptr1': 'fp16', 'out_ptr0': '*fp16', 'y0_numel': 'i32', 'x1_numel': 'i32'},	31	triton_meta={'signature': {'in_ptr0': 'fp16', 'in_ptr1': 'fp16', 'out_ptr0': '*fp16', 'y0_numel': 'i32', 'x1_numel': 'i32'},
32	- 'device': NPUDeviceProperties(type='npu', index=0, multi_processor_count=40, cc='Ascend910B3',	32	+ 'device': NPUDeviceProperties(type='npu', index=0, multi_processor_count=40, cc='Ascend910B3',
33	major=None, regs_per_multiprocessor=None, max_threads_per_multi_processor=None, warp_size=32),	33	major=None, regs_per_multiprocessor=None, max_threads_per_multi_processor=None, warp_size=32),
34	'constants': {}, 'mix_mode': 'aiv'},	34	'constants': {}, 'mix_mode': 'aiv'},
35	- inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_unk_fused_add_0', 'mutated_arg_names': [],	35	+ inductor_meta={'autotune_hints': set(), 'kernel_name': 'triton_unk_fused_add_0', 'mutated_arg_names': [],
36	- 'backend_hash': 'bc71dba4086164e7ac2b0779fa861dbf7467f0265d4a57b8f48cf6dda02b150f', 'split_axis': [0],	36	+ 'backend_hash': 'bc71dba4086164e7ac2b0779fa861dbf7467f0265d4a57b8f48cf6dda02b150f', 'split_axis': [0],
37	- 'tiling_axis': [0, 1], 'axis_names': ['y0', 'x1'], 'low_dims': {1}, 'numof_reduction_axis': 0,	37	+ 'tiling_axis': [0, 1], 'axis_names': ['y0', 'x1'], 'low_dims': {1}, 'numof_reduction_axis': 0,
38	- 'split_axis_dtype': torch.float16, 'dual_reduction': False, 'traced_graph_hash': 'TRACED_GRAPH_HASH',	38	+ 'split_axis_dtype': torch.float16, 'dual_reduction': False, 'traced_graph_hash': 'TRACED_GRAPH_HASH',
39	'traced_graph_dir': 'TRACED_GRAPH_DIR'},	39	'traced_graph_dir': 'TRACED_GRAPH_DIR'},
40	min_elem_per_thread=0	40	min_elem_per_thread=0
41	)	41	)

Mtest/_inductor/test_force_fallback.py+3-3

	@@ -18,10 +18,10 @@ class TestForceFallback(TestUtils):
18	def opoverload_call(self, /, args, *kwargs):	18	def opoverload_call(self, /, args, *kwargs):
19	op_list.append(str(self))	19	op_list.append(str(self))
20	return self._op(args, *kwargs)	20	return self._op(args, *kwargs)
21	-	21	+
22	def run(x, y):	22	def run(x, y):
23	return F.relu(x) + y	23	return F.relu(x) + y
24	-	24	+
25	x = torch.randn(10).npu()	25	x = torch.randn(10).npu()
26	y = torch.randn(10).npu()	26	y = torch.randn(10).npu()
27	g = run(x, y)	27	g = run(x, y)
	@@ -49,7 +49,7 @@ class TestForceFallback(TestUtils):
49	self.assertTrue("aten.relu.default" in op_list)	49	self.assertTrue("aten.relu.default" in op_list)
50	self.assertTrue("aten.add.Tensor" in op_list)	50	self.assertTrue("aten.add.Tensor" in op_list)
51	self.assertEqual(z, g)	51	self.assertEqual(z, g)
52	-	52	+
53	# reset	53	# reset
54	torch_npu._inductor.config.force_fallback_kernel_id = []	54	torch_npu._inductor.config.force_fallback_kernel_id = []
55		55

Mtest/_inductor/test_inductor_fallback_list.py+6-6

Mtest/_inductor/test_lazy_register.py+1-1

	@@ -6,7 +6,7 @@ import torch_npu
6		6
7	class TestLazyRegister(TestUtils):	7	class TestLazyRegister(TestUtils):
8		8
9	-	9	+
10	def test_disale_register_inductor_npu(self):	10	def test_disale_register_inductor_npu(self):
11	torch_npu.utils._dynamo.disable_register_inductor_npu()	11	torch_npu.utils._dynamo.disable_register_inductor_npu()
12		12

Mtest/_inductor/test_mlir_enable.py+1-1

	@@ -6,7 +6,7 @@ import torch_npu
6		6
7		7
8	class TestAdd(TestUtils):	8	class TestAdd(TestUtils):
9	-	9	+
10		10
11	def op_calc(self, first_element, second_element):	11	def op_calc(self, first_element, second_element):
12	result = first_element + second_element	12	result = first_element + second_element

Mtest/_inductor/test_npu_kernel_features.py+5-5

	@@ -10,7 +10,7 @@ class TestNumeList(TestCase):
10	def test_numels(self):	10	def test_numels(self):
11	numel_list = NumelList([2, 3, 4])	11	numel_list = NumelList([2, 3, 4])
12	self.assertEqual(numel_list.numels(), 24)	12	self.assertEqual(numel_list.numels(), 24)
13	-	13	+
14	def test_equality(self):	14	def test_equality(self):
15	numel_list1 = NumelList([2, 3, 4])	15	numel_list1 = NumelList([2, 3, 4])
16	numel_list2 = NumelList([2, 3, 4])	16	numel_list2 = NumelList([2, 3, 4])
	@@ -45,11 +45,11 @@ class TestNumeList(TestCase):
45	numel_list1 = NumelList([2, 3, 5])	45	numel_list1 = NumelList([2, 3, 5])
46	numel_list2 = NumelList([2, 3, 4])	46	numel_list2 = NumelList([2, 3, 4])
47	self.assertTrue(numel_list1 >= numel_list2)	47	self.assertTrue(numel_list1 >= numel_list2)
48	-	48	+
49	def test_modulo(self):	49	def test_modulo(self):
50	numel_list = NumelList([2, 3, 4])	50	numel_list = NumelList([2, 3, 4])
51	self.assertEqual(numel_list % 5, 4)	51	self.assertEqual(numel_list % 5, 4)
52	-	52	+
53	def test_division(self):	53	def test_division(self):
54	numel_list = NumelList([2, 3, 4])	54	numel_list = NumelList([2, 3, 4])
55	self.assertEqual(numel_list / 2, 12.0)	55	self.assertEqual(numel_list / 2, 12.0)
	@@ -63,8 +63,8 @@ class TestNumeList(TestCase):
63	def test_addition(self):	63	def test_addition(self):
64	numel_list = NumelList([2, 3, 4])	64	numel_list = NumelList([2, 3, 4])
65	self.assertEqual(numel_list + 2, 26)	65	self.assertEqual(numel_list + 2, 26)
66	- self.assertEqual(2 + numel_list, 26)	66	+ self.assertEqual(2 + numel_list, 26)
67	-	67	+
68	def test_hash(self):	68	def test_hash(self):
69	# 测试相同内容的hash值相同	69	# 测试相同内容的hash值相同
70	numel_list1 = NumelList([2, 3, 4])	70	numel_list1 = NumelList([2, 3, 4])

Mtest/_inductor/test_opensora_graph1.py+1-1

	@@ -196,7 +196,7 @@ class TestModel(TestUtils):
196	mul_1: "i64[]" = torch.ops.aten.mul.Tensor(primals_3, 2)	196	mul_1: "i64[]" = torch.ops.aten.mul.Tensor(primals_3, 2)
197	mul_2: "i64[]" = torch.ops.aten.mul.Tensor(primals_4, 2)	197	mul_2: "i64[]" = torch.ops.aten.mul.Tensor(primals_4, 2)
198	return [permute, mul, mul_1, mul_2]	198	return [permute, mul, mul_1, mul_2]
199	-	199	+
200	primals_1 = torch.randn((1, 8, 30, 40, 1, 2, 2, 8), device=device_npu, dtype=torch.float32)	200	primals_1 = torch.randn((1, 8, 30, 40, 1, 2, 2, 8), device=device_npu, dtype=torch.float32)
201	primals_2 = torch.tensor((1), device=device_npu, dtype=torch.int64)	201	primals_2 = torch.tensor((1), device=device_npu, dtype=torch.int64)
202	primals_3 = torch.tensor((1), device=device_npu, dtype=torch.int64)	202	primals_3 = torch.tensor((1), device=device_npu, dtype=torch.int64)

Mtest/_inductor/test_rng_prims.py+1-1

	@@ -15,7 +15,7 @@ class TestRNGPrims(TestUtils):
15	def test_default(self):	15	def test_default(self):
16	register_run_and_save_rng_state_op()	16	register_run_and_save_rng_state_op()
17	x = torch.randn(4, 4).to("npu")	17	x = torch.randn(4, 4).to("npu")
18	- args = (x,)	18	+ args = (x,)
19	kwargs = {}	19	kwargs = {}
20	expected_rng_state = torch_npu.npu.get_rng_state()	20	expected_rng_state = torch_npu.npu.get_rng_state()
21	rng_state, out = run_and_save_rng_state(lambda x: x, args, *kwargs)	21	rng_state, out = run_and_save_rng_state(lambda x: x, args, *kwargs)

Mtest/_inductor/test_run_and_save_rng_state.py+1-1

	@@ -40,7 +40,7 @@ class TestRunAndSaveRngState(TestUtils):
40	self.op_calc(like, device, dtype)	40	self.op_calc(like, device, dtype)
41		41
42	self.assertEqual(res1_eager, res2_eager)	42	self.assertEqual(res1_eager, res2_eager)
43	- self.assertTrue(torch.equal(rng_state1_eager, rng_state2_eager))	43	+ self.assertTrue(torch.equal(rng_state1_eager, rng_state2_eager))
44		44
45	instantiate_parametrized_tests(TestRunAndSaveRngState)	45	instantiate_parametrized_tests(TestRunAndSaveRngState)
46		46

Mtest/_inductor/test_where.py+1-1

	@@ -11,7 +11,7 @@ class TestWhere(TestUtils):
11		11
12	@unittest.skip("it takes too long, not supported yet")	12	@unittest.skip("it takes too long, not supported yet")
13	@parametrize('shape', TestUtils._pointwise_demo_shapes)	13	@parametrize('shape', TestUtils._pointwise_demo_shapes)
14	- @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32'])	14	+ @parametrize('dtype', ['float16', 'float32', 'bfloat16', 'int32'])
15	def test_pointwise_cases(self, shape, dtype):	15	def test_pointwise_cases(self, shape, dtype):
16	first_element = self._generate_tensor(shape, dtype)	16	first_element = self._generate_tensor(shape, dtype)
17	second_element = self._generate_tensor(shape, dtype)	17	second_element = self._generate_tensor(shape, dtype)

Mtest/allocator/host/test_expandable_host_allocator.py+2-2

	@@ -45,7 +45,7 @@ class TestHostCachingAllocator(TestCase):
45	def setUpClass(cls):	45	def setUpClass(cls):
46	os.environ['PYTORCH_NPU_ALLOC_CONF'] = 'pin_memory_expandable_segments:True'	46	os.environ['PYTORCH_NPU_ALLOC_CONF'] = 'pin_memory_expandable_segments:True'
47		47
48	-	48	+
49	def test_allocate_with_block_cut(self):	49	def test_allocate_with_block_cut(self):
50	# 申请一个64M的tensor, 会预选绑定80M物理内存	50	# 申请一个64M的tensor, 会预选绑定80M物理内存
51	memory_64m = torch.ones([1024, 1024, 16]).pin_memory()	51	memory_64m = torch.ones([1024, 1024, 16]).pin_memory()
	@@ -130,7 +130,7 @@ class TestHostCachingAllocator(TestCase):
130	npu_output = npu_copy_op_exec(npu_input1, cpu_out2)	130	npu_output = npu_copy_op_exec(npu_input1, cpu_out2)
131	self.assertRtolEqual(cpu_output, npu_output)	131	self.assertRtolEqual(cpu_output, npu_output)
132		132
133	-	133	+
134	def test_event_free(self):	134	def test_event_free(self):
135	tensor = torch.ones([1024, 1024, 16]).npu()	135	tensor = torch.ones([1024, 1024, 16]).npu()
136	tensor_cpu = torch.ones([1024, 1024, 16]).pin_memory()	136	tensor_cpu = torch.ones([1024, 1024, 16]).pin_memory()

Mtest/allocator/test_hostallocator.py+10-10

	@@ -26,20 +26,20 @@ def _collect():
26		26
27		27
28	class TestHostCachingAllocatorBasic(TestCase):	28	class TestHostCachingAllocatorBasic(TestCase):
29	- def test_pin_memory_on_non_blocking_copy(self):	29	+ def test_pin_memory_on_non_blocking_copy(self):
30	t_acc = torch.randn(100).to(torch.accelerator.current_accelerator())	30	t_acc = torch.randn(100).to(torch.accelerator.current_accelerator())
31	t_host = t_acc.to("cpu", non_blocking=True)	31	t_host = t_acc.to("cpu", non_blocking=True)
32	torch.accelerator.synchronize()	32	torch.accelerator.synchronize()
33	self.assertTrue(t_host.is_pinned())	33	self.assertTrue(t_host.is_pinned())
34	self.assertEqual(t_acc.cpu(), t_host)	34	self.assertEqual(t_acc.cpu(), t_host)
35	-	35	+
36	def test_pin_memory_reuse(self):	36	def test_pin_memory_reuse(self):
37	t = torch.FloatTensor([1]).pin_memory()	37	t = torch.FloatTensor([1]).pin_memory()
38	ptr = t.data_ptr()	38	ptr = t.data_ptr()
39	del t	39	del t
40	t_new = torch.FloatTensor([1]).pin_memory()	40	t_new = torch.FloatTensor([1]).pin_memory()
41	self.assertEqual(t_new.data_ptr(), ptr)	41	self.assertEqual(t_new.data_ptr(), ptr)
42	-	42	+
43	def test_to_non_blocking(self):	43	def test_to_non_blocking(self):
44	stream = torch_npu.npu.current_stream()	44	stream = torch_npu.npu.current_stream()
45		45
	@@ -61,7 +61,7 @@ class TestHostCachingAllocatorBasic(TestCase):
61	device="npu" if dst == "cpu" else "cpu",	61	device="npu" if dst == "cpu" else "cpu",
62	pin_memory=True if dst == "npu" else False)	62	pin_memory=True if dst == "npu" else False)
63	_test_to_non_blocking(src, try_non_blocking, dst)	63	_test_to_non_blocking(src, try_non_blocking, dst)
64	-	64	+
65	def test_pin_memory_basic(self):	65	def test_pin_memory_basic(self):
66	a = torch.Tensor([1])	66	a = torch.Tensor([1])
67	b = a.pin_memory()	67	b = a.pin_memory()
	@@ -70,7 +70,7 @@ class TestHostCachingAllocatorBasic(TestCase):
70	self.assertTrue(a.data_ptr() != b.data_ptr())	70	self.assertTrue(a.data_ptr() != b.data_ptr())
71	self.assertTrue(b.data_ptr() != c.data_ptr())	71	self.assertTrue(b.data_ptr() != c.data_ptr())
72	self.assertTrue(b.data_ptr() == d.data_ptr())	72	self.assertTrue(b.data_ptr() == d.data_ptr())
73	-	73	+
74	def test_malloc_copykernel(self):	74	def test_malloc_copykernel(self):
75	a = torch.Tensor([1])	75	a = torch.Tensor([1])
76	b = torch.Tensor([1])	76	b = torch.Tensor([1])
	@@ -119,7 +119,7 @@ class TestHostCachingAllocatorBasic(TestCase):
119		119
120	torch.npu.synchronize()	120	torch.npu.synchronize()
121	self.assertTrue(not errs)	121	self.assertTrue(not errs)
122	-	122	+
123	def test_pin_memory_on_views_and_clones(self):	123	def test_pin_memory_on_views_and_clones(self):
124	base = torch.randn(1024, 1024)	124	base = torch.randn(1024, 1024)
125	view = base[:512, :].pin_memory()	125	view = base[:512, :].pin_memory()
	@@ -157,7 +157,7 @@ class CountingDataset(Dataset):
157		157
158	def __getitem__(self, i):	158	def __getitem__(self, i):
159	return i	159	return i
160	-	160	+
161	def __len__(self):	161	def __len__(self):
162	return self.n	162	return self.n
163		163
	@@ -165,7 +165,7 @@ class CountingDataset(Dataset):
165	class DictDataset(Dataset):	165	class DictDataset(Dataset):
166	def __len__(self):	166	def __len__(self):
167	return 4	167	return 4
168	-	168	+
169	def __getitem__(self, ndx):	169	def __getitem__(self, ndx):
170	return {	170	return {
171	'a_tensor': torch.empty(4, 2).fill_(ndx),	171	'a_tensor': torch.empty(4, 2).fill_(ndx),
	@@ -181,7 +181,7 @@ class StringDataset(Dataset):
181		181
182	def __len__(self):	182	def __len__(self):
183	return len(self.s)	183	return len(self.s)
184	-	184	+
185	def __getitem__(self, ndx):	185	def __getitem__(self, ndx):
186	return (self.s[ndx], ndx)	186	return (self.s[ndx], ndx)
187		187
	@@ -196,7 +196,7 @@ class SimpleCustomBatch:
196	self.inp = self.inp.pin_memory()	196	self.inp = self.inp.pin_memory()
197	self.tgt = self.tgt.pin_memory()	197	self.tgt = self.tgt.pin_memory()
198	return self	198	return self
199	-	199	+
200	def is_pinned(self):	200	def is_pinned(self):
201	return self.inp.is_pinned() and self.tgt.is_pinned()	201	return self.inp.is_pinned() and self.tgt.is_pinned()
202		202

Mtest/custom_ops/test_masked_softmax_with_rel_pos_bias.py+1-1

	@@ -10,7 +10,7 @@ from torch_npu.testing.common_utils import SupportedDevices
10	class TestMaskedSoftmaxWithRelPosBias(TestCase):	10	class TestMaskedSoftmaxWithRelPosBias(TestCase):
11		11
12	def supported_op_exec(self, x, relative_pos_bias, atten_mask):	12	def supported_op_exec(self, x, relative_pos_bias, atten_mask):
13	- # add + add + softmax	13	+ # add + add + softmax
14	y = torch.add(x, atten_mask)	14	y = torch.add(x, atten_mask)
15	y = torch.add(y, relative_pos_bias)	15	y = torch.add(y, relative_pos_bias)
16	softmax_out = torch.nn.functional.softmax(y, dim=-1)	16	softmax_out = torch.nn.functional.softmax(y, dim=-1)

Mtest/custom_ops/test_npu_anti_quant.py+6-6

	@@ -32,12 +32,12 @@ class TestAntiQuant(TestCase):
32	scale = torch.broadcast_to(scale, input_x.shape)	32	scale = torch.broadcast_to(scale, input_x.shape)
33	if offset is None:	33	if offset is None:
34	offset = torch.zeros_like(scale)	34	offset = torch.zeros_like(scale)
35	-	35	+
36	x = input_x.to(torch.float32)	36	x = input_x.to(torch.float32)
37	-	37	+
38	offset_temp = x + offset	38	offset_temp = x + offset
39	output = offset_temp * scale	39	output = offset_temp * scale
40	-	40	+
41	output = output.to(dst_dtype)	41	output = output.to(dst_dtype)
42	return output.cpu().detach()	42	return output.cpu().detach()
43		43
	@@ -57,12 +57,12 @@ class TestAntiQuant(TestCase):
57	[[np.int32, -1, [10, 25]], [np.float32, -1, [200]], [np.float32, -1, [200]], torch.float16, None],	57	[[np.int32, -1, [10, 25]], [np.float32, -1, [200]], [np.float32, -1, [200]], torch.float16, None],
58	[[np.int32, -1, [10, 25]], [np.float32, -1, [200]], [np.float32, -1, [200]], torch.bfloat16, None],	58	[[np.int32, -1, [10, 25]], [np.float32, -1, [200]], [np.float32, -1, [200]], torch.bfloat16, None],
59	]	59	]
60	-	60	+
61	for item in shape_format:	61	for item in shape_format:
62	cpu_input_x, npu_input_x = create_common_tensor(item[0], -127, 127)	62	cpu_input_x, npu_input_x = create_common_tensor(item[0], -127, 127)
63	cpu_scale, npu_scale = create_common_tensor(item[1], -100, 100)	63	cpu_scale, npu_scale = create_common_tensor(item[1], -100, 100)
64	cpu_offset, npu_offset = (None, None) if item[2] is None else create_common_tensor(item[2], -100, 100)	64	cpu_offset, npu_offset = (None, None) if item[2] is None else create_common_tensor(item[2], -100, 100)
65	-	65	+
66	npu_output = self.npu_op_exec(npu_input_x, npu_scale, npu_offset, *item[3:])	66	npu_output = self.npu_op_exec(npu_input_x, npu_scale, npu_offset, *item[3:])
67	custom_output = self.custom_op_exec(cpu_input_x, cpu_scale, cpu_offset, *item[3:])	67	custom_output = self.custom_op_exec(cpu_input_x, cpu_scale, cpu_offset, *item[3:])
68		68
	@@ -77,7 +77,7 @@ class TestAntiQuant(TestCase):
77	shape_format = [	77	shape_format = [
78	[[np.int8, -1, [10, 100]], [np.float32, -1, [100]], [np.float32, -1, [100]], torch.float16, None],	78	[[np.int8, -1, [10, 100]], [np.float32, -1, [100]], [np.float32, -1, [100]], torch.float16, None],
79	]	79	]
80	-	80	+
81	for item in shape_format:	81	for item in shape_format:
82	_, npu_input_x = create_common_tensor(item[0], -127, 127)	82	_, npu_input_x = create_common_tensor(item[0], -127, 127)
83	_, npu_scale = create_common_tensor(item[1], -100, 100)	83	_, npu_scale = create_common_tensor(item[1], -100, 100)

Mtest/custom_ops/test_npu_dtype_cast.py+1-1

	@@ -41,7 +41,7 @@ class TestDtypeCast(TestCase):
41	y = torch_npu.npu_dtype_cast(x, torch.complex128)	41	y = torch_npu.npu_dtype_cast(x, torch.complex128)
42	grad_fn = str(y.grad_fn)	42	grad_fn = str(y.grad_fn)
43	self.assertTrue("NpuDtypeCastBackward" in grad_fn)	43	self.assertTrue("NpuDtypeCastBackward" in grad_fn)
44	-	44	+
45	with self.assertRaisesRegex(RuntimeError, r'grad can be implicitly created'):	45	with self.assertRaisesRegex(RuntimeError, r'grad can be implicitly created'):
46	y.sum().backward()	46	y.sum().backward()
47		47

Mtest/custom_ops/test_npu_stride_copy.py+1-1

	@@ -17,7 +17,7 @@ class TestNpuStrideCopy(TestCase):
17	output = torch_npu.npu_stride_copy(input1, size, stride, storage_offset)	17	output = torch_npu.npu_stride_copy(input1, size, stride, storage_offset)
18	output = output.cpu().numpy()	18	output = output.cpu().numpy()
19	return output	19	return output
20	-	20	+
21	@unittest.skip("skip test_npu_stride_copy now")	21	@unittest.skip("skip test_npu_stride_copy now")
22	def test_npu_stride_copy(self):	22	def test_npu_stride_copy(self):
23	shape_format = [	23	shape_format = [

Mtest/custom_ops/test_resize_.py+1-1

	@@ -21,7 +21,7 @@ class TestResize(TestCase):
21	out_tensor_npu = torch.masked_select(input_data_npu, mask_npu, out=out_tensor_npu)	21	out_tensor_npu = torch.masked_select(input_data_npu, mask_npu, out=out_tensor_npu)
22	out_tensor = torch.masked_select(input_data, mask, out=out_tensor)	22	out_tensor = torch.masked_select(input_data, mask, out=out_tensor)
23	self.assertRtolEqual(out_tensor_npu, out_tensor)	23	self.assertRtolEqual(out_tensor_npu, out_tensor)
24	-	24	+
25	def test_resize_ncdhw(self):	25	def test_resize_ncdhw(self):
26	out_tensor = torch.empty((1, 1, 1, 1, 1), dtype=torch.float16).npu()	26	out_tensor = torch.empty((1, 1, 1, 1, 1), dtype=torch.float16).npu()
27	shape = [25]	27	shape = [25]

Mtest/distributed/elastic/agent/server/test_healthcheckserver_api.py+1-1

	@@ -2,7 +2,7 @@
2	Add validation cases for torch.distributed.elastic.agent.server.health_check_server APIs.	2	Add validation cases for torch.distributed.elastic.agent.server.health_check_server APIs.
3		3
4	1. PyTorch community tests do not cover HealthCheckServer APIs, so this file is added.	4	1. PyTorch community tests do not cover HealthCheckServer APIs, so this file is added.
5	-2. This file validates :	5	+2. This file validates :
6	torch.distributed.elastic.agent.server.health_check_server.HealthCheckServer	6	torch.distributed.elastic.agent.server.health_check_server.HealthCheckServer
7	torch.distributed.elastic.agent.server.health_check_server.HealthCheckServer.start	7	torch.distributed.elastic.agent.server.health_check_server.HealthCheckServer.start
8	torch.distributed.elastic.agent.server.health_check_server.HealthCheckServer.stop	8	torch.distributed.elastic.agent.server.health_check_server.HealthCheckServer.stop

Mtest/distributed/tensor/test_math_ops.py+20-20

	@@ -325,7 +325,7 @@ class TestConv2d(NPUDTensorTestBase):
325	output_tensor = torch_npu.npu_conv2d(input_tensor, weight_tensor, bias, stride, padding, dilation, groups)	325	output_tensor = torch_npu.npu_conv2d(input_tensor, weight_tensor, bias, stride, padding, dilation, groups)
326	output_dtensor = torch_npu.npu_conv2d(input_dtensor, weight_dtensor, bias, stride, padding, dilation, groups)	326	output_dtensor = torch_npu.npu_conv2d(input_dtensor, weight_dtensor, bias, stride, padding, dilation, groups)
327	self.assertEqual(output_dtensor.full_tensor(), output_tensor)	327	self.assertEqual(output_dtensor.full_tensor(), output_tensor)
328	-	328	+
329	@SupportedDevices(['Ascend910B'])	329	@SupportedDevices(['Ascend910B'])
330	@skipIfUnsupportMultiNPU(2)	330	@skipIfUnsupportMultiNPU(2)
331	@with_comms	331	@with_comms
	@@ -431,7 +431,7 @@ class TestConv2d(NPUDTensorTestBase):
431	input_dgrad, weight_dgrad, bias_dgrad = torch_npu.npu_conv2d_backward(input_dtensor, grad_output_dtensor, weight_dtensor, stride, padding, dilation, groups, output_mask)	431	input_dgrad, weight_dgrad, bias_dgrad = torch_npu.npu_conv2d_backward(input_dtensor, grad_output_dtensor, weight_dtensor, stride, padding, dilation, groups, output_mask)
432	self.assertEqual(input_dgrad.full_tensor(), input_grad)	432	self.assertEqual(input_dgrad.full_tensor(), input_grad)
433	self.assertEqual(weight_dgrad.full_tensor(), weight_grad)	433	self.assertEqual(weight_dgrad.full_tensor(), weight_grad)
434	-	434	+
435	@SupportedDevices(['Ascend910B'])	435	@SupportedDevices(['Ascend910B'])
436	@skipIfUnsupportMultiNPU(2)	436	@skipIfUnsupportMultiNPU(2)
437	@with_comms	437	@with_comms
	@@ -461,7 +461,7 @@ class TestConv2d(NPUDTensorTestBase):
461	self.assertEqual(input_dgrad.full_tensor(), input_grad)	461	self.assertEqual(input_dgrad.full_tensor(), input_grad)
462	self.assertEqual(weight_dgrad.full_tensor(), weight_grad)	462	self.assertEqual(weight_dgrad.full_tensor(), weight_grad)
463	self.assertEqual(bias_dgrad.full_tensor(), bias_grad)	463	self.assertEqual(bias_dgrad.full_tensor(), bias_grad)
464	-	464	+
465	@SupportedDevices(['Ascend910B'])	465	@SupportedDevices(['Ascend910B'])
466	@skipIfUnsupportMultiNPU(2)	466	@skipIfUnsupportMultiNPU(2)
467	@with_comms	467	@with_comms
	@@ -490,7 +490,7 @@ class TestConv2d(NPUDTensorTestBase):
490	input_dgrad, weight_dgrad, bias_dgrad = torch_npu.npu_conv2d_backward(input_dtensor, grad_output_dtensor, weight_dtensor, stride, padding, dilation, groups, output_mask)	490	input_dgrad, weight_dgrad, bias_dgrad = torch_npu.npu_conv2d_backward(input_dtensor, grad_output_dtensor, weight_dtensor, stride, padding, dilation, groups, output_mask)
491	self.assertEqual(input_dgrad.full_tensor(), input_grad)	491	self.assertEqual(input_dgrad.full_tensor(), input_grad)
492	self.assertEqual(weight_dgrad.full_tensor(), weight_grad)	492	self.assertEqual(weight_dgrad.full_tensor(), weight_grad)
493	-	493	+
494	@SupportedDevices(['Ascend910B'])	494	@SupportedDevices(['Ascend910B'])
495	@skipIfUnsupportMultiNPU(2)	495	@skipIfUnsupportMultiNPU(2)
496	@with_comms	496	@with_comms
	@@ -572,7 +572,7 @@ class TestGroupedMatmulAdd(NPUDTensorTestBase):
572	torch_npu.npu_grouped_matmul_add_(y, x, weight, group_list, transpose_x=transpose_x, transpose_weight=transpose_weight, group_type=group_type)	572	torch_npu.npu_grouped_matmul_add_(y, x, weight, group_list, transpose_x=transpose_x, transpose_weight=transpose_weight, group_type=group_type)
573	torch_npu.npu_grouped_matmul_add_(y_dtensor, x_dtensor, weight_dtensor, group_list_dtensor, transpose_x=transpose_x, transpose_weight=transpose_weight, group_type=group_type)	573	torch_npu.npu_grouped_matmul_add_(y_dtensor, x_dtensor, weight_dtensor, group_list_dtensor, transpose_x=transpose_x, transpose_weight=transpose_weight, group_type=group_type)
574	self.assertEqual(y_dtensor.full_tensor(), y)	574	self.assertEqual(y_dtensor.full_tensor(), y)
575	-	575	+
576	@SupportedDevices(['Ascend910B'])	576	@SupportedDevices(['Ascend910B'])
577	@skipIfUnsupportMultiNPU(2)	577	@skipIfUnsupportMultiNPU(2)
578	@with_comms	578	@with_comms
	@@ -614,7 +614,7 @@ class TestCrossEntropyLoss(NPUDTensorTestBase):
614	return input_tuple	614	return input_tuple
615	else:	615	else:
616	input_tuple = (x, target, input_dtensor, target_dtensor, mesh)	616	input_tuple = (x, target, input_dtensor, target_dtensor, mesh)
617	-	617	+
618	return input_tuple	618	return input_tuple
619		619
620		620
	@@ -630,7 +630,7 @@ class TestCrossEntropyLoss(NPUDTensorTestBase):
630	self.assertEqual(loss_dtensor.full_tensor(), loss)	630	self.assertEqual(loss_dtensor.full_tensor(), loss)
631	self.assertEqual(log_prob_dtensor.full_tensor(), log_prob)	631	self.assertEqual(log_prob_dtensor.full_tensor(), log_prob)
632		632
633	-	633	+
634	@SupportedDevices(['Ascend910B'])	634	@SupportedDevices(['Ascend910B'])
635	@skipIfUnsupportMultiNPU(2)	635	@skipIfUnsupportMultiNPU(2)
636	@with_comms	636	@with_comms
	@@ -643,7 +643,7 @@ class TestCrossEntropyLoss(NPUDTensorTestBase):
643	self.assertEqual(loss_dtensor.full_tensor(), loss)	643	self.assertEqual(loss_dtensor.full_tensor(), loss)
644	self.assertEqual(log_prob_dtensor.full_tensor(), log_prob)	644	self.assertEqual(log_prob_dtensor.full_tensor(), log_prob)
645		645
646	-	646	+
647	@SupportedDevices(['Ascend910B'])	647	@SupportedDevices(['Ascend910B'])
648	@skipIfUnsupportMultiNPU(2)	648	@skipIfUnsupportMultiNPU(2)
649	@with_comms	649	@with_comms
	@@ -671,7 +671,7 @@ class TestCrossEntropyLoss(NPUDTensorTestBase):
671	self.assertEqual(loss_dtensor.full_tensor(), loss)	671	self.assertEqual(loss_dtensor.full_tensor(), loss)
672	self.assertEqual(log_prob_dtensor.full_tensor(), log_prob)	672	self.assertEqual(log_prob_dtensor.full_tensor(), log_prob)
673		673
674	-	674	+
675	@SupportedDevices(['Ascend910B'])	675	@SupportedDevices(['Ascend910B'])
676	@skipIfUnsupportMultiNPU(2)	676	@skipIfUnsupportMultiNPU(2)
677	@with_comms	677	@with_comms
	@@ -685,14 +685,14 @@ class TestCrossEntropyLoss(NPUDTensorTestBase):
685	loss_dtensor.backward()	685	loss_dtensor.backward()
686	self.assertEqual(input_dtensor.grad.full_tensor(), x.grad)	686	self.assertEqual(input_dtensor.grad.full_tensor(), x.grad)
687		687
688	-	688	+
689	@SupportedDevices(['Ascend910B'])	689	@SupportedDevices(['Ascend910B'])
690	@skipIfUnsupportMultiNPU(2)	690	@skipIfUnsupportMultiNPU(2)
691	@with_comms	691	@with_comms
692	def test_torch_npu_npu_cross_entropy_loss_backward_input_shard0_reduction_is_none(self):	692	def test_torch_npu_npu_cross_entropy_loss_backward_input_shard0_reduction_is_none(self):
693	reductions = ["none", "sum", "mean"]	693	reductions = ["none", "sum", "mean"]
694	x, target, input_dtensor, target_dtensor, mesh = self.generate_data_cross_entropy_loss(8, 8, [Shard(0)], [Shard(0)])	694	x, target, input_dtensor, target_dtensor, mesh = self.generate_data_cross_entropy_loss(8, 8, [Shard(0)], [Shard(0)])
695	-	695	+
696	for re in reductions:	696	for re in reductions:
697	loss, log_prob, _, _ = torch_npu.npu_cross_entropy_loss(x, target, reduction=re)	697	loss, log_prob, _, _ = torch_npu.npu_cross_entropy_loss(x, target, reduction=re)
698	loss_dtensor, log_prob_dtensor, _, _ = torch_npu.npu_cross_entropy_loss(input_dtensor, target_dtensor, reduction=re)	698	loss_dtensor, log_prob_dtensor, _, _ = torch_npu.npu_cross_entropy_loss(input_dtensor, target_dtensor, reduction=re)
	@@ -707,7 +707,7 @@ class TestCrossEntropyLoss(NPUDTensorTestBase):
707	loss.backward()	707	loss.backward()
708	loss_dtensor.backward()	708	loss_dtensor.backward()
709	self.assertEqual(input_dtensor.grad.full_tensor(), x.grad)	709	self.assertEqual(input_dtensor.grad.full_tensor(), x.grad)
710	-	710	+
711	@SupportedDevices(['Ascend910B'])	711	@SupportedDevices(['Ascend910B'])
712	@skipIfUnsupportMultiNPU(2)	712	@skipIfUnsupportMultiNPU(2)
713	@with_comms	713	@with_comms
	@@ -716,7 +716,7 @@ class TestCrossEntropyLoss(NPUDTensorTestBase):
716		716
717	loss, log_prob, _, _ = torch_npu.npu_cross_entropy_loss(x, target, reduction="sum")	717	loss, log_prob, _, _ = torch_npu.npu_cross_entropy_loss(x, target, reduction="sum")
718	loss_dtensor, log_prob_dtensor, _, _ = torch_npu.npu_cross_entropy_loss(input_dtensor, target_dtensor, reduction="sum")	718	loss_dtensor, log_prob_dtensor, _, _ = torch_npu.npu_cross_entropy_loss(input_dtensor, target_dtensor, reduction="sum")
719	-	719	+
720	loss.backward()	720	loss.backward()
721	loss_dtensor.backward()	721	loss_dtensor.backward()
722	self.assertEqual(input_dtensor.grad.full_tensor(), x.grad)	722	self.assertEqual(input_dtensor.grad.full_tensor(), x.grad)
	@@ -741,7 +741,7 @@ class TestRepeatInterleaveSelfInt(NPUDTensorTestBase):
741		741
742	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value, dim=1)	742	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value, dim=1)
743	output = torch.repeat_interleave(input_tensor, repeats_value, dim=1)	743	output = torch.repeat_interleave(input_tensor, repeats_value, dim=1)
744	-	744	+
745	self.assertEqual(output_dtensor.full_tensor(), output)	745	self.assertEqual(output_dtensor.full_tensor(), output)
746		746
747	@SupportedDevices(['Ascend910B'])	747	@SupportedDevices(['Ascend910B'])
	@@ -752,7 +752,7 @@ class TestRepeatInterleaveSelfInt(NPUDTensorTestBase):
752		752
753	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value, dim=1)	753	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value, dim=1)
754	output = torch.repeat_interleave(input_tensor, repeats_value, dim=1)	754	output = torch.repeat_interleave(input_tensor, repeats_value, dim=1)
755	-	755	+
756	self.assertEqual(output_dtensor.full_tensor(), output)	756	self.assertEqual(output_dtensor.full_tensor(), output)
757		757
758	@SupportedDevices(['Ascend910B'])	758	@SupportedDevices(['Ascend910B'])
	@@ -763,7 +763,7 @@ class TestRepeatInterleaveSelfInt(NPUDTensorTestBase):
763		763
764	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value, dim=1)	764	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value, dim=1)
765	output = torch.repeat_interleave(input_tensor, repeats_value, dim=1)	765	output = torch.repeat_interleave(input_tensor, repeats_value, dim=1)
766	-	766	+
767	self.assertEqual(output_dtensor.full_tensor(), output)	767	self.assertEqual(output_dtensor.full_tensor(), output)
768		768
769	@SupportedDevices(['Ascend910B'])	769	@SupportedDevices(['Ascend910B'])
	@@ -774,7 +774,7 @@ class TestRepeatInterleaveSelfInt(NPUDTensorTestBase):
774		774
775	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value)	775	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value)
776	output = torch.repeat_interleave(input_tensor, repeats_value)	776	output = torch.repeat_interleave(input_tensor, repeats_value)
777	-	777	+
778	self.assertEqual(output_dtensor.full_tensor(), output)	778	self.assertEqual(output_dtensor.full_tensor(), output)
779		779
780	@SupportedDevices(['Ascend910B'])	780	@SupportedDevices(['Ascend910B'])
	@@ -785,7 +785,7 @@ class TestRepeatInterleaveSelfInt(NPUDTensorTestBase):
785		785
786	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value, dim=1)	786	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value, dim=1)
787	output = torch.repeat_interleave(input_tensor, repeats_value, dim=1)	787	output = torch.repeat_interleave(input_tensor, repeats_value, dim=1)
788	-	788	+
789	self.assertEqual(output_dtensor.full_tensor(), output)	789	self.assertEqual(output_dtensor.full_tensor(), output)
790		790
791	@SupportedDevices(['Ascend910B'])	791	@SupportedDevices(['Ascend910B'])
	@@ -796,7 +796,7 @@ class TestRepeatInterleaveSelfInt(NPUDTensorTestBase):
796		796
797	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value, dim=1)	797	output_dtensor = torch.repeat_interleave(input_dtensor, repeats_value, dim=1)
798	output = torch.repeat_interleave(input_tensor, repeats_value, dim=1)	798	output = torch.repeat_interleave(input_tensor, repeats_value, dim=1)
799	-	799	+
800	self.assertEqual(output_dtensor.full_tensor(), output)	800	self.assertEqual(output_dtensor.full_tensor(), output)
801		801
802	@SupportedDevices(['Ascend910B'])	802	@SupportedDevices(['Ascend910B'])
	@@ -983,7 +983,7 @@ class TestKLDivLoss(NPUDTensorTestBase):
983		983
984	pred = torch.randn(4, 4, 4, device="npu", requires_grad=True)	984	pred = torch.randn(4, 4, 4, device="npu", requires_grad=True)
985	target = torch.randn(4, 4, 4, device="npu")	985	target = torch.randn(4, 4, 4, device="npu")
986	-	986	+
987	# def test_placement_comb(placements1, placements2):	987	# def test_placement_comb(placements1, placements2):
988	pred_dt = distribute_tensor(pred, mesh, [pred_placement])	988	pred_dt = distribute_tensor(pred, mesh, [pred_placement])
989	# pred_dt = pred_dt	989	# pred_dt = pred_dt

Mtest/distributed/test_allgather.py+2-2

	@@ -73,7 +73,7 @@ class HcclAllGatherTestBase(TestCase):
73	gather_tensor = list()	73	gather_tensor = list()
74	for input_tensor in inputlist:	74	for input_tensor in inputlist:
75	gather_tensor.append(torch.empty_like(input_tensor, device="cpu"))	75	gather_tensor.append(torch.empty_like(input_tensor, device="cpu"))
76	-	76	+
77	for i in range(world_size):	77	for i in range(world_size):
78	p = ctx.Process(	78	p = ctx.Process(
79	target=f,	79	target=f,
	@@ -184,7 +184,7 @@ class HcclAllGatherTest(HcclAllGatherTestBase):
184	for _ in range(dim):	184	for _ in range(dim):
185	shape_list.append(randint(1, max_value))	185	shape_list.append(randint(1, max_value))
186	return create_common_tensor([np.float32, format_list[randint(0, 3)], shape_list], -10, 10)	186	return create_common_tensor([np.float32, format_list[randint(0, 3)], shape_list], -10, 10)
187	-	187	+
188	for world_size in ranks:	188	for world_size in ranks:
189	cpu_excepted_result = list()	189	cpu_excepted_result = list()
190	npu_excepted_result = list()	190	npu_excepted_result = list()

Mtest/distributed/test_device.py+2-2

	@@ -15,7 +15,7 @@ class TestDevice(MultiProcessTestCase):
15	@property	15	@property
16	def world_size(self):	16	def world_size(self):
17	return 1	17	return 1
18	-	18	+
19	def test_event_create(self):	19	def test_event_create(self):
20	a = torch.full((3, 4), float(0), device='npu:0')	20	a = torch.full((3, 4), float(0), device='npu:0')
21	e = torch.npu.Event()	21	e = torch.npu.Event()
	@@ -31,7 +31,7 @@ class TestDevice(MultiProcessTestCase):
31	t.start()	31	t.start()
32	t.join()	32	t.join()
33	self.assertEqual(result[0], 1)	33	self.assertEqual(result[0], 1)
34	-	34	+
35	def test_event_isinstance(self):	35	def test_event_isinstance(self):
36	npu_event = torch.npu.Event()	36	npu_event = torch.npu.Event()
37	self.assertIsInstance(npu_event, torch.npu.Event)	37	self.assertIsInstance(npu_event, torch.npu.Event)

Mtest/distributed/test_device_mesh.py+1-1

	@@ -229,7 +229,7 @@ class DeviceMeshTestF(NPUDTensorTestBase):
229	mesh = torch.arange(self.world_size).to(self.rank)	229	mesh = torch.arange(self.world_size).to(self.rank)
230	with self.assertRaises(ValueError):	230	with self.assertRaises(ValueError):
231	device_mesh = DeviceMesh(self.device_type, mesh)	231	device_mesh = DeviceMesh(self.device_type, mesh)
232	-	232	+
233	@skipIfUnsupportMultiNPU(4)	233	@skipIfUnsupportMultiNPU(4)
234	@with_comms	234	@with_comms
235	def test_get_local_rank(self):	235	def test_get_local_rank(self):

Mtest/distributed/test_distributed.py+1-1

	@@ -671,7 +671,7 @@ class _DistTestBase(object):
671	torch.testing.assert_allclose(running_mean, all_input_var.mean(1))	671	torch.testing.assert_allclose(running_mean, all_input_var.mean(1))
672	torch.testing.assert_allclose(running_var.cpu(), all_input_var.cpu().var(1, unbiased=False))	672	torch.testing.assert_allclose(running_var.cpu(), all_input_var.cpu().var(1, unbiased=False))
673		673
674	- # need more 4 device, less 4 divice there may be accuracy issues	674	+ # need more 4 device, less 4 divice there may be accuracy issues
675	@skipIfUnsupportMultiNPU(4)	675	@skipIfUnsupportMultiNPU(4)
676	def test_DistributedDataParallel_SyncBatchNorm_Diff_Input_Sizes_Running_Value(self):	676	def test_DistributedDataParallel_SyncBatchNorm_Diff_Input_Sizes_Running_Value(self):
677	for bk in [True, False]:	677	for bk in [True, False]:

Mtest/distributed/test_flight_recorder.py+1-1

	@@ -907,7 +907,7 @@ class HcclHeartbeatDumpTest(HCCLTraceTestBase):
907	if self.rank == 0:	907	if self.rank == 0:
908	# sleep for heartbeat dump	908	# sleep for heartbeat dump
909	time.sleep(30)	909	time.sleep(30)
910	-	910	+
911	pg.allreduce(a).wait()	911	pg.allreduce(a).wait()
912		912
913	torch.npu.synchronize(device=device)	913	torch.npu.synchronize(device=device)

Mtest/distributed/test_hccl_shared_buffer.py+5-5

	@@ -71,12 +71,12 @@ class HcclSharedBufferTest(TestCase):
71	self.assertEqual(output, expected,	71	self.assertEqual(output, expected,
72	"rank {} world_size {} dtype {} shape {} Expect receive tensor {} but got {}.".format(	72	"rank {} world_size {} dtype {} shape {} Expect receive tensor {} but got {}.".format(
73	rank, world_size, expected.dtype, expected.shape, expected, output))	73	rank, world_size, expected.dtype, expected.shape, expected, output))
74	-	74	+
75	# For case where we want to examine the memory usage of ccl buffer	75	# For case where we want to examine the memory usage of ccl buffer
76	if mem_diff:	76	if mem_diff:
77	for pg in mem_diff:	77	for pg in mem_diff:
78	used_mem, expected_mem = mem_diff[pg]	78	used_mem, expected_mem = mem_diff[pg]
79	- self.assertLess(used_mem, expected_mem,	79	+ self.assertLess(used_mem, expected_mem,
80	f"Expected memory used to be less than {expected_mem} for {pg}, but got {used_mem}")	80	f"Expected memory used to be less than {expected_mem} for {pg}, but got {used_mem}")
81		81
82	for _ in range(world_size):	82	for _ in range(world_size):
	@@ -90,7 +90,7 @@ class HcclSharedBufferTest(TestCase):
90	op1_expected = 0	90	op1_expected = 0
91	for _ in range(world_size):	91	for _ in range(world_size):
92	op1_expected += inputs	92	op1_expected += inputs
93	-	93	+
94	#dist.ReduceOp.AVG	94	#dist.ReduceOp.AVG
95	expected = 0	95	expected = 0
96	for _ in range(world_size):	96	for _ in range(world_size):
	@@ -222,7 +222,7 @@ class HcclSharedBufferTest(TestCase):
222	del pg1	222	del pg1
223	del pg2	223	del pg2
224		224
225	- hccl_config3 = {"hccl_buffer_name": "subSharedBuffer"} #	225	+ hccl_config3 = {"hccl_buffer_name": "subSharedBuffer"} #
226	options.hccl_config = hccl_config3	226	options.hccl_config = hccl_config3
227	pg3 = dist.new_group(backend='hccl', ranks=ranks, pg_options=options)	227	pg3 = dist.new_group(backend='hccl', ranks=ranks, pg_options=options)
228	dist.all_reduce(input1, group=pg3, async_op=True)	228	dist.all_reduce(input1, group=pg3, async_op=True)
	@@ -348,7 +348,7 @@ class HcclSharedBufferTest(TestCase):
348	p2c.get()	348	p2c.get()
349		349
350	return True	350	return True
351	-	351	+
352	@classmethod	352	@classmethod
353	# pylint:disable=huawei-too-many-arguments	353	# pylint:disable=huawei-too-many-arguments
354	def _test_buffer_memory_with_deleted_pg(cls, rank, input1, world_size, init_pg, c2p, p2c, seq):	354	def _test_buffer_memory_with_deleted_pg(cls, rank, input1, world_size, init_pg, c2p, p2c, seq):

Mtest/distributed/test_nslb.py+1-1

	@@ -31,7 +31,7 @@ class HcclNslbTest(TestCase):
31	dist_group.all_reduce(input1)	31	dist_group.all_reduce(input1)
32		32
33	def _test_multiprocess(self, f, init_pg, input1, world_size, nslb_dir):	33	def _test_multiprocess(self, f, init_pg, input1, world_size, nslb_dir):
34	-	34	+
35	ctx = mp.get_context('spawn')	35	ctx = mp.get_context('spawn')
36		36
37	ps = []	37	ps = []

Mtest/distributed/test_options.py+17-17

Mtest/distributed/test_send_recv.py+1-1

	@@ -228,7 +228,7 @@ class HcclSendRecvDistTest(TestCase):
228	HcclSendRecvDistTest._test_send_recv_dist_with_internal_format_and_offset,	228	HcclSendRecvDistTest._test_send_recv_dist_with_internal_format_and_offset,
229	torch.randn(31, 31),	229	torch.randn(31, 31),
230	HcclSendRecvDistTest._init_dist_hccl)	230	HcclSendRecvDistTest._init_dist_hccl)
231	-	231	+
232	@unittest.skip("Temporarily disabled")	232	@unittest.skip("Temporarily disabled")
233	@skipIfUnsupportMultiNPU(4)	233	@skipIfUnsupportMultiNPU(4)
234	def test_send_recv_hccl_dist_with_p2p(self):	234	def test_send_recv_hccl_dist_with_p2p(self):

Mtest/distributed/test_store_api.py+14-14

	@@ -1,7 +1,7 @@
1	"""	1	"""
2	Add validation cases for torch.distributed APIs on NPU:	2	Add validation cases for torch.distributed APIs on NPU:
3	1. test/distributed/test_store.py from PyTorch community lacks sufficient API validations, so this file is added.	3	1. test/distributed/test_store.py from PyTorch community lacks sufficient API validations, so this file is added.
4	-2. This file validates	4	+2. This file validates
5	torch.distributed.FileStore.path	5	torch.distributed.FileStore.path
6	torch.distributed.Store.__init__	6	torch.distributed.Store.__init__
7	torch.distributed.Store.add	7	torch.distributed.Store.add
	@@ -37,14 +37,14 @@ class TestStoreAPIs(TestCase):
37	"""Test if FileStore actually uses the specified path for data exchange."""	37	"""Test if FileStore actually uses the specified path for data exchange."""
38	with tempfile.TemporaryDirectory() as temp_dir:	38	with tempfile.TemporaryDirectory() as temp_dir:
39	filename = os.path.join(temp_dir, "npu_filestore.txt")	39	filename = os.path.join(temp_dir, "npu_filestore.txt")
40	-	40	+
41	store_master = dist.FileStore(filename, 2)	41	store_master = dist.FileStore(filename, 2)
42	self.assertEqual(store_master.path, filename)	42	self.assertEqual(store_master.path, filename)
43	store_master.set("shared_key", "npu_data")	43	store_master.set("shared_key", "npu_data")
44		44
45	store_worker = dist.FileStore(filename, 2)	45	store_worker = dist.FileStore(filename, 2)
46	val = store_worker.get("shared_key")	46	val = store_worker.get("shared_key")
47	-	47	+
48	self.assertEqual(val, b"npu_data")	48	self.assertEqual(val, b"npu_data")
49		49
50	def test_store_init(self):	50	def test_store_init(self):
	@@ -62,22 +62,22 @@ class TestStoreAPIs(TestCase):
62	"""Test the add operation mathematically on a HashStore."""	62	"""Test the add operation mathematically on a HashStore."""
63	store = dist.HashStore()	63	store = dist.HashStore()
64	key = "test_add_key"	64	key = "test_add_key"
65	-	65	+
66	res1 = store.add(key, 5)	66	res1 = store.add(key, 5)
67	self.assertEqual(res1, 5)	67	self.assertEqual(res1, 5)
68	-	68	+
69	res2 = store.add(key, 10)	69	res2 = store.add(key, 10)
70	self.assertEqual(res2, 15)	70	self.assertEqual(res2, 15)
71	-	71	+
72	self.assertEqual(store.get(key), b"15")	72	self.assertEqual(store.get(key), b"15")
73	-	73	+
74	res3 = store.add(key, -3)	74	res3 = store.add(key, -3)
75	self.assertEqual(res3, 12)	75	self.assertEqual(res3, 12)
76		76
77	def test_store_timeout_behavior(self):	77	def test_store_timeout_behavior(self):
78	"""Test if the timeout property actively interrupts blocking operations."""	78	"""Test if the timeout property actively interrupts blocking operations."""
79	store = dist.HashStore()	79	store = dist.HashStore()
80	-	80	+
81	timeout_seconds = 1	81	timeout_seconds = 1
82	test_timeout = timedelta(seconds=timeout_seconds)	82	test_timeout = timedelta(seconds=timeout_seconds)
83	store.set_timeout(test_timeout)	83	store.set_timeout(test_timeout)
	@@ -92,7 +92,7 @@ class TestStoreAPIs(TestCase):
92	"Timeout" in str(context.exception) or "Wait timeout" in str(context.exception),	92	"Timeout" in str(context.exception) or "Wait timeout" in str(context.exception),
93	f"Exception message does not indicate timeout: {context.exception}"	93	f"Exception message does not indicate timeout: {context.exception}"
94	)	94	)
95	-	95	+
96	self.assertTrue(	96	self.assertTrue(
97	0.8 <= elapsed_time <= 2.5,	97	0.8 <= elapsed_time <= 2.5,
98	f"Actual wait time {elapsed_time:.2f}s did not respect the {timeout_seconds}s timeout."	98	f"Actual wait time {elapsed_time:.2f}s did not respect the {timeout_seconds}s timeout."
	@@ -102,19 +102,19 @@ class TestStoreAPIs(TestCase):
102	"""Test TCPStore host and port by establishing an actual Client-Server connection."""	102	"""Test TCPStore host and port by establishing an actual Client-Server connection."""
103	host = "127.0.0.1"	103	host = "127.0.0.1"
104	port = find_free_port()	104	port = find_free_port()
105	-	105	+
106	server_store = dist.TCPStore(	106	server_store = dist.TCPStore(
107	host_name=host,	107	host_name=host,
108	port=port,	108	port=port,
109	world_size=2,	109	world_size=2,
110	is_master=True,	110	is_master=True,
111	timeout=timedelta(seconds=5),	111	timeout=timedelta(seconds=5),
112	- wait_for_workers=False	112	+ wait_for_workers=False
113	)	113	)
114	-	114	+
115	self.assertEqual(server_store.host, host)	115	self.assertEqual(server_store.host, host)
116	self.assertEqual(server_store.port, port)	116	self.assertEqual(server_store.port, port)
117	-	117	+
118	server_store.set("tcp_key", "tcp_value")	118	server_store.set("tcp_key", "tcp_value")
119		119
120	client_store = dist.TCPStore(	120	client_store = dist.TCPStore(
	@@ -124,7 +124,7 @@ class TestStoreAPIs(TestCase):
124	is_master=False,	124	is_master=False,
125	timeout=timedelta(seconds=5)	125	timeout=timedelta(seconds=5)
126	)	126	)
127	-	127	+
128	client_store.wait(["tcp_key"], timedelta(seconds=5))	128	client_store.wait(["tcp_key"], timedelta(seconds=5))
129	val = client_store.get("tcp_key")	129	val = client_store.get("tcp_key")
130	self.assertEqual(val, b"tcp_value")	130	self.assertEqual(val, b"tcp_value")

Mtest/distributed/test_watchdog.py+2-2

	@@ -40,13 +40,13 @@ class ElasticLaunchTest(TestCase):
40	)	40	)
41	except Exception:	41	except Exception:
42	print("Program fail and exit")	42	print("Program fail and exit")
43	-	43	+
44	end_time = time.time()	44	end_time = time.time()
45	excution_time = end_time - start_time	45	excution_time = end_time - start_time
46	if excution_time > 120:	46	if excution_time > 120:
47	print(f"Excution time using time.time(): {excution_time} seconds")	47	print(f"Excution time using time.time(): {excution_time} seconds")
48	raise RuntimeError("Test case fail")	48	raise RuntimeError("Test case fail")
49	-	49	+
50		50
51	if __name__ == "__main__":	51	if __name__ == "__main__":
52	run_tests()	52	run_tests()

Mtest/distributed/test_with_device.py+1-1

	@@ -14,7 +14,7 @@ class TestWithDevice(TestCase):
14	# -int -> ignore, return -1	14	# -int -> ignore, return -1
15	# future exchangeDevice:	15	# future exchangeDevice:
16	# if < std::numeric_limits<c10::DeviceIndex>::min(), raise error	16	# if < std::numeric_limits<c10::DeviceIndex>::min(), raise error
17	- # else, ignore, return -1	17	+ # else, ignore, return -1
18	for i in [-258, -200.8, -128, -128.8, -127.99, -7, -7.88, -1, -0.2]:	18	for i in [-258, -200.8, -128, -128.8, -127.99, -7, -7.88, -1, -0.2]:
19	s = torch.npu.Stream(i)	19	s = torch.npu.Stream(i)
20	self.assertEqual(s.device_index, 1)	20	self.assertEqual(s.device_index, 1)

Mtest/dynamo/test_minifier.py+5-5

	@@ -51,7 +51,7 @@ inner(torch.randn(20, 20).to("{device}"))
51	self._test_after_dynamo(	51	self._test_after_dynamo(
52	"cuda", "relu_compile_error_TESTING_ONLY", "ReluCompileError"	52	"cuda", "relu_compile_error_TESTING_ONLY", "ReluCompileError"
53	)	53	)
54	-	54	+
55	@requires_npu()	55	@requires_npu()
56	def test_after_dynamo_npu_compile_error(self):	56	def test_after_dynamo_npu_compile_error(self):
57	self._test_after_dynamo(	57	self._test_after_dynamo(
	@@ -63,7 +63,7 @@ inner(torch.randn(20, 20).to("{device}"))
63	self._test_after_dynamo(	63	self._test_after_dynamo(
64	"cuda", "relu_runtime_error_TESTING_ONLY", "ReluRuntimeError"	64	"cuda", "relu_runtime_error_TESTING_ONLY", "ReluRuntimeError"
65	)	65	)
66	-	66	+
67	@requires_npu()	67	@requires_npu()
68	def test_after_dynamo_npu_runtime_error(self):	68	def test_after_dynamo_npu_runtime_error(self):
69	self._test_after_dynamo(	69	self._test_after_dynamo(
	@@ -75,7 +75,7 @@ inner(torch.randn(20, 20).to("{device}"))
75	self._test_after_dynamo(	75	self._test_after_dynamo(
76	"cuda", "relu_accuracy_error_TESTING_ONLY", "AccuracyError"	76	"cuda", "relu_accuracy_error_TESTING_ONLY", "AccuracyError"
77	)	77	)
78	-	78	+
79	@requires_npu()	79	@requires_npu()
80	def test_after_dynamo_npu_accuracy_error(self):	80	def test_after_dynamo_npu_accuracy_error(self):
81	self._test_after_dynamo(	81	self._test_after_dynamo(
	@@ -134,7 +134,7 @@ inner(torch.randn(20, 20, requires_grad=True) + 1)
134	self._test_after_dynamo_backend_passes(	134	self._test_after_dynamo_backend_passes(
135	"cuda", "relu_runtime_error_TESTING_ONLY"	135	"cuda", "relu_runtime_error_TESTING_ONLY"
136	)	136	)
137	-	137	+
138	@requires_npu()	138	@requires_npu()
139	def test_after_dynamo_npu_runtime_backend_passes(self):	139	def test_after_dynamo_npu_runtime_backend_passes(self):
140	self._test_after_dynamo_backend_passes(	140	self._test_after_dynamo_backend_passes(
	@@ -146,7 +146,7 @@ inner(torch.randn(20, 20, requires_grad=True) + 1)
146	self._test_after_dynamo_backend_passes(	146	self._test_after_dynamo_backend_passes(
147	"cuda", "relu_accuracy_error_TESTING_ONLY"	147	"cuda", "relu_accuracy_error_TESTING_ONLY"
148	)	148	)
149	-	149	+
150	@requires_npu()	150	@requires_npu()
151	def test_after_dynamo_npu_accuracy_backend_passes(self):	151	def test_after_dynamo_npu_accuracy_backend_passes(self):
152	self._test_after_dynamo_backend_passes(	152	self._test_after_dynamo_backend_passes(

Mtest/dynamo/test_npugraph_ex.py+2-2

	@@ -44,14 +44,14 @@ class TestNpuGraphEx(TestCase):
44	def custom_compiler(gm: torch.fx.GraphModule, example_inputs):	44	def custom_compiler(gm: torch.fx.GraphModule, example_inputs):
45	compiled_graph = torch.npu.npugraph_ex.compile_fx(gm, example_inputs)	45	compiled_graph = torch.npu.npugraph_ex.compile_fx(gm, example_inputs)
46	return compiled_graph	46	return compiled_graph
47	-	47	+
48	def custom_compiler_with_options(gm: torch.fx.GraphModule, example_inputs):	48	def custom_compiler_with_options(gm: torch.fx.GraphModule, example_inputs):
49	test_kwargs = {	49	test_kwargs = {
50	"clone_input": False	50	"clone_input": False
51	}	51	}
52	compiled_graph = torch.npu.npugraph_ex.compile_fx(gm, example_inputs, test_kwargs)	52	compiled_graph = torch.npu.npugraph_ex.compile_fx(gm, example_inputs, test_kwargs)
53	return compiled_graph	53	return compiled_graph
54	-	54	+
55	def my_backend(gm: torch.fx.GraphModule, example_inputs):	55	def my_backend(gm: torch.fx.GraphModule, example_inputs):
56	return aot_module_simplified(gm, example_inputs, fw_compiler=custom_compiler)	56	return aot_module_simplified(gm, example_inputs, fw_compiler=custom_compiler)
57		57

Mtest/dynamo/test_subclasses.py+1-1

	@@ -904,7 +904,7 @@ class TestNestedTensor(torch._dynamo.test_case.TestCase):
904	@unittest.skipIf(not torch.cuda.is_available(), "requires cuda")	904	@unittest.skipIf(not torch.cuda.is_available(), "requires cuda")
905	def test_basic_autograd_inductor(self):	905	def test_basic_autograd_inductor(self):
906	self._test_autograd("inductor")	906	self._test_autograd("inductor")
907	-	907	+
908	@unittest.skipIf(not torch.npu.is_available(), "requires npu")	908	@unittest.skipIf(not torch.npu.is_available(), "requires npu")
909	def test_basic_autograd_npu_backend(self):	909	def test_basic_autograd_npu_backend(self):
910	npu_backend = torchair.get_npu_backend()	910	npu_backend = torchair.get_npu_backend()

Mtest/dynamo/test_torchair_no_init.py+4-4

Mtest/nn/test_linear_functions.py+1-1

	@@ -20,7 +20,7 @@ class TestLinearFunctions(TestCase):
20	npu_output = F.linear(npu_input, npu_weight)	20	npu_output = F.linear(npu_input, npu_weight)
21		21
22	self.assertRtolEqual(cpu_output.numpy(), npu_output.cpu().numpy())	22	self.assertRtolEqual(cpu_output.numpy(), npu_output.cpu().numpy())
23	-	23	+
24	@unittest.skip("skip test_bilinear now")	24	@unittest.skip("skip test_bilinear now")
25	def test_bilinear(self):	25	def test_bilinear(self):
26	input1 = torch.randn(10, 30)	26	input1 = torch.randn(10, 30)

Mtest/nn/test_loss_functions.py+1-1

	@@ -279,7 +279,7 @@ class TestLossFunctions(TestCase):
279	npu_output = F.hinge_embedding_loss(npu_input, npu_targets)	279	npu_output = F.hinge_embedding_loss(npu_input, npu_targets)
280		280
281	self.assertRtolEqual(cpu_output.detach().numpy(), npu_output.detach().cpu().numpy())	281	self.assertRtolEqual(cpu_output.detach().numpy(), npu_output.detach().cpu().numpy())
282	-	282	+
283	@unittest.skip("skip test_kl_div now")	283	@unittest.skip("skip test_kl_div now")
284	def test_kl_div(self):	284	def test_kl_div(self):
285	input1 = torch.randn(5, 3)	285	input1 = torch.randn(5, 3)

Mtest/nn/test_modules_api.py+10-10

	@@ -16,7 +16,7 @@ class TestNPUModuleAPIs(TestCase):
16	def test_module_device_consistency(self):	16	def test_module_device_consistency(self):
17	"""验证Module设备迁移后参数/缓冲区设备一致"""	17	"""验证Module设备迁移后参数/缓冲区设备一致"""
18	m = nn.Sequential(nn.Linear(10, 20), nn.ReLU(), nn.BatchNorm1d(20)).to(device)	18	m = nn.Sequential(nn.Linear(10, 20), nn.ReLU(), nn.BatchNorm1d(20)).to(device)
19	-	19	+
20	for p in m.parameters():	20	for p in m.parameters():
21	self.assertEqual(p.device, device)	21	self.assertEqual(p.device, device)
22	for b in m.buffers():	22	for b in m.buffers():
	@@ -29,13 +29,13 @@ class TestNPUModuleAPIs(TestCase):
29	m.dict = nn.ModuleDict({"linear": nn.Linear(10, 20), "base": base})	29	m.dict = nn.ModuleDict({"linear": nn.Linear(10, 20), "base": base})
30	m.list = nn.ModuleList([nn.Linear(20, 30), base])	30	m.list = nn.ModuleList([nn.Linear(20, 30), base])
31	m.to(device)	31	m.to(device)
32	-	32	+
33	sd = m.state_dict()	33	sd = m.state_dict()
34	torch.npu.synchronize()	34	torch.npu.synchronize()
35	-	35	+
36	for v in sd.values():	36	for v in sd.values():
37	self.assertEqual(v.device, device)	37	self.assertEqual(v.device, device)
38	-	38	+
39	base2 = nn.Sequential(nn.Linear(10, 20), nn.ReLU(), nn.BatchNorm1d(20))	39	base2 = nn.Sequential(nn.Linear(10, 20), nn.ReLU(), nn.BatchNorm1d(20))
40	m2 = nn.Module()	40	m2 = nn.Module()
41	m2.dict = nn.ModuleDict({"linear": nn.Linear(10, 20), "base": base2})	41	m2.dict = nn.ModuleDict({"linear": nn.Linear(10, 20), "base": base2})
	@@ -43,20 +43,20 @@ class TestNPUModuleAPIs(TestCase):
43	m2.load_state_dict(sd)	43	m2.load_state_dict(sd)
44	m2.to(device)	44	m2.to(device)
45	torch.npu.synchronize()	45	torch.npu.synchronize()
46	-	46	+
47	for (n1, p1), (n2, p2) in zip(m.named_parameters(), m2.named_parameters()):	47	for (n1, p1), (n2, p2) in zip(m.named_parameters(), m2.named_parameters()):
48	self.assertTrue(torch.allclose(p1, p2))	48	self.assertTrue(torch.allclose(p1, p2))
49		49
50	def test_moduledict_operations(self):	50	def test_moduledict_operations(self):
51	"""验证ModuleDict增删/索引/遍历"""	51	"""验证ModuleDict增删/索引/遍历"""
52	m = nn.ModuleDict({"a": nn.Linear(10, 20).to(device)})	52	m = nn.ModuleDict({"a": nn.Linear(10, 20).to(device)})
53	-	53	+
54	self.assertIn("a", m)	54	self.assertIn("a", m)
55	m["b"] = nn.Linear(20, 30).to(device)	55	m["b"] = nn.Linear(20, 30).to(device)
56	self.assertEqual(m["b"].weight.device, device)	56	self.assertEqual(m["b"].weight.device, device)
57	del m["b"]	57	del m["b"]
58	self.assertNotIn("b", m)	58	self.assertNotIn("b", m)
59	-	59	+
60	for sub in m.values():	60	for sub in m.values():
61	for p in sub.parameters():	61	for p in sub.parameters():
62	self.assertEqual(p.device, device)	62	self.assertEqual(p.device, device)
	@@ -64,15 +64,15 @@ class TestNPUModuleAPIs(TestCase):
64	def test_modulelist_operations(self):	64	def test_modulelist_operations(self):
65	"""验证ModuleList索引/新增/删除/遍历"""	65	"""验证ModuleList索引/新增/删除/遍历"""
66	m = nn.ModuleList([nn.Linear(20, 30).to(device), nn.BatchNorm1d(30).to(device)])	66	m = nn.ModuleList([nn.Linear(20, 30).to(device), nn.BatchNorm1d(30).to(device)])
67	-	67	+
68	self.assertEqual(m[0].weight.device, device)	68	self.assertEqual(m[0].weight.device, device)
69	self.assertEqual(m[1].running_mean.device, device)	69	self.assertEqual(m[1].running_mean.device, device)
70	-	70	+
71	m.append(nn.Linear(30, 40).to(device))	71	m.append(nn.Linear(30, 40).to(device))
72	m.insert(0, nn.Linear(10, 20).to(device))	72	m.insert(0, nn.Linear(10, 20).to(device))
73	m.pop(-1) # 修复：指定索引	73	m.pop(-1) # 修复：指定索引
74	m.pop(0) # 修复：指定索引	74	m.pop(0) # 修复：指定索引
75	-	75	+
76	self.assertEqual(len(m), 2)	76	self.assertEqual(len(m), 2)
77	for sub in m:	77	for sub in m:
78	for p in sub.parameters():	78	for p in sub.parameters():

Mtest/nn/test_nonlinear_activation_functions.py+1-1

	@@ -122,7 +122,7 @@ class TestNonLiACFunctions(TestCase):
122	npu_output = F.glu(npu_input)	122	npu_output = F.glu(npu_input)
123		123
124	self.assertRtolEqual(cpu_output.numpy(), npu_output.cpu().numpy())	124	self.assertRtolEqual(cpu_output.numpy(), npu_output.cpu().numpy())
125	-	125	+
126	@unittest.skip("skip test_gelu now")	126	@unittest.skip("skip test_gelu now")
127	def test_gelu(self):	127	def test_gelu(self):
128	input1 = torch.randn(2)	128	input1 = torch.randn(2)

Mtest/nn/test_recurrent_layers.py+1-1

	@@ -14,7 +14,7 @@ class TestRecurrentLayers(TestCase):
14	rnn = nn.RNN(10, 20, 2).npu()	14	rnn = nn.RNN(10, 20, 2).npu()
15	output, hn = rnn(input1, h0)	15	output, hn = rnn(input1, h0)
16	self.assertEqual(output is not None, True)	16	self.assertEqual(output is not None, True)
17	-	17	+
18	@unittest.skip("Temporarily skipping")	18	@unittest.skip("Temporarily skipping")
19	def test_LSTM(self):	19	def test_LSTM(self):
20	input1 = torch.randn(5, 3, 10).npu()	20	input1 = torch.randn(5, 3, 10).npu()

Mtest/nn/test_uninitialized_parameter_cls_to_become.py+2-2

	@@ -9,13 +9,13 @@ torch_npu.npu.set_compile_mode(jit_compile=False)
9	# 修复：将自定义属性设为类属性（确保实例化后必存在）	9	# 修复：将自定义属性设为类属性（确保实例化后必存在）
10	class CustomParameter(torch.nn.Parameter):	10	class CustomParameter(torch.nn.Parameter):
11	custom_attr = "custom_param" # 类属性，所有实例共享，无需__init__赋值	11	custom_attr = "custom_param" # 类属性，所有实例共享，无需__init__赋值
12	-	12	+
13	def __init__(self, data=None, requires_grad=True):	13	def __init__(self, data=None, requires_grad=True):
14	super().__init__(data, requires_grad)	14	super().__init__(data, requires_grad)
15		15
16		16
17	class TestUninitializedParameterClsToBecome(TestCase):	17	class TestUninitializedParameterClsToBecome(TestCase):
18	-	18	+
19	def test_core_functionality_npu(self):	19	def test_core_functionality_npu(self):
20	"""极简验证NPU环境下cls_to_become+materialize核心功能"""	20	"""极简验证NPU环境下cls_to_become+materialize核心功能"""
21	# 1. 创建NPU未初始化参数	21	# 1. 创建NPU未初始化参数

Mtest/nn/test_vision_functions.py+1-1

	@@ -67,7 +67,7 @@ class TestVisionFunctions(TestCase):
67		67
68	def test_affine_grid(self):	68	def test_affine_grid(self):
69	'''	69	'''
70	- Because of the limitation of NPU op, the NPU op will automatically convert the input	70	+ Because of the limitation of NPU op, the NPU op will automatically convert the input
71	fp32 to fp16 for calculation, so the input must be passed data within the representable	71	fp32 to fp16 for calculation, so the input must be passed data within the representable
72	range of fp16.	72	range of fp16.
73	'''	73	'''

Mtest/npu/test_aclgraph_launch_host_func.py+1-1

	@@ -38,7 +38,7 @@ class TestAclgraphLaunchHostFunc(TestCase):
38		38
39	self.capture_stream = torch_npu.npu.Stream()	39	self.capture_stream = torch_npu.npu.Stream()
40	self.graph = torch_npu.npu.NPUGraph()	40	self.graph = torch_npu.npu.NPUGraph()
41	-	41	+
42	torch_npu.npu._subscribe_report(self.capture_stream)	42	torch_npu.npu._subscribe_report(self.capture_stream)
43	a = torch.randn([5, 5]).npu()	43	a = torch.randn([5, 5]).npu()
44	b = torch.randn([5, 5]).npu()	44	b = torch.randn([5, 5]).npu()

Mtest/npu/test_aclgraph_support_blocking.py+3-3

Mtest/npu/test_aclgraph_update.py+14-14

Mtest/npu/test_allocator_multi_thread_prof.py+1-1

	@@ -12,7 +12,7 @@ class RandomDataset(Dataset):
12		12
13	def __len__(self):	13	def __len__(self):
14	return self.len	14	return self.len
15	-	15	+
16	def __getitem__(self, index):	16	def __getitem__(self, index):
17	return self.data[index].clone()	17	return self.data[index].clone()
18		18

Mtest/npu/test_cann_version.py+2-2

	@@ -24,10 +24,10 @@ class TestCANNversion(TestCase):
24	self.assertTrue(is_match, f"The env version is {version_env}. The format of cann version {version} is invalid.")	24	self.assertTrue(is_match, f"The env version is {version_env}. The format of cann version {version} is invalid.")
25	else:	25	else:
26	self.assertTrue(version == "", "When verssion_env < '8.1.RC1', the result of get_cann_version is not right.")	26	self.assertTrue(version == "", "When verssion_env < '8.1.RC1', the result of get_cann_version is not right.")
27	-	27	+
28	version = get_cann_version(module="CAN")	28	version = get_cann_version(module="CAN")
29	self.assertTrue(version == "", "When module is invalid, the result of get_cann_version is not right.")	29	self.assertTrue(version == "", "When module is invalid, the result of get_cann_version is not right.")
30	-	30	+
31	def test_get_driver_version(self):	31	def test_get_driver_version(self):
32	try:	32	try:
33	version = get_cann_version(module="DRIVER")	33	version = get_cann_version(module="DRIVER")

Mtest/npu/test_compatibility.py+5-5

	@@ -12,7 +12,7 @@ import pkgutil
12	import torch	12	import torch
13	from torch.testing._internal.common_utils import TestCase, run_tests	13	from torch.testing._internal.common_utils import TestCase, run_tests
14	from torch._utils_internal import get_file_path_2	14	from torch._utils_internal import get_file_path_2
15	-import torch_npu	15	+import torch_npu
16		16
17		17
18	NOT_IMPORT_LIST = [	18	NOT_IMPORT_LIST = [
	@@ -101,7 +101,7 @@ def is_not_compatibility(base_str, new_str, api_str=None):
101	# case: delete/different default value/different parameter name/different parameter dtype	101	# case: delete/different default value/different parameter name/different parameter dtype
102	if base_diff_params:	102	if base_diff_params:
103	return True	103	return True
104	-	104	+
105	# case: add params	105	# case: add params
106	new_diff_params = set(new_params) - set(base_params)	106	new_diff_params = set(new_params) - set(base_params)
107	# special case	107	# special case
	@@ -260,7 +260,7 @@ class TestPublicApiCompatibility(TestCase):
260	for key, value in base_schema0.items():	260	for key, value in base_schema0.items():
261	if not key.startswith("torch_c_func:") and not key.startswith("torch_npu_public_env:"):	261	if not key.startswith("torch_c_func:") and not key.startswith("torch_npu_public_env:"):
262	base_schema[key] = value	262	base_schema[key] = value
263	-	263	+
264	# load torchair torch_npu_schema.json	264	# load torchair torch_npu_schema.json
265	torchair_schema = {}	265	torchair_schema = {}
266	try:	266	try:
	@@ -270,7 +270,7 @@ class TestPublicApiCompatibility(TestCase):
270	except Exception:	270	except Exception:
271	warnings.warn(	271	warnings.warn(
272	"if you are debugging UT file in clone repo, please recursively update the torchair submodule")	272	"if you are debugging UT file in clone repo, please recursively update the torchair submodule")
273	-	273	+
274	if torchair_schema:	274	if torchair_schema:
275	base_schema.update(torchair_schema)	275	base_schema.update(torchair_schema)
276		276
	@@ -350,7 +350,7 @@ class TestPublicApiCompatibility(TestCase):
350	for func in deleted_apis:	350	for func in deleted_apis:
351	failure_list.append(f"# {func}:")	351	failure_list.append(f"# {func}:")
352	failure_list.append(f" - {func} has been deleted.")	352	failure_list.append(f" - {func} has been deleted.")
353	-	353	+
354	newly_apis = set(now_funcs) - set(base_funcs)	354	newly_apis = set(now_funcs) - set(base_funcs)
355	for func in newly_apis:	355	for func in newly_apis:
356	failure_list.append(f"# {func}:")	356	failure_list.append(f"# {func}:")

Mtest/npu/test_compile_aclgraph_update.py+2-2

	@@ -84,8 +84,8 @@ class TestPAAclgraphUpdate(TestCase):
84		84
85	# 执行注意力计算	85	# 执行注意力计算
86	out = self.ref_masked_attention(	86	out = self.ref_masked_attention(
87	- params_np.query[i:i + 1],	87	+ params_np.query[i:i + 1],
88	- np.stack(keys),	88	+ np.stack(keys),
89	np.stack(values)	89	np.stack(values)
90	)	90	)
91	output[i] = out.reshape(self.num_heads, -1)	91	output[i] = out.reshape(self.num_heads, -1)

Mtest/npu/test_copy.py+15-15

	@@ -14,7 +14,7 @@ class TestCopyKernelMemoryFormat(TestCase):
14	shape_format = [	14	shape_format = [
15	[dtype, 2, shape] for dtype in dtype_list for shape in shape_list	15	[dtype, 2, shape] for dtype in dtype_list for shape in shape_list
16	]	16	]
17	-	17	+
18	for item in shape_format:	18	for item in shape_format:
19	cpu_input, npu_input = create_common_tensor(item, -100, 100)	19	cpu_input, npu_input = create_common_tensor(item, -100, 100)
20	npu_input_copy = cpu_input.npu()	20	npu_input_copy = cpu_input.npu()
	@@ -26,7 +26,7 @@ class TestCopyKernelMemoryFormat(TestCase):
26	shape_format = [	26	shape_format = [
27	[dtype, 2, shape] for dtype in dtype_list for shape in shape_list	27	[dtype, 2, shape] for dtype in dtype_list for shape in shape_list
28	]	28	]
29	-	29	+
30	for item in shape_format:	30	for item in shape_format:
31	cpu_input, npu_input = create_common_tensor(item, -100, 100)	31	cpu_input, npu_input = create_common_tensor(item, -100, 100)
32	cpu_transposed = cpu_input.transpose(-1, -2)	32	cpu_transposed = cpu_input.transpose(-1, -2)
	@@ -40,7 +40,7 @@ class TestCopyKernelMemoryFormat(TestCase):
40	shape_format = [	40	shape_format = [
41	[dtype, 2, shape] for dtype in dtype_list for shape in shape_list	41	[dtype, 2, shape] for dtype in dtype_list for shape in shape_list
42	]	42	]
43	-	43	+
44	for item in shape_format:	44	for item in shape_format:
45	cpu_input, npu_input = create_common_tensor(item, -100, 100)	45	cpu_input, npu_input = create_common_tensor(item, -100, 100)
46	cpu_output = npu_input.cpu()	46	cpu_output = npu_input.cpu()
	@@ -52,7 +52,7 @@ class TestCopyKernelMemoryFormat(TestCase):
52	shape_format = [	52	shape_format = [
53	[dtype, 2, shape] for dtype in dtype_list for shape in shape_list	53	[dtype, 2, shape] for dtype in dtype_list for shape in shape_list
54	]	54	]
55	-	55	+
56	for item in shape_format:	56	for item in shape_format:
57	cpu_input, npu_input = create_common_tensor(item, -100, 100)	57	cpu_input, npu_input = create_common_tensor(item, -100, 100)
58	npu_transposed = npu_input.transpose(-1, -2)	58	npu_transposed = npu_input.transpose(-1, -2)
	@@ -63,14 +63,14 @@ class TestCopyKernelMemoryFormat(TestCase):
63	src_dtype_list = [np.float32, np.float16]	63	src_dtype_list = [np.float32, np.float16]
64	dst_dtype_list = [torch.float16, torch.float32]	64	dst_dtype_list = [torch.float16, torch.float32]
65	shape = [32, 64]	65	shape = [32, 64]
66	-	66	+
67	for src_dtype, dst_dtype in zip(src_dtype_list, dst_dtype_list):	67	for src_dtype, dst_dtype in zip(src_dtype_list, dst_dtype_list):
68	cpu_input = torch.randn(shape, dtype=torch.float32) * 100	68	cpu_input = torch.randn(shape, dtype=torch.float32) * 100
69	cpu_input = cpu_input.to(torch.from_numpy(np.array([])).dtype if src_dtype == np.float32 else torch.float16)	69	cpu_input = cpu_input.to(torch.from_numpy(np.array([])).dtype if src_dtype == np.float32 else torch.float16)
70	-	70	+
71	npu_input = cpu_input.npu()	71	npu_input = cpu_input.npu()
72	npu_output = npu_input.to(dst_dtype)	72	npu_output = npu_input.to(dst_dtype)
73	-	73	+
74	cpu_output = cpu_input.to(dst_dtype)	74	cpu_output = cpu_input.to(dst_dtype)
75	self.assertRtolEqual(npu_output.cpu().numpy(), cpu_output.numpy())	75	self.assertRtolEqual(npu_output.cpu().numpy(), cpu_output.numpy())
76		76
	@@ -80,21 +80,21 @@ class TestCopyKernelMemoryFormat(TestCase):
80	(np.float32, torch.float16),	80	(np.float32, torch.float16),
81	]	81	]
82	shape = [32, 64]	82	shape = [32, 64]
83	-	83	+
84	for src_dtype, dst_dtype in dtype_pairs:	84	for src_dtype, dst_dtype in dtype_pairs:
85	cpu_input, npu_input = create_common_tensor([src_dtype, 0, shape], -100, 100)	85	cpu_input, npu_input = create_common_tensor([src_dtype, 0, shape], -100, 100)
86	cpu_output = npu_input.cpu().to(dst_dtype)	86	cpu_output = npu_input.cpu().to(dst_dtype)
87	-	87	+
88	expected = cpu_input.to(dst_dtype)	88	expected = cpu_input.to(dst_dtype)
89	self.assertRtolEqual(cpu_output.numpy(), expected.numpy())	89	self.assertRtolEqual(cpu_output.numpy(), expected.numpy())
90		90
91	def test_h2d_copy_slice_tensor(self):	91	def test_h2d_copy_slice_tensor(self):
92	shape = [64, 128]	92	shape = [64, 128]
93	cpu_input = torch.randn(shape)	93	cpu_input = torch.randn(shape)
94	-	94	+
95	cpu_slice = cpu_input[10:30, 20:60]	95	cpu_slice = cpu_input[10:30, 20:60]
96	npu_slice = cpu_slice.npu()	96	npu_slice = cpu_slice.npu()
97	-	97	+
98	npu_contiguous = npu_slice.contiguous()	98	npu_contiguous = npu_slice.contiguous()
99	self.assertRtolEqual(npu_contiguous.cpu().numpy(), cpu_slice.contiguous().numpy())	99	self.assertRtolEqual(npu_contiguous.cpu().numpy(), cpu_slice.contiguous().numpy())
100		100
	@@ -102,7 +102,7 @@ class TestCopyKernelMemoryFormat(TestCase):
102	shape = [64, 128]	102	shape = [64, 128]
103	cpu_input = torch.randn(shape)	103	cpu_input = torch.randn(shape)
104	npu_input = cpu_input.npu()	104	npu_input = cpu_input.npu()
105	-	105	+
106	npu_slice = npu_input[10:30, 20:60]	106	npu_slice = npu_input[10:30, 20:60]
107	cpu_slice = npu_slice.cpu()	107	cpu_slice = npu_slice.cpu()
108	self.assertRtolEqual(cpu_slice.numpy(), cpu_input[10:30, 20:60].contiguous().numpy())	108	self.assertRtolEqual(cpu_slice.numpy(), cpu_input[10:30, 20:60].contiguous().numpy())
	@@ -111,7 +111,7 @@ class TestCopyKernelMemoryFormat(TestCase):
111	shape = [1, 64, 1]	111	shape = [1, 64, 1]
112	cpu_input = torch.randn(shape)	112	cpu_input = torch.randn(shape)
113	npu_input = cpu_input.npu()	113	npu_input = cpu_input.npu()
114	-	114	+
115	npu_broadcast = npu_input.expand(4, 64, 128)	115	npu_broadcast = npu_input.expand(4, 64, 128)
116	cpu_output = npu_broadcast.cpu()	116	cpu_output = npu_broadcast.cpu()
117	self.assertRtolEqual(cpu_output.numpy(), cpu_input.expand(4, 64, 128).contiguous().numpy())	117	self.assertRtolEqual(cpu_output.numpy(), cpu_input.expand(4, 64, 128).contiguous().numpy())
	@@ -128,8 +128,8 @@ class TestCopyKernelMemoryFormat(TestCase):
128	shape = [32, 64, 128]	128	shape = [32, 64, 128]
129	cpu_input = torch.randn(shape)	129	cpu_input = torch.randn(shape)
130	npu_input = cpu_input.npu()	130	npu_input = cpu_input.npu()
131	-	131	+
132	- npu_permuted = npu_input.permute(2, 0, 1)	132	+ npu_permuted = npu_input.permute(2, 0, 1)
133	cpu_output = npu_permuted.cpu()	133	cpu_output = npu_permuted.cpu()
134	self.assertRtolEqual(cpu_output.numpy(), cpu_input.permute(2, 0, 1).contiguous().numpy())	134	self.assertRtolEqual(cpu_output.numpy(), cpu_input.permute(2, 0, 1).contiguous().numpy())
135		135

Mtest/npu/test_dlpack.py+17-17

Mtest/npu/test_errcode.py+2-2

	@@ -10,13 +10,13 @@ class TestErrorCode(TestCase):
10	def test_set_per_process_memory_fraction(self):	10	def test_set_per_process_memory_fraction(self):
11	with self.assertRaisesRegex(TypeError, "ERR00002 PTA invalid type"):	11	with self.assertRaisesRegex(TypeError, "ERR00002 PTA invalid type"):
12	torch_npu.npu.set_per_process_memory_fraction(1)	12	torch_npu.npu.set_per_process_memory_fraction(1)
13	-	13	+
14	def test_div(self):	14	def test_div(self):
15	x1 = torch.tensor(1).npu()	15	x1 = torch.tensor(1).npu()
16	x2 = torch.tensor(1).npu()	16	x2 = torch.tensor(1).npu()
17	with self.assertRaisesRegex(RuntimeError, "ERR01001 OPS invalid parameter"):	17	with self.assertRaisesRegex(RuntimeError, "ERR01001 OPS invalid parameter"):
18	torch.div(x1, x2, rounding_mode="test")	18	torch.div(x1, x2, rounding_mode="test")
19	-	19	+
20		20
21	if __name__ == "__main__":	21	if __name__ == "__main__":
22	run_tests()	22	run_tests()

Mtest/npu/test_mstx.py+1-1

	@@ -56,7 +56,7 @@ class TestMstx(TestCase):
56	self.assertEqual("", self.mark_domain)	56	self.assertEqual("", self.mark_domain)
57	torch_npu.npu.mstx.mark("test", stream=1, domain="test")	57	torch_npu.npu.mstx.mark("test", stream=1, domain="test")
58	self.assertEqual("", self.mark_msg)	58	self.assertEqual("", self.mark_msg)
59	- self.assertEqual("", self.mark_domain)	59	+ self.assertEqual("", self.mark_domain)
60		60
61	# valid inputs	61	# valid inputs
62	torch_npu.npu.mstx.mark("test1")	62	torch_npu.npu.mstx.mark("test1")

Mtest/npu/test_multi_devices_single_process.py+2-2

	@@ -102,7 +102,7 @@ class TestOp(TestCase):
102	output = torch.abs(input1)	102	output = torch.abs(input1)
103	output = output.cpu().numpy()	103	output = output.cpu().numpy()
104	return output	104	return output
105	-	105	+
106	def _test_abs(self, device="npu:1"):	106	def _test_abs(self, device="npu:1"):
107	torch.npu.set_device(0)	107	torch.npu.set_device(0)
108	cpu_input = torch.Tensor([1, -2, -10])	108	cpu_input = torch.Tensor([1, -2, -10])
	@@ -161,7 +161,7 @@ class TestOp(TestCase):
161	scale = 1 / 0.0078125	161	scale = 1 / 0.0078125
162	return torch_npu.npu_prompt_flash_attention(	162	return torch_npu.npu_prompt_flash_attention(
163	query, key, value, num_heads=32, input_layout="BNSD", scale_value=scale, pre_tokens=65535, next_tokens=65535, sparse_mode=0)	163	query, key, value, num_heads=32, input_layout="BNSD", scale_value=scale, pre_tokens=65535, next_tokens=65535, sparse_mode=0)
164	-	164	+
165	@SupportedDevices(['Ascend910B'])	165	@SupportedDevices(['Ascend910B'])
166	def _test_npu_prompt_flash_attention(self, device="npu:1"):	166	def _test_npu_prompt_flash_attention(self, device="npu:1"):
167	torch.npu.set_device(0)	167	torch.npu.set_device(0)

Mtest/npu/test_multi_stream_lazy_reclaim.py+7-7

	@@ -25,7 +25,7 @@ def extract_aclrtQueryEventStatus_count(prof_dir):
25		25
26	Args:	26	Args:
27	prof_dir: str, path to the profiling result directory	27	prof_dir: str, path to the profiling result directory
28	-	28	+
29	Returns:	29	Returns:
30	count: int, call count of aclrtQueryEventStatus, 0 if not found	30	count: int, call count of aclrtQueryEventStatus, 0 if not found
31	"""	31	"""
	@@ -166,12 +166,12 @@ class TestMultiStreamLazyReclaim(TestCase):
166	- lazy reclaim mode: Only queries in the following cases:	166	- lazy reclaim mode: Only queries in the following cases:
167	1. No available memory block found (!block_found)	167	1. No available memory block found (!block_found)
168	2. Event queue exceeds threshold kLazyQuerySize (512)	168	2. Event queue exceeds threshold kLazyQuerySize (512)
169	-	169	+
170	Test Method:	170	Test Method:
171	Use multiprocessing to test in two separate processes:	171	Use multiprocessing to test in two separate processes:
172	- Process 1: Enable multi_stream_lazy_reclaim	172	- Process 1: Enable multi_stream_lazy_reclaim
173	- Process 2: Disable multi_stream_lazy_reclaim	173	- Process 2: Disable multi_stream_lazy_reclaim
174	-	174	+
175	Each process sets environment variables independently to ensure configuration takes effect.	175	Each process sets environment variables independently to ensure configuration takes effect.
176	"""	176	"""
177		177
	@@ -197,7 +197,7 @@ class TestMultiStreamLazyReclaim(TestCase):
197	)	197	)
198		198
199	process.start()	199	process.start()
200	- process.join(timeout=300) #	200	+ process.join(timeout=300) #
201		201
202	if process.is_alive():	202	if process.is_alive():
203	process.terminate()	203	process.terminate()
	@@ -214,7 +214,7 @@ class TestMultiStreamLazyReclaim(TestCase):
214		214
215	eager_counts = results["eager"]	215	eager_counts = results["eager"]
216	lazy_counts = results["lazy"]	216	lazy_counts = results["lazy"]
217	-	217	+
218	# Output comparison results	218	# Output comparison results
219	print(f"\n========== Event Query Count Comparison ==========")	219	print(f"\n========== Event Query Count Comparison ==========")
220	print(f"Eager reclaim (multi_stream_lazy_reclaim:False): {eager_counts}")	220	print(f"Eager reclaim (multi_stream_lazy_reclaim:False): {eager_counts}")
	@@ -223,8 +223,8 @@ class TestMultiStreamLazyReclaim(TestCase):
223	# Core validation: aclrtQueryEventStatus call count in lazy mode must be less than eager mode	223	# Core validation: aclrtQueryEventStatus call count in lazy mode must be less than eager mode
224	# This is direct evidence that multi_stream_lazy_reclaim feature is working	224	# This is direct evidence that multi_stream_lazy_reclaim feature is working
225	self.assertLessEqual(	225	self.assertLessEqual(
226	- lazy_counts,	226	+ lazy_counts,
227	- eager_counts,	227	+ eager_counts,
228	f"Lazy reclaim mode should reduce event queries. "	228	f"Lazy reclaim mode should reduce event queries. "
229	f"Eager: {eager_counts}, Lazy: {lazy_counts}. "	229	f"Eager: {eager_counts}, Lazy: {lazy_counts}. "
230	f"If lazy >= eager, the optimization may not be working."	230	f"If lazy >= eager, the optimization may not be working."

Mtest/npu/test_npu_format.py+2-2

	@@ -5,7 +5,7 @@ from torch_npu._C import _weak_ref_tensor
5		5
6		6
7	class TestNPUFormat(TestCase):	7	class TestNPUFormat(TestCase):
8	-	8	+
9	def test_enum_values(self):	9	def test_enum_values(self):
10	"""test the enumeration value"""	10	"""test the enumeration value"""
11	self.assertEqual(torch_npu.Format.NCHW.value, 0)	11	self.assertEqual(torch_npu.Format.NCHW.value, 0)
	@@ -24,7 +24,7 @@ class TestNPUFormat(TestCase):
24	self.assertEqual(fmt2, torch_npu.Format.NHWC)	24	self.assertEqual(fmt2, torch_npu.Format.NHWC)
25		25
26	torch_npu.npu.config.allow_internal_format = True	26	torch_npu.npu.config.allow_internal_format = True
27	-	27	+
28	out3 = torch_npu.npu_format_cast(tensor, torch_npu.Format.FRACTAL_NZ)	28	out3 = torch_npu.npu_format_cast(tensor, torch_npu.Format.FRACTAL_NZ)
29	fmt3 = torch_npu.get_npu_format(out3)	29	fmt3 = torch_npu.get_npu_format(out3)
30	self.assertEqual(fmt3, torch_npu.Format.FRACTAL_NZ)	30	self.assertEqual(fmt3, torch_npu.Format.FRACTAL_NZ)

Mtest/npu/test_public_bindings.py+6-6

	@@ -705,7 +705,7 @@ class TestPublicBindings(TestCase):
705	return f'v{version_list[0]}.{version_list[1]}'	705	return f'v{version_list[0]}.{version_list[1]}'
706	else:	706	else:
707	raise RuntimeError("Invalid torch_npu version.")	707	raise RuntimeError("Invalid torch_npu version.")
708	-	708	+
709	try:	709	try:
710	file_abspath = os.path.abspath(__file__)	710	file_abspath = os.path.abspath(__file__)
711	air_path = 'third_party/torchair/torchair/tests/st/allowlist_for_publicAPI.json'	711	air_path = 'third_party/torchair/torchair/tests/st/allowlist_for_publicAPI.json'
	@@ -716,7 +716,7 @@ class TestPublicBindings(TestCase):
716	except Exception:	716	except Exception:
717	update_allow_dict_torchair = {}	717	update_allow_dict_torchair = {}
718	warnings.warn("if you are debugging UT file in clone repo, please recursively update the torchair submodule")	718	warnings.warn("if you are debugging UT file in clone repo, please recursively update the torchair submodule")
719	-	719	+
720	try:	720	try:
721	file_abspath = os.path.abspath(__file__)	721	file_abspath = os.path.abspath(__file__)
722	op_plugin_path = 'third_party/op-plugin/test/allowlist_for_publicAPI.json'	722	op_plugin_path = 'third_party/op-plugin/test/allowlist_for_publicAPI.json'
	@@ -733,7 +733,7 @@ class TestPublicBindings(TestCase):
733	except Exception as e:	733	except Exception as e:
734	allow_dict_op_plugin = {}	734	allow_dict_op_plugin = {}
735	warnings.warn(f"{e}")	735	warnings.warn(f"{e}")
736	-	736	+
737	with open(get_file_path_2(os.path.dirname(os.path.dirname(__file__)),	737	with open(get_file_path_2(os.path.dirname(os.path.dirname(__file__)),
738	'allowlist_for_publicAPI.json')) as json_file:	738	'allowlist_for_publicAPI.json')) as json_file:
739	# no new entries should be added to this allow_dict.	739	# no new entries should be added to this allow_dict.
	@@ -748,16 +748,16 @@ class TestPublicBindings(TestCase):
748	with open(	748	with open(
749	os.path.join(os.path.dirname(os.path.dirname(__file__)), 'deprecated_apis.json')) as json_file:	749	os.path.join(os.path.dirname(os.path.dirname(__file__)), 'deprecated_apis.json')) as json_file:
750	deprecated_dict = json.load(json_file)	750	deprecated_dict = json.load(json_file)
751	-	751	+
752	if update_allow_dict_torchair:	752	if update_allow_dict_torchair:
753	allow_dict.update(update_allow_dict_torchair)	753	allow_dict.update(update_allow_dict_torchair)
754	-	754	+
755	if allow_dict_op_plugin and "torch_npu" in allow_dict_op_plugin:	755	if allow_dict_op_plugin and "torch_npu" in allow_dict_op_plugin:
756	if "torch_npu" in allow_dict:	756	if "torch_npu" in allow_dict:
757	allow_dict["torch_npu"].extend(allow_dict_op_plugin["torch_npu"])	757	allow_dict["torch_npu"].extend(allow_dict_op_plugin["torch_npu"])
758	else:	758	else:
759	allow_dict.update(allow_dict_op_plugin["torch_npu"])	759	allow_dict.update(allow_dict_op_plugin["torch_npu"])
760	-	760	+
761	def test_module(modname):	761	def test_module(modname):
762	try:	762	try:
763	if "__main__" in modname or \	763	if "__main__" in modname or \

Mtest/npu/test_resnet.py+2-2

	@@ -167,7 +167,7 @@ def train(model, criterion, optimizer, epoch):
167	num += 1	167	num += 1
168	yield (torch.randn([128, 3, 224, 224]).npu() + torch.randint(-2, 2, [128, 3, 224, 224]).npu()).cpu(), \	168	yield (torch.randn([128, 3, 224, 224]).npu() + torch.randint(-2, 2, [128, 3, 224, 224]).npu()).cpu(), \
169	torch.randint(1, 1000, [128])	169	torch.randint(1, 1000, [128])
170	-	170	+
171	# switch to train mode	171	# switch to train mode
172	model.train()	172	model.train()
173		173
	@@ -217,7 +217,7 @@ def validate(model, criterion):
217	1,	217	1,
218	[batch_time, losses, top1, top5],	218	[batch_time, losses, top1, top5],
219	prefix='Test: ')	219	prefix='Test: ')
220	-	220	+
221	def fake_val_data(num):	221	def fake_val_data(num):
222	while num < 5:	222	while num < 5:
223	num += 1	223	num += 1

Mtest/npu/test_sanitizer.py+1-1

	@@ -14,7 +14,7 @@ REPO_ROOT = os.path.abspath(os.path.join(os.path.dirname(__file__), "..", ".."))
14	PYTORCH_INSTALL_PATH = os.path.dirname(os.path.realpath(torch.__file__))	14	PYTORCH_INSTALL_PATH = os.path.dirname(os.path.realpath(torch.__file__))
15	PYTORCH_NPU_INSTALL_PATH = os.path.dirname(os.path.realpath(torch_npu.__file__))	15	PYTORCH_NPU_INSTALL_PATH = os.path.dirname(os.path.realpath(torch_npu.__file__))
16		16
17	-	17	+
18	class TestSanitizer(TestCase):	18	class TestSanitizer(TestCase):
19	def tearDown(self):	19	def tearDown(self):
20	if sanitizer.npu_sanitizer.dispatch is not None:	20	if sanitizer.npu_sanitizer.dispatch is not None:

Mtest/npu/test_sanitizer_record_stream.py+1-1

	@@ -35,7 +35,7 @@ import os
35	import torch	35	import torch
36	import torch.cuda._sanitizer as csan	36	import torch.cuda._sanitizer as csan
37	import torch.distributed as dist	37	import torch.distributed as dist
38	-	38	+
39	import torch_npu	39	import torch_npu
40	from torch_npu.testing.testcase import TestCase, run_tests	40	from torch_npu.testing.testcase import TestCase, run_tests
41		41

Mtest/npu/test_save_async.py+4-4

	@@ -22,7 +22,7 @@ class TestAsyncSave(TestCase):
22	@classmethod	22	@classmethod
23	def tearDownClass(cls):	23	def tearDownClass(cls):
24	PathManager.remove_path_safety(TestAsyncSave.test_save_path)	24	PathManager.remove_path_safety(TestAsyncSave.test_save_path)
25	-	25	+
26	def wait_for_save_completion(self, file_path, timeout_sec=60, poll_interval_sec=0.5):	26	def wait_for_save_completion(self, file_path, timeout_sec=60, poll_interval_sec=0.5):
27	start_time = time.time()	27	start_time = time.time()
28		28
	@@ -43,13 +43,13 @@ class TestAsyncSave(TestCase):
43	save_tensor = torch.rand(1024, dtype=torch.float32).npu()	43	save_tensor = torch.rand(1024, dtype=torch.float32).npu()
44	async_save_path = os.path.join(TestAsyncSave.test_save_path, "async_save_tensor.pt")	44	async_save_path = os.path.join(TestAsyncSave.test_save_path, "async_save_tensor.pt")
45	torch_npu.utils.save_async(save_tensor, async_save_path)	45	torch_npu.utils.save_async(save_tensor, async_save_path)
46	-	46	+
47	if self.wait_for_save_completion(async_save_path):	47	if self.wait_for_save_completion(async_save_path):
48	tensor_async = torch.load(async_save_path, weights_only=False)	48	tensor_async = torch.load(async_save_path, weights_only=False)
49	self.assertEqual(tensor_async, save_tensor)	49	self.assertEqual(tensor_async, save_tensor)
50	else:	50	else:
51	self.assertTrue(False, f"{async_save_path} is not exist!")	51	self.assertTrue(False, f"{async_save_path} is not exist!")
52	-	52	+
53	def test_save_async(self):	53	def test_save_async(self):
54	loss1 = [1.6099495, 1.6099086, 1.6098710]	54	loss1 = [1.6099495, 1.6099086, 1.6098710]
55	loss2 = []	55	loss2 = []
	@@ -78,7 +78,7 @@ class TestAsyncSave(TestCase):
78	loss.backward()	78	loss.backward()
79		79
80	optimerizer.step()	80	optimerizer.step()
81	-	81	+
82	loss2.append(loss)	82	loss2.append(loss)
83	checkpoint = {	83	checkpoint = {
84	"model": model.state_dict(),	84	"model": model.state_dict(),

Mtest/npu/test_storage.py+1-1

	@@ -203,7 +203,7 @@ class TestStorage(TestCase):
203	self.assertEqual(cpu_res.size(), npu_res.size())	203	self.assertEqual(cpu_res.size(), npu_res.size())
204	self.assertEqual(cpu_res, npu_res.cpu())	204	self.assertEqual(cpu_res, npu_res.cpu())
205	self.assertEqual(cpu_res.tolist(), npu_res.cpu().tolist())	205	self.assertEqual(cpu_res.tolist(), npu_res.cpu().tolist())
206	-	206	+
207	@SupportedDevices(['Ascend910B'])	207	@SupportedDevices(['Ascend910B'])
208	def _test_datatype_cast_complex(cpu_storage, npu_storage):	208	def _test_datatype_cast_complex(cpu_storage, npu_storage):
209	dtypes = [	209	dtypes = [

Mtest/npu/test_swapped_memory_allocator.py+13-13

	@@ -39,13 +39,13 @@ class TestNPUSwappedMemoryAllocator(unittest.TestCase):
39	tensors = []	39	tensors = []
40	for i in range(5):	40	for i in range(5):
41	tensor = torch_npu.empty_with_swapped_memory(	41	tensor = torch_npu.empty_with_swapped_memory(
42	- [256 * (i + 1)],	42	+ [256 * (i + 1)],
43	- dtype=torch.float32,	43	+ dtype=torch.float32,
44	device='npu:0'	44	device='npu:0'
45	)	45	)
46	tensor.fill_(float(i))	46	tensor.fill_(float(i))
47	tensors.append(tensor)	47	tensors.append(tensor)
48	-	48	+
49	for t in tensors:	49	for t in tensors:
50	del t	50	del t
51	del tensors	51	del tensors
	@@ -57,34 +57,34 @@ class TestNPUSwappedMemoryAllocator(unittest.TestCase):
57	def test_02_async_operations_and_release(self):	57	def test_02_async_operations_and_release(self):
58	"""Test 2: Async operations followed by release (verify stream sync works)"""	58	"""Test 2: Async operations followed by release (verify stream sync works)"""
59	tensor = torch_npu.empty_with_swapped_memory(	59	tensor = torch_npu.empty_with_swapped_memory(
60	- [512, 512],	60	+ [512, 512],
61	- dtype=torch.float32,	61	+ dtype=torch.float32,
62	device='npu:0'	62	device='npu:0'
63	)	63	)
64	-	64	+
65	tensor.fill_(1.0)	65	tensor.fill_(1.0)
66	-	66	+
67	for i in range(10):	67	for i in range(10):
68	tensor = tensor * 1.1	68	tensor = tensor * 1.1
69	tensor = tensor + i * 0.1	69	tensor = tensor + i * 0.1
70	-	70	+
71	tensor.sqrt_()	71	tensor.sqrt_()
72	-	72	+
73	expected_tensor = torch.empty([512, 512], dtype=torch.float32, device='npu:0')	73	expected_tensor = torch.empty([512, 512], dtype=torch.float32, device='npu:0')
74	expected_tensor.fill_(1.0)	74	expected_tensor.fill_(1.0)
75	for i in range(10):	75	for i in range(10):
76	expected_tensor = expected_tensor * 1.1	76	expected_tensor = expected_tensor * 1.1
77	expected_tensor = expected_tensor + i * 0.1	77	expected_tensor = expected_tensor + i * 0.1
78	expected_tensor = expected_tensor.sqrt()	78	expected_tensor = expected_tensor.sqrt()
79	-	79	+
80	self.assertTrue(torch.allclose(tensor, expected_tensor, rtol=1e-5, atol=1e-5))	80	self.assertTrue(torch.allclose(tensor, expected_tensor, rtol=1e-5, atol=1e-5))
81	-	81	+
82	weak_ref = weakref.ref(tensor)	82	weak_ref = weakref.ref(tensor)
83	del tensor	83	del tensor
84	gc.collect()	84	gc.collect()
85	-	85	+
86	self.assertIsNone(weak_ref())	86	self.assertIsNone(weak_ref())
87	-	87	+
88	torch.npu.empty_cache()	88	torch.npu.empty_cache()
89		89
90		90

Mtest/npu/test_tensors.py+2-2

	@@ -163,10 +163,10 @@ class TestTensor(TestCase):
163	scalar_input = torch.randn(1).npu()	163	scalar_input = torch.randn(1).npu()
164	bool_scalar = scalar_input.to(torch.bool)	164	bool_scalar = scalar_input.to(torch.bool)
165	self.assertTrue(isinstance(bool_scalar.item(), bool))	165	self.assertTrue(isinstance(bool_scalar.item(), bool))
166	-	166	+
167	half_scalar = scalar_input.to(torch.float16)	167	half_scalar = scalar_input.to(torch.float16)
168	self.assertTrue(isinstance(half_scalar.item(), float))	168	self.assertTrue(isinstance(half_scalar.item(), float))
169	-	169	+
170	bf16_scalar = scalar_input.to(torch.bfloat16)	170	bf16_scalar = scalar_input.to(torch.bfloat16)
171	self.assertTrue(isinstance(bf16_scalar.item(), float))	171	self.assertTrue(isinstance(bf16_scalar.item(), float))
172		172

Mtest/npu/test_torch_npu.py+5-5

Mtest/onnx/onnx_test_common.py+1-1

	@@ -428,7 +428,7 @@ def add_decorate_info(
428	# Skip does not apply to this opset	428	# Skip does not apply to this opset
429	continue	429	continue
430	opinfo = ops_mapping.get((decorate_meta.op_name, decorate_meta.variant_name))	430	opinfo = ops_mapping.get((decorate_meta.op_name, decorate_meta.variant_name))
431	- assert opinfo is not None, (	431	+ assert opinfo is not None, (
432	f"Couldn't find OpInfo for {decorate_meta}. Did you need to specify variant_name?")	432	f"Couldn't find OpInfo for {decorate_meta}. Did you need to specify variant_name?")
433	assert decorate_meta.model_type is None, (	433	assert decorate_meta.model_type is None, (
434	f"Tested op: {decorate_meta.op_name} in wrong position! "	434	f"Tested op: {decorate_meta.op_name} in wrong position! "

Mtest/onnx/test_combined_onnx_ops.py+1-1

	@@ -216,7 +216,7 @@ class TestOnnxOps(TestCase):
216		216
217	torch.npu.config.allow_internal_format = True	217	torch.npu.config.allow_internal_format = True
218	torch.npu.set_compile_mode(jit_compile=True)	218	torch.npu.set_compile_mode(jit_compile=True)
219	-	219	+
220	def export_onnx(onnx_model_name):	220	def export_onnx(onnx_model_name):
221	input_ = torch.rand([1, 128, 4, 14, 14]).npu()	221	input_ = torch.rand([1, 128, 4, 14, 14]).npu()
222	model = Model().to("npu")	222	model = Model().to("npu")

Mtest/onnx/test_wrapper_onnx_ops.py+15-15

	@@ -194,7 +194,7 @@ class TestOnnxOps(TestCase):
194	class Model(torch.nn.Module):	194	class Model(torch.nn.Module):
195	def __init__(self):	195	def __init__(self):
196	super(Model, self).__init__()	196	super(Model, self).__init__()
197	-	197	+
198	def forward(self, x):	198	def forward(self, x):
199	return torch_npu.npu_geglu(x)	199	return torch_npu.npu_geglu(x)
200		200
	@@ -203,7 +203,7 @@ class TestOnnxOps(TestCase):
203	model = Model().to("npu")	203	model = Model().to("npu")
204	model(x)	204	model(x)
205	self.onnx_export(model, x, onnx_model_name, ["input"], ["output1", "output2"])	205	self.onnx_export(model, x, onnx_model_name, ["input"], ["output1", "output2"])
206	-	206	+
207	onnx_model_name = "model_npu_geglu.onnx"	207	onnx_model_name = "model_npu_geglu.onnx"
208	export_onnx(onnx_model_name)	208	export_onnx(onnx_model_name)
209	assert(os.path.isfile(os.path.join(TestOnnxOps.test_onnx_path,	209	assert(os.path.isfile(os.path.join(TestOnnxOps.test_onnx_path,
	@@ -1156,7 +1156,7 @@ class TestOnnxOps(TestCase):
1156		1156
1157	def forward(self, sorted_experts):	1157	def forward(self, sorted_experts):
1158	return torch_npu.npu_moe_compute_expert_tokens(sorted_experts=5)	1158	return torch_npu.npu_moe_compute_expert_tokens(sorted_experts=5)
1159	-	1159	+
1160	def export_onnx(onnx_model_name):	1160	def export_onnx(onnx_model_name):
1161	data = list(range(20))	1161	data = list(range(20))
1162	experts = torch.tensor(data, dtype=torch.int32).npu()	1162	experts = torch.tensor(data, dtype=torch.int32).npu()
	@@ -1167,7 +1167,7 @@ class TestOnnxOps(TestCase):
1167	onnx_model_name = "model_moe_compute_expert_tokens.onnx"	1167	onnx_model_name = "model_moe_compute_expert_tokens.onnx"
1168	export_onnx(onnx_model_name)	1168	export_onnx(onnx_model_name)
1169	assert (os.path.isfile(os.path.join(TestOnnxOps.test_onnx_path,	1169	assert (os.path.isfile(os.path.join(TestOnnxOps.test_onnx_path,
1170	- onnx_model_name)))	1170	+ onnx_model_name)))
1171		1171
1172	@unittest.skip("skip now")	1172	@unittest.skip("skip now")
1173	def test_wrapper_npu_mish(self):	1173	def test_wrapper_npu_mish(self):
	@@ -1226,7 +1226,7 @@ class TestOnnxOps(TestCase):
1226	epsilon = 1e-6	1226	epsilon = 1e-6
1227	x = torch_npu.npu_rms_norm(x, gamma, epsilon)	1227	x = torch_npu.npu_rms_norm(x, gamma, epsilon)
1228	return x	1228	return x
1229	-	1229	+
1230	def export_onnx(onnx_model_name):	1230	def export_onnx(onnx_model_name):
1231	x = torch.rand(10, 1024).uniform_(-3, 3).npu().half()	1231	x = torch.rand(10, 1024).uniform_(-3, 3).npu().half()
1232	gamma = torch.rand(1024).uniform_(-3, 3).npu().half()	1232	gamma = torch.rand(1024).uniform_(-3, 3).npu().half()
	@@ -1249,7 +1249,7 @@ class TestOnnxOps(TestCase):
1249	epsilon = 1e-6	1249	epsilon = 1e-6
1250	x = torch_npu.npu_add_rms_norm(x1, x2, gamma, epsilon)	1250	x = torch_npu.npu_add_rms_norm(x1, x2, gamma, epsilon)
1251	return x	1251	return x
1252	-	1252	+
1253	def export_onnx(onnx_model_name):	1253	def export_onnx(onnx_model_name):
1254	x1 = torch.rand(10, 1024).uniform_(-3, 3).npu().half()	1254	x1 = torch.rand(10, 1024).uniform_(-3, 3).npu().half()
1255	x2 = torch.rand(10, 1024).uniform_(-3, 3).npu().half()	1255	x2 = torch.rand(10, 1024).uniform_(-3, 3).npu().half()
	@@ -1318,7 +1318,7 @@ class TestOnnxOps(TestCase):
1318	def forward(self, input_dummy, smooth_scales_dummy):	1318	def forward(self, input_dummy, smooth_scales_dummy):
1319	output, scale = torch_npu.npu_dynamic_quant(input_dummy, smooth_scales=smooth_scales_dummy)	1319	output, scale = torch_npu.npu_dynamic_quant(input_dummy, smooth_scales=smooth_scales_dummy)
1320	return output, scale	1320	return output, scale
1321	-	1321	+
1322	def export_onnx(onnx_model_name):	1322	def export_onnx(onnx_model_name):
1323	input_dummy = torch.rand(4, 1024, 512).uniform_(-3, 3).npu().to(torch.float16)	1323	input_dummy = torch.rand(4, 1024, 512).uniform_(-3, 3).npu().to(torch.float16)
1324	smooth_scales_dummy = torch.rand(512).uniform_(-3, 3).npu().to(torch.float16)	1324	smooth_scales_dummy = torch.rand(512).uniform_(-3, 3).npu().to(torch.float16)
	@@ -1340,7 +1340,7 @@ class TestOnnxOps(TestCase):
1340	def forward(self, input_dummy, smooth_scales_dummy, group_index_dummy):	1340	def forward(self, input_dummy, smooth_scales_dummy, group_index_dummy):
1341	output, scale = torch_npu.npu_dynamic_quant(input_dummy, smooth_scales=smooth_scales_dummy, group_index=group_index_dummy)	1341	output, scale = torch_npu.npu_dynamic_quant(input_dummy, smooth_scales=smooth_scales_dummy, group_index=group_index_dummy)
1342	return output, scale	1342	return output, scale
1343	-	1343	+
1344	def export_onnx(onnx_model_name):	1344	def export_onnx(onnx_model_name):
1345	input_dummy = torch.rand(4, 1024, 512).uniform_(-3, 3).npu().to(torch.float16)	1345	input_dummy = torch.rand(4, 1024, 512).uniform_(-3, 3).npu().to(torch.float16)
1346	group_num = 10	1346	group_num = 10
	@@ -1370,7 +1370,7 @@ class TestOnnxOps(TestCase):
1370	def forward(self, input_dummy, smooth_scales_dummy, group_index_dummy):	1370	def forward(self, input_dummy, smooth_scales_dummy, group_index_dummy):
1371	output, scale, offset = torch_npu.npu_dynamic_quant_asymmetric(input_dummy, smooth_scales=smooth_scales_dummy, group_index=group_index_dummy)	1371	output, scale, offset = torch_npu.npu_dynamic_quant_asymmetric(input_dummy, smooth_scales=smooth_scales_dummy, group_index=group_index_dummy)
1372	return output, scale, offset	1372	return output, scale, offset
1373	-	1373	+
1374	def export_onnx(onnx_model_name):	1374	def export_onnx(onnx_model_name):
1375	input_dummy = torch.rand(4, 1024, 512).uniform_(-3, 3).npu().to(torch.float16)	1375	input_dummy = torch.rand(4, 1024, 512).uniform_(-3, 3).npu().to(torch.float16)
1376	group_num = 10	1376	group_num = 10
	@@ -1436,7 +1436,7 @@ class TestOnnxOps(TestCase):
1436	assert (os.path.isfile(os.path.join(TestOnnxOps.test_onnx_path, onnx_model_name)))	1436	assert (os.path.isfile(os.path.join(TestOnnxOps.test_onnx_path, onnx_model_name)))
1437		1437
1438	@unittest.skip("skip now")	1438	@unittest.skip("skip now")
1439	- def test_wrapper_npu_quantize(self):	1439	+ def test_wrapper_npu_quantize(self):
1440	class Model(torch.nn.Module):	1440	class Model(torch.nn.Module):
1441	def __init__(self):	1441	def __init__(self):
1442	super().__init__()	1442	super().__init__()
	@@ -1461,7 +1461,7 @@ class TestOnnxOps(TestCase):
1461		1461
1462	@unittest.skip("skip now")	1462	@unittest.skip("skip now")
1463	@SupportedDevices(['Ascend910B'])	1463	@SupportedDevices(['Ascend910B'])
1464	- def test_wrapper_npu_group_quant(self):	1464	+ def test_wrapper_npu_group_quant(self):
1465	class Model(torch.nn.Module):	1465	class Model(torch.nn.Module):
1466	def __init__(self):	1466	def __init__(self):
1467	super().__init__()	1467	super().__init__()
	@@ -1488,7 +1488,7 @@ class TestOnnxOps(TestCase):
1488		1488
1489	@unittest.skip("skip now")	1489	@unittest.skip("skip now")
1490	@SupportedDevices(['Ascend910B'])	1490	@SupportedDevices(['Ascend910B'])
1491	- def test_wrapper_npu_moe_finalize_routing(self):	1491	+ def test_wrapper_npu_moe_finalize_routing(self):
1492	class Model(torch.nn.Module):	1492	class Model(torch.nn.Module):
1493	def __init__(self):	1493	def __init__(self):
1494	super().__init__()	1494	super().__init__()
	@@ -1521,7 +1521,7 @@ class TestOnnxOps(TestCase):
1521		1521
1522	@unittest.skip("skip now")	1522	@unittest.skip("skip now")
1523	@SupportedDevices(['Ascend910B'])	1523	@SupportedDevices(['Ascend910B'])
1524	- def test_wrapper_npu_moe_finalize_routing_v2(self):	1524	+ def test_wrapper_npu_moe_finalize_routing_v2(self):
1525	class Model(torch.nn.Module):	1525	class Model(torch.nn.Module):
1526	def __init__(self):	1526	def __init__(self):
1527	super().__init__()	1527	super().__init__()
	@@ -1529,7 +1529,7 @@ class TestOnnxOps(TestCase):
1529	def forward(self, expanded_permuted_rows, skip1, skip2_optional, bias, scales,	1529	def forward(self, expanded_permuted_rows, skip1, skip2_optional, bias, scales,
1530	expanded_src_to_dst_row, expert_for_source_row):	1530	expanded_src_to_dst_row, expert_for_source_row):
1531	return torch_npu.npu_moe_finalize_routing(expanded_permuted_rows, skip1, skip2_optional,	1531	return torch_npu.npu_moe_finalize_routing(expanded_permuted_rows, skip1, skip2_optional,
1532	- bias, scales, expanded_src_to_dst_row,	1532	+ bias, scales, expanded_src_to_dst_row,
1533	expert_for_source_row, drop_pad_mode=1)	1533	expert_for_source_row, drop_pad_mode=1)
1534		1534
1535	def export_onnx(onnx_model_name):	1535	def export_onnx(onnx_model_name):
	@@ -1566,7 +1566,7 @@ class TestOnnxOps(TestCase):
1566	return y	1566	return y
1567		1567
1568	def export_onnx(onnx_model_name):	1568	def export_onnx(onnx_model_name):
1569	- x = torch.tensor([[1, 2, 3, 4, 5], [6, 7, 8, 9, 10]],	1569	+ x = torch.tensor([[1, 2, 3, 4, 5], [6, 7, 8, 9, 10]],
1570	dtype=torch.float16).npu()	1570	dtype=torch.float16).npu()
1571	model = Model().to("npu")	1571	model = Model().to("npu")
1572	model(x)	1572	model(x)

Mtest/optim/test_fused_optimizers.py+3-3

	@@ -103,7 +103,7 @@ class TestFusedOptim(TestCase):
103	if p.grad is not None:	103	if p.grad is not None:
104	self.assertEqual(p.grad, p_clone.grad)	104	self.assertEqual(p.grad, p_clone.grad)
105	self.assertEqual(p.grad, torch.zeros_like(p.grad))	105	self.assertEqual(p.grad, torch.zeros_like(p.grad))
106	-	106	+
107	def test_step(self):	107	def test_step(self):
108	optim_cases = self._create_optimizer_cases(all_cases=True)	108	optim_cases = self._create_optimizer_cases(all_cases=True)
109	num_iters = 10	109	num_iters = 10
	@@ -160,7 +160,7 @@ class TestFusedOptim(TestCase):
160	for p in m.parameters():	160	for p in m.parameters():
161	if p.grad is not None:	161	if p.grad is not None:
162	self.assertEqual(grads_before_unscale[p] / 128, p.grad)	162	self.assertEqual(grads_before_unscale[p] / 128, p.grad)
163	-	163	+
164	@unittest.skip("Temporarily skipping")	164	@unittest.skip("Temporarily skipping")
165	def test_simple_model_train_dynamic(self):	165	def test_simple_model_train_dynamic(self):
166	model = self._create_simple_model()	166	model = self._create_simple_model()
	@@ -191,7 +191,7 @@ class TestFusedOptim(TestCase):
191	scaler_fused.step(opt_fused)	191	scaler_fused.step(opt_fused)
192	scaler_fused.update()	192	scaler_fused.update()
193	self.assertRtolEqual(loss, loss_fused)	193	self.assertRtolEqual(loss, loss_fused)
194	-	194	+
195	@unittest.skip("Temporarily skipping")	195	@unittest.skip("Temporarily skipping")
196	def test_simple_model_train_static(self):	196	def test_simple_model_train_static(self):
197	model = self._create_simple_model()	197	model = self._create_simple_model()

Mtest/profiler/analysis/prof_bean/test_op_mark_bean.py+1-1

	@@ -12,7 +12,7 @@ class TestOpMarkBean(TestCase):
12	def setUpClass(cls):	12	def setUpClass(cls):
13	super().setUpClass()	13	super().setUpClass()
14	cls.samples = cls.generate_samples()	14	cls.samples = cls.generate_samples()
15	-	15	+
16		16
17	@classmethod	17	@classmethod
18	def generate_samples(cls):	18	def generate_samples(cls):

Mtest/profiler/analysis/prof_common_func/test_constant.py+1-1

	@@ -18,7 +18,7 @@ class TestConstant(TestCase):
18	with self.assertRaises((RuntimeError, TypeError)):	18	with self.assertRaises((RuntimeError, TypeError)):
19	convert_ns2us_float(18789.89)	19	convert_ns2us_float(18789.89)
20	self.assertEqual(str(convert_ns2us_float(1459635878536856678)), str(1459635878536856678 / 1000))	20	self.assertEqual(str(convert_ns2us_float(1459635878536856678)), str(1459635878536856678 / 1000))
21	-	21	+
22	def test_convert_ns2us_str(self):	22	def test_convert_ns2us_str(self):
23	self.assertEqual(convert_ns2us_str(float("inf")), "inf")	23	self.assertEqual(convert_ns2us_str(float("inf")), "inf")
24	with self.assertRaises((RuntimeError, TypeError)):	24	with self.assertRaises((RuntimeError, TypeError)):

Mtest/profiler/analysis/prof_common_func/test_trace_event_manager.py+1-1

	@@ -87,7 +87,7 @@ class TestTraceEventManager(TestCase):
87	"tid": 444, "ts": "4.000", "cat": "fwdbwd"}	87	"tid": 444, "ts": "4.000", "cat": "fwdbwd"}
88	]	88	]
89	self.assertEqual(expect, TraceEventManager.create_fwd_flow(events))	89	self.assertEqual(expect, TraceEventManager.create_fwd_flow(events))
90	-	90	+
91	def test_python_event(self):	91	def test_python_event(self):
92	process_id = random.randint(1, 2**64 - 1)	92	process_id = random.randint(1, 2**64 - 1)
93	thread_id = random.randint(1, 2**64 - 1)	93	thread_id = random.randint(1, 2**64 - 1)

Mtest/profiler/test_profiler_tree.py+0-1

	@@ -733,4 +733,3 @@ class TestProfilerTree(TestCase):
733		733
734	if __name__ == "__main__":	734	if __name__ == "__main__":
735	run_tests()	735	run_tests()
736	-

Mtest/sample_torch_npu_run_store.py+2-2

	@@ -16,7 +16,7 @@ class TorchNpuRunStoreSample:
16	self._init_method = f'parallel://{ip}:{port}'	16	self._init_method = f'parallel://{ip}:{port}'
17	self._timeout = timedelta(minutes=1)	17	self._timeout = timedelta(minutes=1)
18	self._key = 'sample_torch_npu_run_store:test_case_001'	18	self._key = 'sample_torch_npu_run_store:test_case_001'
19	-	19	+
20	rendezvous_iterator = rendezvous(	20	rendezvous_iterator = rendezvous(
21	self._init_method, self._current_rank, self._world_size, timeout=self._timeout	21	self._init_method, self._current_rank, self._world_size, timeout=self._timeout
22	)	22	)
	@@ -41,7 +41,7 @@ class TorchNpuRunStoreSample:
41	if timedelta(seconds=(time.time() - start_time)) > self._timeout:	41	if timedelta(seconds=(time.time() - start_time)) > self._timeout:
42	timestamp = datetime.now().strftime('%Y-%m-%d %H:%M:%S.%f')	42	timestamp = datetime.now().strftime('%Y-%m-%d %H:%M:%S.%f')
43	raise RuntimeError(f'[{timestamp}]rank: {self._current_rank} wait all workers ready timeout')	43	raise RuntimeError(f'[{timestamp}]rank: {self._current_rank} wait all workers ready timeout')
44	-	44	+
45	timestamp = datetime.now().strftime('%Y-%m-%d %H:%M:%S.%f')	45	timestamp = datetime.now().strftime('%Y-%m-%d %H:%M:%S.%f')
46	print(f'[{timestamp}] Rank: {self._current_rank} complete store-based barrier for worker count: {alive_count}')	46	print(f'[{timestamp}] Rank: {self._current_rank} complete store-based barrier for worker count: {alive_count}')
47		47

Mtest/test_autocast.py+2-2

	@@ -289,7 +289,7 @@ class TestAutocastNPUfp32(TestCase):
289	with torch.autocast(device_type=device, dtype=torch.float32):	289	with torch.autocast(device_type=device, dtype=torch.float32):
290	b = torch.mm(a, a)	290	b = torch.mm(a, a)
291	self.assertEqual(b.dtype, torch.float32)	291	self.assertEqual(b.dtype, torch.float32)
292	-	292	+
293	def test_autocast_fp32_when_origin_dtype_is_bfloat16(self):	293	def test_autocast_fp32_when_origin_dtype_is_bfloat16(self):
294	device = "npu"	294	device = "npu"
295	a = torch.rand((8, 8), device=device, dtype=torch.bfloat16)	295	a = torch.rand((8, 8), device=device, dtype=torch.bfloat16)
	@@ -303,7 +303,7 @@ class TestAutocastNPUfp32(TestCase):
303	with torch.autocast(device_type=device, dtype=torch.float32):	303	with torch.autocast(device_type=device, dtype=torch.float32):
304	b = torch.mm(a, a)	304	b = torch.mm(a, a)
305	self.assertEqual(b.dtype, torch.float32)	305	self.assertEqual(b.dtype, torch.float32)
306	-	306	+
307	def test_autocast_fp32_when_disabled(self):	307	def test_autocast_fp32_when_disabled(self):
308	device = "npu"	308	device = "npu"
309	a = torch.rand((8, 8), device=device, dtype=torch.bfloat16)	309	a = torch.rand((8, 8), device=device, dtype=torch.bfloat16)

Mtest/test_binary_ufuncs.py+20-20

Mtest/test_fake_tensor.py+1-1

	@@ -1903,7 +1903,7 @@ class TestQuantMatmul(TestCase):
1903	expect_ret_fp16 = torch.randint(-1, 1, (1, 1, 100), dtype=torch.float16).npu()	1903	expect_ret_fp16 = torch.randint(-1, 1, (1, 1, 100), dtype=torch.float16).npu()
1904	bias_fp32 = torch.randint(-1, 1, (1, 1, 100), dtype=torch.float32).npu()	1904	bias_fp32 = torch.randint(-1, 1, (1, 1, 100), dtype=torch.float32).npu()
1905	pertoken_scale = torch.randn(1, dtype=torch.float32).npu()	1905	pertoken_scale = torch.randn(1, dtype=torch.float32).npu()
1906	- res_fp16 = torch_npu.npu_quant_matmul(x1, x2, scale, offset=None, pertoken_scale=pertoken_scale,	1906	+ res_fp16 = torch_npu.npu_quant_matmul(x1, x2, scale, offset=None, pertoken_scale=pertoken_scale,
1907	bias=bias_fp32, output_dtype=torch.float16)	1907	bias=bias_fp32, output_dtype=torch.float16)
1908	self.assertTrue(expect_ret_fp16.shape == res_fp16.shape)	1908	self.assertTrue(expect_ret_fp16.shape == res_fp16.shape)
1909	self.assertTrue(expect_ret_fp16.dtype == res_fp16.dtype)	1909	self.assertTrue(expect_ret_fp16.dtype == res_fp16.dtype)

Mtest/test_indexing.py+2-2

	@@ -128,7 +128,7 @@ class TestIndexing(TestCase):
128		128
129	self.assertRaises(TypeError, delitem)	129	self.assertRaises(TypeError, delitem)
130		130
131	-	131	+
132	@dtypes(torch.half, torch.double)	132	@dtypes(torch.half, torch.double)
133	def test_advancedindex(self, device, dtype):	133	def test_advancedindex(self, device, dtype):
134	# Tests for Integer Array Indexing, Part I - Purely integer array	134	# Tests for Integer Array Indexing, Part I - Purely integer array
	@@ -876,7 +876,7 @@ class TestIndexing(TestCase):
876		876
877	self.assertEqual(output, input_list)	877	self.assertEqual(output, input_list)
878		878
879	-	879	+
880	def test_index_ind_dtype(self, device):	880	def test_index_ind_dtype(self, device):
881	x = torch.randn(4, 4, device=device)	881	x = torch.randn(4, 4, device=device)
882	ind_long = torch.randint(4, (4,), dtype=torch.long, device=device)	882	ind_long = torch.randint(4, (4,), dtype=torch.long, device=device)

Mtest/test_jit_fuser.py+1-1

	@@ -787,7 +787,7 @@ class TestFuser(JitTestCase):
787	FileCheck.check("FusionGroup").run(str(graph))	787	FileCheck.check("FusionGroup").run(str(graph))
788	except RuntimeError as e:	788	except RuntimeError as e:
789	if 'Failed to compile' in e.args[0]:	789	if 'Failed to compile' in e.args[0]:
790	- warnings.warn('CPU fuser test has failed! This is not a hard failure, '	790	+ warnings.warn('CPU fuser test has failed! This is not a hard failure, '
791	'because the kernels sometimes trigger bugs in compilers '	791	'because the kernels sometimes trigger bugs in compilers '
792	'(most notably GCC 7.2).')	792	'(most notably GCC 7.2).')
793	raise unittest.SkipTest('Failed to compile') from e	793	raise unittest.SkipTest('Failed to compile') from e

Mtest/test_multiprocessing_api.py+15-15

Mtest/test_nestedtensor.py+5-5

Mtest/test_npu_pinned_memory_background_threads.py+1-1

	@@ -38,4 +38,4 @@ class TestPinnedMemoryBackgroundThreads(TestCase):
38	self.copy_tensor(ITERS)	38	self.copy_tensor(ITERS)
39		39
40	if __name__ == '__main__':	40	if __name__ == '__main__':
41	- run_tests()	41	+ run_tests()

Mtest/test_npu_sanitizer.py+2-2

	@@ -160,7 +160,7 @@ class TestArgumentHandler(TestCase):
160		160
161	self.assertEqual({out.data_ptr()}, argument_handler.dataptrs_written)	161	self.assertEqual({out.data_ptr()}, argument_handler.dataptrs_written)
162	self.assertEqual({out.data_ptr()}, argument_handler.outputs)	162	self.assertEqual({out.data_ptr()}, argument_handler.outputs)
163	-	163	+
164	def test_equal_reads_inputs_but_no_tensor_output_written(self):	164	def test_equal_reads_inputs_but_no_tensor_output_written(self):
165	"""Data-reading op with non-tensor output should not record tensor writes."""	165	"""Data-reading op with non-tensor output should not record tensor writes."""
166	equal_func = torch.ops.aten.equal.default	166	equal_func = torch.ops.aten.equal.default
	@@ -177,7 +177,7 @@ class TestArgumentHandler(TestCase):
177	self.assertEqual(set(), argument_handler.outputs)	177	self.assertEqual(set(), argument_handler.outputs)
178	self.assertTrue(isinstance(out, bool))	178	self.assertTrue(isinstance(out, bool))
179		179
180	-	180	+
181	class TestRecordStreamHandler(TestCase):	181	class TestRecordStreamHandler(TestCase):
182	def test_erase_stream_removes_recorded_stream(self):	182	def test_erase_stream_removes_recorded_stream(self):
183	"""Communication eraseStream should clear the matching recorded stream."""	183	"""Communication eraseStream should clear the matching recorded stream."""

Mtest/test_reductions.py+16-16

Mtest/test_sort_and_select.py+2-2

	@@ -6,7 +6,7 @@ import numpy as np
6	import torch	6	import torch
7	from torch import nan	7	from torch import nan
8	from torch.testing import make_tensor	8	from torch.testing import make_tensor
9	-from torch.testing._internal.common_dtype import (all_types, all_types_and, floating_types_and,	9	+from torch.testing._internal.common_dtype import (all_types, all_types_and, floating_types_and,
10	integral_types, _dispatch_dtypes)	10	integral_types, _dispatch_dtypes)
11	from torch.testing._internal.common_utils import (TestCase, run_tests, slowTest, skipIfTorchDynamo)	11	from torch.testing._internal.common_utils import (TestCase, run_tests, slowTest, skipIfTorchDynamo)
12	from torch.testing._internal.common_device_type import \	12	from torch.testing._internal.common_device_type import \
	@@ -17,7 +17,7 @@ import torch_npu
17	import torch_npu.testing	17	import torch_npu.testing
18		18
19	SIZE = 100	19	SIZE = 100
20	-all_types_without_double = _dispatch_dtypes((torch.half, torch.float32, torch.uint8,	20	+all_types_without_double = _dispatch_dtypes((torch.half, torch.float32, torch.uint8,
21	torch.int8, torch.int16, torch.int32, torch.int64))	21	torch.int8, torch.int16, torch.int32, torch.int64))
22		22
23		23

Mtest/test_torch_npu_init.py+1-1

	@@ -453,7 +453,7 @@ class TestTorchNpuBootstrap(TestCase):
453	)	453	)
454	"""	454	"""
455	)	455	)
456	-	456	+
457	def test_08_top_level_unsupported_dtype_compatibility(self):	457	def test_08_top_level_unsupported_dtype_compatibility(self):
458	self._run_python(	458	self._run_python(
459	"""	459	"""

Mtest/trans_contiguous/test_tensorto_preserve_format.py+7-7

Mtest/utils/test_torch_npu_logs_filter.py+1-1

	@@ -19,7 +19,7 @@ class TestTorchNpuLogs(TestCase):
19	os.environ['TORCH_NPU_LOGS'] = self.original_torch_npu_logs	19	os.environ['TORCH_NPU_LOGS'] = self.original_torch_npu_logs
20	else:	20	else:
21	del os.environ['TORCH_NPU_LOGS']	21	del os.environ['TORCH_NPU_LOGS']
22	-	22	+
23	if self.original_torch_npu_logs_filter is not None:	23	if self.original_torch_npu_logs_filter is not None:
24	os.environ['TORCH_NPU_LOGS_FILTER'] = self.original_torch_npu_logs_filter	24	os.environ['TORCH_NPU_LOGS_FILTER'] = self.original_torch_npu_logs_filter
25	else:	25	else:

Mtools/flight_recorder/components/types.py+1-1

	@@ -354,7 +354,7 @@ class Op:
354	def __init__(self, event: dict[Any, Any], memberships: dict[str, set[Any]], pg_name: str):	354	def __init__(self, event: dict[Any, Any], memberships: dict[str, set[Any]], pg_name: str):
355		355
356	frames = event.get("frames")	356	frames = event.get("frames")
357	- if not frames:	357	+ if not frames:
358	raise ValueError(self.MISSING_FRAMES_ERR)	358	raise ValueError(self.MISSING_FRAMES_ERR)
359	first_frame = frames[0] if len(frames) > 0 else None	359	first_frame = frames[0] if len(frames) > 0 else None
360	if not first_frame:	360	if not first_frame:

Mtools/flight_recorder/components/utils.py+4-4

Mtorch_npu/__init__.py+1-1

	@@ -58,7 +58,7 @@ def _initialize():
58	# 5. final extension barrier and shutdown hook	58	# 5. final extension barrier and shutdown hook
59	_initialize_runtime_lifecycle()	59	_initialize_runtime_lifecycle()
60		60
61	- # 6. optional runtime features	61	+ # 6. optional runtime features
62	_enable_optional_features()	62	_enable_optional_features()
63		63
64		64

Mtorch_npu/_inductor/__init__.py+3-3

	@@ -34,7 +34,7 @@ def _load_mlir_backend():
34	def _load_dvm_backend():	34	def _load_dvm_backend():
35	from .ascend_npu_ir.ascend_npu_ir.npu import npu_inductor_plugin	35	from .ascend_npu_ir.ascend_npu_ir.npu import npu_inductor_plugin
36	from .dvm import mlir_fusion	36	from .dvm import mlir_fusion
37	-	37	+
38	def _load_triton_backend():	38	def _load_triton_backend():
39	import torch	39	import torch
40	import os	40	import os
	@@ -154,8 +154,8 @@ _BACKEND_LOADERS = {
154	"dvm": _load_dvm_backend,	154	"dvm": _load_dvm_backend,
155	"default": _load_triton_backend,	155	"default": _load_triton_backend,
156	}	156	}
157	-	157	+
158	-	158	+
159	def _load_backend():	159	def _load_backend():
160		160
161	backend = _get_backend()	161	backend = _get_backend()

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/codecache.py+8-8

	@@ -55,7 +55,7 @@ def codegen_subgraph_dump(inds, shapes, strides, dtypes, inds2):
55	codes.append(f'args = new_args')	55	codes.append(f'args = new_args')
56	return '\n'.join(codes)	56	return '\n'.join(codes)
57		57
58	-	58	+
59	def _worker_compile(	59	def _worker_compile(
60	kernel, cc: int, device: torch.device, logger_level=None, extra_env=None	60	kernel, cc: int, device: torch.device, logger_level=None, extra_env=None
61	) -> None:	61	) -> None:
	@@ -79,10 +79,10 @@ def _akg_worker_compile(
79		79
80		80
81	def _load_kernel(	81	def _load_kernel(
82	- kernel_name: str,	82	+ kernel_name: str,
83	- source_code: str,	83	+ source_code: str,
84	- no_more_compile=False,	84	+ no_more_compile=False,
85	- suppress_error=False,	85	+ suppress_error=False,
86	kernel_meta=None,	86	kernel_meta=None,
87	extra_env=None) -> ModuleType:	87	extra_env=None) -> ModuleType:
88	if os.getenv("TORCHINDUCTOR_USE_AKG", "0") == "1":	88	if os.getenv("TORCHINDUCTOR_USE_AKG", "0") == "1":
	@@ -144,7 +144,7 @@ class MulitprocessCompileFuture(CodeCacheFuture):
144	errors.append(e)	144	errors.append(e)
145		145
146	if len(errors) < len(self.futures):	146	if len(errors) < len(self.futures):
147	- kernel = self.kernel = _load_kernel(self.kernel_name, self.source_code,	147	+ kernel = self.kernel = _load_kernel(self.kernel_name, self.source_code,
148	no_more_compile=True, suppress_error=True,	148	no_more_compile=True, suppress_error=True,
149	kernel_meta=self.kernel_meta, extra_env=self.extra_env)	149	kernel_meta=self.kernel_meta, extra_env=self.extra_env)
150	elif self.kernel_meta.get('num_outputs', 0): # All compiles fail and auto fallback	150	elif self.kernel_meta.get('num_outputs', 0): # All compiles fail and auto fallback
	@@ -223,7 +223,7 @@ class CustomAsyncCompile(AsyncCompile):
223	pool.ready_future = pool.submit(AsyncCompile._get_ready) # type: ignore[attr-defined]	223	pool.ready_future = pool.submit(AsyncCompile._get_ready) # type: ignore[attr-defined]
224	_pool_set.add(pool)	224	_pool_set.add(pool)
225	return pool	225	return pool
226	-	226	+
227	def mlir(	227	def mlir(
228	self, kernel_name: str, source_code: str, device_str: str = "npu"	228	self, kernel_name: str, source_code: str, device_str: str = "npu"
229	) -> Union[NPUTritonFuture, ModuleType]:	229	) -> Union[NPUTritonFuture, ModuleType]:
	@@ -241,7 +241,7 @@ class CustomAsyncCompile(AsyncCompile):
241	return NPUTritonFuture(kernel_name, source_code, future)	241	return NPUTritonFuture(kernel_name, source_code, future)
242	else:	242	else:
243	return _load_kernel(kernel_name, source_code)	243	return _load_kernel(kernel_name, source_code)
244	-	244	+
245	def mlir_auto_fallback(	245	def mlir_auto_fallback(
246	self, kernel_name: str, source_code: str, kernel_meta: Dict[str, Any]) -> Callable:	246	self, kernel_name: str, source_code: str, kernel_meta: Dict[str, Any]) -> Callable:
247	_compile_start()	247	_compile_start()

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/config.py+2-2

@@ -44,7 +44,7 @@ def parse_rtol_atol(env_str: str):

44 rtol, atol = None, None 44 rtol, atol = None, None

45 if not env_str.strip(): 45 if not env_str.strip():

46 return rtol, atol 46 return rtol, atol

47 - 47 +

48 parts = [p.strip() for p in env_str.split(",") if p.strip()] 48 parts = [p.strip() for p in env_str.split(",") if p.strip()]

49 for part in parts: 49 for part in parts:

50 match = re.match(r"^(rtol|atol)\s*=s\*([0-9.eE+-]+)$", part, re.IGNORECASE) 50 match = re.match(r"^(rtol|atol)\s*=s\*([0-9.eE+-]+)$", part, re.IGNORECASE)

atomgit-bot5月21日

🟡 Medium Priority

Line 50 正则为 r"^(rtol|atol)\s*=s\*([0-9.eE+-]+)$"。=s\* 匹配字面量 = + 零个或多个 s + 字面量 *。前置已有 \s* 允许等号前空格，但等号后空格不被允许，导致 rtol = 1e-5 这类带空格的合法输入被拒绝，与 warning 提示的格式示例 rtol=1e-6 虽一致，但与 \s* 的设计意图矛盾。极大概率为 =\s* 的笔误（反斜杠错位）。

改动建议

50	- match = re.match(r"^(rtol\|atol)\s=s\([0-9.eE+-]+)$", part, re.IGNORECASE)
50	+ match = re.match(r"^(rtol\|atol)\s=\s([0-9.eE+-]+)$", part, re.IGNORECASE)

应用建议

不准确？

Jingwei Huang

5月21日评论：

@@ -54,7 +54,7 @@ def parse_rtol_atol(env_str: str):

54 f"It should be like 'rtol=1e-6,atol=1e-5'. " 54 f"It should be like 'rtol=1e-6,atol=1e-5'. "

55 ) 55 )

56 continue 56 continue

57 - 57 +

58 key, value_str = match.groups() 58 key, value_str = match.groups()

59 try: 59 try:

60 value = float(value_str) 60 value = float(value_str)

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/codegen/akg.py+1-1

	@@ -15,7 +15,7 @@ class AkgKernel(NpuMetaKernel):
15	def call_kernel(self, name: str, node=None):	15	def call_kernel(self, name: str, node=None):
16	wrapper = V.graph.wrapper_code	16	wrapper = V.graph.wrapper_code
17	call_args = self.get_call_args()	17	call_args = self.get_call_args()
18	-	18	+
19	if len(call_args) > 0:	19	if len(call_args) > 0:
20	wrapper.generate_kernel_call(	20	wrapper.generate_kernel_call(
21	name,	21	name,

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/codegen/cpp_wrapper.py+5-5

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/codegen/meta_kernel.py+25-25

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/inductor_patch/ir.py+1-1

	@@ -450,7 +450,7 @@ ir.View.create = _patch_view_create
450	@classmethod	450	@classmethod
451	def _patch_sliceview_create(	451	def _patch_sliceview_create(
452	cls, x, dim, start, end, step=1, clamp=True, traced_graph=None, node_name=None	452	cls, x, dim, start, end, step=1, clamp=True, traced_graph=None, node_name=None
453	-):	453	+):
454	step = sympy.expand(step)	454	step = sympy.expand(step)
455	assert isinstance(step, sympy.Expr) or step > 0	455	assert isinstance(step, sympy.Expr) or step > 0
456	try:	456	try:

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/inductor_patch/lowering.py+38-38

	@@ -17,13 +17,13 @@ from collections.abc import Iterable, Sequence
17	from typing import Any, Callable, cast, Optional, TYPE_CHECKING, TypeVar, Union	17	from typing import Any, Callable, cast, Optional, TYPE_CHECKING, TypeVar, Union
18	from typing_extensions import ParamSpec	18	from typing_extensions import ParamSpec
19	from typing import (	19	from typing import (
20	- Any,	20	+ Any,
21	- Callable,	21	+ Callable,
22	- Dict,	22	+ Dict,
23	- List,	23	+ List,
24	- Optional,	24	+ Optional,
25	- Set,	25	+ Set,
26	- Tuple,	26	+ Tuple,
27	Union,	27	Union,
28	)	28	)
29	from unittest.mock import patch	29	from unittest.mock import patch
	@@ -230,7 +230,7 @@ def create_sym_inputs(traced_graph: TracedGraph, size: List[Expr]):
230	traced_graph.sym_nodes.update({s_name: new_node})	230	traced_graph.sym_nodes.update({s_name: new_node})
231		231
232		232
233	-def process_ir_constant(inp: ExpandView) -> Union[TracedGraph, int, float]:	233	+def process_ir_constant(inp: ExpandView) -> Union[TracedGraph, int, float]:
234	skip = False	234	skip = False
235	if isinstance(inp.data, IndexingConstant):	235	if isinstance(inp.data, IndexingConstant):
236	dtype = inp.data.dtype	236	dtype = inp.data.dtype
	@@ -268,13 +268,13 @@ def fetch_graphs(inputs: Optional[List[TensorBox]]):
268	input_graphs.append(inp)	268	input_graphs.append(inp)
269	continue	269	continue
270	if isinstance(inp, ExpandView):	270	if isinstance(inp, ExpandView):
271	- inp, skip = process_ir_constant(inp)	271	+ inp, skip = process_ir_constant(inp)
272	if not skip:	272	if not skip:
273	input_graphs.append(inp)	273	input_graphs.append(inp)
274	continue	274	continue
275	name = inp.get_name()	275	name = inp.get_name()
276	traced_graph = inp.get_traced_graph()	276	traced_graph = inp.get_traced_graph()
277	- if traced_graph is not None:	277	+ if traced_graph is not None:
278	input_graphs.append(traced_graph)	278	input_graphs.append(traced_graph)
279	continue	279	continue
280	traced_graph = TracedGraph()	280	traced_graph = TracedGraph()
	@@ -308,7 +308,7 @@ def merge_traced_graphs(input_graphs: List[TracedGraph], origin_fn, node_name, *
308	exist_nodes[node.name] = new_node	308	exist_nodes[node.name] = new_node
309	if node.name in input_graph.sym_nodes:	309	if node.name in input_graph.sym_nodes:
310	new_graph.sym_nodes.update({node.name: new_node})	310	new_graph.sym_nodes.update({node.name: new_node})
311	-	311	+
312	def parse_args(input_graphs, exist_nodes):	312	def parse_args(input_graphs, exist_nodes):
313	args = []	313	args = []
314	for input_graph in input_graphs:	314	for input_graph in input_graphs:
	@@ -324,7 +324,7 @@ def merge_traced_graphs(input_graphs: List[TracedGraph], origin_fn, node_name, *
324	else:	324	else:
325	args.append(input_graph)	325	args.append(input_graph)
326	return args	326	return args
327	-	327	+
328	num_args = len(input_graphs)	328	num_args = len(input_graphs)
329		329
330	for k, v in kwargs.items():	330	for k, v in kwargs.items():
	@@ -1054,8 +1054,8 @@ def _foreach_map(subgraph, args, *kwargs):
1054		1054
1055	assert all(x is not None for x in outputs)	1055	assert all(x is not None for x in outputs)
1056	return outputs	1056	return outputs
1057	-	1057	+
1058	-	1058	+
1059	@register_lowering(torch.ops.npu.npu_dtype_cast_backward, type_promotion_kind=None)	1059	@register_lowering(torch.ops.npu.npu_dtype_cast_backward, type_promotion_kind=None)
1060	@register_lowering(torch.ops.npu.npu_dtype_cast, type_promotion_kind=None)	1060	@register_lowering(torch.ops.npu.npu_dtype_cast, type_promotion_kind=None)
1061	@register_lowering(torch.ops.npu._npu_dtype_cast_backward, type_promotion_kind=None)	1061	@register_lowering(torch.ops.npu._npu_dtype_cast_backward, type_promotion_kind=None)
	@@ -1110,7 +1110,7 @@ def to_device(x: TensorBox, device: torch.device, *, copy=False, non_blocking=Fa
1110	device = decode_device(device)	1110	device = decode_device(device)
1111	if x.get_device() == device:	1111	if x.get_device() == device:
1112	return clone(x) if copy else x	1112	return clone(x) if copy else x
1113	-	1113	+
1114	input_graphs = fetch_graphs([x, device])	1114	input_graphs = fetch_graphs([x, device])
1115	node_name = f'to_device_{next(node_id)}'	1115	node_name = f'to_device_{next(node_id)}'
1116	new_graph = merge_traced_graphs(input_graphs, aten.to.device, node_name, dtype=src_dtype, copy=copy)	1116	new_graph = merge_traced_graphs(input_graphs, aten.to.device, node_name, dtype=src_dtype, copy=copy)
	@@ -1150,7 +1150,7 @@ def register_pointwise(
1150	if override_fn_when_input_bool is not None:	1150	if override_fn_when_input_bool is not None:
1151	override_fn_when_input_bool = ops_wrapper(override_fn_when_input_bool)	1151	override_fn_when_input_bool = ops_wrapper(override_fn_when_input_bool)
1152		1152
1153	- fn = register_fn_to_aten_fn(fn, aten_fn)	1153	+ fn = register_fn_to_aten_fn(fn, aten_fn)
1154		1154
1155	fn = make_pointwise(	1155	fn = make_pointwise(
1156	fn,	1156	fn,
	@@ -1435,7 +1435,7 @@ def repeat(x, repeats):
1435	input_graphs = fetch_graphs([x, repeats])	1435	input_graphs = fetch_graphs([x, repeats])
1436	node_name = f'repeat_{next(node_id)}'	1436	node_name = f'repeat_{next(node_id)}'
1437	new_graph = merge_traced_graphs(input_graphs, aten.repeat, node_name)	1437	new_graph = merge_traced_graphs(input_graphs, aten.repeat, node_name)
1438	-	1438	+
1439	old_size = list(x.get_size())	1439	old_size = list(x.get_size())
1440	if len(repeats) > len(old_size):	1440	if len(repeats) > len(old_size):
1441	old_size = [sympy.S.One] * (len(repeats) - len(old_size)) + old_size	1441	old_size = [sympy.S.One] * (len(repeats) - len(old_size)) + old_size
	@@ -1516,7 +1516,7 @@ def slice_(x, dim=0, start=0, end=2**63, step=1, clamp=True):
1516	input_graphs = fetch_graphs([x.data])	1516	input_graphs = fetch_graphs([x.data])
1517	node_name = f'slice_{next(node_id)}'	1517	node_name = f'slice_{next(node_id)}'
1518	new_graph = merge_traced_graphs(input_graphs, aten.slice, node_name, dim=dim, start=start, end=end, step=step)	1518	new_graph = merge_traced_graphs(input_graphs, aten.slice, node_name, dim=dim, start=start, end=end, step=step)
1519	-	1519	+
1520	return TensorBox(ir.SliceView.create(x.data, dim, start, end, step, clamp=clamp, traced_graph=new_graph, node_name=node_name))	1520	return TensorBox(ir.SliceView.create(x.data, dim, start, end, step, clamp=clamp, traced_graph=new_graph, node_name=node_name))
1521		1521
1522		1522
	@@ -3253,8 +3253,8 @@ def slice_scatter(x, src, dim=0, start=None, end=None, step=1):
3253	input_graphs = fetch_graphs([x, src])	3253	input_graphs = fetch_graphs([x, src])
3254	node_name = f'slice_scatter_{next(node_id)}'	3254	node_name = f'slice_scatter_{next(node_id)}'
3255	new_graph = merge_traced_graphs(input_graphs, aten.slice_scatter, node_name, \	3255	new_graph = merge_traced_graphs(input_graphs, aten.slice_scatter, node_name, \
3256	- dim=dim,	3256	+ dim=dim,
3257	- start=start,	3257	+ start=start,
3258	end=end,	3258	end=end,
3259	step=step)	3259	step=step)
3260	x_loader = x.make_loader()	3260	x_loader = x.make_loader()
	@@ -3337,7 +3337,7 @@ def tensor(data, *, dtype=None, device=None, layout=None, pin_memory=False):
3337	input_graphs = fetch_graphs([data])	3337	input_graphs = fetch_graphs([data])
3338	node_name = f'tensor_{next(node_id)}'	3338	node_name = f'tensor_{next(node_id)}'
3339	new_graph = merge_traced_graphs(input_graphs, torch.tensor, node_name, \	3339	new_graph = merge_traced_graphs(input_graphs, torch.tensor, node_name, \
3340	- dtype=dtype,	3340	+ dtype=dtype,
3341	device='npu',	3341	device='npu',
3342	pin_memory=False)	3342	pin_memory=False)
3343	if isinstance(_unwrap(data), int):	3343	if isinstance(_unwrap(data), int):
	@@ -3496,7 +3496,7 @@ def _full(fill_value, device, dtype, size):
3496		3496
3497	def inner_fn(index):	3497	def inner_fn(index):
3498	return value_loader([])	3498	return value_loader([])
3499	-	3499	+
3500	node_name = f'full_{next(node_id)}'	3500	node_name = f'full_{next(node_id)}'
3501	# [wtd#18] Use the passed device parameter instead of hardcoded 'npu' to avoid device mismatch	3501	# [wtd#18] Use the passed device parameter instead of hardcoded 'npu' to avoid device mismatch
3502	new_graph = merge_traced_graphs([size, fill_value], aten.full.default, node_name, \	3502	new_graph = merge_traced_graphs([size, fill_value], aten.full.default, node_name, \
	@@ -3751,8 +3751,8 @@ def embedding(weight, indices, padding_idx=-1, scale_grad_by_freq=False, sparse=
3751	input_graphs = fetch_graphs([weight, indices])	3751	input_graphs = fetch_graphs([weight, indices])
3752	node_name = f'embedding_{next(node_id)}'	3752	node_name = f'embedding_{next(node_id)}'
3753	new_graph = merge_traced_graphs(input_graphs, aten.embedding, node_name, \	3753	new_graph = merge_traced_graphs(input_graphs, aten.embedding, node_name, \
3754	- padding_idx=padding_idx,	3754	+ padding_idx=padding_idx,
3755	- scale_grad_by_freq=scale_grad_by_freq,	3755	+ scale_grad_by_freq=scale_grad_by_freq,
3756	sparse=sparse)	3756	sparse=sparse)
3757		3757
3758	return Pointwise.create(	3758	return Pointwise.create(
	@@ -3933,7 +3933,7 @@ def _unsafe_index(x, indices):
3933	# We cannot have this lowering as a decomposition as it introduces	3933	# We cannot have this lowering as a decomposition as it introduces
3934	# mutation in the graph, which is bad for Aot Autograd. Aot Autograd runs dead	3934	# mutation in the graph, which is bad for Aot Autograd. Aot Autograd runs dead
3935	# code elimination and common subexpression elimination optimizations, which	3935	# code elimination and common subexpression elimination optimizations, which
3936	-# assume graphs to be side-effect free.	3936	+# assume graphs to be side-effect free.
3937	@register_lowering(aten.index_put)	3937	@register_lowering(aten.index_put)
3938	def index_put(x, indices, values, accumulate=False):	3938	def index_put(x, indices, values, accumulate=False):
3939	return index_put_impl_(	3939	return index_put_impl_(
	@@ -6059,7 +6059,7 @@ def should_not_sum(a, b, keepdims):
6059	if not keepdims:	6059	if not keepdims:
6060	for i in unique_indices:	6060	for i in unique_indices:
6061	del a[i]	6061	del a[i]
6062	- return a	6062	+ return a
6063	return a	6063	return a
6064		6064
6065	def make_reduction(reduction_type: ReductionType, override_return_dtype=None):	6065	def make_reduction(reduction_type: ReductionType, override_return_dtype=None):
	@@ -6082,15 +6082,15 @@ def make_reduction(reduction_type: ReductionType, override_return_dtype=None):
6082	node_name = f'reshape_{next(node_id)}'	6082	node_name = f'reshape_{next(node_id)}'
6083	input_graphs = fetch_graphs([x, new_size])	6083	input_graphs = fetch_graphs([x, new_size])
6084	new_graph = merge_traced_graphs(input_graphs, aten.reshape, node_name)	6084	new_graph = merge_traced_graphs(input_graphs, aten.reshape, node_name)
6085	- else:	6085	+ else:
6086	node_name = f'reduction_{next(node_id)}'	6086	node_name = f'reduction_{next(node_id)}'
6087	input_graphs = fetch_graphs([x, axis if axis is not None else list(range(len(x.get_size())))])	6087	input_graphs = fetch_graphs([x, axis if axis is not None else list(range(len(x.get_size())))])
6088	- new_graph = merge_traced_graphs(input_graphs, reduction_type_to_aten_fn[reduction_type],	6088	+ new_graph = merge_traced_graphs(input_graphs, reduction_type_to_aten_fn[reduction_type],
6089	node_name, keepdim=keepdims)	6089	node_name, keepdim=keepdims)
6090	- result = Reduction.create(reduction_type=reduction_type,	6090	+ result = Reduction.create(reduction_type=reduction_type,
6091	- input_node=x,	6091	+ input_node=x,
6092	- node_name=node_name,	6092	+ node_name=node_name,
6093	- traced_graph=new_graph,	6093	+ traced_graph=new_graph,
6094	**kwargs)	6094	**kwargs)
6095	if isinstance(	6095	if isinstance(
6096	result.data.data, # type: ignore[attr-defined]	6096	result.data.data, # type: ignore[attr-defined]
	@@ -6304,7 +6304,7 @@ def pow(a, b):
6304		6304
6305	def fn(idx):	6305	def fn(idx):
6306	return pow_recursive(loader(idx), b, a.get_dtype())	6306	return pow_recursive(loader(idx), b, a.get_dtype())
6307	-	6307	+
6308	input_graphs = fetch_graphs([a, b])	6308	input_graphs = fetch_graphs([a, b])
6309	node_name = f'pointwise_{next(node_id)}'	6309	node_name = f'pointwise_{next(node_id)}'
6310	new_graph = merge_traced_graphs(input_graphs, aten.pow, node_name)	6310	new_graph = merge_traced_graphs(input_graphs, aten.pow, node_name)
	@@ -6426,7 +6426,7 @@ def mul(a, b):
6426	return logical_and(a, b)	6426	return logical_and(a, b)
6427	else:	6427	else:
6428	fn = ops_wrapper(aten.mul.__name__)	6428	fn = ops_wrapper(aten.mul.__name__)
6429	- fn = register_fn_to_aten_fn(fn, aten.mul)	6429	+ fn = register_fn_to_aten_fn(fn, aten.mul)
6430	return make_pointwise(fn)(a, b)	6430	return make_pointwise(fn)(a, b)
6431		6431
6432		6432
	@@ -6534,13 +6534,13 @@ def split_last_continuous(lst):
6534	n = len(lst)	6534	n = len(lst)
6535	if n == 1:	6535	if n == 1:
6536	return lst, []	6536	return lst, []
6537	-	6537	+
6538	i = n - 2	6538	i = n - 2
6539	while i >= 0:	6539	while i >= 0:
6540	if lst[i] + 1 != lst[i + 1]:	6540	if lst[i] + 1 != lst[i + 1]:
6541	break	6541	break
6542	i -= 1	6542	i -= 1
6543	-	6543	+
6544	start = i + 1	6544	start = i + 1
6545	last_part = lst[start:]	6545	last_part = lst[start:]
6546	remaining = lst[:start]	6546	remaining = lst[:start]
	@@ -6564,7 +6564,7 @@ def sum_(x, axis=None, keepdims=False, *, dtype=None):
6564	fn = make_reduction("sum", override_return_dtype=dtype)	6564	fn = make_reduction("sum", override_return_dtype=dtype)
6565	r = fn(x, axis, keepdims, dtype=dtype)	6565	r = fn(x, axis, keepdims, dtype=dtype)
6566	return r	6566	return r
6567	-	6567	+
6568		6568
6569	fallback_cumsum = fallback_handler(aten.cumsum.default)	6569	fallback_cumsum = fallback_handler(aten.cumsum.default)
6570	fallback_cumprod = fallback_handler(aten.cumprod.default)	6570	fallback_cumprod = fallback_handler(aten.cumprod.default)
	@@ -6799,7 +6799,7 @@ def register_pointwise_numeric_ldf64(op):
6799	type_promotion_kind=ELEMENTWISE_TYPE_PROMOTION_KIND.INT_TO_FLOAT,	6799	type_promotion_kind=ELEMENTWISE_TYPE_PROMOTION_KIND.INT_TO_FLOAT,
6800	use_libdevice_for_f64=True,	6800	use_libdevice_for_f64=True,
6801	)	6801	)
6802	-	6802	+
6803	@register_lowering([aten.neg])	6803	@register_lowering([aten.neg])
6804	def neg(a):	6804	def neg(a):
6805	if a.get_dtype() in (torch.int32, torch.int64):	6805	if a.get_dtype() in (torch.int32, torch.int64):

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/meta_compiler.py+4-4

	@@ -100,10 +100,10 @@ class MetaCompiler:
100	shutil.rmtree(failed_subgraph_dump_path)	100	shutil.rmtree(failed_subgraph_dump_path)
101	shutil.copytree(subgraph_dump_path, failed_subgraph_dump_path)	101	shutil.copytree(subgraph_dump_path, failed_subgraph_dump_path)
102	return failed_subgraph_dump_path	102	return failed_subgraph_dump_path
103	-	103	+
104	def acc_compare_and_dump(self, args, dump_data=True, *kwargs) -> Tuple[Any, bool]:	104	def acc_compare_and_dump(self, args, dump_data=True, *kwargs) -> Tuple[Any, bool]:
105	self.register_fx_fallback(self.kernel_meta)	105	self.register_fx_fallback(self.kernel_meta)
106	-	106	+
107	output, has_acc_error = check_accuracy_mlir(	107	output, has_acc_error = check_accuracy_mlir(
108	*args,	108	*args,
109	kernel_name=self.kernel_name,	109	kernel_name=self.kernel_name,
	@@ -112,7 +112,7 @@ class MetaCompiler:
112	dynamic=self.dynamic,	112	dynamic=self.dynamic,
113	**kwargs	113	**kwargs
114	)	114	)
115	-	115	+
116	if anir_config.fx_subgraph_dump_path:	116	if anir_config.fx_subgraph_dump_path:
117	data = args	117	data = args
118	if dump_data and has_acc_error:	118	if dump_data and has_acc_error:
	@@ -122,7 +122,7 @@ class MetaCompiler:
122	torch.npu.synchronize()	122	torch.npu.synchronize()
123	self.launchers = [self.launchers[0]]	123	self.launchers = [self.launchers[0]]
124	self.is_fallback_kernels = [self.is_fallback_kernels[0]]	124	self.is_fallback_kernels = [self.is_fallback_kernels[0]]
125	-	125	+
126	return output, not has_acc_error	126	return output, not has_acc_error
127		127
128	def compile(self, args, *kwargs):	128	def compile(self, args, *kwargs):

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/mlir_compiler.py+24-24

	@@ -37,9 +37,9 @@ reinterpret_tensor = torch.ops.inductor._reinterpret_tensor
37	global_cache = set()	37	global_cache = set()
38		38
39	class NpuMlirCompiler(MetaCompiler):	39	class NpuMlirCompiler(MetaCompiler):
40	- def __init__(self,	40	+ def __init__(self,
41	- kernel_name: str = '',	41	+ kernel_name: str = '',
42	- multiprocess_compile=False,	42	+ multiprocess_compile=False,
43	no_more_compile=False,	43	no_more_compile=False,
44	kernel_meta=None,	44	kernel_meta=None,
45	autotune=True):	45	autotune=True):
	@@ -127,7 +127,7 @@ class NpuMlirCompiler(MetaCompiler):
127	logger.info(f"[bisheng-compile failed]")	127	logger.info(f"[bisheng-compile failed]")
128	logger.warning(f"Compile error msg: {e.stderr.decode('utf-8')}")	128	logger.warning(f"Compile error msg: {e.stderr.decode('utf-8')}")
129	raise e	129	raise e
130	-	130	+
131	def prepare_launch(self, need_pickle=False):	131	def prepare_launch(self, need_pickle=False):
132	def get_launch_mod(so_path):	132	def get_launch_mod(so_path):
133	spec = importlib.util.spec_from_file_location("__launcher", so_path)	133	spec = importlib.util.spec_from_file_location("__launcher", so_path)
	@@ -169,7 +169,7 @@ class NpuMlirCompiler(MetaCompiler):
169	shutil.copy(cache_mlir_path, os.path.join(anir_config.fx_subgraph_dump_path, \	169	shutil.copy(cache_mlir_path, os.path.join(anir_config.fx_subgraph_dump_path, \
170	str(self.device_index), self.kernel_name))	170	str(self.device_index), self.kernel_name))
171	return cache_mlir_path	171	return cache_mlir_path
172	-	172	+
173	def get_launch_dynamic(self, function, tiling_func, tiling_size):	173	def get_launch_dynamic(self, function, tiling_func, tiling_size):
174	block_dim = anir_config.block_dim	174	block_dim = anir_config.block_dim
175	arg_tiling_device = torch.empty((tiling_size // 8), device='npu', dtype=torch.int64)	175	arg_tiling_device = torch.empty((tiling_size // 8), device='npu', dtype=torch.int64)
	@@ -177,30 +177,30 @@ class NpuMlirCompiler(MetaCompiler):
177	def kernel_call(*args, stream=None):	177	def kernel_call(*args, stream=None):
178	self.launch(block_dim, stream, function, tiling_func, tiling_size, arg_tiling_host, arg_tiling_device, None, None, None, *args)	178	self.launch(block_dim, stream, function, tiling_func, tiling_size, arg_tiling_host, arg_tiling_device, None, None, None, *args)
179	return kernel_call	179	return kernel_call
180	-	180	+
181	def get_launch(self, function):	181	def get_launch(self, function):
182	block_dim = anir_config.block_dim	182	block_dim = anir_config.block_dim
183	def kernel_call(*args, function, stream=None):	183	def kernel_call(*args, function, stream=None):
184	self.launch(block_dim, stream, function, None, None, None, *args)	184	self.launch(block_dim, stream, function, None, None, None, *args)
185		185
186	return functools.partial(kernel_call, function=function)	186	return functools.partial(kernel_call, function=function)
187	-	187	+
188	def get_launch_func(self, cache_kernel_path):	188	def get_launch_func(self, cache_kernel_path):
189	if self.dynamic:	189	if self.dynamic:
190	- function, tiling_func, tiling_size = self.get_host_func_and_tiling_size(self.kernel_name,	190	+ function, tiling_func, tiling_size = self.get_host_func_and_tiling_size(self.kernel_name,
191	- self.kernel_name + '_tiling_function',	191	+ self.kernel_name + '_tiling_function',
192	- self.kernel_name + '_get_tiling_struct_size_function',	192	+ self.kernel_name + '_get_tiling_struct_size_function',
193	cache_kernel_path)	193	cache_kernel_path)
194	return self.get_launch_dynamic(function, tiling_func, tiling_size)	194	return self.get_launch_dynamic(function, tiling_func, tiling_size)
195	else:	195	else:
196	function = load_kernel_binary(self.kernel_name, cache_kernel_path)	196	function = load_kernel_binary(self.kernel_name, cache_kernel_path)
197	return self.get_launch(function)	197	return self.get_launch(function)
198	-	198	+
199	- def register_launcher(self,	199	+ def register_launcher(self,
200	- launcher,	200	+ launcher,
201	- kernel_path=None,	201	+ kernel_path=None,
202	- num_outputs=None,	202	+ num_outputs=None,
203	- disable_dump=False,	203	+ disable_dump=False,
204	auto_fallback=False,	204	auto_fallback=False,
205	is_fallback_kernel=False):	205	is_fallback_kernel=False):
206	if num_outputs:	206	if num_outputs:
	@@ -216,7 +216,7 @@ class NpuMlirCompiler(MetaCompiler):
216	self.fx_subgraph_dump('fallback')	216	self.fx_subgraph_dump('fallback')
217	logger.info(f"register launcher {launcher} {kernel_path} success")	217	logger.info(f"register launcher {launcher} {kernel_path} success")
218		218
219	- def compile_mlir(self,	219	+ def compile_mlir(self,
220	device_info: Tuple[Any],	220	device_info: Tuple[Any],
221	compile_args: List[Any],	221	compile_args: List[Any],
222	logger_level = None) -> Callable[..., None]:	222	logger_level = None) -> Callable[..., None]:
	@@ -237,7 +237,7 @@ class NpuMlirCompiler(MetaCompiler):
237		237
238	logger.info("Start to get cached kernel. Tiling info: " +	238	logger.info("Start to get cached kernel. Tiling info: " +
239	f"tiling_size {tiling_size} ops_reorder {ops_reorder} auto_db {auto_db}")	239	f"tiling_size {tiling_size} ops_reorder {ops_reorder} auto_db {auto_db}")
240	-	240	+
241	if cache_kernel_path is None and self.no_more_compile:	241	if cache_kernel_path is None and self.no_more_compile:
242	raise RuntimeError("Skip compile.")	242	raise RuntimeError("Skip compile.")
243		243
	@@ -248,8 +248,8 @@ class NpuMlirCompiler(MetaCompiler):
248	self.bisheng_compile(named_op_mlir_path, kernel_path, tiling_size=tiling_size,	248	self.bisheng_compile(named_op_mlir_path, kernel_path, tiling_size=tiling_size,
249	ops_reorder=ops_reorder, auto_db=auto_db,	249	ops_reorder=ops_reorder, auto_db=auto_db,
250	extra_command=anir_config.extra_command)	250	extra_command=anir_config.extra_command)
251	-	251	+
252	-	252	+
253	if self.dynamic:	253	if self.dynamic:
254	kernel_path = os.path.join(tmpdir, f"lib{tiling_kernel_name}.so")	254	kernel_path = os.path.join(tmpdir, f"lib{tiling_kernel_name}.so")
255	with open(kernel_path, "rb") as f:	255	with open(kernel_path, "rb") as f:
	@@ -303,7 +303,7 @@ class NpuMlirCompiler(MetaCompiler):
303	self.launch = getattr(mod, "launch")	303	self.launch = getattr(mod, "launch")
304	if self.dynamic:	304	if self.dynamic:
305	self.get_host_func_and_tiling_size = getattr(mod, "get_host_func_and_tiling_size")	305	self.get_host_func_and_tiling_size = getattr(mod, "get_host_func_and_tiling_size")
306	-	306	+
307	launch_func = self.get_launch_func(kernel_path)	307	launch_func = self.get_launch_func(kernel_path)
308	self.register_launcher(launch_func, kernel_path)	308	self.register_launcher(launch_func, kernel_path)
309	return True	309	return True
	@@ -324,7 +324,7 @@ class NpuMlirCompiler(MetaCompiler):
324	and self.kernel_meta.get('is_reduction', False)	324	and self.kernel_meta.get('is_reduction', False)
325	)	325	)
326		326
327	- def precompile(self,	327	+ def precompile(self,
328	device_info: Tuple[Any],	328	device_info: Tuple[Any],
329	suppress_error=False,	329	suppress_error=False,
330	logger_level=None):	330	logger_level=None):
	@@ -403,7 +403,7 @@ class NpuMlirCompiler(MetaCompiler):
403	else clone_preserve_strides(arg) for arg in args[-self.num_outputs:]]	403	else clone_preserve_strides(arg) for arg in args[-self.num_outputs:]]
404	fx_inputs = [clone_preserve_strides(arg) if isinstance(arg, torch.Tensor) else arg for arg in args[:-self.num_outputs]]	404	fx_inputs = [clone_preserve_strides(arg) if isinstance(arg, torch.Tensor) else arg for arg in args[:-self.num_outputs]]
405	fx_inputs = [inp.float() if isinstance(inp, torch.Tensor) and inp.dtype == torch.bfloat16 else inp for inp in fx_inputs]	405	fx_inputs = [inp.float() if isinstance(inp, torch.Tensor) and inp.dtype == torch.bfloat16 else inp for inp in fx_inputs]
406	-	406	+
407	fx_args = fx_inputs + fx_outputs	407	fx_args = fx_inputs + fx_outputs
408	launcher_fx(fx_args, *kwargs)	408	launcher_fx(fx_args, *kwargs)
409		409
	@@ -455,7 +455,7 @@ class NpuMlirCompiler(MetaCompiler):
455	print(f"{self.kernel_name}: Tuning accuracy failed, no valid kernels found, using fallback")	455	print(f"{self.kernel_name}: Tuning accuracy failed, no valid kernels found, using fallback")
456	timings.append([float(1.0), len(self.launchers) - 1])	456	timings.append([float(1.0), len(self.launchers) - 1])
457	return timings	457	return timings
458	-	458	+
459	def autotune_to_one_config(self, args, *kwargs):	459	def autotune_to_one_config(self, args, *kwargs):
460	if any([isinstance(arg, torch.Tensor) and not arg.is_contiguous() for arg in args]):	460	if any([isinstance(arg, torch.Tensor) and not arg.is_contiguous() for arg in args]):
461	print(f'Non contiguous args exists! Kernel name is {self.kernel_name}')	461	print(f'Non contiguous args exists! Kernel name is {self.kernel_name}')

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/npu_decomp.py+5-5

	@@ -81,7 +81,7 @@ def npu_convolution_backward(
81	[output_mask[0], output_mask[1], False],	81	[output_mask[0], output_mask[1], False],
82	)	82	)
83	return (grad_inp, grad_weight, grad_bias)	83	return (grad_inp, grad_weight, grad_bias)
84	-	84	+
85	def npu__softmax_backward_data(	85	def npu__softmax_backward_data(
86	grad_output: torch.Tensor,	86	grad_output: torch.Tensor,
87	output: torch.Tensor,	87	output: torch.Tensor,
	@@ -112,9 +112,9 @@ def npu_rms_norm(
112	output = (x * rsqrt * weight).to(dtype)	112	output = (x * rsqrt * weight).to(dtype)
113	return output, rsqrt	113	return output, rsqrt
114		114
115	-def npu_rms_norm_backward(grad_output: torch.Tensor,	115	+def npu_rms_norm_backward(grad_output: torch.Tensor,
116	- x: torch.Tensor,	116	+ x: torch.Tensor,
117	- weight: torch.Tensor,	117	+ weight: torch.Tensor,
118	rsqrt: torch.Tensor) -> Tuple[torch.Tensor, torch.Tensor]:	118	rsqrt: torch.Tensor) -> Tuple[torch.Tensor, torch.Tensor]:
119	dx = (grad_output * weight - x * rsqrt * (grad_output * weight * x * rsqrt).mean(-1, keepdim=True)) * rsqrt	119	dx = (grad_output * weight - x * rsqrt * (grad_output * weight * x * rsqrt).mean(-1, keepdim=True)) * rsqrt
120	dgamma = (grad_output * x * rsqrt).sum(0, keepdim=False)	120	dgamma = (grad_output * x * rsqrt).sum(0, keepdim=False)
	@@ -140,7 +140,7 @@ def npu_swiglu_backward(grad_output, x, dim=-1):
140	def _rotate_half(x: Tensor) -> Tensor:	140	def _rotate_half(x: Tensor) -> Tensor:
141	x1, x2 = torch.chunk(x, 2, dim=-1)	141	x1, x2 = torch.chunk(x, 2, dim=-1)
142	return torch.cat((-x2, x1), dim=-1)	142	return torch.cat((-x2, x1), dim=-1)
143	-	143	+
144	def npu_rotary_mul(t, cos_, sin_):	144	def npu_rotary_mul(t, cos_, sin_):
145	t = (t * cos_) + (_rotate_half(t) * sin_)	145	t = (t * cos_) + (_rotate_half(t) * sin_)
146	return t	146	return t

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/npu_inductor_plugin.py+8-8

	@@ -131,7 +131,7 @@ def disable_implicit_decomposition():
131	op_override.py_kernels.pop(DispatchKey.Autograd)	131	op_override.py_kernels.pop(DispatchKey.Autograd)
132	if DispatchKey.CompositeImplicitAutograd in op_override.py_kernels:	132	if DispatchKey.CompositeImplicitAutograd in op_override.py_kernels:
133	op_override.py_kernels.pop(DispatchKey.CompositeImplicitAutograd)	133	op_override.py_kernels.pop(DispatchKey.CompositeImplicitAutograd)
134	-	134	+
135		135
136	def _patch_run_node(tracer, node, args, kwargs, nnmodule):	136	def _patch_run_node(tracer, node, args, kwargs, nnmodule):
137	op = node.op	137	op = node.op
	@@ -143,7 +143,7 @@ def _patch_run_node(tracer, node, args, kwargs, nnmodule):
143		143
144	try:	144	try:
145	if op == "call_function":	145	if op == "call_function":
146	- # patch start	146	+ # patch start
147	if 'npu.npu_fusion_attention' in str(node.target):	147	if 'npu.npu_fusion_attention' in str(node.target):
148	if 'actual_seq_qlen' in kwargs:	148	if 'actual_seq_qlen' in kwargs:
149	kwargs['actual_seq_qlen'] = list(kwargs['actual_seq_qlen'])	149	kwargs['actual_seq_qlen'] = list(kwargs['actual_seq_qlen'])
	@@ -190,7 +190,7 @@ disable_implicit_decomposition()
190	torch._dynamo.utils.run_node = _patch_run_node	190	torch._dynamo.utils.run_node = _patch_run_node
191		191
192		192
193	-from torch._dynamo.backends import common	193	+from torch._dynamo.backends import common
194	from torch._dynamo.backends.common import AotAutograd	194	from torch._dynamo.backends.common import AotAutograd
195		195
196	def wrap_compiler(fn):	196	def wrap_compiler(fn):
	@@ -215,13 +215,13 @@ def wrap_aot_autograd(fn):
215		215
216	AotAutograd.__call__ = wrap_aot_autograd(AotAutograd.__call__)	216	AotAutograd.__call__ = wrap_aot_autograd(AotAutograd.__call__)
217		217
218	-# recompute last usage for inductor scheduler	218	+# recompute last usage for inductor scheduler
219	from torch._inductor import scheduler	219	from torch._inductor import scheduler
220	from torch._inductor.scheduler import (	220	from torch._inductor.scheduler import (
221	Dep,	221	Dep,
222	WeakDep,	222	WeakDep,
223	- Scheduler,	223	+ Scheduler,
224	- SchedulerNode,	224	+ SchedulerNode,
225	SchedulerBuffer,	225	SchedulerBuffer,
226	FusedSchedulerNode,	226	FusedSchedulerNode,
227	BaseSchedulerNode,	227	BaseSchedulerNode,
	@@ -432,7 +432,7 @@ def patch_transfer_to_npu():
432	_replace_cuda_to_npu_in_kwargs,	432	_replace_cuda_to_npu_in_kwargs,
433	)	433	)
434		434
435	- def new_wrapper_cuda(module, method):	435	+ def new_wrapper_cuda(module, method):
436	src_method = f"_src_{method}"	436	src_method = f"_src_{method}"
437	if hasattr(getattr(module, method), '__wrapped__'):	437	if hasattr(getattr(module, method), '__wrapped__'):
438	src_func = getattr(module, method).__wrapped__	438	src_func = getattr(module, method).__wrapped__
	@@ -441,7 +441,7 @@ def patch_transfer_to_npu():
441		441
442	setattr(module, src_method, src_func)	442	setattr(module, src_method, src_func)
443	fn = getattr(module, src_method)	443	fn = getattr(module, src_method)
444	-	444	+
445	def decorated(args, *kwargs):	445	def decorated(args, *kwargs):
446	replace_int = fn.__name__ in ['to', 'to_empty']	446	replace_int = fn.__name__ in ['to', 'to_empty']
447	if args:	447	if args:

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/npu_patch_deprecated.py+1-2

	@@ -50,7 +50,7 @@ def _patch_add_ephemeral_timeout_for_all_pgs(timeout: timedelta) -> None:
50	devices = pg._device_types	50	devices = pg._device_types
51	if torch.device("npu") in devices:	51	if torch.device("npu") in devices:
52	backend = pg._get_backend(torch.device("npu"))	52	backend = pg._get_backend(torch.device("npu"))
53	-	53	+
54	distributed_c10d._add_ephemeral_timeout_for_all_pgs = _patch_add_ephemeral_timeout_for_all_pgs	54	distributed_c10d._add_ephemeral_timeout_for_all_pgs = _patch_add_ephemeral_timeout_for_all_pgs
55		55
56	if get_anir_mode() == 'O0':	56	if get_anir_mode() == 'O0':
	@@ -58,4 +58,3 @@ if get_anir_mode() == 'O0':
58	def my_silu_backward(grad_out, self):	58	def my_silu_backward(grad_out, self):
59	# use with some caution: this is only really valid to run in the context of proxy tensor tracing	59	# use with some caution: this is only really valid to run in the context of proxy tensor tracing
60	return NotImplemented	60	return NotImplemented
61	-

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/npu_stream.py+89-89

	@@ -60,7 +60,7 @@ class StreamResgistrator:
60	NPU_EVENTS[tag] = event	60	NPU_EVENTS[tag] = event
61		61
62	def npu_set_stream(	62	def npu_set_stream(
63	- dependency: Sequence[torch.Tensor],	63	+ dependency: Sequence[torch.Tensor],
64	stream_tag: str,	64	stream_tag: str,
65	) -> List[torch.Tensor]:	65	) -> List[torch.Tensor]:
66	stream = NPU_STREAMS[stream_tag]	66	stream = NPU_STREAMS[stream_tag]
	@@ -68,7 +68,7 @@ def npu_set_stream(
68	return dependency	68	return dependency
69		69
70	def npu_set_stream_fake(	70	def npu_set_stream_fake(
71	- dependency: Sequence[torch.Tensor],	71	+ dependency: Sequence[torch.Tensor],
72	stream_tag: str,	72	stream_tag: str,
73	) -> List[torch.Tensor]:	73	) -> List[torch.Tensor]:
74	return dependency	74	return dependency
	@@ -82,7 +82,7 @@ direct_register_custom_op(
82	)	82	)
83		83
84	def npu_event_record(	84	def npu_event_record(
85	- dependency: Sequence[torch.Tensor],	85	+ dependency: Sequence[torch.Tensor],
86	event_tag: str,	86	event_tag: str,
87	stream_tag: str	87	stream_tag: str
88	) -> List[torch.Tensor]:	88	) -> List[torch.Tensor]:
	@@ -92,7 +92,7 @@ def npu_event_record(
92	return dependency	92	return dependency
93		93
94	def npu_event_record_fake(	94	def npu_event_record_fake(
95	- dependency: Sequence[torch.Tensor],	95	+ dependency: Sequence[torch.Tensor],
96	event_tag: str,	96	event_tag: str,
97	stream_tag: str	97	stream_tag: str
98	) -> List[torch.Tensor]:	98	) -> List[torch.Tensor]:
	@@ -107,7 +107,7 @@ direct_register_custom_op(
107	)	107	)
108		108
109	def npu_event_wait(	109	def npu_event_wait(
110	- dependency: Sequence[torch.Tensor],	110	+ dependency: Sequence[torch.Tensor],
111	event_tag: str,	111	event_tag: str,
112	) -> List[torch.Tensor]:	112	) -> List[torch.Tensor]:
113	event = NPU_EVENTS[event_tag]	113	event = NPU_EVENTS[event_tag]
	@@ -115,7 +115,7 @@ def npu_event_wait(
115	return dependency	115	return dependency
116		116
117	def npu_event_wait_fake(	117	def npu_event_wait_fake(
118	- dependency: Sequence[torch.Tensor],	118	+ dependency: Sequence[torch.Tensor],
119	event_tag: str,	119	event_tag: str,
120	) -> List[torch.Tensor]:	120	) -> List[torch.Tensor]:
121	return dependency	121	return dependency
	@@ -129,12 +129,12 @@ direct_register_custom_op(
129	)	129	)
130		130
131	def graph_break(	131	def graph_break(
132	- dependency: Sequence[torch.Tensor],	132	+ dependency: Sequence[torch.Tensor],
133	) -> List[torch.Tensor]:	133	) -> List[torch.Tensor]:
134	return dependency	134	return dependency
135		135
136	def graph_break_fake(	136	def graph_break_fake(
137	- dependency: Sequence[torch.Tensor],	137	+ dependency: Sequence[torch.Tensor],
138	) -> List[torch.Tensor]:	138	) -> List[torch.Tensor]:
139	return dependency	139	return dependency
140		140
	@@ -150,7 +150,7 @@ direct_register_custom_op(
150	)	150	)
151		151
152	def npu_wait_stream(	152	def npu_wait_stream(
153	- dependency: Sequence[torch.Tensor],	153	+ dependency: Sequence[torch.Tensor],
154	stream1_tag: str,	154	stream1_tag: str,
155	stream2_tag: str,	155	stream2_tag: str,
156	) -> List[torch.Tensor]:	156	) -> List[torch.Tensor]:
	@@ -160,7 +160,7 @@ def npu_wait_stream(
160	return dependency	160	return dependency
161		161
162	def npu_wait_stream_fake(	162	def npu_wait_stream_fake(
163	- dependency: Sequence[torch.Tensor],	163	+ dependency: Sequence[torch.Tensor],
164	stream1_tag: str,	164	stream1_tag: str,
165	stream2_tag: str,	165	stream2_tag: str,
166	) -> List[torch.Tensor]:	166	) -> List[torch.Tensor]:
	@@ -186,48 +186,48 @@ def graph_break(
186	inductor_npu_lib = Library("inductor_npu", "FRAGMENT") # noqa	186	inductor_npu_lib = Library("inductor_npu", "FRAGMENT") # noqa
187		187
188	def npu_fusion_attention(	188	def npu_fusion_attention(
189	- query: torch.Tensor,	189	+ query: torch.Tensor,
190	- key: torch.Tensor,	190	+ key: torch.Tensor,
191	- value: torch.Tensor,	191	+ value: torch.Tensor,
192	- head_num: int,	192	+ head_num: int,
193	- input_layout: str,	193	+ input_layout: str,
194	pse: Optional[torch.Tensor] = None,	194	pse: Optional[torch.Tensor] = None,
195	padding_mask: Optional[torch.Tensor] = None,	195	padding_mask: Optional[torch.Tensor] = None,
196	atten_mask: Optional[torch.Tensor] = None,	196	atten_mask: Optional[torch.Tensor] = None,
197	- scale: float = 1.0,	197	+ scale: float = 1.0,
198	- keep_prob: float = 1.0,	198	+ keep_prob: float = 1.0,
199	- pre_tockens: int = 2147483647,	199	+ pre_tockens: int = 2147483647,
200	next_tockens: int = 2147483647,	200	next_tockens: int = 2147483647,
201	- inner_precise: int = 0,	201	+ inner_precise: int = 0,
202	- prefix: Optional[torch.Tensor] = None,	202	+ prefix: Optional[torch.Tensor] = None,
203	- actual_seq_qlen: Optional[torch.Tensor] = None,	203	+ actual_seq_qlen: Optional[torch.Tensor] = None,
204	- actual_seq_kvlen: Optional[torch.Tensor] = None,	204	+ actual_seq_kvlen: Optional[torch.Tensor] = None,
205	sparse_mode: int = 0,	205	sparse_mode: int = 0,
206	- gen_mask_parallel: bool = True,	206	+ gen_mask_parallel: bool = True,
207	sync: bool = False	207	sync: bool = False
208	) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:	208	) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
209	prefix = prefix.tolist() if prefix is not None else prefix	209	prefix = prefix.tolist() if prefix is not None else prefix
210	actual_seq_qlen = actual_seq_qlen.tolist() if actual_seq_qlen is not None else actual_seq_qlen	210	actual_seq_qlen = actual_seq_qlen.tolist() if actual_seq_qlen is not None else actual_seq_qlen
211	actual_seq_kvlen = actual_seq_kvlen.tolist() if actual_seq_kvlen is not None else actual_seq_kvlen	211	actual_seq_kvlen = actual_seq_kvlen.tolist() if actual_seq_kvlen is not None else actual_seq_kvlen
212	attention_score, softmax_max, softmax_sum, softmax_out, seed, offset, numels = torch.ops.npu.npu_fusion_attention(	212	attention_score, softmax_max, softmax_sum, softmax_out, seed, offset, numels = torch.ops.npu.npu_fusion_attention(
213	- query,	213	+ query,
214	- key,	214	+ key,
215	- value,	215	+ value,
216	- head_num,	216	+ head_num,
217	- input_layout,	217	+ input_layout,
218	pse=pse,	218	pse=pse,
219	padding_mask=padding_mask,	219	padding_mask=padding_mask,
220	atten_mask=atten_mask,	220	atten_mask=atten_mask,
221	- scale=scale,	221	+ scale=scale,
222	- keep_prob=keep_prob,	222	+ keep_prob=keep_prob,
223	- pre_tockens=pre_tockens,	223	+ pre_tockens=pre_tockens,
224	next_tockens=next_tockens,	224	next_tockens=next_tockens,
225	- inner_precise=inner_precise,	225	+ inner_precise=inner_precise,
226	- prefix=prefix,	226	+ prefix=prefix,
227	- actual_seq_qlen=actual_seq_qlen,	227	+ actual_seq_qlen=actual_seq_qlen,
228	- actual_seq_kvlen=actual_seq_kvlen,	228	+ actual_seq_kvlen=actual_seq_kvlen,
229	sparse_mode=sparse_mode,	229	sparse_mode=sparse_mode,
230	- gen_mask_parallel=gen_mask_parallel,	230	+ gen_mask_parallel=gen_mask_parallel,
231	sync=sync	231	sync=sync
232	)	232	)
233		233
	@@ -238,24 +238,24 @@ def npu_fusion_attention(
238	return attention_score, softmax_max, softmax_sum, softmax_out, seed, offset, numels	238	return attention_score, softmax_max, softmax_sum, softmax_out, seed, offset, numels
239		239
240	def npu_fusion_attention_fake(	240	def npu_fusion_attention_fake(
241	- query: torch.Tensor,	241	+ query: torch.Tensor,
242	- key: torch.Tensor,	242	+ key: torch.Tensor,
243	- value: torch.Tensor,	243	+ value: torch.Tensor,
244	- head_num: int,	244	+ head_num: int,
245	- input_layout: str,	245	+ input_layout: str,
246	pse: Optional[torch.Tensor] = None,	246	pse: Optional[torch.Tensor] = None,
247	padding_mask: Optional[torch.Tensor] = None,	247	padding_mask: Optional[torch.Tensor] = None,
248	atten_mask: Optional[torch.Tensor] = None,	248	atten_mask: Optional[torch.Tensor] = None,
249	- scale: float = 1.0,	249	+ scale: float = 1.0,
250	- keep_prob: float = 1.0,	250	+ keep_prob: float = 1.0,
251	- pre_tockens: int = 2147483647,	251	+ pre_tockens: int = 2147483647,
252	next_tockens: int = 2147483647,	252	next_tockens: int = 2147483647,
253	- inner_precise: int = 0,	253	+ inner_precise: int = 0,
254	- prefix: Optional[torch.Tensor] = None,	254	+ prefix: Optional[torch.Tensor] = None,
255	- actual_seq_qlen: Optional[torch.Tensor] = None,	255	+ actual_seq_qlen: Optional[torch.Tensor] = None,
256	- actual_seq_kvlen: Optional[torch.Tensor] = None,	256	+ actual_seq_kvlen: Optional[torch.Tensor] = None,
257	sparse_mode: int = 0,	257	sparse_mode: int = 0,
258	- gen_mask_parallel: bool = True,	258	+ gen_mask_parallel: bool = True,
259	sync: bool = False	259	sync: bool = False
260	) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:	260	) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
261	B = query.size(0)	261	B = query.size(0)
	@@ -301,32 +301,32 @@ direct_register_custom_op(
301		301
302	def npu_fusion_attention_grad(	302	def npu_fusion_attention_grad(
303	query: torch.Tensor,	303	query: torch.Tensor,
304	- key: torch.Tensor,	304	+ key: torch.Tensor,
305	value: torch.Tensor,	305	value: torch.Tensor,
306	- dy: torch.Tensor,	306	+ dy: torch.Tensor,
307	- head_num: int,	307	+ head_num: int,
308	- input_layout: str,	308	+ input_layout: str,
309	- *,	309	+ *,
310	- pse: Optional[torch.Tensor] = None,	310	+ pse: Optional[torch.Tensor] = None,
311	- padding_mask: Optional[torch.Tensor] = None,	311	+ padding_mask: Optional[torch.Tensor] = None,
312	atten_mask: Optional[torch.Tensor] = None,	312	atten_mask: Optional[torch.Tensor] = None,
313	- softmax_max: Optional[torch.Tensor] = None,	313	+ softmax_max: Optional[torch.Tensor] = None,
314	- softmax_sum: Optional[torch.Tensor] = None,	314	+ softmax_sum: Optional[torch.Tensor] = None,
315	- softmax_in: Optional[torch.Tensor] = None,	315	+ softmax_in: Optional[torch.Tensor] = None,
316	- attention_in: Optional[torch.Tensor] = None,	316	+ attention_in: Optional[torch.Tensor] = None,
317	scale_value: float = 1.0,	317	scale_value: float = 1.0,
318	- keep_prob: float = 1.0,	318	+ keep_prob: float = 1.0,
319	- pre_tockens: int = 2147483647,	319	+ pre_tockens: int = 2147483647,
320	- next_tockens: int = 2147483647,	320	+ next_tockens: int = 2147483647,
321	- inner_precise: int = 0,	321	+ inner_precise: int = 0,
322	- seed: Optional[torch.Tensor] = None,	322	+ seed: Optional[torch.Tensor] = None,
323	offset: Optional[torch.Tensor] = None,	323	offset: Optional[torch.Tensor] = None,
324	- numels: Optional[torch.Tensor] = None,	324	+ numels: Optional[torch.Tensor] = None,
325	- prefix: Optional[torch.Tensor] = None,	325	+ prefix: Optional[torch.Tensor] = None,
326	actual_seq_qlen: Optional[torch.Tensor] = None,	326	actual_seq_qlen: Optional[torch.Tensor] = None,
327	- actual_seq_kvlen: Optional[torch.Tensor] = None,	327	+ actual_seq_kvlen: Optional[torch.Tensor] = None,
328	sparse_mode: int = 0,	328	sparse_mode: int = 0,
329	- gen_mask_parallel: bool = True,	329	+ gen_mask_parallel: bool = True,
330	sync: bool = False	330	sync: bool = False
331	) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:	331	) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
332	prefix = prefix.tolist() if prefix is not None else prefix	332	prefix = prefix.tolist() if prefix is not None else prefix
	@@ -349,32 +349,32 @@ def npu_fusion_attention_grad(
349		349
350	def npu_fusion_attention_grad_fake(	350	def npu_fusion_attention_grad_fake(
351	query: torch.Tensor,	351	query: torch.Tensor,
352	- key: torch.Tensor,	352	+ key: torch.Tensor,
353	value: torch.Tensor,	353	value: torch.Tensor,
354	- dy: torch.Tensor,	354	+ dy: torch.Tensor,
355	- head_num: int,	355	+ head_num: int,
356	- input_layout: str,	356	+ input_layout: str,
357	- *,	357	+ *,
358	- pse: Optional[torch.Tensor] = None,	358	+ pse: Optional[torch.Tensor] = None,
359	- padding_mask: Optional[torch.Tensor] = None,	359	+ padding_mask: Optional[torch.Tensor] = None,
360	atten_mask: Optional[torch.Tensor] = None,	360	atten_mask: Optional[torch.Tensor] = None,
361	- softmax_max: Optional[torch.Tensor] = None,	361	+ softmax_max: Optional[torch.Tensor] = None,
362	- softmax_sum: Optional[torch.Tensor] = None,	362	+ softmax_sum: Optional[torch.Tensor] = None,
363	- softmax_in: Optional[torch.Tensor] = None,	363	+ softmax_in: Optional[torch.Tensor] = None,
364	- attention_in: Optional[torch.Tensor] = None,	364	+ attention_in: Optional[torch.Tensor] = None,
365	scale_value: float = 1.0,	365	scale_value: float = 1.0,
366	- keep_prob: float = 1.0,	366	+ keep_prob: float = 1.0,
367	- pre_tockens: int = 2147483647,	367	+ pre_tockens: int = 2147483647,
368	- next_tockens: int = 2147483647,	368	+ next_tockens: int = 2147483647,
369	- inner_precise: int = 0,	369	+ inner_precise: int = 0,
370	- seed: Optional[torch.Tensor] = None,	370	+ seed: Optional[torch.Tensor] = None,
371	offset: Optional[torch.Tensor] = None,	371	offset: Optional[torch.Tensor] = None,
372	- numels: Optional[torch.Tensor] = None,	372	+ numels: Optional[torch.Tensor] = None,
373	prefix: Optional[torch.Tensor] = None,	373	prefix: Optional[torch.Tensor] = None,
374	actual_seq_qlen: Optional[torch.Tensor] = None,	374	actual_seq_qlen: Optional[torch.Tensor] = None,
375	- actual_seq_kvlen: Optional[torch.Tensor] = None,	375	+ actual_seq_kvlen: Optional[torch.Tensor] = None,
376	sparse_mode: int = 0,	376	sparse_mode: int = 0,
377	- gen_mask_parallel: bool = True,	377	+ gen_mask_parallel: bool = True,
378	sync: bool = False	378	sync: bool = False
379	) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:	379	) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
380	dq = torch.empty_like(query, dtype=query.dtype, device=query.device).contiguous()	380	dq = torch.empty_like(query, dtype=query.dtype, device=query.device).contiguous()
	@@ -436,7 +436,7 @@ class InductorNpuAttentionFunction(torch.autograd.Function):
436	return (	436	return (
437	grad_query, grad_key, grad_value, None, None, grad_pse, None, None, None, None, None, None, None, None, None,	437	grad_query, grad_key, grad_value, None, None, grad_pse, None, None, None, None, None, None, None, None, None,
438	None, None, None, None, None, None, None, None, None, None, None)	438	None, None, None, None, None, None, None, None, None, None, None)
439	-	439	+
440	def inductor_npu_fusion_attention(query, key, value, head_num, input_layout, pse=None, padding_mask=None,	440	def inductor_npu_fusion_attention(query, key, value, head_num, input_layout, pse=None, padding_mask=None,
441	atten_mask=None, scale=1.0, keep_prob=1.0, pre_tockens=2147483647,	441	atten_mask=None, scale=1.0, keep_prob=1.0, pre_tockens=2147483647,
442	next_tockens=2147483647,	442	next_tockens=2147483647,

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/torch_mlir_patch.py+1-1

	@@ -66,7 +66,7 @@ def _patch_import_stateless_graph(
66	range_constraints = {}	66	range_constraints = {}
67	for nd in graph.find_nodes(	67	for nd in graph.find_nodes(
68	op="placeholder"	68	op="placeholder"
69	- ):	69	+ ):
70	if isinstance(nd.meta['val'], torch.Tensor):	70	if isinstance(nd.meta['val'], torch.Tensor):
71	for s in nd.meta['val'].size():	71	for s in nd.meta['val'].size():
72	if isinstance(s, torch.SymInt):	72	if isinstance(s, torch.SymInt):

Mtorch_npu/_inductor/ascend_npu_ir/ascend_npu_ir/npu/utils.py+55-55

	@@ -68,7 +68,7 @@ def get_device_info(example_inputs) -> Union[Tuple[str, int], None]:
68	for inp in example_inputs:	68	for inp in example_inputs:
69	if isinstance(inp, torch.Tensor):	69	if isinstance(inp, torch.Tensor):
70	return inp.device, inp.device.index	70	return inp.device, inp.device.index
71	-	71	+
72	@functools.lru_cache(None)	72	@functools.lru_cache(None)
73	def _get_ascend_path() -> str:	73	def _get_ascend_path() -> str:
74	path = os.getenv("ASCEND_HOME_PATH", "")	74	path = os.getenv("ASCEND_HOME_PATH", "")
	@@ -102,11 +102,11 @@ def _build_npu_ext(obj_name: str, src_path, src_dir) -> str:
102		102
103	cc_cmd += [f"-I{py_include_dir}"]	103	cc_cmd += [f"-I{py_include_dir}"]
104	torch_npu_root = Path(torch_npu.__file__).resolve().parent	104	torch_npu_root = Path(torch_npu.__file__).resolve().parent
105	-	105	+
106	cpp_common_dir = (	106	cpp_common_dir = (
107	torch_npu_root / "include" / "torch_npu" / "csrc" / "inductor" / "mlir"	107	torch_npu_root / "include" / "torch_npu" / "csrc" / "inductor" / "mlir"
108	)	108	)
109	-	109	+
110	torch_npu_dir = torch_npu_root / "include"	110	torch_npu_dir = torch_npu_root / "include"
111	torch_npu_lib_dir = torch_npu_root / "lib"	111	torch_npu_lib_dir = torch_npu_root / "lib"
112		112
	@@ -158,7 +158,7 @@ import torch._inductor.inductor_prims
158	return model_str	158	return model_str
159		159
160		160
161	-def get_fx_graph_code(code, num_args, method=2, runnable=False, kernel_code='', kernel_name=None):	161	+def get_fx_graph_code(code, num_args, method=2, runnable=False, kernel_code='', kernel_name=None):
162	kernel_header = ''	162	kernel_header = ''
163	kernel_wrapper = ''	163	kernel_wrapper = ''
164	kernel_runner_and_acc_comp = ''	164	kernel_runner_and_acc_comp = ''
	@@ -214,7 +214,7 @@ def get_args():
214	"""	214	"""
215	run_code_template = f"""	215	run_code_template = f"""
216		216
217	-try:	217	+try:
218	args = torch.load(os.path.join(dir_path, "data.pth"))	218	args = torch.load(os.path.join(dir_path, "data.pth"))
219	except Exception as e:	219	except Exception as e:
220	{{{{FAKE_ARGS_PLACEHOLDER}}}}	220	{{{{FAKE_ARGS_PLACEHOLDER}}}}
	@@ -229,16 +229,16 @@ with torch.no_grad():
229	output2 = model(*fx_inputs)	229	output2 = model(*fx_inputs)
230	"""	230	"""
231	code_template = f"""	231	code_template = f"""
232	-import os	232	+import os
233	import torch	233	import torch
234	from torch._inductor.compile_fx import clone_preserve_strides	234	from torch._inductor.compile_fx import clone_preserve_strides
235	from torch._dynamo.testing import rand_strided	235	from torch._dynamo.testing import rand_strided
236	from torch import device	236	from torch import device
237		237
238	import torch_npu	238	import torch_npu
239	-from torch_npu._inductor.ascend_npu_ir.ascend_npu_ir import config as npu_config	239	+from torch_npu._inductor.ascend_npu_ir.ascend_npu_ir import config as npu_config
240	{kernel_header}	240	{kernel_header}
241	-file_path = os.path.abspath(__file__)	241	+file_path = os.path.abspath(__file__)
242	dir_path = os.path.dirname(file_path)	242	dir_path = os.path.dirname(file_path)
243		243
244	{kernel_code}	244	{kernel_code}
	@@ -250,7 +250,7 @@ class GraphModule(torch.nn.Module):
250	{code}	250	{code}
251	model = GraphModule().npu()	251	model = GraphModule().npu()
252		252
253	-{run_code_template if runnable else transformed_code_template}	253	+{run_code_template if runnable else transformed_code_template}
254	{fx_runner if runnable else ''}	254	{fx_runner if runnable else ''}
255	{kernel_runner_and_acc_comp if runnable else ''}	255	{kernel_runner_and_acc_comp if runnable else ''}
256	"""	256	"""
	@@ -270,28 +270,28 @@ def view_to_reshape(gm: torch.fx.GraphModule):
270	op="call_function", target=torch.ops.aten.view.default	270	op="call_function", target=torch.ops.aten.view.default
271	):	271	):
272	nd.target = torch.ops.aten.reshape.default	272	nd.target = torch.ops.aten.reshape.default
273	-	273	+
274	for nd in gm.graph.find_nodes(	274	for nd in gm.graph.find_nodes(
275	op="call_function", target=torch.ops.aten.div.Tensor	275	op="call_function", target=torch.ops.aten.div.Tensor
276	):	276	):
277	if not (isinstance(nd.args[1], torch.fx.node.Node) and \	277	if not (isinstance(nd.args[1], torch.fx.node.Node) and \
278	isinstance(nd.args[1].meta['val'], torch.Tensor)):	278	isinstance(nd.args[1].meta['val'], torch.Tensor)):
279	nd.target = torch.ops.aten.div.Scalar	279	nd.target = torch.ops.aten.div.Scalar
280	-	280	+
281	for nd in gm.graph.find_nodes(	281	for nd in gm.graph.find_nodes(
282	op="call_function", target=torch.ops.aten.add.Tensor	282	op="call_function", target=torch.ops.aten.add.Tensor
283	):	283	):
284	if not (isinstance(nd.args[1], torch.fx.node.Node) and \	284	if not (isinstance(nd.args[1], torch.fx.node.Node) and \
285	isinstance(nd.args[1].meta['val'], torch.Tensor)):	285	isinstance(nd.args[1].meta['val'], torch.Tensor)):
286	nd.target = torch.ops.aten.add.Scalar	286	nd.target = torch.ops.aten.add.Scalar
287	-	287	+
288	for nd in gm.graph.find_nodes(	288	for nd in gm.graph.find_nodes(
289	op="call_function", target=torch.ops.aten.sub.Tensor	289	op="call_function", target=torch.ops.aten.sub.Tensor
290	- ):	290	+ ):
291	if not (isinstance(nd.args[1], torch.fx.node.Node) and \	291	if not (isinstance(nd.args[1], torch.fx.node.Node) and \
292	isinstance(nd.args[1].meta['val'], torch.Tensor)):	292	isinstance(nd.args[1].meta['val'], torch.Tensor)):
293	nd.target = torch.ops.aten.sub.Scalar	293	nd.target = torch.ops.aten.sub.Scalar
294	-	294	+
295	for nd in gm.graph.find_nodes(	295	for nd in gm.graph.find_nodes(
296	op="call_function", target=torch.ops.aten.mul.Tensor	296	op="call_function", target=torch.ops.aten.mul.Tensor
297	):	297	):
	@@ -303,7 +303,7 @@ def view_to_reshape(gm: torch.fx.GraphModule):
303	op="call_function", target=torch.ops.prims.convert_element_type.default	303	op="call_function", target=torch.ops.prims.convert_element_type.default
304	):	304	):
305	nd.target = torch.ops.npu.npu_dtype_cast.default	305	nd.target = torch.ops.npu.npu_dtype_cast.default
306	-	306	+
307	def npu_cast_to_prim_cast(gm: torch.fx.GraphModule):	307	def npu_cast_to_prim_cast(gm: torch.fx.GraphModule):
308	"""	308	"""
309	Replace npu.npu_dtype_cast ops in the GraphModule to prims.convert_element_type ops.	309	Replace npu.npu_dtype_cast ops in the GraphModule to prims.convert_element_type ops.
	@@ -359,7 +359,7 @@ def npu_optimize_fx_graph(gm: torch.fx.GraphModule):
359	gm.graph.erase_node(nd)	359	gm.graph.erase_node(nd)
360	aten_empty_nodes.remove(node0)	360	aten_empty_nodes.remove(node0)
361	gm.graph.erase_node(node0)	361	gm.graph.erase_node(node0)
362	-	362	+
363	gm.recompile()	363	gm.recompile()
364		364
365		365
	@@ -374,7 +374,7 @@ def fold_expand(gm: torch.fx.GraphModule) -> None:
374		374
375	inp0 = node.args[0] if len(node.args) > 0 else None	375	inp0 = node.args[0] if len(node.args) > 0 else None
376	inp1 = node.args[1] if len(node.args) > 1 else None	376	inp1 = node.args[1] if len(node.args) > 1 else None
377	- if (isinstance(inp0, torch.fx.Node) and inp0.op == 'call_function' and	377	+ if (isinstance(inp0, torch.fx.Node) and inp0.op == 'call_function' and
378	inp0.target == torch.ops.aten.expand.default):	378	inp0.target == torch.ops.aten.expand.default):
379	if len(inp0.args) > 0:	379	if len(inp0.args) > 0:
380	expand_input = inp0.args[0]	380	expand_input = inp0.args[0]
	@@ -382,7 +382,7 @@ def fold_expand(gm: torch.fx.GraphModule) -> None:
382	if len(inp0.users) == 0:	382	if len(inp0.users) == 0:
383	graph.erase_node(inp0)	383	graph.erase_node(inp0)
384	changed = True	384	changed = True
385	- elif (isinstance(inp1, torch.fx.Node) and inp1.op == 'call_function' and	385	+ elif (isinstance(inp1, torch.fx.Node) and inp1.op == 'call_function' and
386	inp1.target == torch.ops.aten.expand.default):	386	inp1.target == torch.ops.aten.expand.default):
387	if len(inp1.args) > 0:	387	if len(inp1.args) > 0:
388	expand_input = inp1.args[0]	388	expand_input = inp1.args[0]
	@@ -393,7 +393,7 @@ def fold_expand(gm: torch.fx.GraphModule) -> None:
393	if changed:	393	if changed:
394	graph.lint()	394	graph.lint()
395	graph.eliminate_dead_code()	395	graph.eliminate_dead_code()
396	-	396	+
397	gm.recompile()	397	gm.recompile()
398		398
399		399
	@@ -465,12 +465,12 @@ class MLIRProcessor:
465	def __init__(self, bisheng_install_path: str = None):	465	def __init__(self, bisheng_install_path: str = None):
466	"""	466	"""
467	初始化MLIR处理器	467	初始化MLIR处理器
468	-	468	+
469	:param bisheng_install_path: Bisheng安装路径，默认从环境变量获取	469	:param bisheng_install_path: Bisheng安装路径，默认从环境变量获取
470	"""	470	"""
471	bisheng_install_path = os.getenv('BISHENG_INSTALL_PATH', '')	471	bisheng_install_path = os.getenv('BISHENG_INSTALL_PATH', '')
472	self.bisheng_torch_mlir_path = os.path.join(bisheng_install_path, "bishengir-opt")	472	self.bisheng_torch_mlir_path = os.path.join(bisheng_install_path, "bishengir-opt")
473	-	473	+
474	def extract_function(self, module: ir.Module) -> func_dialect.FuncOp:	474	def extract_function(self, module: ir.Module) -> func_dialect.FuncOp:
475	"""从MLIR模块中提取主函数并添加标记属性"""	475	"""从MLIR模块中提取主函数并添加标记属性"""
476	with module.context:	476	with module.context:
	@@ -479,20 +479,20 @@ class MLIRProcessor:
479	func.attributes["hacc.placeholder"] = ir.UnitAttr.get(func.context)	479	func.attributes["hacc.placeholder"] = ir.UnitAttr.get(func.context)
480	return func	480	return func
481	raise ValueError("No valid FuncOp found in module")	481	raise ValueError("No valid FuncOp found in module")
482	-	482	+
483	def rebuild_mlir_module(self, module_str: str) -> ir.Module:	483	def rebuild_mlir_module(self, module_str: str) -> ir.Module:
484	"""从字符串重新构建MLIR模块"""	484	"""从字符串重新构建MLIR模块"""
485	with ir.Context() as ctx:	485	with ir.Context() as ctx:
486	ctx.allow_unregistered_dialects = True	486	ctx.allow_unregistered_dialects = True
487	torch_mlir.dialects.torch.register_dialect(ctx)	487	torch_mlir.dialects.torch.register_dialect(ctx)
488	return ir.Module.parse(module_str)	488	return ir.Module.parse(module_str)
489	-	489	+
490	def get_signature(self, func: func_dialect.FuncOp) -> tuple:	490	def get_signature(self, func: func_dialect.FuncOp) -> tuple:
491	"""获取函数的签名信息：类型签名、输出数量和张量维度"""	491	"""获取函数的签名信息：类型签名、输出数量和张量维度"""
492	func_type = func.type	492	func_type = func.type
493	signature = {}	493	signature = {}
494	ranks = []	494	ranks = []
495	-	495	+
496	# 处理输入+输出类型	496	# 处理输入+输出类型
497	for i, tensor_type in enumerate(func_type.inputs + func_type.results):	497	for i, tensor_type in enumerate(func_type.inputs + func_type.results):
498	try: # RankedTensorType	498	try: # RankedTensorType
	@@ -506,17 +506,17 @@ class MLIRProcessor:
506	dim_end = type_str.find(']', dim_start)	506	dim_end = type_str.find(']', dim_start)
507	dim_str = type_str[dim_start:dim_end]	507	dim_str = type_str[dim_start:dim_end]
508	ranks.append(dim_str.count(',') + 1 if dim_str else 0)	508	ranks.append(dim_str.count(',') + 1 if dim_str else 0)
509	-	509	+
510	num_outputs = len(func_type.results)	510	num_outputs = len(func_type.results)
511	return signature, num_outputs, ranks	511	return signature, num_outputs, ranks
512	-	512	+
513	- def process_mlir(self,	513	+ def process_mlir(self,
514	- module: Union[str, ir.Module],	514	+ module: Union[str, ir.Module],
515	- get_sig: bool = True,	515	+ get_sig: bool = True,
516	dynamic: bool = False) -> tuple:	516	dynamic: bool = False) -> tuple:
517	"""	517	"""
518	处理MLIR模块的核心方法	518	处理MLIR模块的核心方法
519	-	519	+
520	:param module: MLIR模块字符串或对象	520	:param module: MLIR模块字符串或对象
521	:param get_sig: 是否获取函数签名	521	:param get_sig: 是否获取函数签名
522	:param dynamic: 是否为动态执行模式	522	:param dynamic: 是否为动态执行模式
	@@ -524,7 +524,7 @@ class MLIRProcessor:
524	"""	524	"""
525	if isinstance(module, str):	525	if isinstance(module, str):
526	module = self.rebuild_mlir_module(module)	526	module = self.rebuild_mlir_module(module)
527	-	527	+
528	func = self.extract_function(module)	528	func = self.extract_function(module)
529	kernel_info = None	529	kernel_info = None
530	func_str = str(func)	530	func_str = str(func)
	@@ -539,29 +539,29 @@ class MLIRProcessor:
539	"ranks": ranks,	539	"ranks": ranks,
540	'kernel_hash': module_hash,	540	'kernel_hash': module_hash,
541	}	541	}
542	-	542	+
543	return func_str, kernel_info	543	return func_str, kernel_info
544	-	544	+
545	def get_named_op_str(self,	545	def get_named_op_str(self,
546	module: Union[str, ir.Module],	546	module: Union[str, ir.Module],
547	kernel_name: str,	547	kernel_name: str,
548	dynamic: bool = False) -> Dict[str, Any]:	548	dynamic: bool = False) -> Dict[str, Any]:
549	"""	549	"""
550	获取命名操作格式的MLIR字符串	550	获取命名操作格式的MLIR字符串
551	-	551	+
552	:param module: MLIR模块字符串或对象	552	:param module: MLIR模块字符串或对象
553	:param kernel_name: 内核名称（用于临时文件）	553	:param kernel_name: 内核名称（用于临时文件）
554	:param dynamic: 是否为动态执行模式	554	:param dynamic: 是否为动态执行模式
555	:return: 包含处理结果和签名字典	555	:return: 包含处理结果和签名字典
556	"""	556	"""
557	func_str, sig_dict = self.process_mlir(module, get_sig=True, dynamic=dynamic)	557	func_str, sig_dict = self.process_mlir(module, get_sig=True, dynamic=dynamic)
558	-	558	+
559	cleaned_func = func_str.replace(	559	cleaned_func = func_str.replace(
560	- '"#hfusion.fusion_kind<PURE_ELEMWISE>"',	560	+ '"#hfusion.fusion_kind<PURE_ELEMWISE>"',
561	'#hfusion.fusion_kind<PURE_ELEMWISE>'	561	'#hfusion.fusion_kind<PURE_ELEMWISE>'
562	)	562	)
563	logger.debug(f"原始Linalg方言MLIR:\n{cleaned_func}")	563	logger.debug(f"原始Linalg方言MLIR:\n{cleaned_func}")
564	-	564	+
565	# 执行转换命令	565	# 执行转换命令
566	with tempfile.TemporaryDirectory() as tmpdir:	566	with tempfile.TemporaryDirectory() as tmpdir:
567	torch_mlir_path = os.path.join(tmpdir, f"{kernel_name}.mlir")	567	torch_mlir_path = os.path.join(tmpdir, f"{kernel_name}.mlir")
	@@ -572,33 +572,33 @@ class MLIRProcessor:
572	"--torch-backend-to-named-op-backend-pipeline="	572	"--torch-backend-to-named-op-backend-pipeline="
573	"\"ensure-no-implicit-broadcast=true\" "	573	"\"ensure-no-implicit-broadcast=true\" "
574	f"{torch_mlir_path}")	574	f"{torch_mlir_path}")
575	-	575	+
576	try:	576	try:
577	result = subprocess.check_output(	577	result = subprocess.check_output(
578	cmd, text=True, shell=True	578	cmd, text=True, shell=True
579	)	579	)
580	# 过滤全局定义并更新函数属性	580	# 过滤全局定义并更新函数属性
581	processed_mlir = "\n".join(	581	processed_mlir = "\n".join(
582	- line for line in result.splitlines()	582	+ line for line in result.splitlines()
583	if "ml_program.global" not in line	583	if "ml_program.global" not in line
584	)	584	)
585	-	585	+
586	# 根据模式设置函数属性	586	# 根据模式设置函数属性
587	- func_attr = ("hacc.entry, hacc.function_kind = #hacc.function_kind<HOST>"	587	+ func_attr = ("hacc.entry, hacc.function_kind = #hacc.function_kind<HOST>"
588	- if dynamic else	588	+ if dynamic else
589	"hacc.entry, hacc.function_kind = #hacc.function_kind<DEVICE>")	589	"hacc.entry, hacc.function_kind = #hacc.function_kind<DEVICE>")
590	processed_mlir = processed_mlir.replace("hacc.placeholder", func_attr)	590	processed_mlir = processed_mlir.replace("hacc.placeholder", func_attr)
591	-	591	+
592	# 应用额外的数据类型处理（需实现mlir_match_and_replace_unsupported_dtypes）	592	# 应用额外的数据类型处理（需实现mlir_match_and_replace_unsupported_dtypes）
593	final_mlir = self._replace_unsupported_dtypes(processed_mlir)	593	final_mlir = self._replace_unsupported_dtypes(processed_mlir)
594	logger.debug(f"转换后的NamedOp方言MLIR:\n{final_mlir}")	594	logger.debug(f"转换后的NamedOp方言MLIR:\n{final_mlir}")
595	-	595	+
596	return final_mlir, sig_dict	596	return final_mlir, sig_dict
597	-	597	+
598	except subprocess.CalledProcessError as e:	598	except subprocess.CalledProcessError as e:
599	logger.error(f"命令执行失败: {cmd}\n错误: {e.output}")	599	logger.error(f"命令执行失败: {cmd}\n错误: {e.output}")
600	raise RuntimeError(f"MLIR转换失败: {e.stderr}") from e	600	raise RuntimeError(f"MLIR转换失败: {e.stderr}") from e
601	-	601	+
602	def _replace_unsupported_dtypes(self, mlir_text: str) -> str:	602	def _replace_unsupported_dtypes(self, mlir_text: str) -> str:
603	"""替换不支持的MLIR数据类型"""	603	"""替换不支持的MLIR数据类型"""
604	pattern1 = r"%(\d+) = arith\.truncf %(\w+) : f64 to bf16"	604	pattern1 = r"%(\d+) = arith\.truncf %(\w+) : f64 to bf16"
	@@ -624,8 +624,8 @@ def mlir_match_and_replace_unsupported_dtypes(mlir_text: str) -> str:
624		624
625		625
626	def to_folder(	626	def to_folder(
627	- gm: torch.fx.GraphModule,	627	+ gm: torch.fx.GraphModule,
628	- folder: Union[str, os.PathLike],	628	+ folder: Union[str, os.PathLike],
629	graph_hash: str,	629	graph_hash: str,
630	module_name: str = "FxModule"):	630	module_name: str = "FxModule"):
631	"""Dumps out module to ``folder`` with ``module_name`` so that it can be	631	"""Dumps out module to ``folder`` with ``module_name`` so that it can be
	@@ -726,27 +726,27 @@ def is_fx_dynamic(graph):
726	def replace_placeholders(file_path: str, replacements: dict, placeholder_format: str = r'\{\{(\w+)\}\}') -> None:	726	def replace_placeholders(file_path: str, replacements: dict, placeholder_format: str = r'\{\{(\w+)\}\}') -> None:
727	"""	727	"""
728	替换文件中的占位符	728	替换文件中的占位符
729	-	729	+
730	:param file_path: 文件路径	730	:param file_path: 文件路径
731	:param replacements: 替换字典，如 {'function_body': 'your _code'}	731	:param replacements: 替换字典，如 {'function_body': 'your _code'}
732	:param placeholder_format: 占位符正则表达式（默认匹配{{xxx}}）	732	:param placeholder_format: 占位符正则表达式（默认匹配{{xxx}}）
733	"""	733	"""
734	with open(file_path, 'r', encoding='utf-8') as f:	734	with open(file_path, 'r', encoding='utf-8') as f:
735	content = f.read()	735	content = f.read()
736	-	736	+
737	pattern = re.compile(placeholder_format)	737	pattern = re.compile(placeholder_format)
738	-	738	+
739	def replacer(match: re.Match) -> str:	739	def replacer(match: re.Match) -> str:
740	placeholder = match.group(1)	740	placeholder = match.group(1)
741	replacement = replacements.get(placeholder, match.group(0))	741	replacement = replacements.get(placeholder, match.group(0))
742	-	742	+
743	line_start = content.rfind('\n', 0, match.start()) + 1	743	line_start = content.rfind('\n', 0, match.start()) + 1
744	indent = re.match(r'^\s*', content[line_start:match.start()]).group(0)	744	indent = re.match(r'^\s*', content[line_start:match.start()]).group(0)
745	-	745	+
746	return '\n'.join([indent + line for line in replacement.split('\n')])	746	return '\n'.join([indent + line for line in replacement.split('\n')])
747	-	747	+
748	new_content = pattern.sub(replacer, content)	748	new_content = pattern.sub(replacer, content)
749	-	749	+
750	with open(file_path, 'w', encoding='utf-8') as f:	750	with open(file_path, 'w', encoding='utf-8') as f:
751	f.write(new_content)	751	f.write(new_content)
752		752

Mtorch_npu/_inductor/codegen/cpp_wrapper.py+1-1

	@@ -679,7 +679,7 @@ class CppWrapperNpu(CppWrapperCpu):
679	static_cast<int32_t>(grid_1),	679	static_cast<int32_t>(grid_1),
680	static_cast<int32_t>(grid_2)	680	static_cast<int32_t>(grid_2)
681	}};	681	}};
682	-	682	+
683	uint32_t block_num = grid_0 * grid_1 * grid_2;	683	uint32_t block_num = grid_0 * grid_1 * grid_2;
684	auto arg_ptr = static_cast<void*>(&kernel_args);	684	auto arg_ptr = static_cast<void*>(&kernel_args);
685	auto arg_size = sizeof(kernel_args);	685	auto arg_size = sizeof(kernel_args);

Mtorch_npu/_inductor/codegen/kernel_analysis.py+2-2

Mtorch_npu/_inductor/codegen/scheduling.py+1-1

	@@ -448,7 +448,7 @@ class NPUTritonScheduling(TritonScheduling):
448	split_tiling = SplitTiling(kernel)	448	split_tiling = SplitTiling(kernel)
449	split_tiling.select_split_tiling_axis()	449	split_tiling.select_split_tiling_axis()
450	kernel.load_store_indexing = split_tiling.indexing	450	kernel.load_store_indexing = split_tiling.indexing
451	- # ReductionAnalysis depends on kernel.load_store_indexing	451	+ # ReductionAnalysis depends on kernel.load_store_indexing
452	if kernel.inside_reduction:	452	if kernel.inside_reduction:
453	kernel.reduce_analysis = ReductionAnalysis(kernel)	453	kernel.reduce_analysis = ReductionAnalysis(kernel)
454		454

Mtorch_npu/_inductor/codegen/split_tiling.py+2-2

Mtorch_npu/_inductor/codegen/triton.py+5-5

	@@ -193,7 +193,7 @@ def patch_TritonCSEVariable__init__(
193	super(TritonCSEVariable, self).__init__(name, bounds, dtype, shape=shape)	193	super(TritonCSEVariable, self).__init__(name, bounds, dtype, shape=shape)
194	self.mask_vars: OrderedSet[str] = OrderedSet()	194	self.mask_vars: OrderedSet[str] = OrderedSet()
195	assert dtype is not None, "TritonCSEVariable must have dtype"	195	assert dtype is not None, "TritonCSEVariable must have dtype"
196	-	196	+
197		197
198	@staticmethod	198	@staticmethod
199	def select_index_dtype(node_schedule, numel, reduction_numel):	199	def select_index_dtype(node_schedule, numel, reduction_numel):
	@@ -262,7 +262,7 @@ class IterationRangesEntryNPUIndex(IterationRangesEntry):
262	line = f"{var.name} = {self.codegen_index(dir_index)}"	262	line = f"{var.name} = {self.codegen_index(dir_index)}"
263	self.writeline(line)	263	self.writeline(line)
264		264
265	- # reduction axis	265	+ # reduction axis
266	if self.prefix == 'r':	266	if self.prefix == 'r':
267	if V.kernel.inside_reduction and V.kernel.current_node \	267	if V.kernel.inside_reduction and V.kernel.current_node \
268	and isinstance(V.kernel.current_node, SchedulerNode) \	268	and isinstance(V.kernel.current_node, SchedulerNode) \
	@@ -596,7 +596,7 @@ class NPUIndexTritonKernel(TritonKernel):
596	break	596	break
597	if dim is None:	597	if dim is None:
598	continue	598	continue
599	-	599	+
600	if dim.parent == axis.parent:	600	if dim.parent == axis.parent:
601	dtype = V.graph.get_dtype(node.node.name)	601	dtype = V.graph.get_dtype(node.node.name)
602	should_break_all = True	602	should_break_all = True
	@@ -1110,7 +1110,7 @@ class NPUIndexTritonKernel(TritonKernel):
1110		1110
1111	return None	1111	return None
1112		1112
1113	- # select the golden varlist, from to which to deduce permute, broadcast shape	1113	+ # select the golden varlist, from to which to deduce permute, broadcast shape
1114	def select_golden_varlist(self):	1114	def select_golden_varlist(self):
1115	longest = None	1115	longest = None
1116	maximum_length = 0	1116	maximum_length = 0
	@@ -1416,7 +1416,7 @@ class NPUIndexTritonKernel(TritonKernel):
1416	index_str = indexing.index_str	1416	index_str = indexing.index_str
1417	mask_str = indexing.mask_str	1417	mask_str = indexing.mask_str
1418	line = f"tl.load({var} + ({index_str}), {mask_str}{ep}{other})"	1418	line = f"tl.load({var} + ({index_str}), {mask_str}{ep}{other})"
1419	-	1419	+
1420	if (	1420	if (
1421	dtype in (torch.float16, torch.bfloat16)	1421	dtype in (torch.float16, torch.bfloat16)
1422	and config.triton.codegen_upcast_to_fp32	1422	and config.triton.codegen_upcast_to_fp32

Mtorch_npu/_inductor/codegen/wrapper.py+7-7

	@@ -146,7 +146,7 @@ class NPUWrapperCodeGen(_NPUKernelCodegenMixin, PythonWrapperCodegen):
146	"""	146	"""
147	if not config.benchmark_harness:	147	if not config.benchmark_harness:
148	return None	148	return None
149	-	149	+
150	if npu_config.aot_inductor.debug_kernel:	150	if npu_config.aot_inductor.debug_kernel:
151	return self.add_npu_repro(output)	151	return self.add_npu_repro(output)
152		152
	@@ -169,7 +169,7 @@ class NPUWrapperCodeGen(_NPUKernelCodegenMixin, PythonWrapperCodegen):
169	"print(result)",	169	"print(result)",
170	]	170	]
171	)	171	)
172	-	172	+
173	def add_repro_func(self, output):	173	def add_repro_func(self, output):
174	seen_constants = set()	174	seen_constants = set()
175		175
	@@ -186,7 +186,7 @@ class NPUWrapperCodeGen(_NPUKernelCodegenMixin, PythonWrapperCodegen):
186	sha1 = hashlib.sha1()	186	sha1 = hashlib.sha1()
187	sha1.update(byte)	187	sha1.update(byte)
188	return sha1.hexdigest()	188	return sha1.hexdigest()
189	-	189	+
190	def save_tensor(tensor, path):	190	def save_tensor(tensor, path):
191	dirname = os.path.dirname(path)	191	dirname = os.path.dirname(path)
192	if not os.path.exists(dirname):	192	if not os.path.exists(dirname):
	@@ -235,11 +235,11 @@ class NPUWrapperCodeGen(_NPUKernelCodegenMixin, PythonWrapperCodegen):
235	# these 'global var_name' lines	235	# these 'global var_name' lines
236	output.writeline(f"global {name}")	236	output.writeline(f"global {name}")
237	add_torchbind_input(name, torchbind_obj)	237	add_torchbind_input(name, torchbind_obj)
238	-	238	+
239	call_str = f"call([{', '.join(V.graph.graph_inputs.keys())}])"	239	call_str = f"call([{', '.join(V.graph.graph_inputs.keys())}])"
240	output.writeline(f"fn = lambda: {call_str}")	240	output.writeline(f"fn = lambda: {call_str}")
241	output.writeline("return fn()")	241	output.writeline("return fn()")
242	-	242	+
243	def add_benchmark_func(self, output):	243	def add_benchmark_func(self, output):
244	def add_fake_input(name, shape, stride, device, dtype):	244	def add_fake_input(name, shape, stride, device, dtype):
245	output.writeline(	245	output.writeline(
	@@ -286,7 +286,7 @@ class NPUWrapperCodeGen(_NPUKernelCodegenMixin, PythonWrapperCodegen):
286	value.get_device(),	286	value.get_device(),
287	value.get_dtype(),	287	value.get_dtype(),
288	)	288	)
289	-	289	+
290	call_str = f"repro_run({', '.join(V.graph.graph_inputs.keys())})"	290	call_str = f"repro_run({', '.join(V.graph.graph_inputs.keys())})"
291	output.writeline(f"fn = lambda: {call_str}")	291	output.writeline(f"fn = lambda: {call_str}")
292	output.writeline("return fn()")	292	output.writeline("return fn()")
	@@ -294,7 +294,7 @@ class NPUWrapperCodeGen(_NPUKernelCodegenMixin, PythonWrapperCodegen):
294	def write_prefix(self) -> None:	294	def write_prefix(self) -> None:
295	super().write_prefix()	295	super().write_prefix()
296	if torch_npu.npu.aclnn._use_static_aclnn_kernel:	296	if torch_npu.npu.aclnn._use_static_aclnn_kernel:
297	- self.prefix.do_indent()	297	+ self.prefix.do_indent()
298	with self.prefix.indent():	298	with self.prefix.indent():
299	self.prefix.writeline('global has_initialized')	299	self.prefix.writeline('global has_initialized')
300	self.prefix.writeline('if not has_initialized:')	300	self.prefix.writeline('if not has_initialized:')

Mtorch_npu/_inductor/cpp_builder.py+8-8

@@ -29,7 +29,7 @@ def include_paths(npu: bool = False) -> List[str]:

29

30

Args:

30

Args:

31

npu: If 'True', includes NPU-specific include paths.

31

npu: If 'True', includes NPU-specific include paths.

32

-

32

+

33

Returns:

33

Returns:

34

A list if include path strings.

34

A list if include path strings.

35

"""

35

"""

@@ -88,7 +88,7 @@ def get_cpp_torch_device_options(

88

aot_mode: bool = False,

88

aot_mode: bool = False,

89

compile_only: bool = False,

89

compile_only: bool = False,

90

) -> Tuple[List[str], List[str], List[str], List[str], List[str], List[str], List[str]]:

90

) -> Tuple[List[str], List[str], List[str], List[str], List[str], List[str], List[str]]:

91

-

91

+

92

npu = "npu" == device_type

92

npu = "npu" == device_type

93

94

definations: List[str] = []

94

definations: List[str] = []

@@ -157,14 +157,14 @@ def _get_optimization_cflags(

157

debug_cflags, debug_ldflags = _get_inductor_debug_symbol_cflags()

157

debug_cflags, debug_ldflags = _get_inductor_debug_symbol_cflags()

158

cflags += debug_cflags

158

cflags += debug_cflags

159

ldflags += debug_ldflags

159

ldflags += debug_ldflags

160

-

160

+

161

cflags += _get_ffast_math_flags()

161

cflags += _get_ffast_math_flags()

162

-

162

+

163

if _IS_WINDOWS:

163

if _IS_WINDOWS:

164

pass

164

pass

165

else:

165

else:

166

if sys.platform != "darwin":

166

if sys.platform != "darwin":

167

- # on macos, unknown argument: '-fno-tree-loop-vectorize'

167

+ # on macos, unknown argument: '-fno-tree-loop-vectorize'

168

if _is_gcc(cpp_compiler):

168

if _is_gcc(cpp_compiler):

169

cflags.append("fno-tree-loop-vectorize")

169

cflags.append("fno-tree-loop-vectorize")

atomgit-bot5月21日

🟡 Medium Priority

第 169 行 cflags.append("fno-tree-loop-vectorize") 应为 "-fno-tree-loop-vectorize"，否则 GCC 会将 fno-tree-loop-vectorize 视为源文件名而非编译选项。该行出现于 diff 上下文中，属于已有代码。

改动建议

169	- cflags.append("fno-tree-loop-vectorize")
169	+ cflags.append("-fno-tree-loop-vectorize")

应用建议

不准确？

Jingwei Huang

5月21日评论：

170

# https://stackoverflow.com/questions/65966969/why-does-march-native-not-work-on-apple-m1

170

# https://stackoverflow.com/questions/65966969/why-does-march-native-not-work-on-apple-m1

@@ -176,10 +176,10 @@ def _get_optimization_cflags(

176

cflags.append("march=rv64gc")

176

cflags.append("march=rv64gc")

177

elif platform.machine() == "riscv32":

177

elif platform.machine() == "riscv32":

178

cflags.append("march=rv32gc")

178

cflags.append("march=rv32gc")

179

-

179

+

180

- if config.aot_inductor.enable_lto and _is_clang(cpp_compiler):

180

+ if config.aot_inductor.enable_lto and _is_clang(cpp_compiler):

181

cflags.append("flto=thin")

181

cflags.append("flto=thin")

182

-

182

+

183

return cflags, ldflags

183

return cflags, ldflags

184

185

Mtorch_npu/_inductor/fx_passes/joint_graph.py+1-1

	@@ -11,5 +11,5 @@ def patch_constant_fold_uniform_value():
11	src_func(gm)	11	src_func(gm)
12	if isinstance(gm, torch.fx.GraphModule):	12	if isinstance(gm, torch.fx.GraphModule):
13	gm.graph.eliminate_dead_code()	13	gm.graph.eliminate_dead_code()
14	-	14	+
15	joint_graph.constant_fold_uniform_value = new_constant_fold_uniform_value	15	joint_graph.constant_fold_uniform_value = new_constant_fold_uniform_value

Mtorch_npu/_inductor/lowering.py+3-3

	@@ -201,7 +201,7 @@ def _register_npu_inductor_fallbacks():
201	for op in overload_op_set:	201	for op in overload_op_set:
202	if op in lowerings:	202	if op in lowerings:
203	del lowerings[op]	203	del lowerings[op]
204	-	204	+
205	if npu_config.dump_fx_graph:	205	if npu_config.dump_fx_graph:
206	from .lowering_fx import _register_npu_inductor_fallbacks_fx	206	from .lowering_fx import _register_npu_inductor_fallbacks_fx
207	(squeeze, _validate_dim, div, square, sub, sum_) = _register_npu_inductor_fallbacks_fx(make_reduction)	207	(squeeze, _validate_dim, div, square, sub, sum_) = _register_npu_inductor_fallbacks_fx(make_reduction)
	@@ -356,7 +356,7 @@ def get_nested_attr(obj, attr_path, default=None):
356	return reduce(getattr, attr_path.split('.'), obj)	356	return reduce(getattr, attr_path.split('.'), obj)
357	except AttributeError:	357	except AttributeError:
358	return default	358	return default
359	-	359	+
360		360
361	def _fallback_ops_with_meta():	361	def _fallback_ops_with_meta():
362	"""	362	"""
	@@ -377,7 +377,7 @@ def _fallback_ops_with_meta():
377	name, overload = name_with_overload.rsplit(".", 1)	377	name, overload = name_with_overload.rsplit(".", 1)
378	else:	378	else:
379	name, overload = name_with_overload, "default"	379	name, overload = name_with_overload, "default"
380	-	380	+
381	normalized_path = f"{namespace}.{name}.{overload}"	381	normalized_path = f"{namespace}.{name}.{overload}"
382	op_overload = get_nested_attr(torch.ops, normalized_path)	382	op_overload = get_nested_attr(torch.ops, normalized_path)
383	if not isinstance(op_overload, torch._ops.OpOverload):	383	if not isinstance(op_overload, torch._ops.OpOverload):

Mtorch_npu/_inductor/lowering_fx.py+5-5

	@@ -455,7 +455,7 @@ def create_compile_kwargs(final_kernel, fx_call_args, fx_args):
455	def generate_fx_graph_code(code, kernel_code, kernel_name, compile_kwargs):	455	def generate_fx_graph_code(code, kernel_code, kernel_name, compile_kwargs):
456	code = textwrap.indent(code, ' ')	456	code = textwrap.indent(code, ' ')
457	code_template = f"""	457	code_template = f"""
458	-import os	458	+import os
459	import torch	459	import torch
460	from torch._inductor.compile_fx import clone_preserve_strides	460	from torch._inductor.compile_fx import clone_preserve_strides
461	from torch._dynamo.testing import rand_strided	461	from torch._dynamo.testing import rand_strided
	@@ -514,7 +514,7 @@ del async_compile
514	def run():	514	def run():
515	stream0 = get_raw_stream(0)	515	stream0 = get_raw_stream(0)
516		516
517	-	517	+
518	args = torch.load(os.path.join(dir_path, "data.pth"))	518	args = torch.load(os.path.join(dir_path, "data.pth"))
519		519
520	call_inputs_indices = call_args_mapping[:num_inputs]	520	call_inputs_indices = call_args_mapping[:num_inputs]
	@@ -522,7 +522,7 @@ def run():
522		522
523	args = [arg.npu() if isinstance(arg, torch.Tensor) else arg for arg in args]	523	args = [arg.npu() if isinstance(arg, torch.Tensor) else arg for arg in args]
524		524
525	- fx_args = []	525	+ fx_args = []
526	for idx in call_args_mapping:	526	for idx in call_args_mapping:
527	arg = args[idx]	527	arg = args[idx]
528	if isinstance(arg, torch.Tensor):	528	if isinstance(arg, torch.Tensor):
	@@ -540,7 +540,7 @@ def run():
540	out1 = out1.reshape(out2.shape)	540	out1 = out1.reshape(out2.shape)
541	if idx in non_contiguous_indices['outputs']:	541	if idx in non_contiguous_indices['outputs']:
542	out2.copy_(out1)	542	out2.copy_(out1)
543	- else:	543	+ else:
544	out2.data = out1.data	544	out2.data = out1.data
545		545
546	{kernel_name}.run(*args, stream=stream0)	546	{kernel_name}.run(*args, stream=stream0)
	@@ -728,7 +728,7 @@ def _make_reduction_inner(x, *, axis, keepdims, dtype, override_return_dtype):
728	reduction_ranges=reduced_sizes,	728	reduction_ranges=reduced_sizes,
729	)	729	)
730		730
731	-	731	+
732	def dump_fx_graph_code(code, dump_path, traced_graph_hash):	732	def dump_fx_graph_code(code, dump_path, traced_graph_hash):
733	py_path = os.path.join(dump_path, traced_graph_hash + '.py')	733	py_path = os.path.join(dump_path, traced_graph_hash + '.py')
734	PathManager.check_input_file_path(py_path)	734	PathManager.check_input_file_path(py_path)

Mtorch_npu/_inductor/npu_compare.py+6-6

	@@ -26,7 +26,7 @@ def compare_outputs(
26	continue	26	continue
27	if actual.dtype != expected.dtype:	27	if actual.dtype != expected.dtype:
28	expected = expected.to(actual.dtype)	28	expected = expected.to(actual.dtype)
29	-	29	+
30	tol = tolerances.get(actual.dtype, tolerances["default"])	30	tol = tolerances.get(actual.dtype, tolerances["default"])
31	rtol, atol = tol["rtol"], tol["atol"]	31	rtol, atol = tol["rtol"], tol["atol"]
32	matches = torch.isclose(actual, expected, rtol=rtol, atol=atol, equal_nan=True)	32	matches = torch.isclose(actual, expected, rtol=rtol, atol=atol, equal_nan=True)
	@@ -34,7 +34,7 @@ def compare_outputs(
34	_report_mismatch(idx, actual, expected, matches, rtol, atol, kernel_name)	34	_report_mismatch(idx, actual, expected, matches, rtol, atol, kernel_name)
35	failed_indices.append(idx)	35	failed_indices.append(idx)
36	del matches	36	del matches
37	-	37	+
38	return not failed_indices	38	return not failed_indices
39		39
40		40
	@@ -117,7 +117,7 @@ def check_accuracy_triton(args, launcher, grid, stream, inductor_meta, *kwargs
117	arg = args[idx]	117	arg = args[idx]
118	if isinstance(arg, torch.Tensor):	118	if isinstance(arg, torch.Tensor):
119	fx_args.append(clone_for_accuracy(arg))	119	fx_args.append(clone_for_accuracy(arg))
120	-	120	+
121	fx_graph_call(*fx_args)	121	fx_graph_call(*fx_args)
122		122
123	launcher(args, *kwargs, stream=stream)	123	launcher(args, *kwargs, stream=stream)
	@@ -155,7 +155,7 @@ def check_accuracy_mlir(args, kernel_name, launchers, num_outputs, dynamic, *k
155	args_new = args_new + (arg, arg, 0) + arg.size() + arg.stride()	155	args_new = args_new + (arg, arg, 0) + arg.size() + arg.stride()
156	else:	156	else:
157	args_new = args	157	args_new = args
158	-	158	+
159	output = launcher(args_new, *kwargs)	159	output = launcher(args_new, *kwargs)
160	result = compare_outputs(	160	result = compare_outputs(
161	args[num_inputs:],	161	args[num_inputs:],
	@@ -186,8 +186,8 @@ def _load_fx_model(acc_meta):
186	model = Model()	186	model = Model()
187	acc_meta['_fx_model'] = model	187	acc_meta['_fx_model'] = model
188	return model	188	return model
189	-	189	+
190	-	190	+
191	def check_accuracy_dvm(kobj, acc_meta, kernel_name, args):	191	def check_accuracy_dvm(kobj, acc_meta, kernel_name, args):
192	"""Run DVM kernel then compare outputs against FX graph reference."""	192	"""Run DVM kernel then compare outputs against FX graph reference."""
193	from torch_npu._inductor.ascend_npu_ir.ascend_npu_ir import config as anir_config	193	from torch_npu._inductor.ascend_npu_ir.ascend_npu_ir import config as anir_config

Mtorch_npu/_inductor/npu_device.py+1-1

	@@ -77,7 +77,7 @@ class NewNPUDeviceOpOverrides(NPUDeviceOpOverrides):
77	load_code = """	77	load_code = """
78	static std::unordered_map<std::string, size_t> registered_names;	78	static std::unordered_map<std::string, size_t> registered_names;
79	static std::unordered_map<std::string, std::unique_ptr<size_t>> func_stubs;	79	static std::unordered_map<std::string, std::unique_ptr<size_t>> func_stubs;
80	-	80	+
81	static inline void * loadKernel(	81	static inline void * loadKernel(
82	std::string filePath,	82	std::string filePath,
83	const std::string &&nameFuncMode,	83	const std::string &&nameFuncMode,

Mtorch_npu/_inductor/npu_triton_heuristics.py+8-8

	@@ -212,7 +212,7 @@ class GridExprNpu(GridExpr):
212	class PrecomputedGridNpu(GridNpu):	212	class PrecomputedGridNpu(GridNpu):
213	def __init__(self, , inductor_meta, mode="python", *kwargs):	213	def __init__(self, , inductor_meta, mode="python", *kwargs):
214	super().__init__(inductor_meta=inductor_meta, mode=mode, numels=kwargs.get("numels"))	214	super().__init__(inductor_meta=inductor_meta, mode=mode, numels=kwargs.get("numels"))
215	-	215	+
216	def generate(self, meta: dict[str, int]) -> None:	216	def generate(self, meta: dict[str, int]) -> None:
217	for candidate in self.inductor_meta["precomputed_grids"]:	217	for candidate in self.inductor_meta["precomputed_grids"]:
218	if all(meta.get(k) == v for k, v in candidate["config"].items()):	218	if all(meta.get(k) == v for k, v in candidate["config"].items()):
	@@ -227,7 +227,7 @@ class PrecomputedGridNpu(GridNpu):
227	class FixedGridNpu(GridNpu):	227	class FixedGridNpu(GridNpu):
228	def __init__(self, , inductor_meta, mode="python", *kwargs):	228	def __init__(self, , inductor_meta, mode="python", *kwargs):
229	super().__init__(inductor_meta=inductor_meta, mode=mode, numels=kwargs.get("numels"))	229	super().__init__(inductor_meta=inductor_meta, mode=mode, numels=kwargs.get("numels"))
230	-	230	+
231	@staticmethod	231	@staticmethod
232	def setup_grid_as_args() -> dict[str, Any]:	232	def setup_grid_as_args() -> dict[str, Any]:
233	"""Inductor meta so the launcher takes three extra grid arguments"""	233	"""Inductor meta so the launcher takes three extra grid arguments"""
	@@ -413,7 +413,7 @@ class TritonCompileResultNpu(TritonCompileResult):
413	class NPUCachingAutotuner(CachingAutotuner):	413	class NPUCachingAutotuner(CachingAutotuner):
414	def __init__(	414	def __init__(
415	self,	415	self,
416	- fn,	416	+ fn,
417	triton_meta, # passed directly to triton	417	triton_meta, # passed directly to triton
418	configs,	418	configs,
419	save_cache_hook,	419	save_cache_hook,
	@@ -500,7 +500,7 @@ class NPUCachingAutotuner(CachingAutotuner):
500	compile_meta["num_warps"] = cfg.num_warps	500	compile_meta["num_warps"] = cfg.num_warps
501	compile_meta["num_stages"] = cfg.num_stages	501	compile_meta["num_stages"] = cfg.num_stages
502	compile_meta["debug"] = (	502	compile_meta["debug"] = (
503	- os.getenv("INDUCTOR_ASCEND_DEBUG", 'false').lower() in ('true', '1')	503	+ os.getenv("INDUCTOR_ASCEND_DEBUG", 'false').lower() in ('true', '1')
504	and self.inductor_meta.get("assert_indirect_indexing", True)	504	and self.inductor_meta.get("assert_indirect_indexing", True)
505	and not self.inductor_meta.get("is_hip", False)	505	and not self.inductor_meta.get("is_hip", False)
506	)	506	)
	@@ -717,11 +717,11 @@ class NPUCachingAutotuner(CachingAutotuner):
717	kernel_id = -1	717	kernel_id = -1
718	if kernel_id not in fallback_id:	718	if kernel_id not in fallback_id:
719	return False	719	return False
720	- return True	720	+ return True
721		721
722	if not should_fallback():	722	if not should_fallback():
723	return None	723	return None
724	-	724	+
725	fx_graph_call, _, _, fx_module = get_triton_fx_graph_call(self.inductor_meta)	725	fx_graph_call, _, _, fx_module = get_triton_fx_graph_call(self.inductor_meta)
726	if not fx_graph_call:	726	if not fx_graph_call:
727	return None	727	return None
	@@ -743,7 +743,7 @@ class NPUCachingAutotuner(CachingAutotuner):
743	for arg in fx_args:	743	for arg in fx_args:
744	del arg	744	del arg
745	return True	745	return True
746	-	746	+
747	def debug_kernel_in_run(self, args, launcher, stream, *kwargs):	747	def debug_kernel_in_run(self, args, launcher, stream, *kwargs):
748	'''	748	'''
749	Save tensors for kernel args and outputs before and after kernel execute.	749	Save tensors for kernel args and outputs before and after kernel execute.
	@@ -1283,7 +1283,7 @@ def user_autotune_npu(
1283	inductor_meta=None,	1283	inductor_meta=None,
1284	custom_kernel=False,	1284	custom_kernel=False,
1285	):	1285	):
1286	-	1286	+
1287	if len(configs) == 0:	1287	if len(configs) == 0:
1288	configs = [triton.Config({})]	1288	configs = [triton.Config({})]
1289	else:	1289	else:

Mtorch_npu/asd/_silent_fault_data.py+1-1

	@@ -11,7 +11,7 @@ def parse_thresh(env_var_name, default_value, min_value):
11	thresh = [max(int(value), min_value) for value in thresh]	11	thresh = [max(int(value), min_value) for value in thresh]
12	if thresh[0] <= thresh[1]:	12	if thresh[0] <= thresh[1]:
13	thresh = [int(value) for value in default_value.split(",")]	13	thresh = [int(value) for value in default_value.split(",")]
14	-	14	+
15	return thresh	15	return thresh
16		16
17		17

Mtorch_npu/asd/asd.py+8-8

	@@ -55,7 +55,7 @@ class _SilentFaultDetector:
55	self.high_step = torch.tensor(self.min_step, dtype=torch.int32).npu()	55	self.high_step = torch.tensor(self.min_step, dtype=torch.int32).npu()
56	if grad.dtype == torch.float16:	56	if grad.dtype == torch.float16:
57	if not self.set_loss_scale_flag:	57	if not self.set_loss_scale_flag:
58	- return	58	+ return
59	else:	59	else:
60	grad = grad.float() / self.loss_scale	60	grad = grad.float() / self.loss_scale
61		61
	@@ -66,7 +66,7 @@ class _SilentFaultDetector:
66	self.silent_data_dict[idx] = SilentFaultData()	66	self.silent_data_dict[idx] = SilentFaultData()
67		67
68	sfda = self.silent_data_dict[idx]	68	sfda = self.silent_data_dict[idx]
69	-	69	+
70	if self.global_step <= self.min_step:	70	if self.global_step <= self.min_step:
71	self.step += 1	71	self.step += 1
72	self.global_step = self.step // (len(self.silent_data_dict) + 1)	72	self.global_step = self.step // (len(self.silent_data_dict) + 1)
	@@ -427,13 +427,13 @@ class _MatmulSilentCheck:
427	def parameter_filtering(self):	427	def parameter_filtering(self):
428	self.filter_index = (self.filter_index + 1) % self.filter_interval	428	self.filter_index = (self.filter_index + 1) % self.filter_interval
429	return self.filter_index == 0	429	return self.filter_index == 0
430	-	430	+
431	def register_module_hook(self, module, name):	431	def register_module_hook(self, module, name):
432	self.check_stat[name + "_backward"] = {'avg': 0, 'pre_val': 0, 'step': 0, 'none_zero_step': 0}	432	self.check_stat[name + "_backward"] = {'avg': 0, 'pre_val': 0, 'step': 0, 'none_zero_step': 0}
433	hook = partial(self.module_hook, name=name + "_backward")	433	hook = partial(self.module_hook, name=name + "_backward")
434	self.hook_dict[name + "_backward"] = module.register_full_backward_hook(hook)	434	self.hook_dict[name + "_backward"] = module.register_full_backward_hook(hook)
435	self.registered_modules.append(name)	435	self.registered_modules.append(name)
436	-	436	+
437	def module_hook(self, module, grad_input, grad_output, name):	437	def module_hook(self, module, grad_input, grad_output, name):
438	for _, param in module.named_parameters():	438	for _, param in module.named_parameters():
439	if param.dim() >= 2:	439	if param.dim() >= 2:
	@@ -706,7 +706,7 @@ class _MatmulSilentCheck:
706		706
707	while int(self.store.get('counter2').decode()) < world_size and self.checksum_state_thread_running:	707	while int(self.store.get('counter2').decode()) < world_size and self.checksum_state_thread_running:
708	time.sleep(0.1)	708	time.sleep(0.1)
709	-	709	+
710	if self.rank == 0:	710	if self.rank == 0:
711	self.store.add('checksum_state', 0 - global_state)	711	self.store.add('checksum_state', 0 - global_state)
712	self.store.add('counter', 0 - world_size)	712	self.store.add('counter', 0 - world_size)
	@@ -719,12 +719,12 @@ class _MatmulSilentCheck:
719	state['_lock'] = None	719	state['_lock'] = None
720	state['store'] = None	720	state['store'] = None
721	return state	721	return state
722	-	722	+
723	def __setstate(self, state):	723	def __setstate(self, state):
724	self.__dict__.update(state)	724	self.__dict__.update(state)
725	self.store = None	725	self.store = None
726		726
727	- def _startup(self):	727	+ def _startup(self):
728	if not self.check_thread_running:	728	if not self.check_thread_running:
729	self.check_thread_running = True	729	self.check_thread_running = True
730	self.check_thread = threading.Thread(	730	self.check_thread = threading.Thread(
	@@ -826,7 +826,7 @@ def _matmul_silent_check_decorator(func):
826	matmul_check.init_marks[matmul_check.first_module_id] = True	826	matmul_check.init_marks[matmul_check.first_module_id] = True
827		827
828	tmp = func(self, args, *kwargs)	828	tmp = func(self, args, *kwargs)
829	-	829	+
830	if matmul_check.get_matmul_hook_enable():	830	if matmul_check.get_matmul_hook_enable():
831	if hasattr(self, "matmul_check_outer") and self.matmul_check_outer:	831	if hasattr(self, "matmul_check_outer") and self.matmul_check_outer:
832	matmul_check.init_param()	832	matmul_check.init_param()

Mtorch_npu/csrc/aten/ops/CopyFromAndResizeKernelNpu.cpp+1-1

	@@ -8,7 +8,7 @@ at::Tensor NPUNativeFunctions::_copy_from_and_resize(const at::Tensor& self, con
8	{	8	{
9	TORCH_CHECK(dst.defined(), "dst is undefined", OPS_ERROR(ErrCode::NOT_SUPPORT));	9	TORCH_CHECK(dst.defined(), "dst is undefined", OPS_ERROR(ErrCode::NOT_SUPPORT));
10	TORCH_CHECK(self.defined(), "self is undefined", OPS_ERROR(ErrCode::NOT_SUPPORT));	10	TORCH_CHECK(self.defined(), "self is undefined", OPS_ERROR(ErrCode::NOT_SUPPORT));
11	-	11	+
12	if (dst.numel() == 0) {	12	if (dst.numel() == 0) {
13	dst.resize_as_(self);	13	dst.resize_as_(self);
14	}	14	}

Mtorch_npu/csrc/core/npu/GetCANNInfo.cpp+5-5

	@@ -253,7 +253,7 @@ int64_t VersionV2ToNum(std::string versionStr) {
253	parsed = true;	253	parsed = true;
254	}	254	}
255		255
256	- if (!parsed && tokens.size() == tokenNum4) { // ([0-9]+).([0-9]+).([0-9]+)-alpha, ([0-9]+).([0-9]+).([0-9]+)-beta, ([0-9]+).([0-9]+).([0-9]+)-rc,	256	+ if (!parsed && tokens.size() == tokenNum4) { // ([0-9]+).([0-9]+).([0-9]+)-alpha, ([0-9]+).([0-9]+).([0-9]+)-beta, ([0-9]+).([0-9]+).([0-9]+)-rc,
257	// ([0-9]+).([0-9]+).([0-9]+).alpha([0-9]+), ([0-9]+).([0-9]+).([0-9]+).beta([0-9]+), ([0-9]+).([0-9]+).([0-9]+).rc([0-9]+)	257	// ([0-9]+).([0-9]+).([0-9]+).alpha([0-9]+), ([0-9]+).([0-9]+).([0-9]+).beta([0-9]+), ([0-9]+).([0-9]+).([0-9]+).rc([0-9]+)
258	parsed = true;	258	parsed = true;
259	if (tokens[index3] == "alpha") {	259	if (tokens[index3] == "alpha") {
	@@ -469,7 +469,7 @@ std::string GetCANNVersion(const std::string& module)
469	module_version = version.version;	469	module_version = version.version;
470	CANNVersionCache[module] = module_version;	470	CANNVersionCache[module] = module_version;
471	}	471	}
472	-	472	+
473	if (find_module_v2 != pkgNameV2Map.end()) {	473	if (find_module_v2 != pkgNameV2Map.end()) {
474	char versionStr[ACL_PKG_VERSION_MAX_SIZE] = {0};	474	char versionStr[ACL_PKG_VERSION_MAX_SIZE] = {0};
475	aclError retV2 = c10_npu::acl::AclsysGetVersionStr(const_cast<char*>(module.c_str()), versionStr);	475	aclError retV2 = c10_npu::acl::AclsysGetVersionStr(const_cast<char*>(module.c_str()), versionStr);
	@@ -481,7 +481,7 @@ std::string GetCANNVersion(const std::string& module)
481	module_version = versionStr;	481	module_version = versionStr;
482	CANNVersionCache[module] = module_version;	482	CANNVersionCache[module] = module_version;
483	}	483	}
484	-	484	+
485	return module_version;	485	return module_version;
486	}	486	}
487		487
	@@ -492,7 +492,7 @@ bool IsGteCANNVersion(const std::string version, const std::string module)
492	if (module.compare(unsupportedModule) == 0) {	492	if (module.compare(unsupportedModule) == 0) {
493	TORCH_CHECK(false, "When the module is DRIVER, IsGteCANNVersion is not supported. ", PTA_ERROR(ErrCode::VALUE));	493	TORCH_CHECK(false, "When the module is DRIVER, IsGteCANNVersion is not supported. ", PTA_ERROR(ErrCode::VALUE));
494	}	494	}
495	-	495	+
496	std::vector<std::string> tokensVersion = SplitVersionStr(version);	496	std::vector<std::string> tokensVersion = SplitVersionStr(version);
497	std::vector<std::string> tokensBaseVersion = SplitVersionStr(baseVersion);	497	std::vector<std::string> tokensBaseVersion = SplitVersionStr(baseVersion);
498	if (tokensVersion.size() < tokenNum2 \|\| !isDigits(tokensVersion[index0])) {	498	if (tokensVersion.size() < tokenNum2 \|\| !isDigits(tokensVersion[index0])) {
	@@ -507,7 +507,7 @@ bool IsGteCANNVersion(const std::string version, const std::string module)
507		507
508	std::string currentVersion = GetCANNVersion(module);	508	std::string currentVersion = GetCANNVersion(module);
509	std::vector<std::string> tokensCurrentVersion = SplitVersionStr(currentVersion);	509	std::vector<std::string> tokensCurrentVersion = SplitVersionStr(currentVersion);
510	-	510	+
511	int64_t current_num = 0;	511	int64_t current_num = 0;
512	int64_t boundary_num = 0;	512	int64_t boundary_num = 0;
513	bool isInvalid = false;	513	bool isInvalid = false;

Mtorch_npu/csrc/core/npu/NPUCachingAllocator.cpp+1-1

	@@ -3589,7 +3589,7 @@ public:
3589	TORCH_NPU_MEMORY_LOGE("%s", retmsg.c_str());	3589	TORCH_NPU_MEMORY_LOGE("%s", retmsg.c_str());
3590	TORCH_CHECK_WITH(OutOfMemoryError, false, retmsg.c_str());	3590	TORCH_CHECK_WITH(OutOfMemoryError, false, retmsg.c_str());
3591	}	3591	}
3592	-	3592	+
3593	int device = 0;	3593	int device = 0;
3594	NPU_CHECK_ERROR(c10_npu::GetDevice(&device));	3594	NPU_CHECK_ERROR(c10_npu::GetDevice(&device));
3595	LazySetDevice(device);	3595	LazySetDevice(device);

Mtorch_npu/csrc/core/npu/NPUErrorCodes.h+1-1

	@@ -3,7 +3,7 @@
3	#include <unordered_map>	3	#include <unordered_map>
4		4
5	namespace c10_npu::acl {	5	namespace c10_npu::acl {
6	-	6	+
7	class AclErrorCode {	7	class AclErrorCode {
8	public:	8	public:
9	std::unordered_map<int, std::string> error_code_map = {	9	std::unordered_map<int, std::string> error_code_map = {

Mtorch_npu/csrc/core/npu/NPUPeerToPeerAccess.cpp+1-1

	@@ -50,7 +50,7 @@ bool NpuP2pCtrl::get_p2p_access(int32_t source_dev, int32_t dest_dev, bool& flag
50	// get access source_dev -> dest_dev	50	// get access source_dev -> dest_dev
51	auto &cache_s2d = p2p_access_enabled_cache_[source_dev * num_devices_ + dest_dev];	51	auto &cache_s2d = p2p_access_enabled_cache_[source_dev * num_devices_ + dest_dev];
52	auto &cache_d2s = p2p_access_enabled_cache_[dest_dev * num_devices_ + source_dev];	52	auto &cache_d2s = p2p_access_enabled_cache_[dest_dev * num_devices_ + source_dev];
53	-	53	+
54	if (cache_s2d != P2pStatus::UNKONWN) {	54	if (cache_s2d != P2pStatus::UNKONWN) {
55	return static_cast<bool>(cache_s2d);	55	return static_cast<bool>(cache_s2d);
56	}	56	}

Mtorch_npu/csrc/core/npu/NPUQueue.cpp+1-1

	@@ -607,7 +607,7 @@ void Repository::Enqueue(void *cur_paras)
607	#ifndef BUILD_LIBTORCH	607	#ifndef BUILD_LIBTORCH
608	// double check the current thread hold a Gil lock	608	// double check the current thread hold a Gil lock
609	// and release the GIL to TE op compiler in case the acl thread deadlock.	609	// and release the GIL to TE op compiler in case the acl thread deadlock.
610	- // However, this operator could produce another form of deadlock.	610	+ // However, this operator could produce another form of deadlock.
611	// When thread A deconstract a tensor, it will hold the mutex of deviceCachingAllocator and insert an event into the taskqueue.	611	// When thread A deconstract a tensor, it will hold the mutex of deviceCachingAllocator and insert an event into the taskqueue.
612	// If the taskqueue is full, thead A will run into here and release the GIL.	612	// If the taskqueue is full, thead A will run into here and release the GIL.
613	// Once another thread B get GIL and trigger GC, it may deconstract another tensor	613	// Once another thread B get GIL and trigger GC, it may deconstract another tensor

Mtorch_npu/csrc/core/npu/NPUWorkspaceAllocator.cpp+1-1

	@@ -512,7 +512,7 @@ public:
512	if (c10_npu::option::OptionsManager::CheckForceUncached() &&	512	if (c10_npu::option::OptionsManager::CheckForceUncached() &&
513	(c10_npu::currentStreamCaptureStatus() == c10_npu::CaptureStatus::None)) {	513	(c10_npu::currentStreamCaptureStatus() == c10_npu::CaptureStatus::None)) {
514	return &uncached_delete;	514	return &uncached_delete;
515	- }	515	+ }
516	return &local_raw_delete;	516	return &local_raw_delete;
517	}	517	}
518		518

Mtorch_npu/csrc/core/npu/NpuVariables.cpp+1-1

	@@ -49,7 +49,7 @@ void SetSocVersion(const char* const socVersion)
49	SocVersion curSocVersion = SocVersion::UnsupportedSocVersion;	49	SocVersion curSocVersion = SocVersion::UnsupportedSocVersion;
50	std::string inputVersion = socVersion;	50	std::string inputVersion = socVersion;
51	std::string ascend950 = "Ascend950";	51	std::string ascend950 = "Ascend950";
52	-	52	+
53	auto const& iter = socVersionMap.find(socVersion);	53	auto const& iter = socVersionMap.find(socVersion);
54	if (iter != socVersionMap.end()) {	54	if (iter != socVersionMap.end()) {
55	curSocVersion = iter->second;	55	curSocVersion = iter->second;

Mtorch_npu/csrc/core/npu/register/OptionsManager.cpp+1-1

	@@ -520,7 +520,7 @@ uint32_t OptionsManager::GetAclOpInitMode()
520	} else {	520	} else {
521	acl_op_init_mode_ = (buf_val != nullptr) ? strtol(buf_val, nullptr, 10) : 0;	521	acl_op_init_mode_ = (buf_val != nullptr) ? strtol(buf_val, nullptr, 10) : 0;
522	}	522	}
523	-	523	+
524	std::unordered_map<int32_t, std::string> aclOpInitMode = getAclOpInitMode();	524	std::unordered_map<int32_t, std::string> aclOpInitMode = getAclOpInitMode();
525	if (aclOpInitMode.find(acl_op_init_mode_) == aclOpInitMode.end()) {	525	if (aclOpInitMode.find(acl_op_init_mode_) == aclOpInitMode.end()) {
526	if (default_value_acl_mode && isCannVersionGteBase) {	526	if (default_value_acl_mode && isCannVersionGteBase) {

Mtorch_npu/csrc/distributed/HCCLUtils.cpp+1-1

	@@ -112,7 +112,7 @@ bool isSupportHcclCommName()
112		112
113	HCCLComm::HCCLComm(HcclComm hcclComm) : hcclComm_(hcclComm), hcclAsyncErr_(HCCL_SUCCESS),	113	HCCLComm::HCCLComm(HcclComm hcclComm) : hcclComm_(hcclComm), hcclAsyncErr_(HCCL_SUCCESS),
114	hcclCommType(0), p2pPeer(0) {}	114	hcclCommType(0), p2pPeer(0) {}
115	-	115	+
116	HCCLComm::~HCCLComm()	116	HCCLComm::~HCCLComm()
117	{	117	{
118	destroyHcclComm();	118	destroyHcclComm();

Mtorch_npu/csrc/distributed/Init.cpp+2-2

	@@ -210,7 +210,7 @@ PyObject* c10d_npu_init(PyObject* _unused, PyObject* noargs)
210	throw python_error();	210	throw python_error();
211	}	211	}
212	auto torch_npu_C_m = py::handle(torch_npu_C_module).cast<py::module>();	212	auto torch_npu_C_m = py::handle(torch_npu_C_module).cast<py::module>();
213	-	213	+
214	auto m =	214	auto m =
215	torch_npu_C_m.def_submodule("_distributed_c10d", "distributed c10d bindings");	215	torch_npu_C_m.def_submodule("_distributed_c10d", "distributed c10d bindings");
216	auto module = py::handle(m).cast<py::module>();	216	auto module = py::handle(m).cast<py::module>();
	@@ -473,7 +473,7 @@ PyObject* c10d_npu_init(PyObject* _unused, PyObject* noargs)
473	.def_readwrite("hccl_config", &::c10d_npu::ProcessGroupHCCL::Options::hccl_config)	473	.def_readwrite("hccl_config", &::c10d_npu::ProcessGroupHCCL::Options::hccl_config)
474	.def_readwrite("group_id",	474	.def_readwrite("group_id",
475	&::c10d_npu::ProcessGroupHCCL::Options::group_id);	475	&::c10d_npu::ProcessGroupHCCL::Options::group_id);
476	-	476	+
477	// bind for ProcessGroupLCCL	477	// bind for ProcessGroupLCCL
478	auto processGroupLCCL = intrusive_ptr_no_gil_destructor_class_<::c10d_npu::ProcessGroupLCCL>(	478	auto processGroupLCCL = intrusive_ptr_no_gil_destructor_class_<::c10d_npu::ProcessGroupLCCL>(
479	module, "ProcessGroupLCCL", dist.attr("Backend"))	479	module, "ProcessGroupLCCL", dist.attr("Backend"))

Mtorch_npu/csrc/distributed/ParallelTcpServer.cpp+3-3

	@@ -232,7 +232,7 @@ int ParallelTcpServer::CreateSocket(const std::string host, uint16_t port) noexc
232	if (sockFd >= 0) {	232	if (sockFd >= 0) {
233	return sockFd;	233	return sockFd;
234	}	234	}
235	-	235	+
236	sockFd = CreateSocketWithFamily(host, port, AF_INET6);	236	sockFd = CreateSocketWithFamily(host, port, AF_INET6);
237	if (sockFd >= 0) {	237	if (sockFd >= 0) {
238	return sockFd;	238	return sockFd;
	@@ -298,7 +298,7 @@ int ParallelTcpServer::CreateLocalSocket(const std::string &localSocketPath) noe
298	LOG(ERROR) << "local socket path invalid." << errno << " : " << strerror(errno);	298	LOG(ERROR) << "local socket path invalid." << errno << " : " << strerror(errno);
299	return -1;	299	return -1;
300	}	300	}
301	-	301	+
302	struct sockaddr_un servAddr {};	302	struct sockaddr_un servAddr {};
303	servAddr.sun_family = AF_UNIX;	303	servAddr.sun_family = AF_UNIX;
304	servAddr.sun_path[0] = '\0';	304	servAddr.sun_path[0] = '\0';
	@@ -386,7 +386,7 @@ int ParallelTcpServer::SetBlockSocketTimeout(int fd) noexcept
386	LOG(ERROR) << "set block accept timeout failed " << errno << " : " << strerror(errno);	386	LOG(ERROR) << "set block accept timeout failed " << errno << " : " << strerror(errno);
387	return -1;	387	return -1;
388	}	388	}
389	-	389	+
390	return 0;	390	return 0;
391	}	391	}
392		392

Mtorch_npu/csrc/distributed/ProcessGroupHCCL.cpp+13-13

	@@ -154,7 +154,7 @@ HcclReduceOp getHcclReduceOp(const c10d::ReduceOp reduceOp, at::Tensor& input)
154	// represent a bool (see hcclDataType mapping).	154	// represent a bool (see hcclDataType mapping).
155	return HCCL_REDUCE_MAX;	155	return HCCL_REDUCE_MAX;
156	}	156	}
157	-	157	+
158	if (unsupportedOp.find(reduceOp) != unsupportedOp.end()) {	158	if (unsupportedOp.find(reduceOp) != unsupportedOp.end()) {
159	TORCH_CHECK(false,	159	TORCH_CHECK(false,
160	"Cannot use ReduceOp." + unsupportedOp[reduceOp] + " with HCCL",	160	"Cannot use ReduceOp." + unsupportedOp[reduceOp] + " with HCCL",
	@@ -1289,7 +1289,7 @@ void ProcessGroupHCCL::waitForFutureOrTimeout(
1289	void ProcessGroupHCCL::shutdown()	1289	void ProcessGroupHCCL::shutdown()
1290	{	1290	{
1291	LOG(INFO) << logPrefix() << "Starting to destroy process group, flushing operations.";	1291	LOG(INFO) << logPrefix() << "Starting to destroy process group, flushing operations.";
1292	-	1292	+
1293	if (terminateProcessGroup_.exchange(true)) {	1293	if (terminateProcessGroup_.exchange(true)) {
1294	return;	1294	return;
1295	}	1295	}
	@@ -1377,7 +1377,7 @@ void ProcessGroupHCCL::deleteTCPStoreKey()
1377	}	1377	}
1378		1378
1379	TORCH_NPU_HCCL_LOGI("Delete TCP store key success.");	1379	TORCH_NPU_HCCL_LOGI("Delete TCP store key success.");
1380	-	1380	+
1381	TCPStoreKeyList_.clear();	1381	TCPStoreKeyList_.clear();
1382	}	1382	}
1383		1383
	@@ -1951,7 +1951,7 @@ void ProcessGroupHCCL::logWorkEnd(WorkHCCL& work)
1951		1951
1952	storeError_ = !c10d::traceUpdate(store_, traceKeyEnd_, work.seq_, opTypeToString(work.opType_));	1952	storeError_ = !c10d::traceUpdate(store_, traceKeyEnd_, work.seq_, opTypeToString(work.opType_));
1953	}	1953	}
1954	-	1954	+
1955	std::string ProcessGroupHCCL::createLogPrefix() const	1955	std::string ProcessGroupHCCL::createLogPrefix() const
1956	{	1956	{
1957	if (!pg_desc_.empty() && pg_desc_ != "undefined") {	1957	if (!pg_desc_.empty() && pg_desc_ != "undefined") {
	@@ -2040,7 +2040,7 @@ void ProcessGroupHCCL::Watchdog::runLoop()
2040	auto timenow = std::chrono::steady_clock::now();	2040	auto timenow = std::chrono::steady_clock::now();
2041	bool recordflag = false;	2041	bool recordflag = false;
2042	int kThousandMillis = 1000;	2042	int kThousandMillis = 1000;
2043	-	2043	+
2044	while (!pg_->terminateProcessGroup_.load()) {	2044	while (!pg_->terminateProcessGroup_.load()) {
2045	if (status_save_enable) {	2045	if (status_save_enable) {
2046	checkAndMakePath(status_save_path.c_str(), "Open shared directory failed. Please check whether input path is valid.");	2046	checkAndMakePath(status_save_path.c_str(), "Open shared directory failed. Please check whether input path is valid.");
	@@ -2085,7 +2085,7 @@ void ProcessGroupHCCL::Watchdog::runLoop()
2085	TORCH_NPU_HCCL_LOGI("Find FORCE STOP when runloop setDevice.");	2085	TORCH_NPU_HCCL_LOGI("Find FORCE STOP when runloop setDevice.");
2086	}	2086	}
2087	}	2087	}
2088	-	2088	+
2089	// check NCCL errors first	2089	// check NCCL errors first
2090	if (!pg_->terminateProcessGroup_.load()) {	2090	if (!pg_->terminateProcessGroup_.load()) {
2091	work.checkAndSetException();	2091	work.checkAndSetException();
	@@ -3960,7 +3960,7 @@ c10::intrusive_ptr<c10d::Work> ProcessGroupHCCL::collective(
3960		3960
3961	const std::vector<uint32_t>& ranks = groupRanks();	3961	const std::vector<uint32_t>& ranks = groupRanks();
3962	outfile << "[GLOBAL RANKID]:" << ranks[rank_] << "\n";	3962	outfile << "[GLOBAL RANKID]:" << ranks[rank_] << "\n";
3963	-	3963	+
3964	outfile.close();	3964	outfile.close();
3965	}	3965	}
3966	} else {	3966	} else {
	@@ -4064,7 +4064,7 @@ c10::intrusive_ptr<c10d::Work> ProcessGroupHCCL::collective(
4064	} else {	4064	} else {
4065	c10_npu::NPUGraph::dec_pending_event_queries();	4065	c10_npu::NPUGraph::dec_pending_event_queries();
4066	}	4066	}
4067	-	4067	+
4068	return work;	4068	return work;
4069	}	4069	}
4070		4070
	@@ -4196,7 +4196,7 @@ c10::intrusive_ptr<c10d::Work> ProcessGroupHCCL::collectiveCoalesced(
4196		4196
4197	const std::vector<uint32_t>& ranks = groupRanks();	4197	const std::vector<uint32_t>& ranks = groupRanks();
4198	outfile << "[GLOBAL RANKID]:" << ranks[rank_] << "\n";	4198	outfile << "[GLOBAL RANKID]:" << ranks[rank_] << "\n";
4199	-	4199	+
4200	outfile.close();	4200	outfile.close();
4201	}	4201	}
4202	} else {	4202	} else {
	@@ -4281,7 +4281,7 @@ c10::intrusive_ptr<c10d::Work> ProcessGroupHCCL::collectiveCoalesced(
4281	} else {	4281	} else {
4282	c10_npu::NPUGraph::dec_pending_event_queries();	4282	c10_npu::NPUGraph::dec_pending_event_queries();
4283	}	4283	}
4284	-	4284	+
4285	return work;	4285	return work;
4286	}	4286	}
4287		4287
	@@ -4352,7 +4352,7 @@ c10::intrusive_ptr<c10d::Work> ProcessGroupHCCL::pointToPoint(
4352	"Got device ", device.index(), " but expected ", coalescedDevice_.index());	4352	"Got device ", device.index(), " but expected ", coalescedDevice_.index());
4353	}	4353	}
4354	}	4354	}
4355	-	4355	+
4356	// Verify communicator consistency	4356	// Verify communicator consistency
4357	if (coalescedComm_ == nullptr) {	4357	if (coalescedComm_ == nullptr) {
4358	coalescedComm_ = hcclComms[0];	4358	coalescedComm_ = hcclComms[0];
	@@ -4452,7 +4452,7 @@ c10::intrusive_ptr<c10d::Work> ProcessGroupHCCL::pointToPoint(
4452		4452
4453	const std::vector<uint32_t>& ranks = groupRanks();	4453	const std::vector<uint32_t>& ranks = groupRanks();
4454	outfile << "[GLOBAL RANKID]:" << ranks[rank_] << "\n";	4454	outfile << "[GLOBAL RANKID]:" << ranks[rank_] << "\n";
4455	-	4455	+
4456	outfile.close();	4456	outfile.close();
4457	}	4457	}
4458	} else {	4458	} else {
	@@ -4532,7 +4532,7 @@ c10::intrusive_ptr<c10d::Work> ProcessGroupHCCL::pointToPoint(
4532	// as multi-device per process is deprecated	4532	// as multi-device per process is deprecated
4533	work->numelIn_ = work->numelOut_ = static_cast<size_t>(tensors[i].numel());	4533	work->numelIn_ = work->numelOut_ = static_cast<size_t>(tensors[i].numel());
4534	}	4534	}
4535	-	4535	+
4536	c10_npu::NPUGraph::inc_pending_event_queries();	4536	c10_npu::NPUGraph::inc_pending_event_queries();
4537	if (asyncErrorHandling_ != NoHandling && capture_status == c10_npu::CaptureStatus::None) {	4537	if (asyncErrorHandling_ != NoHandling && capture_status == c10_npu::CaptureStatus::None) {
4538	workEnqueue(work);	4538	workEnqueue(work);

Mtorch_npu/csrc/distributed/ProcessGroupHCCL.hpp+12-12

	@@ -400,7 +400,7 @@ public:
400	// for tests.	400	// for tests.
401	virtual std::exception_ptr checkForHCCLErrors(	401	virtual std::exception_ptr checkForHCCLErrors(
402	const std::vector<std::shared_ptr<HCCLComm>>& hcclComms) const;	402	const std::vector<std::shared_ptr<HCCLComm>>& hcclComms) const;
403	-	403	+
404	friend std::ostream& operator<<(	404	friend std::ostream& operator<<(
405	std::ostream& output,	405	std::ostream& output,
406	const WorkHCCL& workHCCL);	406	const WorkHCCL& workHCCL);
	@@ -452,11 +452,11 @@ public:
452	std::vector<std::pair<c10::weak_intrusive_ptr<c10::StorageImpl>, c10_npu::NPUStream>> recorded_outputs_;	452	std::vector<std::pair<c10::weak_intrusive_ptr<c10::StorageImpl>, c10_npu::NPUStream>> recorded_outputs_;
453		453
454	std::vector<at::Tensor> lazy_destroy_tensors_;	454	std::vector<at::Tensor> lazy_destroy_tensors_;
455	-	455	+
456	// unique id used to tell the trace buffer that this	456	// unique id used to tell the trace buffer that this
457	// work has completed	457	// work has completed
458	c10::optional<uint64_t> trace_id_;	458	c10::optional<uint64_t> trace_id_;
459	-	459	+
460	mutable std::once_flag print_flag;	460	mutable std::once_flag print_flag;
461		461
462	friend class ProcessGroupHCCL;	462	friend class ProcessGroupHCCL;
	@@ -598,7 +598,7 @@ public:
598	{	598	{
599	return std::string(HCCL_BACKEND_NAME);	599	return std::string(HCCL_BACKEND_NAME);
600	}	600	}
601	-	601	+
602	bool supportsCoalescing() const override	602	bool supportsCoalescing() const override
603	{	603	{
604	return true;	604	return true;
	@@ -869,7 +869,7 @@ protected:
869	{	869	{
870	return pg_desc_;	870	return pg_desc_;
871	}	871	}
872	-	872	+
873	void setP2pPeer(int newPeer)	873	void setP2pPeer(int newPeer)
874	{	874	{
875	peer_ = newPeer;	875	peer_ = newPeer;
	@@ -879,7 +879,7 @@ protected:
879	{	879	{
880	return peer_;	880	return peer_;
881	}	881	}
882	-	882	+
883	// In the timeout case and we will dump debug info such as the NCCL flight	883	// In the timeout case and we will dump debug info such as the NCCL flight
884	// recorder to storage. Down the road, if we have more complicated or blocking	884	// recorder to storage. Down the road, if we have more complicated or blocking
885	// operations, we might need to use a side thread to do it.	885	// operations, we might need to use a side thread to do it.
	@@ -895,7 +895,7 @@ protected:
895	// we can dump the debugging information and abort the process.	895	// we can dump the debugging information and abort the process.
896	virtual void heartbeatMonitor();	896	virtual void heartbeatMonitor();
897	#endif	897	#endif
898	-	898	+
899	// Instance of the watchdog thread.	899	// Instance of the watchdog thread.
900	std::unique_ptr<Watchdog> watchdog_;	900	std::unique_ptr<Watchdog> watchdog_;
901	// Function that directly trigger std::abort so that the whole process	901	// Function that directly trigger std::abort so that the whole process
	@@ -982,7 +982,7 @@ protected:
982	std::unordered_map<std::string, std::vector<std::shared_ptr<HCCLComm>>> devHCCLCommMap_;	982	std::unordered_map<std::string, std::vector<std::shared_ptr<HCCLComm>>> devHCCLCommMap_;
983		983
984	std::unordered_set<std::string> reportedErrorComms_;	984	std::unordered_set<std::string> reportedErrorComms_;
985	-	985	+
986	std::unordered_map<int, std::vector<std::string>> p2pSendRecvKeys_;	986	std::unordered_map<int, std::vector<std::string>> p2pSendRecvKeys_;
987		987
988	std::unordered_map<std::string, std::string> devHCCLCommNameMap_;	988	std::unordered_map<std::string, std::string> devHCCLCommNameMap_;
	@@ -1075,7 +1075,7 @@ protected:
1075	// The NPU events used to control task rate to protect streams	1075	// The NPU events used to control task rate to protect streams
1076	std::unordered_map<std::string, std::vector<c10_npu::NPUEvent>>	1076	std::unordered_map<std::string, std::vector<c10_npu::NPUEvent>>
1077	rateCtrlEvents_;	1077	rateCtrlEvents_;
1078	-	1078	+
1079	std::unordered_map<std::string, std::vector<uint64_t>> collectiveCnts_;	1079	std::unordered_map<std::string, std::vector<uint64_t>> collectiveCnts_;
1080		1080
1081	// Device Indexes used for all collectives in this group	1081	// Device Indexes used for all collectives in this group
	@@ -1232,7 +1232,7 @@ private:
1232	Fn fn,	1232	Fn fn,
1233	c10d::OpType opType,	1233	c10d::OpType opType,
1234	bool asyncOp = false);	1234	bool asyncOp = false);
1235	-	1235	+
1236	template <typename Fn, typename PreProcess, typename PostProcess>	1236	template <typename Fn, typename PreProcess, typename PostProcess>
1237	c10::intrusive_ptr<c10d::Work> collective(	1237	c10::intrusive_ptr<c10d::Work> collective(
1238	std::vector<at::Tensor>& input,	1238	std::vector<at::Tensor>& input,
	@@ -1355,7 +1355,7 @@ private:
1355	// Util function to assign timeout to each work.	1355	// Util function to assign timeout to each work.
1356	void assignTimeoutToWork(const c10::intrusive_ptr<ProcessGroupHCCL::WorkHCCL>& work,	1356	void assignTimeoutToWork(const c10::intrusive_ptr<ProcessGroupHCCL::WorkHCCL>& work,
1357	const c10::intrusive_ptr<Options>& option);	1357	const c10::intrusive_ptr<Options>& option);
1358	-	1358	+
1359	void silenceCheck(at::Tensor &input, c10d::OpType opType);	1359	void silenceCheck(at::Tensor &input, c10d::OpType opType);
1360		1360
1361	HcclCommConfig createHcclCommConfigWithOptions();	1361	HcclCommConfig createHcclCommConfigWithOptions();
	@@ -1382,7 +1382,7 @@ private:
1382	c10::optional<at::Tensor> windowMem_;	1382	c10::optional<at::Tensor> windowMem_;
1383		1383
1384	uint32_t cached_aic_num;	1384	uint32_t cached_aic_num;
1385	-	1385	+
1386	uint32_t cached_aiv_num;	1386	uint32_t cached_aiv_num;
1387		1387
1388	};	1388	};

Mtorch_npu/csrc/distributed/StoreClient.cpp+3-3

	@@ -30,7 +30,7 @@
30		30
31	namespace c10d {	31	namespace c10d {
32	namespace torch_npu {	32	namespace torch_npu {
33	-	33	+
34	Client::Client(const std::string host, uint16_t port, const std::chrono::milliseconds timeout) noexcept	34	Client::Client(const std::string host, uint16_t port, const std::chrono::milliseconds timeout) noexcept
35	: host_{ host }, port_{ port }, socketFd_(-1), timeout_{ timeout }	35	: host_{ host }, port_{ port }, socketFd_(-1), timeout_{ timeout }
36	{}	36	{}
	@@ -111,7 +111,7 @@ int Client::Connect() noexcept
111	if (ret >= 0) {	111	if (ret >= 0) {
112	return 0;	112	return 0;
113	}	113	}
114	-	114	+
115	ret = TryConnect(AF_INET6);	115	ret = TryConnect(AF_INET6);
116	if (ret >= 0) {	116	if (ret >= 0) {
117	return 0;	117	return 0;
	@@ -221,7 +221,7 @@ int Client::SyncCall(const StoreMessage &request, StoreMessage &response) noexce
221	if (errno == EINTR) { // interrupted by signal	221	if (errno == EINTR) { // interrupted by signal
222	continue;	222	continue;
223	}	223	}
224	-	224	+
225	LOG(ERROR) << "read data from server(" << host_ << ":" << port_ << ") failed " << errno << " : " <<	225	LOG(ERROR) << "read data from server(" << host_ << ":" << port_ << ") failed " << errno << " : " <<
226	strerror(errno);	226	strerror(errno);
227	return -1;	227	return -1;

Mtorch_npu/csrc/framework/OpCommand.cpp+1-1

	@@ -243,7 +243,7 @@ void OpCommand::RunOpApiV2(const string &op_name, const PROC_FUNC &func, bool sy
243	execParams.customHandler = const_cast<PROC_FUNC*>(&func);	243	execParams.customHandler = const_cast<PROC_FUNC*>(&func);
244		244
245	c10_npu::queue::QueueParas params(c10_npu::queue::EXECUTE_OPAPI_V2, sizeof(ExecuteParasOpApiV2), &execParams);	245	c10_npu::queue::QueueParas params(c10_npu::queue::EXECUTE_OPAPI_V2, sizeof(ExecuteParasOpApiV2), &execParams);
246	-	246	+
247	auto start = std::chrono::steady_clock::now();	247	auto start = std::chrono::steady_clock::now();
248		248
249	c10_npu::enCurrentNPUStream(&params);	249	c10_npu::enCurrentNPUStream(&params);

Mtorch_npu/csrc/framework/OpCommand.h+2-2

	@@ -72,7 +72,7 @@ public:
72	toType);	72	toType);
73	return AddHostTensorInput(cpuTensor, compileType, realDtype, descName);	73	return AddHostTensorInput(cpuTensor, compileType, realDtype, descName);
74	}	74	}
75	-	75	+
76	// IntArrayRef/SmallVector Input, usually hostmemory input, we will do h2d in launch kernel	76	// IntArrayRef/SmallVector Input, usually hostmemory input, we will do h2d in launch kernel
77	OpCommand& Input(const c10::IntArrayRef &dimListRef,	77	OpCommand& Input(const c10::IntArrayRef &dimListRef,
78	at::ScalarType toType = at::kLong,	78	at::ScalarType toType = at::kLong,
	@@ -133,7 +133,7 @@ private:
133	OpCommand& AddTensorInput(at::Tensor &tensor,	133	OpCommand& AddTensorInput(at::Tensor &tensor,
134	at::ScalarType forceScaleType = at::ScalarType::Undefined,	134	at::ScalarType forceScaleType = at::ScalarType::Undefined,
135	const string &descName = "", const string &realData = "");	135	const string &descName = "", const string &realData = "");
136	-	136	+
137	OpCommand& AddTensorInput(const string &str);	137	OpCommand& AddTensorInput(const string &str);
138		138
139	OpCommand& AddHostTensorInput(	139	OpCommand& AddHostTensorInput(

Mtorch_npu/csrc/framework/utils/CalcuOpUtil.cpp+1-1

	@@ -390,7 +390,7 @@ at::ScalarType CalcuOpUtil::ConvertToScalarType(const aclDataType data_type)
390	std::string("aclDataType:") + std::to_string(data_type) + " has not been supported",	390	std::string("aclDataType:") + std::to_string(data_type) + " has not been supported",
391	OPS_ERROR(ErrCode::NOT_SUPPORT))	391	OPS_ERROR(ErrCode::NOT_SUPPORT))
392	}	392	}
393	-	393	+
394	return iter->second;	394	return iter->second;
395	}	395	}
396		396

Mtorch_npu/csrc/inductor/aoti_runner/model_container_runner.h+1-1

	@@ -70,7 +70,7 @@ protected:
70	decltype(&AOTInductorModelContainerGetCallSpec) get_call_spec_func_ { nullptr };	70	decltype(&AOTInductorModelContainerGetCallSpec) get_call_spec_func_ { nullptr };
71	decltype(&AOTInductorModelContainerGetConstantsBlobSize) get_constants_blob_size_func_{nullptr};	71	decltype(&AOTInductorModelContainerGetConstantsBlobSize) get_constants_blob_size_func_{nullptr};
72	decltype(&AOTInductorModelUpdateConstantsFromBlob) update_constants_from_blob_func_{nullptr};	72	decltype(&AOTInductorModelUpdateConstantsFromBlob) update_constants_from_blob_func_{nullptr};
73	-	73	+
74	AOTInductorModelContainerHandle container_handle_ = nullptr;	74	AOTInductorModelContainerHandle container_handle_ = nullptr;
75		75
76	AOTIProxyExecutorHandle proxy_executor_handle_;	76	AOTIProxyExecutorHandle proxy_executor_handle_;

Mtorch_npu/csrc/inductor/aoti_runtime/interface.h+1-1

	@@ -172,7 +172,7 @@ AOTIRuntimeError AOTInductorModelContainerGetConstantsBlobSize(AOTInductorModelC
172		172
173	AOTIRuntimeError AOTInductorModelUpdateConstantsFromBlob(AOTInductorModelContainerHandle container_handle,	173	AOTIRuntimeError AOTInductorModelUpdateConstantsFromBlob(AOTInductorModelContainerHandle container_handle,
174	const uint8_t* weight_blob_ptr);	174	const uint8_t* weight_blob_ptr);
175	-	175	+
176	// Delete an AOTInductorModel created by AOTInductorModelCreate.	176	// Delete an AOTInductorModel created by AOTInductorModelCreate.
177	AOTIRuntimeError AOTInductorModelDelete(AOTInductorModelHandle model_handle);	177	AOTIRuntimeError AOTInductorModelDelete(AOTInductorModelHandle model_handle);
178		178

Mtorch_npu/csrc/inductor/mlir/hacl_rt.h+1-1

	@@ -183,7 +183,7 @@ typedef struct tagRtArgsEx {
183	uint16_t tilingDataOffset; // size to tiling data	183	uint16_t tilingDataOffset; // size to tiling data
184	uint16_t hostInputInfoNum; // 0	184	uint16_t hostInputInfoNum; // 0
185	uint8_t hasTiling; // has tiling	185	uint8_t hasTiling; // has tiling
186	- uint8_t isNoNeedH2DCopy; // not need rtKernelLaunchWithFlag copy tiling from host to device	186	+ uint8_t isNoNeedH2DCopy; // not need rtKernelLaunchWithFlag copy tiling from host to device
187	uint8_t reserved[4];	187	uint8_t reserved[4];
188	} rtArgsEx_t;	188	} rtArgsEx_t;
189		189

Mtorch_npu/csrc/npu/Graph.cpp+1-1

	@@ -789,7 +789,7 @@ void TORCH_NPU_API THNPGraph_init(PyObject* module) {
789	helper.processStringArrayOption(key, item.second.cast<std::vector<std::string>>());	789	helper.processStringArrayOption(key, item.second.cast<std::vector<std::string>>());
790	} else if (py::isinstance<py::int_>(item.second)) {	790	} else if (py::isinstance<py::int_>(item.second)) {
791	helper.processInitOption(key, item.second.cast<int>());	791	helper.processInitOption(key, item.second.cast<int>());
792	- }	792	+ }
793	}	793	}
794	}	794	}
795		795

Mtorch_npu/csrc/npu/Stress_detect.cpp+1-1

	@@ -130,7 +130,7 @@ int StressDetector::perform_stress_detect(int deviceid, int mode, int64_t comm)
130		130
131	// Set task parameters	131	// Set task parameters
132	task_in_progress.store(true);	132	task_in_progress.store(true);
133	-	133	+
134	// Allocate workspace memory	134	// Allocate workspace memory
135	workspaceAddr = nullptr;	135	workspaceAddr = nullptr;
136	uint64_t size = 10;	136	uint64_t size = 10;

Mtorch_npu/csrc/npu/Stress_detect.h+3-3

	@@ -20,17 +20,17 @@ private:
20	static void worker_thread();	20	static void worker_thread();
21		21
22	static int transfer_result(int detectResult);	22	static int transfer_result(int detectResult);
23	-	23	+
24	// Thread for handling the stress detection task	24	// Thread for handling the stress detection task
25	static std::thread stress_detect_thread;	25	static std::thread stress_detect_thread;
26		26
27	// Condition variable and mutex to control the thread	27	// Condition variable and mutex to control the thread
28	static std::condition_variable cv;	28	static std::condition_variable cv;
29	static std::mutex mtx;	29	static std::mutex mtx;
30	-	30	+
31	// Flag to indicate if a task is in progress	31	// Flag to indicate if a task is in progress
32	static std::atomic<bool> task_in_progress;	32	static std::atomic<bool> task_in_progress;
33	-	33	+
34	// Flag to signal the thread to stop	34	// Flag to signal the thread to stop
35	static std::atomic<bool> stop_thread;	35	static std::atomic<bool> stop_thread;
36		36

Mtorch_npu/csrc/profiler/profiler_mgr.cpp+2-2

	@@ -95,7 +95,7 @@ void ProfilerMgr::EnableMsProfiler(uint32_t *deviceIdList, uint32_t deviceNum, a
95	if (profConfig_ == nullptr) {	95	if (profConfig_ == nullptr) {
96	profConfig_ = at_npu::native::AclProfilingCreateConfig(deviceIdList, deviceNum, aicMetrics, nullptr, dataTypeConfig);	96	profConfig_ = at_npu::native::AclProfilingCreateConfig(deviceIdList, deviceNum, aicMetrics, nullptr, dataTypeConfig);
97	}	97	}
98	-	98	+
99	if (profConfig_ == nullptr) {	99	if (profConfig_ == nullptr) {
100	ASCEND_LOGE("Create Prof Config failed.");	100	ASCEND_LOGE("Create Prof Config failed.");
101	return;	101	return;
	@@ -275,7 +275,7 @@ void ProfilerMgr::Stop()
275	StopDataReceiver();	275	StopDataReceiver();
276	profile_memory_.store(false);	276	profile_memory_.store(false);
277	}	277	}
278	-	278	+
279	if (npu_trace_.load()) {	279	if (npu_trace_.load()) {
280	at_npu::native::AclProfilingStop(profConfig_);	280	at_npu::native::AclProfilingStop(profConfig_);
281	auto ret = at_npu::native::AclProfilingDestroyConfig(profConfig_);	281	auto ret = at_npu::native::AclProfilingDestroyConfig(profConfig_);

Mtorch_npu/distributed/nn/functional.py+1-1

	@@ -8,7 +8,7 @@ from torch.distributed import ReduceOp
8	def _allgather_base_backward_hccl(ctx, grad_output):	8	def _allgather_base_backward_hccl(ctx, grad_output):
9	"""	9	"""
10	Backward function for _AllGatherBase that supports HCCL backend.	10	Backward function for _AllGatherBase that supports HCCL backend.
11	-	11	+
12	Original PyTorch implementation only supports NCCL backend.	12	Original PyTorch implementation only supports NCCL backend.
13	This version adds HCCL support for NPU devices.	13	This version adds HCCL support for NPU devices.
14	"""	14	"""

Mtorch_npu/distributed/run.py+1-1

	@@ -15,7 +15,7 @@ def parse_args(args):
15	action=env,	15	action=env,
16	type=str,	16	type=str,
17	default="false",	17	default="false",
18	- help="Turn parallel tcpstore tiered optimization, if true, The agent adds a proxy role,"	18	+ help="Turn parallel tcpstore tiered optimization, if true, The agent adds a proxy role,"
19	"the worker on this node will connect to the server through the proxy.",	19	"the worker on this node will connect to the server through the proxy.",
20	)	20	)
21	return parser.parse_args(args)	21	return parser.parse_args(args)

Mtorch_npu/distributed/tensor/_attention.py+3-3

Mtorch_npu/distributed/tensor/_math_ops.py+4-4

	@@ -402,7 +402,7 @@ def custom_npu_conv2d_strategy(x, weight, bias, stride, padding, dilation, group
402	]	402	]
403	)	403	)
404	acceptable_shardings.append(replicate_strategy)	404	acceptable_shardings.append(replicate_strategy)
405	-	405	+
406	# x layout: (N, Ci, Hi, Wi)	406	# x layout: (N, Ci, Hi, Wi)
407	# weight layout: (Co, Ci/groups, Hk, Wk)	407	# weight layout: (Co, Ci/groups, Hk, Wk)
408	# bias layout: (Co)	408	# bias layout: (Co)
	@@ -500,7 +500,7 @@ def custom_grouped_matmul_add__strategy(y, x, weight, group_list, transpose_x=Tr
500		500
501	def is_tensor_evenly_shardable(shape, spec):	501	def is_tensor_evenly_shardable(shape, spec):
502	"""Check if the shape is evenly shardable according to the spec."""	502	"""Check if the shape is evenly shardable according to the spec."""
503	- # verify parameter validity	503	+ # verify parameter validity
504	if not isinstance(spec, DTensorSpec):	504	if not isinstance(spec, DTensorSpec):
505	raise TypeError(	505	raise TypeError(
506	f"Expected 'spec' to be DTensorSpec instance, got {type(spec).__name__} instead."	506	f"Expected 'spec' to be DTensorSpec instance, got {type(spec).__name__} instead."
	@@ -529,14 +529,14 @@ def custom_cross_entropy_loss_sharding(op_schema: OpSchema):
529	single_mesh_dim_strategies = []	529	single_mesh_dim_strategies = []
530		530
531	args_schema = op_schema.args_schema	531	args_schema = op_schema.args_schema
532	-	532	+
533	input_strategy = args_schema[0] if len(args_schema) > 0 else None	533	input_strategy = args_schema[0] if len(args_schema) > 0 else None
534	target_strategy = args_schema[1] if len(args_schema) > 1 else None	534	target_strategy = args_schema[1] if len(args_schema) > 1 else None
535	weight_strategy = args_schema[2] if len(args_schema) > 2 else None	535	weight_strategy = args_schema[2] if len(args_schema) > 2 else None
536	reduction = args_schema[3] if len(args_schema) > 3 else 'mean'	536	reduction = args_schema[3] if len(args_schema) > 3 else 'mean'
537		537
538	mesh = input_strategy.mesh	538	mesh = input_strategy.mesh
539	-	539	+
540	all_replicate: PlacementList = [	540	all_replicate: PlacementList = [
541	Replicate(), # loss	541	Replicate(), # loss
542	Replicate(), # log_prob	542	Replicate(), # log_prob

Mtorch_npu/dynamo/trace_rule.py+1-1

	@@ -49,7 +49,7 @@ torch_non_c_binding_in_graph_functions_npu = dict.fromkeys(
49	"torch.npu._get_current_allocator",	49	"torch.npu._get_current_allocator",
50	"torch.npu.is_bf16_supported",	50	"torch.npu.is_bf16_supported",
51	"torch.npu.memory._get_current_allocator",	51	"torch.npu.memory._get_current_allocator",
52	-	52	+
53	],	53	],
54	TorchInGraphFunctionVariable,	54	TorchInGraphFunctionVariable,
55	)	55	)

Mtorch_npu/jit/fusion_pass/fast_gelu.py+1-1

	@@ -14,5 +14,5 @@ def fast_gelu_pass(jit_mod):
14	%out = npu::fast_gelu(%x)	14	%out = npu::fast_gelu(%x)
15	return (%out)	15	return (%out)
16	"""	16	"""
17	-	17	+
18	torch._C._jit_pass_custom_pattern_based_rewrite_graph(pattern, replacement, jit_mod.graph)	18	torch._C._jit_pass_custom_pattern_based_rewrite_graph(pattern, replacement, jit_mod.graph)

Mtorch_npu/jit/register_fusion_pattern.py+1-2

	@@ -7,6 +7,5 @@ __all__ = ["optimize"]
7	def optimize(jit_mod):	7	def optimize(jit_mod):
8	if isinstance(jit_mod, torch.jit.ScriptModule):	8	if isinstance(jit_mod, torch.jit.ScriptModule):
9	torch.jit.optimize_for_inference(jit_mod)	9	torch.jit.optimize_for_inference(jit_mod)
10	-	10	+
11	fast_gelu_pass(jit_mod)	11	fast_gelu_pass(jit_mod)
12	-

Mtorch_npu/npu/__init__.py+4-4

	@@ -155,7 +155,7 @@ from torch_npu._init.common.warning_utils import _should_print_warning
155		155
156	import torch_npu	156	import torch_npu
157	from torch_npu.utils._error_code import ErrCode, pta_error, prof_error	157	from torch_npu.utils._error_code import ErrCode, pta_error, prof_error
158	-from .utils import (obfuscation_initialize, obfuscation_calculate, obfuscation_finalize,	158	+from .utils import (obfuscation_initialize, obfuscation_calculate, obfuscation_finalize,
159	synchronize, set_device, current_device, _get_device_index,	159	synchronize, set_device, current_device, _get_device_index,
160	device, device_of, StreamContext, stream, set_stream, current_stream, default_stream, set_sync_debug_mode,	160	device, device_of, StreamContext, stream, set_stream, current_stream, default_stream, set_sync_debug_mode,
161	get_sync_debug_mode, init_dump, current_blas_handle, is_bf16_supported,	161	get_sync_debug_mode, init_dump, current_blas_handle, is_bf16_supported,
	@@ -510,7 +510,7 @@ _cached_device_capability_env = None
510	def get_device_capability(device=None):	510	def get_device_capability(device=None):
511	r"""Query the minor and major data of device.	511	r"""Query the minor and major data of device.
512		512
513	- This function can be configured via the TORCH_NPU_DEVICE_CAPABILITY environment variable.	513	+ This function can be configured via the TORCH_NPU_DEVICE_CAPABILITY environment variable.
514	The format should be "major.minor", e.g., "9.0" or "8.0".	514	The format should be "major.minor", e.g., "9.0" or "8.0".
515		515
516	.. note::	516	.. note::
	@@ -520,7 +520,7 @@ def get_device_capability(device=None):
520	device (torch.device or int, optional): The device parameter has no practical meaning.	520	device (torch.device or int, optional): The device parameter has no practical meaning.
521		521
522	Returns:	522	Returns:
523	- tuple(int, int): the device capability of the device. Returns the tuple(major, minor) configured via	523	+ tuple(int, int): the device capability of the device. Returns the tuple(major, minor) configured via
524	TORCH_NPU_DEVICE_CAPABILITY, or None if TORCH_NPU_DEVICE_CAPABILITY not configured.	524	TORCH_NPU_DEVICE_CAPABILITY, or None if TORCH_NPU_DEVICE_CAPABILITY not configured.
525		525
526	Example:	526	Example:
	@@ -619,7 +619,7 @@ def _get_deterministic_level():
619	torch_npu.npu.set_deterministic_level(level)	619	torch_npu.npu.set_deterministic_level(level)
620	return level	620	return level
621	if level >= 1 and not torch.are_deterministic_algorithms_enabled():	621	if level >= 1 and not torch.are_deterministic_algorithms_enabled():
622	- level = 0	622	+ level = 0
623	torch_npu.npu.set_deterministic_level(level)	623	torch_npu.npu.set_deterministic_level(level)
624	return level	624	return level
625	return level	625	return level

Mtorch_npu/npu/_fft_plan_cache.py+1-1

	@@ -11,7 +11,7 @@ class NPUFFTPlanCache:
11	if name == "max_size":	11	if name == "max_size":
12	return torch_npu._C._npu_get_fft_plan_cache_max_size()	12	return torch_npu._C._npu_get_fft_plan_cache_max_size()
13	raise AttributeError("Unknown attribute " + name)	13	raise AttributeError("Unknown attribute " + name)
14	-	14	+
15	def __setattr__(self, name, value):	15	def __setattr__(self, name, value):
16	if name == "size":	16	if name == "size":
17	raise RuntimeError(".size is a read-only property showing the number of plans currently in the cache.")	17	raise RuntimeError(".size is a read-only property showing the number of plans currently in the cache.")

Mtorch_npu/npu/_format.py+2-2

	@@ -33,11 +33,11 @@ class Format(IntEnum):
33		33
34	def _apply_npu_format_patch():	34	def _apply_npu_format_patch():
35	orig_get_format = torch_npu.get_npu_format	35	orig_get_format = torch_npu.get_npu_format
36	-	36	+
37	def patched_get_format(tensor):	37	def patched_get_format(tensor):
38	"""get the Format type of tensor"""	38	"""get the Format type of tensor"""
39	format_int = orig_get_format(tensor)	39	format_int = orig_get_format(tensor)
40	return Format(format_int)	40	return Format(format_int)
41	-	41	+
42	torch_npu.get_npu_format = patched_get_format	42	torch_npu.get_npu_format = patched_get_format
43	torch_npu.Format = Format	43	torch_npu.Format = Format

Mtorch_npu/npu/aclnn/backends.py+1-1

	@@ -2,7 +2,7 @@ import warnings
2		2
3		3
4	def version():	4	def version():
5	- """Currently, the ACLNN version is not available and does not support it.	5	+ """Currently, the ACLNN version is not available and does not support it.
6	By default, it returns None.	6	By default, it returns None.
7	"""	7	"""
8	warnings.warn("torch.npu.aclnn.version isn't implemented!")	8	warnings.warn("torch.npu.aclnn.version isn't implemented!")

Mtorch_npu/npu/amp/sharded_grad_scaler.py+2-2

	@@ -280,7 +280,7 @@ class _ShardedGradScaler(GradScaler):
280	# Synchronize the detected inf across the ranks	280	# Synchronize the detected inf across the ranks
281	optimizer_state = self._per_optimizer_states[id(optimizer)]	281	optimizer_state = self._per_optimizer_states[id(optimizer)]
282	works = []	282	works = []
283	-	283	+
284	for found_inf in optimizer_state["found_inf_per_device"].values():	284	for found_inf in optimizer_state["found_inf_per_device"].values():
285	if found_inf.device.type == "cpu":	285	if found_inf.device.type == "cpu":
286	found_inf_npu = found_inf.to(self._scale.device)	286	found_inf_npu = found_inf.to(self._scale.device)
	@@ -288,7 +288,7 @@ class _ShardedGradScaler(GradScaler):
288	works.append((work, found_inf, found_inf_npu))	288	works.append((work, found_inf, found_inf_npu))
289	else:	289	else:
290	works.append((dist.all_reduce(found_inf, async_op=True, group=self.process_group), None, None))	290	works.append((dist.all_reduce(found_inf, async_op=True, group=self.process_group), None, None))
291	-	291	+
292	for item in works:	292	for item in works:
293	if item[1] is not None:	293	if item[1] is not None:
294	work, found_inf_cpu, found_inf_npu = item	294	work, found_inf_cpu, found_inf_npu = item

Mtorch_npu/npu/autocast_utils.py+1-1

	@@ -8,7 +8,7 @@ __all__ = ["get_amp_supported_dtype", "is_autocast_enabled", "set_autocast_enabl
8	def get_amp_supported_dtype():	8	def get_amp_supported_dtype():
9	if torch.npu.is_bf16_supported():	9	if torch.npu.is_bf16_supported():
10	return [torch.float16, torch.bfloat16, torch.float32]	10	return [torch.float16, torch.bfloat16, torch.float32]
11	- return [torch.float16, torch.float32]	11	+ return [torch.float16, torch.float32]
12		12
13		13
14	def is_autocast_enabled():	14	def is_autocast_enabled():

Mtorch_npu/npu/graphs.py+12-12

	@@ -229,18 +229,18 @@ def _print_npugraph_tensor_impl(input, tensor_name=None):
229	if device.type == "cpu":	229	if device.type == "cpu":
230	_print_callback_pending(tensor_name, input)	230	_print_callback_pending(tensor_name, input)
231	return	231	return
232	-	232	+
233	if device.type != "npu":	233	if device.type != "npu":
234	return	234	return
235		235
236	device_index = device.index	236	device_index = device.index
237	save_stream = _get_save_tensor_stream(device_index)	237	save_stream = _get_save_tensor_stream(device_index)
238	-	238	+
239	# Record event on the original compute stream before switching	239	# Record event on the original compute stream before switching
240	event1 = torch.npu.Event()	240	event1 = torch.npu.Event()
241	event2 = torch.npu.Event()	241	event2 = torch.npu.Event()
242	event1.record()	242	event1.record()
243	-	243	+
244	with torch.npu.stream(save_stream):	244	with torch.npu.stream(save_stream):
245	# Wait for the original stream to complete before D2H	245	# Wait for the original stream to complete before D2H
246	event1.wait()	246	event1.wait()
	@@ -255,7 +255,7 @@ def _print_npugraph_tensor_impl(input, tensor_name=None):
255	)	255	)
256	# Mark save_stream completion	256	# Mark save_stream completion
257	event2.record()	257	event2.record()
258	-	258	+
259	# Wait for save_stream to complete (back to original stream now)	259	# Wait for save_stream to complete (back to original stream now)
260	event2.wait()	260	event2.wait()
261		261
	@@ -268,19 +268,19 @@ def _save_npugraph_tensor_impl(input, save_path=None, overwrite=False):
268	if device.type == "cpu":	268	if device.type == "cpu":
269	torch.save(input, _build_save_npugraph_tensor_path(save_path, overwrite=overwrite))	269	torch.save(input, _build_save_npugraph_tensor_path(save_path, overwrite=overwrite))
270	return	270	return
271	-	271	+
272	if device.type != "npu":	272	if device.type != "npu":
273	return	273	return
274		274
275	device_index = device.index	275	device_index = device.index
276	save_stream = _get_save_tensor_stream(device_index)	276	save_stream = _get_save_tensor_stream(device_index)
277	final_path = _build_save_npugraph_tensor_path(save_path, device_index, overwrite)	277	final_path = _build_save_npugraph_tensor_path(save_path, device_index, overwrite)
278	-	278	+
279	# Record event on the original compute stream before switching	279	# Record event on the original compute stream before switching
280	event1 = torch.npu.Event()	280	event1 = torch.npu.Event()
281	event2 = torch.npu.Event()	281	event2 = torch.npu.Event()
282	event1.record()	282	event1.record()
283	-	283	+
284	with torch.npu.stream(save_stream):	284	with torch.npu.stream(save_stream):
285	# Wait for the original stream to complete before D2H	285	# Wait for the original stream to complete before D2H
286	event1.wait()	286	event1.wait()
	@@ -295,7 +295,7 @@ def _save_npugraph_tensor_impl(input, save_path=None, overwrite=False):
295	)	295	)
296	# Mark save_stream completion	296	# Mark save_stream completion
297	event2.record()	297	event2.record()
298	-	298	+
299	# Wait for save_stream to complete (back to original stream now)	299	# Wait for save_stream to complete (back to original stream now)
300	event2.wait()	300	event2.wait()
301		301
	@@ -305,19 +305,19 @@ def _save_npugraph_tensor_tensor_list_impl(input, save_path=None, overwrite=Fals
305	if device.type == "cpu":	305	if device.type == "cpu":
306	torch.save(list(input), _build_save_npugraph_tensor_path(save_path, overwrite=overwrite))	306	torch.save(list(input), _build_save_npugraph_tensor_path(save_path, overwrite=overwrite))
307	return	307	return
308	-	308	+
309	if device.type != "npu":	309	if device.type != "npu":
310	return	310	return
311		311
312	device_index = device.index	312	device_index = device.index
313	save_stream = _get_save_tensor_stream(device_index)	313	save_stream = _get_save_tensor_stream(device_index)
314	final_path = _build_save_npugraph_tensor_path(save_path, device_index, overwrite)	314	final_path = _build_save_npugraph_tensor_path(save_path, device_index, overwrite)
315	-	315	+
316	# Record event on the original compute stream before switching	316	# Record event on the original compute stream before switching
317	event1 = torch.npu.Event()	317	event1 = torch.npu.Event()
318	event2 = torch.npu.Event()	318	event2 = torch.npu.Event()
319	event1.record()	319	event1.record()
320	-	320	+
321	with torch.npu.stream(save_stream):	321	with torch.npu.stream(save_stream):
322	# Wait for the original stream to complete before D2H	322	# Wait for the original stream to complete before D2H
323	event1.wait()	323	event1.wait()
	@@ -332,7 +332,7 @@ def _save_npugraph_tensor_tensor_list_impl(input, save_path=None, overwrite=Fals
332	)	332	)
333	# Mark save_stream completion	333	# Mark save_stream completion
334	event2.record()	334	event2.record()
335	-	335	+
336	# Wait for save_stream to complete (back to original stream now)	336	# Wait for save_stream to complete (back to original stream now)
337	event2.wait()	337	event2.wait()
338		338

Mtorch_npu/npu/memory.py+1-1

	@@ -264,7 +264,7 @@ def reset_peak_host_memory_stats():
264	"""	264	"""
265	return torch_npu._C._npu_resetPeakHostMemoryStats()	265	return torch_npu._C._npu_resetPeakHostMemoryStats()
266		266
267	-	267	+
268	def empty_virt_addr_cache():	268	def empty_virt_addr_cache():
269	r"""Light-weight version of empty_cache(). It only unmaps virtual address,	269	r"""Light-weight version of empty_cache(). It only unmaps virtual address,
270	and store the free physical handles for later malloc.	270	and store the free physical handles for later malloc.

Mtorch_npu/npu/npugraph_ex/__init__.py+1-1

	@@ -30,5 +30,5 @@ def register_replacement(search_fn: SearchFn, replace_fn: ReplaceFn, example_inp
30	return npugraph_ex.patterns.pattern_pass_manager.register_replacement(search_fn, replace_fn, example_inputs,	30	return npugraph_ex.patterns.pattern_pass_manager.register_replacement(search_fn, replace_fn, example_inputs,
31	trace_fn=trace_fn, extra_check=extra_check,	31	trace_fn=trace_fn, extra_check=extra_check,
32	search_fn_pattern=search_fn_pattern,	32	search_fn_pattern=search_fn_pattern,
33	- scalar_workaround=scalar_workaround,	33	+ scalar_workaround=scalar_workaround,
34	skip_duplicates=skip_duplicates)	34	skip_duplicates=skip_duplicates)

Mtorch_npu/npu/utils.py+1-1

	@@ -22,7 +22,7 @@ __all__ = ["obfuscation_initialize", "obfuscation_finalize", "obfuscation_calcul
22		22
23		23
24	def obfuscation_initialize(hidden_size, tp_rank, cmd, *, data_type=None, model_obf_seed_id=0, data_obf_seed_id=0, thread_num=4, obf_coefficient=1.0):	24	def obfuscation_initialize(hidden_size, tp_rank, cmd, *, data_type=None, model_obf_seed_id=0, data_obf_seed_id=0, thread_num=4, obf_coefficient=1.0):
25	- return torch_npu.obfuscation_initialize(hidden_size, tp_rank, cmd, data_type=data_type, model_obf_seed_id=model_obf_seed_id,	25	+ return torch_npu.obfuscation_initialize(hidden_size, tp_rank, cmd, data_type=data_type, model_obf_seed_id=model_obf_seed_id,
26	data_obf_seed_id=data_obf_seed_id, thread_num=thread_num, obf_coefficient=obf_coefficient)	26	data_obf_seed_id=data_obf_seed_id, thread_num=thread_num, obf_coefficient=obf_coefficient)
27		27
28		28

Mtorch_npu/optim/npu_fused_adadelta.py+7-7

	@@ -71,13 +71,13 @@ class NpuFusedAdadelta(NpuFusedOptimizerBase):
71	if grad.is_sparse:	71	if grad.is_sparse:
72	raise RuntimeError('NpuFusedAdadelta does not support sparse gradients' +	72	raise RuntimeError('NpuFusedAdadelta does not support sparse gradients' +
73	pta_error(ErrCode.NOT_SUPPORT))	73	pta_error(ErrCode.NOT_SUPPORT))
74	-	74	+
75	self._init_param_state(p)	75	self._init_param_state(p)
76	state = self.state[p]	76	state = self.state[p]
77	step_list.append(state['step'])	77	step_list.append(state['step'])
78	square_avg_list.append(state['square_avg'])	78	square_avg_list.append(state['square_avg'])
79	acc_delta_list.append(state['acc_delta'])	79	acc_delta_list.append(state['acc_delta'])
80	-	80	+
81	combined_step = 0	81	combined_step = 0
82	combined_square_avg = None	82	combined_square_avg = None
83	combined_acc_delta = None	83	combined_acc_delta = None
	@@ -86,7 +86,7 @@ class NpuFusedAdadelta(NpuFusedOptimizerBase):
86	combined_step = step_list[0]	86	combined_step = step_list[0]
87	combined_square_avg = npu_combine_tensors(square_avg_list)	87	combined_square_avg = npu_combine_tensors(square_avg_list)
88	combined_acc_delta = npu_combine_tensors(acc_delta_list)	88	combined_acc_delta = npu_combine_tensors(acc_delta_list)
89	-	89	+
90	combined_state = defaultdict(dict)	90	combined_state = defaultdict(dict)
91	combined_state['step'] = combined_step	91	combined_state['step'] = combined_step
92	combined_state['square_avg'] = combined_square_avg	92	combined_state['square_avg'] = combined_square_avg
	@@ -97,12 +97,12 @@ class NpuFusedAdadelta(NpuFusedOptimizerBase):
97	def _maybe_init_combined_states(self):	97	def _maybe_init_combined_states(self):
98	if self.is_states_combined:	98	if self.is_states_combined:
99	return	99	return
100	-	100	+
101	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]	101	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]
102		102
103	for i, _ in enumerate(self.param_groups):	103	for i, _ in enumerate(self.param_groups):
104	self._combine_group_param_states(i)	104	self._combine_group_param_states(i)
105	-	105	+
106	if not all(value is None for value in self.combined_param_states_indexed_by_group):	106	if not all(value is None for value in self.combined_param_states_indexed_by_group):
107	self.is_states_combined = True	107	self.is_states_combined = True
108		108
	@@ -124,8 +124,8 @@ class NpuFusedAdadelta(NpuFusedOptimizerBase):
124	combined_group_grads = self.combined_grads_indexed_by_group[group_index]	124	combined_group_grads = self.combined_grads_indexed_by_group[group_index]
125	combined_group_param_states = self.combined_param_states_indexed_by_group[group_index]	125	combined_group_param_states = self.combined_param_states_indexed_by_group[group_index]
126		126
127	- for combined_param, combined_grad, combined_param_state in zip(combined_group_params,	127	+ for combined_param, combined_grad, combined_param_state in zip(combined_group_params,
128	- combined_group_grads,	128	+ combined_group_grads,
129	combined_group_param_states):	129	combined_group_param_states):
130	if combined_param is None or combined_grad is None:	130	if combined_param is None or combined_grad is None:
131	continue	131	continue

Mtorch_npu/optim/npu_fused_adam.py+6-6

	@@ -103,7 +103,7 @@ class NpuFusedAdam(NpuFusedOptimizerBase):
103	exp_avg_sq_list.append(state['exp_avg_sq'])	103	exp_avg_sq_list.append(state['exp_avg_sq'])
104	if amsgrad:	104	if amsgrad:
105	max_exp_avg_sq_list.append(state['max_exp_avg_sq'])	105	max_exp_avg_sq_list.append(state['max_exp_avg_sq'])
106	-	106	+
107	combined_step = 0	107	combined_step = 0
108	combined_exp_avg = None	108	combined_exp_avg = None
109	combined_exp_avg_sq = None	109	combined_exp_avg_sq = None
	@@ -114,7 +114,7 @@ class NpuFusedAdam(NpuFusedOptimizerBase):
114	combined_exp_avg = npu_combine_tensors(exp_avg_list)	114	combined_exp_avg = npu_combine_tensors(exp_avg_list)
115	combined_exp_avg_sq = npu_combine_tensors(exp_avg_sq_list)	115	combined_exp_avg_sq = npu_combine_tensors(exp_avg_sq_list)
116	combined_max_exp_avg_sq = npu_combine_tensors(max_exp_avg_sq_list)	116	combined_max_exp_avg_sq = npu_combine_tensors(max_exp_avg_sq_list)
117	-	117	+
118	combined_state = defaultdict(dict)	118	combined_state = defaultdict(dict)
119	combined_state['step'] = combined_step	119	combined_state['step'] = combined_step
120	combined_state['exp_avg'] = combined_exp_avg	120	combined_state['exp_avg'] = combined_exp_avg
	@@ -126,12 +126,12 @@ class NpuFusedAdam(NpuFusedOptimizerBase):
126	def _maybe_init_combined_states(self):	126	def _maybe_init_combined_states(self):
127	if self.is_states_combined:	127	if self.is_states_combined:
128	return	128	return
129	-	129	+
130	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]	130	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]
131		131
132	for i, _ in enumerate(self.param_groups):	132	for i, _ in enumerate(self.param_groups):
133	self._combine_group_param_states(i)	133	self._combine_group_param_states(i)
134	-	134	+
135	if not all(value is None for value in self.combined_param_states_indexed_by_group):	135	if not all(value is None for value in self.combined_param_states_indexed_by_group):
136	self.is_states_combined = True	136	self.is_states_combined = True
137		137
	@@ -154,8 +154,8 @@ class NpuFusedAdam(NpuFusedOptimizerBase):
154	combined_group_grads = self.combined_grads_indexed_by_group[group_index]	154	combined_group_grads = self.combined_grads_indexed_by_group[group_index]
155	combined_group_param_states = self.combined_param_states_indexed_by_group[group_index]	155	combined_group_param_states = self.combined_param_states_indexed_by_group[group_index]
156		156
157	- for combined_param, combined_grad, combined_param_state in zip(combined_group_params,	157	+ for combined_param, combined_grad, combined_param_state in zip(combined_group_params,
158	- combined_group_grads,	158	+ combined_group_grads,
159	combined_group_param_states):	159	combined_group_param_states):
160	if combined_param is None or combined_grad is None:	160	if combined_param is None or combined_grad is None:
161	continue	161	continue

Mtorch_npu/optim/npu_fused_adamp.py+3-3

Mtorch_npu/optim/npu_fused_bert_adam.py+2-2

	@@ -149,12 +149,12 @@ class NpuFusedBertAdam(NpuFusedOptimizerBase):
149	def _maybe_init_combined_states(self):	149	def _maybe_init_combined_states(self):
150	if self.is_states_combined:	150	if self.is_states_combined:
151	return	151	return
152	-	152	+
153	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]	153	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]
154		154
155	for i, _ in enumerate(self.param_groups):	155	for i, _ in enumerate(self.param_groups):
156	self._combine_group_param_states(i)	156	self._combine_group_param_states(i)
157	-	157	+
158	if not all(value is None for value in self.combined_param_states_indexed_by_group):	158	if not all(value is None for value in self.combined_param_states_indexed_by_group):
159	self.is_states_combined = True	159	self.is_states_combined = True
160		160

Mtorch_npu/optim/npu_fused_lamb.py+8-8

	@@ -137,13 +137,13 @@ class NpuFusedLamb(NpuFusedOptimizerBase):
137	if grad.is_sparse:	137	if grad.is_sparse:
138	raise RuntimeError('NpuFusedLamb does not support sparse gradients, '	138	raise RuntimeError('NpuFusedLamb does not support sparse gradients, '
139	'please consider SparseAdam instead.' + pta_error(ErrCode.NOT_SUPPORT))	139	'please consider SparseAdam instead.' + pta_error(ErrCode.NOT_SUPPORT))
140	-	140	+
141	self._init_param_state(p)	141	self._init_param_state(p)
142	state = self.state[p]	142	state = self.state[p]
143	step_list.append(state['step'])	143	step_list.append(state['step'])
144	exp_avg_list.append(state['exp_avg'])	144	exp_avg_list.append(state['exp_avg'])
145	exp_avg_sq_list.append(state['exp_avg_sq'])	145	exp_avg_sq_list.append(state['exp_avg_sq'])
146	-	146	+
147	combined_step = 0	147	combined_step = 0
148	combined_exp_avg = None	148	combined_exp_avg = None
149	combined_exp_avg_sq = None	149	combined_exp_avg_sq = None
	@@ -152,7 +152,7 @@ class NpuFusedLamb(NpuFusedOptimizerBase):
152	combined_step = step_list[0]	152	combined_step = step_list[0]
153	combined_exp_avg = npu_combine_tensors(exp_avg_list)	153	combined_exp_avg = npu_combine_tensors(exp_avg_list)
154	combined_exp_avg_sq = npu_combine_tensors(exp_avg_sq_list)	154	combined_exp_avg_sq = npu_combine_tensors(exp_avg_sq_list)
155	-	155	+
156	combined_state = defaultdict(dict)	156	combined_state = defaultdict(dict)
157	combined_state['step'] = combined_step	157	combined_state['step'] = combined_step
158	combined_state['exp_avg'] = combined_exp_avg	158	combined_state['exp_avg'] = combined_exp_avg
	@@ -163,12 +163,12 @@ class NpuFusedLamb(NpuFusedOptimizerBase):
163	def _maybe_init_combined_states(self):	163	def _maybe_init_combined_states(self):
164	if self.is_states_combined:	164	if self.is_states_combined:
165	return	165	return
166	-	166	+
167	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]	167	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]
168		168
169	for i, _ in enumerate(self.param_groups):	169	for i, _ in enumerate(self.param_groups):
170	self._combine_group_param_states(i)	170	self._combine_group_param_states(i)
171	-	171	+
172	if not all(value is None for value in self.combined_param_states_indexed_by_group):	172	if not all(value is None for value in self.combined_param_states_indexed_by_group):
173	self.is_states_combined = True	173	self.is_states_combined = True
174		174
	@@ -237,8 +237,8 @@ class NpuFusedLamb(NpuFusedOptimizerBase):
237	combined_param_pow.copy_(combined_param.pow(2))	237	combined_param_pow.copy_(combined_param.pow(2))
238	combined_adam_step_pow.copy_(adam_step.pow(2))	238	combined_adam_step_pow.copy_(adam_step.pow(2))
239		239
240	- for param_pow, adam_step_pow, trust_ratio in zip(param_pow_list,	240	+ for param_pow, adam_step_pow, trust_ratio in zip(param_pow_list,
241	- adam_step_pow_list,	241	+ adam_step_pow_list,
242	trust_ratio_list):	242	trust_ratio_list):
243	weight_norm = param_pow.sum().sqrt().clamp(0, 10)	243	weight_norm = param_pow.sum().sqrt().clamp(0, 10)
244	adam_norm = adam_step_pow.sum().sqrt()	244	adam_norm = adam_step_pow.sum().sqrt()
	@@ -253,7 +253,7 @@ class NpuFusedLamb(NpuFusedOptimizerBase):
253	def step(self, closure=None):	253	def step(self, closure=None):
254	if not self.is_params_grads_combined:	254	if not self.is_params_grads_combined:
255	self._maybe_init_combined_params_and_grads()	255	self._maybe_init_combined_params_and_grads()
256	-	256	+
257	if not self.is_states_combined:	257	if not self.is_states_combined:
258	self._maybe_init_combined_states()	258	self._maybe_init_combined_states()
259		259

Mtorch_npu/optim/npu_fused_optim_base.py+5-5

	@@ -29,7 +29,7 @@ class NpuFusedOptimizerBase(Optimizer):
29	return	29	return
30		30
31	self.combined_params_indexed_by_group = len(self.param_groups) * [[]]	31	self.combined_params_indexed_by_group = len(self.param_groups) * [[]]
32	- self.combined_grads_indexed_by_group = len(self.param_groups) * [[]]	32	+ self.combined_grads_indexed_by_group = len(self.param_groups) * [[]]
33		33
34	params_list_each_group = []	34	params_list_each_group = []
35	params_size_each_group = []	35	params_size_each_group = []
	@@ -103,9 +103,9 @@ class NpuFusedOptimizerBase(Optimizer):
103		103
104	self.combined_params_indexed_by_group[group_index] = group_combined_params	104	self.combined_params_indexed_by_group[group_index] = group_combined_params
105	self.combined_grads_indexed_by_group[group_index] = group_combined_grads	105	self.combined_grads_indexed_by_group[group_index] = group_combined_grads
106	-	106	+
107	if not all(value is None for value in self.params_all_group_combined):	107	if not all(value is None for value in self.params_all_group_combined):
108	- self.is_params_grads_combined = True	108	+ self.is_params_grads_combined = True
109		109
110	@torch.no_grad()	110	@torch.no_grad()
111	def step(self, closure=None):	111	def step(self, closure=None):
	@@ -115,7 +115,7 @@ class NpuFusedOptimizerBase(Optimizer):
115		115
116	if not self.is_params_grads_combined:	116	if not self.is_params_grads_combined:
117	self._maybe_init_combined_params_and_grads()	117	self._maybe_init_combined_params_and_grads()
118	-	118	+
119	if not self.is_states_combined:	119	if not self.is_states_combined:
120	self._maybe_init_combined_states()	120	self._maybe_init_combined_states()
121		121
	@@ -135,7 +135,7 @@ class NpuFusedOptimizerBase(Optimizer):
135	if not self.is_params_grads_combined:	135	if not self.is_params_grads_combined:
136	super().zero_grad(set_to_none)	136	super().zero_grad(set_to_none)
137	return	137	return
138	-	138	+
139	for grads_combined_one_dtype in self.grads_all_group_combined:	139	for grads_combined_one_dtype in self.grads_all_group_combined:
140	if grads_combined_one_dtype is None:	140	if grads_combined_one_dtype is None:
141	continue	141	continue

Mtorch_npu/optim/npu_fused_rmsprop.py+2-2

	@@ -129,12 +129,12 @@ class NpuFusedRMSprop(NpuFusedOptimizerBase):
129	def _maybe_init_combined_states(self):	129	def _maybe_init_combined_states(self):
130	if self.is_states_combined:	130	if self.is_states_combined:
131	return	131	return
132	-	132	+
133	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]	133	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]
134		134
135	for i, _ in enumerate(self.param_groups):	135	for i, _ in enumerate(self.param_groups):
136	self._combine_group_param_states(i)	136	self._combine_group_param_states(i)
137	-	137	+
138	if not all(value is None for value in self.combined_param_states_indexed_by_group):	138	if not all(value is None for value in self.combined_param_states_indexed_by_group):
139	self.is_states_combined = True	139	self.is_states_combined = True
140		140

Mtorch_npu/optim/npu_fused_rmsprop_tf.py+2-2

	@@ -129,12 +129,12 @@ class NpuFusedRMSpropTF(NpuFusedOptimizerBase):
129	def _maybe_init_combined_states(self):	129	def _maybe_init_combined_states(self):
130	if self.is_states_combined:	130	if self.is_states_combined:
131	return	131	return
132	-	132	+
133	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]	133	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]
134		134
135	for i, _ in enumerate(self.param_groups):	135	for i, _ in enumerate(self.param_groups):
136	self._combine_group_param_states(i)	136	self._combine_group_param_states(i)
137	-	137	+
138	if not all(value is None for value in self.combined_param_states_indexed_by_group):	138	if not all(value is None for value in self.combined_param_states_indexed_by_group):
139	self.is_states_combined = True	139	self.is_states_combined = True
140		140

Mtorch_npu/optim/npu_fused_sgd.py+3-3

	@@ -117,12 +117,12 @@ class NpuFusedSGD(NpuFusedOptimizerBase):
117	def _maybe_init_combined_states(self):	117	def _maybe_init_combined_states(self):
118	if self.is_states_combined:	118	if self.is_states_combined:
119	return	119	return
120	-	120	+
121	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]	121	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]
122		122
123	for i, _ in enumerate(self.param_groups):	123	for i, _ in enumerate(self.param_groups):
124	self._combine_group_param_states(i)	124	self._combine_group_param_states(i)
125	-	125	+
126	if not all(value is None for value in self.combined_param_states_indexed_by_group):	126	if not all(value is None for value in self.combined_param_states_indexed_by_group):
127	self.is_states_combined = True	127	self.is_states_combined = True
128		128
	@@ -163,7 +163,7 @@ class NpuFusedSGD(NpuFusedOptimizerBase):
163		163
164	combined_param_one_dtype.add_(combined_grad_one_dtype,	164	combined_param_one_dtype.add_(combined_grad_one_dtype,
165	alpha=-group['lr'])	165	alpha=-group['lr'])
166	-	166	+
167	def step(self, closure=None):	167	def step(self, closure=None):
168	ret = super().step(closure)	168	ret = super().step(closure)
169	self._momentum_buffer_already_in_state = True	169	self._momentum_buffer_already_in_state = True

Mtorch_npu/profiler/_profiler_action_controller.py+1-1

	@@ -13,7 +13,7 @@ class ProfActionController:
13	self,	13	self,
14	prof,	14	prof,
15	prof_inst: _ProfInterface,	15	prof_inst: _ProfInterface,
16	- on_trace_ready: Optional[Callable[..., Any]] = None,	16	+ on_trace_ready: Optional[Callable[..., Any]] = None,
17	) -> None:	17	) -> None:
18	self.prof = prof	18	self.prof = prof
19	self.prof_inst = prof_inst	19	self.prof_inst = prof_inst

Mtorch_npu/profiler/_profiler_path_creator.py+1-1

	@@ -41,7 +41,7 @@ class ProfPathCreator:
41	PathManager.check_input_directory_path(dir_path)	41	PathManager.check_input_directory_path(dir_path)
42	self._dir_path = dir_path	42	self._dir_path = dir_path
43	elif dir_name is None:	43	elif dir_name is None:
44	- self._dir_path = dir_name	44	+ self._dir_path = dir_name
45	else:	45	else:
46	print_warn_msg("Invalid parameter dir_name, reset it to default.")	46	print_warn_msg("Invalid parameter dir_name, reset it to default.")
47	self._dir_path = None	47	self._dir_path = None

Mtorch_npu/profiler/analysis/prof_bean/_ge_memory_record_bean.py+1-1

	@@ -5,7 +5,7 @@ from ..prof_common_func._constant import convert_us2ns
5		5
6	__all__ = []	6	__all__ = []
7		7
8	-	8	+
9	class GeMemoryRecordBean(CommonBean):	9	class GeMemoryRecordBean(CommonBean):
10		10
11	def __init__(self, data: dict):	11	def __init__(self, data: dict):

Mtorch_npu/profiler/analysis/prof_bean/_ge_op_memory_bean.py+1-1

	@@ -15,7 +15,7 @@ class GeOpMemoryBean(CommonBean):
15		15
16	@property	16	@property
17	def row(self) -> list:	17	def row(self) -> list:
18	- return [self.name, self.size, self.allocation_time, self.release_time, None,	18	+ return [self.name, self.size, self.allocation_time, self.release_time, None,
19	self.dur, None, self.allocation_total_allocated, self.allocation_total_reserved, None,	19	self.dur, None, self.allocation_total_allocated, self.allocation_total_reserved, None,
20	self.release_total_allocated, self.release_total_reserved, None, None, self.device]	20	self.release_total_allocated, self.release_total_reserved, None, None, self.device]
21		21

Mtorch_npu/profiler/analysis/prof_bean/_memory_use_bean.py+1-1

	@@ -117,7 +117,7 @@ class MemoryUseBean(CommonBean):
117	@property	117	@property
118	def data_type(self) -> int:	118	def data_type(self) -> int:
119	return self._data_type	119	return self._data_type
120	-	120	+
121	@property	121	@property
122	def allocator_type(self) -> int:	122	def allocator_type(self) -> int:
123	return self._allocator_type	123	return self._allocator_type

Mtorch_npu/profiler/analysis/prof_bean/_param_tensor_bean.py+4-4

	@@ -27,19 +27,19 @@ class ParamTensorBean:
27	self._key = self._constant_data[0]	27	self._key = self._constant_data[0]
28	self._module_params = None	28	self._module_params = None
29	self._optimizer_params = None	29	self._optimizer_params = None
30	-	30	+
31	module_params = self._origin_data.get(self.TLV_TYPE_DICT.get(Constant.MODULE_PARAM))	31	module_params = self._origin_data.get(self.TLV_TYPE_DICT.get(Constant.MODULE_PARAM))
32	if module_params is not None:	32	if module_params is not None:
33	self._module_params = [param for param in module_params.split('}')]	33	self._module_params = [param for param in module_params.split('}')]
34	-	34	+
35	optimizer_params = self._origin_data.get(self.TLV_TYPE_DICT.get(Constant.OPTIMIZER_PARAM))	35	optimizer_params = self._origin_data.get(self.TLV_TYPE_DICT.get(Constant.OPTIMIZER_PARAM))
36	if optimizer_params is not None:	36	if optimizer_params is not None:
37	self._optimizer_params = [param for param in optimizer_params.split('}')]	37	self._optimizer_params = [param for param in optimizer_params.split('}')]
38	-	38	+
39	@property	39	@property
40	def key(self) -> int:	40	def key(self) -> int:
41	return self._key	41	return self._key
42	-	42	+
43	@property	43	@property
44	def params(self) -> KeyAndParam:	44	def params(self) -> KeyAndParam:
45	return KeyAndParam(self._key, self._module_params, self._optimizer_params)	45	return KeyAndParam(self._key, self._module_params, self._optimizer_params)

Mtorch_npu/profiler/analysis/prof_bean/_torch_op_bean.py+2-2

	@@ -95,7 +95,7 @@ class TorchOpBean:
95	if self._call_stack is None:	95	if self._call_stack is None:
96	self._call_stack = self._origin_data.get(self.TLV_TYPE_DICT.get(Constant.CALL_STACK), "").replace(";", ";\r\n")	96	self._call_stack = self._origin_data.get(self.TLV_TYPE_DICT.get(Constant.CALL_STACK), "").replace(";", ";\r\n")
97	return self._call_stack	97	return self._call_stack
98	-	98	+
99	@property	99	@property
100	def inputs(self):	100	def inputs(self):
101	if self._inputs is None:	101	if self._inputs is None:
	@@ -121,7 +121,7 @@ class TorchOpBean:
121	@property	121	@property
122	def is_torch_op(self):	122	def is_torch_op(self):
123	return True	123	return True
124	-	124	+
125	def _init_timestamps(self):	125	def _init_timestamps(self):
126	profiler_config = ProfilerConfig()	126	profiler_config = ProfilerConfig()
127	start_syscnt = self._constant_data[TorchOpEnum.START_NS.value]	127	start_syscnt = self._constant_data[TorchOpEnum.START_NS.value]

Mtorch_npu/profiler/analysis/prof_common_func/_cann_package_manager.py+2-2

	@@ -11,10 +11,10 @@ def check_msprof_help_output(search_text: str) -> bool:
11	msprof_path = shutil.which("msprof")	11	msprof_path = shutil.which("msprof")
12	if not msprof_path:	12	if not msprof_path:
13	return False	13	return False
14	-	14	+
15	if not ProfilerPathManager.check_path_permission(msprof_path):	15	if not ProfilerPathManager.check_path_permission(msprof_path):
16	return False	16	return False
17	-	17	+
18	completed_process = subprocess.run([msprof_path, "--help"], capture_output=True, shell=False, text=True)	18	completed_process = subprocess.run([msprof_path, "--help"], capture_output=True, shell=False, text=True)
19	if completed_process.returncode != COMMAND_SUCCESS:	19	if completed_process.returncode != COMMAND_SUCCESS:
20	return False	20	return False

Mtorch_npu/profiler/analysis/prof_common_func/_db_manager.py+4-4

	@@ -33,7 +33,7 @@ class DbManager:
33	def create_connect_db(cls, db_path: str) -> tuple:	33	def create_connect_db(cls, db_path: str) -> tuple:
34	"""	34	"""
35	create and connect database	35	create and connect database
36	- """	36	+ """
37	if os.path.exists(db_path):	37	if os.path.exists(db_path):
38	FileManager.check_db_file_vaild(db_path)	38	FileManager.check_db_file_vaild(db_path)
39	try:	39	try:
	@@ -41,7 +41,7 @@ class DbManager:
41	conn = sqlite3.connect(db_path, timeout=2147483, check_same_thread=False)	41	conn = sqlite3.connect(db_path, timeout=2147483, check_same_thread=False)
42	except sqlite3.Error as err:	42	except sqlite3.Error as err:
43	return EmptyClass("emoty conn"), EmptyClass("empty curs")	43	return EmptyClass("emoty conn"), EmptyClass("empty curs")
44	-	44	+
45	try:	45	try:
46	curs = conn.cursor()	46	curs = conn.cursor()
47	os.chmod(db_path, Constant.FILE_AUTHORITY)	47	os.chmod(db_path, Constant.FILE_AUTHORITY)
	@@ -60,7 +60,7 @@ class DbManager:
60	cur.close()	60	cur.close()
61	except sqlite3.Error as err:	61	except sqlite3.Error as err:
62	raise RuntimeError(f"Falied to close db connection cursor") from err	62	raise RuntimeError(f"Falied to close db connection cursor") from err
63	-	63	+
64	try:	64	try:
65	conn.close()	65	conn.close()
66	except sqlite3.Error as err:	66	except sqlite3.Error as err:
	@@ -162,7 +162,7 @@ class DbManager:
162	except sqlite3.Error as err:	162	except sqlite3.Error as err:
163	print_error_msg("SQLite Error: %s" % " ".join(err.args))	163	print_error_msg("SQLite Error: %s" % " ".join(err.args))
164	return []	164	return []
165	-	165	+
166	@classmethod	166	@classmethod
167	def fetch_one_data(cls, cur: sqlite3.Cursor, sql: str) -> list:	167	def fetch_one_data(cls, cur: sqlite3.Cursor, sql: str) -> list:
168	"""	168	"""

Mtorch_npu/profiler/analysis/prof_common_func/_id_manager.py+2-2

	@@ -9,10 +9,10 @@ class Str2IdManager:
9	def __init__(self) -> None:	9	def __init__(self) -> None:
10	self._str_id_map = {}	10	self._str_id_map = {}
11	self._curr_id = 0	11	self._curr_id = 0
12	-	12	+
13	def set_start_id(self, start_id: int):	13	def set_start_id(self, start_id: int):
14	self._curr_id = start_id	14	self._curr_id = start_id
15	-	15	+
16	def get_id_from_str(self, string: str) -> int:	16	def get_id_from_str(self, string: str) -> int:
17	# 先查询；有性能影响的话直接+1，不查询了	17	# 先查询；有性能影响的话直接+1，不查询了
18	if not string:	18	if not string:

Mtorch_npu/profiler/analysis/prof_common_func/_singleton.py+0-1

	@@ -7,4 +7,3 @@ class Singleton(object):
7	if self._cls not in self._instance:	7	if self._cls not in self._instance:
8	self._instance[self._cls] = self._cls()	8	self._instance[self._cls] = self._cls()
9	return self._instance[self._cls]	9	return self._instance[self._cls]
10	-

Mtorch_npu/profiler/analysis/prof_common_func/_time_range_calculator.py+1-1

	@@ -27,7 +27,7 @@ class TimeRange:
27		27
28		28
29	class CommunicationTimeRange(TimeRange):	29	class CommunicationTimeRange(TimeRange):
30	-	30	+
31	def __init__(self):	31	def __init__(self):
32	super().__init__()	32	super().__init__()
33		33

Mtorch_npu/profiler/analysis/prof_config/_parser_config.py+1-1

	@@ -27,7 +27,7 @@ from ..prof_view._memory_view_parser import MemoryViewParser
27	from ..prof_view._integrate_parser import IntegrateParser	27	from ..prof_view._integrate_parser import IntegrateParser
28	from ..prof_view._communication_parser import CommunicationParser	28	from ..prof_view._communication_parser import CommunicationParser
29	from ..prof_view._memory_timeline_parser import MemoryTimelineParser	29	from ..prof_view._memory_timeline_parser import MemoryTimelineParser
30	-from ..prof_view.prof_db_parse._db_parser import DbParser	30	+from ..prof_view.prof_db_parse._db_parser import DbParser
31	from ..prof_view.prepare_parse._fwk_pre_parser import (	31	from ..prof_view.prepare_parse._fwk_pre_parser import (
32	TracePreParser,	32	TracePreParser,
33	TreeBuildParser,	33	TreeBuildParser,

Mtorch_npu/profiler/analysis/prof_parse/_event_tree_parser.py+39-39

	@@ -111,11 +111,11 @@ class _TensorMetadata:
111	self.device_type = device_type	111	self.device_type = device_type
112	self.device_index = device_index	112	self.device_index = device_index
113	self.id = _IDReference(None, None)	113	self.id = _IDReference(None, None)
114	-	114	+
115	@property	115	@property
116	def tensor_id(self) -> Optional[int]:	116	def tensor_id(self) -> Optional[int]:
117	return self.id.tensor_id	117	return self.id.tensor_id
118	-	118	+
119	@property	119	@property
120	def allocation_id(self) -> Optional[int]:	120	def allocation_id(self) -> Optional[int]:
121	return self.id.allocation_id	121	return self.id.allocation_id
	@@ -125,7 +125,7 @@ def parse_tensor_metadata(tensor_str: str) -> Optional[_TensorMetadata]:
125	parts = tensor_str.split(';')	125	parts = tensor_str.split(';')
126	if len(parts) != TensorEnum.NUM_FIELDS.value:	126	if len(parts) != TensorEnum.NUM_FIELDS.value:
127	return None	127	return None
128	-	128	+
129	impl = int(parts[TensorEnum.TENSOR_IMPL.value], BASE_16)	129	impl = int(parts[TensorEnum.TENSOR_IMPL.value], BASE_16)
130	ptr = int(parts[TensorEnum.STORAGE_PTR.value], BASE_16) if parts[TensorEnum.STORAGE_PTR.value] else None	130	ptr = int(parts[TensorEnum.STORAGE_PTR.value], BASE_16) if parts[TensorEnum.STORAGE_PTR.value] else None
131	dtype = parts[TensorEnum.DTYPE.value]	131	dtype = parts[TensorEnum.DTYPE.value]
	@@ -193,7 +193,7 @@ class _ExtraFields_TorchOp:
193	self.scope = bean.scope	193	self.scope = bean.scope
194	self.forward_tid = bean.args.get(Constant.FORWARD_THREAD_ID, -1)	194	self.forward_tid = bean.args.get(Constant.FORWARD_THREAD_ID, -1)
195	self.sequence_num = bean.args.get(Constant.SEQUENCE_NUMBER, -1)	195	self.sequence_num = bean.args.get(Constant.SEQUENCE_NUMBER, -1)
196	-	196	+
197	types_string = bean.args.get(Constant.INPUT_DTYPES, None)	197	types_string = bean.args.get(Constant.INPUT_DTYPES, None)
198	tensors_string = bean.inputs.get(Constant.INPUT_TENSORS, None)	198	tensors_string = bean.inputs.get(Constant.INPUT_TENSORS, None)
199	tensorlists_string = bean.inputs.get(Constant.INPUT_TENSORLISTS, None)	199	tensorlists_string = bean.inputs.get(Constant.INPUT_TENSORLISTS, None)
	@@ -211,11 +211,11 @@ class _ExtraFields_Allocation:
211	self.device_type = bean.device_type	211	self.device_type = bean.device_type
212	self.device_index = bean.device_index	212	self.device_index = bean.device_index
213	self.id = _IDReference(None, None)	213	self.id = _IDReference(None, None)
214	-	214	+
215	@property	215	@property
216	def tensor_id(self) -> Optional[int]:	216	def tensor_id(self) -> Optional[int]:
217	return self.id.tensor_id	217	return self.id.tensor_id
218	-	218	+
219	@property	219	@property
220	def allocation_id(self) -> Optional[int]:	220	def allocation_id(self) -> Optional[int]:
221	return self.id.allocation_id	221	return self.id.allocation_id
	@@ -242,7 +242,7 @@ def parse_module_param(param_str: str) -> Optional[_ModuleParam]:
242	param_list = param_str.strip().split(')')	242	param_list = param_str.strip().split(')')
243	if len(param_list) != ModuleParamEnum.NUM_FIELDS.value:	243	if len(param_list) != ModuleParamEnum.NUM_FIELDS.value:
244	return None	244	return None
245	-	245	+
246	name = param_list[ModuleParamEnum.NAME.value]	246	name = param_list[ModuleParamEnum.NAME.value]
247	tensor = parse_tensor_metadata(param_list[ModuleParamEnum.METADATA.value])	247	tensor = parse_tensor_metadata(param_list[ModuleParamEnum.METADATA.value])
248	grad = (parse_tensor_metadata(param_list[ModuleParamEnum.GRAD.value])	248	grad = (parse_tensor_metadata(param_list[ModuleParamEnum.GRAD.value])
	@@ -261,7 +261,7 @@ def parse_state_param(state_str: str) -> Optional[List[Tuple[str, _TensorMetadat
261	return None	261	return None
262	state_pairs.append(tuple([state_pair[StateParamEnum.NAME.value],	262	state_pairs.append(tuple([state_pair[StateParamEnum.NAME.value],
263	parse_tensor_metadata(state_pair[StateParamEnum.METADATA.value])]))	263	parse_tensor_metadata(state_pair[StateParamEnum.METADATA.value])]))
264	-	264	+
265	return state_pairs	265	return state_pairs
266		266
267		267
	@@ -269,7 +269,7 @@ def parse_optimizer_param(param_str: str) -> Optional[_OptimizerParam]:
269	param_list = param_str.strip().split(')')	269	param_list = param_str.strip().split(')')
270	if len(param_list) != OptimizerParamEnum.NUM_FIELDS.value:	270	if len(param_list) != OptimizerParamEnum.NUM_FIELDS.value:
271	return None	271	return None
272	-	272	+
273	tensor = parse_tensor_metadata(param_list[OptimizerParamEnum.METADATA.value])	273	tensor = parse_tensor_metadata(param_list[OptimizerParamEnum.METADATA.value])
274	grad = (parse_tensor_metadata(param_list[OptimizerParamEnum.GRAD.value])	274	grad = (parse_tensor_metadata(param_list[OptimizerParamEnum.GRAD.value])
275	if param_list[OptimizerParamEnum.GRAD.value] else None)	275	if param_list[OptimizerParamEnum.GRAD.value] else None)
	@@ -316,7 +316,7 @@ class _ProfilerEvent:
316	self.tid = bean.tid	316	self.tid = bean.tid
317	self.start_time_ns = bean.ts	317	self.start_time_ns = bean.ts
318	self.extra_fields = _ExtraFields_PyCall(bean)	318	self.extra_fields = _ExtraFields_PyCall(bean)
319	-	319	+
320	@property	320	@property
321	def name(self) -> str:	321	def name(self) -> str:
322	if self.tag == _EventType.TorchOp:	322	if self.tag == _EventType.TorchOp:
	@@ -326,7 +326,7 @@ class _ProfilerEvent:
326	elif self.tag == _EventType.PyCall:	326	elif self.tag == _EventType.PyCall:
327	return self.extra_fields.name	327	return self.extra_fields.name
328	return ""	328	return ""
329	-	329	+
330	@property	330	@property
331	def end_time_ns(self) -> int:	331	def end_time_ns(self) -> int:
332	if self.tag == _EventType.TorchOp:	332	if self.tag == _EventType.TorchOp:
	@@ -336,7 +336,7 @@ class _ProfilerEvent:
336	elif self.tag == _EventType.PyCall:	336	elif self.tag == _EventType.PyCall:
337	return self.extra_fields.end_time_ns	337	return self.extra_fields.end_time_ns
338	return -1	338	return -1
339	-	339	+
340	def __lt__(self, other: '_ProfilerEvent') -> bool:	340	def __lt__(self, other: '_ProfilerEvent') -> bool:
341	return self.end_time_ns < other.end_time_ns	341	return self.end_time_ns < other.end_time_ns
342		342
	@@ -362,17 +362,17 @@ def push_event(event: _ProfilerEvent,
362	if event.finished:	362	if event.finished:
363	print_error_msg("Error when building tree: the event finished.")	363	print_error_msg("Error when building tree: the event finished.")
364	return False	364	return False
365	-	365	+
366	parent = thread_event.get(event.tid)	366	parent = thread_event.get(event.tid)
367	if parent is None:	367	if parent is None:
368	fwd_tid = event.extra_fields.forward_tid if event.tag == _EventType.TorchOp else 0	368	fwd_tid = event.extra_fields.forward_tid if event.tag == _EventType.TorchOp else 0
369	if fwd_tid:	369	if fwd_tid:
370	parent = thread_event.get(fwd_tid)	370	parent = thread_event.get(fwd_tid)
371	-	371	+
372	if parent is not None:	372	if parent is not None:
373	event.parent = parent	373	event.parent = parent
374	parent.children.append(event)	374	parent.children.append(event)
375	-	375	+
376	if event.end_time_ns > event.start_time_ns:	376	if event.end_time_ns > event.start_time_ns:
377	thread_event[event.tid] = event	377	thread_event[event.tid] = event
378	unfinished_events.put((event.end_time_ns, event))	378	unfinished_events.put((event.end_time_ns, event))
	@@ -380,20 +380,20 @@ def push_event(event: _ProfilerEvent,
380	else:	380	else:
381	if not mark_finished(event):	381	if not mark_finished(event):
382	return False	382	return False
383	-	383	+
384	return True	384	return True
385		385
386		386
387	def pop_event(event: _ProfilerEvent, thread_event: Dict[int, _ProfilerEvent]) -> bool:	387	def pop_event(event: _ProfilerEvent, thread_event: Dict[int, _ProfilerEvent]) -> bool:
388	if event.finished:	388	if event.finished:
389	return True	389	return True
390	-	390	+
391	tid = event.tid	391	tid = event.tid
392	cur_event = thread_event.get(tid)	392	cur_event = thread_event.get(tid)
393	if cur_event is None:	393	if cur_event is None:
394	print_error_msg("Error when building tree: current event is none.")	394	print_error_msg("Error when building tree: current event is none.")
395	return False	395	return False
396	-	396	+
397	while cur_event != event:	397	while cur_event != event:
398	if not mark_finished(cur_event):	398	if not mark_finished(cur_event):
399	return False	399	return False
	@@ -401,13 +401,13 @@ def pop_event(event: _ProfilerEvent, thread_event: Dict[int, _ProfilerEvent]) ->
401	print_error_msg("Error when building tree: current event's parent is None.")	401	print_error_msg("Error when building tree: current event's parent is None.")
402	return False	402	return False
403	cur_event = cur_event.parent	403	cur_event = cur_event.parent
404	-	404	+
405	if not mark_finished(event):	405	if not mark_finished(event):
406	return False	406	return False
407	thread_event.pop(tid, None)	407	thread_event.pop(tid, None)
408	if event.parent:	408	if event.parent:
409	thread_event[tid] = event.parent	409	thread_event[tid] = event.parent
410	-	410	+
411	return True	411	return True
412		412
413		413
	@@ -423,7 +423,7 @@ def build_event_tree(sorted_events: List[_ProfilerEvent]) -> None:
423	return	423	return
424	if not push_event(ev, thread_event, unfinished_events):	424	if not push_event(ev, thread_event, unfinished_events):
425	return	425	return
426	-	426	+
427	# Cleanup remaining exit events.	427	# Cleanup remaining exit events.
428	while not unfinished_events.empty():	428	while not unfinished_events.empty():
429	_, top_event = unfinished_events.get()	429	_, top_event = unfinished_events.get()
	@@ -453,7 +453,7 @@ def get_tensor_info(sorted_events: List[_ProfilerEvent]) -> List[_RawTensorInfo]
453	elif ev.tag == _EventType.PyCall:	453	elif ev.tag == _EventType.PyCall:
454	if ev.extra_fields.key is None or ev.extra_fields.key in seen_pycalls:	454	if ev.extra_fields.key is None or ev.extra_fields.key in seen_pycalls:
455	continue	455	continue
456	-	456	+
457	seen_pycalls.add(ev.extra_fields.key)	457	seen_pycalls.add(ev.extra_fields.key)
458	if ev.extra_fields.module_parameters is not None:	458	if ev.extra_fields.module_parameters is not None:
459	for p in ev.extra_fields.module_parameters:	459	for p in ev.extra_fields.module_parameters:
	@@ -477,17 +477,17 @@ def get_tensor_info(sorted_events: List[_ProfilerEvent]) -> List[_RawTensorInfo]
477	_RawTensorInfo(t.impl, t.ptr, t.device_type, t.device_index, False, t.id)	477	_RawTensorInfo(t.impl, t.ptr, t.device_type, t.device_index, False, t.id)
478	for _, t in p.state	478	for _, t in p.state
479	)	479	)
480	-	480	+
481	return tensors	481	return tensors
482		482
483		483
484	# Assign Allocation ID for each Storage, and ID for each tensor.	484	# Assign Allocation ID for each Storage, and ID for each tensor.
485	-# A tensor has a unique id, but it can have multiple allocation IDs,	485	+# A tensor has a unique id, but it can have multiple allocation IDs,
486	-# because the tensor might use memory multiple times.	486	+# because the tensor might use memory multiple times.
487	def calculate_unique_id(sorted_events: List[_ProfilerEvent]):	487	def calculate_unique_id(sorted_events: List[_ProfilerEvent]):
488	# Step 1: Flatten events to a uniform representation	488	# Step 1: Flatten events to a uniform representation
489	tensors = get_tensor_info(sorted_events)	489	tensors = get_tensor_info(sorted_events)
490	-	490	+
491	# Step 2: Assign Allocation IDs for Storage	491	# Step 2: Assign Allocation IDs for Storage
492	counter: int = 0	492	counter: int = 0
493	storage_map: Dict[Tuple[int, int, int], int] = {}	493	storage_map: Dict[Tuple[int, int, int], int] = {}
	@@ -499,11 +499,11 @@ def calculate_unique_id(sorted_events: List[_ProfilerEvent]):
499	t.id_ref.allocation_id = storage_map[key]	499	t.id_ref.allocation_id = storage_map[key]
500	if t.is_free:	500	if t.is_free:
501	storage_map.pop(key, None)	501	storage_map.pop(key, None)
502	-	502	+
503	# Step 3: Handle allocation events which we cannot prove are for Tensor storage	503	# Step 3: Handle allocation events which we cannot prove are for Tensor storage
504	tensor_set = {t.id_ref.allocation_id for t in tensors if t.impl is not None}	504	tensor_set = {t.id_ref.allocation_id for t in tensors if t.impl is not None}
505	tensors = [t for t in tensors if t.id_ref.allocation_id in tensor_set]	505	tensors = [t for t in tensors if t.id_ref.allocation_id in tensor_set]
506	-	506	+
507	# Step 4: Assign tensor IDs using allocation IDs	507	# Step 4: Assign tensor IDs using allocation IDs
508	id_map: Dict[int, int] = {}	508	id_map: Dict[int, int] = {}
509	counter = 0	509	counter = 0
	@@ -511,7 +511,7 @@ def calculate_unique_id(sorted_events: List[_ProfilerEvent]):
511	if t.id_ref.allocation_id not in id_map:	511	if t.id_ref.allocation_id not in id_map:
512	id_map[t.id_ref.allocation_id] = counter	512	id_map[t.id_ref.allocation_id] = counter
513	counter += 1	513	counter += 1
514	-	514	+
515	# Step 5: Write back to Tensor IDs	515	# Step 5: Write back to Tensor IDs
516	for t in tensors:	516	for t in tensors:
517	if t.id_ref.allocation_id not in id_map:	517	if t.id_ref.allocation_id not in id_map:
	@@ -523,7 +523,7 @@ def calculate_unique_id(sorted_events: List[_ProfilerEvent]):
523	class EventTree:	523	class EventTree:
524	def __init__(self, profiler_path: str):	524	def __init__(self, profiler_path: str):
525	self.profiler_path = profiler_path	525	self.profiler_path = profiler_path
526	-	526	+
527	self.events: List[_ProfilerEvent] = []	527	self.events: List[_ProfilerEvent] = []
528	self.fetch_op_events(FwkFileParser(self.profiler_path))	528	self.fetch_op_events(FwkFileParser(self.profiler_path))
529	self.fetch_allocation_events(FwkFileParser(self.profiler_path))	529	self.fetch_allocation_events(FwkFileParser(self.profiler_path))
	@@ -538,7 +538,7 @@ class EventTree:
538	op_bean_list: List[TorchOpBean] = fwk_file_parser.get_file_data_by_tag(FileTag.TORCH_OP)	538	op_bean_list: List[TorchOpBean] = fwk_file_parser.get_file_data_by_tag(FileTag.TORCH_OP)
539	if not op_bean_list:	539	if not op_bean_list:
540	return	540	return
541	-	541	+
542	op_events = [_ProfilerEvent(op_bean) for op_bean in op_bean_list]	542	op_events = [_ProfilerEvent(op_bean) for op_bean in op_bean_list]
543		543
544	# Connected Autograd info to the top level annotation	544	# Connected Autograd info to the top level annotation
	@@ -548,22 +548,22 @@ class EventTree:
548	and op_events[i].extra_fields.name.startswith("autograd::engine::evaluate_function: ")):	548	and op_events[i].extra_fields.name.startswith("autograd::engine::evaluate_function: ")):
549	op_events[i].extra_fields.sequence_num = op_events[i + 1].extra_fields.sequence_num	549	op_events[i].extra_fields.sequence_num = op_events[i + 1].extra_fields.sequence_num
550	op_events[i].extra_fields.forward_tid = op_events[i + 1].extra_fields.forward_tid	550	op_events[i].extra_fields.forward_tid = op_events[i + 1].extra_fields.forward_tid
551	-	551	+
552	self.events.extend(op_events)	552	self.events.extend(op_events)
553	-	553	+
554	def fetch_allocation_events(self, fwk_file_parser: FwkFileParser) -> None:	554	def fetch_allocation_events(self, fwk_file_parser: FwkFileParser) -> None:
555	mem_bean_list: List[MemoryUseBean] = fwk_file_parser.get_file_data_by_tag(FileTag.MEMORY)	555	mem_bean_list: List[MemoryUseBean] = fwk_file_parser.get_file_data_by_tag(FileTag.MEMORY)
556	if not mem_bean_list:	556	if not mem_bean_list:
557	return	557	return
558	-	558	+
559	mem_events = [	559	mem_events = [
560	_ProfilerEvent(mem_bean)	560	_ProfilerEvent(mem_bean)
561	for mem_bean in mem_bean_list	561	for mem_bean in mem_bean_list
562	if mem_bean.data_type != _AllocEventType.BLOCK_FREE.value	562	if mem_bean.data_type != _AllocEventType.BLOCK_FREE.value
563	]	563	]
564	-	564	+
565	self.events.extend(mem_events)	565	self.events.extend(mem_events)
566	-	566	+
567	def fetch_pycall_events(self, fwk_file_parser: FwkFileParser) -> None:	567	def fetch_pycall_events(self, fwk_file_parser: FwkFileParser) -> None:
568	trace_hash_data = fwk_file_parser.get_file_data_by_tag(FileTag.PYTHON_TRACER_HASH)	568	trace_hash_data = fwk_file_parser.get_file_data_by_tag(FileTag.PYTHON_TRACER_HASH)
569	func_call_data = fwk_file_parser.get_file_data_by_tag(FileTag.PYTHON_TRACER_FUNC)	569	func_call_data = fwk_file_parser.get_file_data_by_tag(FileTag.PYTHON_TRACER_FUNC)
	@@ -573,18 +573,18 @@ class EventTree:
573	pycall_bean_list = python_trace_parser.get_pycall_data()	573	pycall_bean_list = python_trace_parser.get_pycall_data()
574	if not pycall_bean_list:	574	if not pycall_bean_list:
575	return	575	return
576	-	576	+
577	pycall_events = [_ProfilerEvent(pycall_bean) for pycall_bean in pycall_bean_list]	577	pycall_events = [_ProfilerEvent(pycall_bean) for pycall_bean in pycall_bean_list]
578		578
579	self.events.extend(pycall_events)	579	self.events.extend(pycall_events)
580	-	580	+
581	def validate_events(self) -> None:	581	def validate_events(self) -> None:
582	for ev in self.sorted_events:	582	for ev in self.sorted_events:
583	# Check the time of events is right	583	# Check the time of events is right
584	if ev.start_time_ns > ev.end_time_ns:	584	if ev.start_time_ns > ev.end_time_ns:
585	print_error_msg(f"Error in {ev.name}: {ev.start_time_ns} > {ev.end_time_ns}.")	585	print_error_msg(f"Error in {ev.name}: {ev.start_time_ns} > {ev.end_time_ns}.")
586	return	586	return
587	-	587	+
588	# Check the inputs in TorchOp	588	# Check the inputs in TorchOp
589	if ev.tag == _EventType.TorchOp:	589	if ev.tag == _EventType.TorchOp:
590	for i in ev.extra_fields.inputs:	590	for i in ev.extra_fields.inputs:

Mtorch_npu/profiler/analysis/prof_parse/_fwk_cann_relation_parser.py+1-1

	@@ -68,7 +68,7 @@ class FwkCANNRelationParser:
68	if not step_node_list:	68	if not step_node_list:
69	self.logger.warning("Get step range failed, the step node list is empty.")	69	self.logger.warning("Get step range failed, the step node list is empty.")
70	return []	70	return []
71	-	71	+
72	# Gather flow events start time in each step node	72	# Gather flow events start time in each step node
73	if not FwkFileParser(self._profiler_path).has_task_queue_data():	73	if not FwkFileParser(self._profiler_path).has_task_queue_data():
74	acl_start_time_list = sorted(list(kernel_dict.keys()))	74	acl_start_time_list = sorted(list(kernel_dict.keys()))

Mtorch_npu/profiler/analysis/prof_parse/_python_trace_parser.py+7-7

	@@ -10,10 +10,10 @@ __all__ = []
10	MODULE_NAME_DELIMITER = "######"	10	MODULE_NAME_DELIMITER = "######"
11		11
12		12
13	-class TraceTag(Enum):	13	+class TraceTag(Enum):
14	kPy_Call = 0	14	kPy_Call = 0
15	kPy_Return = 1	15	kPy_Return = 1
16	- kC_Call = 2	16	+ kC_Call = 2
17	kC_Return = 3	17	kC_Return = 3
18		18
19		19
	@@ -59,7 +59,7 @@ class PyTraceEvent:
59	@property	59	@property
60	def parent_id(self):	60	def parent_id(self):
61	return self._parent_id	61	return self._parent_id
62	-	62	+
63	@parent_id.setter	63	@parent_id.setter
64	def parent_id(self, parent_id):	64	def parent_id(self, parent_id):
65	self._parent_id = parent_id	65	self._parent_id = parent_id
	@@ -96,7 +96,7 @@ class PyTraceEvent:
96	@property	96	@property
97	def dur(self):	97	def dur(self):
98	return self._end_time - self._start_time	98	return self._end_time - self._start_time
99	-	99	+
100	@property	100	@property
101	def params(self):	101	def params(self):
102	return self._params	102	return self._params
	@@ -152,7 +152,7 @@ class PythonTraceParser:
152	trace_api_data[i] = [event.ts, event.ts + event.dur, contact_2num(event.pid, event.tid), None,	152	trace_api_data[i] = [event.ts, event.ts + event.dur, contact_2num(event.pid, event.tid), None,
153	str2id_manager.get_id_from_str(event.name), None, None, None, None, None, ApiType.PYTHON_TRACE]	153	str2id_manager.get_id_from_str(event.name), None, None, None, None, None, ApiType.PYTHON_TRACE]
154	return trace_api_data	154	return trace_api_data
155	-	155	+
156	def get_pycall_data(self) -> list:	156	def get_pycall_data(self) -> list:
157	self._gen_param_map()	157	self._gen_param_map()
158	return self._gen_python_trace_event_data()	158	return self._gen_python_trace_event_data()
	@@ -225,11 +225,11 @@ class PythonTraceParser:
225		225
226	def _gen_hash_map(self):	226	def _gen_hash_map(self):
227	self._hash_map = {hash_bean.key: hash_bean.value for hash_bean in self._hash_data}	227	self._hash_map = {hash_bean.key: hash_bean.value for hash_bean in self._hash_data}
228	-	228	+
229	def _gen_param_map(self):	229	def _gen_param_map(self):
230	if self._param_data is not None:	230	if self._param_data is not None:
231	self._param_map = {param_bean.key: param_bean.params for param_bean in self._param_data}	231	self._param_map = {param_bean.key: param_bean.params for param_bean in self._param_data}
232	-	232	+
233	@staticmethod	233	@staticmethod
234	def get_module_info_from_value(name: str, module_name_counter: dict, module_uid_map: dict):	234	def get_module_info_from_value(name: str, module_name_counter: dict, module_uid_map: dict):
235	module_name, module_uid = name.split(MODULE_NAME_DELIMITER)	235	module_name, module_uid = name.split(MODULE_NAME_DELIMITER)

Mtorch_npu/profiler/analysis/prof_view/_memory_timeline_parser.py+60-60

Mtorch_npu/profiler/analysis/prof_view/prof_db_parse/_basic_db_parser.py+2-2

	@@ -66,7 +66,7 @@ class BasicDbParser(BaseParser):
66	continue	66	continue
67	return file_path	67	return file_path
68	return ""	68	return ""
69	-	69	+
70	def create_ascend_db(self):	70	def create_ascend_db(self):
71	if not TorchDb().create_connect_db():	71	if not TorchDb().create_connect_db():
72	raise RuntimeError(f"Failed to connect to db file: {TorchDb().get_db_path()}")	72	raise RuntimeError(f"Failed to connect to db file: {TorchDb().get_db_path()}")
	@@ -89,7 +89,7 @@ class BasicDbParser(BaseParser):
89	rank_device_pairs.append([rank_id, device_id])	89	rank_device_pairs.append([rank_id, device_id])
90	TorchDb().insert_data_into_table(DbConstant.TABLE_RANK_DEVICE_MAP,	90	TorchDb().insert_data_into_table(DbConstant.TABLE_RANK_DEVICE_MAP,
91	rank_device_pairs)	91	rank_device_pairs)
92	-	92	+
93	def save_host_info_to_db(self):	93	def save_host_info_to_db(self):
94	if TorchDb().judge_table_exist(DbConstant.TABLE_HOST_INFO):	94	if TorchDb().judge_table_exist(DbConstant.TABLE_HOST_INFO):
95	return	95	return

Mtorch_npu/profiler/analysis/prof_view/prof_db_parse/_communication_db_parser.py+2-2

	@@ -97,7 +97,7 @@ class CommunicationDbParser(CommunicationParser):
97		97
98	def generate_view(self) -> None:	98	def generate_view(self) -> None:
99	self.generate_communication_db()	99	self.generate_communication_db()
100	-	100	+
101	def generate_communication_db(self):	101	def generate_communication_db(self):
102	db_files = CANNFileParser(self._profiler_path).get_file_list_by_type(CANNDataEnum.ANALYSIS_DB)	102	db_files = CANNFileParser(self._profiler_path).get_file_list_by_type(CANNDataEnum.ANALYSIS_DB)
103	if not db_files:	103	if not db_files:
	@@ -186,7 +186,7 @@ class CommunicationDbParser(CommunicationParser):
186	op_info.get(self.BANDWIDTH_GB_S), step, op_type, hccl_op_name	186	op_info.get(self.BANDWIDTH_GB_S), step, op_type, hccl_op_name
187	])	187	])
188	return res_data	188	return res_data
189	-	189	+
190	step_op_dict = {}	190	step_op_dict = {}
191	for data in matrix_data:	191	for data in matrix_data:
192	op_name = \	192	op_name = \

Mtorch_npu/profiler/analysis/prof_view/prof_db_parse/_memory_db_parser.py+4-4

Mtorch_npu/profiler/analysis/prof_view/prof_db_parse/_trace_step_time_db_parser.py+3-3

	@@ -140,7 +140,7 @@ class TraceStepTimeDbParser(BaseParser):
140	return	140	return
141	if TorchDb().judge_table_exist(DbConstant.TABLE_COMPUTE_TASK_INFO):	141	if TorchDb().judge_table_exist(DbConstant.TABLE_COMPUTE_TASK_INFO):
142	sql = """	142	sql = """
143	- SELECT	143	+ SELECT
144	STRING_IDS.value,	144	STRING_IDS.value,
145	task.startNs,	145	task.startNs,
146	task.endNs,	146	task.endNs,
	@@ -163,14 +163,14 @@ class TraceStepTimeDbParser(BaseParser):
163	connectionId	163	connectionId
164	FROM COMMUNICATION_OP c	164	FROM COMMUNICATION_OP c
165	)	165	)
166	- SELECT	166	+ SELECT
167	comm.opName,	167	comm.opName,
168	comm.startNs,	168	comm.startNs,
169	comm.endNs,	169	comm.endNs,
170	t.deviceId	170	t.deviceId
171	FROM comm_info comm	171	FROM comm_info comm
172	JOIN (	172	JOIN (
173	- SELECT	173	+ SELECT
174	connectionId,	174	connectionId,
175	deviceId	175	deviceId
176	FROM TASK	176	FROM TASK

Mtorch_npu/testing/common_utils.py+2-2

	@@ -206,7 +206,7 @@ class SupportedDevices:
206	reason = f"Only run on {repr(self.supported_devices)}, current device is {device_name}."	206	reason = f"Only run on {repr(self.supported_devices)}, current device is {device_name}."
207	raise unittest.SkipTest(reason)	207	raise unittest.SkipTest(reason)
208	return fn(slf, args, *kwargs)	208	return fn(slf, args, *kwargs)
209	-	209	+
210	return dep_fn	210	return dep_fn
211		211
212		212
	@@ -214,7 +214,7 @@ class SkipIfNotGteCANNVersion:
214	def __init__(self, base_version, module="CANN"):	214	def __init__(self, base_version, module="CANN"):
215	self.base_version = base_version	215	self.base_version = base_version
216	self.module = module	216	self.module = module
217	-	217	+
218	def __call__(self, fn):	218	def __call__(self, fn):
219	@wraps(fn)	219	@wraps(fn)
220	def func(slf, args, *kwargs):	220	def func(slf, args, *kwargs):

Mtorch_npu/testing/decorator.py+3-3

	@@ -81,7 +81,7 @@ def gen_ops_testcase(cls, func, name, keys, value, op_info):
81		81
82	def gen_op_input(testcase, func, op_info):	82	def gen_op_input(testcase, func, op_info):
83	data = {	83	data = {
84	- 'dtype': func.dtypes if hasattr(func, "dtypes") else op_info.dtypesIfNPU,	84	+ 'dtype': func.dtypes if hasattr(func, "dtypes") else op_info.dtypesIfNPU,
85	'npu_format': func.formats if hasattr(func, "formats") else op_info.formats	85	'npu_format': func.formats if hasattr(func, "formats") else op_info.formats
86	}	86	}
87		87
	@@ -98,7 +98,7 @@ def instantiate_ops_tests(op_db):
98		98
99	def wrapper(cls):	99	def wrapper(cls):
100	testcases = [x for x in dir(cls) if x.startswith('test_')]	100	testcases = [x for x in dir(cls) if x.startswith('test_')]
101	- for testcase in testcases:	101	+ for testcase in testcases:
102	if hasattr(cls, testcase):	102	if hasattr(cls, testcase):
103	func = getattr(cls, testcase)	103	func = getattr(cls, testcase)
104	for op_info in op_db:	104	for op_info in op_db:
	@@ -112,7 +112,7 @@ def instantiate_ops_tests(op_db):
112	delattr(cls, testcase)	112	delattr(cls, testcase)
113		113
114	return cls	114	return cls
115	-	115	+
116	return wrapper	116	return wrapper
117		117
118		118

Mtorch_npu/testing/testcase.py+4-4

	@@ -154,7 +154,7 @@ class TestCase(expecttest.TestCase):
154	self.assertEqual(tc._values(), t._values())	154	self.assertEqual(tc._values(), t._values())
155		155
156	return tg	156	return tg
157	-	157	+
158	def assertRtolEqual(self, x, y, prec=1.e-4, prec16=1.e-3, auto_trans_dtype=False, message=None):	158	def assertRtolEqual(self, x, y, prec=1.e-4, prec16=1.e-3, auto_trans_dtype=False, message=None):
159		159
160	def _assertRtolEqual(x, y, prec, prec16, message):	160	def _assertRtolEqual(x, y, prec, prec16, message):
	@@ -198,7 +198,7 @@ class TestCase(expecttest.TestCase):
198	self.fail("result error!")	198	self.fail("result error!")
199	return	199	return
200	x = x.detach().cpu().numpy()	200	x = x.detach().cpu().numpy()
201	- y = y.detach().cpu().numpy()	201	+ y = y.detach().cpu().numpy()
202	elif isinstance(x, Number) and isinstance(y, Number):	202	elif isinstance(x, Number) and isinstance(y, Number):
203	x = np.array(x)	203	x = np.array(x)
204	y = np.array(y)	204	y = np.array(y)
	@@ -208,7 +208,7 @@ class TestCase(expecttest.TestCase):
208	self.fail("shape error")	208	self.fail("shape error")
209	if (x.dtype != y.dtype):	209	if (x.dtype != y.dtype):
210	self.fail("dtype error")	210	self.fail("dtype error")
211	- dtype_list = [np.bool_, np.uint16, np.int16, np.int32, np.float16,	211	+ dtype_list = [np.bool_, np.uint16, np.int16, np.int32, np.float16,
212	np.float32, np.int8, np.uint8, np.int64, np.float64]	212	np.float32, np.int8, np.uint8, np.int64, np.float64]
213	if x.dtype not in dtype_list:	213	if x.dtype not in dtype_list:
214	self.fail("required dtype in [np.bool_, np.uint16, np.int16, " +	214	self.fail("required dtype in [np.bool_, np.uint16, np.int16, " +
	@@ -502,7 +502,7 @@ class TestCase(expecttest.TestCase):
502	def run(self, result=None):	502	def run(self, result=None):
503	# run test to precompile operators	503	# run test to precompile operators
504	super(TestCase, self).run(result)	504	super(TestCase, self).run(result)
505	-	505	+
506	if PERF_TEST_ENABLE:	506	if PERF_TEST_ENABLE:
507	performanceResult = TestResult()	507	performanceResult = TestResult()
508	startTime = time.perf_counter()	508	startTime = time.perf_counter()

Mtorch_npu/utils/_dynamo.py+1-1

	@@ -187,7 +187,7 @@ def patch_inductor_wrapper():
187	or torch._inductor.config.npu_backend == "mlir"	187	or torch._inductor.config.npu_backend == "mlir"
188	):	188	):
189	os.environ["TORCHINDUCTOR_NPU_BACKEND"] = "mlir"	189	os.environ["TORCHINDUCTOR_NPU_BACKEND"] = "mlir"
190	- device_id = torch_npu.npu.current_device()	190	+ device_id = torch_npu.npu.current_device()
191	torch_npu._C._recovery_all_npu_stream(device_id)	191	torch_npu._C._recovery_all_npu_stream(device_id)
192		192
193	elif (	193	elif (

Mtorch_npu/utils/_graph_tree.py+1-1

	@@ -228,7 +228,7 @@ def npugraphify_impl(
228		228
229	else:	229	else:
230	copy_indices = [	230	copy_indices = [
231	- idx	231	+ idx
232	for idx in range(len(static_inputs))	232	for idx in range(len(static_inputs))
233	if idx not in static_input_idxs	233	if idx not in static_input_idxs
234	]	234	]

Mtorch_npu/utils/_inductor.py+5-5

	@@ -49,8 +49,8 @@ def patch_register_philox_rand():
49	def get_register_philox_rand_patch():	49	def get_register_philox_rand_patch():
50	name = "philox_rand"	50	name = "philox_rand"
51	schema = "(SymInt[] size, Tensor seed, Tensor offset, int[]? stride, Device? device=None, ScalarType? dtype=None) -> (Tensor, Tensor)" # noqa: B950	51	schema = "(SymInt[] size, Tensor seed, Tensor offset, int[]? stride, Device? device=None, ScalarType? dtype=None) -> (Tensor, Tensor)" # noqa: B950
52	-	52	+
53	-	53	+
54	def _philox_rand_meta(	54	def _philox_rand_meta(
55	shape: torch.Size,	55	shape: torch.Size,
56	seed: torch.Tensor,	56	seed: torch.Tensor,
	@@ -66,7 +66,7 @@ def patch_register_philox_rand():
66	offset = philox_rand_offset_meta(shape)	66	offset = philox_rand_offset_meta(shape)
67	return (random_values, offset)	67	return (random_values, offset)
68		68
69	-	69	+
70	def _philox_rand(	70	def _philox_rand(
71	shape: torch.Size,	71	shape: torch.Size,
72	seed: torch.Tensor,	72	seed: torch.Tensor,
	@@ -80,13 +80,13 @@ def patch_register_philox_rand():
80	else:	80	else:
81	devices = [device]	81	devices = [device]
82		82
83	- with torch.random.fork_rng(devices, device_type="npu"):	83	+ with torch.random.fork_rng(devices, device_type="npu"):
84	CUDARngStateHelper.set_torch_state_tensor(seed, offset)	84	CUDARngStateHelper.set_torch_state_tensor(seed, offset)
85	random_values = torch.rand(shape, device=device, dtype=dtype)	85	random_values = torch.rand(shape, device=device, dtype=dtype)
86		86
87	return random_values, philox_rand_offset(shape)	87	return random_values, philox_rand_offset(shape)
88		88
89	-	89	+
90	register_rng_prim(	90	register_rng_prim(
91	name=name,	91	name=name,
92	schema=schema,	92	schema=schema,

Mtorch_npu/utils/_npu_meta_registration.py+1-1

	@@ -87,7 +87,7 @@ def patch_torch_inductor_decompositions():
87	don't accidentally overwrite unrelated inductor decompositions.	87	don't accidentally overwrite unrelated inductor decompositions.
88	'''	88	'''
89	import torch._inductor.decomposition as inductor_decomposition	89	import torch._inductor.decomposition as inductor_decomposition
90	-	90	+
91	for op_overload in inductor_decomp_table:	91	for op_overload in inductor_decomp_table:
92	if op_overload in npu_meta_table:	92	if op_overload in npu_meta_table:
93	inductor_decomposition.decompositions[op_overload] = npu_meta_table[op_overload]	93	inductor_decomposition.decompositions[op_overload] = npu_meta_table[op_overload]

Mtorch_npu/utils/_step.py+5-5

	@@ -37,7 +37,7 @@ class PerfDumpState:
37	if sub_module != module:	37	if sub_module != module:
38	module_list.append(sub_module)	38	module_list.append(sub_module)
39	self.module_dict[module] = module_list	39	self.module_dict[module] = module_list
40	-	40	+
41	def is_child_module(self, module):	41	def is_child_module(self, module):
42	for item in self.module_dict.items():	42	for item in self.module_dict.items():
43	if module in item[1]:	43	if module in item[1]:
	@@ -72,7 +72,7 @@ def _validate_path(path):
72	return True	72	return True
73	else:	73	else:
74	return False	74	return False
75	-	75	+
76		76
77	def _get_perf_dump_path():	77	def _get_perf_dump_path():
78	perf_dump_path = os.environ.get("PERF_DUMP_PATH")	78	perf_dump_path = os.environ.get("PERF_DUMP_PATH")
	@@ -85,7 +85,7 @@ def _get_perf_dump_path():
85	def delete_pref_pt_logs(perf_dump_path, device_id):	85	def delete_pref_pt_logs(perf_dump_path, device_id):
86	log_pattern = os.path.join(perf_dump_path, f"perf_pt__{device_id}.log")	86	log_pattern = os.path.join(perf_dump_path, f"perf_pt__{device_id}.log")
87	log_files = glob.glob(log_pattern)	87	log_files = glob.glob(log_pattern)
88	-	88	+
89	for log_file in log_files:	89	for log_file in log_files:
90	if os.path.islink(log_file):	90	if os.path.islink(log_file):
91	continue	91	continue
	@@ -101,9 +101,9 @@ def _get_uuid():
101		101
102	if master_addr is None or master_port is None:	102	if master_addr is None or master_port is None:
103	return "127.0.0.1_8888"	103	return "127.0.0.1_8888"
104	-	104	+
105	return master_addr + "_" + master_port	105	return master_addr + "_" + master_port
106	-	106	+
107		107
108	def _setup_logger(name, path):	108	def _setup_logger(name, path):
109	logger = logging.getLogger(name)	109	logger = logging.getLogger(name)

Mtorch_npu/utils/dlpack.py+10-10

Mtorch_npu/utils/flops_count.py+3-3

	@@ -7,10 +7,10 @@ __all__ = []
7	class _FlopsCounter:	7	class _FlopsCounter:
8	def __init__(self, ):	8	def __init__(self, ):
9	self.flop_count_instance = torch_npu._C._flops_count._FlopCountContext.GetInstance()	9	self.flop_count_instance = torch_npu._C._flops_count._FlopCountContext.GetInstance()
10	-	10	+
11	def __enter__(self):	11	def __enter__(self):
12	self.count_enable()	12	self.count_enable()
13	-	13	+
14	def __exit__(self):	14	def __exit__(self):
15	self.count_disable()	15	self.count_disable()
16		16
	@@ -20,7 +20,7 @@ class _FlopsCounter:
20	def stop(self):	20	def stop(self):
21	self.flop_count_instance.disable()	21	self.flop_count_instance.disable()
22	self.flop_count_instance.reset()	22	self.flop_count_instance.reset()
23	-	23	+
24	def pause(self):	24	def pause(self):
25	self.flop_count_instance.pause()	25	self.flop_count_instance.pause()
26		26

Mtorch_npu/utils/profiler.py+4-4

	@@ -26,7 +26,7 @@ class Profile(object):
26	save_path: str = "./npu_profiling",	26	save_path: str = "./npu_profiling",
27	profile_type: str = None,	27	profile_type: str = None,
28	use_npu=True,	28	use_npu=True,
29	- record_shape: bool = True,	29	+ record_shape: bool = True,
30	experimental_config: Optional[_ExperimentalConfig] = torch_npu.profiler._ExperimentalConfig(	30	experimental_config: Optional[_ExperimentalConfig] = torch_npu.profiler._ExperimentalConfig(
31	profiler_level=torch_npu.profiler.ProfilerLevel.Level2	31	profiler_level=torch_npu.profiler.ProfilerLevel.Level2
32	),	32	),
	@@ -72,9 +72,9 @@ class Profile(object):
72	raise ValueError("Args '%s' invaild, expect args '%s' ." % (kwargs.keys(), ascend_profiler_args_set) +	72	raise ValueError("Args '%s' invaild, expect args '%s' ." % (kwargs.keys(), ascend_profiler_args_set) +
73	prof_error(ErrCode.VALUE))	73	prof_error(ErrCode.VALUE))
74	self.prof = torch_npu.profiler.profile(	74	self.prof = torch_npu.profiler.profile(
75	- on_trace_ready=torch_npu.profiler.tensorboard_trace_handler(self.save_path),	75	+ on_trace_ready=torch_npu.profiler.tensorboard_trace_handler(self.save_path),
76	- experimental_config=self.experimental_config,	76	+ experimental_config=self.experimental_config,
77	- record_shapes=self.record_shape,	77	+ record_shapes=self.record_shape,
78	**kwargs	78	**kwargs
79	)	79	)
80		80

Mtorch_npu/utils/tensor_methods.py+1-1

	@@ -70,7 +70,7 @@ class _NPUTensortypeCache(object):
70	def _npu_type(self, dtype=None, non_blocking=False, **kwargs):	70	def _npu_type(self, dtype=None, non_blocking=False, **kwargs):
71	if dtype is None:	71	if dtype is None:
72	return self.type_raw(dtype, non_blocking, **kwargs)	72	return self.type_raw(dtype, non_blocking, **kwargs)
73	-	73	+
74	_NPUTensortypeCache.tensortype_list_dict_init()	74	_NPUTensortypeCache.tensortype_list_dict_init()
75	if isinstance(dtype, str) and dtype in _NPUTensortypeCache.get_tensortype_dict():	75	if isinstance(dtype, str) and dtype in _NPUTensortypeCache.get_tensortype_dict():
76	tensortype_class = _NPUTensortypeCache.get_tensortype_dict()[dtype]	76	tensortype_class = _NPUTensortypeCache.get_tensortype_dict()[dtype]

Mtorchnpugen/__init__.py+1-1

	@@ -22,6 +22,6 @@ def _write_if_changed_security(self, filename: str, contents: str) -> None:
22		22
23	def apply_codegen_patches():	23	def apply_codegen_patches():
24	torchgen.gen.FileManager._write_if_changed = _write_if_changed_security	24	torchgen.gen.FileManager._write_if_changed = _write_if_changed_security
25	-	25	+
26		26
27	apply_codegen_patches()	27	apply_codegen_patches()

Mtorchnpugen/autograd/gen_autograd_functions.py+1-1

	@@ -47,7 +47,7 @@ def gen_autograd_functions_python(
47	infos,	47	infos,
48	key_fn=lambda info: info.name,	48	key_fn=lambda info: info.name,
49	base_env={	49	base_env={
50	- "generated_comment":	50	+ "generated_comment":
51	f"@ generated from {fm.template_dir_for_comments()}/python_functions.cpp",	51	f"@ generated from {fm.template_dir_for_comments()}/python_functions.cpp",
52	},	52	},
53	env_callable=lambda info: {	53	env_callable=lambda info: {

Mtorchnpugen/autograd/gen_variable_type.py+2-2

	@@ -56,7 +56,7 @@ def gen_variable_type(
56	template_path: str,	56	template_path: str,
57	) -> None:	57	) -> None:
58	"""Generate VariableType.cpp body	58	"""Generate VariableType.cpp body
59	-	59	+
60	Generate variable type definition for torch and npu method here.	60	Generate variable type definition for torch and npu method here.
61	"""	61	"""
62	fm = FileManager(install_dir=out, template_dir=template_path, dry_run=False)	62	fm = FileManager(install_dir=out, template_dir=template_path, dry_run=False)
	@@ -80,7 +80,7 @@ def gen_variable_type_head(
80	fns_with_diff_infos: List[NativeFunctionWithDifferentiabilityInfo],	80	fns_with_diff_infos: List[NativeFunctionWithDifferentiabilityInfo],
81	template_path: str,	81	template_path: str,
82	) -> None:	82	) -> None:
83	-	83	+
84	"""Generate VariableType.h body	84	"""Generate VariableType.h body
85	"""	85	"""
86	fm = FileManager(install_dir=out, template_dir=template_path, dry_run=False)	86	fm = FileManager(install_dir=out, template_dir=template_path, dry_run=False)

Mtorchnpugen/autograd/utils.py+1-1

	@@ -46,7 +46,7 @@ def parse_derivatives(
46	# original code logic	46	# original code logic
47	derivatives_path = str(Path(autograd_dir).parents[1].joinpath(	47	derivatives_path = str(Path(autograd_dir).parents[1].joinpath(
48	f'third_party/op-plugin/op_plugin/config/v{VERSION_PART[0]}r{VERSION_PART[1]}/derivatives.yaml'	48	f'third_party/op-plugin/op_plugin/config/v{VERSION_PART[0]}r{VERSION_PART[1]}/derivatives.yaml'
49	- ))	49	+ ))
50		50
51	differentiability_infos, _ = load_derivatives(	51	differentiability_infos, _ = load_derivatives(
52	derivatives_path, native_functions_path, tags_path)	52	derivatives_path, native_functions_path, tags_path)

Mtorchnpugen/codegen_ops_info.py+1-1

	@@ -15,7 +15,7 @@ from torchnpugen.utils import PathManager
15	project_path = Path(os.path.dirname(__file__)).parent	15	project_path = Path(os.path.dirname(__file__)).parent
16	op_plugin_info_path = os.path.realpath(os.path.join(	16	op_plugin_info_path = os.path.realpath(os.path.join(
17	project_path,	17	project_path,
18	- f'third_party/op-plugin/test/test_v{VERSION_PART[0]}r{VERSION_PART[1]}_ops',	18	+ f'third_party/op-plugin/test/test_v{VERSION_PART[0]}r{VERSION_PART[1]}_ops',
19	"unsupported_ops_info.yaml"))	19	"unsupported_ops_info.yaml"))
20	torch_npu_info_path = os.path.realpath(os.path.join(project_path, "test", "unsupported_ops_info.yaml"))	20	torch_npu_info_path = os.path.realpath(os.path.join(project_path, "test", "unsupported_ops_info.yaml"))
21		21

Mtorchnpugen/templates/npu_testing_utils.py+1-1

	@@ -5,7 +5,7 @@ from torch.testing._internal.common_methods_invocations import op_db, python_ref
5	from torch.testing._internal.opinfo.core import DecorateInfo	5	from torch.testing._internal.opinfo.core import DecorateInfo
6		6
7	"""	7	"""
8	-strategy: Due to the restriction of NPU operators.	8	+strategy: Due to the restriction of NPU operators.
9	patch the data classes to avoid unsupported cases.	9	patch the data classes to avoid unsupported cases.
10	"""	10	"""
11		11

	@@ -7,25 +7,25 @@ import os
7	os.environ["NPU_INDUCTOR_FALLBACK_LIST"] = "aten.div,aten.add.Tensor"	7	os.environ["NPU_INDUCTOR_FALLBACK_LIST"] = "aten.div,aten.add.Tensor"
8		8
9	class TestFallback(TestUtils):	9	class TestFallback(TestUtils):
10	-	10	+
11	def add_op(self, x, y):	11	def add_op(self, x, y):
12	return x / y	12	return x / y
13	-	13	+
14	def test_add_fallback_detection(self):	14	def test_add_fallback_detection(self):
15	-	15	+
16	compiled_add = torch.compile(self.add_op, backend="inductor")	16	compiled_add = torch.compile(self.add_op, backend="inductor")
17	-	17	+
18	x = torch.randn(4, 4, dtype=torch.float32).to("npu")	18	x = torch.randn(4, 4, dtype=torch.float32).to("npu")
19	y = torch.randn(4, 4, dtype=torch.float32).to("npu")	19	y = torch.randn(4, 4, dtype=torch.float32).to("npu")
20		20
21	_ , codes = run_and_get_code(compiled_add, x, y)	21	_ , codes = run_and_get_code(compiled_add, x, y)
22		22
23	self.assertTrue('unk_fused_div' not in codes[0])	23	self.assertTrue('unk_fused_div' not in codes[0])
24	-	24	+
25	def test_add_fallback_detection_mlir(self):	25	def test_add_fallback_detection_mlir(self):
26		26
27	compiled_add = torch.compile(self.add_op, backend="inductor", options={"npu_backend": "mlir"})	27	compiled_add = torch.compile(self.add_op, backend="inductor", options={"npu_backend": "mlir"})
28	-	28	+
29	x = torch.randn(4, 4, dtype=torch.float32).to("npu")	29	x = torch.randn(4, 4, dtype=torch.float32).to("npu")
30	y = torch.randn(4, 4, dtype=torch.float32).to("npu")	30	y = torch.randn(4, 4, dtype=torch.float32).to("npu")
31		31

	@@ -59,7 +59,7 @@ class OptionsTest(TestCase):
59	1. HCCL config is correctly applied to default process group	59	1. HCCL config is correctly applied to default process group
60	2. New process group inherits the correct HCCL configuration	60	2. New process group inherits the correct HCCL configuration
61	3. all_reduce executes successfully on NPU	61	3. all_reduce executes successfully on NPU
62	-	62	+
63	Args:	63	Args:
64	rank (int): Current process rank	64	rank (int): Current process rank
65	ranks (list[int]): List of ranks in the process group	65	ranks (list[int]): List of ranks in the process group
	@@ -86,32 +86,32 @@ class OptionsTest(TestCase):
86	# 4. Move input tensor to target NPU device (specify rank to avoid device conflict)	86	# 4. Move input tensor to target NPU device (specify rank to avoid device conflict)
87	input1 = input1.npu(rank)	87	input1 = input1.npu(rank)
88	test_case.assertEqual(	88	test_case.assertEqual(
89	- input1.device,	89	+ input1.device,
90	- torch.device(f'npu:{rank}'),	90	+ torch.device(f'npu:{rank}'),
91	"Tensor not correctly moved to target NPU device"	91	"Tensor not correctly moved to target NPU device"
92	)	92	)
93		93
94	# 5. Execute all_reduce on default process group and validate configuration	94	# 5. Execute all_reduce on default process group and validate configuration
95	dist.all_reduce(input1)	95	dist.all_reduce(input1)
96	-	96	+
97	# Get default process group's HCCL backend	97	# Get default process group's HCCL backend
98	default_pg = c10d._get_default_group()._get_backend(torch.device(f'npu:{rank}'))	98	default_pg = c10d._get_default_group()._get_backend(torch.device(f'npu:{rank}'))
99	-	99	+
100	# Validate full HCCL configuration	100	# Validate full HCCL configuration
101	test_case.assertEqual(	101	test_case.assertEqual(
102	- default_pg.options.hccl_config,	102	+ default_pg.options.hccl_config,
103	cls.HCCL_DEFAULT_CONFIG,	103	cls.HCCL_DEFAULT_CONFIG,
104	"Default process group HCCL config does not match expected configuration"	104	"Default process group HCCL config does not match expected configuration"
105	)	105	)
106	-	106	+
107	# Validate individual config items (add default value to avoid KeyError)	107	# Validate individual config items (add default value to avoid KeyError)
108	test_case.assertEqual(	108	test_case.assertEqual(
109	- default_pg.options.hccl_config.get("hccl_exec_timeout", -1),	109	+ default_pg.options.hccl_config.get("hccl_exec_timeout", -1),
110	500,	110	500,
111	"hccl_exec_timeout config value mismatch"	111	"hccl_exec_timeout config value mismatch"
112	)	112	)
113	test_case.assertEqual(	113	test_case.assertEqual(
114	- default_pg.options.hccl_config.get("hccl_algo", ""),	114	+ default_pg.options.hccl_config.get("hccl_algo", ""),
115	"allreduce=level0:NA;level1:ring/allgather=level0:NA;level1:H-D_R",	115	"allreduce=level0:NA;level1:ring/allgather=level0:NA;level1:H-D_R",
116	"hccl_algo config value mismatch"	116	"hccl_algo config value mismatch"
117	)	117	)
	@@ -119,22 +119,22 @@ class OptionsTest(TestCase):
119	# 6. Create new process group with same options and validate configuration	119	# 6. Create new process group with same options and validate configuration
120	new_pg = dist.new_group(backend='hccl', ranks=ranks, pg_options=options)	120	new_pg = dist.new_group(backend='hccl', ranks=ranks, pg_options=options)
121	test_case.assertTrue(new_pg is not None, "Failed to create new HCCL process group")	121	test_case.assertTrue(new_pg is not None, "Failed to create new HCCL process group")
122	-	122	+
123	# Validate new process group's HCCL configuration	123	# Validate new process group's HCCL configuration
124	new_pg_backend = new_pg._get_backend(torch.device(f'npu:{rank}'))	124	new_pg_backend = new_pg._get_backend(torch.device(f'npu:{rank}'))
125	test_case.assertEqual(	125	test_case.assertEqual(
126	- new_pg_backend.options.hccl_config,	126	+ new_pg_backend.options.hccl_config,
127	cls.HCCL_DEFAULT_CONFIG,	127	cls.HCCL_DEFAULT_CONFIG,
128	"New process group HCCL config does not match expected configuration"	128	"New process group HCCL config does not match expected configuration"
129	)	129	)
130	-	130	+
131	# Execute all_reduce on new process group	131	# Execute all_reduce on new process group
132	dist.all_reduce(input1, group=new_pg)	132	dist.all_reduce(input1, group=new_pg)
133		133
134	except Exception as e:	134	except Exception as e:
135	# Capture exceptions and mark test as failed	135	# Capture exceptions and mark test as failed
136	test_case.fail(f"Test execution failed with error: {str(e)}")	136	test_case.fail(f"Test execution failed with error: {str(e)}")
137	-	137	+
138	finally:	138	finally:
139	# 7. Clean up resources to prevent memory leaks	139	# 7. Clean up resources to prevent memory leaks
140	# Destroy custom process group if created	140	# Destroy custom process group if created
	@@ -156,7 +156,7 @@ class OptionsTest(TestCase):
156	with test_case.assertRaises(RuntimeError) as cm:	156	with test_case.assertRaises(RuntimeError) as cm:
157	OptionsTest._init_dist_hccl(rank, options, world_size)	157	OptionsTest._init_dist_hccl(rank, options, world_size)
158	dist.all_reduce(input1)	158	dist.all_reduce(input1)
159	-	159	+
160	test_case.assertTrue(error_expect in str(cm.exception),	160	test_case.assertTrue(error_expect in str(cm.exception),
161	f"Expected error messages '{error_expect}' not found in actual error: {str(cm.exception)}")	161	f"Expected error messages '{error_expect}' not found in actual error: {str(cm.exception)}")
162		162
	@@ -207,16 +207,16 @@ class OptionsTest(TestCase):
207	exceed_length = max_length + 100	207	exceed_length = max_length + 100
208	long_algo_str = "a" * exceed_length	208	long_algo_str = "a" * exceed_length
209	hccl_config = {"hccl_algo": long_algo_str}	209	hccl_config = {"hccl_algo": long_algo_str}
210	-	210	+
211	options = torch_npu._C._distributed_c10d.ProcessGroupHCCL.Options()	211	options = torch_npu._C._distributed_c10d.ProcessGroupHCCL.Options()
212	options.hccl_config = hccl_config	212	options.hccl_config = hccl_config
213	input1 = input1.npu()	213	input1 = input1.npu()
214	-	214	+
215	test_case = TestCase()	215	test_case = TestCase()
216	try:	216	try:
217	OptionsTest._init_dist_hccl(rank, options, world_size)	217	OptionsTest._init_dist_hccl(rank, options, world_size)
218	dist.all_reduce(input1)	218	dist.all_reduce(input1)
219	-	219	+
220	default_pg = c10d._get_default_group()._get_backend(torch.device('npu'))	220	default_pg = c10d._get_default_group()._get_backend(torch.device('npu'))
221	actual_algo = default_pg.options.hccl_config.get("hccl_algo", "")	221	actual_algo = default_pg.options.hccl_config.get("hccl_algo", "")
222	test_case.assertEqual(len(actual_algo), max_length - 1,	222	test_case.assertEqual(len(actual_algo), max_length - 1,

	@@ -38,7 +38,7 @@ class TestTorchairNoInit(TestCase):
38	for m in sys.modules:	38	for m in sys.modules:
39	if hasattr(sys.modules[m], '_attr_test_hasattr'):	39	if hasattr(sys.modules[m], '_attr_test_hasattr'):
40	setattr(sys.modules[m], '_attr_test_hasattr', 1)	40	setattr(sys.modules[m], '_attr_test_hasattr', 1)
41	-	41	+
42	torchair = sys.modules.get('torchair', None)	42	torchair = sys.modules.get('torchair', None)
43	self.assertTrue(torchair is not None)	43	self.assertTrue(torchair is not None)
44	self.assertTrue(not hasattr(torchair, '_attr_test_hasattr'))	44	self.assertTrue(not hasattr(torchair, '_attr_test_hasattr'))
	@@ -48,13 +48,13 @@ class TestTorchairNoInit(TestCase):
48	for m in sys.modules.values():	48	for m in sys.modules.values():
49	if getattr(m, '__warningregistry__', None):	49	if getattr(m, '__warningregistry__', None):
50	m.__warningregistry__ = {}	50	m.__warningregistry__ = {}
51	-	51	+
52	self._check_torchair_no_init()	52	self._check_torchair_no_init()
53	-	53	+
54	def test_attribute_error(self):	54	def test_attribute_error(self):
55	torchair = sys.modules.get('torchair', None)	55	torchair = sys.modules.get('torchair', None)
56	self.assertTrue(torchair is not None)	56	self.assertTrue(torchair is not None)
57	- with self.assertRaisesRegex(AttributeError,	57	+ with self.assertRaisesRegex(AttributeError,
58	"Try to get torchair's attr `get_npu_backend` before torchair is initialized."):	58	"Try to get torchair's attr `get_npu_backend` before torchair is initialized."):
59	torchair.get_npu_backend()	59	torchair.get_npu_backend()
60	self._check_torchair_no_init()	60	self._check_torchair_no_init()

	@@ -50,7 +50,7 @@ class TestIFAAclgraphUpdateSupportBlocking(TestCase):
50	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535, workspace=workspace,	50	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535, workspace=workspace,
51	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length, out=[output, softmax_lse])	51	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length, out=[output, softmax_lse])
52	handle = torch.npu.graph_task_group_end(stream)	52	handle = torch.npu.graph_task_group_end(stream)
53	-	53	+
54	with torch.npu.stream(update_stream):	54	with torch.npu.stream(update_stream):
55	torch.npu.graph_task_update_begin(update_stream, handle)	55	torch.npu.graph_task_update_begin(update_stream, handle)
56	torch_npu.npu_fused_infer_attention_score.out(	56	torch_npu.npu_fused_infer_attention_score.out(
	@@ -91,7 +91,7 @@ class TestIFAAclgraphUpdateSupportBlocking(TestCase):
91	torch_npu.npu_fused_infer_attention_score.out(	91	torch_npu.npu_fused_infer_attention_score.out(
92	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535, workspace=workspace,	92	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535, workspace=workspace,
93	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length, out=[output, softmax_lse])	93	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length, out=[output, softmax_lse])
94	-	94	+
95	g.update(cpu_update_input=[{"actual_seq_lengths": length_new}])	95	g.update(cpu_update_input=[{"actual_seq_lengths": length_new}])
96	g.replay()	96	g.replay()
97	self.assertEqual(output.cpu(), res_src[0].cpu())	97	self.assertEqual(output.cpu(), res_src[0].cpu())
	@@ -121,7 +121,7 @@ class TestIFAAclgraphUpdateSupportBlocking(TestCase):
121	output, softmax_lse = torch_npu.npu_fused_infer_attention_score(	121	output, softmax_lse = torch_npu.npu_fused_infer_attention_score(
122	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535,	122	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535,
123	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length)	123	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length)
124	-	124	+
125	g.update(cpu_update_input=[{"actual_seq_lengths": length_new}])	125	g.update(cpu_update_input=[{"actual_seq_lengths": length_new}])
126	g.replay()	126	g.replay()
127	self.assertEqual(output.cpu(), res_src[0].cpu())	127	self.assertEqual(output.cpu(), res_src[0].cpu())

	@@ -51,7 +51,7 @@ class TestIFAAclgraphUpdate(TestCase):
51	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535, workspace=workspace,	51	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535, workspace=workspace,
52	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length, out=[output, softmax_lse])	52	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length, out=[output, softmax_lse])
53	handle = torch.npu.graph_task_group_end(stream)	53	handle = torch.npu.graph_task_group_end(stream)
54	-	54	+
55	with torch.npu.stream(update_stream):	55	with torch.npu.stream(update_stream):
56	torch.npu.graph_task_update_begin(update_stream, handle)	56	torch.npu.graph_task_update_begin(update_stream, handle)
57	torch_npu.npu_fused_infer_attention_score.out(	57	torch_npu.npu_fused_infer_attention_score.out(
	@@ -100,7 +100,7 @@ class TestIFAAclgraphUpdate(TestCase):
100	torch_npu.npu_fused_infer_attention_score.out(	100	torch_npu.npu_fused_infer_attention_score.out(
101	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535, workspace=workspace,	101	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535, workspace=workspace,
102	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length, out=[output, softmax_lse])	102	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length, out=[output, softmax_lse])
103	-	103	+
104	g.update(cpu_update_input=[{"actual_seq_lengths": length_new}])	104	g.update(cpu_update_input=[{"actual_seq_lengths": length_new}])
105	g.replay()	105	g.replay()
106	self.assertEqual(output.cpu(), res_src[0].cpu())	106	self.assertEqual(output.cpu(), res_src[0].cpu())
	@@ -130,7 +130,7 @@ class TestIFAAclgraphUpdate(TestCase):
130	output, softmax_lse = torch_npu.npu_fused_infer_attention_score(	130	output, softmax_lse = torch_npu.npu_fused_infer_attention_score(
131	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535,	131	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535,
132	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length)	132	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length)
133	-	133	+
134	g.update(cpu_update_input=[{"actual_seq_lengths": length_new}])	134	g.update(cpu_update_input=[{"actual_seq_lengths": length_new}])
135	g.replay()	135	g.replay()
136	self.assertEqual(output.cpu(), res_src[0].cpu())	136	self.assertEqual(output.cpu(), res_src[0].cpu())
	@@ -164,7 +164,7 @@ class TestIFAAclgraphUpdate(TestCase):
164	torch_npu.npu_fused_infer_attention_score_v2.out(	164	torch_npu.npu_fused_infer_attention_score_v2.out(
165	query, key, value, num_query_heads=32, input_layout="BNSD", softmax_scale=scale, pre_tokens=65535, workspace=workspace,	165	query, key, value, num_query_heads=32, input_layout="BNSD", softmax_scale=scale, pre_tokens=65535, workspace=workspace,
166	next_tokens=65535, return_softmax_lse=False, actual_seq_qlen=length, out=[output, softmax_lse])	166	next_tokens=65535, return_softmax_lse=False, actual_seq_qlen=length, out=[output, softmax_lse])
167	-	167	+
168	g.update(cpu_update_input=[{"actual_seq_lengths": length_new}])	168	g.update(cpu_update_input=[{"actual_seq_lengths": length_new}])
169	g.replay()	169	g.replay()
170	self.assertEqual(output.cpu(), res_src[0].cpu())	170	self.assertEqual(output.cpu(), res_src[0].cpu())
	@@ -193,7 +193,7 @@ class TestIFAAclgraphUpdate(TestCase):
193	output, softmax_lse = torch_npu.npu_fused_infer_attention_score_v2(	193	output, softmax_lse = torch_npu.npu_fused_infer_attention_score_v2(
194	query, key, value, num_query_heads=32, input_layout="BNSD", softmax_scale=scale, pre_tokens=65535,	194	query, key, value, num_query_heads=32, input_layout="BNSD", softmax_scale=scale, pre_tokens=65535,
195	next_tokens=65535, return_softmax_lse=False, actual_seq_qlen=length)	195	next_tokens=65535, return_softmax_lse=False, actual_seq_qlen=length)
196	-	196	+
197	g.update(cpu_update_input=[{"actual_seq_qlen": length_new}])	197	g.update(cpu_update_input=[{"actual_seq_qlen": length_new}])
198	g.replay()	198	g.replay()
199	self.assertEqual(output.cpu(), res_src[0].cpu())	199	self.assertEqual(output.cpu(), res_src[0].cpu())
	@@ -233,7 +233,7 @@ class TestIFAAclgraphUpdate(TestCase):
233	query, key, value, num_query_heads=32, input_layout="BNSD", softmax_scale=scale, pre_tokens=65535, workspace=workspace,	233	query, key, value, num_query_heads=32, input_layout="BNSD", softmax_scale=scale, pre_tokens=65535, workspace=workspace,
234	next_tokens=65535, return_softmax_lse=False, actual_seq_qlen=length, out=[output, softmax_lse])	234	next_tokens=65535, return_softmax_lse=False, actual_seq_qlen=length, out=[output, softmax_lse])
235	handle = torch.npu.graph_task_group_end(stream)	235	handle = torch.npu.graph_task_group_end(stream)
236	-	236	+
237	with torch.npu.stream(update_stream):	237	with torch.npu.stream(update_stream):
238	torch.npu.graph_task_update_begin(update_stream, handle)	238	torch.npu.graph_task_update_begin(update_stream, handle)
239	torch_npu.npu_fused_infer_attention_score_v2.out(	239	torch_npu.npu_fused_infer_attention_score_v2.out(
	@@ -309,7 +309,7 @@ class TestIFAAclgraphUpdate(TestCase):
309	torch.nn.Dropout(p=0.2),	309	torch.nn.Dropout(p=0.2),
310	torch.nn.Linear(H, D_out),	310	torch.nn.Linear(H, D_out),
311	torch.nn.Dropout(p=0.1)).npu()	311	torch.nn.Dropout(p=0.1)).npu()
312	-	312	+
313	static_input = torch.randn(N, D_in, device='npu')	313	static_input = torch.randn(N, D_in, device='npu')
314	s = torch.npu.Stream()	314	s = torch.npu.Stream()
315	s.wait_stream(torch.npu.current_stream())	315	s.wait_stream(torch.npu.current_stream())
	@@ -365,7 +365,7 @@ class TestIFAAclgraphUpdate(TestCase):
365	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535, workspace=workspace,	365	query, key, value, num_heads=32, input_layout="BNSD", scale=scale, pre_tokens=65535, workspace=workspace,
366	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length, out=[output, softmax_lse])	366	next_tokens=65535, softmax_lse_flag=False, actual_seq_lengths=length, out=[output, softmax_lse])
367	handle = torch.npu.graph_task_group_end(stream)	367	handle = torch.npu.graph_task_group_end(stream)
368	-	368	+
369	with torch.npu.stream(update_stream):	369	with torch.npu.stream(update_stream):
370	torch.npu.graph_task_update_begin(update_stream, handle)	370	torch.npu.graph_task_update_begin(update_stream, handle)
371	torch_npu.npu_fused_infer_attention_score.out(	371	torch_npu.npu_fused_infer_attention_score.out(
	@@ -412,7 +412,7 @@ class TestIFAAclgraphUpdate(TestCase):
412	query, key, value, num_query_heads=32, input_layout="BNSD", softmax_scale=scale, pre_tokens=65535, workspace=workspace,	412	query, key, value, num_query_heads=32, input_layout="BNSD", softmax_scale=scale, pre_tokens=65535, workspace=workspace,
413	next_tokens=65535, return_softmax_lse=False, actual_seq_qlen=length, out=[output, softmax_lse])	413	next_tokens=65535, return_softmax_lse=False, actual_seq_qlen=length, out=[output, softmax_lse])
414	handle = torch.npu.graph_task_group_end(stream)	414	handle = torch.npu.graph_task_group_end(stream)
415	-	415	+
416	with torch.npu.stream(update_stream):	416	with torch.npu.stream(update_stream):
417	torch.npu.graph_task_update_begin(update_stream, handle)	417	torch.npu.graph_task_update_begin(update_stream, handle)
418	torch_npu.npu_fused_infer_attention_score_v2.out(	418	torch_npu.npu_fused_infer_attention_score_v2.out(
	@@ -503,8 +503,8 @@ class TestPAAclgraphUpdate(TestCase):
503		503
504	# 执行注意力计算	504	# 执行注意力计算
505	out = self.ref_masked_attention(	505	out = self.ref_masked_attention(
506	- params_np.query[i:i + 1],	506	+ params_np.query[i:i + 1],
507	- np.stack(keys),	507	+ np.stack(keys),
508	np.stack(values)	508	np.stack(values)
509	)	509	)
510	output[i] = out.reshape(self.num_heads, -1)	510	output[i] = out.reshape(self.num_heads, -1)
	@@ -536,11 +536,11 @@ class TestPAAclgraphUpdate(TestCase):
536		536
537	def atb_paged_attention(self, params):	537	def atb_paged_attention(self, params):
538	torch_npu._npu_paged_attention(	538	torch_npu._npu_paged_attention(
539	- query=params.query,	539	+ query=params.query,
540	- key_cache=params.key_cache,	540	+ key_cache=params.key_cache,
541	value_cache=params.value_cache,	541	value_cache=params.value_cache,
542	num_kv_heads=self.kv_heads,	542	num_kv_heads=self.kv_heads,
543	- num_heads=self.num_heads,	543	+ num_heads=self.num_heads,
544	scale_value=self.scale,	544	scale_value=self.scale,
545	block_table=params.block_table,	545	block_table=params.block_table,
546	context_lens=params.context_lens,	546	context_lens=params.context_lens,

	@@ -18,13 +18,13 @@ class TestDLPack(TestCase):
18	original = torch.randint(-10, 10, (2, 3, 4), dtype=dtype, device=device)	18	original = torch.randint(-10, 10, (2, 3, 4), dtype=dtype, device=device)
19	else:	19	else:
20	original = torch.randn(2, 3, 4, dtype=dtype, device=device)	20	original = torch.randn(2, 3, 4, dtype=dtype, device=device)
21	-	21	+
22	# Convert to dlpack	22	# Convert to dlpack
23	dlpack_tensor = to_dlpack(original)	23	dlpack_tensor = to_dlpack(original)
24	-	24	+
25	# Convert back to torch_npu tensor	25	# Convert back to torch_npu tensor
26	restored = from_dlpack(dlpack_tensor)	26	restored = from_dlpack(dlpack_tensor)
27	-	27	+
28	# Verify the roundtrip	28	# Verify the roundtrip
29	self.assertEqual(original, restored)	29	self.assertEqual(original, restored)
30	self.assertEqual(original.dtype, restored.dtype)	30	self.assertEqual(original.dtype, restored.dtype)
	@@ -42,13 +42,13 @@ class TestDLPack(TestCase):
42	(2, 2, 2, 2), # 4D tensor	42	(2, 2, 2, 2), # 4D tensor
43	(1, 1, 1, 1, 1) # 5D tensor	43	(1, 1, 1, 1, 1) # 5D tensor
44	]	44	]
45	-	45	+
46	for shape in shapes:	46	for shape in shapes:
47	with self.subTest(shape=shape):	47	with self.subTest(shape=shape):
48	original = torch.randn(shape, dtype=dtype, device=device)	48	original = torch.randn(shape, dtype=dtype, device=device)
49	dlpack_tensor = to_dlpack(original)	49	dlpack_tensor = to_dlpack(original)
50	restored = from_dlpack(dlpack_tensor)	50	restored = from_dlpack(dlpack_tensor)
51	-	51	+
52	self.assertEqual(original, restored)	52	self.assertEqual(original, restored)
53	self.assertEqual(original.shape, restored.shape)	53	self.assertEqual(original.shape, restored.shape)
54		54
	@@ -58,20 +58,20 @@ class TestDLPack(TestCase):
58	# Test contiguous tensor	58	# Test contiguous tensor
59	original_contiguous = torch.randn(4, 4, dtype=dtype, device=device)	59	original_contiguous = torch.randn(4, 4, dtype=dtype, device=device)
60	self.assertTrue(original_contiguous.is_contiguous())	60	self.assertTrue(original_contiguous.is_contiguous())
61	-	61	+
62	dlpack_tensor = to_dlpack(original_contiguous)	62	dlpack_tensor = to_dlpack(original_contiguous)
63	restored = from_dlpack(dlpack_tensor)	63	restored = from_dlpack(dlpack_tensor)
64	-	64	+
65	self.assertEqual(original_contiguous, restored)	65	self.assertEqual(original_contiguous, restored)
66	self.assertTrue(restored.is_contiguous())	66	self.assertTrue(restored.is_contiguous())
67	-	67	+
68	# Test non-contiguous tensor (transpose)	68	# Test non-contiguous tensor (transpose)
69	original_non_contiguous = original_contiguous.t()	69	original_non_contiguous = original_contiguous.t()
70	self.assertFalse(original_non_contiguous.is_contiguous())	70	self.assertFalse(original_non_contiguous.is_contiguous())
71	-	71	+
72	dlpack_tensor = to_dlpack(original_non_contiguous)	72	dlpack_tensor = to_dlpack(original_non_contiguous)
73	restored = from_dlpack(dlpack_tensor)	73	restored = from_dlpack(dlpack_tensor)
74	-	74	+
75	self.assertEqual(original_non_contiguous, restored)	75	self.assertEqual(original_non_contiguous, restored)
76	self.assertEqual(original_non_contiguous.stride(), restored.stride())	76	self.assertEqual(original_non_contiguous.stride(), restored.stride())
77		77
	@@ -80,14 +80,14 @@ class TestDLPack(TestCase):
80	"""Test that dlpack shares memory with original tensor"""	80	"""Test that dlpack shares memory with original tensor"""
81	original = torch.randn(3, 3, dtype=dtype, device=device)	81	original = torch.randn(3, 3, dtype=dtype, device=device)
82	original_data_ptr = original.data_ptr()	82	original_data_ptr = original.data_ptr()
83	-	83	+
84	# Convert to dlpack and back	84	# Convert to dlpack and back
85	dlpack_tensor = to_dlpack(original)	85	dlpack_tensor = to_dlpack(original)
86	restored = from_dlpack(dlpack_tensor)	86	restored = from_dlpack(dlpack_tensor)
87	-	87	+
88	# Check if memory is shared (data_ptr should be the same)	88	# Check if memory is shared (data_ptr should be the same)
89	self.assertEqual(original_data_ptr, restored.data_ptr())	89	self.assertEqual(original_data_ptr, restored.data_ptr())
90	-	90	+
91	# Modify original tensor and check if restored tensor is also modified	91	# Modify original tensor and check if restored tensor is also modified
92	original.fill_(42.0)	92	original.fill_(42.0)
93	self.assertEqual(original, restored)	93	self.assertEqual(original, restored)
	@@ -98,20 +98,20 @@ class TestDLPack(TestCase):
98	original = torch.randn(3, 4, dtype=dtype, device=device)	98	original = torch.randn(3, 4, dtype=dtype, device=device)
99	dlpack_tensor = to_dlpack(original)	99	dlpack_tensor = to_dlpack(original)
100	restored = from_dlpack(dlpack_tensor)	100	restored = from_dlpack(dlpack_tensor)
101	-	101	+
102	self.assertEqual(original, restored)	102	self.assertEqual(original, restored)
103	self.assertEqual(original.dtype, restored.dtype)	103	self.assertEqual(original.dtype, restored.dtype)
104	-	104	+
105	@Dtypes(torch.float)	105	@Dtypes(torch.float)
106	def test_dlpack_cpu(self, dtype, device="cpu"):	106	def test_dlpack_cpu(self, dtype, device="cpu"):
107	"""Test that dlpack shares memory with original cpu tensor"""	107	"""Test that dlpack shares memory with original cpu tensor"""
108	original = torch.randn(3, 3, dtype=dtype, device=device)	108	original = torch.randn(3, 3, dtype=dtype, device=device)
109	original_data_ptr = original.data_ptr()	109	original_data_ptr = original.data_ptr()
110	-	110	+
111	# Convert to dlpack and back	111	# Convert to dlpack and back
112	dlpack_tensor = to_dlpack(original)	112	dlpack_tensor = to_dlpack(original)
113	restored = from_dlpack(dlpack_tensor)	113	restored = from_dlpack(dlpack_tensor)
114	-	114	+
115	# Check if memory is shared (data_ptr should be the same)	115	# Check if memory is shared (data_ptr should be the same)
116	self.assertEqual(original_data_ptr, restored.data_ptr())	116	self.assertEqual(original_data_ptr, restored.data_ptr())
117		117

	@@ -95,14 +95,14 @@ class TorchNPUDeviceTestCase(TestCase):
95	os.environ["TORCH_NPU_DEVICE_CAPABILITY"] = "9.0"	95	os.environ["TORCH_NPU_DEVICE_CAPABILITY"] = "9.0"
96	res = torch_npu.npu.get_device_capability()	96	res = torch_npu.npu.get_device_capability()
97	self.assertEqual(res, (9, 0))	97	self.assertEqual(res, (9, 0))
98	-	98	+
99	os.environ["TORCH_NPU_DEVICE_CAPABILITY"] = "8.0"	99	os.environ["TORCH_NPU_DEVICE_CAPABILITY"] = "8.0"
100	res = torch_npu.npu.get_device_capability(0)	100	res = torch_npu.npu.get_device_capability(0)
101	self.assertEqual(res, (8, 0))	101	self.assertEqual(res, (8, 0))
102	device = torch_npu.npu.device("npu")	102	device = torch_npu.npu.device("npu")
103	res = torch_npu.npu.get_device_capability(device)	103	res = torch_npu.npu.get_device_capability(device)
104	self.assertEqual(res, (8, 0))	104	self.assertEqual(res, (8, 0))
105	-	105	+
106	os.environ["TORCH_NPU_DEVICE_CAPABILITY"] = "8.a"	106	os.environ["TORCH_NPU_DEVICE_CAPABILITY"] = "8.a"
107	res = torch_npu.npu.get_device_capability()	107	res = torch_npu.npu.get_device_capability()
108	self.assertEqual(res, None)	108	self.assertEqual(res, None)
	@@ -118,7 +118,7 @@ class TorchNPUDeviceTestCase(TestCase):
118	torch_npu.npu.synchronize()	118	torch_npu.npu.synchronize()
119	after_free_memory, after_total_memory = torch_npu.npu.mem_get_info(0)	119	after_free_memory, after_total_memory = torch_npu.npu.mem_get_info(0)
120	self.assertEqual(before_total_memory, after_total_memory)	120	self.assertEqual(before_total_memory, after_total_memory)
121	-	121	+
122	@unittest.skip("CANN doesn't support now.")	122	@unittest.skip("CANN doesn't support now.")
123	def test_set_device_res_limit(self):	123	def test_set_device_res_limit(self):
124	ans_dict = {'cube_core_num': 12, 'vector_core_num': 24}	124	ans_dict = {'cube_core_num': 12, 'vector_core_num': 24}
	@@ -301,7 +301,7 @@ class TorchNPUApiTestCase(TestCase):
301	end_event.record()	301	end_event.record()
302	res = start_event.elapsed_time(end_event)	302	res = start_event.elapsed_time(end_event)
303	self.assertIsInstance(res, float)	303	self.assertIsInstance(res, float)
304	-	304	+
305	def test_npu_event_recorded_time(self):	305	def test_npu_event_recorded_time(self):
306	event_1 = torch_npu.npu.Event(enable_timing=True)	306	event_1 = torch_npu.npu.Event(enable_timing=True)
307	event_1.record()	307	event_1.record()
	@@ -426,7 +426,7 @@ print(f"{{r1}}, {{r2}}")
426	self.fail("Expected exception not raised for negative timeout value")	426	self.fail("Expected exception not raised for negative timeout value")
427	except Exception as e:	427	except Exception as e:
428	self.assertIn("can't convert negative value to unsigned int", str(e), f"{e}")	428	self.assertIn("can't convert negative value to unsigned int", str(e), f"{e}")
429	-	429	+
430	try:	430	try:
431	torch_npu.npu.set_op_timeout_ms(2**32)	431	torch_npu.npu.set_op_timeout_ms(2**32)
432	self.fail("Expected exception not raised for large timeout value")	432	self.fail("Expected exception not raised for large timeout value")

	@@ -1604,7 +1604,7 @@ class TestBinaryUfuncs(TestCase):
1604	for tensor in tensors:	1604	for tensor in tensors:
1605	self._test_pow(base, tensor)	1605	self._test_pow(base, tensor)
1606		1606
1607	-	1607	+
1608	def test_cuda_tensor_pow_scalar_tensor(self, device):	1608	def test_cuda_tensor_pow_scalar_tensor(self, device):
1609	cuda_tensors = [	1609	cuda_tensors = [
1610	torch.randn((3, 3), device=device),	1610	torch.randn((3, 3), device=device),
	@@ -1618,7 +1618,7 @@ class TestBinaryUfuncs(TestCase):
1618	for base, exp in product(cuda_tensors, scalar_tensors):	1618	for base, exp in product(cuda_tensors, scalar_tensors):
1619	self._test_pow(base, exp)	1619	self._test_pow(base, exp)
1620		1620
1621	-	1621	+
1622	def test_cpu_tensor_pow_cuda_scalar_tensor(self, device):	1622	def test_cpu_tensor_pow_cuda_scalar_tensor(self, device):
1623	cuda_tensors = [	1623	cuda_tensors = [
1624	torch.tensor(5.0, device="privatuse1"),	1624	torch.tensor(5.0, device="privatuse1"),
	@@ -1633,7 +1633,7 @@ class TestBinaryUfuncs(TestCase):
1633	base = torch.tensor(3.0, device="cpu")	1633	base = torch.tensor(3.0, device="cpu")
1634	self._test_pow(base, exp)	1634	self._test_pow(base, exp)
1635		1635
1636	-	1636	+
1637	@dtypes(torch.complex64, torch.complex128)	1637	@dtypes(torch.complex64, torch.complex128)
1638	def test_pow_cuda_complex_extremal_failing(self, device, dtype):	1638	def test_pow_cuda_complex_extremal_failing(self, device, dtype):
1639	t = torch.tensor(complex(-1.0, float("inf")), dtype=dtype, device=device)	1639	t = torch.tensor(complex(-1.0, float("inf")), dtype=dtype, device=device)
	@@ -1642,7 +1642,7 @@ class TestBinaryUfuncs(TestCase):
1642	cpu_out = t.cpu().pow(2)	1642	cpu_out = t.cpu().pow(2)
1643	self.assertEqual(cpu_out, cuda_out)	1643	self.assertEqual(cpu_out, cuda_out)
1644		1644
1645	-	1645	+
1646	@skipIfTorchDynamo()	1646	@skipIfTorchDynamo()
1647	@dtypes(*all_types_and_complex_and(torch.half))	1647	@dtypes(*all_types_and_complex_and(torch.half))
1648	def test_complex_scalar_pow_tensor(self, device, dtype):	1648	def test_complex_scalar_pow_tensor(self, device, dtype):
	@@ -1669,7 +1669,7 @@ class TestBinaryUfuncs(TestCase):
1669	self._test_pow(base, first_exp)	1669	self._test_pow(base, first_exp)
1670	self._test_pow(base, second_exp)	1670	self._test_pow(base, second_exp)
1671		1671
1672	-	1672	+
1673	@skipMeta	1673	@skipMeta
1674	def test_pow_scalar_type_promotion(self, device):	1674	def test_pow_scalar_type_promotion(self, device):
1675	# Test against a scalar and non-scalar input	1675	# Test against a scalar and non-scalar input
	@@ -1837,7 +1837,7 @@ class TestBinaryUfuncs(TestCase):
1837	_scalar_helper(lambda a, b: math.floor(a / b), operator.floordiv)	1837	_scalar_helper(lambda a, b: math.floor(a / b), operator.floordiv)
1838	_scalar_helper(lambda a, b: math.floor(a / b), torch.floor_divide)	1838	_scalar_helper(lambda a, b: math.floor(a / b), torch.floor_divide)
1839		1839
1840	-	1840	+
1841	@skipIfTorchDynamo("Not a suitable test for TorchDynamo")	1841	@skipIfTorchDynamo("Not a suitable test for TorchDynamo")
1842	def test_div_and_floordiv_script_vs_python(self, device):	1842	def test_div_and_floordiv_script_vs_python(self, device):
1843	# Creates jitted functions of two tensors	1843	# Creates jitted functions of two tensors
	@@ -1908,7 +1908,7 @@ class TestBinaryUfuncs(TestCase):
1908	# See issue gh-52387	1908	# See issue gh-52387
1909	self.assertEqual(5 // a, scripted_rfloordiv_scalar(a_t))	1909	self.assertEqual(5 // a, scripted_rfloordiv_scalar(a_t))
1910		1910
1911	-	1911	+
1912	@skipIfTorchDynamo("Not a suitable test for TorchDynamo")	1912	@skipIfTorchDynamo("Not a suitable test for TorchDynamo")
1913	def test_idiv_and_ifloordiv_vs_python(self, device):	1913	def test_idiv_and_ifloordiv_vs_python(self, device):
1914	def _wrapped_idiv_tensor(a, b):	1914	def _wrapped_idiv_tensor(a, b):
	@@ -2239,7 +2239,7 @@ class TestBinaryUfuncs(TestCase):
2239	torch.ones(1, device=device, dtype=dtypes[0]),	2239	torch.ones(1, device=device, dtype=dtypes[0]),
2240	)	2240	)
2241		2241
2242	-	2242	+
2243	def test_maximum_minimum_cross_device(self, device):	2243	def test_maximum_minimum_cross_device(self, device):
2244	a = torch.tensor((1, 2, -1))	2244	a = torch.tensor((1, 2, -1))
2245	b = torch.tensor((3, 0, 4), device=device)	2245	b = torch.tensor((3, 0, 4), device=device)
	@@ -2822,7 +2822,7 @@ class TestBinaryUfuncs(TestCase):
2822	expected = np.hypot(input[0].cpu().numpy(), input[1].cpu().numpy())	2822	expected = np.hypot(input[0].cpu().numpy(), input[1].cpu().numpy())
2823	self.assertEqual(actual, expected, exact_dtype=False)	2823	self.assertEqual(actual, expected, exact_dtype=False)
2824		2824
2825	-	2825	+
2826	@dtypes(torch.uint8, torch.int8, torch.int16, torch.int32, torch.int64)	2826	@dtypes(torch.uint8, torch.int8, torch.int16, torch.int32, torch.int64)
2827	def test_gcd(self, device, dtype):	2827	def test_gcd(self, device, dtype):
2828	# Tests gcd(0, 0), gcd(0, a) cases	2828	# Tests gcd(0, 0), gcd(0, a) cases
	@@ -2847,7 +2847,7 @@ class TestBinaryUfuncs(TestCase):
2847	expected = np.gcd(a.cpu().numpy(), b.cpu().numpy())	2847	expected = np.gcd(a.cpu().numpy(), b.cpu().numpy())
2848	self.assertEqual(actual, expected)	2848	self.assertEqual(actual, expected)
2849		2849
2850	-	2850	+
2851	@dtypes(torch.int16, torch.int32, torch.int64)	2851	@dtypes(torch.int16, torch.int32, torch.int64)
2852	def test_lcm(self, device, dtype):	2852	def test_lcm(self, device, dtype):
2853	# Tests lcm(0, 0), lcm(0, a) cases	2853	# Tests lcm(0, 0), lcm(0, a) cases
	@@ -2864,7 +2864,7 @@ class TestBinaryUfuncs(TestCase):
2864	expected = np.lcm(a.cpu().numpy(), b.cpu().numpy())	2864	expected = np.lcm(a.cpu().numpy(), b.cpu().numpy())
2865	self.assertEqual(actual, expected, exact_dtype=False)	2865	self.assertEqual(actual, expected, exact_dtype=False)
2866		2866
2867	-	2867	+
2868	@dtypes(torch.float32, torch.float64)	2868	@dtypes(torch.float32, torch.float64)
2869	def test_nextafter(self, device, dtype):	2869	def test_nextafter(self, device, dtype):
2870	# Test special cases	2870	# Test special cases
	@@ -2888,7 +2888,7 @@ class TestBinaryUfuncs(TestCase):
2888	expected = np.nextafter(a.cpu().numpy(), b.cpu().numpy())	2888	expected = np.nextafter(a.cpu().numpy(), b.cpu().numpy())
2889	self.assertEqual(actual, expected, atol=0, rtol=0)	2889	self.assertEqual(actual, expected, atol=0, rtol=0)
2890		2890
2891	-	2891	+
2892	@dtypes(torch.bfloat16)	2892	@dtypes(torch.bfloat16)
2893	def test_nextafter_bfloat16(self, device, dtype):	2893	def test_nextafter_bfloat16(self, device, dtype):
2894	nan = float("nan")	2894	nan = float("nan")
	@@ -3108,7 +3108,7 @@ class TestBinaryUfuncs(TestCase):
3108	) # all casts to complex128 are safe	3108	) # all casts to complex128 are safe
3109	compare_with_numpy_bin_op(torch_op, numpy_op, a, b, out=out)	3109	compare_with_numpy_bin_op(torch_op, numpy_op, a, b, out=out)
3110		3110
3111	-	3111	+
3112	@dtypes(torch.int8, torch.int16, torch.int32, torch.int64)	3112	@dtypes(torch.int8, torch.int16, torch.int32, torch.int64)
3113	def test_signed_shift(self, device, dtype):	3113	def test_signed_shift(self, device, dtype):
3114	"Ensure that signed integer bit shifting works as expected."	3114	"Ensure that signed integer bit shifting works as expected."
	@@ -3124,7 +3124,7 @@ class TestBinaryUfuncs(TestCase):
3124	self.assertEqual(a >> 1, expected_r)	3124	self.assertEqual(a >> 1, expected_r)
3125	self.compare_with_numpy(lambda x: x >> 1, lambda x: np.right_shift(x, 1), a)	3125	self.compare_with_numpy(lambda x: x >> 1, lambda x: np.right_shift(x, 1), a)
3126		3126
3127	-	3127	+
3128	@dtypes(*get_all_int_dtypes())	3128	@dtypes(*get_all_int_dtypes())
3129	def test_shift_limits(self, device, dtype):	3129	def test_shift_limits(self, device, dtype):
3130	"Ensure that integer bit shifting works as expected with out-of-limits shift values."	3130	"Ensure that integer bit shifting works as expected with out-of-limits shift values."
	@@ -3167,7 +3167,7 @@ class TestBinaryUfuncs(TestCase):
3167	exact_dtype=exact_dtype, msg=f">> {shift}"	3167	exact_dtype=exact_dtype, msg=f">> {shift}"
3168	)	3168	)
3169		3169
3170	-	3170	+
3171	@dtypes(	3171	@dtypes(
3172	*list(	3172	*list(
3173	product(	3173	product(
	@@ -3236,7 +3236,7 @@ class TestBinaryUfuncs(TestCase):
3236	):	3236	):
3237	input.heaviside_(values)	3237	input.heaviside_(values)
3238		3238
3239	-	3239	+
3240	def test_heaviside_cross_device(self, device):	3240	def test_heaviside_cross_device(self, device):
3241	x = torch.tensor([-9, 5, 0, 6, -2, 2], device=device)	3241	x = torch.tensor([-9, 5, 0, 6, -2, 2], device=device)
3242	y = torch.tensor(0)	3242	y = torch.tensor(0)
	@@ -3405,7 +3405,7 @@ class TestBinaryUfuncs(TestCase):
3405	expected = start + weight * (end - start)	3405	expected = start + weight * (end - start)
3406	self.assertEqual(expected, actual)	3406	self.assertEqual(expected, actual)
3407		3407
3408	-	3408	+
3409	@dtypes(torch.half, torch.bfloat16)	3409	@dtypes(torch.half, torch.bfloat16)
3410	def test_lerp_lowp(self, device, dtype):	3410	def test_lerp_lowp(self, device, dtype):
3411	xvals = (0.0, -30000.0)	3411	xvals = (0.0, -30000.0)
	@@ -3646,7 +3646,7 @@ class TestBinaryUfuncs(TestCase):
3646	lambda: torch.add(m1, m1, out=m2),	3646	lambda: torch.add(m1, m1, out=m2),
3647	)	3647	)
3648		3648
3649	-	3649	+
3650	def test_addsub_half_tensor(self, device):	3650	def test_addsub_half_tensor(self, device):
3651	x = torch.tensor([60000.0], dtype=torch.half, device=device)	3651	x = torch.tensor([60000.0], dtype=torch.half, device=device)
3652	for op, y, alpha in (	3652	for op, y, alpha in (
	@@ -3879,7 +3879,7 @@ class TestBinaryUfuncs(TestCase):
3879	def test_cumulative_trapezoid(self, device):	3879	def test_cumulative_trapezoid(self, device):
3880		3880
3881	import scipy.integrate	3881	import scipy.integrate
3882	-	3882	+
3883	if hasattr(scipy.integrate, "cumulative_trapezoid"):	3883	if hasattr(scipy.integrate, "cumulative_trapezoid"):
3884	_scipy_cumulative_trapezoid = scipy.integrate.cumulative_trapezoid	3884	_scipy_cumulative_trapezoid = scipy.integrate.cumulative_trapezoid
3885	else: # Older version of SciPy uses a different name	3885	else: # Older version of SciPy uses a different name
	@@ -4352,7 +4352,7 @@ class TestBinaryUfuncs(TestCase):
4352	x = make_tensor((2, 3, 4), dtype=x_dtype, device=device)	4352	x = make_tensor((2, 3, 4), dtype=x_dtype, device=device)
4353	test_helper(x, q)	4353	test_helper(x, q)
4354		4354
4355	-	4355	+
4356	@dtypes(	4356	@dtypes(
4357	torch.chalf,	4357	torch.chalf,
4358	)	4358	)

	@@ -50,22 +50,22 @@ class TestMultiprocessingAPIs(TestCase):
50	mp.set_start_method(method, force=True)	50	mp.set_start_method(method, force=True)
51	current_method = mp.get_start_method()	51	current_method = mp.get_start_method()
52	self.assertEqual(current_method, method)	52	self.assertEqual(current_method, method)
53	-	53	+
54	# Verify that the child process uses the correct start method	54	# Verify that the child process uses the correct start method
55	queue = mp.SimpleQueue()	55	queue = mp.SimpleQueue()
56	process = mp.Process(target=_worker, args=(queue,))	56	process = mp.Process(target=_worker, args=(queue,))
57	process.start()	57	process.start()
58	process.join()	58	process.join()
59	-	59	+
60	# Get the start method from the child process	60	# Get the start method from the child process
61	self.assertFalse(queue.empty(), "Queue should contain the start method")	61	self.assertFalse(queue.empty(), "Queue should contain the start method")
62	child_method = queue.get()	62	child_method = queue.get()
63	- self.assertEqual(child_method, method,	63	+ self.assertEqual(child_method, method,
64	f"Child process should use {method} start method")	64	f"Child process should use {method} start method")
65	-	65	+
66	# Verify that the main process context has not changed	66	# Verify that the main process context has not changed
67	self.assertEqual(mp.get_start_method(), method)	67	self.assertEqual(mp.get_start_method(), method)
68	-	68	+
69		69
70	def test_value(self):	70	def test_value(self):
71	"""Test Value API"""	71	"""Test Value API"""
	@@ -184,49 +184,49 @@ class TestMultiprocessingAPIs(TestCase):
184	"""Test torch.multiprocessing.reductions.init_reductions and reduce_tensor APIs"""	184	"""Test torch.multiprocessing.reductions.init_reductions and reduce_tensor APIs"""
185	# Test init_reductions - verify it doesn't raise any exception	185	# Test init_reductions - verify it doesn't raise any exception
186	mp.reductions.init_reductions()	186	mp.reductions.init_reductions()
187	-	187	+
188	# Test reduce_tensor directly for CPU tensor	188	# Test reduce_tensor directly for CPU tensor
189	# Create a simple CPU tensor	189	# Create a simple CPU tensor
190	tensor = torch.tensor([1, 2, 3, 4])	190	tensor = torch.tensor([1, 2, 3, 4])
191	reduced = mp.reductions.reduce_tensor(tensor)	191	reduced = mp.reductions.reduce_tensor(tensor)
192	-	192	+
193	# Verify the reduced form is a tuple with expected structure	193	# Verify the reduced form is a tuple with expected structure
194	self.assertIsInstance(reduced, tuple)	194	self.assertIsInstance(reduced, tuple)
195	self.assertEqual(len(reduced), 2)	195	self.assertEqual(len(reduced), 2)
196	-	196	+
197	# Try to reconstruct the tensor	197	# Try to reconstruct the tensor
198	constructor, args = reduced	198	constructor, args = reduced
199	reconstructed = constructor(*args)	199	reconstructed = constructor(*args)
200	-	200	+
201	# Verify reconstruction worked	201	# Verify reconstruction worked
202	self.assertTrue(torch.equal(tensor, reconstructed))	202	self.assertTrue(torch.equal(tensor, reconstructed))
203	self.assertEqual(tensor.device, reconstructed.device)	203	self.assertEqual(tensor.device, reconstructed.device)
204	self.assertEqual(tensor.dtype, reconstructed.dtype)	204	self.assertEqual(tensor.dtype, reconstructed.dtype)
205	-	205	+
206	# Test with NPU tensor if available	206	# Test with NPU tensor if available
207	if torch.npu.is_available():	207	if torch.npu.is_available():
208	# Create a simple NPU tensor	208	# Create a simple NPU tensor
209	npu_tensor = torch.tensor([1, 2, 3, 4], device='npu:0')	209	npu_tensor = torch.tensor([1, 2, 3, 4], device='npu:0')
210	-	210	+
211	# Test reduce_tensor for NPU tensor	211	# Test reduce_tensor for NPU tensor
212	reduced_npu = mp.reductions.reduce_tensor(npu_tensor)	212	reduced_npu = mp.reductions.reduce_tensor(npu_tensor)
213	self.assertIsInstance(reduced_npu, tuple)	213	self.assertIsInstance(reduced_npu, tuple)
214	self.assertEqual(len(reduced_npu), 2)	214	self.assertEqual(len(reduced_npu), 2)
215	-	215	+
216	# Verify reconstruction for NPU tensor	216	# Verify reconstruction for NPU tensor
217	constructor_npu, args_npu = reduced_npu	217	constructor_npu, args_npu = reduced_npu
218	reconstructed_npu = constructor_npu(*args_npu)	218	reconstructed_npu = constructor_npu(*args_npu)
219	self.assertTrue(torch.equal(npu_tensor.cpu(), reconstructed_npu.cpu()))	219	self.assertTrue(torch.equal(npu_tensor.cpu(), reconstructed_npu.cpu()))
220	-	220	+
221	def test_reductions_invalid_input(self):	221	def test_reductions_invalid_input(self):
222	"""Test reduction APIs with invalid inputs"""	222	"""Test reduction APIs with invalid inputs"""
223	# Test reduce_tensor with invalid input	223	# Test reduce_tensor with invalid input
224	with self.assertRaises(Exception):	224	with self.assertRaises(Exception):
225	mp.reductions.reduce_tensor(None)	225	mp.reductions.reduce_tensor(None)
226	-	226	+
227	with self.assertRaises(Exception):	227	with self.assertRaises(Exception):
228	mp.reductions.reduce_tensor("not a tensor")	228	mp.reductions.reduce_tensor("not a tensor")
229	-	229	+
230	# Test init_reductions multiple times (should be safe)	230	# Test init_reductions multiple times (should be safe)
231	mp.reductions.init_reductions()	231	mp.reductions.init_reductions()
232	mp.reductions.init_reductions()	232	mp.reductions.init_reductions()

	@@ -84,13 +84,13 @@ class TestNestedTensor(TestCase):
84	self.assertEqual(len(nt_as_list), len(nt_list))	84	self.assertEqual(len(nt_as_list), len(nt_list))
85	self.assertEqual(nt_as_list[0], nt_list[0])	85	self.assertEqual(nt_as_list[0], nt_list[0])
86	self.assertEqual(nt_as_list[1], nt_list[1])	86	self.assertEqual(nt_as_list[1], nt_list[1])
87	-	87	+
88	def test_unbind_and_asnested_int64(self):	88	def test_unbind_and_asnested_int64(self):
89	a = torch.tensor([[1, 2, 3], [4, 5, 6]])	89	a = torch.tensor([[1, 2, 3], [4, 5, 6]])
90	b = torch.tensor([[7, 8], [10, 11]])	90	b = torch.tensor([[7, 8], [10, 11]])
91	self._test_unbind_case(a, b)	91	self._test_unbind_case(a, b)
92	self._test_asnested_case(a, b)	92	self._test_asnested_case(a, b)
93	-	93	+
94	def test_unbind_and_asnested_float32(self):	94	def test_unbind_and_asnested_float32(self):
95	a = torch.tensor([[1, 2, 3], [4, 5, 6]], dtype=torch.float32)	95	a = torch.tensor([[1, 2, 3], [4, 5, 6]], dtype=torch.float32)
96	b = torch.tensor([[7, 8], [10, 11]], dtype=torch.float32)	96	b = torch.tensor([[7, 8], [10, 11]], dtype=torch.float32)
	@@ -101,7 +101,7 @@ class TestNestedTensor(TestCase):
101	a = torch.tensor([[], []])	101	a = torch.tensor([[], []])
102	b = torch.tensor([[], [], []])	102	b = torch.tensor([[], [], []])
103	self._test_unbind_case(a, b)	103	self._test_unbind_case(a, b)
104	- self._test_asnested_case(a, b)	104	+ self._test_asnested_case(a, b)
105		105
106	def test_default_options_nested_tensor(self):	106	def test_default_options_nested_tensor(self):
107	default_nested_tensor = torch.nested.nested_tensor([], device="npu:0")	107	default_nested_tensor = torch.nested.nested_tensor([], device="npu:0")
	@@ -111,12 +111,12 @@ class TestNestedTensor(TestCase):
111	self.assertEqual(default_nested_tensor.layout, default_tensor.layout)	111	self.assertEqual(default_nested_tensor.layout, default_tensor.layout)
112	self.assertEqual(default_nested_tensor.dim(), default_tensor.dim())	112	self.assertEqual(default_nested_tensor.dim(), default_tensor.dim())
113	self.assertEqual(default_nested_tensor.requires_grad, default_tensor.requires_grad)	113	self.assertEqual(default_nested_tensor.requires_grad, default_tensor.requires_grad)
114	-	114	+
115	def test_nested_tensor_errsize(self):	115	def test_nested_tensor_errsize(self):
116	nt = torch.nested.nested_tensor([torch.tensor([[1, 2, 3], [4, 5, 6]]).npu(), torch.tensor([[7, 8], [10, 11], [12, 13]]).npu()])	116	nt = torch.nested.nested_tensor([torch.tensor([[1, 2, 3], [4, 5, 6]]).npu(), torch.tensor([[7, 8], [10, 11], [12, 13]]).npu()])
117	self.assertEqual(nt.size(0), 2)	117	self.assertEqual(nt.size(0), 2)
118	self.assertRaisesRegex(RuntimeError,	118	self.assertRaisesRegex(RuntimeError,
119	- "Given dimension 1 is irregular and does not have a size",	119	+ "Given dimension 1 is irregular and does not have a size",
120	lambda: nt.size(1),	120	lambda: nt.size(1),
121	)	121	)
122		122

	@@ -993,7 +993,7 @@ class TestReductions(TestCase):
993	# Check whether the returned values are the mode	993	# Check whether the returned values are the mode
994	self.assertTrue((values == v).all().item())	994	self.assertTrue((values == v).all().item())
995		995
996	-	996	+
997	@dtypes(*all_types_and(torch.half, torch.bfloat16))	997	@dtypes(*all_types_and(torch.half, torch.bfloat16))
998	def test_mode_large(self, device, dtype):	998	def test_mode_large(self, device, dtype):
999	# i should be less than (d - 2) / 2	999	# i should be less than (d - 2) / 2
	@@ -1063,7 +1063,7 @@ class TestReductions(TestCase):
1063	test_for_dtypes(torch.int32, torch.int32, torch.float32, indices_err)	1063	test_for_dtypes(torch.int32, torch.int32, torch.float32, indices_err)
1064	test_for_dtypes(torch.float32, torch.float32, torch.float64, indices_err)	1064	test_for_dtypes(torch.float32, torch.float32, torch.float64, indices_err)
1065		1065
1066	-	1066	+
1067	def test_mode_wrong_device(self, device):	1067	def test_mode_wrong_device(self, device):
1068	# CPU Input Tensor	1068	# CPU Input Tensor
1069	x = torch.ones(2)	1069	x = torch.ones(2)
	@@ -1212,7 +1212,7 @@ class TestReductions(TestCase):
1212	@dtypes(torch.float, torch.double)	1212	@dtypes(torch.float, torch.double)
1213	def test_amax(self, device, dtype):	1213	def test_amax(self, device, dtype):
1214	self._test_minmax_helper(torch.amax, np.amax, device, dtype)	1214	self._test_minmax_helper(torch.amax, np.amax, device, dtype)
1215	-	1215	+
1216	@dtypes(torch.float, torch.double)	1216	@dtypes(torch.float, torch.double)
1217	@dtypesIfPRIVATEUSE1(torch.half, torch.float, torch.bfloat16)	1217	@dtypesIfPRIVATEUSE1(torch.half, torch.float, torch.bfloat16)
1218	def test_aminmax(self, device, dtype):	1218	def test_aminmax(self, device, dtype):
	@@ -1463,7 +1463,7 @@ class TestReductions(TestCase):
1463	torch.sum(x, (2, 1), out=res2)	1463	torch.sum(x, (2, 1), out=res2)
1464	self.assertEqual(res1, res2)	1464	self.assertEqual(res1, res2)
1465		1465
1466	-	1466	+
1467	@dtypes(torch.float16, torch.float32)	1467	@dtypes(torch.float16, torch.float32)
1468	def test_prod_gpu(self, device, dtype):	1468	def test_prod_gpu(self, device, dtype):
1469	x = torch.tensor([2, 3, 6, 9, 8], dtype=dtype, device=device)	1469	x = torch.tensor([2, 3, 6, 9, 8], dtype=dtype, device=device)
	@@ -1729,7 +1729,7 @@ class TestReductions(TestCase):
1729	# So we must skip this as well.	1729	# So we must skip this as well.
1730	if dtype == torch.uint8:	1730	if dtype == torch.uint8:
1731	exact_dtype = False	1731	exact_dtype = False
1732	-	1732	+
1733	# Investigate why the output is not close to numpy.	1733	# Investigate why the output is not close to numpy.
1734	atol, rtol = self._get_relaxed_tolerances_for(dtype)	1734	atol, rtol = self._get_relaxed_tolerances_for(dtype)
1735		1735
	@@ -1743,7 +1743,7 @@ class TestReductions(TestCase):
1743	self._test_sum_reduction_vs_numpy(torch.sum, np.sum, device, dtype, with_extremal=True)	1743	self._test_sum_reduction_vs_numpy(torch.sum, np.sum, device, dtype, with_extremal=True)
1744	self._test_sum_reduction_vs_numpy(torch.sum, np.sum, device, dtype, with_keepdim=True)	1744	self._test_sum_reduction_vs_numpy(torch.sum, np.sum, device, dtype, with_keepdim=True)
1745		1745
1746	-	1746	+
1747	@dtypes(*set(all_types_and(torch.half)) - {torch.uint8})	1747	@dtypes(*set(all_types_and(torch.half)) - {torch.uint8})
1748	def test_nansum_vs_numpy(self, device, dtype):	1748	def test_nansum_vs_numpy(self, device, dtype):
1749	self._test_sum_reduction_vs_numpy(torch.nansum, np.nansum, device, dtype)	1749	self._test_sum_reduction_vs_numpy(torch.nansum, np.nansum, device, dtype)
	@@ -1977,7 +1977,7 @@ class TestReductions(TestCase):
1977	op(x, dim=dim)	1977	op(x, dim=dim)
1978		1978
1979	# update this test to comapre against NumPy	1979	# update this test to comapre against NumPy
1980	-	1980	+
1981	def test_var(self, device):	1981	def test_var(self, device):
1982	cpu_tensor = torch.randn(2, 3, 3)	1982	cpu_tensor = torch.randn(2, 3, 3)
1983	device_tensor = cpu_tensor.to(device)	1983	device_tensor = cpu_tensor.to(device)
	@@ -1992,7 +1992,7 @@ class TestReductions(TestCase):
1992	device_tensor = cpu_tensor.to(device)	1992	device_tensor = cpu_tensor.to(device)
1993	self.assertEqual(device_tensor.var(), cpu_tensor.var())	1993	self.assertEqual(device_tensor.var(), cpu_tensor.var())
1994		1994
1995	- # update this test to compare against NumPy	1995	+ # update this test to compare against NumPy
1996	def test_var_large_input(self, device):	1996	def test_var_large_input(self, device):
1997	# Large, not-nice input	1997	# Large, not-nice input
1998	cpu_tensor = torch.randn(2 * 32 * 1024 + 1, 2, 67)	1998	cpu_tensor = torch.randn(2 * 32 * 1024 + 1, 2, 67)
	@@ -2001,7 +2001,7 @@ class TestReductions(TestCase):
2001	self.assertEqual(cpu_tensor.var(2), device_tensor.var(2))	2001	self.assertEqual(cpu_tensor.var(2), device_tensor.var(2))
2002		2002
2003	# update this to compare against NumPy instead of CPU	2003	# update this to compare against NumPy instead of CPU
2004	-	2004	+
2005	@dtypes(torch.double)	2005	@dtypes(torch.double)
2006	def test_sum_noncontig(self, device, dtype):	2006	def test_sum_noncontig(self, device, dtype):
2007	x = torch.randn(1, 75, 57, 20, dtype=dtype, device=device).permute(0, 3, 1, 2)	2007	x = torch.randn(1, 75, 57, 20, dtype=dtype, device=device).permute(0, 3, 1, 2)
	@@ -2046,7 +2046,7 @@ class TestReductions(TestCase):
2046	torch.sum(x, dim=[0], dtype=torch.float32, out=y)	2046	torch.sum(x, dim=[0], dtype=torch.float32, out=y)
2047		2047
2048	# Assert for illegal dtype would not be raised on XLA	2048	# Assert for illegal dtype would not be raised on XLA
2049	-	2049	+
2050	def test_minmax_illegal_dtype(self, device):	2050	def test_minmax_illegal_dtype(self, device):
2051	x = torch.randn(5, 5, dtype=torch.float32, device=device)	2051	x = torch.randn(5, 5, dtype=torch.float32, device=device)
2052	valid_values = torch.empty(5, dtype=torch.float32, device=device)	2052	valid_values = torch.empty(5, dtype=torch.float32, device=device)
	@@ -2266,7 +2266,7 @@ class TestReductions(TestCase):
2266	expected = fn(y, 1, keepdim=False)	2266	expected = fn(y, 1, keepdim=False)
2267	self.assertEqual(x[:, 1], expected, msg=f'{fn_name} with out= kwarg')	2267	self.assertEqual(x[:, 1], expected, msg=f'{fn_name} with out= kwarg')
2268		2268
2269	-	2269	+
2270	@largeTensorTest('10GB')	2270	@largeTensorTest('10GB')
2271	def test_reduction_split(self, device):	2271	def test_reduction_split(self, device):
2272	# Test reduction when there is a 32bit-indexing split	2272	# Test reduction when there is a 32bit-indexing split
	@@ -2275,7 +2275,7 @@ class TestReductions(TestCase):
2275	expect = input_[0] + input_[1] + input_[2] + input_[3] + input_[4]	2275	expect = input_[0] + input_[1] + input_[2] + input_[3] + input_[4]
2276	self.assertEqual(result, expect)	2276	self.assertEqual(result, expect)
2277		2277
2278	-	2278	+
2279	@dtypes(torch.half, torch.float, torch.double, torch.bfloat16)	2279	@dtypes(torch.half, torch.float, torch.double, torch.bfloat16)
2280	def test_reduction_vectorize_along_input_corner(self, device, dtype):	2280	def test_reduction_vectorize_along_input_corner(self, device, dtype):
2281	# 1D case: sum	2281	# 1D case: sum
	@@ -2373,7 +2373,7 @@ class TestReductions(TestCase):
2373	self.assertEqual(xs1[j].item(), size[1] - i)	2373	self.assertEqual(xs1[j].item(), size[1] - i)
2374	self.assertEqual(xs2[j].item(), size[1] - i)	2374	self.assertEqual(xs2[j].item(), size[1] - i)
2375		2375
2376	-	2376	+
2377	@dtypes(torch.half, torch.float, torch.double, torch.bfloat16)	2377	@dtypes(torch.half, torch.float, torch.double, torch.bfloat16)
2378	def test_reduction_vectorize_along_output(self, device, dtype):	2378	def test_reduction_vectorize_along_output(self, device, dtype):
2379	def run_test(input_):	2379	def run_test(input_):
	@@ -2397,7 +2397,7 @@ class TestReductions(TestCase):
2397	run_test(torch.zeros(64, 61, dtype=dtype, device=device))	2397	run_test(torch.zeros(64, 61, dtype=dtype, device=device))
2398	run_test(torch.zeros(64, 1, dtype=dtype, device=device))	2398	run_test(torch.zeros(64, 1, dtype=dtype, device=device))
2399		2399
2400	-	2400	+
2401	def test_argminmax_large_axis(self, device):	2401	def test_argminmax_large_axis(self, device):
2402	# Regression test for gh-32863	2402	# Regression test for gh-32863
2403	x = torch.zeros(2**31, device=device, dtype=torch.int8)	2403	x = torch.zeros(2**31, device=device, dtype=torch.int8)
	@@ -2541,7 +2541,7 @@ class TestReductions(TestCase):
2541	self.assertEqual(a[:, ::2, :].nanmedian(-1)[0], torch.tensor([[0, 4], [6, 10]], device=device))	2541	self.assertEqual(a[:, ::2, :].nanmedian(-1)[0], torch.tensor([[0, 4], [6, 10]], device=device))
2542		2542
2543		2543
2544	-	2544	+
2545	@dtypes(torch.float, torch.double)	2545	@dtypes(torch.float, torch.double)
2546	def test_quantile(self, device, dtype):	2546	def test_quantile(self, device, dtype):
2547	# Generate some random test cases	2547	# Generate some random test cases
	@@ -3571,7 +3571,7 @@ as the input tensor excluding its innermost dimension'):
3571		3571
3572	self.assertEqual(actual, expected, msg, exact_dtype=exact_dtype)	3572	self.assertEqual(actual, expected, msg, exact_dtype=exact_dtype)
3573		3573
3574	-	3574	+
3575	@largeTensorTest("8GB")	3575	@largeTensorTest("8GB")
3576	@dtypes(torch.half, torch.chalf, torch.bfloat16)	3576	@dtypes(torch.half, torch.chalf, torch.bfloat16)
3577	def test_reductions_large_half_tensors(self, device, dtype):	3577	def test_reductions_large_half_tensors(self, device, dtype):

	@@ -70,7 +70,7 @@ class TestTensorToPreserveFormat(TestCase):
70	"""H2D: transpose 2D non-contiguous -> preserve_format preserves stride"""	70	"""H2D: transpose 2D non-contiguous -> preserve_format preserves stride"""
71	cpu_t = torch.randn(4, 6).t()	71	cpu_t = torch.randn(4, 6).t()
72	self.assertFalse(cpu_t.is_contiguous())	72	self.assertFalse(cpu_t.is_contiguous())
73	-	73	+
74		74
75	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)	75	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)
76	self._verify_preserve_format_result(cpu_t, npu_t, "H2D-transpose_2d")	76	self._verify_preserve_format_result(cpu_t, npu_t, "H2D-transpose_2d")
	@@ -96,7 +96,7 @@ class TestTensorToPreserveFormat(TestCase):
96	"""H2D: slice dim0 non-contiguous -> preserve_format preserves stride"""	96	"""H2D: slice dim0 non-contiguous -> preserve_format preserves stride"""
97	cpu_t = torch.randn(8, 5)[::2]	97	cpu_t = torch.randn(8, 5)[::2]
98	self.assertFalse(cpu_t.is_contiguous())	98	self.assertFalse(cpu_t.is_contiguous())
99	-	99	+
100		100
101	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)	101	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)
102	expec_strides = self._get_dense_strides(cpu_t)	102	expec_strides = self._get_dense_strides(cpu_t)
	@@ -107,7 +107,7 @@ class TestTensorToPreserveFormat(TestCase):
107	"""H2D: slice dim1 non-contiguous -> preserve_format preserves stride"""	107	"""H2D: slice dim1 non-contiguous -> preserve_format preserves stride"""
108	cpu_t = torch.randn(4, 8)[:, ::2]	108	cpu_t = torch.randn(4, 8)[:, ::2]
109	self.assertFalse(cpu_t.is_contiguous())	109	self.assertFalse(cpu_t.is_contiguous())
110	-	110	+
111		111
112	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)	112	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)
113	self._verify_preserve_format_result(cpu_t, npu_t, "H2D-slice_dim1")	113	self._verify_preserve_format_result(cpu_t, npu_t, "H2D-slice_dim1")
	@@ -119,14 +119,14 @@ class TestTensorToPreserveFormat(TestCase):
119	"""H2D: narrow non-contiguous -> preserve_format preserves stride"""	119	"""H2D: narrow non-contiguous -> preserve_format preserves stride"""
120	cpu_t = torch.randn(6, 8).narrow(1, 1, 5)	120	cpu_t = torch.randn(6, 8).narrow(1, 1, 5)
121	self.assertFalse(cpu_t.is_contiguous())	121	self.assertFalse(cpu_t.is_contiguous())
122	-	122	+
123		123
124	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)	124	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)
125	self._verify_preserve_format_result(cpu_t, npu_t, "H2D-narrow")	125	self._verify_preserve_format_result(cpu_t, npu_t, "H2D-narrow")
126	expec_strides = self._get_dense_strides(cpu_t)	126	expec_strides = self._get_dense_strides(cpu_t)
127	self.assertEqual(expec_strides, npu_t.stride(),	127	self.assertEqual(expec_strides, npu_t.stride(),
128	"H2D-narrow: stride not preserved")	128	"H2D-narrow: stride not preserved")
129	-	129	+
130	def test_h2d_select(self):	130	def test_h2d_select(self):
131	"""H2D: select (reduces dim) non-contiguous -> preserve_format preserves stride"""	131	"""H2D: select (reduces dim) non-contiguous -> preserve_format preserves stride"""
132	cpu_t = torch.randn(3, 5, 4).select(1, 2)	132	cpu_t = torch.randn(3, 5, 4).select(1, 2)
	@@ -164,7 +164,7 @@ class TestTensorToPreserveFormat(TestCase):
164	"""H2D: expand 3D (stride=0) -> preserve_format falls back to suggest_memory_format"""	164	"""H2D: expand 3D (stride=0) -> preserve_format falls back to suggest_memory_format"""
165	cpu_t = torch.randn(2, 1, 4).expand(2, 3, 4)	165	cpu_t = torch.randn(2, 1, 4).expand(2, 3, 4)
166	self.assertFalse(cpu_t.is_contiguous())	166	self.assertFalse(cpu_t.is_contiguous())
167	-	167	+
168		168
169	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)	169	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)
170	self._verify_preserve_format_result(cpu_t, npu_t, "H2D-expand_3d")	170	self._verify_preserve_format_result(cpu_t, npu_t, "H2D-expand_3d")
	@@ -174,7 +174,7 @@ class TestTensorToPreserveFormat(TestCase):
174	def test_h2d_as_strided_overlap(self):	174	def test_h2d_as_strided_overlap(self):
175	"""H2D: as_strided with overlap -> preserve_format falls back to suggest_memory_format"""	175	"""H2D: as_strided with overlap -> preserve_format falls back to suggest_memory_format"""
176	cpu_t = torch.randn(12).as_strided((3, 3), (4, 1))	176	cpu_t = torch.randn(12).as_strided((3, 3), (4, 1))
177	-	177	+
178		178
179	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)	179	npu_t = cpu_t.to("npu", memory_format=torch.preserve_format)
180	self._verify_preserve_format_result(cpu_t, npu_t, "H2D-as_strided_overlap")	180	self._verify_preserve_format_result(cpu_t, npu_t, "H2D-as_strided_overlap")

	@@ -158,10 +158,10 @@ def check_size_alltoall(alltoall_cases: list[dict[str, Any]]) -> tuple[bool, int
158	for e in alltoall_cases:	158	for e in alltoall_cases:
159	input_sizes = e.get("input_sizes", [])	159	input_sizes = e.get("input_sizes", [])
160	output_sizes = e.get("output_sizes", [])	160	output_sizes = e.get("output_sizes", [])
161	-	161	+
162	if input_sizes and len(input_sizes) > 0:	162	if input_sizes and len(input_sizes) > 0:
163	input_numel += math.prod(input_sizes[0])	163	input_numel += math.prod(input_sizes[0])
164	-	164	+
165	if output_sizes and len(output_sizes) > 0:	165	if output_sizes and len(output_sizes) > 0:
166	output_numel += math.prod(output_sizes[0])	166	output_numel += math.prod(output_sizes[0])
167	return input_numel != output_numel, input_numel, output_numel	167	return input_numel != output_numel, input_numel, output_numel
	@@ -330,7 +330,7 @@ def just_print_entries(
330	entry = all_entries[rank].pop(0)	330	entry = all_entries[rank].pop(0)
331	process_group = entry.get("process_group", [None, None])	331	process_group = entry.get("process_group", [None, None])
332	pg_name = _pg_guids.get((process_group[0], rank))	332	pg_name = _pg_guids.get((process_group[0], rank))
333	-	333	+
334	if (	334	if (
335	args.pg_filters is None	335	args.pg_filters is None
336	or process_group[1] in args.pg_filters	336	or process_group[1] in args.pg_filters
	@@ -356,7 +356,7 @@ def check_no_missing_dump_files(entries: dict[int, Any], memberships: list[Membe
356	dumps_ranks = {int(key) for key in entries.keys()}	356	dumps_ranks = {int(key) for key in entries.keys()}
357	except (ValueError, TypeError) as e:	357	except (ValueError, TypeError) as e:
358	raise ValueError(f"Cannot extract rank from entries keys. Invalid key value encountered: {e}") from e	358	raise ValueError(f"Cannot extract rank from entries keys. Invalid key value encountered: {e}") from e
359	-	359	+
360	missing_ranks = all_ranks - dumps_ranks	360	missing_ranks = all_ranks - dumps_ranks
361	if missing_ranks:	361	if missing_ranks:
362	raise ValueError(	362	raise ValueError(

	@@ -104,18 +104,18 @@ static void _launch(void* func, void* tiling_func, int64_t tiling_size, void* ar
104	// only 1D parallelization is supported for NPU	104	// only 1D parallelization is supported for NPU
105	// Pointer type becomes flattend 1-D Memref tuple: base_ptr, data_ptr, offset, shape, stride	105	// Pointer type becomes flattend 1-D Memref tuple: base_ptr, data_ptr, offset, shape, stride
106	// base_ptr offset shape and stride are not used, arbitrarily set for now	106	// base_ptr offset shape and stride are not used, arbitrarily set for now
107	-	107	+
108	if (tiling_size == 0) {{	108	if (tiling_size == 0) {{
109	auto launch_call = [func, tiling_func, tiling_size, arg_tiling_host, arg_tiling_device, gridX, stream, {', '.join(f"arg{i}" + ("" if "torch." in ty else f", arg_allocate{i}, offset{i}" +(', ' if ranks[i] > 0 else '') + ', '.join(f"sizes{i}_{rank}" for rank in range(ranks[i])) + (', ' if ranks[i] > 0 else '') + ', '.join(f"strides{i}_{rank}" for rank in range(ranks[i]))) for i, ty in signature.items())}]() {{	109	auto launch_call = [func, tiling_func, tiling_size, arg_tiling_host, arg_tiling_device, gridX, stream, {', '.join(f"arg{i}" + ("" if "torch." in ty else f", arg_allocate{i}, offset{i}" +(', ' if ranks[i] > 0 else '') + ', '.join(f"sizes{i}_{rank}" for rank in range(ranks[i])) + (', ' if ranks[i] > 0 else '') + ', '.join(f"strides{i}_{rank}" for rank in range(ranks[i]))) for i, ty in signature.items())}]() {{
110	struct __attribute__((packed)) {{	110	struct __attribute__((packed)) {{
111	-	111	+
112	{' '.join(f'{_ty_to_cpp(ty)} arg{i} __attribute__((aligned({4 if ty[0] != "" and ty[-2:] != "64" else 8}))); ' + ('' if "torch." in ty else f'{_ty_to_cpp(ty)} arg_allocate{i} __attribute__((aligned({4 if ty[0] != "" and ty[-2:] != "64" else 8}))); {_ty_to_cpp(ty)} offset{i} __attribute__((aligned(8))); ' + ' '.join(f'{_ty_to_cpp(ty)} sizes{i}_{rank} __attribute__((aligned(8)));' for rank in range(ranks[i])) + ' ' + ' '.join(f'{_ty_to_cpp(ty)} strides{i}_{rank} __attribute__((aligned(8)));' for rank in range(ranks[i]))) for i, ty in signature.items())}	112	{' '.join(f'{_ty_to_cpp(ty)} arg{i} __attribute__((aligned({4 if ty[0] != "" and ty[-2:] != "64" else 8}))); ' + ('' if "torch." in ty else f'{_ty_to_cpp(ty)} arg_allocate{i} __attribute__((aligned({4 if ty[0] != "" and ty[-2:] != "64" else 8}))); {_ty_to_cpp(ty)} offset{i} __attribute__((aligned(8))); ' + ' '.join(f'{_ty_to_cpp(ty)} sizes{i}_{rank} __attribute__((aligned(8)));' for rank in range(ranks[i])) + ' ' + ' '.join(f'{_ty_to_cpp(ty)} strides{i}_{rank} __attribute__((aligned(8)));' for rank in range(ranks[i]))) for i, ty in signature.items())}
113		113
114	}} args = {{	114	}} args = {{
115	{', '.join(f"static_cast<{_ty_to_cpp(ty)}>(arg{i})" + ("" if "torch." in ty else f", static_cast<{_ty_to_cpp(ty)}>(arg_allocate{i}), static_cast<{_ty_to_cpp(ty)}>(offset{i})"+ (', ' if ranks[i] > 0 else '') + ', '.join(f"static_cast<{_ty_to_cpp(ty)}>(sizes{i}_{rank})" for rank in range(ranks[i])) + (', ' if ranks[i] > 0 else '') + ', '.join(f"static_cast<{_ty_to_cpp(ty)}>(strides{i}_{rank})" for rank in range(ranks[i]))) for i, ty in signature.items())}	115	{', '.join(f"static_cast<{_ty_to_cpp(ty)}>(arg{i})" + ("" if "torch." in ty else f", static_cast<{_ty_to_cpp(ty)}>(arg_allocate{i}), static_cast<{_ty_to_cpp(ty)}>(offset{i})"+ (', ' if ranks[i] > 0 else '') + ', '.join(f"static_cast<{_ty_to_cpp(ty)}>(sizes{i}_{rank})" for rank in range(ranks[i])) + (', ' if ranks[i] > 0 else '') + ', '.join(f"static_cast<{_ty_to_cpp(ty)}>(strides{i}_{rank})" for rank in range(ranks[i]))) for i, ty in signature.items())}
116		116
117	}};	117	}};
118	-	118	+
119	rtError_t ret = common_launch_dyn(const_cast<char>("{kernel_name}"), func, tiling_func, tiling_size, arg_tiling_host, arg_tiling_device, gridX, static_cast<void >(&args), sizeof(args), stream);	119	rtError_t ret = common_launch_dyn(const_cast<char>("{kernel_name}"), func, tiling_func, tiling_size, arg_tiling_host, arg_tiling_device, gridX, static_cast<void >(&args), sizeof(args), stream);
120	return ret;	120	return ret;
121	}};	121	}};
	@@ -128,7 +128,7 @@ static void _launch(void* func, void* tiling_func, int64_t tiling_size, void* ar
128	void* strides_tiling = (void*)1;	128	void* strides_tiling = (void*)1;
129	auto launch_call = [func, tiling_func, tiling_size, arg_tiling_host, arg_tiling_device, gridX, stream, {', '.join(f"arg{i}" + ("" if "torch." in ty else f", arg_allocate{i}, offset{i}" + (', ' if ranks[i] > 0 else '') + ', '.join(f"sizes{i}_{rank}" for rank in range(ranks[i])) + (', ' if ranks[i] > 0 else '') + ', '.join(f"strides{i}_{rank}" for rank in range(ranks[i]))) for i, ty in signature.items())}, key_tiling, offset_tiling, sizes_tiling, strides_tiling]() {{	129	auto launch_call = [func, tiling_func, tiling_size, arg_tiling_host, arg_tiling_device, gridX, stream, {', '.join(f"arg{i}" + ("" if "torch." in ty else f", arg_allocate{i}, offset{i}" + (', ' if ranks[i] > 0 else '') + ', '.join(f"sizes{i}_{rank}" for rank in range(ranks[i])) + (', ' if ranks[i] > 0 else '') + ', '.join(f"strides{i}_{rank}" for rank in range(ranks[i]))) for i, ty in signature.items())}, key_tiling, offset_tiling, sizes_tiling, strides_tiling]() {{
130	struct __attribute__((packed)) {{	130	struct __attribute__((packed)) {{
131	-	131	+
132	{' '.join(f'{_ty_to_cpp(ty)} arg{i} __attribute__((aligned({4 if ty[0] != "" and ty[-2:] != "64" else 8}))); ' + ('' if "torch." in ty else f'{_ty_to_cpp(ty)} arg_allocate{i} __attribute__((aligned({4 if ty[0] != "" and ty[-2:] != "64" else 8}))); {_ty_to_cpp(ty)} offset{i} __attribute__((aligned(8))); ' + ' '.join(f'{_ty_to_cpp(ty)} sizes{i}_{rank} __attribute__((aligned(8)));' for rank in range(ranks[i])) + ' ' + ' '.join(f'{_ty_to_cpp(ty)} strides{i}_{rank} __attribute__((aligned(8)));' for rank in range(ranks[i]))) for i, ty in signature.items())}	132	{' '.join(f'{_ty_to_cpp(ty)} arg{i} __attribute__((aligned({4 if ty[0] != "" and ty[-2:] != "64" else 8}))); ' + ('' if "torch." in ty else f'{_ty_to_cpp(ty)} arg_allocate{i} __attribute__((aligned({4 if ty[0] != "" and ty[-2:] != "64" else 8}))); {_ty_to_cpp(ty)} offset{i} __attribute__((aligned(8))); ' + ' '.join(f'{_ty_to_cpp(ty)} sizes{i}_{rank} __attribute__((aligned(8)));' for rank in range(ranks[i])) + ' ' + ' '.join(f'{_ty_to_cpp(ty)} strides{i}_{rank} __attribute__((aligned(8)));' for rank in range(ranks[i]))) for i, ty in signature.items())}
133		133
134	void* key_tiling __attribute__((aligned(8)));	134	void* key_tiling __attribute__((aligned(8)));
	@@ -143,7 +143,7 @@ static void _launch(void* func, void* tiling_func, int64_t tiling_size, void* ar
143		143
144	(void)(&key_tiling), arg_tiling_host, arg_tiling_device, static_cast<void>(offset_tiling), static_cast<void>(sizes_tiling), static_cast<void>(strides_tiling)	144	(void)(&key_tiling), arg_tiling_host, arg_tiling_device, static_cast<void>(offset_tiling), static_cast<void>(sizes_tiling), static_cast<void>(strides_tiling)
145	}};	145	}};
146	-	146	+
147	rtError_t ret = common_launch_dyn(const_cast<char>("{kernel_name}"), func, tiling_func, tiling_size, arg_tiling_host, arg_tiling_device, gridX, static_cast<void >(&args), sizeof(args), stream);	147	rtError_t ret = common_launch_dyn(const_cast<char>("{kernel_name}"), func, tiling_func, tiling_size, arg_tiling_host, arg_tiling_device, gridX, static_cast<void >(&args), sizeof(args), stream);
148	return ret;	148	return ret;
149	}};	149	}};

	@@ -59,7 +59,7 @@ id_iter = count()
59		59
60		60
61	class NpuTritonKernel(TritonKernel):	61	class NpuTritonKernel(TritonKernel):
62	- def __init__(self,	62	+ def __init__(self,
63	tiling: Dict[str, sympy.Expr],	63	tiling: Dict[str, sympy.Expr],
64	min_elem_per_thread=0,	64	min_elem_per_thread=0,
65	optimize_mask=True,	65	optimize_mask=True,
	@@ -77,7 +77,7 @@ class NpuTritonKernel(TritonKernel):
77	@staticmethod	77	@staticmethod
78	def inductor_meta_common():	78	def inductor_meta_common():
79	return {}	79	return {}
80	-	80	+
81	def call_kernel(self, call_args, name: str):	81	def call_kernel(self, call_args, name: str):
82	wrapper = V.graph.wrapper_code	82	wrapper = V.graph.wrapper_code
83	for call_arg in call_args:	83	for call_arg in call_args:
	@@ -170,14 +170,14 @@ def create_fx_from_snodes_by_traced_graph(snodes: List[scheduler.SchedulerNode],
170		170
171	def runnable_gm(*args):	171	def runnable_gm(*args):
172	return torch.fx.Interpreter(gm).run(*args)	172	return torch.fx.Interpreter(gm).run(*args)
173	- with V.graph.fake_mode:	173	+ with V.graph.fake_mode:
174	gm = make_fx(runnable_gm)(*inputs)	174	gm = make_fx(runnable_gm)(*inputs)
175	view_to_reshape(gm)	175	view_to_reshape(gm)
176	- non_contiguous_indices["outputs"] = [i + num_inputs	176	+ non_contiguous_indices["outputs"] = [i + num_inputs
177	for i, call_output in enumerate(call_outputs)	177	for i, call_output in enumerate(call_outputs)
178	if not V.graph.try_get_buffer(call_output).layout.is_contiguous()]	178	if not V.graph.try_get_buffer(call_output).layout.is_contiguous()]
179	- return (gm, call_args, {"num_outputs": num_outputs,	179	+ return (gm, call_args, {"num_outputs": num_outputs,
180	- "non_contiguous_indices": non_contiguous_indices,	180	+ "non_contiguous_indices": non_contiguous_indices,
181	"mutated_indices": mutated_indices, })	181	"mutated_indices": mutated_indices, })
182		182
183		183
	@@ -192,10 +192,10 @@ class NpuMetaKernel(Kernel):
192	self._gm = gm	192	self._gm = gm
193	self._gm_with_prim_cast = self.build_gm_with_prim_cast(gm)	193	self._gm_with_prim_cast = self.build_gm_with_prim_cast(gm)
194	self._is_dynamic = is_fx_dynamic(self._gm)	194	self._is_dynamic = is_fx_dynamic(self._gm)
195	-	195	+
196	if anir_config.online_acc_comp:	196	if anir_config.online_acc_comp:
197	modify_gm_for_acc_comp(self._gm)	197	modify_gm_for_acc_comp(self._gm)
198	-	198	+
199	self._snodes = snodes	199	self._snodes = snodes
200	self._call_args = call_args	200	self._call_args = call_args
201	self.non_contiguous_indices = non_contiguous_indices	201	self.non_contiguous_indices = non_contiguous_indices
	@@ -216,7 +216,7 @@ class NpuMetaKernel(Kernel):
216	V.graph.device_ops.import_get_raw_stream_as("get_raw_stream")	216	V.graph.device_ops.import_get_raw_stream_as("get_raw_stream")
217	)	217	)
218	)	218	)
219	-	219	+
220	def build_gm_with_prim_cast(self, gm):	220	def build_gm_with_prim_cast(self, gm):
221	return npu_cast_to_prim_cast(gm)	221	return npu_cast_to_prim_cast(gm)
222		222
	@@ -254,12 +254,12 @@ class NpuMetaKernel(Kernel):
254	def call_kernel(self, name: str, node=None):	254	def call_kernel(self, name: str, node=None):
255	wrapper = V.graph.wrapper_code	255	wrapper = V.graph.wrapper_code
256	call_args = self.get_call_args()	256	call_args = self.get_call_args()
257	-	257	+
258	for call_arg in call_args:	258	for call_arg in call_args:
259	if call_arg.startswith('_uwu'):	259	if call_arg.startswith('_uwu'):
260	expression = map_strings_to_operators(call_arg)	260	expression = map_strings_to_operators(call_arg)
261	wrapper.writeline(f'{call_arg} = {expression}')	261	wrapper.writeline(f'{call_arg} = {expression}')
262	-	262	+
263	if len(call_args) > 0:	263	if len(call_args) > 0:
264	wrapper.generate_kernel_call(name, call_args)	264	wrapper.generate_kernel_call(name, call_args)
265		265
	@@ -296,7 +296,7 @@ class NpuMetaScheduling(SIMDScheduling):
296	def define_kernel(self, src_code, mlir_kernel, traced_graph, mode=None):	296	def define_kernel(self, src_code, mlir_kernel, traced_graph, mode=None):
297	if mode is None:	297	if mode is None:
298	mode = anir_config._get_compile_mode()	298	mode = anir_config._get_compile_mode()
299	-	299	+
300	wrapper = V.graph.wrapper_code	300	wrapper = V.graph.wrapper_code
301		301
302	kernel_key = (src_code, tuple(mlir_kernel.non_contiguous_indices))	302	kernel_key = (src_code, tuple(mlir_kernel.non_contiguous_indices))
	@@ -338,9 +338,9 @@ class NpuMetaScheduling(SIMDScheduling):
338	kernel_meta.update(extra_kernel_meta)	338	kernel_meta.update(extra_kernel_meta)
339		339
340	wrapper.src_to_kernel[kernel_key] = kernel_name	340	wrapper.src_to_kernel[kernel_key] = kernel_name
341	-	341	+
342	subs_name = kernel_name if config.triton.unique_kernel_names else f"{self._get_kernel_prefix()}_"	342	subs_name = kernel_name if config.triton.unique_kernel_names else f"{self._get_kernel_prefix()}_"
343	-	343	+
344	compile_wrapper = IndentedBuffer()	344	compile_wrapper = IndentedBuffer()
345	metadata_comment = ""	345	metadata_comment = ""
346		346
	@@ -390,7 +390,7 @@ class NpuMetaScheduling(SIMDScheduling):
390		390
391	def _handle_auto_fallback_mode(self, compile_wrapper, src_code, name, subs_name, meta, wrapper, metadata_comment, mlir_kernel=None):	391	def _handle_auto_fallback_mode(self, compile_wrapper, src_code, name, subs_name, meta, wrapper, metadata_comment, mlir_kernel=None):
392	_basename, _, kernel_path = get_path(code_hash(src_code.strip()), "py")	392	_basename, _, kernel_path = get_path(code_hash(src_code.strip()), "py")
393	-	393	+
394	compile_wrapper.writeline(f"async_compile.{self._get_compile_api()}({subs_name!r}, '''")	394	compile_wrapper.writeline(f"async_compile.{self._get_compile_api()}({subs_name!r}, '''")
395	compile_wrapper.splice(src_code, strip=True)	395	compile_wrapper.splice(src_code, strip=True)
396	compile_wrapper.writeline(f"''', kernel_meta={meta})")	396	compile_wrapper.writeline(f"''', kernel_meta={meta})")
	@@ -399,7 +399,7 @@ class NpuMetaScheduling(SIMDScheduling):
399		399
400	origins, detailed_origins = get_kernel_metadata(mlir_kernel._snodes, wrapper)	400	origins, detailed_origins = get_kernel_metadata(mlir_kernel._snodes, wrapper)
401	metadata_comment += "\n" + origins + "\n" + detailed_origins	401	metadata_comment += "\n" + origins + "\n" + detailed_origins
402	-	402	+
403	wrapper.define_kernel(name, compile_wrapper.getvalue(), metadata_comment)	403	wrapper.define_kernel(name, compile_wrapper.getvalue(), metadata_comment)
404		404
405	if metrics.is_metric_table_enabled("kernel_metadata"):	405	if metrics.is_metric_table_enabled("kernel_metadata"):
	@@ -417,15 +417,15 @@ class NpuMetaScheduling(SIMDScheduling):
417	return "auto_fallback"	417	return "auto_fallback"
418		418
419	def _dump_fx_graph_for_fallback(self, mlir_kernel, device, graph_hash, kernel_name, compile_code):	419	def _dump_fx_graph_for_fallback(self, mlir_kernel, device, graph_hash, kernel_name, compile_code):
420	-	420	+
421	cache_root = os.getenv("TORCHINDUCTOR_CACHE_DIR")	421	cache_root = os.getenv("TORCHINDUCTOR_CACHE_DIR")
422	dump_path = os.path.join(	422	dump_path = os.path.join(
423	- cache_root,	423	+ cache_root,
424	- anir_config.traced_graph_cache or "traced_graph_cache",	424	+ anir_config.traced_graph_cache or "traced_graph_cache",
425	- str(device.index),	425	+ str(device.index),
426	graph_hash	426	graph_hash
427	)	427	)
428	-	428	+
429	if not os.path.exists(dump_path):	429	if not os.path.exists(dump_path):
430	os.makedirs(dump_path, exist_ok=True)	430	os.makedirs(dump_path, exist_ok=True)
431	to_folder(mlir_kernel._gm, dump_path, graph_hash=graph_hash, module_name=graph_hash)	431	to_folder(mlir_kernel._gm, dump_path, graph_hash=graph_hash, module_name=graph_hash)
	@@ -433,11 +433,11 @@ class NpuMetaScheduling(SIMDScheduling):
433	if anir_config.fx_subgraph_dump_path is not None:	433	if anir_config.fx_subgraph_dump_path is not None:
434	subgraph_dump_path = os.path.join(anir_config.fx_subgraph_dump_path, str(device.index), kernel_name)	434	subgraph_dump_path = os.path.join(anir_config.fx_subgraph_dump_path, str(device.index), kernel_name)
435	os.makedirs(subgraph_dump_path, exist_ok=True)	435	os.makedirs(subgraph_dump_path, exist_ok=True)
436	-	436	+
437	num_args = len(mlir_kernel._gm.code.split('forward(', )[1].split(')')[0].split(', ')) - 1	437	num_args = len(mlir_kernel._gm.code.split('forward(', )[1].split(')')[0].split(', ')) - 1
438	fx_graph_code = get_fx_graph_code(mlir_kernel._gm.code, num_args, runnable=False, kernel_code=compile_code, kernel_name=kernel_name)	438	fx_graph_code = get_fx_graph_code(mlir_kernel._gm.code, num_args, runnable=False, kernel_code=compile_code, kernel_name=kernel_name)
439	runnable_fx_graph_code = get_fx_graph_code(mlir_kernel._gm.code, num_args, runnable=True, kernel_code=compile_code, kernel_name=kernel_name)	439	runnable_fx_graph_code = get_fx_graph_code(mlir_kernel._gm.code, num_args, runnable=True, kernel_code=compile_code, kernel_name=kernel_name)
440	-	440	+
441	with open(os.path.join(subgraph_dump_path, f'{kernel_name}.py'), 'w') as f:	441	with open(os.path.join(subgraph_dump_path, f'{kernel_name}.py'), 'w') as f:
442	f.write(fx_graph_code)	442	f.write(fx_graph_code)
443	with open(os.path.join(subgraph_dump_path, f'runnable_{kernel_name}.py'), 'w') as f:	443	with open(os.path.join(subgraph_dump_path, f'runnable_{kernel_name}.py'), 'w') as f:
	@@ -519,10 +519,10 @@ class NpuMetaScheduling(SIMDScheduling):
519		519
520	def codegen_node(self, node: Union[scheduler.SchedulerNode, object]):	520	def codegen_node(self, node: Union[scheduler.SchedulerNode, object]):
521	nodes: List[scheduler.SchedulerNode] = node.get_nodes()	521	nodes: List[scheduler.SchedulerNode] = node.get_nodes()
522	-	522	+
523	_, (numel, rnumel) = max(nodes, key=lambda x: int(x.is_reduction())).group	523	_, (numel, rnumel) = max(nodes, key=lambda x: int(x.is_reduction())).group
524		524
525	node_schedule = self.generate_node_schedule(nodes, numel, rnumel)	525	node_schedule = self.generate_node_schedule(nodes, numel, rnumel)
526	kernel_features = SIMDKernelFeatures(node_schedule, numel, rnumel)	526	kernel_features = SIMDKernelFeatures(node_schedule, numel, rnumel)
527	-	527	+
528	return self.codegen_node_schedule(kernel_features, nodes)	528	return self.codegen_node_schedule(kernel_features, nodes)

	@@ -27,7 +27,7 @@ class IndexAnalysis:
27	for key, coeff in self.index.as_coefficients_dict().items()	27	for key, coeff in self.index.as_coefficients_dict().items()
28	if not isinstance(key, sympy.Integer)	28	if not isinstance(key, sympy.Integer)
29	]	29	]
30	- # sort by stride	30	+ # sort by stride
31	self.var_stride.sort(key=lambda x: x[1])	31	self.var_stride.sort(key=lambda x: x[1])
32	# only contains tiing axis var	32	# only contains tiing axis var
33	self.var_list = tuple([x[0] for x in self.var_stride if x[0] in self.tiling_axis])	33	self.var_list = tuple([x[0] for x in self.var_stride if x[0] in self.tiling_axis])
	@@ -289,7 +289,7 @@ class ReductionAnalysis:
289	j = len(tiling_axis) - 1	289	j = len(tiling_axis) - 1
290	# remove all low_dims from tiling_axis	290	# remove all low_dims from tiling_axis
291	# all axis before ahead of j are high-orders	291	# all axis before ahead of j are high-orders
292	- # then following is reduced dim	292	+ # then following is reduced dim
293	ranges = [x for x in reduction.ranges if x > 1]	293	ranges = [x for x in reduction.ranges if x > 1]
294	for i in reversed(low_dims):	294	for i in reversed(low_dims):
295	len_axis = tiling_axis[j].length	295	len_axis = tiling_axis[j].length

	@@ -16,7 +16,7 @@ from ..config import num_vector_core, log
16	class SplitTiling:	16	class SplitTiling:
17	def __init__(self, kernel: TritonKernel):	17	def __init__(self, kernel: TritonKernel):
18	self.kernel = kernel	18	self.kernel = kernel
19	- self.indexing = [] # load and store indexing among all scheduler nodes	19	+ self.indexing = [] # load and store indexing among all scheduler nodes
20	kernel.sorted_axis = [x for x in kernel.range_tree_nodes.values()]	20	kernel.sorted_axis = [x for x in kernel.range_tree_nodes.values()]
21	kernel.sorted_axis.sort(reverse=True, key=self.key)	21	kernel.sorted_axis.sort(reverse=True, key=self.key)
22	for i, dim in enumerate(kernel.sorted_axis):	22	for i, dim in enumerate(kernel.sorted_axis):
	@@ -114,7 +114,7 @@ class SplitTiling:
114	for i, x in enumerate(self.kernel.split_axis):	114	for i, x in enumerate(self.kernel.split_axis):
115	x.split_order = i	115	x.split_order = i
116		116
117	- # Tiling 原则1：load / store 中索引表达式的中的低维轴都要成为tiling 轴.	117	+ # Tiling 原则1：load / store 中索引表达式的中的低维轴都要成为tiling 轴.
118	# Tiling 原则2：对于规约算子，规约轴要成为tiling轴。	118	# Tiling 原则2：对于规约算子，规约轴要成为tiling轴。
119	# Tiling 原则3: 多维规约，只有规约轴可以被选择为tiling轴	119	# Tiling 原则3: 多维规约，只有规约轴可以被选择为tiling轴
120	# Tiling 原则4: tiling轴要覆盖 total numel 的 80%	120	# Tiling 原则4: tiling轴要覆盖 total numel 的 80%

	@@ -426,7 +426,7 @@ def npu_fusion_attention_grad_v3_strategy(query, key, value, dy, head_num, input
426	None, None, None, None, None, # others	426	None, None, None, None, None, # others
427	None if seed is None else Replicate(), # seed	427	None if seed is None else Replicate(), # seed
428	None if offset is None else Replicate(), # offset	428	None if offset is None else Replicate(), # offset
429	- None,	429	+ None,
430	None if actual_seq_qlen is None else Replicate(), # actual_seq_qlen	430	None if actual_seq_qlen is None else Replicate(), # actual_seq_qlen
431	None if actual_seq_kvlen is None else Replicate(), # actual_seq_kvlen	431	None if actual_seq_kvlen is None else Replicate(), # actual_seq_kvlen
432	None, None, None, None, # others	432	None, None, None, None, # others
	@@ -477,7 +477,7 @@ def npu_fusion_attention_grad_v3_strategy(query, key, value, dy, head_num, input
477	None, None, None, None, None, # others	477	None, None, None, None, None, # others
478	None if seed is None else Replicate(), # seed	478	None if seed is None else Replicate(), # seed
479	None if offset is None else Replicate(), # offset	479	None if offset is None else Replicate(), # offset
480	- None,	480	+ None,
481	None if actual_seq_qlen is None else Replicate(), # actual_seq_qlen	481	None if actual_seq_qlen is None else Replicate(), # actual_seq_qlen
482	None if actual_seq_kvlen is None else Replicate(), # actual_seq_kvlen	482	None if actual_seq_kvlen is None else Replicate(), # actual_seq_kvlen
483	None, None, None, None, # others	483	None, None, None, None, # others
	@@ -520,7 +520,7 @@ def npu_fusion_attention_grad_v3_strategy(query, key, value, dy, head_num, input
520	None, None, None, None, None, # others	520	None, None, None, None, None, # others
521	None if seed is None else Replicate(), # seed	521	None if seed is None else Replicate(), # seed
522	None if offset is None else Replicate(), # offset	522	None if offset is None else Replicate(), # offset
523	- None,	523	+ None,
524	None if actual_seq_qlen is None else Replicate(), # actual_seq_qlen	524	None if actual_seq_qlen is None else Replicate(), # actual_seq_qlen
525	None if actual_seq_kvlen is None else Replicate(), # actual_seq_kvlen	525	None if actual_seq_kvlen is None else Replicate(), # actual_seq_kvlen
526	None, None, None, None, # others	526	None, None, None, None, # others

	@@ -176,12 +176,12 @@ class NpuFusedAdamP(NpuFusedOptimizerBase):
176	def _maybe_init_combined_states(self):	176	def _maybe_init_combined_states(self):
177	if self.is_states_combined:	177	if self.is_states_combined:
178	return	178	return
179	-	179	+
180	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]	180	self.combined_param_states_indexed_by_group = len(self.param_groups) * [None]
181		181
182	for i, _ in enumerate(self.param_groups):	182	for i, _ in enumerate(self.param_groups):
183	self._combine_group_param_states(i)	183	self._combine_group_param_states(i)
184	-	184	+
185	if not all(value is None for value in self.combined_param_states_indexed_by_group):	185	if not all(value is None for value in self.combined_param_states_indexed_by_group):
186	self.is_states_combined = True	186	self.is_states_combined = True
187		187
	@@ -260,7 +260,7 @@ class NpuFusedAdamP(NpuFusedOptimizerBase):
260	def step(self, closure=None):	260	def step(self, closure=None):
261	if not self.is_params_grads_combined:	261	if not self.is_params_grads_combined:
262	self._maybe_init_combined_params_and_grads()	262	self._maybe_init_combined_params_and_grads()
263	-	263	+
264	if not self.is_states_combined:	264	if not self.is_states_combined:
265	self._maybe_init_combined_states()	265	self._maybe_init_combined_states()
266		266

	@@ -77,7 +77,7 @@ class MemoryDbParser(BaseParser):
77	pta_ge_record_list[MemoryRecordTableRow.STREAM_PTR.value] = cur_record[MemoryRecordTableRow.STREAM_PTR.value] if cur_record[MemoryRecordTableRow.STREAM_PTR.value] \	77	pta_ge_record_list[MemoryRecordTableRow.STREAM_PTR.value] = cur_record[MemoryRecordTableRow.STREAM_PTR.value] if cur_record[MemoryRecordTableRow.STREAM_PTR.value] \
78	else last_record_data[MemoryRecordTableRow.STREAM_PTR.value]	78	else last_record_data[MemoryRecordTableRow.STREAM_PTR.value]
79	return [cur_record, pta_ge_record_list]	79	return [cur_record, pta_ge_record_list]
80	-	80	+
81	def run(self, deps_data: dict):	81	def run(self, deps_data: dict):
82	self.logger.info("MemoryDbParser start.")	82	self.logger.info("MemoryDbParser start.")
83	try:	83	try:
	@@ -94,7 +94,7 @@ class MemoryDbParser(BaseParser):
94	return Constant.FAIL, None	94	return Constant.FAIL, None
95	self.logger.info("MemoryDbParser finish.")	95	self.logger.info("MemoryDbParser finish.")
96	return Constant.SUCCESS, None	96	return Constant.SUCCESS, None
97	-	97	+
98	def init_db_connect(self):	98	def init_db_connect(self):
99	if not TorchDb().create_connect_db():	99	if not TorchDb().create_connect_db():
100	raise RuntimeError(f"Failed to connect to db file: {TorchDb().get_db_path()}")	100	raise RuntimeError(f"Failed to connect to db file: {TorchDb().get_db_path()}")
	@@ -197,7 +197,7 @@ class MemoryDbParser(BaseParser):
197	memory_bean.total_allocated_for_db, memory_bean.total_reserved_for_db,	197	memory_bean.total_allocated_for_db, memory_bean.total_reserved_for_db,
198	memory_bean.total_active_for_db, memory_bean.stream_ptr,	198	memory_bean.total_active_for_db, memory_bean.stream_ptr,
199	self.device_index if self.device_index != -1 else memory_bean.device_index])	199	self.device_index if self.device_index != -1 else memory_bean.device_index])
200	-	200	+
201	def get_pta_ge_record_list(self):	201	def get_pta_ge_record_list(self):
202	"""	202	"""
203	ge records are to be sorted firstly and pta records are already sorted,	203	ge records are to be sorted firstly and pta records are already sorted,
	@@ -245,7 +245,7 @@ class MemoryDbParser(BaseParser):
245	def save_strings_id(self):	245	def save_strings_id(self):
246	TorchDb().create_table_with_headers(DbConstant.TABLE_STRING_IDS, TableColumnsManager.TableColumns.get(DbConstant.TABLE_STRING_IDS))	246	TorchDb().create_table_with_headers(DbConstant.TABLE_STRING_IDS, TableColumnsManager.TableColumns.get(DbConstant.TABLE_STRING_IDS))
247	TorchDb().insert_data_into_table(DbConstant.TABLE_STRING_IDS, Str2IdManager().get_all_string_2_id_data())	247	TorchDb().insert_data_into_table(DbConstant.TABLE_STRING_IDS, Str2IdManager().get_all_string_2_id_data())
248	-	248	+
249	def save_memory_data_to_db(self):	249	def save_memory_data_to_db(self):
250	self.get_ge_memory_data()	250	self.get_ge_memory_data()
251	self.save_memory_record_data_to_db()	251	self.save_memory_record_data_to_db()

	@@ -21,11 +21,11 @@ def _from_dlpack(ext_tensor) -> 'torch.Tensor':
21	def _apply_dlpack_patch():	21	def _apply_dlpack_patch():
22	"""Patch torch.utils.dlpack and torch.utils to use torch_npu implementation for NPU tensors"""	22	"""Patch torch.utils.dlpack and torch.utils to use torch_npu implementation for NPU tensors"""
23	import torch.utils.dlpack as torch_dlpack	23	import torch.utils.dlpack as torch_dlpack
24	-	24	+
25	# Store original functions	25	# Store original functions
26	_original_to_dlpack = torch_dlpack.to_dlpack	26	_original_to_dlpack = torch_dlpack.to_dlpack
27	_original_from_dlpack = torch_dlpack.from_dlpack	27	_original_from_dlpack = torch_dlpack.from_dlpack
28	-	28	+
29	def create_patched_to_dlpack(module_name):	29	def create_patched_to_dlpack(module_name):
30	"""Create a patched to_dlpack function with proper __module__ attribute"""	30	"""Create a patched to_dlpack function with proper __module__ attribute"""
31	def patched_to_dlpack(tensor):	31	def patched_to_dlpack(tensor):
	@@ -35,7 +35,7 @@ def _apply_dlpack_patch():
35	return _original_to_dlpack(tensor)	35	return _original_to_dlpack(tensor)
36	patched_to_dlpack.__module__ = module_name	36	patched_to_dlpack.__module__ = module_name
37	return patched_to_dlpack	37	return patched_to_dlpack
38	-	38	+
39	def create_patched_from_dlpack(module_name):	39	def create_patched_from_dlpack(module_name):
40	"""Create a patched from_dlpack function with proper __module__ attribute"""	40	"""Create a patched from_dlpack function with proper __module__ attribute"""
41	def patched_from_dlpack(ext_tensor):	41	def patched_from_dlpack(ext_tensor):
	@@ -48,35 +48,35 @@ def _apply_dlpack_patch():
48	return _original_from_dlpack(ext_tensor)	48	return _original_from_dlpack(ext_tensor)
49	patched_from_dlpack.__module__ = module_name	49	patched_from_dlpack.__module__ = module_name
50	return patched_from_dlpack	50	return patched_from_dlpack
51	-	51	+
52	# Apply patches to torch.utils.dlpack	52	# Apply patches to torch.utils.dlpack
53	torch_dlpack.to_dlpack = create_patched_to_dlpack('torch.utils.dlpack')	53	torch_dlpack.to_dlpack = create_patched_to_dlpack('torch.utils.dlpack')
54	torch_dlpack.from_dlpack = create_patched_from_dlpack('torch.utils.dlpack')	54	torch_dlpack.from_dlpack = create_patched_from_dlpack('torch.utils.dlpack')
55	-	55	+
56	# Also patch torch.utils.to_dlpack and torch.utils.from_dlpack if they exist	56	# Also patch torch.utils.to_dlpack and torch.utils.from_dlpack if they exist
57	if hasattr(torch.utils, 'to_dlpack'):	57	if hasattr(torch.utils, 'to_dlpack'):
58	_original_torch_utils_to_dlpack = torch.utils.to_dlpack	58	_original_torch_utils_to_dlpack = torch.utils.to_dlpack
59	torch.utils.to_dlpack = create_patched_to_dlpack('torch.utils')	59	torch.utils.to_dlpack = create_patched_to_dlpack('torch.utils')
60	-	60	+
61	if hasattr(torch.utils, 'from_dlpack'):	61	if hasattr(torch.utils, 'from_dlpack'):
62	_original_torch_utils_from_dlpack = torch.utils.from_dlpack	62	_original_torch_utils_from_dlpack = torch.utils.from_dlpack
63	torch.utils.from_dlpack = create_patched_from_dlpack('torch.utils')	63	torch.utils.from_dlpack = create_patched_from_dlpack('torch.utils')
64	-	64	+
65	# Also patch torch.from_dlpack and torch.to_dlpack if they exist	65	# Also patch torch.from_dlpack and torch.to_dlpack if they exist
66	if hasattr(torch, 'from_dlpack'):	66	if hasattr(torch, 'from_dlpack'):
67	_original_torch_from_dlpack = torch.from_dlpack	67	_original_torch_from_dlpack = torch.from_dlpack
68	torch.from_dlpack = create_patched_from_dlpack('torch')	68	torch.from_dlpack = create_patched_from_dlpack('torch')
69	-	69	+
70	if hasattr(torch, 'to_dlpack'):	70	if hasattr(torch, 'to_dlpack'):
71	_original_torch_to_dlpack = torch.to_dlpack	71	_original_torch_to_dlpack = torch.to_dlpack
72	torch.to_dlpack = create_patched_to_dlpack('torch')	72	torch.to_dlpack = create_patched_to_dlpack('torch')
73	-	73	+
74	# Add to_dlpack to torch.__all__ if it exists, otherwise create it	74	# Add to_dlpack to torch.__all__ if it exists, otherwise create it
75	if not hasattr(torch, '__all__'):	75	if not hasattr(torch, '__all__'):
76	torch.__all__ = []	76	torch.__all__ = []
77	if 'to_dlpack' not in torch.__all__:	77	if 'to_dlpack' not in torch.__all__:
78	torch.__all__.append('to_dlpack')	78	torch.__all__.append('to_dlpack')
79	-	79	+
80	# Also ensure from_dlpack is in torch.__all__ if it exists	80	# Also ensure from_dlpack is in torch.__all__ if it exists
81	if hasattr(torch, 'from_dlpack'):	81	if hasattr(torch, 'from_dlpack'):
82	if not hasattr(torch, '__all__'):	82	if not hasattr(torch, '__all__'):